From 79b9dedbb415eec06df8586073d9f5bbb7584832 Mon Sep 17 00:00:00 2001 From: Joseph Greathouse Date: Wed, 8 Nov 2023 11:19:56 -0600 Subject: [PATCH] SWDEV-431560 - use XCC_ID in __smid() for gfx941 and gfx942 __smid() needs to use both HW_ID and XCC_ID for gfx940, gfx941, and gfx942. Previously, we only did this for gfx940 and thus XCC_ID was incorrectly not passed back on the other two architectures. Change-Id: I9fb13b6cef3280e15463443a180174629d03f8b2 [ROCm/clr commit: a0f29b454c20dd7fc0ae88e0d3da2137355ba44d] --- .../include/hip/amd_detail/amd_device_functions.h | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h index 19f39e16ed..37ef53ff28 100644 --- a/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/projects/clr/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -926,7 +926,7 @@ int __syncthreads_or(int predicate) PIPE_ID 7:6 Pipeline from which the wave was dispatched. CU_ID 11:8 Compute Unit the wave is assigned to. SH_ID 12 Shader Array (within an SE) the wave is assigned to. - SE_ID 15:13 Shader Engine the wave is assigned to for gfx908, gfx90a, gfx940 + SE_ID 15:13 Shader Engine the wave is assigned to for gfx908, gfx90a, gfx940-942 14:13 Shader Engine the wave is assigned to for Vega. TG_ID 19:16 Thread-group ID VM_ID 23:20 Virtual Memory ID @@ -955,7 +955,7 @@ int __syncthreads_or(int predicate) #if (defined(__gfx908__) || defined(__gfx90a__) || \ defined(__GFX11__)) #define HW_ID_SE_ID_SIZE 3 -#else //4 SEs/XCC for gfx940 +#else //4 SEs/XCC for gfx940-942 #define HW_ID_SE_ID_SIZE 2 #endif #if (defined(__GFX10__) || defined(__GFX11__)) @@ -966,7 +966,7 @@ int __syncthreads_or(int predicate) #define HW_ID_SE_ID_OFFSET 13 #endif -#if (defined(__gfx940__)) +#if (defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)) #define XCC_ID 20 #define XCC_ID_XCC_ID_SIZE 4 #define XCC_ID_XCC_ID_OFFSET 0 @@ -1004,7 +1004,7 @@ unsigned __smid(void) unsigned sa_id = __builtin_amdgcn_s_getreg( GETREG_IMMED(HW_ID_SA_ID_SIZE - 1, HW_ID_SA_ID_OFFSET, HW_ID)); #else - #if defined(__gfx940__) + #if (defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)) unsigned xcc_id = __builtin_amdgcn_s_getreg( GETREG_IMMED(XCC_ID_XCC_ID_SIZE - 1, XCC_ID_XCC_ID_OFFSET, XCC_ID)); #endif @@ -1017,7 +1017,7 @@ unsigned __smid(void) temp = (temp << HW_ID_WGP_ID_SIZE) | wgp_id; return temp; //TODO : CU Mode impl - #elif defined(__gfx940__) + #elif (defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__)) unsigned temp = xcc_id; temp = (temp << HW_ID_SE_ID_SIZE) | se_id; temp = (temp << HW_ID_CU_ID_SIZE) | cu_id;