Skip to content

GitLab

Commit edfcdec3 authored by Peter Eastman's avatar Peter Eastman
Browse files

Bug fixes

parent 8647eaad
Hide whitespace changes
Inline Side-by-side
Showing with 7 additions and 3 deletions
platforms/cuda/src/CudaKernels.cpp
View file @ edfcdec3
...@@ -1623,6 +1623,8 @@ CudaCalcNonbondedForceKernel::~CudaCalcNonbondedForceKernel() { ...@@ -1623,6 +1623,8 @@ CudaCalcNonbondedForceKernel::~CudaCalcNonbondedForceKernel() {
delete sort; delete sort;
if (fft != NULL) if (fft != NULL)
delete fft; delete fft;
if (dispersionFft != NULL)
delete dispersionFft;
if (pmeio != NULL) if (pmeio != NULL)
delete pmeio; delete pmeio;
if (hasInitializedFFT) { if (hasInitializedFFT) {
...@@ -1915,6 +1917,8 @@ void CudaCalcNonbondedForceKernel::initialize(const System& system, const Nonbon ...@@ -1915,6 +1917,8 @@ void CudaCalcNonbondedForceKernel::initialize(const System& system, const Nonbon
if (useCudaFFT) { if (useCudaFFT) {
cufftSetStream(fftForward, pmeStream); cufftSetStream(fftForward, pmeStream);
cufftSetStream(fftBackward, pmeStream); cufftSetStream(fftBackward, pmeStream);
cufftSetStream(dispersionFftForward, pmeStream);
cufftSetStream(dispersionFftBackward, pmeStream);
} }
CHECK_RESULT(cuEventCreate(&pmeSyncEvent, CU_EVENT_DISABLE_TIMING), "Error creating event for NonbondedForce"); CHECK_RESULT(cuEventCreate(&pmeSyncEvent, CU_EVENT_DISABLE_TIMING), "Error creating event for NonbondedForce");
int recipForceGroup = force.getReciprocalSpaceForceGroup(); int recipForceGroup = force.getReciprocalSpaceForceGroup();
...@@ -2174,7 +2178,7 @@ double CudaCalcNonbondedForceKernel::execute(ContextImpl& context, bool includeF ...@@ -2174,7 +2178,7 @@ double CudaCalcNonbondedForceKernel::execute(ContextImpl& context, bool includeF
cufftExecR2C(dispersionFftForward, (float*) directPmeGrid->getDevicePointer(), (float2*) reciprocalPmeGrid->getDevicePointer()); cufftExecR2C(dispersionFftForward, (float*) directPmeGrid->getDevicePointer(), (float2*) reciprocalPmeGrid->getDevicePointer());
} }
else { else {
fft->execFFT(*directPmeGrid, *reciprocalPmeGrid, true); dispersionFft->execFFT(*directPmeGrid, *reciprocalPmeGrid, true);
} }
if (includeEnergy) { if (includeEnergy) {
...@@ -2196,7 +2200,7 @@ double CudaCalcNonbondedForceKernel::execute(ContextImpl& context, bool includeF ...@@ -2196,7 +2200,7 @@ double CudaCalcNonbondedForceKernel::execute(ContextImpl& context, bool includeF
cufftExecC2R(dispersionFftBackward, (float2*) reciprocalPmeGrid->getDevicePointer(), (float*) directPmeGrid->getDevicePointer()); cufftExecC2R(dispersionFftBackward, (float2*) reciprocalPmeGrid->getDevicePointer(), (float*) directPmeGrid->getDevicePointer());
} }
else { else {
fft->execFFT(*reciprocalPmeGrid, *directPmeGrid, false); dispersionFft->execFFT(*reciprocalPmeGrid, *directPmeGrid, false);
} }
void* interpolateArgs[] = {&cu.getPosq().getDevicePointer(), &cu.getForce().getDevicePointer(), &directPmeGrid->getDevicePointer(), cu.getPeriodicBoxSizePointer(), void* interpolateArgs[] = {&cu.getPosq().getDevicePointer(), &cu.getForce().getDevicePointer(), &directPmeGrid->getDevicePointer(), cu.getPeriodicBoxSizePointer(),
......
platforms/cuda/src/kernels/pme.cu
View file @ edfcdec3
...@@ -238,7 +238,7 @@ gridEvaluateEnergy(real2* __restrict__ halfcomplex_pmeGrid, mixed* __restrict__ ...@@ -238,7 +238,7 @@ gridEvaluateEnergy(real2* __restrict__ halfcomplex_pmeGrid, mixed* __restrict__
#endif #endif
energy += eterm*(grid.x*grid.x + grid.y*grid.y); energy += eterm*(grid.x*grid.x + grid.y*grid.y);
} }
#ifdef USE_PME_STREAM #if defined(USE_PME_STREAM) && !defined(USE_LJPME)
energyBuffer[blockIdx.x*blockDim.x+threadIdx.x] = 0.5f*energy; energyBuffer[blockIdx.x*blockDim.x+threadIdx.x] = 0.5f*energy;
#else #else
energyBuffer[blockIdx.x*blockDim.x+threadIdx.x] += 0.5f*energy; energyBuffer[blockIdx.x*blockDim.x+threadIdx.x] += 0.5f*energy;
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment