Also pinnend neighbor-struct memory, added additional performance measurements, added nvprof result to logbook

This commit is contained in:
Maximilian Gaul 2021-12-18 15:58:56 +01:00
parent c2bfa3ca3f
commit 134e3f4b78
2 changed files with 9 additions and 5 deletions

View File

@ -1,5 +1,6 @@
END=32 END=32
for ((i=1;i<=END;i++)); do for ((i=1;i<=END;i++)); do
output=$(eval "NUM_THREADS=$i ./MDBench-NVCC -n 50") output=$(eval "NUM_THREADS=$i ./MDBench-NVCC -n 50")
echo "$output" | grep 'atom updates per second' | sed 's/[^0-9.]//g' | awk '{print $1"e6"}' echo -n "$i,"
echo "$output" | grep 'atom updates per second' | sed 's/[^0-9.,]//g' | awk '{print $1"e6"}'
done done

View File

@ -26,6 +26,7 @@
#include <neighbor.h> #include <neighbor.h>
#include <parameter.h> #include <parameter.h>
#include <allocate.h>
#include <atom.h> #include <atom.h>
#define SMALL 1.0e-6 #define SMALL 1.0e-6
@ -174,10 +175,12 @@ void buildNeighbor(Atom *atom, Neighbor *neighbor)
/* extend atom arrays if necessary */ /* extend atom arrays if necessary */
if(nall > nmax) { if(nall > nmax) {
nmax = nall; nmax = nall;
if(neighbor->numneigh) free(neighbor->numneigh); if(neighbor->numneigh) cudaFreeHost(neighbor->numneigh);
if(neighbor->neighbors) free(neighbor->neighbors); if(neighbor->neighbors) cudaFreeHost(neighbor->neighbors);
neighbor->numneigh = (int*) malloc(nmax * sizeof(int)); checkCUDAError( "buildNeighbor numneigh", cudaMallocHost((void**)&(neighbor->numneigh), nmax * sizeof(int)) );
neighbor->neighbors = (int*) malloc(nmax * neighbor->maxneighs * sizeof(int*)); checkCUDAError( "buildNeighbor neighbors", cudaMallocHost((void**)&(neighbor->neighbors), nmax * neighbor->maxneighs * sizeof(int*)) );
// neighbor->numneigh = (int*) malloc(nmax * sizeof(int));
// neighbor->neighbors = (int*) malloc(nmax * neighbor->maxneighs * sizeof(int*));
} }
/* bin local & ghost atoms */ /* bin local & ghost atoms */