From 72e4599acc64b65f360efc7151d5c3440f78c0bc Mon Sep 17 00:00:00 2001 From: Maximilian Gaul Date: Sat, 1 Jan 2022 12:56:42 +0100 Subject: [PATCH] Copy neighbour lists only when reneighbouring happens, added measurements + logbook update --- src/force.cu | 6 ++++-- src/main.c | 1 - 2 files changed, 4 insertions(+), 3 deletions(-) diff --git a/src/force.cu b/src/force.cu index 4575035..0e14356 100644 --- a/src/force.cu +++ b/src/force.cu @@ -179,8 +179,10 @@ double computeForce( checkCUDAError( "c_atom.x memcpy", cudaMemcpy(c_atom.x, atom->x, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice) ); - checkCUDAError( "c_neigh_numneigh memcpy", cudaMemcpy(c_neigh_numneigh, neighbor->numneigh, sizeof(int) * Nlocal, cudaMemcpyHostToDevice) ); - checkCUDAError( "c_neighs memcpy", cudaMemcpy(c_neighs, neighbor->neighbors, sizeof(int) * Nlocal * neighbor->maxneighs, cudaMemcpyHostToDevice) ); + if(reneighbourHappenend) { + checkCUDAError( "c_neigh_numneigh memcpy", cudaMemcpy(c_neigh_numneigh, neighbor->numneigh, sizeof(int) * Nlocal, cudaMemcpyHostToDevice) ); + checkCUDAError( "c_neighs memcpy", cudaMemcpy(c_neighs, neighbor->neighbors, sizeof(int) * Nlocal * neighbor->maxneighs, cudaMemcpyHostToDevice) ); + } const int num_threads_per_block = num_threads; // this should be multiple of 32 as operations are performed at the level of warps const int num_blocks = ceil((float)Nlocal / (float)num_threads_per_block); diff --git a/src/main.c b/src/main.c index cb8a4f3..9767bd8 100644 --- a/src/main.c +++ b/src/main.c @@ -279,7 +279,6 @@ int main(int argc, char** argv) initialIntegrate(¶m, &atom); const bool doReneighbour = (n + 1) % param.every == 0; - const bool doesReneighbourNextRound = (n + 2) % param.every == 0; if(doReneighbour) { timer[NEIGH] += reneighbour(¶m, &atom, &neighbor);