From 6e5ab031c450dee87d31a7950924ba828b4e3f01 Mon Sep 17 00:00:00 2001 From: dsapienza <177992@studenti.unimore.it> Date: Sat, 25 Apr 2020 02:47:31 +0200 Subject: [PATCH] Deformable batch works Signed-off-by: Davide Sapienza --- include/tkDNN/kernels.h | 2 +- include/tkDNN/pluginsRT/DeformableConvRT.h | 45 ++++++++++++---------- src/DeformConv2d.cpp | 4 +- src/kernels/deformable_conv.cu | 12 +++--- 4 files changed, 35 insertions(+), 28 deletions(-) diff --git a/include/tkDNN/kernels.h b/include/tkDNN/kernels.h index afcc168..16f07d1 100644 --- a/include/tkDNN/kernels.h +++ b/include/tkDNN/kernels.h @@ -41,7 +41,7 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, const int stride_h, const int stride_w, const int pad_h, const int pad_w, const int dilation_h, const int dilation_w, - const int deformable_group, + const int deformable_group, const int batch_id, 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 = cudaStream_t(0)); diff --git a/include/tkDNN/pluginsRT/DeformableConvRT.h b/include/tkDNN/pluginsRT/DeformableConvRT.h index 4804c03..bff6370 100644 --- a/include/tkDNN/pluginsRT/DeformableConvRT.h +++ b/include/tkDNN/pluginsRT/DeformableConvRT.h @@ -86,26 +86,26 @@ public: dnnType *output_conv = (dnnType*)reinterpret_cast(inputs[1]); // 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 - activationSIGMOIDForward(mask, mask, chunk_dim); - - // deformable convolution - dcnV2CudaForward(stat, handle, - srcData, data_d, - bias2_d, ones_d1, - offset, mask, - reinterpret_cast(outputs[0]), ones_d2, - kh, kw, - sh, sw, - ph, pw, - 1, 1, - deformableGroup, - i_n, i_c, i_h, i_w, - o_n, o_c, o_h, o_w, - chunk_dim); - + for(int b=0; b(outputs[0]), ones_d2, + kh, kw, + sh, sw, + ph, pw, + 1, 1, + deformableGroup, b, + i_n, i_c, i_h, i_w, + o_n, o_c, o_h, o_w, + chunk_dim); + } return 0; } @@ -185,6 +185,11 @@ public: dnnType * offset; dnnType * mask; dnnType *ones_d2; + // dnnType *input_n; + // dnnType *offset_n; + // dnnType *mask_n; + // dnnType *output_n; + tk::dnn::DeformConv2d *defRT; }; diff --git a/src/DeformConv2d.cpp b/src/DeformConv2d.cpp index 06bba7f..b161a22 100644 --- a/src/DeformConv2d.cpp +++ b/src/DeformConv2d.cpp @@ -102,13 +102,13 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) { dcnV2CudaForward(stat, handle, srcData, this->data_d, this->bias2_d, ones_d1, - offset, mask, + offset, mask, dstData, ones_d2, this->kernelH, this->kernelW, this->strideH, this->strideW, this->paddingH, this->paddingW, 1, 1, - this->deformableGroup, + this->deformableGroup, 0, //batch_id for cudnn is set to 0 (no batch) 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, chunk_dim); diff --git a/src/kernels/deformable_conv.cu b/src/kernels/deformable_conv.cu index e46c579..592c538 100644 --- a/src/kernels/deformable_conv.cu +++ b/src/kernels/deformable_conv.cu @@ -241,12 +241,13 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, const int stride_h, const int stride_w, const int pad_h, const int pad_w, const int dilation_h, const int dilation_w, - const int deformable_group, + const int deformable_group, const int batch_id, 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 chunk_dim, cudaStream_t stream) { // stat and handle have be moved out to preserve 2 - 6 milliseconds every 100. + const int batch = batch_id; const int channels = in_c; const int height = in_h; const int width = in_w; @@ -265,13 +266,14 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, stat = cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_N, n, m, k, &alpha, ones, k, bias, k, - &beta, output, n); + &beta, output + batch * out_c * out_h * out_w, n); if (stat != CUBLAS_STATUS_SUCCESS) FatalError("CUBLAS initialization failed\n"); modulatedDeformableIm2colCuda(stream, - input, offset, - mask, + input + batch * channels * height * width, + offset,// + b * 2 * int((float)chunk_dim / batch), + mask,// + b * int((float)chunk_dim / batch), 1, channels, height, width, height_out, width_out, deformable_group, columns); // modulatedDeformableIm2colCudaGeneralVersion(stream, @@ -290,7 +292,7 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle, stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, n, m, k, &alpha, columns, n, weight, k, - &beta, output, n); + &beta, output + batch * out_c * out_h * out_w, n); if (stat != CUBLAS_STATUS_SUCCESS) FatalError("CUBLAS initialization failed\n");