moebiusband/MD-Bench
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

♻️ Fixing refactoring step

Browse Source
This commit is contained in:
Martin Bauernfeind 2022-06-23 19:32:09 +02:00
parent 62cfc22856
commit 7f068a6959
2 changed files with 12 additions and 12 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

20
src/force.cu
View File

@ -40,7 +40,7 @@ extern "C" {
// cuda kernel // cuda kernel
__global__ void calc_force( __global__ void calc_force(
Atom a, Atom *a,
MD_FLOAT cutforcesq, MD_FLOAT sigma6, MD_FLOAT epsilon, MD_FLOAT cutforcesq, MD_FLOAT sigma6, MD_FLOAT epsilon,
int Nlocal, int neigh_maxneighs, int *neigh_neighbors, int *neigh_numneigh) { int Nlocal, int neigh_maxneighs, int *neigh_neighbors, int *neigh_numneigh) {
@ -49,7 +49,7 @@ __global__ void calc_force(
return; return;
} }
Atom *atom = &a; Atom *atom = a;
const int numneighs = neigh_numneigh[i]; const int numneighs = neigh_numneigh[i];
@ -139,7 +139,7 @@ int get_num_threads() {
return num_threads; return num_threads;
} }
void cuda_final_integrate(bool doReneighbour, Parameter *param, Atom *atom) { void cuda_final_integrate(bool doReneighbour, Parameter *param, Atom *atom, Atom *c_atom) {
const int Nlocal = atom->Nlocal; const int Nlocal = atom->Nlocal;
const int num_threads = get_num_threads(); const int num_threads = get_num_threads();
@ -157,7 +157,7 @@ void cuda_final_integrate(bool doReneighbour, Parameter *param, Atom *atom) {
} }
} }
void cuda_initial_integrate(bool doReneighbour, Parameter *param, Atom *atom) { void cuda_initial_integrate(bool doReneighbour, Parameter *param, Atom *atom, Atom *c_atom) {
const int Nlocal = atom->Nlocal; const int Nlocal = atom->Nlocal;
const int num_threads = get_num_threads(); const int num_threads = get_num_threads();
@ -194,11 +194,11 @@ double computeForce(
const int num_threads = get_num_threads(); const int num_threads = get_num_threads();
c_atom.Natoms = atom->Natoms; c_atom->Natoms = atom->Natoms;
c_atom.Nlocal = atom->Nlocal; c_atom->Nlocal = atom->Nlocal;
c_atom.Nghost = atom->Nghost; c_atom->Nghost = atom->Nghost;
c_atom.Nmax = atom->Nmax; c_atom->Nmax = atom->Nmax;
c_atom.ntypes = atom->ntypes; c_atom->ntypes = atom->ntypes;
/* /*
int nDevices; int nDevices;
@ -219,7 +219,7 @@ double computeForce(
cudaProfilerStart(); cudaProfilerStart();
checkCUDAError( "c_atom.x memcpy", cudaMemcpy(c_atom.x, atom->x, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice) ); checkCUDAError( "c_atom->x memcpy", cudaMemcpy(c_atom->x, atom->x, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice) );
if(reneighbourHappenend) { if(reneighbourHappenend) {
checkCUDAError( "c_neighbor->numneigh memcpy", cudaMemcpy(c_neighbor->numneigh, neighbor->numneigh, sizeof(int) * Nlocal, cudaMemcpyHostToDevice) ); checkCUDAError( "c_neighbor->numneigh memcpy", cudaMemcpy(c_neighbor->numneigh, neighbor->numneigh, sizeof(int) * Nlocal, cudaMemcpyHostToDevice) );

4
src/main.c
View File

@ -310,7 +310,7 @@ int main(int argc, char** argv)
const bool doReneighbour = (n + 1) % param.every == 0; const bool doReneighbour = (n + 1) % param.every == 0;
cuda_initial_integrate(doReneighbour, &param, &atom); cuda_initial_integrate(doReneighbour, &param, &atom, &c_atom);
if(doReneighbour) { if(doReneighbour) {
timer[NEIGH] += reneighbour(&param, &atom, &neighbor); timer[NEIGH] += reneighbour(&param, &atom, &neighbor);
@ -328,7 +328,7 @@ int main(int argc, char** argv)
#endif #endif
} }
cuda_final_integrate(doReneighbour, &param, &atom); cuda_final_integrate(doReneighbour, &param, &atom, &c_atom);
if(!((n + 1) % param.nstat) && (n+1) < param.ntimes) { if(!((n + 1) % param.nstat) && (n+1) < param.ntimes) {
computeThermo(n + 1, &param, &atom); computeThermo(n + 1, &param, &atom);