Skip to content

perf(navi32): enable Origami catalogs for TN HHS/BBS types #11305

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
vuksan314 merged 8 commits into from
Aug 31, 2026
Merged

perf(navi32): enable Origami catalogs for TN HHS/BBS types #11305

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
2055c4c
perf(navi32): enable Origami on TN and serve all four TN types from o…
Aug 26, 2026
7dc114f
Merge branch 'develop' into vmijovic/navi32_tn_origami
vuksan314 Aug 26, 2026
3a380ae
Merge branch 'develop' into vmijovic/navi32_tn_origami
vuksan314 Aug 27, 2026
d01d51f
fix(origami): expose gfx1101 in the nanobind architecture_t enum
Aug 28, 2026
62fbac2
refactor(hipblaslt): move navi32 TN Prediction logic into navi32/Orig…
Aug 28, 2026
5efb5bd
test(origami): add gfx1101 to Python ranking regression baselines
Aug 28, 2026
fb65e7e
Merge branch 'develop' into vmijovic/navi32_tn_origami
vuksan314 Aug 28, 2026
3373e36
Merge branch 'develop' into vmijovic/navi32_tn_origami
bethune-bryant Aug 31, 2026
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
12,991 changes: 0 additions & 12,991 deletions ...ic/asm_full/navi32/GridBased/navi32_Cijk_Alik_Bljk_BBS_BH_Bias_AuxB_HAS_SAV_UserArgs.yaml
View file
Open in desktop

This file was deleted.

12,991 changes: 0 additions & 12,991 deletions ...e/Logic/asm_full/navi32/GridBased/navi32_Cijk_Alik_Bljk_BBS_BH_Bias_HAS_SAV_UserArgs.yaml
View file
Open in desktop

This file was deleted.

14,674 changes: 0 additions & 14,674 deletions ...ic/asm_full/navi32/GridBased/navi32_Cijk_Alik_Bljk_HHS_BH_Bias_AuxH_HAS_SAV_UserArgs.yaml
View file
Open in desktop

This file was deleted.

14,674 changes: 0 additions & 14,674 deletions ...e/Logic/asm_full/navi32/GridBased/navi32_Cijk_Alik_Bljk_HHS_BH_Bias_HAS_SAV_UserArgs.yaml
View file
Open in desktop

This file was deleted.

31,150 changes: 31,150 additions & 0 deletions ...ogic/asm_full/navi32/Origami/navi32_Cijk_Alik_Bljk_BBS_BH_Bias_AuxB_HAS_SAV_UserArgs.yaml
View file
Open in desktop

Large diffs are not rendered by default.

31,150 changes: 31,150 additions & 0 deletions ...ile/Logic/asm_full/navi32/Origami/navi32_Cijk_Alik_Bljk_BBS_BH_Bias_HAS_SAV_UserArgs.yaml
View file
Open in desktop

Large diffs are not rendered by default.

31,150 changes: 31,150 additions & 0 deletions ...ogic/asm_full/navi32/Origami/navi32_Cijk_Alik_Bljk_HHS_BH_Bias_AuxH_HAS_SAV_UserArgs.yaml
View file
Open in desktop

Large diffs are not rendered by default.

31,150 changes: 31,150 additions & 0 deletions ...ile/Logic/asm_full/navi32/Origami/navi32_Cijk_Alik_Bljk_HHS_BH_Bias_HAS_SAV_UserArgs.yaml
View file
Open in desktop

Large diffs are not rendered by default.

