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

Deformable batch works

Browse Source 
Signed-off-by: Davide Sapienza <sapienza.dav@gmail.com>
This commit is contained in:
dsapienza
2020-04-25 02:47:31 +02:00
parent 13c9dc6620
commit 6e5ab031c4
4 changed files with 35 additions and 28 deletions
Download Patch File Download Diff File Expand all files Collapse all files
include/tkDNN/kernels.h
+1 -1
View File
@@ -41,7 +41,7 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
const int stride_h, const int stride_w, const int stride_h, const int stride_w,
const int pad_h, const int pad_w, const int pad_h, const int pad_w,
const int dilation_h, const int dilation_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 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 out_n, const int out_c, const int out_h, const int out_w,
const int dst_dim, cudaStream_t stream = cudaStream_t(0)); const int dst_dim, cudaStream_t stream = cudaStream_t(0));
include/tkDNN/pluginsRT/DeformableConvRT.h
+25 -20
View File
@@ -86,26 +86,26 @@ public:
dnnType *output_conv = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]); dnnType *output_conv = (dnnType*)reinterpret_cast<const dnnType*>(inputs[1]);
// split conv2d outputs into offset to mask // split conv2d outputs into offset to mask
checkCuda(cudaMemcpy(offset, output_conv, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); for(int b=0; b<batchSize; b++) {
checkCuda(cudaMemcpy(mask, output_conv + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice)); checkCuda(cudaMemcpy(offset, output_conv + b * 3 * chunk_dim, 2*chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
// kernel sigmoide checkCuda(cudaMemcpy(mask, output_conv + b * 3 * chunk_dim + 2*chunk_dim, chunk_dim*sizeof(dnnType), cudaMemcpyDeviceToDevice));
activationSIGMOIDForward(mask, mask, chunk_dim); // kernel sigmoide
activationSIGMOIDForward(mask, mask, chunk_dim);
// deformable convolution // deformable convolution
dcnV2CudaForward(stat, handle, dcnV2CudaForward(stat, handle,
srcData, data_d, srcData, data_d,
bias2_d, ones_d1, bias2_d, ones_d1,
offset, mask, offset, mask,
reinterpret_cast<dnnType*>(outputs[0]), ones_d2, reinterpret_cast<dnnType*>(outputs[0]), ones_d2,
kh, kw, kh, kw,
sh, sw, sh, sw,
ph, pw, ph, pw,
1, 1, 1, 1,
deformableGroup, deformableGroup, b,
i_n, i_c, i_h, i_w, i_n, i_c, i_h, i_w,
o_n, o_c, o_h, o_w, o_n, o_c, o_h, o_w,
chunk_dim); chunk_dim);
}
return 0; return 0;
} }
@@ -185,6 +185,11 @@ public:
dnnType * offset; dnnType * offset;
dnnType * mask; dnnType * mask;
dnnType *ones_d2; dnnType *ones_d2;
// dnnType *input_n;
// dnnType *offset_n;
// dnnType *mask_n;
// dnnType *output_n;
tk::dnn::DeformConv2d *defRT; tk::dnn::DeformConv2d *defRT;
}; };
src/DeformConv2d.cpp
+1 -1
View File
@@ -108,7 +108,7 @@ dnnType* DeformConv2d::infer(dataDim_t &dim, dnnType* srcData) {
this->strideH, this->strideW, this->strideH, this->strideW,
this->paddingH, this->paddingW, this->paddingH, this->paddingW,
1, 1, 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, 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, this->output_dim.n, this->output_dim.c, this->output_dim.h, this->output_dim.w,
chunk_dim); chunk_dim);
src/kernels/deformable_conv.cu
+7 -5
View File
@@ -241,12 +241,13 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
const int stride_h, const int stride_w, const int stride_h, const int stride_w,
const int pad_h, const int pad_w, const int pad_h, const int pad_w,
const int dilation_h, const int dilation_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 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 out_n, const int out_c, const int out_h, const int out_w,
const int chunk_dim, cudaStream_t stream) const int chunk_dim, cudaStream_t stream)
{ {
// stat and handle have be moved out to preserve 2 - 6 milliseconds every 100. // 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 channels = in_c;
const int height = in_h; const int height = in_h;
const int width = in_w; 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, 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); &beta, output + batch * out_c * out_h * out_w, n);
if (stat != CUBLAS_STATUS_SUCCESS) if (stat != CUBLAS_STATUS_SUCCESS)
FatalError("CUBLAS initialization failed\n"); FatalError("CUBLAS initialization failed\n");
modulatedDeformableIm2colCuda(stream, modulatedDeformableIm2colCuda(stream,
input, offset, input + batch * channels * height * width,
mask, offset,// + b * 2 * int((float)chunk_dim / batch),
mask,// + b * int((float)chunk_dim / batch),
1, channels, height, width, 1, channels, height, width,
height_out, width_out, deformable_group, columns); height_out, width_out, deformable_group, columns);
// modulatedDeformableIm2colCudaGeneralVersion(stream, // modulatedDeformableIm2colCudaGeneralVersion(stream,
@@ -290,7 +292,7 @@ void dcnV2CudaForward(cublasStatus_t stat, cublasHandle_t handle,
stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, stat = cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N,
n, m, k, &alpha, n, m, k, &alpha,
columns, n, weight, k, columns, n, weight, k,
&beta, output, n); &beta, output + batch * out_c * out_h * out_w, n);
if (stat != CUBLAS_STATUS_SUCCESS) if (stat != CUBLAS_STATUS_SUCCESS)
FatalError("CUBLAS initialization failed\n"); FatalError("CUBLAS initialization failed\n");