Bump minimal version of __INTEL_LLVM_COMPILER needed to enable group_load/group_store

Author: oleksandr-pavlyk

dpctl/tensor/libtensor/include/kernels/copy_as_contiguous.hpp

@@ -77,12 +77,12 @@ class CopyAsCContigFunctor
         using dpctl::tensor::type_utils::is_complex;
         if constexpr (!enable_sg_loadstore || is_complex<T>::value) {
             const std::uint16_t sgSize =
-                ndit.get_sub_group().get_local_range()[0];
+                ndit.get_sub_group().get_max_local_range()[0];
             const std::size_t gid = ndit.get_global_linear_id();
 
             // start = (gid / sgSize) * sgSize * elems_per_wi + (gid % sgSize)
             // gid % sgSize == gid - (gid / sgSize) * sgSize
-            const std::size_t elems_per_sg = sgSize * elems_per_wi;
+            const std::uint16_t elems_per_sg = sgSize * elems_per_wi;
             const std::size_t start =
                 (gid / sgSize) * (elems_per_sg - sgSize) + gid;
             const std::size_t end = std::min(nelems, start + elems_per_sg);
@@ -98,20 +98,21 @@ class CopyAsCContigFunctor
             const size_t base =
                 elems_per_wi * (ndit.get_group(0) * ndit.get_local_range(0) +
                                 sg.get_group_id()[0] * sgSize);
+            const std::uint16_t elems_per_sg = elems_per_wi * sgSize;
 
-            if (base + elems_per_wi * sgSize < nelems) {
-                sycl::vec<T, vec_sz> dst_vec;
-
+            if (base + elems_per_sg < nelems) {
 #pragma unroll
                 for (std::uint8_t it = 0; it < elems_per_wi; it += vec_sz) {
+                    // it == vec_id * vec_sz, for  0 <= vec_id < n_vecs
                     const size_t block_start_id = base + it * sgSize;
                     auto dst_multi_ptr = sycl::address_space_cast<
                         sycl::access::address_space::global_space,
                         sycl::access::decorated::yes>(&dst_p[block_start_id]);
 
                     const size_t elem_id0 = block_start_id + sg.get_local_id();
+                    sycl::vec<T, vec_sz> dst_vec;
 #pragma unroll
-                    for (std::uint8_t k = 0; k < vec_sz; k++) {
+                    for (std::uint8_t k = 0; k < vec_sz; ++k) {
                         const size_t elem_id = elem_id0 + k * sgSize;
                         const ssize_t src_offset = src_indexer(elem_id);
                         dst_vec[k] = src_p[src_offset];

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

@@ -428,7 +428,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_COMIPLER > 20250000u)
+#if defined(__INTEL_LLVM_COMPILER) && (__INTEL_LLVM_COMPILER >= 20250100u)
 #define USE_GROUP_LOAD_STORE 1
 #else
 #define USE_GROUP_LOAD_STORE 0
@@ -450,7 +450,7 @@ 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::group_load(sg, m_ptr, x, ls_ns::data_placement_blocked);
     return x;
 #else
     return sg.load<vec_sz>(m_ptr);
@@ -466,7 +466,7 @@ 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::group_load(sg, m_ptr, x, ls_ns::data_placement_blocked);
     return x;
 #else
     return sg.load(m_ptr);
@@ -486,7 +486,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);
+    static_assert(std::is_same_v<VecT, ElementType>);
+    ls_ns::group_store(sg, val, m_ptr, ls_ns::data_placement_blocked);
     return;
 #else
     sg.store<vec_sz>(m_ptr, val);
@@ -506,7 +507,7 @@ 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::group_store(sg, val, m_ptr, ls_ns::data_placement_blocked);
     return;
 #else
     sg.store(m_ptr, val);