From a49cbb1580c21d4028eb76003af3b5779be8e869 Mon Sep 17 00:00:00 2001 From: mshivama <47909405+mshivama@users.noreply.github.com> Date: Fri, 24 Jan 2020 16:23:28 +0530 Subject: [PATCH] SWDEV-220503: this_grid().thread_rank() gives incorrect result (#1808) * fix a minor bug while computing this.grid()::thread_rank() [ROCm/clr commit: 0605426049798cf42f124a6bde1f8699700c75d3] --- .../include/hip/hcc_detail/hip_cooperative_groups_helper.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h index b74d16d23b..9738448d94 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_cooperative_groups_helper.h @@ -106,7 +106,7 @@ __CG_STATIC_QUALIFIER__ uint32_t size() { } __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { - // Compute global id of the workgroup to which the current threads belongs to + // 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) + @@ -115,7 +115,7 @@ __CG_STATIC_QUALIFIER__ uint32_t thread_rank() { // Compute total number of threads being passed to reach current workgroup // within grid uint32_t num_threads_till_current_workgroup = - (uint32_t)(blkIdx * (hipBlockIdx_x * hipBlockIdx_y * hipBlockIdx_z)); + (uint32_t)(blkIdx * (hipBlockDim_x * hipBlockDim_y * hipBlockDim_z)); // Compute thread local rank within current workgroup uint32_t local_thread_rank =