From a98271b0817adbae934054bd48320704049f0239 Mon Sep 17 00:00:00 2001 From: Wang Huan Date: Fri, 18 Apr 2025 07:37:01 +0000 Subject: [PATCH 1/2] refine forrange --- paddle/phi/kernels/funcs/for_range.h | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/paddle/phi/kernels/funcs/for_range.h b/paddle/phi/kernels/funcs/for_range.h index 7ff8ac7177a3de..fab56a238d075f 100644 --- a/paddle/phi/kernels/funcs/for_range.h +++ b/paddle/phi/kernels/funcs/for_range.h @@ -60,7 +60,7 @@ __global__ static void ForRangeElemwiseOp(Function func, unsigned int limit) { template __global__ static void ForRangeElemwiseOpLargeSize(Function func, size_t limit) { - size_t idx = + ` size_t idx = static_cast(blockIdx.x) * static_cast(blockDim.x) + static_cast(threadIdx.x); if (idx < limit) { @@ -89,7 +89,8 @@ struct ForRange { if (grid_size == 1) { ForRangeElemwiseOpGridIsOne<<<1, block_size, 0, dev_ctx_.stream()>>>( func); - } else if (limit_ > std::numeric_limits::max()) { + } else if (block_size * grid_size > + std::numeric_limits::max()) { ForRangeElemwiseOpLargeSize<< Date: Fri, 18 Apr 2025 07:38:22 +0000 Subject: [PATCH 2/2] refine forrange --- paddle/phi/kernels/funcs/for_range.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/paddle/phi/kernels/funcs/for_range.h b/paddle/phi/kernels/funcs/for_range.h index fab56a238d075f..227b1cdf504903 100644 --- a/paddle/phi/kernels/funcs/for_range.h +++ b/paddle/phi/kernels/funcs/for_range.h @@ -60,7 +60,7 @@ __global__ static void ForRangeElemwiseOp(Function func, unsigned int limit) { template __global__ static void ForRangeElemwiseOpLargeSize(Function func, size_t limit) { - ` size_t idx = + size_t idx = static_cast(blockIdx.x) * static_cast(blockDim.x) + static_cast(threadIdx.x); if (idx < limit) {