amd-agi/geak/tuning-hipblaslt

Tuning hipBLASLt

Read ../tuning-core/SKILL.md first. Build the client per ../env-setup/hipblaslt_bench_build.md — no package ships it.

hipBLASLt is the default GEMM backend under torch on ROCm, so it is usually what you are implicitly benchmarking against. Unlike Triton, you do not author a config space: hipBLASLt ships thousands of pre-compiled solutions and tuning means selecting among them. The whole job is (1) state the problem exactly as the workload states it, (2) race, (3) carry the winner's identity forward correctly.

Scope. Selection is all of this skill. The layer beneath it — tensilelite/ in the hipBLASLt source, which generates and assembles the solutions in the first place — is not covered. That ordering is deliberate: selection is a day's work with immediate payoff, adding a Tensile kernel is a different kind of project. But on MI355 it is worth knowing the deeper layer exists, because §7 shows the shipped selection logic is now mostly analytical, and §3b shows how little racing buys on top of it.

Measured on gfx942 / MI300X and gfx950 / MI355X, ROCm 7.2.2, bf16. Numbers are labelled per part; where only one is given it is gfx942. Build with -a gfx950 for MI355 (../env-setup/hipblaslt_bench_build.md) — ~10 min, and the resulting binary races 2085 solutions rather than gfx942's 1231.

1. Capture the problem, do not reconstruct it

A hipBLASLt "problem" is more than M/N/K: it includes transposes, leading dimensions, strides, dtypes for A/B/C/D, compute type, bias, and epilogue. Tune the wrong one and you have tuned nothing.

Get it from the workload rather than writing it by hand:

export HIPBLASLT_LOG_MASK=32          # log every matmul as a runnable bench command
python3 your_workload.py 2> hipblaslt.log
grep -o '\-\-api_method.*' hipblaslt.log | sort | uniq -c | sort -rn | head

Each line is a complete hipblaslt-bench invocation. Replay it verbatim.

The reason for the emphasis: a torch matmul of (512,1024) × (1024,2048) dumps -m 2048 -n 512 -k 1024. M and N are swapped relative to the torch call, because the library sees the column-major transposed form. Hand-reconstructing "M=512, N=2048" tunes a different problem that happens to have the same FLOP count. Copy the captured command.

The sort | uniq -c matters as much as the capture: it ranks problems by call count, which is your tuning priority order. Tuning the shape that appears 40,000 times beats tuning the one that appears twice, whatever their individual timings.

2. Race all solutions

hipblaslt-bench -m 4096 -n 4096 -k 4096 -r bf16_r --transA T --transB N \
    --algo_method all --print_kernel_info -i 20 -j 5

Output ends with a Winner: block: a CSV row (Gflops, GB/s, us), a --Solution index:, and the solution/kernel names. On the 4096³ bf16 case above the winner was index 205610, MT256x224x64, 633505 Gflops (633 TFLOPS).

The solution name is readable and worth reading. MT256x224x64 is the macro-tile; MI16x16x1 the MFMA instruction shape; ISA942 the target arch; GSU/SK fields describe split-K. Comparing the winners across regimes tells you why one won — the 4096³ winner takes a 256×224 tile, while decode-shaped problems win with tiles like MT16x16x512.

The ISA field is also the reason a solution index never crosses parts. The same command on gfx950 / MI355X:

Two changes worth separating. The tile got deeper (224 → 256), which is a statement about the part. The index changed entirely, which is not — these are separately compiled kernel libraries, and 205610 on gfx950 either names something unrelated or nothing at all. Carry the tile across parts as a hint; re-derive the index.

3. The bracket number is not the solution index

The most expensive small trap here.

Every result line is prefixed with a bracket number — [138], [1357]. That is the enumeration position in this run's output, not an identifier. Feeding it back fails:

--algo_method index --solution_index 138
# error: NO solution found! ... testing_matmul.hpp:2815

The real identifier comes only from --print_kernel_info:

--Solution index: 205610

which replays correctly:

hipblaslt-bench -m 4096 -n 4096 -k 4096 -r bf16_r --transA T --transB N \
    --algo_method index --solution_index 205610 -i 50 -j 10
# Is supported 1 / Total solutions: 1   -> 598810 Gflops

Worse than failing loudly: the bracket number drifts between identical runs (the same winner appeared as [99] and [66] across two runs of one command). So a bracket number recorded from one run can silently name a different solution later. If you did not run with --print_kernel_info, you do not have a result you can deploy — rerun.

