Properly set properties of group_load/group_store to striped

Doing so exactly recovers the behavior of sub_group::load<vec_sz>,
sub_group::store<vec_sz> and eliminates warnings with 2025.1 and SYCLOS.

With this change, enable use of group_load, group_store for 2025.0
compiler and newer

Author: oleksandr-pavlyk

dpctl/tensor/libtensor/include/utils/sycl_utils.hpp

@@ -482,7 +482,7 @@ struct Identity<Op, T, std::enable_if_t<UseBuiltInIdentity<Op, T>::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(__INTEL_LLVM_COMPILER) && (__INTEL_LLVM_COMPILER >= 20250000u)
 #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<ElementType>;
     sycl::vec<ValueT, vec_sz> 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<vec_sz>(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<ElementType>;
     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<VecT, ElementType>);
-    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<vec_sz>(m_ptr, val);
@@ -561,7 +564,8 @@ sub_group_store(const sycl::sub_group &sg,
                 sycl::multi_ptr<ElementType, Space, DecorateAddress> 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);