diff --git a/dpctl/tensor/libtensor/include/utils/sycl_utils.hpp b/dpctl/tensor/libtensor/include/utils/sycl_utils.hpp index 53ecbbbe09..deda9cd810 100644 --- a/dpctl/tensor/libtensor/include/utils/sycl_utils.hpp +++ b/dpctl/tensor/libtensor/include/utils/sycl_utils.hpp @@ -482,7 +482,7 @@ struct Identity::value>> SYCL_EXT_ONEAPI_GROUP_LOAD_STORE #define USE_GROUP_LOAD_STORE 1 #else -#if defined(__INTEL_LLVM_COMPILER) && (__INTEL_LLVM_COMPILER > 20250100u) +#if defined(__LIBSYCL_MAJOR_VERSION) && (__LIBSYCL_MAJOR_VERSION >= 8u) #define USE_GROUP_LOAD_STORE 1 #else #define USE_GROUP_LOAD_STORE 0 @@ -504,7 +504,8 @@ auto sub_group_load(const sycl::sub_group &sg, #if (USE_GROUP_LOAD_STORE) using ValueT = typename std::remove_cv_t; sycl::vec x{}; - ls_ns::group_load(sg, m_ptr, x, ls_ns::data_placement_blocked); + constexpr auto striped = ls_ns::properties{ls_ns::data_placement_striped}; + ls_ns::group_load(sg, m_ptr, x, striped); return x; #else return sg.load(m_ptr); @@ -520,7 +521,8 @@ auto sub_group_load(const sycl::sub_group &sg, #if (USE_GROUP_LOAD_STORE) using ValueT = typename std::remove_cv_t; ValueT x{}; - ls_ns::group_load(sg, m_ptr, x, ls_ns::data_placement_blocked); + constexpr auto striped = ls_ns::properties{ls_ns::data_placement_striped}; + ls_ns::group_load(sg, m_ptr, x, striped); return x; #else return sg.load(m_ptr); @@ -541,7 +543,8 @@ sub_group_store(const sycl::sub_group &sg, { #if (USE_GROUP_LOAD_STORE) static_assert(std::is_same_v); - ls_ns::group_store(sg, val, m_ptr, ls_ns::data_placement_blocked); + constexpr auto striped = ls_ns::properties{ls_ns::data_placement_striped}; + ls_ns::group_store(sg, val, m_ptr, striped); return; #else sg.store(m_ptr, val); @@ -561,7 +564,8 @@ sub_group_store(const sycl::sub_group &sg, sycl::multi_ptr m_ptr) { #if (USE_GROUP_LOAD_STORE) - ls_ns::group_store(sg, val, m_ptr, ls_ns::data_placement_blocked); + constexpr auto striped = ls_ns::properties{ls_ns::data_placement_striped}; + ls_ns::group_store(sg, val, m_ptr, striped); return; #else sg.store(m_ptr, val);