[SYCL][CUDA] Revert #6386 and remove c++17 usage. (#6400)

Reverts #6386 so that experimental/builtins.hpp is included in sycl.hpp. This avoids users being required to include this header when using printf or experimental math functions.

Corresponding update here: intel/llvm-test-suite#1076.

Author: JackAKirk

sycl/include/CL/sycl.hpp

@@ -61,6 +61,7 @@
 #include <sycl/ext/oneapi/backend/level_zero.hpp>
 #endif
 #include <sycl/ext/oneapi/device_global/properties.hpp>
+#include <sycl/ext/oneapi/experimental/builtins.hpp>
 #include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
 #include <sycl/ext/oneapi/filter_selector.hpp>
 #include <sycl/ext/oneapi/group_algorithm.hpp>

sycl/include/sycl/ext/oneapi/experimental/builtins.hpp

@@ -27,7 +27,10 @@
 #endif
 
 __SYCL_INLINE_NAMESPACE(cl) {
-namespace sycl::ext::oneapi::experimental {
+namespace sycl {
+namespace ext {
+namespace oneapi {
+namespace experimental {
 namespace detail {
 template <size_t N>
 uint32_t to_uint32_t(sycl::marray<bfloat16, N> x, size_t start) {
@@ -144,7 +147,7 @@ sycl::marray<bfloat16, N> fabs(sycl::marray<bfloat16, N> x) {
     std::memcpy(&res[i * 2], &partial_res, sizeof(uint32_t));
   }
 
-  if constexpr (N % 2) {
+  if (N % 2) {
     res[N - 1] = bfloat16::from_bits(__clc_fabs(x[N - 1].raw()));
   }
   return res;
@@ -179,7 +182,7 @@ sycl::marray<bfloat16, N> fmin(sycl::marray<bfloat16, N> x,
     std::memcpy(&res[i * 2], &partial_res, sizeof(uint32_t));
   }
 
-  if constexpr (N % 2) {
+  if (N % 2) {
     res[N - 1] =
         bfloat16::from_bits(__clc_fmin(x[N - 1].raw(), y[N - 1].raw()));
   }
@@ -217,7 +220,7 @@ sycl::marray<bfloat16, N> fmax(sycl::marray<bfloat16, N> x,
     std::memcpy(&res[i * 2], &partial_res, sizeof(uint32_t));
   }
 
-  if constexpr (N % 2) {
+  if (N % 2) {
     res[N - 1] =
         bfloat16::from_bits(__clc_fmax(x[N - 1].raw(), y[N - 1].raw()));
   }
@@ -257,7 +260,7 @@ sycl::marray<bfloat16, N> fma(sycl::marray<bfloat16, N> x,
     std::memcpy(&res[i * 2], &partial_res, sizeof(uint32_t));
   }
 
-  if constexpr (N % 2) {
+  if (N % 2) {
     res[N - 1] = bfloat16::from_bits(
         __clc_fma(x[N - 1].raw(), y[N - 1].raw(), z[N - 1].raw()));
   }
@@ -271,7 +274,10 @@ sycl::marray<bfloat16, N> fma(sycl::marray<bfloat16, N> x,
 #endif // defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
 }
 
-} // namespace sycl::ext::oneapi::experimental
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
+} // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
 
 #undef __SYCL_CONSTANT_AS

sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcore.hpp

@@ -10,8 +10,11 @@
 #include <sycl/ext/oneapi/experimental/bfloat16.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
-namespace sycl::ext::oneapi {
-namespace experimental::matrix {
+namespace sycl {
+namespace ext {
+namespace oneapi {
+namespace experimental {
+namespace matrix {
 
 enum class matrix_use { a, b, accumulator };
 
@@ -169,7 +172,8 @@ joint_matrix_fill(Group sg,
 #endif // defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
 }
 
-} // namespace experimental::matrix
+} // namespace matrix
+} // namespace experimental
 
 namespace detail {
 
@@ -199,6 +203,7 @@ constexpr int get_layout_id<
   return 1;
 }
 
+#if __cplusplus >= 201703L // if constexpr usage
 template <typename S, typename T,
           sycl::ext::oneapi::experimental::matrix::matrix_use Use,
           size_t NumRows, size_t NumCols,
@@ -379,6 +384,7 @@ struct joint_matrix_load_impl<
     }
   }
 };
+#endif // __cplusplus >= 201703L
 
 template <typename T, size_t NumRows, size_t NumCols,
           sycl::ext::oneapi::experimental::matrix::matrix_layout Layout,
@@ -391,6 +397,7 @@ struct joint_matrix_store_impl {
         multi_ptr<T, Space> dst, size_t stride);
 };
 
+#if __cplusplus >= 201703L // if constexpr usage
 template <typename T, size_t NumRows, size_t NumCols,
           sycl::ext::oneapi::experimental::matrix::matrix_layout Layout,
           access::address_space Space>
@@ -454,6 +461,7 @@ struct joint_matrix_store_impl<
     }
   }
 };
+#endif // __cplusplus >= 201703L
 
 template <typename T1, typename T2, std::size_t M, std::size_t K, std::size_t N,
           sycl::ext::oneapi::experimental::matrix::matrix_layout LayoutA,
@@ -510,6 +518,7 @@ constexpr int get_layout_pair_id<
   return 3;
 }
 
+#if __cplusplus >= 201703L // if constexpr usage
 template <typename T1, typename T2, std::size_t M, std::size_t K, std::size_t N,
           sycl::ext::oneapi::experimental::matrix::matrix_layout LayoutA,
           sycl::ext::oneapi::experimental::matrix::matrix_layout LayoutB,
@@ -675,10 +684,12 @@ struct joint_matrix_mad_impl<
     return D;
   }
 };
+#endif // __cplusplus >= 201703L
 
 } // namespace detail
 
-namespace experimental::matrix {
+namespace experimental {
+namespace matrix {
 
 template <typename Group, typename S, typename T, matrix_use Use,
           size_t NumRows, size_t NumCols, matrix_layout Layout,
@@ -766,6 +777,9 @@ float round_to_tf32(float a) {
 #endif // defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
 }
 
-} // namespace experimental::matrix
-} // namespace sycl::ext::oneapi
+} // namespace matrix
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
+} // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/test/basic_tests/built-ins.cpp

@@ -7,7 +7,6 @@
 // Hits an assertion with AMD:
 // XFAIL: hip_amd
 
-#include <sycl/ext/oneapi/experimental/builtins.hpp>
 #include <sycl/sycl.hpp>
 
 #include <cassert>

sycl/test/extensions/experimental-printf.cpp

@@ -16,7 +16,6 @@
 // CHECK: Constant [[#TYPE]] [[#CONST:]]
 // CHECK: ExtInst [[#]] [[#]] [[#]] printf [[#]] [[#CONST]]
 
-#include <sycl/ext/oneapi/experimental/builtins.hpp>
 #include <sycl/sycl.hpp>
 
 #ifdef __SYCL_DEVICE_ONLY__