Minor improvements to SYCL utils and related clean-up

- ewald/pme_gpu_types_host.h: Remove unused forward declaration.
- gpu_utils/devicebuffer_sycl.h: Allow casting our "blackhole" class to
  nullptr.
- gpu_utils/sycl_kernel_utils.h: Make subGroupBarrier usable with any
  group dimension.

In preparation for #3927.

diff --git a/src/gromacs/ewald/pme_gpu_types_host.h b/src/gromacs/ewald/pme_gpu_types_host.h
index 055e61e5d8c0ec6bcb6a0389208e2cfb427b46f7..0c6d7bb951b9e96bbc4ad66388c8513c3ff5042c 100644 (file)
--- a/src/gromacs/ewald/pme_gpu_types_host.h
+++ b/src/gromacs/ewald/pme_gpu_types_host.h
@@ -62,11 +62,6 @@
 #include "pme_gpu_settings.h"
 #include "pme_gpu_staging.h"
 
-namespace gmx
-{
-class PmeDeviceBuffers;
-} // namespace gmx
-
 #if GMX_GPU
 struct PmeGpuSpecific;
 #else

diff --git a/src/gromacs/gpu_utils/devicebuffer_sycl.h b/src/gromacs/gpu_utils/devicebuffer_sycl.h
index 97a3d53cdf38fd56bb0624afc863b53e5a6178e9..cb3277b82f4c12e60f4dd74d16042ba37aed1bad 100644 (file)
--- a/src/gromacs/gpu_utils/devicebuffer_sycl.h
+++ b/src/gromacs/gpu_utils/devicebuffer_sycl.h
@@ -195,6 +195,8 @@ struct EmptyClassThatIgnoresConstructorArguments
     [[maybe_unused]] EmptyClassThatIgnoresConstructorArguments(Args&&... /*args*/)
     {
     }
+    //! Allow casting to nullptr
+    constexpr operator std::nullptr_t() const { return nullptr; }
 };
 } // namespace gmx::internal
 

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);