fix: LFM2.5 paged turbo4 on gfx803 (warp-width, sub-128 head_dim, parallel partials)

[kmbandy]

Summary

Makes LFM2.5-8B-A1B paged turbo4 work correctly on gfx803 (RX 480), and fixes two latent bugs in the mt_paged_attention turbo path that affected it. Validated end-to-end on the real service config (512K ctx, --parallel 4, turbo4 paged KV + --kv-tiered 25,25,50 + semantic index) — coherent, correct generation at ~65 tok/s. Resolves MAD-298.

Three independent root causes (one commit each):

1. 62f4f99ae — pin turbo4 cooperative warp-collectives to WARP_SIZE (gfx803 corruption)

The turbo4 cooperative kernels (mt_scatter_kv_turbo4_0_kernel, decode_coop_stage_turbo4, the tile decode, and the set-rows turbo4 quant) are written for 32-lane warps (WARP_SIZE=32, DECODE_NUM_WARPS=NUM_THREADS/32, lane=tid%32) but issued bare 3-arg __shfl_*_sync, which the HIP shim maps to width=warpSize. On gfx803 (GCN) warpSize=64, so two logical 32-lane warps share one 64-wide hardware wave:

• nibble-pack __shfl_sync(.., my_nibble, lane^1) read the partner from the wrong half of the wave for threads 32–63 / 96–127 → ~25% of packed K/V nibbles corrupted.
• dequant norm broadcast __shfl_sync(.., norm_f, 0) pulled lane-0's norm across the whole 64-wide wave → logical warp 1 used logical warp 0's per-block norm.

Fix: pass WARP_SIZE explicitly (4-arg form) on every warp-collective in these kernels. No-op on CUDA (warpSize already 32); restores correctness on gfx803.

2. fb556b194 — support sub-128 head_dim for the paged turbo cache (LFM2.5 head_dim 64)

Paged turbo quantizes 128-element blocks (QK_TURBO=128); LFM2.5's head_dim 64 is smaller than one block and the dispatch rejected it. Zero-pad q/k/v to 128 in build_attn for the paged turbo path (pad before the F16 cast since ggml_pad needs F32; slice the output back to head_dim) and size the paged cache to the padded head_dim. Mirrors the non-paged turbo path. Padded Q dims are zero (contribute nothing to QK), padded V output dims are sliced off.

3. 93a1e6a11 — size paged-decode partials by total_q_tokens, not avg_q_len (arch-independent)

launch_paged_attn_decode sizes the partials scratch and strides every head/seq/chunk partial offset by max_q_len, which the dispatch set to avg_q_len = total_q_tokens / num_seqs. With --parallel N and one active request (1 active seq + N−1 idle slots) this floors to 0 (1/4 = 0):

• partials_n = n_heads * num_seqs * num_chunks * 0 * (HS+2) = 0 → empty buffer
• partial_chunk_base = (...) * max_q_len(0) * (HS+2) = 0 for all heads/seqs/chunks → every partial write collapses onto offset 0 of a zero-sized buffer → OOB + total collision → token salad.

Only surfaces at --parallel > 1 (the decode path). Manifested as token-salad on LFM2.5 paged turbo4 at the service's --parallel 4 (clean at --parallel 1); q8_0 only escaped by luck of allocation layout (the OOB is UB). Fix: max_q_len = total_q_tokens (a safe upper bound on any single seq's q_len; gate-capped ≤ 8). The decode gate was already fixed to key off total_q_tokens for the same flooring reason — this carries that correction through to the buffer sizing it forgot.

How it was found

The isolated decode harness (llama-gpu/tests/test_turbo4_decode_full.cu, -DHS_TEST=128) passed at HS=128 while production was broken — its test norms varied only slowly across tokens (0.4 + 0.0007*tok), masking bug #1's norm contamination. Reproduced only with strongly-varying norms (rel-L2 0.72 → 0.0003 with the fix). Bisection: q8_0/f16 paged coherent but turbo4 garbage (→ turbo-specific); --parallel 1 coherent but --parallel 4 garbage (→ bug #3); MAD_PAGEDATTN_PROBE=verbose showed avg_q_len=0 at parallel 4. The harness norm fill is now a regression guard.

Testing

• gfx803 isolated harness: scatter + decode (q_len 1/2/4, GQA dims, multi-chunk) all PASS at HS=128/256 after fixes.
• Live service llama-server-lfm25-8b (RX 480, full tiered config): coherent + correct, e.g. "List three primary colors:" → "red, yellow, blue…", ~65 tok/s.
• Both builds compile: CUDA build-army (sm_61) and ROCm build-rocm-gfx803.
• Mneme embedder (qwen3-06b) restarted on the new binary — valid 2048-dim embeddings.

Notes for review / merge

• Commits 1 and 3 are no-ops on CUDA; commit 3 is arch-independent and likely also bears on MAD-288 (omnicoder paged-turbo4 CUDA corruption under --parallel>1) — see that ticket's comment.
• Targets fork master (not upstream).