From 266fcbe0aed3e566c167ea8de5114f62c428c013 Mon Sep 17 00:00:00 2001 From: wangguanzhong Date: Mon, 6 Sep 2021 21:11:34 +0800 Subject: [PATCH] support double in deformable conv (#35330) * support double in deformable conv * add double for dcn v2 --- paddle/fluid/operators/deformable_conv_op.cu | 9 ++++-- .../fluid/operators/deformable_conv_v1_op.cc | 6 ++-- .../fluid/operators/deformable_conv_v1_op.cu | 9 ++++-- .../unittests/test_deformable_conv_op.py | 31 ++++++++++++++++++- .../unittests/test_deformable_conv_v1_op.py | 10 +++++- 5 files changed, 55 insertions(+), 10 deletions(-) diff --git a/paddle/fluid/operators/deformable_conv_op.cu b/paddle/fluid/operators/deformable_conv_op.cu index 0a771627e06..67f5ee332ee 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 dfba2070aac..090d8a1fab0 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 a865766f9ad..e399a1fafdb 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 13624d189f7..45a23231945 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 769f05b0fcd..e8b18d601af 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(): -- GitLab