SWDEV-540609 - hipGraph capture of hipExtModuleLaunchKernel assumed grid size was a multiple of block size (#694)

[ROCm/clr commit: 32e027e66c]
This commit is contained in:
Brzak, Branislav
2025-07-09 12:35:40 +02:00
committed by GitHub
parent 433c25eab0
commit afa60c590c
+2 -3
View File
@@ -271,8 +271,7 @@ hipError_t ihipExtLaunchKernel(hipStream_t stream, hipFunction_t f, uint32_t glo
nodeParams.func = f;
nodeParams.blockDim = dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ);
nodeParams.extra = extra;
nodeParams.gridDim = dim3(globalWorkSizeX / localWorkSizeX, globalWorkSizeY / localWorkSizeY,
globalWorkSizeZ / localWorkSizeZ);
nodeParams.gridDim = dim3(globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ);
nodeParams.kernelParams = kernelParams;
nodeParams.sharedMemBytes = sharedMemBytes;
@@ -309,7 +308,7 @@ hipError_t capturehipExtLaunchKernel(hipStream_t& stream, const void*& hostFunct
"[hipGraph] Current capture node ExtLaunchKernel on stream : %p", stream);
return ihipExtLaunchKernel(
stream, reinterpret_cast<hipFunction_t>(const_cast<void*>(hostFunction)),
gridDim.x * blockDim.x, gridDim.y * blockDim.y, gridDim.z * blockDim.z, blockDim.x,
gridDim.x, gridDim.y, gridDim.z, blockDim.x,
blockDim.y, blockDim.z, sharedMemBytes, args, nullptr, startEvent, stopEvent, flags);
}