From 04fd98930b53e587f95f3ba5dc7f5999472cde00 Mon Sep 17 00:00:00 2001
From: sweetsky0901 <work@yq01-idl-gpu-online20.yq01.baidu.com>
Date: Mon, 20 Nov 2017 18:24:19 +0800
Subject: [PATCH] for code review 6

---
 paddle/operators/math/maxouting.cc |  2 --
 paddle/operators/math/maxouting.cu |  3 +--
 paddle/operators/maxout_op.cu.cc   | 12 +++++-------
 3 files changed, 6 insertions(+), 11 deletions(-)

diff --git a/paddle/operators/math/maxouting.cc b/paddle/operators/math/maxouting.cc
index bcd4da612..e5168ce7a 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 0a8afbbac..7c698577b 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 5ee431cb2..a5823fba6 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>);
-- 
GitLab