Repository navigation
perf(navi32): enable Origami catalogs for TN HHS/BBS types - #11305
Merged
Merged
Conversation
…ne catalog
Three changes, all on the TN (Alik_Bljk) family for gfx1101:
1. LibraryType GridBased -> Prediction on all four TN catalogs. Origami is
compiled into every hipBLASLt build but never selects while a GridBased
table is present: Properties.py:68 ranks GridBasedMatching(3) above
PredictionMatching(2). Enablement is this field, not a flag.
2. One lean 100-solution HHS catalog now backs BBS, AuxH and AuxB. The four
TN ProblemTypes are separated by exactly two orthogonal edits to element
[4] and nothing else -- the ISA, the grid and all 100 solution dicts are
identical. The two edits are:
plain -> Aux UseE false -> true (1 field)
HHS -> BBS 7 dtype fields 4 -> 7, and
BiasDataTypeList [0,4] -> [0,7] (8 fields)
Rewriting only DataType/DataTypeA/DataTypeB and leaving DestDataType,
MacDataTypeA/B and DataTypeE at 4 yields an HSS-like hybrid that still
loads and builds, so all eight must move together.
3. gfx1101 registered in origami's arch switches (hardware.cpp) and given
calibrated memory constants (hardware.hpp). The registration is not
cosmetic: without a gfx1101 case, cus_per_multiProcessorCount falls to
default and returns 1, so Origami reads multiProcessorCount (27 WGPs) as
27 physical CUs instead of 54 and every occupancy and oversubscription
term is computed against half the machine. has_MALL() has no default
case at all, so an unlisted arch runs off the end of a non-void function.
The constants themselves are measured on this part rather
measured on this part rather than inherited from gfx1100. Navi 32 halves
the bus (384-bit -> 192-bit)
at essentially the same memory clock, and the model scales bandwidth by
clock only, so the inherited values were ~2.4x optimistic on DRAM and
~2.7x on MALL.
Build gate, all four types: 100 solutions requested, 100 serialized, 0
assembler errors, 0 overflowedResources. Tensile regenerates the runtime key
from the edited ProblemType and emits BB_BB_...Type_BB and the _Aux_ segment,
which independently confirms the edits landed.
Correctness: hipblaslt-bench -v over 60 stratified shapes per ProblemType,
norm_error against per-row rtol. HHS, BBS and AuxB pass 60/60. AuxH passes
59/60; the one miss (10x46x1155, norm 5.71e-05 vs rtol 1e-06) reproduces
identically on the pre-existing per-ProblemType catalog, so it is a
tolerance artifact rather than a regression introduced here.
Measured on Radeon RX 7700 XT, 500 shapes x 3 reps, geomean of per-shape
us ratios against the catalog shipped on develop, A/A noise floor 2.6-2.7%:
catalog only + Origami
TN HHS 1.160 1.212
TN BBS 1.170 1.213
Aux timings not yet collected.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #11305 +/- ##
===========================================
+ Coverage 69.39% 69.63% +0.24%
===========================================
Files 2841 2802 -39
Lines 467531 461755 -5776
Branches 68912 68072 -840
===========================================
- Hits 324434 321520 -2914
+ Misses 119499 116938 -2561
+ Partials 23598 23297 -301
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
Contributor
|
On the four TN catalogs (e.g. (Posting as a general comment because the YAML diff is too large for an inline comment.) |
added 2 commits
August 28, 2026 13:37
The C++ architecture_t gained gfx1101 but bindings.cpp was not updated, so the Python package could not name the arch and Navi 32 calibration/tests on the Python side had no way to construct a gfx1101 hardware_t.
…ami/ The four TN catalogs declare LibraryType: Prediction but lived under navi32/GridBased/. The loader keys on the field rather than the path so it worked, but the location was misleading. Origami/ is the existing convention for Prediction logic (gfx950/gfx950/Origami/, gfx1250/gfx1250/Origami/). Pure renames; file contents and names are unchanged.
Register gfx1101 in the HARDWARE fixture so test_ranking_regression.py covers it alongside the other supported architectures, and add the generated ranking baseline. gfx1101 exposes WMMA for Half/BFloat16/Int8/Int4 only, so 8 of the 20 dtype/transpose combinations run and the remaining 12 skip. Fixture values are those of the RX 7700 XT the gfx1101 constants were microbenchmarked on: 54 CU, 4 MB L2, 2276 MHz. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Contributor
Author
created origami folder and moved logic files there |
bethune-bryant
approved these changes
Aug 28, 2026
bethune-bryant
left a comment
bethune-bryant
left a comment
Contributor
There was a problem hiding this comment.
This looks good to me!
Thank you @vuksan314 !
vuksan314
enabled auto-merge (squash)
August 28, 2026 17:33
Contributor
|
Updated the branch to pick up develop (including #11442), which fixes the unrelated |
Contributor
TheRock Submodule Bump ActivityNewest first
|
vuksan314
added a commit
that referenced
this pull request
Oct 12, 2026
…13387) JIRA ID: AIRADSW-483 ## Motivation On gfx1100 (Radeon RX 7900 XTX, "navi31"), hipBLASLt chooses the GEMM kernel for fp16 and bf16 NN matrix multiplies from GridBased libraries. A GridBased library is a lookup table that maps tuned problem sizes to kernels; sizes it was not tuned for fall back to the nearest entry. The NN tables are thin (471 rows each, against about 9 700 for TN), so most NN sizes take a fallback kernel. This PR replaces the four NN tables with Origami Prediction libraries, the approach navi32 adopted in #11305. A Prediction library holds a short list of kernels, and the Origami analytical model ranks them for each problem at run time. On fresh problem sizes held out of tuning, the new libraries are faster in geomean (the mean of per-shape speedups) for every type: - fp16 NN: 34% faster off the stock table, 17% faster on the stock table's own rows. - bf16 NN: 41% faster. - Aux (UseE) NN: 13–17% faster. - On a blind 194-shape benchmark set that was never used for tuning, fp16 NN wins 19 of 20 categories: geomean 1.334, +1.54M GFLOPS net. The exceptions are listed under Known limits. The TN libraries are not touched by this PR. ## Technical Details Directory: `projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/navi31/`. Each new file in `Origami/` replaces the `GridBased/` file of the same name. Both files describe the same problem type, so they cannot be in one library together. There are no code changes; the last commit adds one CHANGELOG line. | file (`navi31_Cijk_…_UserArgs.yaml`) | types | before | after | |---|---|---|---| | `Ailk_Bljk_HHS_BH_Bias_HAS_SAV` | fp16 in/out, fp32 compute | 70 kernels + 471-row table | 42 kernels | | `Ailk_Bljk_BBS_BH_Bias_HAS_SAV` | bf16 | 73 kernels + 471-row table | 42 (port of HHS-NN) | | `Ailk_Bljk_HHS_BH_Bias_AuxH_HAS_SAV` | fp16, also writes E | 70 kernels + 471-row table | 39 (port, StreamK→0) | | `Ailk_Bljk_BBS_BH_Bias_AuxB_HAS_SAV` | bf16, also writes E | 73 kernels + 471-row table | 39 (port, StreamK→0) | **Ports.** - The bf16 and Aux files are the fp16 catalog with the whole problem type copied from the stock file of the target type. - Aux means UseE: the GEMM also writes the pre-activation output E. With UseE, StreamK=3 kernels fail code generation on this tree, so those members were converted to StreamK=0. One converted member duplicated an existing one and was removed (40 → 39). - The Aux files are based on the 40-kernel HHS-NN catalog with the member-13 rollback. They do not include the narrow-N update, whose new StreamK=3 kernel cannot be used with UseE. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
This was referenced Oct 21, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
JIRA ID : ROCM-1682
Enable Origami on Navi32 and add TN HHS(BBS)
performance vs current gridbased
| catalog | predictor | tiny | small | medium | large | ALL |
| old | GridBased | 1.000 | 1.000 | 1.000 | 1.000 | 1.000 |
| new | Origami | 0.979 | 1.160 | 1.305 | 1.456 | 1.212 |
| catalog | predictor | gemv | skinny_M | skinny_N | rect | square | ALL |
| old | GridBased | 1.000 | 1.000 | 1.000 | 1.000 | 1.000 | 1.000 |
| new | Origami | 1.338 | 1.176 | 1.184 | 1.182 | 1.189 | 1.212 |