Merge branch 'cnet' of https://github.com/ceccocats/tkDNN into cnet

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;

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,

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<dnnType*>(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;

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,

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) {

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;