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

Trying to debug segfault if cudaMemcpy is limited to neighbour list update

Browse Source
This commit is contained in:
Maximilian Gaul 2021-12-25 15:36:08 +01:00
parent 0ea0587442
commit 8009b54113
3 changed files with 79 additions and 21 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

63
src/force.cu
View File

@ -99,7 +99,9 @@ __global__ void calc_force(
extern "C" { extern "C" {
bool initialized = false; bool initialized = false;
Atom c_atom; static Atom c_atom;
int *c_neighs;
int *c_neigh_numneigh;
double computeForce( double computeForce(
bool reneighbourHappenend, bool reneighbourHappenend,
@ -117,9 +119,12 @@ double computeForce(
MD_FLOAT sigma6 = param->sigma6; MD_FLOAT sigma6 = param->sigma6;
MD_FLOAT epsilon = param->epsilon; MD_FLOAT epsilon = param->epsilon;
#endif #endif
printf("-1\r\n");
cudaProfilerStart(); cudaProfilerStart();
printf("0\r\n");
for(int i = 0; i < Nlocal; i++) { for(int i = 0; i < Nlocal; i++) {
fx[i] = 0.0; fx[i] = 0.0;
fy[i] = 0.0; fy[i] = 0.0;
@ -140,6 +145,8 @@ double computeForce(
c_atom.Nmax = atom->Nmax; c_atom.Nmax = atom->Nmax;
c_atom.ntypes = atom->ntypes; c_atom.ntypes = atom->ntypes;
printf("0.1\r\n");
/* /*
int nDevices; int nDevices;
cudaGetDeviceCount(&nDevices); cudaGetDeviceCount(&nDevices);
@ -160,6 +167,8 @@ double computeForce(
// HINT: Run with cuda-memcheck ./MDBench-NVCC in case of error // HINT: Run with cuda-memcheck ./MDBench-NVCC in case of error
// HINT: Only works for data layout = AOS!!! // HINT: Only works for data layout = AOS!!!
printf("1\r\n");
if(!initialized) { if(!initialized) {
checkCUDAError( "c_atom.x malloc", cudaMalloc((void**)&(c_atom.x), sizeof(MD_FLOAT) * atom->Nmax * 3) ); checkCUDAError( "c_atom.x malloc", cudaMalloc((void**)&(c_atom.x), sizeof(MD_FLOAT) * atom->Nmax * 3) );
checkCUDAError( "c_atom.fx malloc", cudaMalloc((void**)&(c_atom.fx), sizeof(MD_FLOAT) * Nlocal) ); checkCUDAError( "c_atom.fx malloc", cudaMalloc((void**)&(c_atom.fx), sizeof(MD_FLOAT) * Nlocal) );
@ -170,25 +179,31 @@ double computeForce(
checkCUDAError( "c_atom.sigma6 malloc", cudaMalloc((void**)&(c_atom.sigma6), sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes) ); checkCUDAError( "c_atom.sigma6 malloc", cudaMalloc((void**)&(c_atom.sigma6), sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes) );
checkCUDAError( "c_atom.cutforcesq malloc", cudaMalloc((void**)&(c_atom.cutforcesq), sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes) ); checkCUDAError( "c_atom.cutforcesq malloc", cudaMalloc((void**)&(c_atom.cutforcesq), sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes) );
initialized = true; checkCUDAError( "c_neighs malloc", cudaMalloc((void**)&c_neighs, sizeof(int) * Nlocal * neighbor->maxneighs) );
checkCUDAError( "c_neigh_numneigh malloc", cudaMalloc((void**)&c_neigh_numneigh, sizeof(int) * Nlocal) );
} }
checkCUDAError( "c_atom.x memcpy", cudaMemcpy(c_atom.x, atom->x, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice) ); printf("2\r\n");
checkCUDAError( "c_atom.fx memcpy", cudaMemcpy(c_atom.fx, fx, sizeof(MD_FLOAT) * Nlocal, cudaMemcpyHostToDevice) );
checkCUDAError( "c_atom.fy memcpy", cudaMemcpy(c_atom.fy, fy, sizeof(MD_FLOAT) * Nlocal, cudaMemcpyHostToDevice) );
checkCUDAError( "c_atom.fz memcpy", cudaMemcpy(c_atom.fz, fz, sizeof(MD_FLOAT) * Nlocal, cudaMemcpyHostToDevice) );
checkCUDAError( "c_atom.type memcpy", cudaMemcpy(c_atom.type, atom->type, sizeof(int) * atom->Nmax, cudaMemcpyHostToDevice) );
checkCUDAError( "c_atom.epsilon memcpy", cudaMemcpy(c_atom.epsilon, atom->epsilon, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice) );
checkCUDAError( "c_atom.sigma6 memcpy", cudaMemcpy(c_atom.sigma6, atom->sigma6, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice) );
checkCUDAError( "c_atom.cutforcesq memcpy", cudaMemcpy(c_atom.cutforcesq, atom->cutforcesq, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice) );
int *c_neighs; if(reneighbourHappenend || !initialized) {
checkCUDAError( "c_neighs malloc", cudaMalloc((void**)&c_neighs, sizeof(int) * Nlocal * neighbor->maxneighs) ); checkCUDAError( "c_atom.x memcpy", cudaMemcpy(c_atom.x, atom->x, sizeof(MD_FLOAT) * atom->Nmax * 3, cudaMemcpyHostToDevice) );
checkCUDAError( "c_neighs memcpy", cudaMemcpy(c_neighs, neighbor->neighbors, sizeof(int) * Nlocal * neighbor->maxneighs, cudaMemcpyHostToDevice) ); checkCUDAError( "c_atom.fx memcpy", cudaMemcpy(c_atom.fx, fx, sizeof(MD_FLOAT) * Nlocal, cudaMemcpyHostToDevice) );
checkCUDAError( "c_atom.fy memcpy", cudaMemcpy(c_atom.fy, fy, sizeof(MD_FLOAT) * Nlocal, cudaMemcpyHostToDevice) );
checkCUDAError( "c_atom.fz memcpy", cudaMemcpy(c_atom.fz, fz, sizeof(MD_FLOAT) * Nlocal, cudaMemcpyHostToDevice) );
checkCUDAError( "c_atom.type memcpy", cudaMemcpy(c_atom.type, atom->type, sizeof(int) * atom->Nmax, cudaMemcpyHostToDevice) );
checkCUDAError( "c_atom.epsilon memcpy", cudaMemcpy(c_atom.epsilon, atom->epsilon, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice) );
checkCUDAError( "c_atom.sigma6 memcpy", cudaMemcpy(c_atom.sigma6, atom->sigma6, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice) );
checkCUDAError( "c_atom.cutforcesq memcpy", cudaMemcpy(c_atom.cutforcesq, atom->cutforcesq, sizeof(MD_FLOAT) * atom->ntypes * atom->ntypes, cudaMemcpyHostToDevice) );
int *c_neigh_numneigh; checkCUDAError( "c_neigh_numneigh memcpy", cudaMemcpy(c_neigh_numneigh, neighbor->numneigh, sizeof(int) * Nlocal, cudaMemcpyHostToDevice) );
checkCUDAError( "c_neigh_numneigh malloc", cudaMalloc((void**)&c_neigh_numneigh, sizeof(int) * Nlocal) ); checkCUDAError( "c_neighs memcpy", cudaMemcpy(c_neighs, neighbor->neighbors, sizeof(int) * Nlocal * neighbor->maxneighs, cudaMemcpyHostToDevice) );
checkCUDAError( "c_neigh_numneigh memcpy", cudaMemcpy(c_neigh_numneigh, neighbor->numneigh, sizeof(int) * Nlocal, cudaMemcpyHostToDevice) ); }
printf("3\r\n");
printf("4\r\n");
printf("5\r\n");
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_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); const int num_blocks = ceil((float)Nlocal / (float)num_threads_per_block);
@ -201,10 +216,16 @@ double computeForce(
checkCUDAError( "PeekAtLastError", cudaPeekAtLastError() ); checkCUDAError( "PeekAtLastError", cudaPeekAtLastError() );
checkCUDAError( "DeviceSync", cudaDeviceSynchronize() ); checkCUDAError( "DeviceSync", cudaDeviceSynchronize() );
printf("6\r\n");
// copy results in c_atom.fx/fy/fz to atom->fx/fy/fz // copy results in c_atom.fx/fy/fz to atom->fx/fy/fz
cudaMemcpy(atom->fx, c_atom.fx, sizeof(MD_FLOAT) * Nlocal, cudaMemcpyDeviceToHost); if(reneighbourHappenend) {
cudaMemcpy(atom->fy, c_atom.fy, sizeof(MD_FLOAT) * Nlocal, cudaMemcpyDeviceToHost); cudaMemcpy(atom->fx, c_atom.fx, sizeof(MD_FLOAT) * Nlocal, cudaMemcpyDeviceToHost);
cudaMemcpy(atom->fz, c_atom.fz, sizeof(MD_FLOAT) * Nlocal, cudaMemcpyDeviceToHost); cudaMemcpy(atom->fy, c_atom.fy, sizeof(MD_FLOAT) * Nlocal, cudaMemcpyDeviceToHost);
cudaMemcpy(atom->fz, c_atom.fz, sizeof(MD_FLOAT) * Nlocal, cudaMemcpyDeviceToHost);
}
printf("7\r\n");
/* /*
cudaFree(c_atom.x); cudaFree(c_atom.x);
@ -215,13 +236,15 @@ double computeForce(
cudaFree(c_atom.cutforcesq); cudaFree(c_atom.cutforcesq);
*/ */
cudaFree(c_neighs); cudaFree(c_neigh_numneigh); // cudaFree(c_neighs); cudaFree(c_neigh_numneigh);
cudaProfilerStop(); cudaProfilerStop();
LIKWID_MARKER_STOP("force"); LIKWID_MARKER_STOP("force");
double E = getTimeStamp(); double E = getTimeStamp();
initialized = true;
return E-S; return E-S;
} }
} }

14
src/main.c
View File

@ -111,13 +111,20 @@ double reneighbour(
{ {
double S, E; double S, E;
printf("10.1\r\n");
S = getTimeStamp(); S = getTimeStamp();
LIKWID_MARKER_START("reneighbour"); LIKWID_MARKER_START("reneighbour");
printf("10.2\r\n");
updateAtomsPbc(atom, param); updateAtomsPbc(atom, param);
printf("10.3\r\n");
setupPbc(atom, param); setupPbc(atom, param);
printf("10.4\r\n");
updatePbc(atom, param); updatePbc(atom, param);
printf("10.5\r\n");
//sortAtom(atom); //sortAtom(atom);
buildNeighbor(atom, neighbor); buildNeighbor(atom, neighbor);
printf("10.6\r\n");
LIKWID_MARKER_STOP("reneighbour"); LIKWID_MARKER_STOP("reneighbour");
E = getTimeStamp(); E = getTimeStamp();
@ -279,6 +286,11 @@ int main(int argc, char** argv)
initialIntegrate(&param, &atom); initialIntegrate(&param, &atom);
const bool doReneighbour = (n + 1) % param.every == 0; const bool doReneighbour = (n + 1) % param.every == 0;
const bool doesReneighbourNextRound = (n + 2) % param.every == 0;
printf("Run %d does reneighbour: %d\r\n", n, doReneighbour);
printf("10\r\n");
if(doReneighbour) { if(doReneighbour) {
timer[NEIGH] += reneighbour(&param, &atom, &neighbor); timer[NEIGH] += reneighbour(&param, &atom, &neighbor);
@ -286,6 +298,8 @@ int main(int argc, char** argv)
updatePbc(&atom, &param); updatePbc(&atom, &param);
} }
printf("11\r\n");
if(param.force_field == FF_EAM) { if(param.force_field == FF_EAM) {
timer[FORCE] += computeForceEam(&eam, &param, &atom, &neighbor, &stats, 0, n + 1); timer[FORCE] += computeForceEam(&eam, &param, &atom, &neighbor, &stats, 0, n + 1);
} else { } else {

23
src/neighbor.c
View File

@ -172,6 +172,8 @@ void buildNeighbor(Atom *atom, Neighbor *neighbor)
{ {
int nall = atom->Nlocal + atom->Nghost; int nall = atom->Nlocal + atom->Nghost;
printf("nall: %d, nmax: %d\r\n", nall, nmax);
/* extend atom arrays if necessary */ /* extend atom arrays if necessary */
if(nall > nmax) { if(nall > nmax) {
nmax = nall; nmax = nall;
@ -183,10 +185,14 @@ void buildNeighbor(Atom *atom, Neighbor *neighbor)
// neighbor->neighbors = (int*) malloc(nmax * neighbor->maxneighs * sizeof(int*)); // neighbor->neighbors = (int*) malloc(nmax * neighbor->maxneighs * sizeof(int*));
} }
printf("10.5.1\r\n");
/* bin local & ghost atoms */ /* bin local & ghost atoms */
binatoms(atom); binatoms(atom);
int resize = 1; int resize = 1;
printf("10.5.2\r\n");
/* loop over each atom, storing neighbors */ /* loop over each atom, storing neighbors */
while(resize) { while(resize) {
int new_maxneighs = neighbor->maxneighs; int new_maxneighs = neighbor->maxneighs;
@ -224,6 +230,7 @@ void buildNeighbor(Atom *atom, Neighbor *neighbor)
#else #else
const MD_FLOAT cutoff = cutneighsq; const MD_FLOAT cutoff = cutneighsq;
#endif #endif
if( rsq <= cutoff ) { if( rsq <= cutoff ) {
neighptr[n++] = j; neighptr[n++] = j;
} }
@ -315,9 +322,14 @@ int coord2bin(MD_FLOAT xin, MD_FLOAT yin, MD_FLOAT zin)
void binatoms(Atom *atom) void binatoms(Atom *atom)
{ {
printf("10.5.1.1\r\n");
int nall = atom->Nlocal + atom->Nghost; int nall = atom->Nlocal + atom->Nghost;
int resize = 1; int resize = 1;
printf("10.5.1.2\r\n");
printf("nall: %d, atom->Nmax: %d\r\n", nall, atom->Nmax);
while(resize > 0) { while(resize > 0) {
resize = 0; resize = 0;
@ -325,8 +337,13 @@ void binatoms(Atom *atom)
bincount[i] = 0; bincount[i] = 0;
} }
printf("10.5.1.3\r\n");
for(int i = 0; i < nall; i++) { for(int i = 0; i < nall; i++) {
int ibin = coord2bin(atom_x(i), atom_y(i), atom_z(i)); MD_FLOAT x = atom_x(i);
MD_FLOAT y = atom_y(i);
MD_FLOAT z = atom_z(i);
int ibin = coord2bin(x, y, z);
if(bincount[ibin] < atoms_per_bin) { if(bincount[ibin] < atoms_per_bin) {
int ac = bincount[ibin]++; int ac = bincount[ibin]++;
@ -336,11 +353,15 @@ void binatoms(Atom *atom)
} }
} }
printf("10.5.1.4\r\n");
if(resize) { if(resize) {
free(bins); free(bins);
atoms_per_bin *= 2; atoms_per_bin *= 2;
bins = (int*) malloc(mbins * atoms_per_bin * sizeof(int)); bins = (int*) malloc(mbins * atoms_per_bin * sizeof(int));
} }
printf("10.5.1.5\r\n");
} }
} }