[schmittcn]

Problem and result

The rocBLAS tables in tools/hip cover gfx1030 with rocBLAS 5.6.0.8d1ae90e (ROCm 10.0.0). #3343 adds 5.4.0.ad84daa81c
(TheRock 7.13). Arch Linux packages ROCm 7.2.4 (rocblas 7.2.4-3). Its rocBLAS reports 5.2.0.dabb6df2b98-dirty, so
an engine built against the system ROCm on Arch finds no table for its build and runs rocBLAS's default kernels.

This PR adds that table. It was made with tools/hip/tune_rocblas on an RX 6900 XT against that rocBLAS, and
docs/AMD_HIP.md now names it. With it, prompts below 1,152 tokens read 12-18% faster on the GPU timeline. Prompts of
1.3K-1.5K tokens read 2-4% faster, and prompts of 11K-12K tokens 4-9% faster. No code changes: setup already points
STRATA_ROCBLAS_TUNING at tools/hip on gfx103x, and the engine takes the file of its own build.

Related discussion

#3343, the same kind of table for TheRock's rocBLAS 5.4.0. This PR is independent of it, and the two files do not
overlap. #981 is the table mechanism.

Approach

LD_LIBRARY_PATH=/opt/rocm/lib build-hip/tune_rocblas --tuning-out tools/hip/gfx1030-rocblas-5.2.0.dabb6df2b98-dirty.txt
with the defaults. It wrote 112 rows: 78 with a solution, 34 default. Each solution it kept matched the default
kernel's output (relative L2 about 2e-5, no padding writes). The full tuner output, the raw timing lines, the answers
and the sanitized config are in bench/results/2026-10-24-rdna2-rocblas-table-rocm72/. The measurement follows the
method of bench/results/2026-10-05-rdna2-rocblas-table.

Validation

Environment:

• RX 6900 XT 16 GB (gfx1030), PCIe 4.0 x16, Ryzen 9 5950X, 64 GB, Arch Linux, kernel 7.2.8.
• ROCm 7.2.4 from Arch's packages, lib_dirs /opt/rocm/lib.
• Engine 0.1.48.1, a HIP build of 7ddc17a3. Its only local changes are in src/qwen35, which the measured model does
  not use.
• Model: Qwen3.8-Flash-Next GSQ-RCO Coder IQ1_M, with setup's config of this PC (--adapt-every 100000 in both arms).

Passed:

• A/B on the GPU timeline: one server per arm, the same binary. The only difference is STRATA_ROCBLAS_TUNING
  (unset or tools/hip). Each arm got 12 prompts (685-12,164 tokens), each from its own offset of
  docs/DETAILS.md, with temperature 0 and max_tokens 8. Each arm ran three times as a fresh server. The GPU
  timeline comes from strata prefill timing. Mean change:
  • 685 / 776 / 977 / 985 / 1,084 tokens: -16.4 / -12.9 / -18.0 / -15.3 / -12.1%.
  • 1,264 / 1,353 / 1,510 / 1,543 tokens: -1.6 / -3.6 / -3.7 / -1.9%.
  • 10,917 / 11,396 / 12,164 tokens: -7.6 / -8.6 / -4.2%.
  • The runs are in the report's table. The largest spread within one arm is about 200 ms.
• Solutions used: the engine log shows the table taken (rocBLAS tuning enabled (112 rows, 78 with a solution, gfx1030, rocBLAS 5.2.0.dabb6df2b98-dirty)). Each on-arm run reports tuned_launches=6268 fallbacks=0.
• Determinism: each arm gives the same 12 answers when repeated (runs 2 and 3 compared).

Failed:

• Answers between the arms: 11 of 12 are the same. One (1,089 tokens) differs at its 6th reasoning token,
  "request:" against "request.". This is the FP16 rounding of a different kernel; hip: a rocBLAS solution table for the FP16 prompt GEMMs on gfx103x (the GDN projection 7x below 1,152 tokens; short prompts -10 to -14%) #981's report saw identical answers
  on its four prompts.

Not run:

• Other gfx103x cards (gfx1031 / gfx1032).
• Other distributions' ROCm 7.2.4 builds, which may report another rocBLAS version string. The engine then refuses the
  table and says so in its log.
• Decode, which does not use these GEMMs.

