From a85592bcbf837c6d33c528e1dfea380ed6912d42 Mon Sep 17 00:00:00 2001 From: Zhong Hui Date: Sat, 26 Sep 2020 19:43:52 +0800 Subject: [PATCH] fix cpplint error for the autmic max/min fix cpplint error for the autmic max/min --- .../fluid/operators/math/segment_pooling.cu | 17 +++++++------ paddle/fluid/platform/cuda_primitives.h | 24 +++++++++---------- 2 files changed, 20 insertions(+), 21 deletions(-) diff --git a/paddle/fluid/operators/math/segment_pooling.cu b/paddle/fluid/operators/math/segment_pooling.cu index bb2b6db100..37155fa184 100644 --- a/paddle/fluid/operators/math/segment_pooling.cu +++ b/paddle/fluid/operators/math/segment_pooling.cu @@ -12,13 +12,12 @@ WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the License for the specific language governing permissions and limitations under the License. */ -#include "paddle/fluid/operators/elementwise/elementwise_div_op.h" +#include #include "paddle/fluid/operators/gather.cu.h" #include "paddle/fluid/operators/math/math_function.h" #include "paddle/fluid/operators/math/segment_pooling.h" #include "paddle/fluid/platform/cuda_primitives.h" #include "paddle/fluid/platform/gpu_launch_param_config.h" -#include "paddle/fluid/platform/macros.h" namespace paddle { namespace operators { @@ -100,7 +99,7 @@ __global__ void SegmentOpsKernel(const Index* segment_ids, const T* input, CUDA_KERNEL_LOOP(stripe_index, h.total_stripe_count) { Index segment_offset, dim_index_base, actual_height; Index inner_dim_size = h.inner_dim_size; - h.calculate(stripe_index, segment_offset, dim_index_base, actual_height); + h.calculate(stripe_index, &segment_offset, &dim_index_base, &actual_height); T minmax = pool.initial(); Index first_segment_id = segment_ids[dim_index_base]; @@ -154,7 +153,7 @@ __global__ void SegmentIndexGradKernel(const Index* segment_ids, const T* input, T* in_grad, Helper h) { CUDA_KERNEL_LOOP(stripe_index, h.total_stripe_count) { Index segment_offset, dim_index_base, actual_height; - h.calculate(stripe_index, segment_offset, dim_index_base, actual_height); + h.calculate(stripe_index, &segment_offset, &dim_index_base, &actual_height); for (Index j = 0; j < actual_height; j++) { Index current_segment_id = segment_ids[dim_index_base + j]; @@ -217,11 +216,11 @@ class ArrangeHelper { total_stripe_count = inner_dim_size * input_outer_dim_num_stripe; } - DEVICE inline void calculate(T stripe_index, T& segment_offset, - T& dim_index_base, T& actual_height) { - segment_offset = stripe_index % inner_dim_size; - dim_index_base = stripe_index / inner_dim_size * DimTileSize; - actual_height = min(DimTileSize, input_length_size - dim_index_base); + DEVICE inline void calculate(T stripe_index, T* segment_offset, + T* dim_index_base, T* actual_height) { + *segment_offset = stripe_index % inner_dim_size; + *dim_index_base = stripe_index / inner_dim_size * DimTileSize; + *actual_height = min(DimTileSize, input_length_size - *dim_index_base); } }; diff --git a/paddle/fluid/platform/cuda_primitives.h b/paddle/fluid/platform/cuda_primitives.h index a5dd19d436..4d9673e964 100644 --- a/paddle/fluid/platform/cuda_primitives.h +++ b/paddle/fluid/platform/cuda_primitives.h @@ -137,12 +137,12 @@ USE_CUDA_ATOMIC(Max, unsigned int); #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350 USE_CUDA_ATOMIC(Max, unsigned long long int); // NOLINT #else -CUDA_ATOMIC_WRAPPER(Max, unsigned long long int) { +CUDA_ATOMIC_WRAPPER(Max, unsigned long long int) { // NOLINT if (*address >= val) { return; } - unsigned long long int old = *address, assumed; + unsigned long long int old = *address, assumed; // NOLINT do { assumed = old; @@ -169,7 +169,7 @@ CUDA_ATOMIC_WRAPPER(Max, float) { return; } - int *const address_as_i = (int *)address; + int *const address_as_i = reinterpret_cast(address); int old = *address_as_i, assumed; do { @@ -187,9 +187,9 @@ CUDA_ATOMIC_WRAPPER(Max, double) { return; } - unsigned long long int *const address_as_ull = - (unsigned long long int *)address; - unsigned long long int old = *address_as_ull, assumed; + unsigned long long int *const address_as_ull = // NOLINT + reinterpret_cast(address); // NOLINT + unsigned long long int old = *address_as_ull, assumed; // NOLINT do { assumed = old; @@ -209,12 +209,12 @@ USE_CUDA_ATOMIC(Min, unsigned int); #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350 USE_CUDA_ATOMIC(Min, unsigned long long int); // NOLINT #else -CUDA_ATOMIC_WRAPPER(Min, unsigned long long int) { +CUDA_ATOMIC_WRAPPER(Min, unsigned long long int) { // NOLINT if (*address <= val) { return; } - unsigned long long int old = *address, assumed; + unsigned long long int old = *address, assumed; // NOLINT do { assumed = old; @@ -241,7 +241,7 @@ CUDA_ATOMIC_WRAPPER(Min, float) { return; } - int *const address_as_i = (int *)address; + int *const address_as_i = reinterpret_cast(address); int old = *address_as_i, assumed; do { @@ -259,9 +259,9 @@ CUDA_ATOMIC_WRAPPER(Min, double) { return; } - unsigned long long int *const address_as_ull = - (unsigned long long int *)address; - unsigned long long int old = *address_as_ull, assumed; + unsigned long long int *const address_as_ull = // NOLINT + reinterpret_cast(address); // NOLINT + unsigned long long int old = *address_as_ull, assumed; // NOLINT do { assumed = old; -- GitLab