mmr/tkDNN
Archived
1
0
Fork 0
Code Issues 46 Pull Requests 12 Releases 4 Activity

Fix the Deformable convolution code sintax.

Browse Source 
Signed-off-by: Davide Sapienza <sapienza.dav@gmail.com>
This commit is contained in:
Davide Sapienza
2020-04-08 11:14:34 +02:00
parent 5562f599a6
commit 0e5c90634f
4 changed files with 50 additions and 106 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/tkDNN/kernels.h
+1 -14
View File
@@ -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,
include/tkDNN/pluginsRT/DeformableConvRT.h
+6 -26
View File
@@ -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;
};
src/DeformConv2d.cpp
+18 -22
View File
@@ -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; i<height_ones*width_ones; i++)
aus1[i]=1.0f;
checkCuda( cudaMemcpy(ones_d1, aus1, (height_ones*width_ones)*sizeof(dnnType), cudaMemcpyHostToDevice) );
checkCuda( cudaFreeHost(aus1) );
ones_h1[i]=1.0f;
checkCuda( cudaMemcpy(ones_d1, ones_h1, (height_ones*width_ones)*sizeof(dnnType), cudaMemcpyHostToDevice) );
checkCuda( cudaFreeHost(ones_h1) );
checkCuda( cudaMalloc(&ones_d2, dim_ones*sizeof(dnnType)) );
dnnType *aus2;
checkCuda( cudaMallocHost(&aus2, dim_ones*sizeof(dnnType)) );
dnnType *ones_h2;
checkCuda( cudaMallocHost(&ones_h2, dim_ones*sizeof(dnnType)) );
for(int i=0; i<dim_ones; i++)
aus2[i]=1.0f;
checkCuda( cudaMemcpy(ones_d2, aus2, (dim_ones)*sizeof(dnnType), cudaMemcpyHostToDevice) );
checkCuda( cudaFreeHost(aus2) );
ones_h2[i]=1.0f;
checkCuda( cudaMemcpy(ones_d2, ones_h2, (dim_ones)*sizeof(dnnType), cudaMemcpyHostToDevice) );
checkCuda( cudaFreeHost(ones_h2) );
checkCuda( cudaDeviceSynchronize() );
}
@@ -57,8 +55,7 @@ DeformConv2d::DeformConv2d( Network *net, int out_ch, int deformable_group, int
std::string d_fname_weights, std::string fname_weights, bool batchnorm) :
LayerWgs(net, net->getOutputDim().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,
src/kernels/deformable_conv.cu
+25 -44
View File
@@ -1,6 +1,8 @@
#include <cstdio>
#include <algorithm>
#include <cstring>
#include <string>
#include <iostream>
#include "kernels.h"
#include <errno.h>
@@ -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<float>(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<float>(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
<<<GET_BLOCKS(num_kernels), CUDA_NUM_THREADS,
0, stream>>>(
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");
}