Change sets 4-5: chunked SSD Mamba-2 prefill + stream-k tuning for CDNA2 (+23.5% pp16384)

[davetha]

Two prefill optimisations for gfx90a, measured on Nemotron-3-Super-120B-A12B (i1-Q4_K_M, 80 GiB) across 2× MI210.

Result

llama-bench, -b 4096 -ub 2048 -fa 1 -sm layer -ctk q8_0 -ctv q8_0 -t 24 -r 2

build	pp4096	pp16384	vs base
upstream 67b9b0e	1366	1775	—
+ SSD on CDNA	1624.07 ± 0.81	2114.21 ± 0.73	+19.1%
+ stream-k off	1683.57 ± 0.40	2192.26 ± 0.84	+23.5%

Gap to vLLM+AITER (4,070 t/s at 16k) narrows from 2.29× to 1.86×. Not closed.

4. Chunked SSD Mamba-2 prefill (+19%)

Upstream PR #22675 replaced the sequential SSM scan with a State-Space-Duality formulation — per-chunk intra-chunk output and chunk-final state become batched FP16 GEMMs — then gated it to NVIDIA Turing+. The author states it "does not affect ... HIP": a scoping decision, not a technical limit.

On gfx90a that scan was the single largest prefill kernel: rocprofv3 put ssm_scan_f32_group at 22.8% of GPU time (1759 ms / 7703 ms), scalar FP32, zero matrix-core use. Admitting HIP to the guards and adding CDNA to the capability test drops it to 1.0% and moves the work onto the 181 TFLOPS v_mfma_f32_16x16x16f16 path.

It compiled for HIP with no changes to the GEMM calls — every cuBLAS symbol already had a hipBLAS alias in vendors/hip.h. Gated to CDNA, not blanket AMD; RDNA's WMMA path is unvalidated. USE_CUB left HIP-excluded (hipCUB collides with ggml's __trap macro).

5. Disable stream-k for K-quants (+3.7%)

Every CASE in mmq-config-cdna.cuh sets stream_k = true, inherited across all CDNA generations. PR #26199 found stream-k "helps a lot for Dense models and hurts MoE models" when retuning RDNA. This is MoE, and MMQ is 53.3% of prefill after change set 4.

Stated plainly: that evidence is RDNA3.5/RDNA4, not CDNA2. This was an experiment that paid off, not a transferred result — an earlier experiment on similar reasoning (extending CDNA3's rocBLAS carve-out to CDNA2) came back 6.5% slower and was discarded.

Verification

• test-backend-ops -o SSM_SCAN 7/7 per device, including the 256/512/300-token multi-chunk shapes that exercise the SSD path.
• test-backend-ops -o MUL_MAT_ID 790/790 per device.
• Generated tokens read at temperature 0 on a 240-token prompt. The SSD path chains batched GEMMs with beta=1 accumulation; a wrong stride doesn't crash, it yields fluent wrong text — so throughput alone is not evidence.
• Patches verified with git apply --check against 67b9b0e, reproducing the modified-files/ copies byte-for-byte.

Pre-existing multi-GPU fault (documented, not introduced)

The second sequential llama-server request faults on GPU node-1 across two cards. Bisected:

build	GPUs	result
unpatched 67b9b0e	2	faults on request 2
+ SSD	2	faults
+ both	2	faults
+ SSD	1	3 requests clean

The unpatched baseline faults identically, so it is unrelated to these changes. llama-bench never surfaces it, pointing at state reuse between requests rather than a prefill kernel — throughput numbers unaffected. Not root-caused.

Note on base

Change sets 4-5 are against upstream ggml-org/llama.cpp 67b9b0e, a different base from change sets 1-3 (TurboQuant c26cbdf). They have not been tested on the TurboQuant fork and ssm-scan.cu changed substantially upstream in the interim, so patches/04-* will likely need rebasing to apply there. The two groups touch disjoint files.

Branch ssd-cdna2 on this repo carries the same work as a full llama.cpp tree with real history, if that is easier to build from.