feat(v4/cuda): the CUDA tier wired into the DeepSeek V4 engine — generic kernels for any sm_80+, Windows runtime DLL loader, Linux CUDA=1

[dcutugno]

Stacked on #772 (@ZacharyZcR's kernels-as-a-tier, rescoped head), on #988 and on #1053 — until those merge, the diff against dev includes them. The three commits that are this PR: dc54ce3 (kernels/loader), c259de0 (engine wiring), cee7149 (docs).

This is the engine side #772 deliberately left out, plus the kernels the wiring needed. It follows the direction set in #772's review: our own kernels, small consumer cards first, every stage CPU-canonical with a per-stage CPU fallback. The vLLM/FlashInfer files are gone (the same removal as #772's rescope), DeepGEMM stays an opt-in build (DEEPGEMM_HOME; a small follow-up PR offers to vendor the sm120 headers, separable).

What it is

• Engine tier (COLI_V4_UNIT_GPU, #ifdef COLI_V4_GPU_TIER): all state stays in host memory; the tier accelerates stage by stage and falls back to the CPU reference for anything it declines. Dense/attention mirrors in VRAM (6.3 GB) with GPU dense matvecs in decode; a VRAM-sized fp4 expert mirror cache for decode (DSV4_CUDA_EXPERT_MIRRORS is an upper bound, growth stops when DSV4_CUDA_VRAM_RESERVE_MB is left; in-place refill); fused routed+shared MoE, sparse attention on the device KV, GPU indexer.
• Prefill (COLI_CUDA_ATTN_BATCH=1): the whole-chunk attention block on the GPU — compressor/indexer projections, sparse attention on a persistent device KV ring (indexed kernel past index_topk), wo, mHC — plus batched indexer scoring (bitwise the CPU loop: __fadd_rn/__fmul_rn, MinGW GCC does not contract FMA) with a bitwise fp8 reference-matmul replica for the query projection.
• Prefill MoE (COLI_CUDA_MOE_BATCH=1): route on device, refill a transient per-layer VRAM expert bank route-aware in pin-slot-sized groups (V4_MOE_REFILL_GROUP), lookups of group g+1 overlapped with the uploads of group g, page-locked host slabs; expert-grouped GEMM rows; a lookup that loses the pin-slot race is retried sequentially instead of failing the layer to the CPU union.
• Windows: MinGW engine + MSVC/nvcc DLLs, two flavours picked at start-up by dsv4_cuda_backend_arch_ok: coli_cuda_dsv4.dll (generic kernels, CUDA_ARCH=portable = sm_80/86/89/90/120 + PTX, any RTX 30/40/50, A/H) and coli_cuda_dsv4_dg.dll (DeepGEMM sm_120a, opt-in). backend_loader_dsv4.c resolves every mandatory symbol up front; COLI_DSV4_DLL forces one; banner [DSV4 CUDA] backend=... (generic|deepgemm-sm120).
• Linux: make -f Makefile.deepseek-v4 deepseek-v4 CUDA=1 [CUDA_ARCH=native|portable|sm_XX] [DEEPGEMM=1 DEEPGEMM_HOME=...] [NO_TC=1] links the tier directly (libcuda from the toolkit stubs + rpath); NO_TC=1 compiles out the opt-in cuBLASLt MXFP8 path for toolkits older than 12.8 (Jetson). Verified on Ubuntu 22.04/WSL2, RTX 5080, CUDA 13.3.1 + 12.6.85 — CPU-only, generic, DeepGEMM builds; the 12.6 compile fails on exactly the four cuBLASLt 12.8 symbols without the gate.
• CLI: --gpu/--vram for the V4 engine, coli doctor sees the DLL, V4_MTP_GPU opt-in; kernel tests c/tests/test_dsv4_*_cuda.c (batched sparse attention: max diff 1e-3 = one bf16 ulp vs the CPU reference; decode kernels through the loader ABI).

Numbers (RTX 5080 16 GB, 2× NVMe mirror, 32 GB RAM)

	before	after
3324-token prefill	271 s	90 s (generic kernels: 225 s, text identical)
8.3k-token opencode first turn	~10 min	~4 min once, then 6–9 s per session/turn (with #1053's checkpoints)
decode at 3k context	0.6 tok/s	1.6–1.76 tok/s

Per-card settings for 6/8/10–12/16/24/32 GB are in docs/deepseek-v4.md (VRAM budget: 6.3 GB dense mirrors, 2.2 GB prefill bank, ~8 MB per decode mirror; free VRAM sizes the cache at run time). Every GPU stage was accepted only when its greedy text matched the CPU reference on the same prompt; what "identical" can and cannot mean across kernels/cache states is written up in the docs' Validation section.

Overlap

#1031 (@rafpigna) touches the same area with a different design (single DLL, per-call safetensors reads for dense-on-GPU, .coli_kv single-session persistence). Happy to reconcile with whichever lands first; the numbers above are the comparison point.

🤖 Generated with Claude Code

[JustVugg]

Two of the things your stack was carrying have landed in dev:

• DeepSeek V4: the CUDA kernels, as a tier rather than an engine (rebase of #772) #1054 — @ZacharyZcR's kernels-as-a-tier (DeepSeek V4: the CUDA kernels, as a tier rather than an engine #772), rebased. The only conflict was NVCCFLAGS in c/Makefile: his $(NVCC_STD) parameterisation against dev's -ftz=false pin from CUDA: warp-per-row fmt=8 (fp8-e4m3) kernels, reference-mirroring accumulation #1037 and the glibc C23 compat block. All three are kept.
• feat(v4): dual-SSD mirror (COLI_MODEL_MIRROR) for the DeepSeek V4 expert store #988 — your dual-SSD mirror, green and clean.

So #1055, #1053 and #1056 no longer need to carry either of them. A rebase on current dev should collapse #1055 to roughly its own three commits (dc54ce3, c259de0, cee7149), which is the point at which it becomes reviewable as the engine wiring rather than as a 10k-line diff.

One thing worth flagging before you rebase: c/Makefile moved twice today. Beyond the NVCCFLAGS resolution above, #1037 landed the fp8 warp-per-row kernels. If your CUDA build touches the same block, that's where the conflict will be — everything else in the merges was additive.

I have deliberately not merged the smaller PRs that touch c/deepseek_v4.c or c/colibri.c yet (#1017, #1035, #1038, #827), precisely so your rebase lands against a stable base instead of chasing a moving one. Tell me when you've pushed and I'll sequence the rest around you.

On the overlap with #1031: your numbers (3324-token prefill 271 s → 90 s, decode at 3k context 0.6 → 1.6–1.76 tok/s) are a different order of magnitude from what that PR measures (+4–5% on a 10.7 GB card), and you offered to reconcile with whichever lands first. That's the right spirit and it's noted — the decision on which carries the tier is the maintainer's, and it'll be made on the rebased diff rather than the stacked one.

Also: #1056 vendors third-party headers and adds THIRD_PARTY_NOTICES.md, which is the correct way to do it — thank you for splitting it out as separable rather than folding it into the tier PR.

[dcutugno]

Rebased all three on current dev (03e8677, i.e. with #1054, #988, #1023, #1037, #827…) and force-pushed:

• feat(v4): prefix checkpoints (memory + disk), resumable 4096-token prefill segments, cancellable prefill, gateway prefix hint #1053 → d44b096: its own 4 commits, 9 files (+1.6k). Only conflicts were docs (took the per-engine tool matrix wording from docs(api): map tool calling support by engine #1033) and fix(v4): pack hot pinned experts outside state->mutex (issue #900) #1023's lock-free rows16 packing, where the COLI_V4_ROWS16=0 valve now lives in a shared hot_rows16_wanted() used by both hot_pack_slot_locked and the new hot_pack_slot_prepare.
• feat(v4/cuda): the CUDA tier wired into the DeepSeek V4 engine — generic kernels for any sm_80+, Windows runtime DLL loader, Linux CUDA=1 #1055 → 056b920: 3 commits on top of feat(v4): prefix checkpoints (memory + disk), resumable 4096-token prefill segments, cancellable prefill, gateway prefix hint #1053 (0a82f07 kernels/loader/def/Makefile targets, 8426107 engine wiring, 056b920 docs), 34 files (+7.1k, of which +5.4k is the .cu/.c and the rest tests, loader, docs). c/Makefile: I kept your $(NVCC_STD) / -ftz=false / compute_121 resolution untouched and only added the two DLL targets, the loader test and the DeepGEMM include paths on top; backend_cuda_dsv4.cu supersedes DeepSeek V4: the CUDA kernels, as a tier rather than an engine (rebase of #772) #1054's file (same rescope, our kernels on top). fix(v4): pack hot pinned experts outside state->mutex (issue #900) #1023's hot_pack_slot_prepare gets the same && !store->gpu guard the hit path has (no rows16 repack while experts run from their fp4 mirror).
• build(v4/cuda): vendor the DeepGEMM sm120 headers — cuda-dsv4-dg-dll / DEEPGEMM=1 need no external checkout #1056 → e77b492: 1 commit on top of feat(v4/cuda): the CUDA tier wired into the DeepSeek V4 engine — generic kernels for any sm_80+, Windows runtime DLL loader, Linux CUDA=1 #1055.

Re-verified after the rebase: #1053 CPU-only build; #1055 Windows engine + both DLLs via the make targets, 826-token prompt / 48 greedy tokens → same text as before the rebase (md5 d4e3f23d3d68), generic DLL forced also OK; #1056 make cuda-dsv4-dg-dll with the vendored default. Ready for the rest to be sequenced around it.

[JustVugg]

Merged into dev. DeepSeek V4 now has a CUDA tier in the tree — the kernels from #1054 finally have something calling them.

What decided it, since I did not read all 7,161 lines and won't pretend otherwise:

• It is compile-gated. 46 references to COLI_V4_GPU_TIER; the default build is untouched, so nobody who doesn't ask for it can be broken by it.
• The blast radius is confined to V4. No other engine's .c and none of the shared headers (st.h, json.h, quant.h, tok.h) are modified. expert_store.h, native_quant.h and tensor.h take small additive changes only.
• You extended the gate rather than routing around it — test_deepseek_v4_tiny.py +66. A GPU tier that adds to the token-exactness check instead of exempting itself from it is the version of this PR I was willing to take.
• 19/19 green, and it shrank from 10,677 to 7,161 lines once feat(v4): dual-SSD mirror (COLI_MODEL_MIRROR) for the DeepSeek V4 expert store #988 and DeepSeek V4: the CUDA kernels, as a tier rather than an engine (rebase of #772) #1054 landed, which is exactly why the stack was worth unwinding in order.

Also worth recording: your design decision to keep every stage CPU-canonical with a per-stage fallback, and to accept a GPU stage only when its greedy text matched the CPU reference on the same prompt, is the reason this could be merged on a compile flag rather than after a month of review. The Validation section in the docs — being explicit about what "identical" can and cannot mean across kernels and cache states — is better than most projects manage.

Three follow-ups, none blocking:

1. feat(v4): prefix checkpoints (memory + disk), resumable 4096-token prefill segments, cancellable prefill, gateway prefix hint #1053 still needs your call on V4_PREFIX_CKPT_DISK defaulting on. Everything else there is verified and ready — I checked the V4_PREFILL_CHUNK 64→128 raise myself, including that __m256 sums[] grew with it, and it's correct.
2. build(v4/cuda): vendor the DeepGEMM sm120 headers — cuda-dsv4-dg-dll / DEEPGEMM=1 need no external checkout #1056 (vendoring CUTLASS + DeepGEMM) is a scope decision for @JustVugg rather than a technical one: ~35k lines of third-party headers in-tree, against not needing an external checkout for DEEPGEMM=1. For what it's worth you did the licensing correctly — cutlass/LICENSE.txt, deep_gemm/LICENSE and THIRD_PARTY_NOTICES.md are all present, which is the part people usually get wrong.
3. feat(dsv4): DeepSeek V4 Flash support — CUDA/VRAM expert tiering, DSML tool use, cross-session KV persistence #1031 overlaps this area with a different design; it now needs reconciling against what just landed, and you'd already offered.

Thanks for splitting this the way you did. Landing the kernels with their tests first, then the wiring, made both reviewable — that sequencing was your suggestion and it was the right one.