support double in deformable conv

Author: jerrywgz

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<float>);
+                       ops::DeformableConvV1CPUKernel<float>,
+                       ops::DeformableConvV1CPUKernel<double>);
 REGISTER_OP_CPU_KERNEL(deformable_conv_v1_grad,
-                       ops::DeformableConvV1GradCPUKernel<float>);
+                       ops::DeformableConvV1GradCPUKernel<float>,
+                       ops::DeformableConvV1GradCPUKernel<double>);

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<T> {
 namespace ops = paddle::operators;
 
 REGISTER_OP_CUDA_KERNEL(deformable_conv_v1,
-                        ops::DeformableConvV1CUDAKernel<float>);
+                        ops::DeformableConvV1CUDAKernel<float>,
+                        ops::DeformableConvV1CUDAKernel<double>);
 REGISTER_OP_CUDA_KERNEL(deformable_conv_v1_grad,
-                        ops::DeformableConvV1GradCUDAKernel<float>);
+                        ops::DeformableConvV1GradCUDAKernel<float>,
+                        ops::DeformableConvV1GradCUDAKernel<double>);

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 @@ def init_dilation(self):
     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 @@ def init_group(self):
         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():