diff --git a/src/ComputeForce.cu b/src/ComputeForce.cu index d636626c70f25ab3cc4872507428d07cc862432e..dec829c8f63e5d0f3c35a3da168746e0664259ae 100644 --- a/src/ComputeForce.cu +++ b/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; diff --git a/src/GPUManager.h b/src/GPUManager.h index 624791c42a54c6f85a4e3cbcf01f40cbe16c685e..1d3335242d431e4b751e6264ee87599f4df809b5 100644 --- a/src/GPUManager.h +++ b/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 diff --git a/src/TabulatedPotential.h b/src/TabulatedPotential.h index 7658702b2b96d1595ea7292972aded9ae7af1ca5..4a0f0fcff48a94d9ca02ccc78b8afa0ab5202cb5 100644 --- a/src/TabulatedPotential.h +++ b/src/TabulatedPotential.h @@ -122,7 +122,6 @@ public: gpuErrchk(cudaFree(dev_ptr)); gpuErrchk(cudaFree(tmp.v0)); tmp.v0 = NULL; - dev_ptr->v0 = NULL; }