Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Nov 11, 2025
+70 −1
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions amd/comgr/src/comgr-isa-metadata.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand Down
2 changes: 2 additions & 0 deletions amd/comgr/test/get_data_isa_name_test.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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},
Expand Down
45 changes: 45 additions & 0 deletions amd/device-libs/ockl/src/cluster.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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;
}
2 changes: 1 addition & 1 deletion amd/device-libs/ockl/src/dots.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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

Expand Down
10 changes: 10 additions & 0 deletions amd/device-libs/oclc/src/isa_version_1250.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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;
10 changes: 10 additions & 0 deletions amd/device-libs/oclc/src/isa_version_1251.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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;