turbo3 KV multi-slot TG regression on NVIDIA Volta sm_72 — prompt-length dependent

[minhtcai]

Summary

turbo3 KV cache shows a clean +12% TG improvement over q8_0 in single-stream tests on Jetson Xavier AGX (sm_72, CUDA 11.4), but regresses -12% to -22% under --parallel >= 4 with longer prompts. The regression scales linearly with prompt length. q8_0 shows no equivalent regression.

Hardware / build context

• Jetson Xavier AGX (sm_72 Volta, 30 GB unified memory)
• CUDA 11.4.315, gcc 9.4
• Branch: feature/triattention-scoring at f9a308d
• Build: cmake -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=72 -DGGML_NATIVE=OFF -DCMAKE_BUILD_TYPE=Release — clean, all turbo fattn-vec instances compiled

Reproduction

Two binaries built from the same source. Same model (Qwen3.5-0.8B-Q4_K_M) for the small-model isolation, Qwen3.6-35B-A3B-UD-Q4_K_XL for the production-size validation.

Single-stream (llama-bench -p 1024 -n 128 -r 3)

KV (K/V)	TG t/s
q8_0 / q8_0	44.33
turbo3 / turbo3	49.66 (+12%)

turbo3 clearly wins — Sparse V skip path is active and helping.

Multi-slot (llama-server --parallel 8, single client, same prompt)

KV (K/V)	TG t/s
q8_0 / q8_0	44.62
turbo3 / turbo3	39.18 (-12%)

Same hardware, same model, identical request — only --parallel 8 and KV-type differ.

Production-scale (Qwen3.6-35B-A3B Q4_K_XL, --parallel 8, --mlock, c=524288)

Prompt size	KV	TG t/s
2 tokens	turbo3	15.52
200 tokens	turbo3	15.12
562 tokens	turbo3	14.15
1.7K tokens	turbo3	12.01
(any)	turbo2 [the legacy 8-bit alias upstream]	~15

Linear degradation with prompt length under multi-slot. Single-stream the same model holds 15+ across prompt sizes.

