diff --git a/src/gromacs/mdlib/settle_gpu_internal_sycl.cpp b/src/gromacs/mdlib/settle_gpu_internal_sycl.cpp index 8a0be6d33aa8061ef2cd873e2e225930411f7b35..924e10ab654ad56b9b7c49caa136cf369fa21f5f 100644 --- a/src/gromacs/mdlib/settle_gpu_internal_sycl.cpp +++ b/src/gromacs/mdlib/settle_gpu_internal_sycl.cpp @@ -309,7 +309,7 @@ auto settleKernel(sycl::handler& c if constexpr (computeVirial) { // This is to ensure that all threads saved the data before reduction starts - subGroupBarrier(itemIdx); + itemIdx.barrier(fence_space::local_space); constexpr int blockSize = sc_workGroupSize; const int subGroupSize = itemIdx.get_sub_group().get_max_local_range()[0]; // Reduce up to one virial per thread block @@ -330,6 +330,10 @@ auto settleKernel(sycl::handler& c } } if (dividedAt > subGroupSize / 2) + { + itemIdx.barrier(fence_space::local_space); + } + else { subGroupBarrier(itemIdx); }