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: a0f29b454c]

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;