diff --git a/paddle/fluid/operators/deformable_conv_op.cu b/paddle/fluid/operators/deformable_conv_op.cu index 0a771627e060f44cc19fb897bdc4c82bf74a74ec..67f5ee332eeb2f16356b7274bf5462543c03553d 100644 --- a/paddle/fluid/operators/deformable_conv_op.cu +++ b/paddle/fluid/operators/deformable_conv_op.cu @@ -126,7 +126,8 @@ __global__ void ModulatedDeformableCol2imGpuKernel( DmcnGetGradientWeight(cur_inv_h_data, cur_inv_w_data, cur_h + dy, cur_w + dx, height, width); - atomicAdd(grad_im + cur_bottom_grad_pos, weight * cur_top_grad); + platform::CudaAtomicAdd(grad_im + cur_bottom_grad_pos, + weight * cur_top_grad); } } } @@ -748,6 +749,8 @@ namespace ops = paddle::operators; using CUDA = paddle::platform::CUDADeviceContext; REGISTER_OP_CUDA_KERNEL(deformable_conv, - ops::DeformableConvCUDAKernel); + ops::DeformableConvCUDAKernel, + ops::DeformableConvCUDAKernel); REGISTER_OP_CUDA_KERNEL(deformable_conv_grad, - ops::DeformableConvGradCUDAKernel); + ops::DeformableConvGradCUDAKernel, + ops::DeformableConvGradCUDAKernel); diff --git a/paddle/fluid/operators/deformable_conv_v1_op.cc b/paddle/fluid/operators/deformable_conv_v1_op.cc index dfba2070aac77a07c789592fd9097de3326b0d67..090d8a1fab0b94ba2871f536360bf4112ae9383d 100644 --- a/paddle/fluid/operators/deformable_conv_v1_op.cc +++ b/paddle/fluid/operators/deformable_conv_v1_op.cc @@ -307,6 +307,8 @@ REGISTER_OPERATOR(deformable_conv_v1, ops::DeformableConvV1Op, REGISTER_OPERATOR(deformable_conv_v1_grad, ops::DeformableConvV1GradOp); REGISTER_OP_CPU_KERNEL(deformable_conv_v1, - ops::DeformableConvV1CPUKernel); + ops::DeformableConvV1CPUKernel, + ops::DeformableConvV1CPUKernel); REGISTER_OP_CPU_KERNEL(deformable_conv_v1_grad, - ops::DeformableConvV1GradCPUKernel); + ops::DeformableConvV1GradCPUKernel, + ops::DeformableConvV1GradCPUKernel); diff --git a/paddle/fluid/operators/deformable_conv_v1_op.cu b/paddle/fluid/operators/deformable_conv_v1_op.cu index a865766f9adbbe2e4a3c994d774438dff731a732..e399a1fafdb71dacd2b4097f30684948fbbd7432 100644 --- a/paddle/fluid/operators/deformable_conv_v1_op.cu +++ b/paddle/fluid/operators/deformable_conv_v1_op.cu @@ -99,7 +99,8 @@ __global__ void DeformableCol2imCUDAKernel( DmcnGetGradientWeight(cur_inv_h_data, cur_inv_w_data, cur_h + dy, cur_w + dx, height, width); - atomicAdd(grad_im + cur_bottom_grad_pos, weight * cur_top_grad); + platform::CudaAtomicAdd(grad_im + cur_bottom_grad_pos, + weight * cur_top_grad); } } } @@ -604,6 +605,8 @@ class DeformableConvV1GradCUDAKernel : public framework::OpKernel { namespace ops = paddle::operators; REGISTER_OP_CUDA_KERNEL(deformable_conv_v1, - ops::DeformableConvV1CUDAKernel); + ops::DeformableConvV1CUDAKernel, + ops::DeformableConvV1CUDAKernel); REGISTER_OP_CUDA_KERNEL(deformable_conv_v1_grad, - ops::DeformableConvV1GradCUDAKernel); + ops::DeformableConvV1GradCUDAKernel, + ops::DeformableConvV1GradCUDAKernel); diff --git a/python/paddle/fluid/tests/unittests/test_deformable_conv_op.py b/python/paddle/fluid/tests/unittests/test_deformable_conv_op.py index 13624d189f72b61f1e042d0353e594add08a5ce7..45a23231945ece4247b5e5f1b9eaa63f8c33f964 100644 --- a/python/paddle/fluid/tests/unittests/test_deformable_conv_op.py +++ b/python/paddle/fluid/tests/unittests/test_deformable_conv_op.py @@ -111,7 +111,7 @@ def dconv_im2col_gemm(input, offset, mask, filter, group, conv_param): class TestModulatedDeformableConvOp(OpTest): def setUp(self): self.op_type = "deformable_conv" - self.dtype = np.float32 + self.init_type() self.init_group() self.init_dilation() self.init_test_case() @@ -183,6 +183,9 @@ class TestModulatedDeformableConvOp(OpTest): def init_group(self): self.groups = 1 + def init_type(self): + self.dtype = np.float32 + class TestWithStride(TestModulatedDeformableConvOp): def init_test_case(self): @@ -258,6 +261,32 @@ class TestWithGroup(TestModulatedDeformableConvOp): self.groups = 2 +class TestWithDouble(TestModulatedDeformableConvOp): + def init_type(self): + self.dtype = np.float64 + + def init_test_case(self): + self.pad = [1, 1] + self.stride = [1, 1] + self.dilations = [1, 1] + self.input_size = [2, 6, 4, 4] # NCHW + assert np.mod(self.input_size[1], self.groups) == 0 + f_c = self.input_size[1] // self.groups + self.filter_size = [4, f_c, 3, 3] + self.im2col_step = 1 + self.deformable_groups = 1 + offset_c = 2 * self.deformable_groups * self.filter_size[ + 2] * self.filter_size[3] + mask_c = self.deformable_groups * self.filter_size[ + 2] * self.filter_size[3] + self.offset_size = [ + self.input_size[0], offset_c, self.input_size[2], self.input_size[3] + ] + self.mask_size = [ + self.input_size[0], mask_c, self.input_size[2], self.input_size[3] + ] + + class TestModulatedDeformableConvInvalidInput(unittest.TestCase): def test_error(self): def test_invalid_input(): diff --git a/python/paddle/fluid/tests/unittests/test_deformable_conv_v1_op.py b/python/paddle/fluid/tests/unittests/test_deformable_conv_v1_op.py index 769f05b0fcd598af2a35a51bf04b4ab657d5a829..e8b18d601afae649ba6af49230f41bc0465a8959 100644 --- a/python/paddle/fluid/tests/unittests/test_deformable_conv_v1_op.py +++ b/python/paddle/fluid/tests/unittests/test_deformable_conv_v1_op.py @@ -108,7 +108,7 @@ def dconv_im2col_gemm(input, offset, filter, group, conv_param): class TestModulatedDeformableConvOp(OpTest): def setUp(self): self.op_type = "deformable_conv_v1" - self.dtype = np.float32 + self.init_type() self.init_group() self.init_dilation() self.init_test_case() @@ -177,6 +177,9 @@ class TestModulatedDeformableConvOp(OpTest): def init_group(self): self.groups = 1 + def init_type(self): + self.dtype = np.float32 + class TestWithStride(TestModulatedDeformableConvOp): def init_test_case(self): @@ -253,6 +256,11 @@ class TestWithGroup(TestModulatedDeformableConvOp): self.groups = 2 +class TestWithDouble(TestModulatedDeformableConvOp): + def init_type(self): + self.dtype = np.float64 + + class TestModulatedDeformableConvV1InvalidInput(unittest.TestCase): def test_error(self): def test_invalid_input():