Rename NBAtomData into NBAtomDataGpu

NBAtomData can only be present in GPU code, hence it should has
Gpu prefix in its name.

diff --git a/src/gromacs/nbnxm/cuda/nbnxm_cuda.cu b/src/gromacs/nbnxm/cuda/nbnxm_cuda.cu
index 190721e2d8b5878a2f7cdeb3fc6d3039fcc31205..4869756afef30f3444f30c726773d4498623e367 100644 (file)
--- a/src/gromacs/nbnxm/cuda/nbnxm_cuda.cu
+++ b/src/gromacs/nbnxm/cuda/nbnxm_cuda.cu
@@ -122,7 +122,7 @@ namespace Nbnxm
 constexpr static int c_bufOpsThreadsPerBlock = 128;
 
 /*! Nonbonded kernel function pointer type */
-typedef void (*nbnxn_cu_kfunc_ptr_t)(const NBAtomData, const NBParamGpu, const gpu_plist, bool);
+typedef void (*nbnxn_cu_kfunc_ptr_t)(const NBAtomDataGpu, const NBParamGpu, const gpu_plist, bool);
 
 /*********************************/
 
@@ -443,7 +443,7 @@ static inline int calc_shmem_required_nonbonded(const int               num_thre
  */
 void gpu_launch_kernel(NbnxmGpu* nb, const gmx::StepWorkload& stepWork, const InteractionLocality iloc)
 {
-    NBAtomData*         adat         = nb->atdat;
+    NBAtomDataGpu*      adat         = nb->atdat;
     NBParamGpu*         nbp          = nb->nbparam;
     gpu_plist*          plist        = nb->plist[iloc];
     Nbnxm::GpuTimers*   timers       = nb->timers;
@@ -564,7 +564,7 @@ static inline int calc_shmem_required_prune(const int num_threads_z)
 
 void gpu_launch_kernel_pruneonly(NbnxmGpu* nb, const InteractionLocality iloc, const int numParts)
 {
-    NBAtomData*         adat         = nb->atdat;
+    NBAtomDataGpu*      adat         = nb->atdat;
     NBParamGpu*         nbp          = nb->nbparam;
     gpu_plist*          plist        = nb->plist[iloc];
     Nbnxm::GpuTimers*   timers       = nb->timers;
@@ -722,7 +722,7 @@ void nbnxn_gpu_x_to_nbat_x(const Nbnxm::Grid&        grid,
 {
     GMX_ASSERT(nb, "Need a valid nbnxn_gpu object");
 
-    NBAtomData* adat = nb->atdat;
+    NBAtomDataGpu* adat = nb->atdat;
 
     const int                  numColumns      = grid.numColumns();
     const int                  cellOffset      = grid.cellOffset();

diff --git a/src/gromacs/nbnxm/cuda/nbnxm_cuda_kernel.cuh b/src/gromacs/nbnxm/cuda/nbnxm_cuda_kernel.cuh
index 2b5a29778eacdb98b1bbd1eb9324f9c51b402a0c..49297a2f0af9f154276b9800e74178db9ccb6e46 100644 (file)
--- a/src/gromacs/nbnxm/cuda/nbnxm_cuda_kernel.cuh
+++ b/src/gromacs/nbnxm/cuda/nbnxm_cuda_kernel.cuh
@@ -160,7 +160,7 @@ __launch_bounds__(THREADS_PER_BLOCK)
         __global__ void NB_KERNEL_FUNC_NAME(nbnxn_kernel, _F_cuda)
 #    endif /* CALC_ENERGIES */
 #endif     /* PRUNE_NBL */
-                (const NBAtomData atdat, const NBParamGpu nbparam, const Nbnxm::gpu_plist plist, bool bCalcFshift)
+                (const NBAtomDataGpu atdat, const NBParamGpu nbparam, const Nbnxm::gpu_plist plist, bool bCalcFshift)
 #ifdef FUNCTION_DECLARATION_ONLY
                         ; /* Only do function declaration, omit the function body. */
 #else

diff --git a/src/gromacs/nbnxm/cuda/nbnxm_cuda_kernel_pruneonly.cu b/src/gromacs/nbnxm/cuda/nbnxm_cuda_kernel_pruneonly.cu
index 2ff980a6b0cf0dd32f13d521a98316a6a765698d..34de55d1948f95d8bc8523c0a706c23ee556e9f4 100644 (file)
--- a/src/gromacs/nbnxm/cuda/nbnxm_cuda_kernel_pruneonly.cu
+++ b/src/gromacs/nbnxm/cuda/nbnxm_cuda_kernel_pruneonly.cu
@@ -39,7 +39,7 @@
 #ifndef FUNCTION_DECLARATION_ONLY
 /* Instantiate external template functions */
 template __global__ void
-nbnxn_kernel_prune_cuda<false>(const NBAtomData, const NBParamGpu, const Nbnxm::gpu_plist, int, int);
+nbnxn_kernel_prune_cuda<false>(const NBAtomDataGpu, const NBParamGpu, const Nbnxm::gpu_plist, int, int);
 template __global__ void
-nbnxn_kernel_prune_cuda<true>(const NBAtomData, const NBParamGpu, const Nbnxm::gpu_plist, int, int);
+nbnxn_kernel_prune_cuda<true>(const NBAtomDataGpu, const NBParamGpu, const Nbnxm::gpu_plist, int, int);
 #endif

diff --git a/src/gromacs/nbnxm/cuda/nbnxm_cuda_kernel_pruneonly.cuh b/src/gromacs/nbnxm/cuda/nbnxm_cuda_kernel_pruneonly.cuh
index c5c55e667f6531969c4f6c770f5342578463585d..bcc3dd1b09d641a76ef6408d775c8e5686a2364f 100644 (file)
--- a/src/gromacs/nbnxm/cuda/nbnxm_cuda_kernel_pruneonly.cuh
+++ b/src/gromacs/nbnxm/cuda/nbnxm_cuda_kernel_pruneonly.cuh
@@ -104,7 +104,7 @@
  */
 template<bool haveFreshList>
 __launch_bounds__(THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP) __global__
-        void nbnxn_kernel_prune_cuda(const NBAtomData       atdat,
+        void nbnxn_kernel_prune_cuda(const NBAtomDataGpu    atdat,
                                      const NBParamGpu       nbparam,
                                      const Nbnxm::gpu_plist plist,
                                      int                    numParts,
@@ -115,9 +115,9 @@ __launch_bounds__(THREADS_PER_BLOCK, MIN_BLOCKS_PER_MP) __global__
 // Add extern declarations so each translation unit understands that
 // there will be a definition provided.
 extern template __global__ void
-nbnxn_kernel_prune_cuda<true>(const NBAtomData, const NBParamGpu, const Nbnxm::gpu_plist, int, int);
+nbnxn_kernel_prune_cuda<true>(const NBAtomDataGpu, const NBParamGpu, const Nbnxm::gpu_plist, int, int);
 extern template __global__ void
-nbnxn_kernel_prune_cuda<false>(const NBAtomData, const NBParamGpu, const Nbnxm::gpu_plist, int, int);
+nbnxn_kernel_prune_cuda<false>(const NBAtomDataGpu, const NBParamGpu, const Nbnxm::gpu_plist, int, int);
 #else
 {
 

diff --git a/src/gromacs/nbnxm/cuda/nbnxm_cuda_types.h b/src/gromacs/nbnxm/cuda/nbnxm_cuda_types.h
index f3d1bb5acc5d1a9a6ed04b55cd63631fbf3e1a36..91464b8010efaa9e88f52c75118998ca109695e2 100644 (file)
--- a/src/gromacs/nbnxm/cuda/nbnxm_cuda_types.h
+++ b/src/gromacs/nbnxm/cuda/nbnxm_cuda_types.h
@@ -80,7 +80,7 @@ struct NbnxmGpu
     bool bNonLocalStreamDoneMarked = false;
 
     /*! \brief atom data */
-    NBAtomData* atdat = nullptr;
+    NBAtomDataGpu* atdat = nullptr;
     /*! \brief array of atom indices */
     int* atomIndices = nullptr;
     /*! \brief size of atom indices */

diff --git a/src/gromacs/nbnxm/gpu_common_utils.h b/src/gromacs/nbnxm/gpu_common_utils.h
index 43da174d41ae9088f96e81d0fbe8876ea634e822..a3baa742d56e15f0f2e35c390e19863a5893c7ff 100644 (file)
--- a/src/gromacs/nbnxm/gpu_common_utils.h
+++ b/src/gromacs/nbnxm/gpu_common_utils.h
@@ -79,7 +79,7 @@ static inline bool canSkipNonbondedWork(const NbnxmGpu& nb, InteractionLocality
  * \param[in] atomLocality Atom locality specifier
  * \returns Range of indexes for selected locality.
  */
-static inline gmx::Range<int> getGpuAtomRange(const NBAtomData* atomData, const AtomLocality atomLocality)
+static inline gmx::Range<int> getGpuAtomRange(const NBAtomDataGpu* atomData, const AtomLocality atomLocality)
 {
     assert(atomData);
 

diff --git a/src/gromacs/nbnxm/gpu_types_common.h b/src/gromacs/nbnxm/gpu_types_common.h
index 8fa52b7f9eb29a5aa1b797c2c1a5c84261ca852d..74e2a02d07e35acb9643952e9cb13fd4aca55056 100644 (file)
--- a/src/gromacs/nbnxm/gpu_types_common.h
+++ b/src/gromacs/nbnxm/gpu_types_common.h
@@ -92,7 +92,7 @@ struct NBStagingData
 /** \internal
  * \brief Nonbonded atom data - both inputs and outputs.
  */
-struct NBAtomData
+struct NBAtomDataGpu
 {
     //! number of atoms
     int numAtoms;

diff --git a/src/gromacs/nbnxm/nbnxm_gpu_data_mgmt.cpp b/src/gromacs/nbnxm/nbnxm_gpu_data_mgmt.cpp
index 534d2b706d04b502015c4ce61332ba0a227aa290..bd0284963688a2c17524504dd4cc4207adb13514 100644 (file)
--- a/src/gromacs/nbnxm/nbnxm_gpu_data_mgmt.cpp
+++ b/src/gromacs/nbnxm/nbnxm_gpu_data_mgmt.cpp
@@ -258,7 +258,7 @@ void init_timings(gmx_wallclock_gpu_nbnxn_t* t)
 }
 
 /*! \brief Initialize \p atomdata first time; it only gets filled at pair-search. */
-static void initAtomdataFirst(NBAtomData*          atomdata,
+static void initAtomdataFirst(NBAtomDataGpu*       atomdata,
                               int                  numTypes,
                               const DeviceContext& deviceContext,
                               const DeviceStream&  localStream)
@@ -354,7 +354,7 @@ NbnxmGpu* gpu_init(const gmx::DeviceStreamManager& deviceStreamManager,
 {
     auto* nb                              = new NbnxmGpu();
     nb->deviceContext_                    = &deviceStreamManager.context();
-    nb->atdat                             = new NBAtomData;
+    nb->atdat                             = new NBAtomDataGpu;
     nb->nbparam                           = new NBParamGpu;
     nb->plist[InteractionLocality::Local] = new Nbnxm::gpu_plist;
     if (bLocalAndNonlocal)
@@ -428,7 +428,7 @@ NbnxmGpu* gpu_init(const gmx::DeviceStreamManager& deviceStreamManager,
 
 void gpu_upload_shiftvec(NbnxmGpu* nb, const nbnxn_atomdata_t* nbatom)
 {
-    NBAtomData*         adat        = nb->atdat;
+    NBAtomDataGpu*      adat        = nb->atdat;
     const DeviceStream& localStream = *nb->deviceStreams[InteractionLocality::Local];
 
     /* only if we have a dynamic box */
@@ -532,7 +532,7 @@ void gpu_init_atomdata(NbnxmGpu* nb, const nbnxn_atomdata_t* nbat)
 {
     bool                 bDoTime       = nb->bDoTime;
     Nbnxm::GpuTimers*    timers        = bDoTime ? nb->timers : nullptr;
-    NBAtomData*          atdat         = nb->atdat;
+    NBAtomDataGpu*       atdat         = nb->atdat;
     const DeviceContext& deviceContext = *nb->deviceContext_;
     const DeviceStream&  localStream   = *nb->deviceStreams[InteractionLocality::Local];
 
@@ -630,7 +630,7 @@ void gpu_init_atomdata(NbnxmGpu* nb, const nbnxn_atomdata_t* nbat)
 
 void gpu_clear_outputs(NbnxmGpu* nb, bool computeVirial)
 {
-    NBAtomData*         adat        = nb->atdat;
+    NBAtomDataGpu*      adat        = nb->atdat;
     const DeviceStream& localStream = *nb->deviceStreams[InteractionLocality::Local];
     // Clear forces
     clearDeviceBufferAsync(&adat->f, 0, nb->atdat->numAtoms, localStream);
@@ -777,7 +777,7 @@ void gpu_launch_cpyback(NbnxmGpu*                nb,
                "beginning of the copy back function.");
 
     /* extract the data */
-    NBAtomData*         adat         = nb->atdat;
+    NBAtomDataGpu*      adat         = nb->atdat;
     Nbnxm::GpuTimers*   timers       = nb->timers;
     bool                bDoTime      = nb->bDoTime;
     const DeviceStream& deviceStream = *nb->deviceStreams[iloc];
@@ -921,7 +921,7 @@ void gpu_copy_xq_to_gpu(NbnxmGpu* nb, const nbnxn_atomdata_t* nbatom, const Atom
 
     const InteractionLocality iloc = atomToInteractionLocality(atomLocality);
 
-    NBAtomData*         adat         = nb->atdat;
+    NBAtomDataGpu*      adat         = nb->atdat;
     gpu_plist*          plist        = nb->plist[iloc];
     Nbnxm::GpuTimers*   timers       = nb->timers;
     const DeviceStream& deviceStream = *nb->deviceStreams[iloc];
@@ -1104,8 +1104,8 @@ void gpu_free(NbnxmGpu* nb)
     delete nb->timers;
     sfree(nb->timings);
 
-    NBAtomData* atdat   = nb->atdat;
-    NBParamGpu* nbparam = nb->nbparam;
+    NBAtomDataGpu* atdat   = nb->atdat;
+    NBParamGpu*    nbparam = nb->nbparam;
 
     /* Free atdat */
     freeDeviceBuffer(&(nb->atdat->xq));

diff --git a/src/gromacs/nbnxm/opencl/nbnxm_ocl.cpp b/src/gromacs/nbnxm/opencl/nbnxm_ocl.cpp
index d14aad3da145f2cb4084fba130e4d345f8860922..3bde29600811dddcd8e7ad6ef88b37553bac18b5 100644 (file)
--- a/src/gromacs/nbnxm/opencl/nbnxm_ocl.cpp
+++ b/src/gromacs/nbnxm/opencl/nbnxm_ocl.cpp
@@ -508,7 +508,7 @@ static void fillin_ocl_structures(NBParamGpu* nbp, cl_nbparam_params_t* nbparams
  */
 void gpu_launch_kernel(NbnxmGpu* nb, const gmx::StepWorkload& stepWork, const Nbnxm::InteractionLocality iloc)
 {
-    NBAtomData*         adat         = nb->atdat;
+    NBAtomDataGpu*      adat         = nb->atdat;
     NBParamGpu*         nbp          = nb->nbparam;
     gpu_plist*          plist        = nb->plist[iloc];
     Nbnxm::GpuTimers*   timers       = nb->timers;
@@ -679,7 +679,7 @@ static inline int calc_shmem_required_prune(const int num_threads_z)
  */
 void gpu_launch_kernel_pruneonly(NbnxmGpu* nb, const InteractionLocality iloc, const int numParts)
 {
-    NBAtomData*         adat         = nb->atdat;
+    NBAtomDataGpu*      adat         = nb->atdat;
     NBParamGpu*         nbp          = nb->nbparam;
     gpu_plist*          plist        = nb->plist[iloc];
     Nbnxm::GpuTimers*   timers       = nb->timers;

diff --git a/src/gromacs/nbnxm/opencl/nbnxm_ocl_types.h b/src/gromacs/nbnxm/opencl/nbnxm_ocl_types.h
index 95558805bde83a9788b6adb7d8359a6548e382ab..3037663842250d1b2022f6fdddb2f1d0c7090db0 100644 (file)
--- a/src/gromacs/nbnxm/opencl/nbnxm_ocl_types.h
+++ b/src/gromacs/nbnxm/opencl/nbnxm_ocl_types.h
@@ -167,7 +167,7 @@ struct NbnxmGpu
     bool bNonLocalStreamDoneMarked = false;
 
     //! atom data
-    NBAtomData* atdat = nullptr;
+    NBAtomDataGpu* atdat = nullptr;
     //! parameters required for the non-bonded calc.
     NBParamGpu* nbparam = nullptr;
     //! pair-list data structures (local and non-local)

diff --git a/src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel.cpp b/src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel.cpp
index 44fad58ac31819828e49b10f91193a83f3238573..1916010e23bba1ed42ae6f39c9f9435b843b570c 100644 (file)
--- a/src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel.cpp
+++ b/src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel.cpp
@@ -987,7 +987,7 @@ cl::sycl::event chooseAndLaunchNbnxmKernel(bool          doPruneNBL,
 
 void launchNbnxmKernel(NbnxmGpu* nb, const gmx::StepWorkload& stepWork, const InteractionLocality iloc)
 {
-    NBAtomData*         adat         = nb->atdat;
+    NBAtomDataGpu*      adat         = nb->atdat;
     NBParamGpu*         nbp          = nb->nbparam;
     gpu_plist*          plist        = nb->plist[iloc];
     const bool          doPruneNBL   = (plist->haveFreshList && !nb->didPrune[iloc]);

diff --git a/src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel_pruneonly.cpp b/src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel_pruneonly.cpp
index d784de3a2cf3369cdeb37fd092b632f4b5a5d22e..a2cc1f8a0ddc41f140ad364ff6e0fc5c1b5f52d2 100644 (file)
--- a/src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel_pruneonly.cpp
+++ b/src/gromacs/nbnxm/sycl/nbnxm_sycl_kernel_pruneonly.cpp
@@ -264,7 +264,7 @@ void launchNbnxmKernelPruneOnly(NbnxmGpu*                 nb,
                                 const int                 part,
                                 const int                 numSciInPart)
 {
-    NBAtomData*         adat          = nb->atdat;
+    NBAtomDataGpu*      adat          = nb->atdat;
     NBParamGpu*         nbp           = nb->nbparam;
     gpu_plist*          plist         = nb->plist[iloc];
     const bool          haveFreshList = plist->haveFreshList;

diff --git a/src/gromacs/nbnxm/sycl/nbnxm_sycl_types.h b/src/gromacs/nbnxm/sycl/nbnxm_sycl_types.h
index 6a82823b377ca6ae7afeb2dae9d6f8fe8ee8612b..a0d8c914fce480b364c7807570a787863b9d3988 100644 (file)
--- a/src/gromacs/nbnxm/sycl/nbnxm_sycl_types.h
+++ b/src/gromacs/nbnxm/sycl/nbnxm_sycl_types.h
@@ -72,7 +72,7 @@ struct NbnxmGpu
     /*! \brief true indicates that the nonlocal_done event was marked */
     bool bNonLocalStreamDoneMarked = false;
     /*! \brief atom data */
-    NBAtomData* atdat = nullptr;
+    NBAtomDataGpu* atdat = nullptr;
 
     // Data for GPU-side coordinate conversion between integrator and NBNXM
     /*! \brief array of atom indices */