[SYCL] Prepare nd_item barrier & mem_fence removal

sycl/include/sycl/nd_item.hpp

@@ -196,6 +196,7 @@ template <int Dimensions = 1> class nd_item {
                                 get_offset());
   }
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   __SYCL2020_DEPRECATED("use sycl::group_barrier() free function instead")
   void barrier([[maybe_unused]] access::fence_space accessSpace =
                    access::fence_space::global_and_local) const {
@@ -230,6 +231,7 @@ template <int Dimensions = 1> class nd_item {
     __spirv_MemoryBarrier(__spv::Scope::Workgroup, flags);
 #endif
   }
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES
 
   /// Asynchronously copies a number of elements specified by \p numElements
   /// from the source pointed by \p src to destination pointed by \p dest

sycl/include/sycl/reduction.hpp

@@ -27,6 +27,7 @@
 #include <sycl/exception_list.hpp>                  // for queue_impl
 #include <sycl/group.hpp>                           // for workGroup...
 #include <sycl/group_algorithm.hpp>                 // for reduce_ov...
+#include <sycl/group_barrier.hpp>
 #include <sycl/handler.hpp>                         // for handler
 #include <sycl/id.hpp>                              // for getDeline...
 #include <sycl/kernel.hpp>                          // for auto_name
@@ -1644,7 +1645,7 @@ struct NDRangeReduction<
 
           // Ensure item 0 is finished with LocalReds before next iteration
           if (E != NElements - 1) {
-            NDIt.barrier();
+            group_barrier(NDIt.get_group());
           }
         }
 
@@ -1884,7 +1885,7 @@ template <> struct NDRangeReduction<reduction::strategy::basic> {
 
           // Ensure item 0 is finished with LocalReds before next iteration
           if (E != NElements - 1) {
-            NDIt.barrier();
+            group_barrier(NDIt.get_group());
           }
         }
       });
@@ -1988,7 +1989,7 @@ template <> struct NDRangeReduction<reduction::strategy::basic> {
 
               // Ensure item 0 is finished with LocalReds before next iteration
               if (E != NElements - 1) {
-                NDIt.barrier();
+                group_barrier(NDIt.get_group());
               }
             }
           });
@@ -2240,7 +2241,7 @@ void reduCGFuncImplArrayHelper(nd_item<Dims> NDIt, LocalAccT LocalReds,
 
     // Ensure item 0 is finished with LocalReds before next iteration
     if (E != NElements - 1) {
-      NDIt.barrier();
+      group_barrier(NDIt.get_group());
     }
   }
 }
@@ -2437,7 +2438,7 @@ void reduAuxCGFuncImplArrayHelper(nd_item<Dims> NDIt, size_t LID, size_t GID,
 
     // Ensure item 0 is finished with LocalReds before next iteration
     if (E != NElements - 1) {
-      NDIt.barrier();
+      group_barrier(NDIt.get_group());
     }
   }
 }

sycl/test-e2e/ESIMD/accessor_global.cpp

@@ -6,6 +6,7 @@
 #include <iostream>
 
 #include "esimd_test_utils.hpp"
+#include <sycl/group_barrier.hpp>
 
 using namespace sycl;
 using namespace sycl::ext::intel::esimd;
@@ -42,7 +43,7 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
            block_store(Ptr + GID * VL, Values);
          }
 