17 changes: 17 additions & 0 deletions shared/origami/include/origami/hardware.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,6 +59,7 @@ class ORIGAMI_EXPORT hardware_t {
gfx1200,
gfx1201,
gfx1100,
gfx1101,
Comment thread
vuksan314 marked this conversation as resolved.
Resolved
gfx1150,
gfx1151,
gfx1152,
Expand All @@ -80,6 +81,7 @@ class ORIGAMI_EXPORT hardware_t {
if (str == "gfx1200") return architecture_t::gfx1200;
if (str == "gfx1201") return architecture_t::gfx1201;
if (str == "gfx1100") return architecture_t::gfx1100;
if (str == "gfx1101") return architecture_t::gfx1101;
if (str == "gfx1150") return architecture_t::gfx1150;
if (str == "gfx1151") return architecture_t::gfx1151;
if (str == "gfx1152") return architecture_t::gfx1152;
Expand All @@ -102,6 +104,7 @@ class ORIGAMI_EXPORT hardware_t {
case architecture_t::gfx1200: return "gfx1200";
case architecture_t::gfx1201: return "gfx1201";
case architecture_t::gfx1100: return "gfx1100";
case architecture_t::gfx1101: return "gfx1101";
case architecture_t::gfx1150: return "gfx1150";
case architecture_t::gfx1151: return "gfx1151";
case architecture_t::gfx1152: return "gfx1152";
Expand Down Expand Up @@ -199,6 +202,9 @@ class ORIGAMI_EXPORT hardware_t {
return {5.74, 1.21875121875121875122 * 2.41, 0.464, 2, std::make_tuple(0, 0.17, 0), 1.5};
case architecture_t::gfx1100:
return {7.12, 1.21875121875121875122 * 3.48, 0.732, 2, std::make_tuple(0, 0.11, 0), 1.5};
case architecture_t::gfx1101:
// AMD Navi 32
return {4.86133, 1.21875121875121875122 * 1.28305, 0.420583, 2, std::make_tuple(0, 0.13169, 0), 1.5};
case architecture_t::gfx1150:
// AMD Strix Point iGPU
return {1.497, NO_MALL_AVAILABLE, 0.077, 16, std::make_tuple(0, 0.18, 0), 1.5};
Expand Down Expand Up @@ -480,6 +486,17 @@ class ORIGAMI_EXPORT hardware_t {
// I4
{matrix_instruction(16, 16, 16, data_type_t::Int4), 16}, // v_wmma_i32_16x16x16_iu4
}},
{architecture_t::gfx1101,
{
// F16
{matrix_instruction(16, 16, 16, data_type_t::Half), 32}, // v_wmma_f32_16x16x16_f16/v_wmma_f16_16x16x16_f16
// BF16
{matrix_instruction(16, 16, 16, data_type_t::BFloat16), 32}, // v_wmma_f32_16x16x16_bf16/v_wmma_bf16_16x16x16_bf16
// I8
{matrix_instruction(16, 16, 16, data_type_t::Int8), 32}, // v_wmma_i32_16x16x16_iu8
// I4
{matrix_instruction(16, 16, 16, data_type_t::Int4), 16}, // v_wmma_i32_16x16x16_iu4
}},
{architecture_t::gfx1150,
{
// F16
Expand Down
1 change: 1 addition & 0 deletions shared/origami/python/src/origami/bindings.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,7 @@ NB_MODULE(origami, m) {
.value("gfx1200", hardware_t::architecture_t::gfx1200)
.value("gfx1201", hardware_t::architecture_t::gfx1201)
.value("gfx1100", hardware_t::architecture_t::gfx1100)
.value("gfx1101", hardware_t::architecture_t::gfx1101)
.value("gfx1150", hardware_t::architecture_t::gfx1150)
.value("gfx1151", hardware_t::architecture_t::gfx1151)
.value("gfx1152", hardware_t::architecture_t::gfx1152)
Expand Down
1,610 changes: 1,610 additions & 0 deletions shared/origami/python/tests/baselines/rankings/gfx1101.yaml
View file
Open in desktop

Large diffs are not rendered by default.

3 changes: 3 additions & 0 deletions shared/origami/python/tests/helpers.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,9 @@
"gfx1100": origami.get_hardware_for_arch(
origami.architecture_t.gfx1100, 96, 64 * 1024, 512 * 1024, 6 * 1024 * 1024, 2500000
),
"gfx1101": origami.get_hardware_for_arch(
origami.architecture_t.gfx1101, 54, 64 * 1024, 512 * 1024, 4 * 1024 * 1024, 2276000
),
"gfx1200": origami.get_hardware_for_arch(
origami.architecture_t.gfx1200, 32, 128 * 1024, 512 * 1024, 4 * 1024 * 1024, 2700000
),
Expand Down
4 changes: 4 additions & 0 deletions shared/origami/src/origami/hardware.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -107,6 +107,7 @@ namespace {
size_t cus_per_multiProcessorCount(hardware_t::architecture_t arch) {
switch (arch) {
case hardware_t::architecture_t::gfx1100: // RDNA3
case hardware_t::architecture_t::gfx1101: // RDNA3 (Navi 32)
case hardware_t::architecture_t::gfx1150: // RDNA3.5 (Strix)
case hardware_t::architecture_t::gfx1151:
case hardware_t::architecture_t::gfx1152:
Expand Down Expand Up @@ -217,6 +218,7 @@ size_t hardware_t::get_default_num_xcds(architecture_t arch) {
case architecture_t::gfx1200: return 1;
case architecture_t::gfx1201: return 1;
case architecture_t::gfx1100: return 1;
case architecture_t::gfx1101: return 1;
case architecture_t::gfx1150: return 1;
case architecture_t::gfx1151: return 1;
case architecture_t::gfx1152: return 1;
Expand Down Expand Up @@ -299,6 +301,7 @@ bool hardware_t::has_MALL() const {
case architecture_t::gfx1200:
case architecture_t::gfx1201:
case architecture_t::gfx1100:
case architecture_t::gfx1101:
case architecture_t::gfx1151:
return true;
case architecture_t::gfx1150:
Expand All @@ -320,6 +323,7 @@ bool hardware_t::has_native_TF32() const {
case architecture_t::gfx1200:
case architecture_t::gfx1201:
case architecture_t::gfx1100:
case architecture_t::gfx1101:
case architecture_t::gfx1150:
case architecture_t::gfx1151:
case architecture_t::gfx1152:
Expand Down
Loading