fix(cuda): turbo KV fused-MMA — use_sparse off-by-one and missing SMEM swizzle

[wilky2005]

Two independent bugs made the fused-MMA TurboQuant KV path (-ctk/-ctv turbo4) unusable: it crashed on load, and after fixing the crash it produced garbage. Both are fixed here; the path is now correct and ~34% faster than q8_0.

Bug 1 — crash on load: use_sparse omitted (off-by-one)

ggml/src/ggml-cuda/fattn-mma-turbo.cuh: the single launch_fattn(...) call omitted the use_sparse argument, so warp_size_host (int 32) landed in the use_sparse bool slot and evaluated to true. With n_kv_max = 0 (op param 4) that trips GGML_ASSERT(n_kv_max > 0) in fattn-common.cuh during warmup decode.

Fix: pass /*use_sparse=*/false before warp_size_host, mirroring the f16 reference call (... stream_k, use_sparse, warp_size_host).

Bug 2 — garbage output: missing SMEM swizzle in the turbo tile loaders

ggml/src/ggml-cuda/fattn-mma-f16.cuh: flash_attn_ext_turbo4/3/2_load_tile wrote K/V tiles linearly (tile_KV[row*stride_tile + c]) while load_ldmatrix<..., swz=true> reads them through the XOR swizzle bytes_rc<stride_tile>(row, col). When swizzle is enabled (nbatch_2 >= 32 && nbatch_2 % 32 == 0, e.g. head-dim 256 -> nbatch_K2 = 128) the un-swizzled write is permuted on read, scrambling attention and producing garbage (the model misread "1724" as "1010").

Fix: add a bool swz template parameter and write through bytes_rc when swz is set, matching flash_attn_ext_f16_load_tile. Call sites pass swz_K / swz_V.

Result

On Ornith-1.5-35B-A3B (qwen35moe, head-dim 256, RTX 3080 Laptop 16 GB):

• correctness 9/9 (was 0/9) on arithmetic probes
• ~43.6 t/s vs 32.6 t/s for q8_0 (+34%), with ~0.9 GB less VRAM

Summary by CodeRabbit

• Bug Fixes
  • Improved CUDA flash attention compatibility for turbo-quantized key/value data.
  • Corrected memory layout handling to support swizzled data paths consistently.
  • Fixed turbo flash attention launch configuration for reliable execution across supported settings.

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: defaults

Review profile: CHILL

Plan: Advanced

Run ID: 0859a9f8-b9ab-45e1-bc40-5cf0e2bc601e

📥 Commits

Reviewing files that changed from the base of the PR and between 27c54b4 and 1fdcce6.

📒 Files selected for processing (2)

• ggml/src/ggml-cuda/fattn-mma-f16.cuh
• ggml/src/ggml-cuda/fattn-mma-turbo.cuh

Included review availability: Your plan provides up to 2 included reviews per hour; 1 remains after this review.

---

📝 Walkthrough

Walkthrough

Changes

CUDA flash-attention integration

Layer / File(s)	Summary
Turbo K/V tile swizzle wiring ggml/src/ggml-cuda/fattn-mma-f16.cuh	K and V turbo tile loader calls now receive swz_K and swz_V.
Turbo launcher argument update ggml/src/ggml-cuda/fattn-mma-turbo.cuh	The turbo launcher now passes use_sparse=false to launch_fattn.

Priority: ➖ Normal

Estimated code review effort: 2 (Simple) | ~10 minutes

Change: Bug fix

Suggested reviewers: johannesgaessler, thecodacus

Merge Risk: ⚪ Minimal · up to 1fdcc

No actionable correctness risk remains from this change; the CUDA turbo swizzle and launcher updates are ready to merge.

🚥 Pre-merge checks | ✅ 4 | ❌ 1

❌ Failed checks (1 warning)

Check name	Status	Explanation	Resolution
Description check	⚠️ Warning	The description clearly explains the two bugs, fixes, and reported results, but it omits the required ## Requirements section and its contributing-guidelines agreement and AI usage disclosure.	Add the required ## Requirements section. Confirm agreement with the contributing guidelines and include the required AI usage disclosure, including whether AI was used and how.

✅ Passed checks (4 passed)

Check name	Status	Explanation
Title check	✅ Passed	The title clearly identifies both primary fixes: the use_sparse argument issue and the missing shared-memory swizzle in the CUDA TurboQuant fused-MMA path.
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check. Docstring coverage is scoped to functions touched by this diff. Analyzed 0 functions across 0…
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.

• Fix all pre-merge checks with AI

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

• X
• Mastodon
• Reddit
• LinkedIn

_{Comment @coderabbitai help to get the list of available commands.}