From 2503eba1735c510f3f9454b536399762fbd2ec11 Mon Sep 17 00:00:00 2001 From: Davide Sapienza Date: Thu, 6 Feb 2020 23:09:12 +0100 Subject: [PATCH] Fix memory leak This commit moves cublasCreate out from dcn_v2_cuda_forward to save some milliseconds and it adds cublasDestroy (cause of memory leak). Signed-off-by: Davide Sapienza --- include/tkDNN/Layer.h | 3 +++ include/tkDNN/kernels.h | 3 ++- include/tkDNN/pluginsRT/DeformableConvRT.h | 18 ++++++++++++++++-- src/DeformConv2d.cpp | 11 +++++++++-- src/NetworkRT.cpp | 2 +- src/kernels/deformable_conv.cu | 12 +++--------- 6 files changed, 34 insertions(+), 15 deletions(-) diff --git a/include/tkDNN/Layer.h b/include/tkDNN/Layer.h index 7716393..f9da061 100644 --- a/include/tkDNN/Layer.h +++ b/include/tkDNN/Layer.h @@ -229,6 +229,9 @@ public: dnnType *offset, *mask; dnnType *output_conv; + cublasStatus_t stat; + cublasHandle_t handle; + protected: cudnnTensorDescriptor_t biasTensorDesc; diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index d7d5d05..94c5dcb 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -38,7 +38,8 @@ void modulated_deformable_im2col_cuda(cudaStream_t stream, const int dilation_h, const int dilation_w, const int deformable_group, float *data_col); -void dcn_v2_cuda_forward(float *input, float *weight, +void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle, + float *input, float *weight, float *bias, float *ones, float *offset, float *mask, float *output, float *columns, diff --git a/include/tkDNN/pluginsRT/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h index dca1020..f4d0b30 100644 --- a/include/tkDNN/pluginsRT/DeformableConvRT.h +++ b/include/tkDNN/pluginsRT/DeformableConvRT.h @@ -52,10 +52,21 @@ public: checkCuda( cudaMemcpy(mask, deformable->mask, sizeof(dnnType)*chunk_dim, cudaMemcpyDeviceToDevice) ); checkCuda( cudaMemcpy(ones_d2, deformable->ones_d2, sizeof(dnnType)*dim_ones, cudaMemcpyDeviceToDevice) ); } + stat = cublasCreate(&handle); + if (stat != CUBLAS_STATUS_SUCCESS) { + printf ("CUBLAS initialization failed\n"); + return; + } } ~DeformableConvRT(){ - + checkCuda( cudaFree(data_d) ); + checkCuda( cudaFree(bias2_d) ); + checkCuda( cudaFree(ones_d1) ); + checkCuda( cudaFree(offset) ); + checkCuda( cudaFree(mask) ); + checkCuda( cudaFree(ones_d2) ); + cublasDestroy(handle); } int getNbOutputs() const override { @@ -100,7 +111,8 @@ public: activationSIGMOIDForward(mask, mask, chunk_dim); // deformable convolution - dcn_v2_cuda_forward(srcData, data_d, + dcn_v2_cuda_forward(stat, handle, + srcData, data_d, bias2_d, ones_d1, offset, mask, reinterpret_cast(outputs[0]), ones_d2, @@ -172,6 +184,8 @@ public: free(aus); } + cublasStatus_t stat; + cublasHandle_t handle; int i_n, i_c, i_h, i_w; int o_n, o_c, o_h, o_w; int size; diff --git a/src/DeformConv2d.cpp b/src/DeformConv2d.cpp index 957ef24..8cf1a6d 100644 --- a/src/DeformConv2d.cpp +++ b/src/DeformConv2d.cpp @@ -9,6 +9,11 @@ namespace tk { namespace dnn { void DeformConv2d::initCUDNN() { + stat = cublasCreate(&handle); + if (stat != CUBLAS_STATUS_SUCCESS) { + printf ("CUBLAS initialization failed\n"); + return; + } checkCUDNN( cudnnCreateTensorDescriptor(&biasTensorDesc) ); checkCUDNN( cudnnSetTensor4dDescriptor(biasTensorDesc, net->tensorFormat, net->dataType, @@ -84,6 +89,7 @@ DeformConv2d::~DeformConv2d() { checkCuda( cudaFree(offset) ); checkCuda( cudaFree(mask) ); checkCuda( cudaFree(output_conv) ); + cublasDestroy(handle); } dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) { @@ -95,9 +101,10 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) { checkCuda(cudaMemcpy(mask, output_conv + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); // kernel sigmoide activationSIGMOIDForward(mask, mask, chunk_dim); - + // deformable convolution - dcn_v2_cuda_forward(srcData, this->data_d, + dcn_v2_cuda_forward(stat, handle, + srcData, this->data_d, this->bias2_d, ones_d1, offset, mask, dstData, ones_d2, diff --git a/src/NetworkRT.cpp b/src/NetworkRT.cpp index 2b7377d..5162d48 100644 --- a/src/NetworkRT.cpp +++ b/src/NetworkRT.cpp @@ -453,7 +453,7 @@ ILayer* NetworkRT::convert_layer(ITensor *input, DeformConv2d *l) { IPluginLayer *lRT = networkRT->addPlugin(inputs, 2, *plugin); checkNULL(lRT); lRT->setName( ("Deformable" + std::to_string(l->id)).c_str() ); - + delete(inputs); // batchnorm void *bias_b, *power_b, *mean_b, *variance_b, *scales_b; if(dtRT == DataType::kHALF) { diff --git a/src/kernels/deformable_conv.cu b/src/kernels/deformable_conv.cu index cb9ded0..0579620 100644 --- a/src/kernels/deformable_conv.cu +++ b/src/kernels/deformable_conv.cu @@ -138,7 +138,8 @@ void modulated_deformable_im2col_cuda(cudaStream_t stream, } -void dcn_v2_cuda_forward(float *input, float *weight, +void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle, + float *input, float *weight, float *bias, float *ones, float *offset, float *mask, float *output, float *columns, @@ -151,14 +152,7 @@ void dcn_v2_cuda_forward(float *input, float *weight, const int out_n, const int out_c, const int out_h, const int out_w, const int chunk_dim, cudaStream_t stream) { - cublasStatus_t stat; - cublasHandle_t handle; - stat = cublasCreate(&handle); - if (stat != CUBLAS_STATUS_SUCCESS) { - printf ("CUBLAS initialization failed\n"); - return; - } - + // stat and handle have be moved out to preserve 2 - 6 milliseconds every 100. const int channels = in_c; const int height = in_h; const int width = in_w;