Fix the inference operation of the deformable convolutional layer.

This commit removes the malloc operation in the inference
method and adds the sigmoid kernel.

Signed-off-by: Davide Sapienza <sapienza.dav@gmail.com>

include/tkDNN/Layer.h

@@ -219,7 +219,12 @@ public:
     int kernelH, kernelW, strideH, strideW, paddingH, paddingW;
 protected:
 
+    dnnType *ones_d1;
+    dnnType *ones_d2;
     cudnnTensorDescriptor_t biasTensorDesc;
+    int chunk_dim;
+    dnnType *offset, *mask;
+    dnnType *output_conv;
 
     void initCUDNN();
 

include/tkDNN/kernels.h

@@ -6,6 +6,7 @@
 void activationELUForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
 void activationLEAKYForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
 void activationLOGISTICForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
+void activationSIGMOIDForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream = cudaStream_t(0));
 
 void fill(dnnType* data, int size, dnnType val, cudaStream_t stream = cudaStream_t(0));
 

src/DeformConv2d.cpp

@@ -14,9 +14,36 @@ void DeformConv2d::initCUDNN() {
                                            net->tensorFormat, net->dataType,
                                            1, output_dim.c, 1, 1) );
 
-        checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc,
+    checkCUDNN( cudnnSetTensor4dDescriptor(dstTensorDesc,
                                     net->tensorFormat, net->dataType, output_dim.n, output_dim.c, output_dim.h, output_dim.w));
 
