fix(cuda): enable hybrid Q2_K/Q4_K GGUF load on decode-perf-tuning

[rrstesiak]

Loading the DeepSeek V4 Flash hybrid Q2_K/Q4_K GGUF on GB10 hung during model load: when the engine reached the final Q4_K layers, ds4-server hung on the CUDA call and the client received a generic CUDA error. This patch makes the hybrid GGUF load end-to-end and decode at ~20 t/s on this branch on fresh ctx (at the GB10 LPDDR5X bandwidth roofline).

Root cause:
routed_moe_launch hard-rejected the q4k MoE path unless
n_tokens == 1u && n_expert == 6u (single-token decode, 6 active
experts). The hybrid GGUF triggers q4k_path during load with
shapes outside that envelope; the rejected launch left the CUDA
stream in a bad state and the load hung.

Changes:

• Loosen the q4k_path gate; allow multi-token and wider-expert shapes through. The conditional body is deliberately empty as a marker: if a future perf regression appears for wider shapes, reinstating return 0 forces the slower IQ2_XXS-only fallback.

• Add a graph-capture-safe handshake (cudaEventRecord + cudaStreamWaitEvent) between ds4_mmq_stream_for_call() and ds4_current_stream() before the decode-vec MoE SwiGLU. cudaStreamWaitEvent records as a DAG edge inside capture, preserving cross-stream parallelism.

• Add a belt-and-braces cudaDeviceSynchronize() after moe_pair in the prefill path. Without it, very-large-prompt loads intermittently hit CUDA memory errors. Root cause is not yet pinned down; the heavier sync is intentional pending a proper repro.

Impact:

• Hybrid Q2_K/Q4_K GGUF now loads end-to-end on GB10/Spark.
• Decode runs at ~20 t/s short-ctx with CUDA graph capture enabled (near the documented ~20.5 t/s LPDDR5X roofline).
• antirez/ds4 upstream supports the hybrid quant correctness- wise but runs significantly slower on Spark; that delta is not investigated in this patch.

Thanks to @Entrpi for the decode-perf-tuning branch — this fix is a small refinement on top of that work. Built and tested on a single DGX Spark; not in a Military Industrial Complex. 284B parameters cranking through the LPDDR5X.

Tested on GB10 / Blackwell (SM121) only; CUDA 13.0.88, decode-perf-tuning branch at 5625a99. Ada / Hopper / Grace-Hopper / GB200 untested.

Test evidence

Per CONTRIBUTING.md: correctness exercised via the in-branch test runner (ds4-eval, the renamed ds4_test on decode-perf-tuning); speed via ds4-bench with the canonical speed-bench/promessi_sposi.txt prompt at the documented sweep parameters (--ctx-start 2048 --ctx-max 65536 --step-incr 2048 --gen-tokens 128). Pure Q2 control run included as CSV below; hybrid run failed at the 2048 frontier with q8 fp16 cache budget exhausted (pre-existing, see Known limitation).

Test	Model	Result
ds4-eval --questions 5	hybrid, this branch	3/4 PASSED (case 4 = pre-existing intermittent)
ds4-eval --questions 5	hybrid, parent 5625a99	❌ case 1 prefill fails (cannot load hybrid at all)
ds4-eval --questions 4	pure Q2, parent 5625a99	3/4 PASSED (same case 4 intermittent)
ds4-bench canonical	pure Q2, this branch	✅ clean through 30720, gen 17.4 → 6.4 t/s, crash at 32768 frontier with illegal memory access
ds4-bench canonical	hybrid, this branch	❌ prefill fails at 2048 frontier (q8 fp16 cache budget exhausted)
smoke test (~20 token prompt, 64 gen)	hybrid, this branch	✅ ~19–20 t/s, no CUDA errors

Pure Q2 canonical bench CSV (this branch)

ctx_tokens,prefill_tokens,prefill_tps,gen_tokens,gen_tps,gen_tps_ss,first_token_sec,kvcache_bytes
2048,2048,71.23,128,17.44,17.46,0.0666,52184476
4096,2048,379.84,128,15.55,15.55,0.0647,80373148
6144,2048,377.50,128,14.04,14.04,0.0716,108561820
8192,2048,379.90,128,12.78,12.78,0.0787,136750492
10240,2048,374.22,128,11.72,11.72,0.0857,164939164
12288,2048,369.63,128,10.82,10.82,0.0927,193127836
14336,2048,366.11,128,10.06,10.06,0.0999,221316508
16384,2048,363.26,128,9.38,9.38,0.1068,249505180
18432,2048,360.11,128,8.80,8.80,0.1140,277693852
20480,2048,356.25,128,8.28,8.28,0.1211,305882524
22528,2048,354.29,128,7.81,7.81,0.1284,334071196
24576,2048,350.12,128,7.40,7.40,0.1357,362259868
26624,2048,348.43,128,7.02,7.02,0.1425,390448540
28672,2048,344.17,128,6.68,6.68,0.1499,418637212
30720,2048,338.07,128,6.39,6.39,0.1570,446825884

Known limitation

Large-context prefill (≥2048 tokens via ds4-bench) with the hybrid GGUF remains a branch-level limitation. Not a regression introduced by this patch — the parent commit (5625a99) cannot prefill the hybrid model at even 200 tokens. This patch extends the working range from zero to ~2000+ tokens, sufficient for interactive chat use of the hybrid quant. Large-prefill support is out of scope here.