Fix bugs to switch on O3 flag for GPU code

XiaohuaZhangLLNL · XiaohuaZhangLLNL · commit 8a52dba4def2 · 2021-02-17T17:03:20.000-08:00
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -95,7 +95,6 @@ if(USE_GPU)
     else()
         target_compile_options(ddcMD-gpu PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-gencode arch=compute_70,code=sm_70 -gencode=arch=compute_60,code=sm_60>)
     endif()
-    # GPU code is debug only -TODO
     target_compile_options(ddcMD-gpu PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-g>)
     #link_directories("/usr/tce/packages/cuda/cuda-10.1.243/lib64")
     #target_compile_options(ddcMD-gpu PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-gencode arch=compute_70,code=sm_70>)
diff --git a/arch/sierra.mk b/arch/sierra.mk
@@ -67,7 +67,7 @@ FFLAGS_DEBUG = -g -DLINUX_DEBUG -pg -cpp -w
 FFLAGS_PROF = -O3 -g
 FFLAGS_OMP = $(FFLAGS_OPT)
 
-NVCCFLAGS_OPT  = $(NVCCFLAGS_BASE)  -g 
+NVCCFLAGS_OPT  = $(NVCCFLAGS_BASE)  -O3 
 NVCCFLAGS_DEBUG  = $(NVCCFLAGS_BASE)  -g -G 
 NVCCFLAGS_PROF  = $(NVCCFLAGS_BASE)   -g -pg 
 
diff --git a/arch/summit.mk b/arch/summit.mk
@@ -69,7 +69,7 @@ FFLAGS_DEBUG = -g -DLINUX_DEBUG -pg -cpp -w
 FFLAGS_PROF = -O3 -g
 FFLAGS_OMP = $(FFLAGS_OPT)
 
-NVCCFLAGS_OPT  = $(NVCCFLAGS_BASE)  -g 
+NVCCFLAGS_OPT  = $(NVCCFLAGS_BASE)  -O3 
 NVCCFLAGS_DEBUG  = $(NVCCFLAGS_BASE)  -g -G 
 NVCCFLAGS_PROF  = $(NVCCFLAGS_BASE)   -g -pg 
 
diff --git a/arch/toss3.mk b/arch/toss3.mk
@@ -50,7 +50,7 @@ ifeq ($(USE_GPU), 1)
 
   NVCC = $(CUDAHOME)/bin/nvcc -std=c++11
 
-  NVCCFLAGS_OPT  = $(NVCCFLAGS_BASE)  -g
+  NVCCFLAGS_OPT  = $(NVCCFLAGS_BASE)  -O3
   NVCCFLAGS_DEBUG  = $(NVCCFLAGS_BASE)  -g -G
   NVCCFLAGS_PROF  = $(NVCCFLAGS_BASE)   -g -pg
 
diff --git a/src/gpuMemUtils.cu b/src/gpuMemUtils.cu
@@ -159,6 +159,7 @@ void allocGPUState(COLLECTION *collection, int nIon)
     gnlist->binCountsg = NULL;
     gnlist->binCountsg2= NULL;
     gnlist->binHeadsg  = NULL;
+    gnlist->binHeadsgSave  = gnlist->binHeadsg;
     gnlist->partial_sumsg = NULL;
     nvtxRangePop();
 
@@ -208,15 +209,16 @@ void allocGPUnbins(GPUNLIST *gnlist, const int nBinsTot, const int blocksize)
         if(gnlist->binCountsg2 != NULL){
             CUDA_SAFE_CALL(cudaFree(gnlist->binCountsg2);)
         }
