Make LeapFrogGpu less tied to CUDA

- Extract common parts of LeapFrogGpu from .cuh into .h

- Make it use DeviceBuffers instead of raw pointers. For CUDA build,
they are the same, so no changes to other parts of the code necessary.

diff --git a/src/gromacs/mdlib/leapfrog_gpu.cu b/src/gromacs/mdlib/leapfrog_gpu.cu
index f34c51db4599da69a0bdedc1c2730acf73d4843c..b4c08e0e350850531aa6c3adf358368eac849e0d 100644 (file)
--- a/src/gromacs/mdlib/leapfrog_gpu.cu
+++ b/src/gromacs/mdlib/leapfrog_gpu.cu
@@ -46,7 +46,7 @@
  */
 #include "gmxpre.h"
 
-#include "leapfrog_gpu.cuh"
+#include "leapfrog_gpu.h"
 
 #include <assert.h>
 #include <stdio.h>
@@ -75,30 +75,6 @@ constexpr static int c_threadsPerBlock = 256;
 //! Maximum number of threads in a block (for __launch_bounds__)
 constexpr static int c_maxThreadsPerBlock = c_threadsPerBlock;
 
-/*! \brief Sets the number of different temperature coupling values
- *
- *  This is needed to template the kernel
- *  \todo Unify with similar enum in CPU update module
- */
-enum class NumTempScaleValues
-{
-    None,    //!< No temperature coupling
-    Single,  //!< Single T-scaling value (one group)
-    Multiple //!< Multiple T-scaling values, need to use T-group indices
-};
-
-/*! \brief Different variants of the Parrinello-Rahman velocity scaling
- *
- *  This is needed to template the kernel
- *  \todo Unify with similar enum in CPU update module
- */
-enum class VelocityScalingType
-{
-    None,     //!< Do not apply velocity scaling (not a PR-coupling run or step)
-    Diagonal, //!< Apply velocity scaling using a diagonal matrix
-    Full      //!< Apply velocity scaling using a full matrix
-};
-
 /*! \brief Main kernel for Leap-Frog integrator.
  *
  *  The coordinates and velocities are updated on the GPU. Also saves the intermediate values of the coordinates for
@@ -261,10 +237,10 @@ inline auto selectLeapFrogKernelPtr(bool                doTemperatureScaling,
     return kernelPtr;
 }
 
-void LeapFrogGpu::integrate(const float3*                     d_x,
-                            float3*                           d_xp,
-                            float3*                           d_v,
-                            const float3*                     d_f,
+void LeapFrogGpu::integrate(const DeviceBuffer<float3>        d_x,
+                            DeviceBuffer<float3>              d_xp,
+                            DeviceBuffer<float3>              d_v,
+                            const DeviceBuffer<float3>        d_f,
                             const real                        dt,
                             const bool                        doTemperatureScaling,
                             gmx::ArrayRef<const t_grp_tcstat> tcstat,

diff --git a/src/gromacs/mdlib/leapfrog_gpu.cuh b/src/gromacs/mdlib/leapfrog_gpu.h
similarity index 82%
rename from src/gromacs/mdlib/leapfrog_gpu.cuh
rename to src/gromacs/mdlib/leapfrog_gpu.h
index cec20f751d0aa52f4a73c0bbac94282b38797cb3..7554e6e793dacd482f391ac23728aeabe9e78e99 100644 (file)
--- a/src/gromacs/mdlib/leapfrog_gpu.cuh
+++ b/src/gromacs/mdlib/leapfrog_gpu.h
@@ -41,10 +41,16 @@
  * \ingroup module_mdlib
  * \inlibraryapi
  */
-#ifndef GMX_MDLIB_LEAPFROG_GPU_CUH
-#define GMX_MDLIB_LEAPFROG_GPU_CUH
+#ifndef GMX_MDLIB_LEAPFROG_GPU_H
+#define GMX_MDLIB_LEAPFROG_GPU_H
+
+#include "config.h"
+
+#if GMX_GPU_CUDA
+#    include "gromacs/gpu_utils/devicebuffer.cuh"
+#    include "gromacs/gpu_utils/gputraits.cuh"
+#endif
 
-#include "gromacs/gpu_utils/gputraits.cuh"
 #include "gromacs/gpu_utils/hostallocator.h"
 #include "gromacs/pbcutil/pbc.h"
 #include "gromacs/pbcutil/pbc_aiuc.h"