Failed hypothesis (sharing so it isn't repeated)

I hypothesized the Sparse V skip in fattn-vec.cuh lines 349/395 was the cause — if (dominated) { continue; } triggers warp divergence, and the code already disables this on HIP/RDNA3 (commit ae9814cd5, "branch overhead exceeds dequant savings"). I patched the guard to also disable it on __CUDA_ARCH__ < 800:

#if !defined(GGML_USE_HIP) && (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 800)

Result was the opposite of what I expected. Disabling Sparse V on Volta caused single-stream turbo3 TG to drop from 49.66 → 35.06 (-29%, ±8.81), while q8_0 held at 44.11. So Sparse V is what gives turbo3 its single-stream win on Volta — turning it off forces full V dequant on every position. Reverted.

This means the multi-slot regression is not caused by Sparse V divergence cost. The fact that single-stream turbo3 runs the Sparse V path more aggressively (more skips at long context) yet still beats q8_0 single-stream suggests the divergence cost itself is fine on sm_72.

Remaining hypotheses (not tested — would need profiler access)

1. kv_unified=false per-slot KV layout interaction with turbo3's structured rotation per row. Multi-slot KV cells are in disjoint sub-buffers, possibly hurting cache locality more for turbo3's WHT-rotated reads than q8_0's simpler row layout.
2. K-side attention sharpening (α=1.0357 per the log) applied per-slot — wondering if it's redundant work in the multi-slot path.
3. WHT rotation matrix in turbo-wht.cu — maybe per-slot init not amortized, or shared-memory contention scales worse than q8_0 dequant.

I don't have nsight-compute on a Jetson + CUDA 11.4 toolchain, so I can't profile to pin which kernel is the actual hot path.

Asks

1. Could you confirm whether you've tested --parallel >= 4 with turbo3 on any hardware? README benchmarks (TURBOQUANT.md) appear to be single-stream RX 7900 XTX.
2. If you have profiling data on Ampere with multi-slot, does turbo3 show similar prompt-length-dependent TG degradation, or is this Volta-specific?
3. Suggested investigation path? Happy to run any nsys/ncu traces if you have a profiling-capable build to send.

Single-stream turbo3 works great on Volta — the underlying kernels are fine. The issue is specifically multi-slot serving with growing KV cache, which is the production deployment target. Reproduces 100% on this hardware.

Why this matters for Jetson deployment

Production prod uses --parallel 8 with c=524288 to serve multiple users. Switching q8_0 -> turbo3 would free ~6 GB of unified memory (5.12× compression vs f16), but the TG regression on long conversations cancels the upside. If we can fix the multi-slot path, this becomes a clean win for edge deployment on 30 GB unified-memory Jetson hardware.

[minhtcai]

Profiler-Driven Diagnosis Update

Got nsys (Nsight Systems 2023.2.4) running on Xavier AGX after the install pull from nsight-systems-2023.2.4 package on JetPack r35.4. Key findings:

Per-call kernel cost (multi-slot 4-way concurrent, 1k prompt + 64 gen, 0.8B turbo3 Q4_K_M)

Kernel	Single-stream avg	Multi-slot avg	Ratio
flash_attn_ext_vec<256, 1, turbo3, turbo3, false>	86,388 ns	2,212,299 ns	25.6×
flash_attn_tile<256, 256, 32, 1>	3,449,951 ns	3,193,779 ns	0.93×
dequantize_block_q4_K<half>	230,767 ns	231,381 ns	1.00×
k_set_rows_turbo3<long, 128>	17,979 ns	37,183 ns	2.07×
k_turbo_wht_f32<0, 128>	19,552 ns	52,911 ns	2.71×
mul_mat_vec_q<14, 4> (multi only)	n/a	691,845 ns	n/a

The attention vec kernel is 25× slower per call under multi-slot. Per-token-position normalized: pre-fix ratio 5.1× (multi 295,028 ns/token-pos vs single 57,817 ns/token-pos).

Root cause (partial)

fattn-common.cuh:1426 heuristic gating KV_max scan:

if (mask && ne11_logical % FATTN_KQ_STRIDE == 0 && (Q->ne[1] >= 1024 || Q->ne[3] > 1)) {

During multi-slot decode with cont-batching:

• Q->ne[1] = number of active slots (1-8), not 1024
• Q->ne[3] = 1 (slots aren't a separate batch dim)

→ KV_max scan skipped → all attention blocks scan full unified KV (ne11) instead of just their slot's range.

Fix attempt

Patch:

if (mask && ne11_logical % FATTN_KQ_STRIDE == 0 && (Q->ne[1] > 1 || Q->ne[3] > 1)) {

Result

Metric	Pre-fix	Post-fix	Δ
Single-stream turbo3	49.66 t/s	49.86 t/s	+0.4%
4-way multi-slot turbo3	14.7 t/s/slot	16.27 t/s/slot	+10.7%
flash_attn_ext_vec avg	2,242,798 ns	2,212,299 ns	-1.4%
flash_attn_mask_to_KV_max calls	0 (gated off)	384+54 = 438	now firing

The fix works as predicted (KV_max kernel now firing in multi-slot decode), but the kernel cost only dropped 1.4% per call. Wall-time +10.7% comes mostly from removed redundancy elsewhere, not the K-scan reduction.

Why the fix is undersized

Hypothesis: the dominant bottleneck is template specialization, not K-scan length:

• flash_attn_ext_vec only has ncols=1 and ncols=2 template instances (see fattn-vec.cuh:592).
• During multi-slot decode with 4 active slots, dispatcher uses ncols=1 and launches gridDim.x=4 blocks (one per slot).
• Each block independently iterates over K positions; blocks must all finish before kernel completes.
• With ncols=4 or ncols=8 instance (like the mma kernel has), 4-8 slots could share K-loads and amortize per-position work.

For Volta sm_72 (no MMA), the mma path isn't available, so the vec kernel's missing ncols>2 instances are an unfilled gap.

Production status

Still on TheTom turbo2 8-bit at 15.1 t/s baseline. Patched binary preserved at /mnt/sdcard/triatt-test/build/bin/llama-server for further experiments. Post-fix turbo3 at +10.7% over pre-fix turbo3 multi-slot but still ~2× behind TheTom 8-bit on the 35B-A3B prod model — the real comparison would need a big-model bench, but small-model trends suggest turbo3 won't beat TheTom 8-bit even with the fix.

Asks

1. Is there a planned ncols=4/ncols=8 instance for fattn-vec to support multi-slot decode without falling back to mma?
2. Is the KV_max gate intentionally conservative for low ne01 to avoid scan overhead, and if so, what's the empirical break-even?
3. Worth a PR for the gate broadening even though gain is modest?

[minhtcai]

Update: ncols=4/8 dispatcher experiment — negative result

Following up on the kernel-design hypothesis: tried adding ncols=4 and ncols=8 instances to the fattn-vec dispatcher to batch multi-slot decode in a single launch. Result: slight regression on Volta sm_72, not a win.

Patch attempted

fattn-vec.cuh::ggml_cuda_flash_attn_ext_vec_case dispatcher changed from binary Q->ne[1]==1 ? ncols=1 : ncols=2 to ladder:

if (q_rows == 1)      → ncols=1
if (q_rows <= 2)      → ncols=2
if (q_rows <= 4)      → ncols=4
else                  → ncols=8

Used std::integral_constant + generic lambda to instantiate without combinatorial repetition. Built clean (-Wpedantic quiet aside from pre-existing warnings).

Bench

Qwen3.5-0.8B Q4_K_M, turbo3 KV, 4 concurrent reqs of 1k prompt + 64 gen, --parallel 8:

Setup	Single-stream	4-way TG/slot	4-way aggregate
pre-any-fix	49.66	14.70	58.80
KV_max gate fix only	49.86	16.27	65.10
+ ncols=4/8 dispatcher	42.43	15.73	62.93

ncols=4/8 underperformed ncols=1 across the board — even single-stream regressed (42 vs 49). Likely cause: register pressure. ncols=4 means VKQ[4][...], KQ_max[4], KQ_sum[4], Q_reg[4][...] — 4× the per-thread state. On sm_72 (65k registers/SM, __launch_bounds__(N, 1)), the kernel was already at min occupancy, and the extra registers spilled to local memory.

Reverted

Patch rolled back. Production keeps the KV_max gate fix from before, which gave a real +10.7% on multi-slot.

Conclusion

The fattn-vec ncols=2 template was the upper limit before register spill on Volta. To actually batch slots into one kernel launch on sm_72 you'd probably need a separate kernel that's more parallel-friendly (e.g., the fattn-mma kernel's approach with explicit register tiling), not just a wider ncols on the existing vec kernel. That's beyond what we can do without a Volta tensor-core path.

Status

• Issue Request: publish TriAttention calibration script (TRIA v2/v3 stats.bin producer) #10 (calibration script) ✅ working tool published, see comment above
• Issue turbo3 KV multi-slot TG regression on NVIDIA Volta sm_72 — prompt-length dependent #11 (multi-slot regression) — KV_max gate fix gives +10.7%; ncols dispatch route exhausted; deeper fix would need new kernel design

If you have plans for a Volta-friendly multi-slot kernel let me know and I'll test on Xavier. Otherwise multi-slot turbo3 on Xavier is bounded at ~16 t/s/slot for 4-way concurrent on the small model.

[minhtcai]

Verification on production model — KV_max gate fix is a no-op on hybrid SSM+attention

Followed up by re-verifying the KV_max gate fix on the actual production model (Qwen3.6-35B-A3B qwen35moe, NOT the 0.8B used in earlier benches). Surprised by the result.

Setup

• Hardware: Jetson Xavier AGX, sm_72 Volta, 30 GB unified, JetPack r35.4 / CUDA 11.4
• Model: Qwen3.6-35B-A3B-UD-Q4_K_XL.gguf (35B-A3B MoE, Q4_K_M)
• Server: --parallel 8 -c 32768 --cache-type-k turbo2 --cache-type-v turbo2 -fa on
• Patched binary: libggml-cuda.so.0 rebuilt with KV_max gate patch (Q->ne[1] > 1 || Q->ne[3] > 1), fattn-vec.cuh reverted to upstream ncols=1/2 dispatcher
• Prompt: 24-tok chat, 256-tok decode

Result

Setup	Single-stream	4-way TG/slot	4-way aggregate
Production (TurboQuant turbo2, no patch)	15.74 t/s	7.20 t/s	28.79 t/s
Patched (KV_max gate fix only)	15.51 t/s	7.08 t/s	28.31 t/s
Delta	-1.5%	-1.7%	-1.7%

Within noise — actually slightly negative. The +10.7% multi-slot win seen on Qwen3.5-0.8B (pure-attention, dense) does NOT transfer to this model.

Why it doesn't transfer

Qwen3.6-35B-A3B is a hybrid model:

• qwen35moe.full_attention_interval = 4 → only every 4th layer is full attention (10 of 40)
• The other 30 layers are Mamba SSM (recurrent state buffer = 502 MiB vs KV buffer = 136 MiB)

The fattn-vec KV_max optimization only touches the attention path. Even a 25× per-call speedup on attention layers buys little end-to-end when:

1. SSM state propagation dominates wall-time (75% of layers)
2. KV is tiny (-c 32768 / 8 slots = 4096 tok/slot × 10 attention layers × 512 dim KV at turbo2 → ~17 MiB total active scan range, well below the FATTN_KQ_STRIDE-aligned scan even without KV_max)

Profiler showed GPU at 97% utilization in both builds — the attention kernel just isn't the bottleneck on this workload.

Where the patch does help

The earlier 0.8B bench was on a dense pure-attention model with full attention on all layers. There the KV_max gate fix did land +10.7% multi-slot. So the patch is correct and useful for dense attention Volta deployments — just not for this specific Qwen3.5/3.6 MoE+SSM hybrid which is what production runs.

Conclusion for prod

Production stays on the unpatched TurboQuant build. Patched binary preserved at /mnt/sdcard/triatt-test/build/bin/llama-server for future testing on dense models. Closing this verification thread; happy to leave the issue open if upstream wants the patch for non-hybrid Volta deployments.

cc: anyone running TurboQuant on dense attention models on Volta — the KV_max gate fix at fattn-common.cuh:1426 (Q->ne[1] >= 1024 || Q->ne[3] > 1 → Q->ne[1] > 1 || Q->ne[3] > 1) gives +10% multi-slot decode for free.