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

Fix a bug in the deformable kernel.

Browse Source 
There was a wrong variable initialization.

Fixes: 51ffcb1f50 ("Optimize deformable kernel")

Signed-off-by: Davide Sapienza <sapienza.dav@gmail.com>
This commit is contained in:
Davide Sapienza
2020-02-11 16:28:17 +01:00
parent 4c2d2a7965
commit 97b88ef52d
1 changed files with 9 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/kernels/deformable_conv.cu
+9 -8
View File
@@ -67,16 +67,15 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n,
const int s_col = height_col * width_col;
const int s_col2 = 2 * s_col;
float *data_col_ptr = data_col + ((c_col * batch_size + b_col) * height_col + h_col) * width_col + w_col;
const int first_member = w_col + width_col * h_col;
// float *data_col_ptr = data_col + ((c_col * batch_size + b_col) * height_col + h_col) * width_col + w_col;
float *data_col_ptr = data_col + first_member + s_col * (c_col * batch_size + b_col);
//const float* data_im_ptr = data_im + ((b_col * num_channels + c_im) * height + h_in) * width + w_in;
const float *data_im_ptr = data_im + (b_col * num_channels + c_im) * height * width;
const int add_ptr = (b_col * deformable_group + deformable_group_index) * kk * s_col;
const float *data_offset_ptr = data_offset + add_ptr + add_ptr;
const float *data_mask_ptr = data_mask + add_ptr;
const int first_member = w_col + width_col * h_col;
float val = static_cast<float>(0);
#pragma unroll
for (int i = 0; i < 3; ++i)
{
@@ -99,6 +98,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel(const int n,
const float h_im = offset_h + h_in + i;
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)
{
//const float map_h = i * dilation_h + offset_h;
@@ -149,16 +149,16 @@ __global__ void modulated_deformable_im2col_gpu_kernel2(const int n,
const int s_col = height_col * width_col;
const int s_col2 = 2 * s_col;
float *data_col_ptr = data_col + ((c_col * batch_size + b_col) * height_col + h_col) * width_col + w_col;
const int first_member = w_col + width_col * h_col;
// float *data_col_ptr = data_col + ((c_col * batch_size + b_col) * height_col + h_col) * width_col + w_col;
float *data_col_ptr = data_col + first_member + s_col * (c_col * batch_size + b_col);
//const float* data_im_ptr = data_im + ((b_col * num_channels + c_im) * height + h_in) * width + w_in;
const float *data_im_ptr = data_im + (b_col * num_channels + c_im) * height * width;
const int add_ptr = (b_col * deformable_group + deformable_group_index) * kk * s_col;
const float *data_offset_ptr = data_offset + add_ptr + add_ptr;
const float *data_mask_ptr = data_mask + add_ptr;
const int first_member = w_col + width_col * h_col;
float val = static_cast<float>(0);
#pragma unroll
for (int i = 0; i < kernel_h; ++i)
{
@@ -181,6 +181,7 @@ __global__ void modulated_deformable_im2col_gpu_kernel2(const int n,
const float h_im = offset_h + h_in + i * dilation_h;
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)
{
//const float map_h = i * dilation_h + offset_h;