From e3f344ba9963f6341de3f287d968dd76562027fb Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Mon, 17 Aug 2020 12:52:40 -0400 Subject: [PATCH] SWDEV-248057 - fix the calculation of allGridSize used in multi_grid_group() API Change-Id: Ib470094e28dcacaa4769dc5c7ab08924f5b7fa41 [ROCm/hip commit: 4f400bc5e90fb60f11f2a558aa678bb8f96d0454] --- projects/hip/rocclr/hip_module.cpp | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/projects/hip/rocclr/hip_module.cpp b/projects/hip/rocclr/hip_module.cpp index 051578dc0a..0ddeaedded 100755 --- a/projects/hip/rocclr/hip_module.cpp +++ b/projects/hip/rocclr/hip_module.cpp @@ -461,14 +461,17 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL hipError_t result = hipErrorUnknown; uint64_t allGridSize = 0; + uint32_t blockDims = 0; std::vector mgpu_list(numDevices); for (int i = 0; i < numDevices; ++i) { const hipLaunchParams& launch = launchParamsList[i]; - allGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z; + blockDims = launch.blockDim.x * launch.blockDim.y * launch.blockDim.z; + allGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z * + blockDims; // Make sure block dimensions are valid - if (0 == launch.blockDim.x * launch.blockDim.y * launch.blockDim.z) { + if (0 == blockDims) { return hipErrorInvalidConfiguration; } if (launch.stream != nullptr) {