Changed code slightly in attempt to mitigate impact of main loop when single GPU is used

src/ComputeForce.cu

@@ -252,8 +252,18 @@ ComputeForce::~ComputeForce() {
 	gpuErrchk(cudaFree(tableAlpha_d));
 	gpuErrchk(cudaFree(tableRad6_d));
 	
-	for (int j = 0; j < numParts * numParts; ++j)
-		delete tablePot[j];
+	for (int i = 0; i < numParts; ++i) {
+	    for (int j = i; j < numParts; ++j) {
+		int ind = i+j*numParts;
+		if (tablePot[ind] != NULL) {
+		    for (std::size_t g = 0; g < gpuman.gpus.size(); ++g) {
+			gpuman.use(g);
+			tablePot_addr[g][ind]->free_from_cuda(tablePot_addr[g][ind]);
+		    }
+		    delete tablePot[ind];
+		}
+	    }
+	}
 	delete[] tablePot;
 	for (auto& tpa : tablePot_addr) delete[] tpa;
 

src/GPUManager.h

@@ -6,6 +6,8 @@
 #include <cuda.h>
 #include <cuda_runtime_api.h>
 
+#include "useful.h"
+
 // #ifdef USE_NCCL
 #include <nccl.h>
 #define NCCLCHECK(cmd) do {					\
@@ -18,7 +20,12 @@
   } while(0)
 // #endif
 
-#include "useful.h"
+#ifndef gpuErrchk
+#define delgpuErrchk
+#define gpuErrchk(code) { if ((code) != cudaSuccess) {					       \
+	fprintf(stderr,"CUDA Error: %s %s %d\n", cudaGetErrorString(code), __FILE__, __LINE__); \
+    }}
+#endif
 
 #define NUMSTREAMS 8
 
@@ -94,7 +101,18 @@ public:
 	static void use(int gpu_id);
 
 	static void sync(int gpu_id);
-	static void sync();
+	static void sync() {
+	    if (gpus.size() > 1) {
+		int curr;
+		gpuErrchk( cudaGetDevice(&curr) );
+		for (auto it = gpus.begin(); it != gpus.end(); ++it) {
+		    gpuErrchk( cudaSetDevice(it->id) );
+		    gpuErrchk( cudaDeviceSynchronize() );
+		}
+		gpuErrchk( cudaSetDevice(curr) );
+	    } else gpuErrchk( cudaDeviceSynchronize() );
+	}
+
 
 	// current
 	// @return the current GPU a thread is using
@@ -140,5 +158,9 @@ public:
     }
     
 };
+#ifndef delgpuErrchk
+#undef  delgpuErrchk
+#undef  gpuErrchk
+#endif
 
 #endif

src/TabulatedPotential.h

@@ -122,7 +122,6 @@ public:
 	gpuErrchk(cudaFree(dev_ptr));
 	gpuErrchk(cudaFree(tmp.v0));
 	tmp.v0 = NULL;
-	dev_ptr->v0 = NULL;
     }