docs: rocBLAS on gfx1151 — both runtime routes closed, and why

[glennneuber]

Closes out the question I raised at the end of the MMQ vs hipBLAS A/B: could rocBLAS offer better kernels on gfx1151? It was worth asking. The answer is no, and the reasons are more useful than the answer.

All measured on gfx1151, build 0.34.2-dynres-f67b1aef, ROCm 7.2.4. Noise floor ±0.4%.

1. These GEMMs never reach should_use_mmq

ggml_cuda_mul_mat bails to cuBLAS at its first branch:

if (bad_padding_clear || src1->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
    ggml_cuda_mul_mat_cublas(ctx, src0, src1, dst);
    return;   // bypasses mmvf, mmf, mmvq and mmq entirely
}

Every captured call has --d_type f16_r, so every one takes this exit.

This corrects how the earlier A/B should be read. GGML_CUDA_FORCE_CUBLAS only changes what should_use_mmq returns, so it cannot move any of these GEMMs — they were on rocBLAS in both arms. That is why dense gemma4:31b came out at +2.1% instead of showing an MMQ win: a large share of its prefill never had a choice of path. The A/B is still valid for what it measured; it simply never covered this.

2. It is not a vision-tower story

I had been calling this "the tower". Measured, it isn't:

model	calls	total TFLOP	tower	LM	other
qwen3.8:27b	240	23.86	—	17.04	6.82
gemma4:31b	243	19.52	10.79	8.71	0.02
gemma4:26b-a4b	263	10.99	10.79	0.18	0.02
nemotron3:33b	287	10.98	10.71	—	0.27
qwen3.6:35b-a3b	532	7.16	6.71	0.34	0.12

Attribution uses each model's own GGUF dimensions. qwen3.8 reports no vision dims at all, so its bf16 tower work is reported as "other" rather than assigned to a tower whose size the model never states — an early draft guessed 1152/4304 by analogy with gemma4 and would have published that guess as fact.

qwen3.6 sending the least work here is consistent with it losing the most (−35.6%) when MMQ was disabled: MMQ does nearly all of its matmul work.

3. Both runtime routes are closed

Three cold repeats on gemma4:31b, one per distinct image, first request after start discarded:

arm	document	finetext	chart	vs baseline
baseline	155.1	154.9	153.3	—
ROCBLAS_USE_HIPBLASLT=1	153.7	153.7	152.4	−0.8%
GGML_CUDA_CUBLAS_COMPUTE_TYPE=f32	79.6	79.7	79.6	−48.6%

hipBLASLt is inert, and for a packaging reason. Upstream calls it better than Tensile on CDNA3. Here it does nothing — because neither the ROCm 7.2.4 image nor the ROCm 10 image ships a single gfx1151 hipBLASLt kernel file, despite the .so referencing TensileLibrary_*_gfx1151. Revisiting means shipping kernels, not flipping a flag.

Compute type is already the fast one by 2×. Every call runs compute_type f16_r. Forcing FP32 accumulate halves prefill, tightly (79.6 / 79.7 / 79.6) and in line with what RDNA3.5 WMMA predicts. There is no headroom to recover.

What this rules out

• Swapping in a better rocBLAS kernel via a runtime knob — both knobs inert or strictly worse.
• An upstream PR on the Q4_K default: return true threshold — the A/B measured that path at −0.1% on dense qwen3.8:27b.

Still open

1. Does the 2× from FP16 accumulate cost accuracy? FP32 accumulate is the more accurate option, so the win is only free if scores match. Suite running now; result will be posted here.
2. Could these skip rocBLAS entirely? The bailout triggers on dst->type != F32. An F16-output path through mmf would avoid it for the tower shapes — an upstream design question, and the one place real headroom might be.

Contents

file	what
docs/maxusai/rocblas-on-gfx1151.md	the writeup
vision-suite/gemm_shapes_gfx1151.tsv	1565 captured GEMM calls, with provenance header
vision-suite/summarize_gemm_shapes.py	generates the tables (ADR 0012 rule 8)
vision-suite/shape_inventory.sh	the capture script

The −48.6% doubles as a regression sentinel: if prefill on this host ever halves, check compute_type before anything else.

🤖 Generated with Claude Code

[glennneuber]

Correction — finding #1 retracted

Pushed in 1a106a47. The PR description above still has the original wording; this comment is the correct version.

What I claimed: every captured GEMM takes the dst->type != F32 bailout in ggml_cuda_mul_mat, inferred from --d_type f16_r appearing on every call. From that I concluded GGML_CUDA_FORCE_CUBLAS could not move any of them, and that this explained dense gemma4:31b measuring +2.1%.

Why it's wrong: d_type is set by prefer_f32_output, not by dst->type:

if (compute_type == GGML_TYPE_F16) {
    prefer_f32_output = cc == GGML_CUDA_CC_VOLTA || GGML_CUDA_CC_IS_RDNA4(cc) || GGML_CUDA_CC_IS_CDNA(cc);
}

gfx1151 is RDNA3_5 — RDNA3, not RDNA4 — so prefer_f32_output is false and F16 output is what this architecture gets regardless of dst->type. ggml_cuda_mul_mat_cublas has two callers (the bailout at the first branch, the fallthrough at the last line) and the ROCBLAS_LAYER log records neither. The tower GEMMs do take the bailout, since they run in an F16 graph; the LM ones are undetermined.

