From: Artem Zhmurov Date: Sat, 3 Apr 2021 09:51:17 +0000 (+0300) Subject: Rename NBAtomData into NBAtomDataGpu X-Git-Url: http://biod.pnpi.spb.ru/gitweb/?a=commitdiff_plain;h=69edbe0861e80aec899af9696dd34d9b86d938c9;p=alexxy%2Fgromacs.git 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 190721e2d8..4869756afe 100644 --- 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 2b5a29778e..49297a2f0a 100644 --- 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 2ff980a6b0..34de55d194 100644 --- 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(const NBAtomData, const NBParamGpu, const Nbnxm::gpu_plist, int, int); +nbnxn_kernel_prune_cuda(const NBAtomDataGpu, const NBParamGpu, const Nbnxm::gpu_plist, int, int); template __global__ void -nbnxn_kernel_prune_cuda(const NBAtomData, const NBParamGpu, const Nbnxm::gpu_plist, int, int); +nbnxn_kernel_prune_cuda(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 c5c55e667f..bcc3dd1b09 100644 --- 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 __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(const NBAtomData, const NBParamGpu, const Nbnxm::gpu_plist, int, int); +nbnxn_kernel_prune_cuda(const NBAtomDataGpu, const NBParamGpu, const Nbnxm::gpu_plist, int, int); extern template __global__ void -nbnxn_kernel_prune_cuda(const NBAtomData, const NBParamGpu, const Nbnxm::gpu_plist, int, int); +nbnxn_kernel_prune_cuda(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 f3d1bb5acc..91464b8010 100644 --- 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 43da174d41..a3baa742d5 100644 --- 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 getGpuAtomRange(const NBAtomData* atomData, const AtomLocality atomLocality) +static inline gmx::Range 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 8fa52b7f9e..74e2a02d07 100644 --- 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 534d2b706d..bd02849636 100644 --- 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 d14aad3da1..3bde296008 100644 --- 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 95558805bd..3037663842 100644 --- 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 44fad58ac3..1916010e23 100644 --- 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 d784de3a2c..a2cc1f8a0d 100644 --- 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 6a82823b37..a0d8c914fc 100644 --- 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 */