DeepSeek V4: the CUDA kernels, as a tier rather than an engine - #772
DeepSeek V4: the CUDA kernels, as a tier rather than an engine#772ZacharyZcR wants to merge 1 commit into
Conversation
e535c9b to
f6368c0
Compare
b9c4160 to
26b42fc
Compare
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.
|
How much VRAM and RAM does this take to load/use? I have 32 GB RAM and 12 GB vram, this will this fit in my card/ram given that this is much smaller than GLM 5.2 correct? |
|
Short answer: no, this one will not fit β and the reason is worth spelling out, because "smaller than GLM-5.2" points the wrong way here. Why parameter count is not the gateThis backend keeps the 145.3 GiB checkpoint fully resident in VRAM. The measured host in this PR is 6Γ RTX 5090 (192 GB VRAM total). There is no streaming path in it: every weight goes through So the gate is not "how many parameters", it is "does this path have an offload tier". DeepSeek V4 Flash has fewer parameters than GLM-5.2 and a much higher hardware floor, because GLM-5.2 runs on colibrΓ¬'s tiered streaming engine and this backend does not. That inversion is confusing and it is on me for not stating it in the PR description β I will add a hardware-requirements line. What 32 GB RAM + 12 GB VRAM does runGLM-5.2 744B, on the normal tiered path β this is the configuration the engine is actually built for:
Being straight with you about speed, since that matters more than whether it boots: on the project's reference box (WSL2, 12 cores, 25 GB RAM, ~1 GB/s drive) a cold cache measures 0.05β0.1 tok/s β roughly ten to twenty seconds per token. It gets substantially better as If you want something responsive on that machine rather than something impressive, OLMoE 7B is the one to start with. Happy to help you get GLM-5.2 configured if you want to try it β the disk space is the real prerequisite, not the RAM. |
|
Can the DeepSeek engine not stream via the same method as GLM? Is that a limitation of the architecture of the model or just not implemented at this time? |
|
Requiring the full amount of VRAM to run this kind of defeats the entire purpose of Colibri which is to make Giant models run on mid tier consumer hardware, if that's the case why would someone use this over llama.cpp? Colibri's entire intent as far as I understand is to make these models more accessible. This sounds like just a normal inference engine which doesn't meet the intent of Colibri. @JustVugg |
|
the other PR you cross-posted to does what you want, afaict. |
Maybe it's me that doesn't understand too much about it, but the whole Colibri engine "reason to exists" is not to serve BIG models using all memory tiers (RAM - VRAM - NVME) as a one big tier? What the reason to use Colibri if you have 6x5090 with 192GB VRAM? llama.cpp should work with much less efforts and ready out of the box after installation. I can't get the point of this PR. |
|
@jazir555 @rafpigna β you are asking the right question, and it should be answered by me rather than by the PR author. ColibrΓ¬'s promise is a hierarchy, not a threshold. The README states it as a core technique:
A backend that will not start unless 145.3 GiB fits in VRAM turns that into exactly one threshold. Someone with a 3090 or a 4070 does not get a slower model here β they get nothing at all. On this project, running out of fast memory is meant to cost you tokens per second, never the ability to run the model. That is the whole reason colibrΓ¬ exists, and it is why I agree with the objection in this thread. So there is one blocking request, and it is not negotiable: this backend needs an offload path. Every other GPU backend in the tree β CUDA, Metal, Vulkan β is a tier: it holds the hot part and falls back for the rest. This one holds everything or refuses to start. Where I think this should go, and it is not "close it"#165 and this PR are each missing exactly what the other has. @DrewZt's #165 is the streaming machinery β expert store, prefetch, resource plan β and it contains no CUDA at all: not one kernel. This PR is the CUDA kernels, and it has no tier. Neither one alone gives anybody DeepSeek V4 on hardware they own. Put together, they are the same shape as every other model here: @ZacharyZcR, you described that pattern yourself earlier in this thread β you told @jazir555 that a 12 GB card "is still useful as the hot-expert VRAM tier". That is the answer for DeepSeek too, and your kernels are the part that does not exist yet. I would much rather have that than either PR merged alone. Which is why the file layout matters right now, and not later: This PR and #165 both create The optional build pathsA policy question I want to settle now, because it will come up again with other backends.
There is also a concrete defect in one of them: override NVCCFLAGS := -O3 ... -gencode=arch=compute_120f,code=sm_120f ...That My position, applied uniformly and not only here: a build path CI never compiles is marked unsupported in the docs, and neither blocks nor gates anything. And One housekeeping item#773 is this PR β same 46 files, the same 25 commit SHAs, plus one commit. Please close one or mark #773 explicitly as stacked; as it stands, two reviewers can read the same 6,700 lines twice without realising it. To answer @jazir555's original question plainly: 32 GB RAM + 12 GB VRAM runs GLM-5.2 744B today on the tiered path. That is the bar a DeepSeek V4 backend should be aiming at, and between these two PRs the pieces to clear it already exist. |
|
@JustVugg β I accept the blocking request, and @jazir555 and @rafpigna were right to push on it. A backend that refuses to start below 145.3 GiB is a threshold, and this project's claim is a hierarchy. I do not want to argue the exception. Accepting the split, and I would rather have it than this PRYour read is correct and it is the cheaper path for everyone: #165 has the streaming machinery and no CUDA; this has the CUDA and no tier. Composed, it is the shape every other model here already has β the engine streams from disk, So the scope of this PR becomes the kernels and nothing else, and #165 owns the engine. Concretely, so it is not just agreement in principle:
@DrewZt, tell me if you would rather own the boundary yourself; I will fit the kernels to whatever entry points your expert store wants rather than the other way round. The residency-planning side is yours and I do not intend to duplicate it. I will rework this branch on that basis rather than push more onto the current shape. If you would prefer it closed and reopened clean once the split is done, say so and I will close it β a 6,700-line PR whose scope just changed is not a thing anyone should have to re-review in place.
|
c8810c1 to
374d2d6
Compare
This is JustVugg#772 rescoped. It was an engine plus its kernels and it would not start below 145.3 GiB of VRAM, which makes it a threshold -- the one thing this project does not ship. @jazir555 and @rafpigna said so, @JustVugg agreed, and the request was not negotiable: a GPU backend here holds the hot part and falls back for the rest. The split @JustVugg proposed is the right one. JustVugg#165 landed the DeepSeek V4 engine with the streaming machinery and no CUDA at all; this was the CUDA with no tier. Neither runs the model on hardware anyone owns. So the engine file belongs to JustVugg#165 and this keeps only the kernels: backend_cuda_dsv4*.{cu,h} dense matmul, batched attention, routed MoE dsv4_mhc.h, dsv4_quant.h the CPU-side formats the kernels consume tests/ three GPU tests, two pure-CPU unit tests tools/ oracles and probes the kernels are checked against Dropped from the old branch: c/deepseek_v4.c (JustVugg#165 owns it), and the c/coli, c/openai_server.py, c/tok.h and docs/api.md edits, which belong to whichever engine lands rather than to a kernel PR. The kernels are self-contained by construction -- backend_cuda_dsv4.h includes only <stdint.h>, and the .cu only its own header and the CUDA runtime -- so they build and their CPU tests run with no engine present. That is what makes them usable as a tier: JustVugg#165 calls in, nothing calls out. Makefile: the kernel rules only. deepseek_v4$(EXE) is gone; the .o rules, dsv4-cuda-test and the two header unit tests remain. DSV4_CUDA_OBJ is empty unless CUDA=1. VLLM_MHC/DEEPGEMM/FLASHINFER are marked UNSUPPORTED in place -- eight configurations, none compiled by CI, each needing an external checkout -- and DEEPGEMM's flag line appends instead of replacing, so it no longer discards $(CUDA_GENCODE) and -ccbin. Verified: colibri, deepseek-v4, cuda-test and bench-omp-grain all still resolve; dsv4-cuda-test resolves; test_dsv4_mhc and test_dsv4_quant build and pass on CPU. The GPU tests need a device and are compile-checked only.
|
Done β force-pushed the rescope rather than opening a new PR, so the discussion above stays attached to it. Say the word if you would still rather have it closed and reopened clean. What changed: the engine is gone. That also clears the collision you flagged. Both PRs created On the blocking request. I am not claiming this satisfies it. A tier is not an offload path until something drives it, and the thing that drives it is #165's expert store. What this PR does is make that possible: the kernels are self-contained ( Judged as "does this let someone with a 3090 run DeepSeek V4", this PR alone still does not β and should not be merged as if it did. Judged as "is this the missing GPU half of #165", that is what I have tried to make it. Both Makefile items are in, applied the way you framed them:
#773 is closed. It was this branch's 25 commits plus one, over the same 46 files. @DrewZt β the offer stands and it is not a formality: if your expert store wants a different call shape than Verified: |
This is JustVugg#772 rescoped. It was an engine plus its kernels and it would not start below 145.3 GiB of VRAM, which makes it a threshold -- the one thing this project does not ship. @jazir555 and @rafpigna said so, @JustVugg agreed, and the request was not negotiable: a GPU backend here holds the hot part and falls back for the rest. The split @JustVugg proposed is the right one. JustVugg#165 landed the DeepSeek V4 engine with the streaming machinery and no CUDA at all; this was the CUDA with no tier. Neither runs the model on hardware anyone owns. So the engine file belongs to JustVugg#165 and this keeps only the kernels: backend_cuda_dsv4*.{cu,h} dense matmul, batched attention, routed MoE dsv4_mhc.h, dsv4_quant.h the CPU-side formats the kernels consume tests/ three GPU tests, two pure-CPU unit tests tools/ oracles and probes the kernels are checked against Dropped from the old branch: c/deepseek_v4.c (JustVugg#165 owns it), and the c/coli, c/openai_server.py, c/tok.h and docs/api.md edits, which belong to whichever engine lands rather than to a kernel PR. The kernels are self-contained by construction -- backend_cuda_dsv4.h includes only <stdint.h>, and the .cu only its own header and the CUDA runtime -- so they build and their CPU tests run with no engine present. That is what makes them usable as a tier: JustVugg#165 calls in, nothing calls out. Makefile: the kernel rules only. deepseek_v4$(EXE) is gone; the .o rules, dsv4-cuda-test and the two header unit tests remain. DSV4_CUDA_OBJ is empty unless CUDA=1. VLLM_MHC/DEEPGEMM/FLASHINFER are marked UNSUPPORTED in place -- eight configurations, none compiled by CI, each needing an external checkout -- and DEEPGEMM's flag line appends instead of replacing, so it no longer discards $(CUDA_GENCODE) and -ccbin. Verified: colibri, deepseek-v4, cuda-test and bench-omp-grain all still resolve; dsv4-cuda-test resolves; test_dsv4_mhc and test_dsv4_quant build and pass on CPU. The GPU tests need a device and are compile-checked only.
374d2d6 to
280f8b3
Compare
First of all: thanks for all the amazing job to @JustVugg and also to you for the DeepseekV4 / CUDA tier. I just want to make clear that I'm just an actual Colibri user and as anyone here my first test was with GLM5.2 that is a MASSIVE model, just to feel that "dream" come true: no limits in the model you can run, if you have disk space to download a 4TB model, you can run it, also on a 4GB VRAM + 16 GB RAM (or less), the drawback is that it will be painfully slow, but it will run! So since my GLM5.2 test was ok, but for me almost unusable at less than 1 tok/s, I started to search if anyone asked to support a "mid-tier" model like the last DeepSeek V4 Flash, and found your PR. So when I had read the original PR I was a little bit "selfish", thinking "why block the model support on how much VRAM you have, when Colibri doesnt block a 1 TB (Kimi K3) to run also if you have no vram at all? Why I will not able to run DS4 Flash if I have only 10GB VRAM that I want to use to have a micro-speed up? So I was really honest when I wrote "Maybe it's me that cant understand" :) So please forgive me if it seemed my question was rude, aggressive or provocative, It was not my means at all! Anyway I'm really happy that all the Colibri users now and in the future will be able to use your contribution to the project, like me, because I have a CUDA GPU with 10GB VRAM, that will not make the model so much faster like a cloud API or 5x5090, but for sure will be probably a little bit faster than running totally on CPU/RAM. Hope to see this PR merged soon :) |
|
I have a decision on this, and I want to give you the reasoning rather than a verdict, because it is 4,940 lines and most of them are yours. The decisionCUDA in colibri has to be ours. The two vLLM-derived files are the part I am declining: The other 3,200 β Why, and it is not about licencesThe argument for vendoring is real and I am not pretending otherwise: those kernels are battle-tested, they are in-tree rather than a It is not. vLLM's kernels are written for the regime opposite to ours. They assume the weights are already in VRAM and they minimise compute time. Compute has never been colibri's bottleneck. From the v1.5.0 regression report in #856, on a 256 GB box: Zeroing the matmul entirely buys 15%. The disk path is where this engine lives or dies, and it is the one thing vLLM does not solve, because it never has to. So importing them makes us worse at being colibri β 1,723 lines nobody here can maintain, debug or evolve β while leaving us worse than vLLM at being vLLM, because they improve those kernels full time. Anyone with the VRAM to hold everything resident does not need us. They already have vLLM, and that is the correct choice for them. And there is a second requirement that points the same wayCUDA here has to work on small cards and scale across several of them. A 16 GB consumer GPU streaming experts from disk, and a box with four of them sharing a tier, are the two shapes our users actually have β see #687, #767, #885, all on single consumer cards. Kernels written for datacenter parts with full residency are not tuned for either, and we would not be able to retune them. The shape I want is one you already shipped#819 landed yesterday: MXFP4 decode on CUDA for Kimi K3. Eighty-three lines, written by you, verified on an RTX 4070 against That is the model. Small, targeted, ours, and checked against a reference we already trust. Your test earned its place on the first run by catching a What I am askingSplit it. Send the DeepSeek V4 CUDA tier without the two The oracles and probes I want regardless. I am aware this is the largest PR you have sent and that a third of it is being turned down after it was written. That is on me for not saying this when you opened it eight days ago instead of now. If it helps: nothing else you have sent has been declined, four of your PRs went in yesterday alone, and this is a direction call rather than a judgement of the work. |
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.
Rescoped. This was a DeepSeek V4 engine plus its CUDA kernels, and it would not start below 145.3 GiB of VRAM. @jazir555 and @rafpigna pointed out that a backend which refuses to run on hardware people own is the opposite of what colibrì is for; @JustVugg agreed and made an offload path a blocking requirement. That was right, and this PR no longer tries to be an engine.
#165 landed the DeepSeek V4 engine β streaming machinery, expert store, resource plan, and no CUDA at all. This PR is the other half: the kernels, with no engine.
What this contains
backend_cuda_dsv4*.{cu,h}dsv4_mhc.h,dsv4_quant.htests/tools/THIRD_PARTY_NOTICES.mdDropped:
c/deepseek_v4.cβ #165 owns that file β plus thec/coli,c/openai_server.py,c/tok.handdocs/api.mdedits. Those belong to whichever engine lands, not to a kernel PR. That also removes the file collision @JustVugg warned about: the two PRs each createdc/deepseek_v4.cwith different contents, and #165 got there first.6,765 lines β 4,940, and the 25 commits are collapsed to one, because reviewing the old history against the new scope would have wasted everyone's time.
Why it works as a tier
The kernels are self-contained by construction.
backend_cuda_dsv4.hincludes only<stdint.h>; the.cuincludes only its own header and the CUDA runtime. Nothing here reaches into an engine, so #165 calls in and nothing calls out β which is the shape every other GPU backend in the tree already has (colibri.cstreams from disk,backend_cuda.cuholds the hot experts).DSV4_CUDA_OBJis empty unlessCUDA=1, so a CPU-only build is unaffected.Makefile
Kernel rules only.
deepseek_v4$(EXE)is gone; the.orules,dsv4-cuda-testand the two header unit tests remain.Two fixes carried over from the review:
VLLM_MHC/DEEPGEMM/FLASHINFERare marked UNSUPPORTED in place. Eight build configurations, CI compiles none of them, each needs an external checkout and one wants Blackwell. As @JustVugg framed it, that is a statement about verification rather than quality β they gate nothing, and the default build is the one to report bugs against.DEEPGEMM=1appends instead of replacing. The oldoverride NVCCFLAGS :=discarded$(CUDA_GENCODE)β silently swapping anyCUDA_ARCHfor a hardcodedsm_120fβ and-ccbin, which is how a host whose defaultg++is too new for CUDA names a supported host compiler. C++20 now rides on anNVCC_STDvariable instead.Verification
colibri,deepseek-v4,cuda-test,bench-omp-grainall still resolve β dev's targets are untoucheddsv4-cuda-testresolvestests/test_dsv4_mhcandtests/test_dsv4_quantbuild and pass on CPU (dsv4-mhc: OK,dsv4-quant: OK)NVCCFLAGSchecked acrossCUDA_ARCH=native/portable/sm_90, with and withoutDEEPGEMM=1: correct standard, architecture preserved, no duplicate-stdOpen question for @DrewZt
The entry points are mine to fit to yours, not the other way round. If the expert store wants a different call shape than what
backend_cuda_dsv4.hexposes today, say so and I will change the header β the residency planning is yours and I am not going to duplicate it.