From 2660ea379dbf21e3caa55062a976f16a1cc72df4 Mon Sep 17 00:00:00 2001 From: whs Date: Fri, 4 Sep 2020 17:50:46 +0800 Subject: [PATCH] Fix cuda kernel of affine grid (#27003) test=develop --- paddle/fluid/operators/affine_grid_op.cu | 16 ++++++++-------- .../fluid/tests/unittests/test_affine_grid_op.py | 1 - 2 files changed, 8 insertions(+), 9 deletions(-) diff --git a/paddle/fluid/operators/affine_grid_op.cu b/paddle/fluid/operators/affine_grid_op.cu index 7aaaa0002c5..58b56bdcf56 100644 --- a/paddle/fluid/operators/affine_grid_op.cu +++ b/paddle/fluid/operators/affine_grid_op.cu @@ -62,11 +62,11 @@ __global__ void affine_grid_kernel(const int count, int n, int out_h, int out_w, int theta_offset = n * 6; // 2 * 3; // affine from (h_coor, w_coor) to (x, y) - output[index * 2] = theta[theta_offset] * h_coor + - theta[theta_offset + 1] * w_coor + + output[index * 2] = theta[theta_offset] * w_coor + + theta[theta_offset + 1] * h_coor + theta[theta_offset + 2]; - output[index * 2 + 1] = theta[theta_offset + 3] * h_coor + - theta[theta_offset + 4] * w_coor + + output[index * 2 + 1] = theta[theta_offset + 3] * w_coor + + theta[theta_offset + 4] * h_coor + theta[theta_offset + 5]; } } @@ -86,13 +86,13 @@ __global__ void affine_grid_grad_kernel(const int count, int n, int out_h, int theta_offset = n * 6; // 2 * 3; T out_grad_x = out_grad[index * 2]; - platform::CudaAtomicAdd(theta_grad + theta_offset, out_grad_x * h_coor); - platform::CudaAtomicAdd(theta_grad + theta_offset + 1, out_grad_x * w_coor); + platform::CudaAtomicAdd(theta_grad + theta_offset, out_grad_x * w_coor); + platform::CudaAtomicAdd(theta_grad + theta_offset + 1, out_grad_x * h_coor); platform::CudaAtomicAdd(theta_grad + theta_offset + 2, out_grad_x); T out_grad_y = out_grad[index * 2 + 1]; - platform::CudaAtomicAdd(theta_grad + theta_offset + 3, out_grad_y * h_coor); - platform::CudaAtomicAdd(theta_grad + theta_offset + 4, out_grad_y * w_coor); + platform::CudaAtomicAdd(theta_grad + theta_offset + 3, out_grad_y * w_coor); + platform::CudaAtomicAdd(theta_grad + theta_offset + 4, out_grad_y * h_coor); platform::CudaAtomicAdd(theta_grad + theta_offset + 5, out_grad_y); } } diff --git a/python/paddle/fluid/tests/unittests/test_affine_grid_op.py b/python/paddle/fluid/tests/unittests/test_affine_grid_op.py index 55612d71a17..d3e990ca13e 100644 --- a/python/paddle/fluid/tests/unittests/test_affine_grid_op.py +++ b/python/paddle/fluid/tests/unittests/test_affine_grid_op.py @@ -49,7 +49,6 @@ class TestAffineGridOp(OpTest): self.initTestCase() self.op_type = "affine_grid" theta = np.random.randint(1, 3, self.theta_shape).astype("float32") - theta = np.ones(self.theta_shape).astype("float32") self.inputs = {'Theta': theta} self.attrs = { "use_cudnn": self.use_cudnn, -- GitLab