feat: add DeepSeek V4 long-context inference by ZacharyZcR · Pull Request #773 · JustVugg/colibri

ZacharyZcR · 2026-08-02T07:05:34Z

Summary

Stacked on #772. This PR contains only the long-context allocation foundation:

store the 128-token FlashInfer SWA cache as a ring instead of allocating by maximum context
cap sparse-attention scratch space at the model top-k width
accept up to 32,768 prompt + output tokens
avoid allocating the duplicate CPU compressed cache in the all-GPU path

The branch was restacked after #772 was rebased; its delta is now limited to c/backend_cuda_dsv4.cu and c/deepseek_v4.c.

Correctness boundary

the exact sparse path is validated only through 2K tokens, where all compressed candidates fit in index_topk=512
32K allocation support is not a claim of correct 32K inference
contexts beyond 2K remain incomplete until the checkpoint Lightning Indexer is connected

Verification

targeted long-context tests passed
local CPU build passed
the base DeepSeek V4: the CUDA kernels, as a tier rather than an engine #772 backend passed a full CUDA 13 / SM120 build on the 6x RTX 5090 host

Status

Draft and dependent on #772. Keep it stacked until the base backend merges; then rebase and implement the FP8 Lightning Indexer query/key cache, MQA logits, top-512 selection, and real long-context A/B validation.

