fix(cuda): KVarN decode combine kernel crashes at ≥768K context (missing shared-mem opt-in)

[chimpera]

Problem

Any KVarN-KV preset (--cache-type-k/-v kvarn* + --flash-attn on) hard-crashes the server at contexts above ~768K (e.g. a 1M --ctx-size build dies at the prefill→decode boundary):

CUDA error: invalid argument
  in function ggml_cuda_fattn_kvarn_decode_launch
    at ggml/src/ggml-cuda/fattn-mma-kvarn-decode.cuh:784

Prefill completes, then the server aborts on the first decode step. Not OOM and not a quality issue — the answer never gets to return.

Root cause

The KVarN decode combine reduction kernel holds n_splits partials in dynamic shared memory, where n_splits = ceil(n_kv / SPLIT_TOKENS) and SPLIT_TOKENS == 64. That shared-mem request grows linearly with context:

n_kv	n_splits	dynamic shared mem
524288 (512K)	8192	32 KB (ok)
655360 (640K)	10240	40 KB (ok)
786432 (768K)	12288	48 KB = 49152 B ← default ceiling
899731 (900K)	14058	~55 KB (over)
1048576 (1M)	16384	64 KB (over)

CUDA's default per-block dynamic-shared-memory limit is 48 KB (49152 B); requesting more requires opting the kernel in via cudaFuncAttributeMaxDynamicSharedMemorySize. The sibling MMA kernels in fattn-mma-kvarn-case.cuh already do this (lines 424/538/715/812/817), but the combine kernel in fattn-mma-kvarn-decode.cuh was missed — so at n_kv ≥ 786432 the launch is rejected with cudaErrorInvalidConfiguration ("invalid argument"). Threshold matches exactly: 786432 = 12288 × 64 = 49152 / 4.

Fix

Add the same cudaFuncSetAttribute opt-in for the combine kernel before its launch, clamped to the device's cudaDevAttrMaxSharedMemoryPerBlockOptin and guarded by GGML_USE_MUSA, mirroring the existing pattern. One localized change in fattn-mma-kvarn-decode.cuh.

Verification

On an RTX 5090 (sm_120), a 900K-context needle-retrieval request that hard-crashed before the patch now completes: prefill ~449 tok/s, correct retrieval, finish=stop, server survives past the prefill→decode boundary. No regression at small context (32K–640K unchanged).

The patched file is identical between main and v0.4.3, so this applies cleanly to either; targeted at v0.4.3 since that's where it was developed/tested — feel free to retarget to main.

🤖 Generated with Claude Code