diff --git a/sycl/include/CL/sycl.hpp b/sycl/include/CL/sycl.hpp index 10dcf00b94d22..9b517ced4b659 100644 --- a/sycl/include/CL/sycl.hpp +++ b/sycl/include/CL/sycl.hpp @@ -61,6 +61,7 @@ #include #endif #include +#include #include #include #include diff --git a/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp b/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp index 9e93e688f181a..57c1b0c58fc2c 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/builtins.hpp +++ b/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 uint32_t to_uint32_t(sycl::marray x, size_t start) { @@ -144,7 +147,7 @@ sycl::marray fabs(sycl::marray 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 fmin(sycl::marray 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 fmax(sycl::marray 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 fma(sycl::marray 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 fma(sycl::marray 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 diff --git a/sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcore.hpp b/sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcore.hpp index 6de66d0f8590c..cf53bec8f943c 100644 --- a/sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcore.hpp +++ b/sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcore.hpp @@ -10,8 +10,11 @@ #include __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 = 201703L template dst, size_t stride); }; +#if __cplusplus >= 201703L // if constexpr usage template @@ -454,6 +461,7 @@ struct joint_matrix_store_impl< } } }; +#endif // __cplusplus >= 201703L template = 201703L // if constexpr usage template = 201703L } // namespace detail -namespace experimental::matrix { +namespace experimental { +namespace matrix { template #include #include diff --git a/sycl/test/extensions/experimental-printf.cpp b/sycl/test/extensions/experimental-printf.cpp index d64e78ee4883e..4c9269e54495a 100644 --- a/sycl/test/extensions/experimental-printf.cpp +++ b/sycl/test/extensions/experimental-printf.cpp @@ -16,7 +16,6 @@ // CHECK: Constant [[#TYPE]] [[#CONST:]] // CHECK: ExtInst [[#]] [[#]] [[#]] printf [[#]] [[#CONST]] -#include #include #ifdef __SYCL_DEVICE_ONLY__