From 2fd8de1749b2006eabcf31dbd622f1e8452b12df Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 22 May 2019 19:20:09 +0000 Subject: [PATCH] Fix bug in __smid not setting correct size The SZ field should minus by 1 since SIZE range is 1..32. Also add comments that results may vary. [ROCm/clr commit: 2b11a8bf0c865d9b144494173425955465a348a6] --- .../hipamd/include/hip/hcc_detail/device_functions.h | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h index e1f81e4547..808ed216e4 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h @@ -1009,14 +1009,20 @@ void __syncthreads() #define GETREG_IMMED(SZ,OFF,REG) (SZ << 11) | (OFF << 6) | REG +/* + __smid returns the wave's assigned Compute Unit and Shader Engine. + The Compute Unit, CU_ID returned in bits 3:0, and Shader Engine, SE_ID in bits 5:4. + Note: the results vary over time. + SZ minus 1 since SIZE is 1-based. +*/ __device__ inline unsigned __smid(void) { unsigned cu_id = __builtin_amdgcn_s_getreg( - GETREG_IMMED(HW_ID_CU_ID_SIZE, HW_ID_CU_ID_OFFSET, HW_ID)); + GETREG_IMMED(HW_ID_CU_ID_SIZE-1, HW_ID_CU_ID_OFFSET, HW_ID)); unsigned se_id = __builtin_amdgcn_s_getreg( - GETREG_IMMED(HW_ID_SE_ID_SIZE, HW_ID_SE_ID_OFFSET, HW_ID)); + GETREG_IMMED(HW_ID_SE_ID_SIZE-1, HW_ID_SE_ID_OFFSET, HW_ID)); /* Each shader engine has 16 CU */ return (se_id << HW_ID_CU_ID_SIZE) + cu_id;