From 488838e762902f907f71e6b2d44411f0c83b27f8 Mon Sep 17 00:00:00 2001 From: Konstantin Zhuravlyov Date: Thu, 6 Nov 2025 11:52:58 -0500 Subject: [PATCH] [device-libs][comgr] - Add gfx1250 and gfx1251 support --- amd/comgr/src/comgr-isa-metadata.def | 2 + amd/comgr/test/get_data_isa_name_test.c | 2 + amd/device-libs/ockl/src/cluster.cl | 45 ++++++++++++++++++++ amd/device-libs/ockl/src/dots.cl | 2 +- amd/device-libs/oclc/src/isa_version_1250.cl | 10 +++++ amd/device-libs/oclc/src/isa_version_1251.cl | 10 +++++ 6 files changed, 70 insertions(+), 1 deletion(-) create mode 100644 amd/device-libs/ockl/src/cluster.cl create mode 100644 amd/device-libs/oclc/src/isa_version_1250.cl create mode 100644 amd/device-libs/oclc/src/isa_version_1251.cl diff --git a/amd/comgr/src/comgr-isa-metadata.def b/amd/comgr/src/comgr-isa-metadata.def index 16a52f84fa8d2..3a54c3c2669bd 100644 --- a/amd/comgr/src/comgr-isa-metadata.def +++ b/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) diff --git a/amd/comgr/test/get_data_isa_name_test.c b/amd/comgr/test/get_data_isa_name_test.c index 655e0d5ae8ff5..3a2818337e008 100644 --- a/amd/comgr/test/get_data_isa_name_test.c +++ b/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}, diff --git a/amd/device-libs/ockl/src/cluster.cl b/amd/device-libs/ockl/src/cluster.cl new file mode 100644 index 0000000000000..c214f3f401213 --- /dev/null +++ b/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; +} diff --git a/amd/device-libs/ockl/src/dots.cl b/amd/device-libs/ockl/src/dots.cl index 480941b5bbaad..f7838764e695e 100644 --- a/amd/device-libs/ockl/src/dots.cl +++ b/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 diff --git a/amd/device-libs/oclc/src/isa_version_1250.cl b/amd/device-libs/oclc/src/isa_version_1250.cl new file mode 100644 index 0000000000000..fdd13d634c833 --- /dev/null +++ b/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; diff --git a/amd/device-libs/oclc/src/isa_version_1251.cl b/amd/device-libs/oclc/src/isa_version_1251.cl new file mode 100644 index 0000000000000..1c79b6b14ca1d --- /dev/null +++ b/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;