From 5145e7cfcd8a0be3efaa19ad7311e039148a3108 Mon Sep 17 00:00:00 2001 From: Paul Bauer Date: Thu, 29 Oct 2020 17:47:33 +0100 Subject: [PATCH] Fix CUDA code issues Remove unused template. Fix variable initialization order and remove obsolete variable. Fix missing newline. --- src/gromacs/domdec/gpuhaloexchange_impl.cu | 2 +- src/gromacs/domdec/gpuhaloexchange_impl.cuh | 2 - src/gromacs/gpu_utils/device_stream.cu | 2 +- .../nbnxm/cuda/nbnxm_buffer_ops_kernels.cuh | 50 ------------------- 4 files changed, 2 insertions(+), 54 deletions(-) diff --git a/src/gromacs/domdec/gpuhaloexchange_impl.cu b/src/gromacs/domdec/gpuhaloexchange_impl.cu index 9efebf6904..e7045d8b2a 100644 --- a/src/gromacs/domdec/gpuhaloexchange_impl.cu +++ b/src/gromacs/domdec/gpuhaloexchange_impl.cu @@ -464,7 +464,6 @@ GpuHaloExchange::Impl::Impl(gmx_domdec_t* dd, int pulse, gmx_wallcycle* wcycle) : dd_(dd), - dimIndex_(dimIndex), sendRankX_(dd->neighbor[dimIndex][1]), recvRankX_(dd->neighbor[dimIndex][0]), sendRankF_(dd->neighbor[dimIndex][0]), @@ -475,6 +474,7 @@ GpuHaloExchange::Impl::Impl(gmx_domdec_t* dd, deviceContext_(deviceContext), localStream_(localStream), nonLocalStream_(nonLocalStream), + dimIndex_(dimIndex), pulse_(pulse), wcycle_(wcycle) { diff --git a/src/gromacs/domdec/gpuhaloexchange_impl.cuh b/src/gromacs/domdec/gpuhaloexchange_impl.cuh index 761938a013..5dd619a343 100644 --- a/src/gromacs/domdec/gpuhaloexchange_impl.cuh +++ b/src/gromacs/domdec/gpuhaloexchange_impl.cuh @@ -204,8 +204,6 @@ private: int dimIndex_ = 0; //! The pulse corresponding to this halo exchange instance int pulse_ = 0; - //! Number of zones. Always 1 for 1-D case. - const int nzone_ = 1; //! The wallclock counter gmx_wallcycle* wcycle_ = nullptr; //! The atom offset for receive (x) or send (f) for dimension index and pulse corresponding to this halo exchange instance diff --git a/src/gromacs/gpu_utils/device_stream.cu b/src/gromacs/gpu_utils/device_stream.cu index 0e07b00b27..cc1f879862 100644 --- a/src/gromacs/gpu_utils/device_stream.cu +++ b/src/gromacs/gpu_utils/device_stream.cu @@ -117,4 +117,4 @@ void DeviceStream::synchronize() const gmx::formatString("cudaStreamSynchronize failed (CUDA error %d: %s).", stat, cudaGetErrorString(stat)) .c_str()); -} \ No newline at end of file +} diff --git a/src/gromacs/nbnxm/cuda/nbnxm_buffer_ops_kernels.cuh b/src/gromacs/nbnxm/cuda/nbnxm_buffer_ops_kernels.cuh index db3fb4a939..9a9ffc6c1c 100644 --- a/src/gromacs/nbnxm/cuda/nbnxm_buffer_ops_kernels.cuh +++ b/src/gromacs/nbnxm/cuda/nbnxm_buffer_ops_kernels.cuh @@ -117,53 +117,3 @@ static __global__ void nbnxn_gpu_x_to_nbat_x_kernel(int numColumns, } } } - -/*! \brief CUDA kernel to sum up the force components - * - * \tparam accumulateForce If the initial forces in \p gm_fTotal should be saved. - * \tparam addPmeForce Whether the PME force should be added to the total. - * - * \param[in] gm_fNB Non-bonded forces in nbnxm format. - * \param[in] gm_fPme PME forces. - * \param[in,out] gm_fTotal Force buffer to be reduced into. - * \param[in] cell Cell index mapping. - * \param[in] atomStart Start atom index. - * \param[in] numAtoms Number of atoms. - */ -template -static __global__ void nbnxn_gpu_add_nbat_f_to_f_kernel(const float3* __restrict__ gm_fNB, - const float3* __restrict__ gm_fPme, - float3* gm_fTotal, - const int* __restrict__ gm_cell, - const int atomStart, - const int numAtoms) -{ - - /* map particle-level parallelism to 1D CUDA thread and block index */ - const int threadIndex = blockIdx.x * blockDim.x + threadIdx.x; - - /* perform addition for each particle*/ - if (threadIndex < numAtoms) - { - - const int i = gm_cell[atomStart + threadIndex]; - float3* gm_fDest = &gm_fTotal[atomStart + threadIndex]; - float3 temp; - - if (accumulateForce) - { - temp = *gm_fDest; - temp += gm_fNB[i]; - } - else - { - temp = gm_fNB[i]; - } - if (addPmeForce) - { - temp += gm_fPme[atomStart + threadIndex]; - } - *gm_fDest = temp; - } - return; -} -- 2.22.0