diff --git a/sycl/include/sycl/nd_item.hpp b/sycl/include/sycl/nd_item.hpp index 708fb3c7b1c2f..62eb7806e4517 100644 --- a/sycl/include/sycl/nd_item.hpp +++ b/sycl/include/sycl/nd_item.hpp @@ -196,6 +196,7 @@ template 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 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 diff --git a/sycl/include/sycl/reduction.hpp b/sycl/include/sycl/reduction.hpp index d4a6e19353cdb..064fbd5b512b0 100644 --- a/sycl/include/sycl/reduction.hpp +++ b/sycl/include/sycl/reduction.hpp @@ -27,6 +27,7 @@ #include // for queue_impl #include // for workGroup... #include // for reduce_ov... +#include #include // for handler #include // for getDeline... #include // 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 { // 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 { // 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 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 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()); } } } diff --git a/sycl/test-e2e/ESIMD/accessor_global.cpp b/sycl/test-e2e/ESIMD/accessor_global.cpp index 77a294934066e..d1888da618bac 100644 --- a/sycl/test-e2e/ESIMD/accessor_global.cpp +++ b/sycl/test-e2e/ESIMD/accessor_global.cpp @@ -6,6 +6,7 @@ #include #include "esimd_test_utils.hpp" +#include 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++) { diff --git a/sycl/test-e2e/ESIMD/accessor_local.cpp b/sycl/test-e2e/ESIMD/accessor_local.cpp index 5d01748c0b644..5c0cfc8596e24 100644 --- a/sycl/test-e2e/ESIMD/accessor_local.cpp +++ b/sycl/test-e2e/ESIMD/accessor_local.cpp @@ -6,6 +6,7 @@ #include #include "esimd_test_utils.hpp" +#include 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++) { diff --git a/sycl/test-e2e/ESIMD/local_accessor_block_load_store.cpp b/sycl/test-e2e/ESIMD/local_accessor_block_load_store.cpp index 7aa73e7c3a3ff..1b418b0928ce2 100644 --- a/sycl/test-e2e/ESIMD/local_accessor_block_load_store.cpp +++ b/sycl/test-e2e/ESIMD/local_accessor_block_load_store.cpp @@ -12,6 +12,7 @@ #include #include "esimd_test_utils.hpp" +#include using namespace sycl; using namespace sycl::ext::intel::esimd; @@ -53,7 +54,7 @@ template 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++) { diff --git a/sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp b/sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp index 90ff147db43eb..c88f5d7017b98 100644 --- a/sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp +++ b/sycl/test-e2e/ESIMD/local_accessor_copy_to_from.cpp @@ -14,6 +14,7 @@ #include #include "esimd_test_utils.hpp" +#include using namespace sycl; using namespace sycl::ext::intel::esimd; @@ -55,7 +56,7 @@ template bool test(queue q) { simd 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++) { diff --git a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp index 858abbd76c8d5..ac40646bece39 100644 --- a/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp +++ b/sycl/test-e2e/ESIMD/lsc/lsc_local_accessor_block_load_store.cpp @@ -14,6 +14,7 @@ #include #include "../esimd_test_utils.hpp" +#include using namespace sycl; using namespace sycl::ext::intel::esimd; @@ -52,7 +53,7 @@ bool test(queue Q, uint32_t LocalRange, uint32_t GlobalRange) { simd 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++) { diff --git a/sycl/test-e2e/ESIMD/slm_block_load_store.cpp b/sycl/test-e2e/ESIMD/slm_block_load_store.cpp index 5aecae59a8ae0..4433602cefb94 100644 --- a/sycl/test-e2e/ESIMD/slm_block_load_store.cpp +++ b/sycl/test-e2e/ESIMD/slm_block_load_store.cpp @@ -17,6 +17,7 @@ #include #include "esimd_test_utils.hpp" +#include using namespace sycl; using namespace sycl::ext::intel::esimd; @@ -55,7 +56,7 @@ template bool test(queue Q) { simd 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++) { diff --git a/sycl/test-e2e/Graph/Update/work_group_static_memory_with_dyn_cgf_and_dyn_params.cpp b/sycl/test-e2e/Graph/Update/work_group_static_memory_with_dyn_cgf_and_dyn_params.cpp index 545d56028307f..49a220f5ae9b3 100644 --- a/sycl/test-e2e/Graph/Update/work_group_static_memory_with_dyn_cgf_and_dyn_params.cpp +++ b/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 +#include 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; diff --git a/sycl/test-e2e/GroupAlgorithm/barrier.cpp b/sycl/test-e2e/GroupAlgorithm/barrier.cpp index 4c9bbe06fc099..d6c85823fe430 100644 --- a/sycl/test-e2e/GroupAlgorithm/barrier.cpp +++ b/sycl/test-e2e/GroupAlgorithm/barrier.cpp @@ -6,6 +6,7 @@ #include #include #include +#include #include 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(), @@ -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);