-         Item.barrier();
+         group_barrier(Item.get_group());
 
          if (LID == 0) {
            for (int LID = 0; LID < LocalRange; LID++) {

sycl/test-e2e/ESIMD/accessor_local.cpp

@@ -6,6 +6,7 @@
 #include <iostream>
 
 #include "esimd_test_utils.hpp"
+#include <sycl/group_barrier.hpp>
 
 using namespace sycl;
 using namespace sycl::ext::intel::esimd;
@@ -52,7 +53,7 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
            slm_block_store(LocalAccOffset + LID * VL * sizeof(T), ValuesToSLM);
          }
 
-         Item.barrier();
+         group_barrier(Item.get_group());
 
          if (LID == 0) {
            for (int LID = 0; LID < LocalRange; LID++) {

sycl/test-e2e/ESIMD/local_accessor_block_load_store.cpp

@@ -12,6 +12,7 @@
 #include <iostream>
 
 #include "esimd_test_utils.hpp"
+#include <sycl/group_barrier.hpp>
 
 using namespace sycl;
 using namespace sycl::ext::intel::esimd;
@@ -53,7 +54,7 @@ template <typename T, int VL, int Align = 16> bool test(queue Q) {
          block_store(LocalAcc, Align + LID * VL * sizeof(T), ValuesToSLM,
                      AlignTag);
 
-         Item.barrier();
+         group_barrier(Item.get_group());
 
          if (LID == 0) {
            for (int LID = 0; LID < LocalRange; LID++) {

sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp

@@ -14,6 +14,7 @@
 #include <iostream>
 
 #include "esimd_test_utils.hpp"
+#include <sycl/group_barrier.hpp>
 
 using namespace sycl;
 using namespace sycl::ext::intel::esimd;
@@ -55,7 +56,7 @@ template <typename T, unsigned VL> bool test(queue q) {
          simd<T, VL> ValuesToSLM(GID * 100, 1);
          ValuesToSLM.copy_to(LocalAcc, LID * VL * sizeof(T));
 
-         Item.barrier();
+         group_barrier(Item.get_group());
 
          if (LID == 0) {
            for (int LID = 0; LID < LocalRange; LID++) {

sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp

@@ -14,6 +14,7 @@
 #include <iostream>
 
 #include "../esimd_test_utils.hpp"
+#include <sycl/group_barrier.hpp>
 
 using namespace sycl;
 using namespace sycl::ext::intel::esimd;
@@ -52,7 +53,7 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) {
          simd<T, VL> ValuesToSLM = IntValues;
          lsc_block_store(LocalAcc, LID * VL * sizeof(T), ValuesToSLM);
 
-         Item.barrier();
+         group_barrier(Item.get_group());
 
          if (LID == 0) {
            for (int LID = 0; LID < LocalRange; LID++) {

sycl/test-e2e/ESIMD/slm_block_load_store.cpp

@@ -17,6 +17,7 @@
 #include <iostream>
 
 #include "esimd_test_utils.hpp"
+#include <sycl/group_barrier.hpp>
 
 using namespace sycl;
 using namespace sycl::ext::intel::esimd;
@@ -55,7 +56,7 @@ template <typename T, int VL, int Align = 16> bool test(queue Q) {
        simd<T, VL> ValuesToSLM = IntValues;
        slm_block_store(Align + LID * VL * sizeof(T), ValuesToSLM, AlignTag);
 
-       Item.barrier();
+       group_barrier(Item.get_group());
 
        if (LID == 0) {
          for (int LID = 0; LID < LocalRange; LID++) {

sycl/test-e2e/Graph/Update/work_group_static_memory_with_dyn_cgf_and_dyn_params.cpp

@@ -12,6 +12,7 @@
 
 #include "../graph_common.hpp"
 #include <sycl/ext/oneapi/work_group_static.hpp>
+#include <sycl/group_barrier.hpp>
 
 constexpr size_t WgSize = 32;
 
@@ -35,7 +36,7 @@ int main() {
     CGH.parallel_for(nd_range({Size}, {WgSize}), [=](nd_item<1> Item) {
       LocalIDBuff[Item.get_local_linear_id()] = Item.get_local_linear_id();
 
-      Item.barrier();
+      group_barrier(Item.get_group());
 
       // Check that the memory is accessible from other work-items
       size_t LocalIdx = Item.get_local_linear_id() ^ 1;
@@ -49,7 +50,7 @@ int main() {
     CGH.parallel_for(nd_range({Size}, {WgSize}), [=](nd_item<1> Item) {
       LocalIDBuff[Item.get_local_linear_id()] = Item.get_local_linear_id();
 
-      Item.barrier();
+      group_barrier(Item.get_group());
 
       // Check that the memory is accessible from other work-items
       size_t LocalIdx = Item.get_local_linear_id() ^ 1;

sycl/test-e2e/GroupAlgorithm/barrier.cpp

@@ -6,6 +6,7 @@
 #include <iostream>
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
+#include <sycl/group_barrier.hpp>
 #include <vector>
 
 using namespace sycl;
@@ -31,7 +32,7 @@ void basic() {
         if (idx < 2) {
           loc_barrier[idx].initialize(N);
         }
-        item.barrier(access::fence_space::local_space);
+        group_barrier(item.get_group());
         for (int i = 0; i < N; i++) {
           int val = loc[idx];
           barrier::arrival_token arr = loc_barrier[0].arrive();
@@ -78,7 +79,7 @@ void interface() {
         if (idx == 1) {
           loc_barrier[1].initialize(N * N);
         }
-        item.barrier(access::fence_space::local_space);
+        group_barrier(item.get_group());
 
         item.async_work_group_copy(
             loc.get_multi_ptr<access::decorated::yes>(),
@@ -93,12 +94,12 @@ void interface() {
           *reused_barrier_space = loc[0];
           loc[0] = 0;
         }
-        item.barrier(access::fence_space::local_space);
+        group_barrier(item.get_group());
         if (idx == 1) {
           int *reused_barrier_space = (int *)(void *)loc_barrier.get_pointer();
           loc[0] = *reused_barrier_space;
         }
-        item.barrier(access::fence_space::local_space);
+        group_barrier(item.get_group());
         if (idx == 0) {
           loc_barrier[0].initialize(N);
         }
@@ -119,7 +120,7 @@ void interface() {
         arr = loc_barrier[1].arrive();
         test1_acc[idx] = loc_barrier[1].test_wait(arr);
         arr = loc_barrier[1].arrive();
-        item.barrier(access::fence_space::local_space);
+        group_barrier(item.get_group());
         test2_acc[idx] = loc_barrier[1].test_wait(arr);
         loc_barrier[1].wait(arr);