Skip to content

hip: gfx1100 hipBLASLt 1.5.0 tuning table (ROCm 10.2 nightly) — closes the gfx1100 100500 gap - #755

Closed
xyzzing wants to merge 1 commit into
Niko1221:mainfrom
xyzzing:table-gfx1100-hipblaslt-100500
Closed

xyzzing wants to merge 1 commit into
Niko1221:mainfrom
xyzzing:table-gfx1100-hipblaslt-100500

Conversation

@xyzzing

@xyzzing xyzzing commented Oct 4, 2026

Copy link
Copy Markdown
Contributor

Summary

Ship tools/hip/gfx1100-hipblaslt-100500.txt — a hipBLASLt tuning table for the RX 7900 XTX calibrated against hipBLASLt 1.5.0 (version number 100500), the same gap #481 closed for gfx1201. This is an engine-tree companion to the results-only #745 (its second commit measures exactly this table; per-run data lives there).

Why: the engine refuses a table whose version doesn't match the installed hipBLASLt, so on a ROCm 10.2 nightly (hipBLASLt 1.5.0) the shipped gfx1100-hipblaslt-100100/100200.txt are refused and prefill silently runs on plain hipBLAS. Every gfx1100 nightly user is on the untuned path today.

What: the same 16 dense GEMM geometries as the other gfx1100 tables (bf16 + f16, T = 4096 and 8192 → 32 rows), calibrated with tune_hipblaslt run against that exact nightly library (ROCm 10.2.0a20261003 SDK, libamdhip64.so.7.17.26392; every candidate gated at rel-L2 ≤ 1e-4 / max-abs ≤ 1e-2 / no padding writes; each row is its case's measured best). setup picks it up by name once the installed hipBLASLt reports 100500 — no config change needed.

Measured (0.1.38-lineage nightly binary of this tree, Qwen3.8-Flash-Next IQ3_S greedy, arms differing only in STRATA_HIPBLASLT_TUNING, median of 3 clean runs each):

prompt tokens with table plain hipBLAS gain
131071 1687 tok/s 926 tok/s 1.82x

Decode unchanged (62.3 vs 60.7 tok/s, overlapping ranges — the table is prefill-only, as designed). Per-run data, calibration command and tuner log: bench/results/2026-10-04-community-rx7900xtx-hipblaslt-100500/ in #745.

docs/AMD_HIP.md: one entry in the shipped-tables list, mirroring the gfx1201 1.5.0 entry (including the caveat that the version number is the only thing the engine can check — recalibrate with tune_hipblaslt for a different 1.5.0 build).

Checklist

…0a nightly (hipBLASLt 1.5.0)

tools/hip/gfx1100-hipblaslt-100500.txt: the same 16 dense GEMM geometries as
the other gfx1100 tables (bf16 and f16, T=4096 and T=8192, 32 rows), calibrated
on an RX 7900 XTX with tune_hipblaslt (relative L2 <= 1e-4, max abs <= 1e-2, no
padding writes) against a ROCm 10.2.0a20261003 nightly SDK
(libamdhip64.so.7.17.26392), hipBLASLt 1.5.0 (version number 100500).

The engine refuses a table of another hipBLASLt version, so on this ROCm the
shipped gfx1100 tables (100100, 100200) are refused and prefill runs on plain
hipBLAS - the same gap the gfx1201 1.5.0 table (Niko1221#481) closed for the R9700.
setup picks this one by name once the installed hipBLASLt reports 100500.

Measured on a 0.1.38-lineage nightly build of this tree, Qwen3.8-Flash-Next
IQ3_S greedy, arms differing only in STRATA_HIPBLASLT_TUNING (median of 3
clean runs each):

| prompt tokens | with table | plain hipBLAS | gain |
| 131071        | 1687 tok/s | 926 tok/s     | 1.82x |

decode unchanged (62.3 vs 60.7 tok/s, overlapping ranges). Per-run data,
the calibration command and the tuner log: Niko1221#745
(bench/results/2026-10-04-community-rx7900xtx-hipblaslt-100500/).

docs/AMD_HIP.md: the shipped-tables list carries the entry.
@Niko1221

Niko1221 commented Oct 6, 2026

Copy link
Copy Markdown
Owner

Merged for 0.1.40 (gfx1100 hipBLASLt table 100500).

@Niko1221 Niko1221 closed this Oct 6, 2026
Niko1221 pushed a commit that referenced this pull request Oct 7, 2026
…5 (PR #917)

Interactive Q2_0 and UD-IQ4_XS chats against llama.cpp Vulkan, then controlled A-E runs of the gfx1151 hipBLASLt tuning table
(250.9 to 528.9 prompt tok/s), dense MMQ (+0.2%) and the shared-expert decode stream (+2.5%, 8 requests). The build is
0.1.39 plus then-unmerged PRs (#895, #820, a tuning table adapted from #755), not the 0.1.40 release; the index line in
COMMUNITY_BENCHMARKS.md says so and notes that 0.1.40 turned the shared-expert stream off on HIP (#826).
0xrjman pushed a commit to 0xrjman/Strata that referenced this pull request Oct 7, 2026
…5 (PR Niko1221#917)

Interactive Q2_0 and UD-IQ4_XS chats against llama.cpp Vulkan, then controlled A-E runs of the gfx1151 hipBLASLt tuning table
(250.9 to 528.9 prompt tok/s), dense MMQ (+0.2%) and the shared-expert decode stream (+2.5%, 8 requests). The build is
0.1.39 plus then-unmerged PRs (Niko1221#895, Niko1221#820, a tuning table adapted from Niko1221#755), not the 0.1.40 release; the index line in
COMMUNITY_BENCHMARKS.md says so and notes that 0.1.40 turned the shared-expert stream off on HIP (Niko1221#826).
xyzzing added a commit to xyzzing/Strata that referenced this pull request Oct 13, 2026
… — median of 3 clean cells per tier, the shipped config

Engine 0.1.43.1 (refresh of the 0.1.38-era row), ROCm 10.2 nightly, the tuned gfx1100
hipBLASLt-100500 table (Niko1221#755), Ryzen 9 7900X / 96 GB, one-shot greedy runs, 256 generated
tokens, median of 3 clean cells per tier (the recorded methodology). Decode now carries
STRATA_ROUTE_TAIL_SKIP=7 (adopted after a task-battery-gated A/B; answers differ slightly,
=0 restores) — the decode row moves 59.2-65.1 -> 71.4-81.3 tok/s; prefill 32K-128K moves
1,494-1,641 -> 2,117-2,149 tok/s.
xyzzing added a commit to xyzzing/Strata that referenced this pull request Oct 13, 2026
… — median of 3 clean cells per tier, the shipped config

Engine 0.1.43.1 (refresh of the 0.1.38-era row), ROCm 10.2 nightly, the tuned gfx1100
hipBLASLt-100500 table (Niko1221#755), Ryzen 9 7900X / 96 GB, one-shot greedy runs, 256 generated
tokens, median of 3 clean cells per tier (the recorded methodology). Decode now carries
STRATA_ROUTE_TAIL_SKIP=7 (adopted after a task-battery-gated A/B; answers differ slightly,
=0 restores) — the decode row moves 59.2-65.1 -> 71.4-81.3 tok/s; prefill 32K-128K moves
1,494-1,641 -> 2,117-2,149 tok/s.
xyzzing added a commit to xyzzing/Strata that referenced this pull request Oct 13, 2026
… — median of 3 clean cells per tier, the shipped config

Engine 0.1.43.1 (refresh of the 0.1.38-era row), ROCm 10.2 nightly, the tuned gfx1100
hipBLASLt-100500 table (Niko1221#755), Ryzen 9 7900X / 96 GB, one-shot greedy runs, 256 generated
tokens, median of 3 clean cells per tier (the recorded methodology). Decode now carries
STRATA_ROUTE_TAIL_SKIP=7 (adopted after a task-battery-gated A/B; answers differ slightly,
=0 restores) — the decode row moves 59.2-65.1 -> 71.4-81.3 tok/s; prefill 32K-128K moves
1,494-1,641 -> 2,117-2,149 tok/s.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants