From d5a8c3fbf9a4d0553dc3927cf755ac652a9aca8c Mon Sep 17 00:00:00 2001 From: Julia Jiang Date: Tue, 7 Jun 2022 19:41:02 -0400 Subject: [PATCH] SWDEV-340754 - Update coordinates in coperatative helper Change-Id: Ibf4e833d10ea5c317900e83f020c9c23f55f3951 [ROCm/clr commit: 30c43056c1003e1c6864170805e7ab4070ad13bd] --- .../hip_cooperative_groups_helper.h | 24 +++++++++---------- 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h index eaee1a715a..e8610e2279 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/amd_detail/hip_cooperative_groups_helper.h @@ -114,23 +114,23 @@ __CG_STATIC_QUALIFIER__ void sync() { __ockl_multi_grid_sync(); } namespace grid { __CG_STATIC_QUALIFIER__ uint32_t size() { - return (uint32_t)((hipBlockDim_z * hipGridDim_z) * (hipBlockDim_y * hipGridDim_y) * - (hipBlockDim_x * hipGridDim_x)); + return (uint32_t)((blockDim.z * gridDim.z) * (blockDim.y * gridDim.y) * + (blockDim.x * gridDim.x)); } __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { // Compute global id of the workgroup to which the current thread belongs to - uint32_t blkIdx = (uint32_t)((hipBlockIdx_z * hipGridDim_y * hipGridDim_x) + - (hipBlockIdx_y * hipGridDim_x) + (hipBlockIdx_x)); + uint32_t blkIdx = (uint32_t)((blockIdx.z * gridDim.y * gridDim.x) + + (blockIdx.y * gridDim.x) + (blockIdx.x)); // Compute total number of threads being passed to reach current workgroup // within grid uint32_t num_threads_till_current_workgroup = - (uint32_t)(blkIdx * (hipBlockDim_x * hipBlockDim_y * hipBlockDim_z)); + (uint32_t)(blkIdx * (blockDim.x * blockDim.y * blockDim.z)); // Compute thread local rank within current workgroup - uint32_t local_thread_rank = (uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) + - (hipThreadIdx_y * hipBlockDim_x) + (hipThreadIdx_x)); + uint32_t local_thread_rank = (uint32_t)((threadIdx.z * blockDim.y * blockDim.x) + + (threadIdx.y * blockDim.x) + (threadIdx.x)); return (num_threads_till_current_workgroup + local_thread_rank); } @@ -148,20 +148,20 @@ __CG_STATIC_QUALIFIER__ void sync() { __ockl_grid_sync(); } namespace workgroup { __CG_STATIC_QUALIFIER__ dim3 group_index() { - return (dim3((uint32_t)hipBlockIdx_x, (uint32_t)hipBlockIdx_y, (uint32_t)hipBlockIdx_z)); + return (dim3((uint32_t)blockIdx.x, (uint32_t)blockIdx.y, (uint32_t)blockIdx.z)); } __CG_STATIC_QUALIFIER__ dim3 thread_index() { - return (dim3((uint32_t)hipThreadIdx_x, (uint32_t)hipThreadIdx_y, (uint32_t)hipThreadIdx_z)); + return (dim3((uint32_t)threadIdx.x, (uint32_t)threadIdx.y, (uint32_t)threadIdx.z)); } __CG_STATIC_QUALIFIER__ uint32_t size() { - return ((uint32_t)(hipBlockDim_x * hipBlockDim_y * hipBlockDim_z)); + return ((uint32_t)(blockDim.x * blockDim.y * blockDim.z)); } __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { - return ((uint32_t)((hipThreadIdx_z * hipBlockDim_y * hipBlockDim_x) + - (hipThreadIdx_y * hipBlockDim_x) + (hipThreadIdx_x))); + return ((uint32_t)((threadIdx.z * blockDim.y * blockDim.x) + + (threadIdx.y * blockDim.x) + (threadIdx.x))); } __CG_STATIC_QUALIFIER__ bool is_valid() {