c/Makefile appears in most open PRs, and any two that each append a target name to the single shared .PHONY conflict on that line by construction -- which is exactly what happens between this branch and the DeepSeek V4 work (JustVugg#772/JustVugg#773), and between this and the other two lint branches. The TEST_RULES block above already documents this failure mode for TEST_BINS and fixes it by deriving the list instead of hand-maintaining one line. .PHONY had the same problem and not the same treatment. make accumulates multiple .PHONY declarations, so a new target now adds a line rather than editing a line everyone else is also editing.

c/Makefile appears in most open PRs, and any two that each append a target name to the single shared .PHONY conflict on that line by construction. That is exactly what happens between this branch, the env-registry branch, the sanitizer branch and the DeepSeek V4 work (JustVugg#772/JustVugg#773) -- four branches, one line, four conflicts that are pure bookkeeping. The TEST_RULES block above already documents this failure mode for TEST_BINS and fixes it by deriving the list rather than hand-maintaining a line. .PHONY had the same problem and never got the same treatment. make accumulates multiple .PHONY declarations, so adding a target now means adding a line instead of editing a line every other branch is also editing.

The section added here credits the research the expert-placement and routing work draws on. It does not credit the code -- the containers, tokenizers, codecs and compat layers this engine reimplements or links against, several of which are named in the sources and nowhere else. Every entry points at where it is used, because a credit nobody can check is decoration: safetensors c/st.h, 13 references, including its fp8 and I64 dtypes tiktoken c/tok.h reimplements byte_pair_encode exactly llama.cpp grammar.h:14 (GBNF subset + set-of-stacks PDA), colibri.c:1956 (Metal newBufferWithBytesNoCopy residency) vLLM colibri.c:5903 -- where the final norm lands vs the LM head transformers tools/oracle-requirements.txt, the token-exact CI oracle DietGPU 6 references in Makefile, the COLI_ANS compressed tier rocWMMA backend_gpu_compat.h maps nvcuda::wmma onto it, which is what lets one .cu compile for both vendors Deliberately NOT listed: SGLang and FlashAttention. Both are obvious names to put in a file like this, and neither appears anywhere in the tree -- no reference, no reimplementation, no comparison. Listing them would be flattering and false. FlashInfer, DeepGEMM and TileLang do appear, but only on the DeepSeek V4 branches (JustVugg#772/JustVugg#773); they belong in this list when that work lands, not before.

ZacharyZcR · 2026-08-05T03:31:42Z

Closing this in favour of #772, as @JustVugg asked in #772.

He is right that this is the same PR. #773 carries #772's 25 commits verbatim plus one — ed563b4 Add DeepSeek V4 long-context cache foundation — over the same 46 files. Two reviewers could have read the same 6,700 lines without either noticing, and that is my fault for opening both.

#772 is the one with the review on it, so that is the one that stays. The long-context commit is not lost: it goes back up on its own once #772's shape is settled, which is not a small caveat right now — the blocking request there is an offload path, and the file layout is likely to move so the engine file belongs to #165's streaming implementation. A long-context cache built on top of the current layout would need redoing anyway.

Better to rebuild it on whatever #772 becomes than to keep a stale superset open.

The section added here credits the research the expert-placement and routing work draws on. It does not credit the code -- the containers, tokenizers, codecs and compat layers this engine reimplements or links against, several of which are named in the sources and nowhere else. Every entry points at where it is used, because a credit nobody can check is decoration: safetensors c/st.h, 13 references, including its fp8 and I64 dtypes tiktoken c/tok.h reimplements byte_pair_encode exactly llama.cpp grammar.h:14 (GBNF subset + set-of-stacks PDA), colibri.c:1956 (Metal newBufferWithBytesNoCopy residency) vLLM colibri.c:5903 -- where the final norm lands vs the LM head transformers tools/oracle-requirements.txt, the token-exact CI oracle DietGPU 6 references in Makefile, the COLI_ANS compressed tier rocWMMA backend_gpu_compat.h maps nvcuda::wmma onto it, which is what lets one .cu compile for both vendors Deliberately NOT listed: SGLang and FlashAttention. Both are obvious names to put in a file like this, and neither appears anywhere in the tree -- no reference, no reimplementation, no comparison. Listing them would be flattering and false. FlashInfer, DeepGEMM and TileLang do appear, but only on the DeepSeek V4 branches (JustVugg#772/JustVugg#773); they belong in this list when that work lands, not before.

ZacharyZcR force-pushed the feat/deepseek-v4-long-context branch 9 times, most recently from 205ac4e to a601640 Compare August 2, 2026 21:48

JustVugg added discussion Proposta / discussione aperta, non un task model-support Supporto a nuovi modelli labels Aug 2, 2026

ZacharyZcR force-pushed the feat/deepseek-v4-long-context branch from a601640 to ae1f9e6 Compare August 2, 2026 22:29

ZacharyZcR marked this pull request as ready for review August 2, 2026 22:31

ZacharyZcR added 17 commits August 3, 2026 06:37

feat: add native DeepSeek V4 Flash inference

877d5d7

feat: accelerate DeepSeek V4 attention

fe2adbb

fix: enable sparse MLA for tensor parallel heads

c7c53bc

feat: reuse vLLM fused QKV cache kernel

0857fa6

fix: compile FlashInfer DeepSeek backend

2bd7d59

fix: bound sparse MLA to the sliding window

ba863f1

feat: connect compressed KV to sparse MLA

f3cabcb

fix: exclude sliding window from compressed MLA cache

ce1a465

fix: size sparse MLA splits to attention window

6a1b491

feat: serve DeepSeek V4 chat and tool calls

ba030c5

fix: preload DeepSeek V4 weights in serve mode

e2e147b

fix: suppress DeepSeek V4 stop token

f598ee8

fix: enable DeepSeek V4 service fast path

98bcdec

fix: retain all C4A compressed candidates

27d93f5

fix: avoid double quantizing compressed KV

74f6440

fix: skip empty compressed attention splits

bf9a613

fix: retain every compressor partial state

617a3c1

ZacharyZcR added 9 commits August 3, 2026 06:37

feat: expose batched DeepSeek V4 prefill kernels

474b634

feat: add batched mHC prefill state path

9336be9

feat: batch initial mHC prefill

e4be850

feat: add batched DeepGEMM projections

d61c2e1

feat: add batched sparse attention core

7cbd882

feat: batch grouped attention output

3988431

feat: add batched routed MoE

dc0e35d

feat: add layer-major batched prefill

26b42fc

Add DeepSeek V4 long-context cache foundation

ed563b4

ZacharyZcR force-pushed the feat/deepseek-v4-long-context branch from ae1f9e6 to ed563b4 Compare August 2, 2026 22:40

JustVugg mentioned this pull request Aug 4, 2026

DeepSeek V4: the CUDA kernels, as a tier rather than an engine #772

Open

ZacharyZcR closed this Aug 5, 2026

Navigation Menu

feat: add DeepSeek V4 long-context inference - #773

feat: add DeepSeek V4 long-context inference#773
ZacharyZcR wants to merge 26 commits into
JustVugg:devfrom
ZacharyZcR:feat/deepseek-v4-long-context

ZacharyZcR commented Aug 2, 2026 •

edited

Loading

Uh oh!

ZacharyZcR commented Aug 5, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

ZacharyZcR commented Aug 2, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Correctness boundary

Verification

Status

Uh oh!

ZacharyZcR commented Aug 5, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

ZacharyZcR commented Aug 2, 2026 •

edited

Loading