Confirmed on gfx950: the Winner: block is prefixed [1910] while its --Solution index: is 441281. The two are not related, and note that --Solution index: is printed after the CSV row it belongs to, so a parser that reads forward from the bracket line pairs every row with the previous row's index. (One place this bites: ../env-setup/hipblaslt_bench_build.md used to state the bracket is the solution index. It is not; that has been corrected.)

Replay the index standalone. On gfx950 this is not a formality.

Always close the loop: replay the winner by index and confirm it reproduces the winning time.

On gfx942, 5.5% is run-to-run spread and the race number is usable. On gfx950 a fifth of the reported throughput is not spread — it is the race inflating itself. Racing 2085 solutions back-to-back keeps the clocks boosted and the caches warm for the whole run; a single solution replayed alone starts from a colder state, which is also the state your workload starts from.

Deploy against the replayed number. If you size a decision on the raced figure on MI355 you will over-predict by ~20%, and if you compare a raced hipBLASLt number against a non-raced number from another backend you will conclude hipBLASLt wins by more than it does. This is the same effect as ../tuning-core/measurement.md Rule 6b, arriving through a vendor tool rather than your own loop.

3b. What the race is actually worth on MI355

The number you care about is not the race figure. It is how much the raced winner beats the solution the library would have chosen on its own — --algo_method heuristic, the default, which is what your workload gets if you tune nothing. Both sides have to be replayed, and on gfx950 they have to be interleaved. Four bf16 shapes, -i 50 -j 10, five alternating rounds each, medians:

(TFLOPS. "claimed" = the race's own Gflops against a separately-run default, i.e. what you would have written down had you not replayed.)

Two things to take from this. The apparent 19–40% uplift is almost entirely measurement artefact; the real range is −2% to +14%. And on one of four shapes the raced winner is not better than the default — the race dutifully reported a winner, because a race always has a winner, and adopting it would have cost 2%. --algo_method all does not tell you whether tuning was worth doing. Only the interleaved A/B against the heuristic does.

That the shipped default holds up this well on MI355 is not luck, and §7 explains where it comes from.

4. Rotating buffers, or you will tune the cache

A benchmark loop reuses the same A and B buffers every iteration. After the first, they are resident in cache — and the timing measures a memory hierarchy your real workload never sees.

Same solution, same shape, 1024³ bf16 (small enough to fit in cache):

Inflation from cache reuse alone. Not measurement noise — a systematically wrong answer that flatters every small shape and can reorder which solution appears to win.

The magnitude is very different between parts, and the direction of the surprise is worth noting: on gfx942 the cache is the dominant measurement error at this shape and on gfx950 it is nearly negligible, while the clock error (§3, 20%) is far larger there. So keep using --rotating on both — it costs nothing and the effect is real on both — but do not assume the error budget you learned on one part describes the other. Re-measure which artefact dominates before deciding what to control for.

Use --rotating with a working set larger than last-level cache for any shape whose operands fit in cache. Large shapes are less affected but the flag costs nothing. This is the concrete case of the cache-invalidation rule in ../tuning-core/measurement.md.

5. Deploying the winner

An index in your notebook changes nothing. Three ways to make it real:

• Through aiter — run the aiter gradlib tuner, which races hipBLASLt against the other

backends and, when hipBLASLt wins, records libtype=hipblaslt plus the solution index in a tuned CSV that the runtime reads. This is the path that matters for serving. There is no flag to force the backend — --libtype is an output column, not an input, and documenting it as an option here was wrong. See ../tuning-aiter/ §4, including where the CSV must be written.

• Through torch TunableOp — torch.cuda.tunable tunes and persists selections for torch's

own matmuls, writing a CSV whose metadata includes GCN_ARCH_NAME.

• Directly in application code — via the hipBLASLt extension API, when you own the call

site.

In all three, finish with ../tuning-core/engagement_verification.md. A solution index that is fast in hipblaslt-bench proves the solution exists and is fast; it says nothing about whether your workload selects it.

6. Sanity-check across tools

Independent agreement is the cheapest validation available. On 4096³ bf16 the aiter tuner independently selected an MT256x224x64 hipBLASLt solution — the same macro-tile family this bench chose. Two tools with separate search procedures landing on the same tile shape is strong evidence the result is real rather than an artifact of one harness.

When tools disagree by more than the noise floor, do not average them. One is measuring something you did not intend — check layout and transpose first (../tuning-core/measurement.md, rule 6).

6b. MX (MXFP4 / MXFP8): everything in §2–§3 stops working