@@ -58,6 +64,31 @@ struct t_grp_tcstat;
 namespace gmx
 {
 
+
+/*! \brief Sets the number of different temperature coupling values
+ *
+ *  This is needed to template the kernel
+ *  \todo Unify with similar enum in CPU update module
+ */
+enum class NumTempScaleValues
+{
+    None,    //!< No temperature coupling
+    Single,  //!< Single T-scaling value (one group)
+    Multiple //!< Multiple T-scaling values, need to use T-group indices
+};
+
+/*! \brief Different variants of the Parrinello-Rahman velocity scaling
+ *
+ *  This is needed to template the kernel
+ *  \todo Unify with similar enum in CPU update module
+ */
+enum class VelocityScalingType
+{
+    None,     //!< Do not apply velocity scaling (not a PR-coupling run or step)
+    Diagonal, //!< Apply velocity scaling using a diagonal matrix
+    Full      //!< Apply velocity scaling using a full matrix
+};
+
 class LeapFrogGpu
 {
 
@@ -86,10 +117,10 @@ public:
      * \param[in]     dtPressureCouple         Period between pressure coupling steps
      * \param[in]     prVelocityScalingMatrix  Parrinello-Rahman velocity scaling matrix
      */
-    void integrate(const float3*                     d_x,
-                   float3*                           d_xp,
-                   float3*                           d_v,
-                   const float3*                     d_f,
+    void integrate(const DeviceBuffer<float3>        d_x,
+                   DeviceBuffer<float3>              d_xp,
+                   DeviceBuffer<float3>              d_v,
+                   const DeviceBuffer<float3>        d_f,
                    const real                        dt,
                    const bool                        doTemperatureScaling,
                    gmx::ArrayRef<const t_grp_tcstat> tcstat,
@@ -127,7 +158,7 @@ private:
     int numAtoms_;
 
     //! 1/mass for all atoms (GPU)
-    real* d_inverseMasses_;
+    DeviceBuffer<float> d_inverseMasses_;
     //! Current size of the reciprocal masses array
     int numInverseMasses_ = -1;
     //! Maximum size of the reciprocal masses array
@@ -141,7 +172,7 @@ private:
      */
     gmx::HostVector<float> h_lambdas_;
     //! Device-side temperature scaling factors
-    float* d_lambdas_;
+    DeviceBuffer<float> d_lambdas_;
     //! Current size of the array with temperature scaling factors (lambdas)
     int numLambdas_ = -1;
     //! Maximum size of the array with temperature scaling factors (lambdas)
@@ -149,7 +180,7 @@ private:
 
 
     //! Array that maps atom index onto the temperature scaling group to get scaling parameter
-    unsigned short* d_tempScaleGroups_;
+    DeviceBuffer<unsigned short> d_tempScaleGroups_;
     //! Current size of the temperature coupling groups array
     int numTempScaleGroups_ = -1;
     //! Maximum size of the temperature coupling groups array

diff --git a/src/gromacs/mdlib/tests/CMakeLists.txt b/src/gromacs/mdlib/tests/CMakeLists.txt
index c1f6592e1422a2b37b222d5c111ab935b5f8acf7..faab9297da7c13969657a20a519c9f57dc426c61 100644 (file)
--- a/src/gromacs/mdlib/tests/CMakeLists.txt
+++ b/src/gromacs/mdlib/tests/CMakeLists.txt
@@ -44,7 +44,6 @@ gmx_add_unit_test(MdlibUnitTest mdlib-test HARDWARE_DETECTION
         leapfrog.cpp
         leapfrogtestdata.cpp
         leapfrogtestrunners.cpp
-        leapfrogtestrunners_gpu.cpp
         settle.cpp
         settletestdata.cpp
         settletestrunners.cpp
@@ -52,6 +51,8 @@ gmx_add_unit_test(MdlibUnitTest mdlib-test HARDWARE_DETECTION
         simulationsignal.cpp
         updategroups.cpp
         updategroupscog.cpp
+    GPU_CPP_SOURCE_FILES
+        leapfrogtestrunners_gpu.cpp
     CUDA_CU_SOURCE_FILES
         constrtestrunners.cu
         settletestrunners.cu

diff --git a/src/gromacs/mdlib/tests/leapfrogtestrunners_gpu.cpp b/src/gromacs/mdlib/tests/leapfrogtestrunners_gpu.cpp
index 9046085434f70d543e53406d991ef65e46aea0d6..5783b663327dd9f0682bee188177ddf1c051729f 100644 (file)
--- a/src/gromacs/mdlib/tests/leapfrogtestrunners_gpu.cpp
+++ b/src/gromacs/mdlib/tests/leapfrogtestrunners_gpu.cpp
@@ -50,7 +50,7 @@
 
 #if GMX_GPU_CUDA
 #    include "gromacs/gpu_utils/devicebuffer.cuh"
-#    include "gromacs/mdlib/leapfrog_gpu.cuh"
+#    include "gromacs/mdlib/leapfrog_gpu.h"
 #endif
 
 #include "gromacs/hardware/device_information.h"

diff --git a/src/gromacs/mdlib/update_constrain_gpu_impl.cu b/src/gromacs/mdlib/update_constrain_gpu_impl.cu
index ac92d8f36977461c41e0092f0336b3e57757b558..8d7a4bba153eed5da15f960f9a61871a283f7b38 100644 (file)
--- a/src/gromacs/mdlib/update_constrain_gpu_impl.cu
+++ b/src/gromacs/mdlib/update_constrain_gpu_impl.cu
@@ -62,7 +62,7 @@
 #include "gromacs/gpu_utils/devicebuffer.h"
 #include "gromacs/gpu_utils/gputraits.cuh"
 #include "gromacs/gpu_utils/vectype_ops.cuh"
-#include "gromacs/mdlib/leapfrog_gpu.cuh"
+#include "gromacs/mdlib/leapfrog_gpu.h"
 #include "gromacs/mdlib/lincs_gpu.cuh"
 #include "gromacs/mdlib/settle_gpu.cuh"
 #include "gromacs/mdlib/update_constrain_gpu.h"

diff --git a/src/gromacs/mdlib/update_constrain_gpu_impl.h b/src/gromacs/mdlib/update_constrain_gpu_impl.h
index a191cc85b7fe9aa14ef7141f1d4e25b7a65599a7..9ee067791318d5d70ee2c34c3029711ad06220ad 100644 (file)
--- a/src/gromacs/mdlib/update_constrain_gpu_impl.h
+++ b/src/gromacs/mdlib/update_constrain_gpu_impl.h
@@ -49,7 +49,7 @@
 #include "gmxpre.h"
 
 #include "gromacs/gpu_utils/gpueventsynchronizer.cuh"
-#include "gromacs/mdlib/leapfrog_gpu.cuh"
+#include "gromacs/mdlib/leapfrog_gpu.h"
 #include "gromacs/mdlib/lincs_gpu.cuh"
 #include "gromacs/mdlib/settle_gpu.cuh"
 #include "gromacs/mdlib/update_constrain_gpu.h"