That distinction matters rather than being pedantic: if the LM GEMMs arrive by fallthrough, then should_use_mmq did decline them and this morning's A/B had a lever on them after all.

How it gets settled: diff the rocBLAS inventory between the MMQ build and the FORCE_CUBLAS build. Present in both → bailout. Only under FORCE_CUBLAS → MMQ was handling it. Queued behind the accuracy suite currently on the GPU.

What replaces it — and it's a better finding

The arch list is the real result. RDNA4 and CDNA get F16 inputs with F32 accumulate. rocBLAS ships HPA kernels for gfx1151 (TensileLibrary_Type_*_HPA_*_gfx1151). gfx1151 is excluded and accumulates in F16 — for BF16 too, via !GGML_CUDA_CC_IS_RDNA3(cc).

This is not the same knob as GGML_CUDA_CUBLAS_COMPUTE_TYPE=f32 measured in the PR body, which converts the inputs to F32 as well and halves throughput. F16-in/F32-accumulate usually costs little on hardware with HPA support and buys accuracy. Whether gfx1151's exclusion is deliberate tuning or an untested default is now the most promising open question here — and unlike the two runtime knobs, it needs a build.

Everything else in the PR body stands: the hipBLASLt packaging finding, the −48.6% compute-type measurement, the shape inventory, and the regression sentinel.

[glennneuber]

Results in — three questions closed, one bigger one opened

Pushed in 2295baaf.

1. Which caller — settled, and the retraction partly un-retracts

Diffing the rocBLAS inventory between the MMQ and FORCE_CUBLAS builds:

model	in BOTH builds	added by FORCE_CUBLAS
gemma4:31b	19.52 TFLOP	44.91 TFLOP
qwen3.8:27b	23.86 TFLOP	65.61 TFLOP
nemotron3:33b	10.98 TFLOP	11.65 TFLOP over 724 distinct shapes

Zero MMQ-build shapes are missing under FORCE_CUBLAS — the flag only ever adds work to rocBLAS. So the conclusion I retracted above (that FORCE_CUBLAS cannot move these GEMMs) is correct after all. The reasoning that produced it was still invalid; this is measured.

2. FP32 accumulate costs 2× and buys nothing

Full suite under GGML_CUDA_CUBLAS_COMPUTE_TYPE=f32, 0 failures, 5/5 models. Scores unchanged everywhere except one fine-text cell on gemma4:31b (9px, 4/4 → 3/4) which moved the wrong way. So −48.6% for no measurable accuracy. The f16 default is correct on both axes, not just speed.

3. prefer_f32_output for RDNA3.5: the exclusion is correct tuning

model	gate off	gate on	Δ
gemma4:31b	154.2	107.7	−30.2%
qwen3.8:27b	243.4	162.3	−33.3%

Gate-off reproduced the production baseline (154.2 vs 154.9), so the build is production plus the gate and the numbers hold. Cheaper than full FP32 but nowhere near free — and since more precision bought no accuracy in (2), there's no reason to pay for it. Scores under the gate are running.

The largest block of work isn't in this document at all

There is no 9900×9900 shape in any of the 11,812 captured GEMM calls. Tower attention never touches rocBLAS — it runs on flash-attention kernels. For gemma4:31b:

block	TFLOP / image request
tower attention (QK^T + AV)	12.2
tower projections	10.8
LM GEMMs	8.7

The biggest single block of vision prefill is one that none of the findings in this PR apply to. Three facts make it the lead:

• flash attention is enabled in production and in the bench harness — those kernels really are running this
• rocwmma has zero hits in the shipped libggml-hip.so; GGML_HIP_ROCWMMA_FATTN is off (upstream default)
• rocWMMA headers and cmake config are in the production build base, so the alternative is buildable

Queued: the rocWMMA build, and a diagnostic run with FA off — which forces tower attention through rocBLAS as explicit mul_mat and is the only way to get 9900×9900 into the inventory. I expect that one to be slower (materialising the matrix is what FA exists to avoid, ~339 GB of traffic over 27 blocks), but the ratio calibrates how much upside rocWMMA can possibly have before spending a build on it.

[glennneuber]

Brought up to date with main (b11081) in 93e2f2d. It had conflicted since the v0.34.4 fold.

• SPEC. Main kept H23 and H24 free for this PR and added H25 after H22, so the conflict was placement only. H23 and H24 now sit between H22 and H25, and the conformance table gains their rows. H23's example now says production ran q8_0 in 2026-09, since production has run f16 since ADR 0043.
• Compat 805 no longer applied at b11081. fattn.cu changed above flash_attn_mask_to_sparse_indices, and merging as it stood would have failed every build at the patch step. It is regenerated against the pinned source, and its added lines are unchanged. At b11081 the AMD WMMA branch it explains still excludes head dims 40 and 72 and caps at 256.
  • CI's patch check (cmake -S llama/server …) applies every compat patch: 805, then 903 and 908 after it.
  • The patched fattn.cu builds as a HIP object for gfx1151, in Dockerfile.rocm's toolchain stage (ROCm 7.2.4). The object carries the GGML_CUDA_FATTN_FORCE_TILE gate.
• kv-cache-type-per-model.md: the q8_0 → f16 cost note now says only the flash-attention cost remains in production.
• Unchanged: the findings are dated measurements on 0.34.2-dynres-f67b1aef (b10969). This PR does not re-measure them at b11081.

amd-server/rocm-gfx1151