diff --git a/Src/Base/AMReX_GpuLaunch.H b/Src/Base/AMReX_GpuLaunch.H index c0d2c9a741b..c1870d2ef58 100644 --- a/Src/Base/AMReX_GpuLaunch.H +++ b/Src/Base/AMReX_GpuLaunch.H @@ -157,7 +157,12 @@ namespace Gpu { makeExecutionConfig (Long N) noexcept { ExecutionConfig ec(dim3{}, dim3{}); - ec.numBlocks.x = (std::max(N,Long(1)) + MT - 1) / MT; + Long numBlocks = (std::max(N,Long(1)) + MT - 1) / MT; + // ensure that blockDim.x*gridDim.x does not overflow + numBlocks = std::min(numBlocks, Long(std::numeric_limits::max()/MT)); + // ensure that the maximum grid size of 2^31-1 won't be exceeded + numBlocks = std::min(numBlocks, Long(std::numeric_limits::max())); + ec.numBlocks.x = numBlocks; ec.numThreads.x = MT; AMREX_ASSERT(MT % Gpu::Device::warp_size == 0); return ec; @@ -167,11 +172,7 @@ namespace Gpu { ExecutionConfig makeExecutionConfig (const Box& box) noexcept { - ExecutionConfig ec(dim3{}, dim3{}); - ec.numBlocks.x = (std::max(box.numPts(),Long(1)) + MT - 1) / MT; - ec.numThreads.x = MT; - AMREX_ASSERT(MT % Gpu::Device::warp_size == 0); - return ec; + return makeExecutionConfig(box.numPts()); } #endif