diff --git a/paddle/operators/math/maxouting.cc b/paddle/operators/math/maxouting.cc
index bcd4da612c841e241b63e62ce19e7c5f4a7db2d5..e5168ce7afd4139475afa6edd5999b9974407c9b 100644
--- a/paddle/operators/math/maxouting.cc
+++ b/paddle/operators/math/maxouting.cc
@@ -85,11 +85,9 @@ public:
           int output_idx = blen + clen + f;
           for (int g = 0; g < groups && continue_match; ++g) {
               int input_idx = input_idx0 + fea_size * g;
-              input_grad_data[input_idx] = 0;
               if (input_data[input_idx] == output_data[output_idx]) {
                 input_grad_data[input_idx] += output_grad_data[output_idx];
                 continue_match = false;
-                break;
               }
           }
         }
diff --git a/paddle/operators/math/maxouting.cu b/paddle/operators/math/maxouting.cu
index 0a8afbbaca3a793b62bc37fd5997a0f6213661f5..7c698577b8a8258a58ba9a2b6c675457b2458a5b 100644
--- a/paddle/operators/math/maxouting.cu
+++ b/paddle/operators/math/maxouting.cu
@@ -69,8 +69,7 @@ __global__ void KernelMaxoutGrad(
         }
       }
       if (max_index != -1) {
-        // atomic add
-        platform::CudaAtomicAdd(input_grad + max_index, output_grad[index]);
+        input_grad[max_index] += output_grad[index];
       }
     }
 }
diff --git a/paddle/operators/maxout_op.cu.cc b/paddle/operators/maxout_op.cu.cc
index 5ee431cb26cde12a620808fb07dbd26843c64bc5..a5823fba6848a0d42a743c90d7d683e3e4ae4422 100644
--- a/paddle/operators/maxout_op.cu.cc
+++ b/paddle/operators/maxout_op.cu.cc
@@ -15,13 +15,11 @@
 #include "paddle/operators/maxout_op.h"
 
 namespace ops = paddle::operators;
-REGISTER_OP_GPU_KERNEL(maxout, ops::MaxOutKernel<paddle::platform::GPUPlace,
-                       float>);
-REGISTER_OP_GPU_KERNEL(maxout, ops::MaxOutKernel<paddle::platform::GPUPlace,
-                       double>);
+REGISTER_OP_GPU_KERNEL(maxout,
+                       ops::MaxOutKernel<paddle::platform::GPUPlace, float>,
+                       ops::MaxOutKernel<paddle::platform::GPUPlace, double>);
 REGISTER_OP_GPU_KERNEL(maxout_grad,
                        ops::MaxOutGradKernel<paddle::platform::GPUPlace,
-                       float>);
-REGISTER_OP_GPU_KERNEL(maxout_grad,
+                        float>,
                        ops::MaxOutGradKernel<paddle::platform::GPUPlace,
-                       double>);
+                        double>);