[device-libs][comgr] - Add gfx1250 and gfx1251 support (#553)

Author: kzhuravl

amd/comgr/src/comgr-isa-metadata.def

@@ -72,6 +72,8 @@ HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx1152",          false, false, EF_AMDGPU_MAC
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx1153",          false, false, EF_AMDGPU_MACH_AMDGCN_GFX1153,         true,  true, 65536,  32,  4,   40, 1024,  106, 800, 106,   16, 1024, 256)
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx1200",          false, false, EF_AMDGPU_MACH_AMDGCN_GFX1200,         true,  true, 65536,  32,  4,   40, 1024,  106, 800, 106,   24, 1536, 256)
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx1201",          false, false, EF_AMDGPU_MACH_AMDGCN_GFX1201,         true,  true, 65536,  32,  4,   40, 1024,  106, 800, 106,   24, 1536, 256)
+HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx1250",          true,  true,  EF_AMDGPU_MACH_AMDGCN_GFX1250,         true, false, 327680, 64,  4,   40, 1024,  106, 800, 106,   16, 1024, 1024)
+HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx1251",          true,  true,  EF_AMDGPU_MACH_AMDGCN_GFX1251,         true, false, 327680, 64,  4,   40, 1024,  106, 800, 106,   16, 1024, 1024)
 
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx9-generic",     false,  true, EF_AMDGPU_MACH_AMDGCN_GFX9_GENERIC,    true,  true, 65536,  32,  4,   40, 1024,   16, 800, 102,    4, 256, 256)
 HANDLE_ISA("amdgcn-amd-amdhsa-", "gfx9-4-generic",   true,   true, EF_AMDGPU_MACH_AMDGCN_GFX9_4_GENERIC,  true, false, 65536,  32,  4,   40, 1024,   16, 800, 102,    8, 512, 512)

amd/comgr/test/get_data_isa_name_test.c

@@ -77,6 +77,8 @@ static isa_features_t IsaFeatures[] = {
   {"amdgcn-amd-amdhsa--gfx1153",         false,  false,  false},
   {"amdgcn-amd-amdhsa--gfx1200",         false,  false,  false},
   {"amdgcn-amd-amdhsa--gfx1201",         false,  false,  false},
+  {"amdgcn-amd-amdhsa--gfx1250",         false,  false,  false},
+  {"amdgcn-amd-amdhsa--gfx1251",         false,  false,  false},
 
   {"amdgcn-amd-amdhsa--gfx9-generic",    false,  true,   true},
   {"amdgcn-amd-amdhsa--gfx9-4-generic",  true,   true,   true},

amd/device-libs/ockl/src/cluster.cl

@@ -0,0 +1,45 @@
+/*===--------------------------------------------------------------------------
+ *                   ROCm Device Libraries
+ *
+ * This file is distributed under the University of Illinois Open Source
+ * License. See LICENSE.TXT for details.
+ *===------------------------------------------------------------------------*/
+
+#include "oclc.h"
+#include "ockl.h"
+
+__attribute__((target("cumode,gfx1250-insts"), const)) uint
+__ockl_cluster_num_workgroups(int dim)
+{
+    switch (dim) {
+    case 0:
+        return __builtin_amdgcn_cluster_workgroup_max_id_x() + 1;
+    case 1:
+        return __builtin_amdgcn_cluster_workgroup_max_id_y() + 1;
+    case 2:
+        return __builtin_amdgcn_cluster_workgroup_max_id_z() + 1;
+    default:
+        return 1;
+    }
+}
+
+__attribute__((target("cumode,gfx1250-insts"), const)) uint
+__ockl_cluster_workgroup_id(int dim)
+{
+    switch (dim) {
+    case 0:
+        return __builtin_amdgcn_cluster_workgroup_id_x();
+    case 1:
+        return __builtin_amdgcn_cluster_workgroup_id_y();
+    case 2:
+        return __builtin_amdgcn_cluster_workgroup_id_z();
+    default:
+        return 0;
+    }
+}
+
+__attribute__((target("cumode,gfx1250-insts"), const)) uint
+__ockl_cluster_flat_num_workgroups(void)
+{
+    return __builtin_amdgcn_cluster_workgroup_max_flat_id() + 1;
+}

amd/device-libs/ockl/src/dots.cl

@@ -64,7 +64,7 @@ __attribute__((target("dot8-insts"), const)) static uint amdgcn_sudot8(bool as,
                         return __builtin_amdgcn_sudot8(true , a, true , b, c, true );
 }
 
-#define SWDOT __oclc_ISA_version < 9006 || __oclc_ISA_version == 9009 || __oclc_ISA_version == 10100
+#define SWDOT __oclc_ISA_version < 9006 || __oclc_ISA_version == 9009 || __oclc_ISA_version == 10100 || __oclc_ISA_version == 12500 || __oclc_ISA_version == 12501
 #define SWIDOT2 __oclc_ISA_version < 9006 || __oclc_ISA_version == 9009 || __oclc_ISA_version == 10100 || __oclc_ISA_version >= 11000
 #define SUDOT __oclc_ISA_version >= 11000
 

amd/device-libs/oclc/src/isa_version_1250.cl

@@ -0,0 +1,10 @@
+/*===--------------------------------------------------------------------------
+ *                   ROCm Device Libraries
+ *
+ * This file is distributed under the University of Illinois Open Source
+ * License. See LICENSE.TXT for details.
+ *===------------------------------------------------------------------------*/
+
+#include "oclc.h"
+
+const __constant int __oclc_ISA_version = 12500;

amd/device-libs/oclc/src/isa_version_1251.cl

@@ -0,0 +1,10 @@
+/*===--------------------------------------------------------------------------
+ *                   ROCm Device Libraries
+ *
+ * This file is distributed under the University of Illinois Open Source
+ * License. See LICENSE.TXT for details.
+ *===------------------------------------------------------------------------*/
+
+#include "oclc.h"
+
+const __constant int __oclc_ISA_version = 12501;