Minor improvements to SYCL utils and related clean-up

[alexxy/gromacs.git] / src / gromacs / gpu_utils / sycl_kernel_utils.h

diff --git a/src/gromacs/gpu_utils/sycl_kernel_utils.h b/src/gromacs/gpu_utils/sycl_kernel_utils.h
index 9a3c041ef8867aeb58c3198b05df0e3c8610e74e..544c9e2b13236021dda8eeef37919d77eb25bf2c 100644 (file)
--- a/src/gromacs/gpu_utils/sycl_kernel_utils.h
+++ b/src/gromacs/gpu_utils/sycl_kernel_utils.h
@@ -73,7 +73,8 @@ static inline T atomicLoad(T& val)
  * Equivalent with CUDA's \c syncwarp(c_cudaFullWarpMask).
  *
  */
-static inline void subGroupBarrier(const cl::sycl::nd_item<1> itemIdx)
+template<int Dim>
+static inline void subGroupBarrier(const cl::sycl::nd_item<Dim> itemIdx)
 {
 #if GMX_SYCL_HIPSYCL
     cl::sycl::group_barrier(itemIdx.get_sub_group(), cl::sycl::memory_scope::sub_group);