+    const int height_ones = (preconv->input_dim.h + 2 * this->paddingH - (1 * (this->kernelH - 1) + 1)) / this->strideH + 1;
+    const int width_ones = (preconv->input_dim.w + 2 * this->paddingW - (1 * (this->kernelW - 1) + 1)) / this->strideW + 1;
+    const int dim_ones = preconv->input_dim.c * this->kernelH * this->kernelW * 1 * height_ones * width_ones;
+
+    int dst_dim = preconv->output_dim.tot();
+    if (dst_dim % 3 != 0 )
+        std::cout<<"take attention\n\n";
+    chunk_dim = dst_dim/3;
+    checkCuda(cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType)));
+    checkCuda(cudaMalloc(&mask, chunk_dim*sizeof(dnnType)));
+    
+    // kernel ones
+    
+    cudaMallocHost(&ones_d1, (height_ones*width_ones)*sizeof(dnnType));
+    float aus1[height_ones*width_ones];
+    for(int i=0; i<height_ones*width_ones; i++)
+        aus1[i]=1.0f;
+    cudaMemcpy(ones_d1, aus1, (height_ones*width_ones)*sizeof(dnnType), cudaMemcpyHostToDevice);
+    cudaDeviceSynchronize();
+    
+    cudaMallocHost(&ones_d2, dim_ones*sizeof(dnnType));
+    float aus2[dim_ones];
+    for(int i=0; i<dim_ones; i++)
+        aus2[i]=1.0f;
+    cudaMemcpy(ones_d2, aus2, (dim_ones)*sizeof(dnnType), cudaMemcpyHostToDevice);
+    cudaDeviceSynchronize();
+
 }
 
 DeformConv2d::DeformConv2d( Network *net, int out_ch, int deformable_group, int kernelH, int kernelW,
@@ -51,93 +78,25 @@ DeformConv2d::~DeformConv2d() {
 
     checkCUDNN( cudnnDestroyTensorDescriptor(biasTensorDesc) );
     checkCuda( cudaFree(dstData) );
-}
+    checkCuda( cudaFreeHost(ones_d1) );
-
+    checkCuda( cudaFreeHost(ones_d2) );
-void Conv2dToChunk(int dim, dnnType* srcData, dnnType* offset, dnnType* mask)
+    checkCuda( cudaFree(offset) );
-{
+    checkCuda( cudaFree(mask) );
-    // std::cout<<"9\n";
+    checkCuda( cudaFree(output_conv) );
-    // cudaMemcpyFromArray(offset, (const struct cudaArray *)srcData, 0, 2*dim.tot()/3, dim.tot()/3, cudaMemcpyDeviceToHost); 
-    // checkCuda(cudaMemcpyFromArray(offset, (const struct cudaArray *)srcData, 0, 0, 2*dim, cudaMemcpyDeviceToDevice)); 
-    checkCuda(cudaMemcpy(offset, srcData, 2*dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); 
-    
-    // std::cout<<"9a\n";
-    cudaDeviceSynchronize();
-    // cudaMemcpyFromArray(mask, (const struct cudaArray *)srcData, 2*dim.tot()/3, dim.tot(), dim.tot()/3, cudaMemcpyDeviceToHost); 
-    // checkCuda(cudaMemcpyFromArray(mask, (const struct cudaArray *)srcData, 0, 2*dim, dim, cudaMemcpyDeviceToDevice)); 
-    checkCuda(cudaMemcpy(mask, srcData + 2*dim, dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); 
-    // std::cout<<"9b\n";
 }
 
 dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) {
-    dnnType *input;
-    checkCuda(cudaMalloc(&input, dim.tot()*sizeof(dnnType)));
-    checkCuda(cudaMemcpy(input, srcData, dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToDevice));
-    cudaDeviceSynchronize();
-    srcData = preconv->infer(dim, srcData);
-    dim = preconv->output_dim;
-
-    //split to chank
-    dnnType *offset, *mask;
-    int dst_dim = dim.tot();
-    if (dst_dim % 3 != 0 )
-        std::cout<<"take attention\n\n";
-    int chunk_dim = dst_dim/3;
-    checkCuda(cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType)));
-    checkCuda(cudaMalloc(&mask, chunk_dim*sizeof(dnnType)));
-    cudaDeviceSynchronize();
-    
-    Conv2dToChunk(chunk_dim, srcData, offset, mask);
 
+    // conv2d 
+    output_conv = preconv->infer(dim, srcData);
+    // split conv2d outputs into offset to mask
+    checkCuda(cudaMemcpy(offset, output_conv, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); 
+    checkCuda(cudaMemcpy(mask, output_conv + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); 
     // kernel sigmoide
-    dnnType *vec;
+    activationSIGMOIDForward(mask, mask, chunk_dim);
-    vec = new dnnType[chunk_dim];
-    cudaDeviceSynchronize();
-    cudaMemcpy(vec, mask, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToHost);
-    cudaDeviceSynchronize();
-    for(int i=0; i<chunk_dim; i++){
-        // std::cout<<i<<" -- "<<vec[i]<<", ";
-        vec[i] = 1.0f / (1.0f + exp(-vec[i]));
-        // std::cout<<vec[i]<<std::endl;
-    }
-       
-    cudaDeviceSynchronize();
-    cudaMemcpy(mask, vec, chunk_dim*sizeof(dnnType), cudaMemcpyHostToDevice);
-    cudaDeviceSynchronize();
-    free(vec);
-
-    // dnnType *tmp;
-    // cudaMallocHost(&tmp, chunk_dim*sizeof(dnnType));
-    // cudaMemcpy(tmp, mask, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToHost);
-    // std::cout<<"conv2d output before chunking"<<std::endl;
-    // for (size_t i = 0; i < chunk_dim; i++)
-    // {
-    //     std::cout<<i<<" -- "<<tmp[i]<<", ";
-    // }
-    // std::cout<<std::endl;
-    // cudaFreeHost(tmp);
  
-
+    // deformable convolution
-    const int height_ones = (preconv->input_dim.h + 2 * this->paddingH - (1 * (this->kernelH - 1) + 1)) / this->strideH + 1;
+    dcn_v2_cuda_forward(srcData, this->data_d,
-    const int width_ones = (preconv->input_dim.w + 2 * this->paddingW - (1 * (this->kernelW - 1) + 1)) / this->strideW + 1;
-    const int dim_ones = preconv->input_dim.c * this->kernelH * this->kernelW * 1 * height_ones * width_ones;
-    
-    // kernel ones
-    dnnType *ones_d1;
-    cudaMallocHost(&ones_d1, (height_ones*width_ones)*sizeof(dnnType));
-    float aus1[height_ones*width_ones];
-    for(int i=0; i<height_ones*width_ones; i++)
-        aus1[i]=1.0f;
-    cudaMemcpy(ones_d1, aus1, (height_ones*width_ones)*sizeof(dnnType), cudaMemcpyHostToDevice);
-    cudaDeviceSynchronize();
-    dnnType *ones_d2;
-    cudaMallocHost(&ones_d2, dim_ones*sizeof(dnnType));
-    float aus2[dim_ones];
-    for(int i=0; i<dim_ones; i++)
-        aus2[i]=1.0f;
-    cudaMemcpy(ones_d2, aus2, (dim_ones)*sizeof(dnnType), cudaMemcpyHostToDevice);
-    cudaDeviceSynchronize();
-
-    dcn_v2_cuda_forward(input, this->data_d,
                          this->bias2_d, ones_d1,
                          offset, mask,
                          dstData, ones_d2,
@@ -148,44 +107,18 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) {
                          this->deformableGroup, 
                          preconv->input_dim.n, preconv->input_dim.c, preconv->input_dim.h, preconv->input_dim.w,
                          this->output_dim.n, this->output_dim.c, this->output_dim.h, this->output_dim.w,
-                         dst_dim);
+                         chunk_dim);
-
-    cudaFree(offset);
-    cudaFree(mask);
-    cudaFree(input);
-    cudaFreeHost(ones_d1);
-    cudaFreeHost(ones_d2);
-    
-
-    // dnnType *aus3;
-    // cudaMallocHost(&aus3, 256*7*7*sizeof(dnnType));
-    // cudaMemcpy(aus3, dstData, (256*7*7)*sizeof(dnnType), cudaMemcpyDeviceToHost);
-    // checkCuda(cudaDeviceSynchronize());
-    // std::cout<<"OutDim:\n";
-    // this->output_dim.print();
-    // std::cout<<"\n\n\nprint dstData: \n";
-    // for (int i = 0 ; i < 256*7*7; i++){
-    //     if(i==294)
-    //         std::cout<<"\n\n\n";
-    //     std::cout<<aus3[i]<<" ";
-    // }
-    // std::cout<<"\n";
-    // cudaFreeHost(aus3);
-
-    std::cout<<"srcData BN:\n";
-    printDeviceVector(64, dstData);
 
     dnnType alpha = dnnType(1);
     dnnType beta  = dnnType(0);
     if(!batchnorm) {
-        // // // bias
+        // bias
         alpha = dnnType(1);
         beta  = dnnType(1);
         checkCUDNN( cudnnAddTensor(net->cudnnHandle,
                                    &alpha, biasTensorDesc, bias_d,
                                    &beta, dstTensorDesc, dstData) );
     } else {
-        std::cout<<"LOL\n";
         alpha = dnnType(1);
         beta  = dnnType(0);
         checkCUDNN( cudnnBatchNormalizationForwardInference(net->cudnnHandle,
@@ -195,9 +128,8 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) {
                                                 scales_d, bias_d, mean_d, variance_d,
                                                 CUDNN_BN_MIN_EPSILON) );
     }
+
     //update data dimensions
-    std::cout<<"dstData BN:\n";
-    printDeviceVector(64, dstData);
     dim = output_dim;
     return dstData;
 }

src/kernels/activation_sigmoid.cu

@@ -0,0 +1,31 @@
+#include "kernels.h"
+
+__device__ 
+__forceinline__ 
+double sigmoid (double a)
+{
+    return 1.0 / (1.0 + exp (-a));
+}
+
+
+__global__
+void activation_sigmoid(dnnType *input, dnnType *output, int size) {
+
+    int stride = gridDim.x * blockDim.x;
+    int tid = blockDim.x * blockIdx.x + threadIdx.x;
+    for (int i = tid; i < size; i += stride) {
+        output[i] = sigmoid (input[i]);
+    }
+ }
+
+
+/**
+    ELU activation function
+*/
+void activationSIGMOIDForward(dnnType* srcData, dnnType* dstData, int size, cudaStream_t stream)
+{
+    int blocks = (size+255)/256;
+    int threads = 256;
+    
+    activation_sigmoid<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
+}

src/kernels/deformable_conv.cu

@@ -149,10 +149,8 @@ void dcn_v2_cuda_forward(float *input, float *weight,
                          const int deformable_group,
                          const int in_n, const int in_c, const int in_h, const int in_w, 
                          const int out_n, const int out_c, const int out_h, const int out_w,
-                         const int dst_dim, cudaStream_t stream)
+                         const int chunk_dim, cudaStream_t stream)
-{
+{  
-  checkCuda(cudaDeviceSynchronize());
-  cudaError_t cudaStat;    
   cublasStatus_t stat;
   cublasHandle_t handle;
   stat = cublasCreate(&handle);
@@ -160,83 +158,50 @@ void dcn_v2_cuda_forward(float *input, float *weight,
       printf ("CUBLAS initialization failed\n");
       return;
   }
-  checkCuda(cudaDeviceSynchronize());
+
-  
-  const int batch = in_n;
   const int channels = in_c;
   const int height = in_h;
   const int width = in_w;
 
+
   const int channels_out = out_c;
-  const int channels_kernel = in_c;
-  const int kernel_h_ = kernel_h;
-  const int kernel_w_ = kernel_w;
 
   const int height_out = (height + 2 * pad_h - (dilation_h * (kernel_h - 1) + 1)) / stride_h + 1;
   const int width_out = (width + 2 * pad_w - (dilation_w * (kernel_w - 1) + 1)) / stride_w + 1;
-
-  float *input_n;
-  cudaMalloc(&input_n, (in_n*in_c*in_h*in_w)*sizeof(float));
-  cudaMemcpy(input_n, input, (in_n*in_c*in_h*in_w)*sizeof(float), cudaMemcpyDeviceToDevice);
-  checkCuda(cudaDeviceSynchronize());
-
-  float *offset_n;
-  cudaMalloc(&offset_n, ((dst_dim/3)*2)*sizeof(float));
-  cudaMemcpy(offset_n, offset, ((dst_dim/3)*2)*sizeof(float), cudaMemcpyDeviceToDevice);
-  checkCuda(cudaDeviceSynchronize());
-
-  float *mask_n;
-  cudaMalloc(&mask_n, (dst_dim/3)*sizeof(float));
-  cudaMemcpy(mask_n, mask, (dst_dim/3)*sizeof(float), cudaMemcpyDeviceToDevice);
-  checkCuda(cudaDeviceSynchronize());
-
-  float *output_n;
-  checkCuda(cudaMalloc(&output_n, (channels_out*height_out*width_out)*sizeof(float)));
-  checkCuda(cudaDeviceSynchronize());
   
-  long m_ = channels_out;
+  long m = channels_out;
-  long n_ = height_out * width_out;
+  long n = height_out * width_out;
-  long k_ = 1;
+  long k = 1;
   float alpha = 1.0;
   float beta = 0.0;
-  checkCuda(cudaDeviceSynchronize());
+  
   stat = cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, 
-              n_, m_, k_, &alpha, 
+              n, m, k, &alpha, 
-              ones, k_, bias, k_, 
+              ones, k, bias, k, 
-              &beta, output_n, n_);
+              &beta, output, n);
   if (stat != CUBLAS_STATUS_SUCCESS) {
       printf ("CUBLAS initialization failed\n");
       return ;
   }
 
-  checkCuda(cudaDeviceSynchronize());
-
   modulated_deformable_im2col_cuda(stream,
-                                    input_n, offset_n,
+                                    input, offset,
-                                    mask_n,
+                                    mask,
                                     1, channels, height, width,
                                     height_out, width_out, kernel_h, kernel_w,
                                     pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w,
                                     deformable_group, columns);
-  checkCuda(cudaDeviceSynchronize());
+  
-
   //(k * m)  x  (m * n)
   // Y = WC
-  long m = channels_out;
+  k = channels * kernel_h * kernel_w;
-  long n = height_out * width_out;
-  long k = channels * kernel_h * kernel_w;
-
-  alpha = 1.0;
   beta = 1.0;
   
   stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, 
               n, m, k, &alpha, 
               columns, n, weight, k, 
-              &beta, output_n, n);
+              &beta, output, n);
 
-  cudaMemcpy(output, output_n, (n*m)*sizeof(float), cudaMemcpyDeviceToDevice);  
-  checkCuda(cudaDeviceSynchronize());
-            
   if (stat != CUBLAS_STATUS_SUCCESS) {
       printf ("CUBLAS initialization failed\n");
       return ;