Limits

One machine, one model. The table is valid for this one rocBLAS build only.

• Measured on one card, where these dense FP16 GEMMs are the prompt path. On a two-card layer split with AMD's
  Ubuntu ROCm 7.2.4, a table of this kind loaded and was used but made no end-to-end difference at 4K and 16K
  prompts (expert streaming and the hand-off dominate there; see the comment by @DylanFromMassey).
• AMD's Ubuntu packages of the same ROCm 7.2.4 report 5.2.0.dabb6df2b9, so this file does not load there; that
  build needs its own table. docs/AMD_HIP.md now says so and recommends pointing STRATA_ROCBLAS_TUNING at the
  directory (setup's default) so each GPU of a split takes its own file.

🤖 Generated with Claude Code

[DylanFromMassey]

One more data point, and it confirms the caveat in your "Not run" list.

Box: RX 6900 XT (gfx1030), system ROCm 7.2.4 from AMD's Ubuntu repo (not Arch), lib_dirs /opt/rocm/lib, engine 0.1.50, two cards with a layer split. Its rocBLAS reports

5.2.0.dabb6df2b9

— i.e. the same 5.2.0 base, a different build string from Arch's 5.2.0.dabb6df2b98-dirty. So our two files are not interchangeable: the engine keys the table to the exact string, so mine will be refused on your Arch build and yours on mine, each with a rocBLAS version mismatch line in the log (which is the good failure mode — it says so rather than running slow).

Our table: gfx1030-rocblas-5.2.0.dabb6df2b9.txt, 112 rows, 75 with a solution (yours: 78/112 — close enough that I suspect the same 14 dense FP16 shapes × T = 256…8192 sweep), made with

HIP_VISIBLE_DEVICES=0 build-hip/tune_rocblas --tuning-out tools/hip/gfx1030-rocblas-5.2.0.dabb6df2b9.txt

and it loads: rocBLAS tuning enabled (112 rows, 75 with a solution, gfx1030, rocBLAS 5.2.0.dabb6df2b9).

Where I differ from you on the effect, and why it matters for what people should expect. In isolation the table is real — one shape measured 6.9× faster with the tuned solution, counters proving use (tuned_launches=259 fallbacks=0). But end-to-end on this box it bought nothing: 4K prefill 593.4 vs 597.9 tok/s and 16K 792.1 vs 806.9 without it. The reason is the workload shape, not the table: this box serves as a layer split, so the gfx1030 stage holds layers 0–24 (or 25–47 depending on order) and its prompt time is dominated by expert streaming and the hand-off, not by dense GEMMs. Your −12 to −18 % at sub-1 152-token prompts is measured on a single-card helper tier where those GEMMs are the prompt path — a genuinely different regime, and I think both results are correct.

So: worth keeping for correctness and headroom, not a prefill feature. One deployment note that cost me time — install it as a directory, not a single file:

mkdir -p tools/hip/rocblas-f16o-5.2.0
cp tools/hip/gfx1030-rocblas-5.2.0.dabb6df2b9.txt tools/hip/gfx1101-rocblas-5.2.0.dabb6df2b9.txt tools/hip/rocblas-f16o-5.2.0/
# STRATA_ROCBLAS_TUNING=tools/hip/rocblas-f16o-5.2.0

STRATA_ROCBLAS_TUNING accepts a directory and then each stage loads <arch>-rocblas-<version>.txt for itself. With a single file path, two different-family stages are handed the same table and one of them is refused.

[schmittcn]

Thanks, that's a useful second data point.

• Version string: agreed; the engine keys on the full rocBLAS version (rocblas_tuning.hpp), so Ubuntu's 5.2.0.dabb6df2b9 and Arch's 5.2.0.dabb6df2b98-dirty need separate files. Yours would fit as its own PR next to this one; the files don't overlap.
• Effect: I agree both results hold. My numbers are one card, where these dense FP16 GEMMs are the prompt path; on your layer split, expert streaming and the hand-off dominate. The PR's Limits and the AMD_HIP.md line now say so.
• Directory: setup already sets STRATA_ROCBLAS_TUNING to the directory tools/hip, so the default picks per-arch files; the single-file trap only applies when it's set by hand. AMD_HIP.md now recommends the directory.