From 71798f5ec547ca2241d77e9a3db60b7bf0f15a6b Mon Sep 17 00:00:00 2001 From: Martin Bauernfeind Date: Tue, 5 Jul 2022 00:54:11 +0200 Subject: [PATCH] :bug: Fixed aforementioned correctness issue by deleting a superflous cudaMemcpy in computeForce() that was overwriting correct data with incorrect data --- src/force.cu | 5 ----- src/main.c | 1 + src/pbc.cu | 4 +++- 3 files changed, 4 insertions(+), 6 deletions(-) diff --git a/src/force.cu b/src/force.cu index e007b44..0f06325 100644 --- a/src/force.cu +++ b/src/force.cu @@ -193,11 +193,6 @@ double computeForce( cudaProfilerStart(); - - if(!reneighbourHappenend) { - checkCUDAError( "computeForce c_atom->x memcpy", cudaMemcpy(c_atom->x, atom->x, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice) ); - } - const int num_blocks = ceil((float)Nlocal / (float)num_threads_per_block); double S = getTimeStamp(); diff --git a/src/main.c b/src/main.c index 19efda9..1cdc09f 100644 --- a/src/main.c +++ b/src/main.c @@ -363,6 +363,7 @@ int main(int argc, char** argv) cuda_final_integrate(doReneighbour, ¶m, &atom, &c_atom, num_threads_per_block); if(!((n + 1) % param.nstat) && (n+1) < param.ntimes) { + checkCUDAError("computeThermo atom->x memcpy back", cudaMemcpy(atom.x, c_atom.x, atom.Nmax * sizeof(MD_FLOAT) * 3, cudaMemcpyDeviceToHost) ); computeThermo(n + 1, ¶m, &atom); } diff --git a/src/pbc.cu b/src/pbc.cu index fc32144..565d0c8 100644 --- a/src/pbc.cu +++ b/src/pbc.cu @@ -131,7 +131,9 @@ void updatePbc_cuda(Atom *atom, Parameter *param, Atom *c_atom, bool doReneighbo * MD_FLOAT xprd, MD_FLOAT yprd, MD_FLOAT zprd) * */ computePbcUpdate<<>>(*c_atom, c_PBCx, c_PBCy, c_PBCz, xprd, yprd, zprd); - checkCUDAError( "updatePbc atom->x memcpy back", cudaMemcpy(atom->x, c_atom->x, atom->Nmax * sizeof(MD_FLOAT) * 3, cudaMemcpyDeviceToHost) ); + if(doReneighbor){ + checkCUDAError( "updatePbc atom->x memcpy back", cudaMemcpy(atom->x, c_atom->x, atom->Nmax * sizeof(MD_FLOAT) * 3, cudaMemcpyDeviceToHost) ); + } } /* relocate atoms that have left domain according