From afa60c590c19a6d3cbdab24dbcb90c4948e86b9c Mon Sep 17 00:00:00 2001 From: "Brzak, Branislav" Date: Wed, 9 Jul 2025 12:35:40 +0200 Subject: [PATCH] SWDEV-540609 - hipGraph capture of hipExtModuleLaunchKernel assumed grid size was a multiple of block size (#694) [ROCm/clr commit: 32e027e66cc95fe60f1c2cee9af89f9f205360d8] --- projects/clr/hipamd/src/hip_graph.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index 42499f29c1..ce367c2bc1 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -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(const_cast(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); }