From 5be71cf9b735eeceef4e191a0970a6e5be09041e Mon Sep 17 00:00:00 2001 From: Davide Sapienza Date: Wed, 8 Apr 2020 11:14:34 +0200 Subject: [PATCH] Fix the Deformable convolution code sintax. Signed-off-by: Davide Sapienza --- include/tkDNN/kernels.h | 15 +---- include/tkDNN/pluginsRT/DeformableConvRT.h | 32 ++-------- src/DeformConv2d.cpp | 40 ++++++------- src/kernels/deformable_conv.cu | 69 ++++++++-------------- 4 files changed, 50 insertions(+), 106 deletions(-) diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index 13fda1f..afcc168 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -32,20 +32,7 @@ void upsampleForward(dnnType *srcData, dnnType *dstData, void float2half(float *srcData, __half *dstData, int size, const cudaStream_t stream = cudaStream_t(0)); -// void modulated_deformable_im2col_cuda(cudaStream_t stream, -// const float *data_im, const float *data_offset, const float *data_mask, -// const int batch_size, const int channels, const int height_im, const int width_im, -// const int height_col, const int width_col, const int kernel_h, const int kenerl_w, -// const int pad_h, const int pad_w, const int stride_h, const int stride_w, -// const int dilation_h, const int dilation_w, -// const int deformable_group, float *data_col); -void modulated_deformable_im2col_cuda(cudaStream_t stream, - const float *data_im, const float *data_offset, const float *data_mask, - const int batch_size, const int channels, const int height_im, const int width_im, - const int height_col, const int width_col, - const int deformable_group, float *data_col); - -void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle, +void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, float *input, float *weight, float *bias, float *ones, float *offset, float *mask, diff --git a/include/tkDNN/pluginsRT/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h index f4d0b30..4804c03 100644 --- a/include/tkDNN/pluginsRT/DeformableConvRT.h +++ b/include/tkDNN/pluginsRT/DeformableConvRT.h @@ -12,12 +12,6 @@ public: int o_n, int o_c, int o_h, int o_w, tk::dnn::DeformConv2d *deformable = nullptr) { this->chunk_dim = chunk_dim; - // int dst_dim = conv_dim.tot(); - // std::cout<<"conv_dim: \n"; - // conv_dim.print(); - // if (dst_dim % 3 != 0 ) - // std::cout<<"take attention\n\n"; - // this->chunk_dim = dst_dim/3; this->kh = kh; this->kw = kw; this->sh = sh; @@ -53,13 +47,11 @@ public: 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; - } + if (stat != CUBLAS_STATUS_SUCCESS) + FatalError("CUBLAS initialization failed\n"); } - ~DeformableConvRT(){ + ~DeformableConvRT() { checkCuda( cudaFree(data_d) ); checkCuda( cudaFree(bias2_d) ); checkCuda( cudaFree(ones_d1) ); @@ -77,24 +69,13 @@ public: return DimsCHW{defRT->output_dim.c, defRT->output_dim.h, defRT->output_dim.w}; } - void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { - // i_n = 1; - // i_c = inputDims[0].d[0]; - // i_h = inputDims[0].d[1]; - // i_w = inputDims[0].d[2]; - // o_n = 1; - // o_c = outputDims[0].d[0]; - // o_h = outputDims[0].d[1]; - // o_w = outputDims[0].d[2]; - } + void configure(const Dims* inputDims, int nbInputs, const Dims* outputDims, int nbOutputs, int maxBatchSize) override { } int initialize() override { - return 0; } - virtual void terminate() override { - } + virtual void terminate() override { } virtual size_t getWorkspaceSize(int maxBatchSize) const override { return 0; @@ -111,7 +92,7 @@ public: activationSIGMOIDForward(mask, mask, chunk_dim); // deformable convolution - dcn_v2_cuda_forward(stat, handle, + dcnV2CudaForward(stat, handle, srcData, data_d, bias2_d, ones_d1, offset, mask, @@ -205,6 +186,5 @@ public: dnnType * mask; dnnType *ones_d2; - tk::dnn::DeformConv2d *defRT; }; diff --git a/src/DeformConv2d.cpp b/src/DeformConv2d.cpp index 8cf1a6d..06bba7f 100644 --- a/src/DeformConv2d.cpp +++ b/src/DeformConv2d.cpp @@ -10,10 +10,9 @@ namespace tk { namespace dnn { void DeformConv2d::initCUDNN() { stat = cublasCreate(&handle); - if (stat != CUBLAS_STATUS_SUCCESS) { - printf ("CUBLAS initialization failed\n"); - return; - } + if (stat != CUBLAS_STATUS_SUCCESS) + FatalError("CUBLAS initialization failed\n"); + checkCUDNN( cudnnCreateTensorDescriptor(&biasTensorDesc) ); checkCUDNN( cudnnSetTensor4dDescriptor(biasTensorDesc, net->tensorFormat, net->dataType, @@ -27,28 +26,27 @@ void DeformConv2d::initCUDNN() { 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"; + if( dst_dim % 3 != 0 ) + FatalError("DeformConv2d: the Conv2d output is not divisible by three"); chunk_dim = dst_dim/3; checkCuda( cudaMalloc(&offset, 2*chunk_dim*sizeof(dnnType))); checkCuda( cudaMalloc(&mask, chunk_dim*sizeof(dnnType))); // kernel ones - checkCuda( cudaMalloc(&ones_d1, (height_ones*width_ones)*sizeof(dnnType)) ); - dnnType *aus1; - checkCuda( cudaMallocHost(&aus1, (height_ones*width_ones)*sizeof(dnnType)) ); + dnnType *ones_h1; + checkCuda( cudaMallocHost(&ones_h1, (height_ones*width_ones)*sizeof(dnnType)) ); for(int i=0; igetOutputDim().c, out_ch, kernelH, kernelW, 1, - d_fname_weights, batchnorm, true){ - + d_fname_weights, batchnorm, true) { this->out_ch = out_ch; this->deformableGroup = deformable_group; this->kernelH = kernelH; @@ -81,7 +78,6 @@ DeformConv2d::DeformConv2d( Network *net, int out_ch, int deformable_group, int } DeformConv2d::~DeformConv2d() { - checkCUDNN( cudnnDestroyTensorDescriptor(biasTensorDesc) ); checkCuda( cudaFree(dstData) ); checkCuda( cudaFree(ones_d1) ); @@ -96,14 +92,14 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) { // conv2d output_conv = preconv->infer(dim, srcData); - // split conv2d outputs into offset to mask + // split conv2d outputs into offset and 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 activationSIGMOIDForward(mask, mask, chunk_dim); // deformable convolution - dcn_v2_cuda_forward(stat, handle, + dcnV2CudaForward(stat, handle, srcData, this->data_d, this->bias2_d, ones_d1, offset, mask, diff --git a/src/kernels/deformable_conv.cu b/src/kernels/deformable_conv.cu index 672655f..e46c579 100644 --- a/src/kernels/deformable_conv.cu +++ b/src/kernels/deformable_conv.cu @@ -1,6 +1,8 @@ #include #include #include +#include +#include #include "kernels.h" #include @@ -17,8 +19,7 @@ inline int GET_BLOCKS(const int N) __device__ float dmcn_im2col_bilinear(const float *bottom_data, const int data_width, - const int height, const int width, float h, float w) -{ + const int height, const int width, float h, float w) { int h_low = floor(h); int w_low = floor(w); int h_high = h_low + 1; @@ -44,8 +45,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n, const int height, const int width, const int batch_size, const int num_channels, const int deformable_group, const int height_col, const int width_col, - float *data_col) -{ + float *data_col) { CUDA_KERNEL_LOOP(index, n) { //If n is a power of 2, ( i / n ) is equivalent to ( i ≫ log2 n ) and ( i % n ) is equivalent to ( i & n - 1 ). @@ -77,11 +77,9 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n, const float *data_mask_ptr = data_mask + add_ptr; #pragma unroll - for (int i = 0; i < 3; ++i) - { + for (int i = 0; i < 3; ++i) { #pragma unroll - for (int j = 0; j < 3; ++j) - { + for (int j = 0; j < 3; ++j) { const int iter_member = (i * 3 + j); // const int data_offset_h_ptr = ((2 * (i * kernel_w + j)) * height_col + h_col) * width_col + w_col; const int data_offset_h_ptr = first_member + s_col2 * iter_member; @@ -99,8 +97,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n, const float w_im = offset_w + w_in + j; //if (h_im >= 0 && w_im >= 0 && h_im < height && w_im < width) { float val = static_cast(0); - if (h_im < height && w_im < width && h_im > -1 && w_im > -1) - { + if (h_im < height && w_im < width && h_im > -1 && w_im > -1) { //const float map_h = i * dilation_h + offset_h; //const float map_w = j * dilation_w + offset_w; //const int cur_height = height - h_in; @@ -116,7 +113,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n, } } -__global__ void modulated_deformable_im2col_gpu_kernel2(const int n, +__global__ void modulated_deformable_im2col_gpu_kernel_general_version(const int n, const float *data_im, const float *data_offset, const float *data_mask, const int height, const int width, const int kernel_h, const int kernel_w, const int pad_h, const int pad_w, @@ -125,12 +122,10 @@ __global__ void modulated_deformable_im2col_gpu_kernel2(const int n, const int channel_per_deformable_group, const int batch_size, const int num_channels, const int deformable_group, const int height_col, const int width_col, - float *data_col) -{ + float *data_col) { CUDA_KERNEL_LOOP(index, n) { //If n is a power of 2, ( i / n ) is equivalent to ( i ≫ log2 n ) and ( i % n ) is equivalent to ( i & n - 1 ). - // printf("--- %d %d %d %d %d %d %d %d\n",kernel_h, kernel_w, pad_h, pad_w, stride_h, stride_w, dilation_h, dilation_w); const int ind_on_w = index / width_col; const int ind_on_w_on_h = ind_on_w / height_col; const int kk = kernel_h * kernel_w; @@ -160,11 +155,9 @@ __global__ void modulated_deformable_im2col_gpu_kernel2(const int n, const float *data_mask_ptr = data_mask + add_ptr; #pragma unroll - for (int i = 0; i < kernel_h; ++i) - { + for (int i = 0; i < kernel_h; ++i) { #pragma unroll - for (int j = 0; j < kernel_w; ++j) - { + for (int j = 0; j < kernel_w; ++j) { const int iter_member = (i * kernel_w + j); // const int data_offset_h_ptr = ((2 * (i * kernel_w + j)) * height_col + h_col) * width_col + w_col; const int data_offset_h_ptr = first_member + s_col2 * iter_member; @@ -182,8 +175,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel2(const int n, const float w_im = offset_w + w_in + j * dilation_w; //if (h_im >= 0 && w_im >= 0 && h_im < height && w_im < width) { float val = static_cast(0); - if (h_im < height && w_im < width && h_im > -1 && w_im > -1) - { + if (h_im < height && w_im < width && h_im > -1 && w_im > -1) { //const float map_h = i * dilation_h + offset_h; //const float map_w = j * dilation_w + offset_w; //const int cur_height = height - h_in; @@ -199,8 +191,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel2(const int n, } } - -void modulated_deformable_im2col_cuda(cudaStream_t stream, +void modulatedDeformableIm2colCuda(cudaStream_t stream, const float* data_im, const float* data_offset, const float* data_mask, const int batch_size, const int channels, const int height_im, const int width_im, const int height_col, const int width_col, @@ -216,13 +207,10 @@ void modulated_deformable_im2col_cuda(cudaStream_t stream, cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) - { - printf("error in modulated_deformable_im2col_cuda: %s\n", cudaGetErrorString(err)); - } - + FatalError("error in modulatedDeformableIm2colCuda: " + std::string(cudaGetErrorString(err)) + "\n"); } -void modulated_deformable_im2col_cuda2(cudaStream_t stream, +void modulatedDeformableIm2colCudaGeneralVersion(cudaStream_t stream, const float* data_im, const float* data_offset, const float* data_mask, const int batch_size, const int channels, const int height_im, const int width_im, const int height_col, const int width_col, const int kernel_h, const int kenerl_w, @@ -232,7 +220,7 @@ void modulated_deformable_im2col_cuda2(cudaStream_t stream, // num_axes should be smaller than block size const int channel_per_deformable_group = channels / deformable_group; const int num_kernels = channels * batch_size * height_col * width_col; - modulated_deformable_im2col_gpu_kernel2 + modulated_deformable_im2col_gpu_kernel_general_version <<>>( num_kernels, data_im, data_offset, data_mask, height_im, width_im, kernel_h, kenerl_w, @@ -241,13 +229,10 @@ void modulated_deformable_im2col_cuda2(cudaStream_t stream, cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) - { - printf("error in modulated_deformable_im2col_cuda: %s\n", cudaGetErrorString(err)); - } - + FatalError("error in modulatedDeformableIm2colCudaGeneralVersion: " + std::string(cudaGetErrorString(err)) + "\n"); } -void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle, +void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, float *input, float *weight, float *bias, float *ones, float *offset, float *mask, @@ -266,7 +251,6 @@ void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle, const int height = in_h; const int width = in_w; - const int channels_out = out_c; const int height_out = (height + 2 * pad_h - (dilation_h * (kernel_h - 1) + 1)) / stride_h + 1; @@ -282,17 +266,15 @@ void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle, n, m, k, &alpha, ones, k, bias, k, &beta, output, n); - if (stat != CUBLAS_STATUS_SUCCESS) { - printf ("CUBLAS initialization failed\n"); - return ; - } + if (stat != CUBLAS_STATUS_SUCCESS) + FatalError("CUBLAS initialization failed\n"); - modulated_deformable_im2col_cuda(stream, + modulatedDeformableIm2colCuda(stream, input, offset, mask, 1, channels, height, width, height_out, width_out, deformable_group, columns); - // modulated_deformable_im2col_cuda2(stream, + // modulatedDeformableIm2colCudaGeneralVersion(stream, // input, offset, // mask, // 1, channels, height, width, @@ -310,8 +292,7 @@ void dcn_v2_cuda_forward(cublasStatus_t stat, cublasHandle_t handle, columns, n, weight, k, &beta, output, n); - if (stat != CUBLAS_STATUS_SUCCESS) { - printf ("CUBLAS initialization failed\n"); - return ; - } + if (stat != CUBLAS_STATUS_SUCCESS) + FatalError("CUBLAS initialization failed\n"); + }