From 038883fd04a7fc7caee138196864b0b9b810efdb Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E6=9D=8E=E5=AD=A3?= <2042519524@qq.com> Date: Tue, 20 Jul 2021 21:26:31 +0800 Subject: [PATCH] Fix cast op that can not cast the arrays that the size of arrays is beyond int32 (#34209) * fix cast --- paddle/fluid/operators/cast_op.cu | 3 ++- paddle/fluid/platform/gpu_launch_config.h | 2 +- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/paddle/fluid/operators/cast_op.cu b/paddle/fluid/operators/cast_op.cu index 1ac110b3caf..0beb2291060 100644 --- a/paddle/fluid/operators/cast_op.cu +++ b/paddle/fluid/operators/cast_op.cu @@ -40,7 +40,8 @@ __global__ void VecCastCUDAKernel(const InT* in, const int64_t N, OutT* out) { int64_t idx = blockDim.x * blockIdx.x + threadIdx.x; using LoadT = AlignedVector; using StoreT = AlignedVector; - for (int i = idx * VecSize; i < N; i += blockDim.x * gridDim.x * VecSize) { + for (int64_t i = idx * VecSize; i < N; + i += blockDim.x * gridDim.x * VecSize) { InT in_vec[VecSize]; LoadT* in_value = reinterpret_cast(&in_vec); *in_value = *reinterpret_cast(&in[i]); diff --git a/paddle/fluid/platform/gpu_launch_config.h b/paddle/fluid/platform/gpu_launch_config.h index 4da91b4e764..a8226241906 100644 --- a/paddle/fluid/platform/gpu_launch_config.h +++ b/paddle/fluid/platform/gpu_launch_config.h @@ -41,7 +41,7 @@ struct GpuLaunchConfig { }; inline GpuLaunchConfig GetGpuLaunchConfig1D( - const platform::CUDADeviceContext& context, int element_count, + const platform::CUDADeviceContext& context, int64_t element_count, #ifdef PADDLE_WITH_HIP // HIP will throw GPU memory access fault if threads > 256 int max_threads = 256) { -- GitLab