From c631fddd323fe28ad27c9f8c1ee407d480921ca8 Mon Sep 17 00:00:00 2001 From: Andrey Alekseenko Date: Thu, 9 Sep 2021 23:43:06 +0300 Subject: [PATCH] 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. --- src/gromacs/ewald/pme_gpu_types_host.h | 5 ----- src/gromacs/gpu_utils/devicebuffer_sycl.h | 2 ++ src/gromacs/gpu_utils/sycl_kernel_utils.h | 3 ++- 3 files changed, 4 insertions(+), 6 deletions(-) diff --git a/src/gromacs/ewald/pme_gpu_types_host.h b/src/gromacs/ewald/pme_gpu_types_host.h index 055e61e5d8..0c6d7bb951 100644 --- 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 97a3d53cdf..cb3277b82f 100644 --- 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 9a3c041ef8..544c9e2b13 100644 --- 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 +static inline void subGroupBarrier(const cl::sycl::nd_item itemIdx) { #if GMX_SYCL_HIPSYCL cl::sycl::group_barrier(itemIdx.get_sub_group(), cl::sycl::memory_scope::sub_group); -- 2.22.0