From 97b88ef52d96da859b4ad53527357bbde735a624 Mon Sep 17 00:00:00 2001 From: Davide Sapienza Date: Tue, 11 Feb 2020 16:28:17 +0100 Subject: [PATCH] Fix a bug in the deformable kernel. There was a wrong variable initialization. Fixes: 51ffcb1f50e3c9aa2b2 ("Optimize deformable kernel") Signed-off-by: Davide Sapienza --- src/kernels/deformable_conv.cu | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/src/kernels/deformable_conv.cu b/src/kernels/deformable_conv.cu index 88e62e9..672655f 100644 --- a/src/kernels/deformable_conv.cu +++ b/src/kernels/deformable_conv.cu @@ -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(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(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(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(0); if (h_im < height && w_im < width && h_im > -1 && w_im > -1) { //const float map_h = i * dilation_h + offset_h;