-        if(gnlist->binHeadsg != NULL){
-            CUDA_SAFE_CALL(cudaFree(gnlist->binHeadsg);)
+        if(gnlist->binHeadsgSave != NULL){
+            CUDA_SAFE_CALL(cudaFree(gnlist->binHeadsgSave);)
         }
         
         // allocate 2X memory for bins/partial sums
         int nBinsTot2=nBinsTot*2;
         gpu_allocator(gnlist->binCountsg, nBinsTot2);
         gpu_allocator(gnlist->binCountsg2, nBinsTot2);
         gpu_allocator(gnlist->binHeadsg, nBinsTot2);
+        gnlist->binHeadsgSave = gnlist->binHeadsg;
 
         accParms->totbins=nBinsTot2;       
 
@@ -314,6 +316,7 @@ void deallocGPU(SYSTEM *sys, int n)
     CUDA_SAFE_CALL(cudaFree((gnlist->binCountsg)));
     cudaFree((gnlist->binCountsg2));
     cudaFree((gnlist->binHeadsg));
+    gnlist->binHeadsg=NULL;
     cudaFree((gnlist->listIdsg));
     CUDA_SAFE_CALL(cudaFree((gnlist->partial_sumsg)));
 
diff --git a/src/gpu_allocator.hpp b/src/gpu_allocator.hpp
@@ -4,7 +4,49 @@
 #include <cuda_runtime_api.h>
 #include <cuda.h>
 
+#include <execinfo.h>
+#include <string.h>
+#include <errno.h>
+#include <unistd.h>
+#include <stdlib.h>
+
+static void full_write(int fd, const char *buf, size_t len)
+{
+        while (len > 0) {
+                ssize_t ret = write(fd, buf, len);
+
+                if ((ret == -1) && (errno != EINTR))
+                        break;
+
+                buf += (size_t) ret;
+                len -= (size_t) ret;
+        }
+}
+
+inline void print_backtrace(void)
+{
+        static const char start[] = "BACKTRACE ------------\n";
+        static const char end[] = "----------------------\n";
+
+        void *bt[1024];
+        int bt_size;
+        char **bt_syms;
+        int i;
+
+        bt_size = backtrace(bt, 1024);
+        bt_syms = backtrace_symbols(bt, bt_size);
+        full_write(STDERR_FILENO, start, strlen(start));
+        for (i = 1; i < bt_size; i++) {
+                size_t len = strlen(bt_syms[i]);
+                full_write(STDERR_FILENO, bt_syms[i], len);
+                full_write(STDERR_FILENO, "\n", 1);
+        }
+        full_write(STDERR_FILENO, end, strlen(end));
+    free(bt_syms);
+}
+
 template<typename T> void gpu_allocator(T &ptrref, long long int n) {
+  //print_backtrace();
   CUDA_SAFE_CALL( cudaMalloc((void **) &ptrref, sizeof(*ptrref) * n);)
 }
 
diff --git a/src/gpunlist.h b/src/gpunlist.h
@@ -25,7 +25,7 @@ typedef struct gpunlist_st {
     double *scratch, *scratch1, *scratch2, *scratch3, *scratch4; //scratch space
     double *partialCOMx, *partialCOMy, *partialCOMz;
     double *virCorx, *virCory, *virCorz;
-    int *r_backg, *r_backbg, *binHeadsg, *binCountsg, *binCountsg2, *listIdsg, *partial_sumsg;
+    int *r_backg, *r_backbg, *binHeadsg, *binHeadsgSave, *binCountsg, *binCountsg2, *listIdsg, *partial_sumsg;
     double *minsg, *lensg;
     int *nbinsg, *nbrIds, *nbrIdsx, *nbrIdsy, *nbrIdsz;
     int numNbrs, nBinsTot;
diff --git a/src/pairProcessGPU.cu b/src/pairProcessGPU.cu
@@ -3,6 +3,7 @@
 #include <thrust/execution_policy.h>
 #include <thrust/pair.h>
 #include <nvToolsExt.h>
+#include "ddcMalloc.h"
 #include "pairProcessGPU.h"
 #include "nlistGPU.h"
 #include "cudaUtils.h"
@@ -32,7 +33,7 @@
 #define TPP_EVAL 1
 #define TPP_EVALp 1
 #define  CHECK_1(x)     //first layer of Debugging, summary of binning data
-#define CHECK_2(x)   //second layer of debugging, print all binning data
+#define CHECK_2(x)    //second layer of debugging, print all binning data
 #define CHECK_3(x)   //not used rn
 #define CHECK(x)   //final layer of debugging, print all particle data
 #define CHECKB (x) 
@@ -1153,7 +1154,7 @@ void binParticlesGPU(SYSTEM * sys, double rcut)
     double *minsg, *lensg;
     minsg = gnlist->minsg;
     lensg = gnlist->lensg;
-
+  
     //determine number of bins necessary based on cutoff
     CHECK_1(
             printf("box len %f %f %f \n", lens[0], lens[1], lens[2]);
@@ -1166,7 +1167,6 @@ void binParticlesGPU(SYSTEM * sys, double rcut)
     //binL[0] = lens[0] / nbins[0];
     //binL[1] = lens[1] / nbins[1];
     //binL[2] = lens[2] / nbins[2];
-
     int nBinsTot = nbins[0] * nbins[1] * nbins[2];
     gnlist->nBinsTot = nBinsTot;
     //int tBins = nbins[0]*nbins[1]*nbins[2]; 
@@ -1267,44 +1267,63 @@ void binParticlesGPU(SYSTEM * sys, double rcut)
         //27 if 1d
         // 9 if 2d
         // 3 if 1d
-        int* x_nbr;
-    int* y_nbr;
-    int* z_nbr;
+        int* x_nbr=NULL;
+    int* y_nbr=NULL;
+    int* z_nbr=NULL;
     int num_x_nbr = 3;
     int num_y_nbr = 3;
     int num_z_nbr = 3;
     if (nbins[0] > 1)
     {
-        int x_nbr1 [3] = {-1, 0, 1};
-        x_nbr = x_nbr1;
+        //int x_nbr1 [3] = {-1, 0, 1};
+        //x_nbr = x_nbr1;
+        x_nbr=(int *)ddcMalloc(num_x_nbr*sizeof(int));
+        x_nbr[0]=-1;
+        x_nbr[1]=0;
+        x_nbr[2]=1;
     }
     else
     {
-        int x_nbr1 [1] = {0};
-        x_nbr = x_nbr1;
+        //int x_nbr1 [1] = {0};
+        //x_nbr = x_nbr1;
         num_x_nbr = 1;
+        x_nbr=(int *)ddcMalloc(num_x_nbr*sizeof(int));
+        x_nbr[0]=0;
     }
     if (nbins[1] > 1)
     {
-        int y_nbr1 [3] = {-1, 0, 1}; //{-nbins[0],0,nbins[0]};
-        y_nbr = y_nbr1;
+        //int y_nbr1 [3] = {-1, 0, 1}; //{-nbins[0],0,nbins[0]};
+        //y_nbr = y_nbr1;
+        y_nbr=(int *)ddcMalloc(num_y_nbr*sizeof(int));
+        y_nbr[0]=-1;
+        y_nbr[1]=0;
+        y_nbr[2]=1;
+
     }
     else
     {
-        int y_nbr1 [1] = {0};
-        y_nbr = y_nbr1;
+        //int y_nbr1 [1] = {0};
+        //y_nbr = y_nbr1;
         num_y_nbr = 1;
+        y_nbr=(int *)ddcMalloc(num_y_nbr*sizeof(int));
+        y_nbr[0]=0;
     }
     if (nbins[2] > 1)
     {
-        int z_nbr1 [3] = {-1, 0, 1}; //{-nbins[0]*nbins[1],0,nbins[0]*nbins[1]};
-        z_nbr = z_nbr1;
+        //int z_nbr1 [3] = {-1, 0, 1}; //{-nbins[0]*nbins[1],0,nbins[0]*nbins[1]};
+        //z_nbr = z_nbr1;
+        z_nbr=(int *)ddcMalloc(num_z_nbr*sizeof(int));
+        z_nbr[0]=-1;
+        z_nbr[1]=0;
+        z_nbr[2]=1;
     }
     else
     {
-        int z_nbr1 [1] = {0};
-        z_nbr = z_nbr1;
+        //int z_nbr1 [1] = {0};
+        //z_nbr = z_nbr1;
         num_z_nbr = 1;
+        z_nbr=(int *)ddcMalloc(num_z_nbr*sizeof(int));
+        z_nbr[0]=0;
     }
 
     CHECK_1(
@@ -1351,6 +1370,12 @@ void binParticlesGPU(SYSTEM * sys, double rcut)
     gpu_memcpy_host2device(gnlist->nbrIdsy, y_nbr, num_y_nbr);
     gpu_memcpy_host2device(gnlist->nbrIdsz, z_nbr, num_z_nbr);
     gpu_memcpy_host2device(gnlist->numNbrsxyz, num_nbrsxyz, 3);
+
+
+    // Free the memory of pointers
+    ddcFree(x_nbr);
+    ddcFree(y_nbr);
+    ddcFree(z_nbr);
 }
 
 //
