From 9ebe930a3cba8349d31f16c610f69a05048ff484 Mon Sep 17 00:00:00 2001 From: zhiqiu Date: Thu, 9 Dec 2021 16:46:49 +0800 Subject: [PATCH] fix int32 overflow in cuda kernel loop --- paddle/fluid/operators/label_smooth_op.cu | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/paddle/fluid/operators/label_smooth_op.cu b/paddle/fluid/operators/label_smooth_op.cu index 33ae35a81f8481..c94a37f03f2b72 100644 --- a/paddle/fluid/operators/label_smooth_op.cu +++ b/paddle/fluid/operators/label_smooth_op.cu @@ -21,8 +21,7 @@ template __global__ void LabelSmoothRunOriginKernel(const int N, const float epsilon, const int label_dim, const T* src, T* dst) { - int idx = blockDim.x * blockIdx.x + threadIdx.x; - for (; idx < N; idx += blockDim.x * gridDim.x) { + CUDA_KERNEL_LOOP(idx, N) { dst[idx] = static_cast(1 - epsilon) * src[idx] + static_cast(epsilon / label_dim); } @@ -32,8 +31,7 @@ template __global__ void LabelSmoothRunDistKernel(const int N, const float epsilon, const int dist_numel, const T* src, const T* dist_data, T* dst) { - int idx = blockDim.x * blockIdx.x + threadIdx.x; - for (; idx < N; idx += blockDim.x * gridDim.x) { + CUDA_KERNEL_LOOP(idx, N) { int dist_idx = idx % dist_numel; dst[idx] = static_cast(1 - epsilon) * src[idx] + static_cast(epsilon) * dist_data[dist_idx]; @@ -43,8 +41,7 @@ __global__ void LabelSmoothRunDistKernel(const int N, const float epsilon, template __global__ void LabelSmoothGradRunKernel(const int N, const float epsilon, const T* src, T* dst) { - int idx = blockDim.x * blockIdx.x + threadIdx.x; - for (; idx < N; idx += blockDim.x * gridDim.x) { + CUDA_KERNEL_LOOP(idx, N) { dst[idx] = static_cast(1 - epsilon) * src[idx]; } }