Skip to content

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

Open
ZacharyZcR wants to merge 1 commit into
JustVugg:devfrom
ZacharyZcR:feat/deepseek-v4-flash
Open

DeepSeek V4: the CUDA kernels, as a tier rather than an engine#772
ZacharyZcR wants to merge 1 commit into
JustVugg:devfrom
ZacharyZcR:feat/deepseek-v4-flash

Conversation

@ZacharyZcR

@ZacharyZcR ZacharyZcR commented Aug 2, 2026

Copy link
Copy Markdown
Contributor

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} dense matmul, batched attention, routed MoE, optional vLLM/FlashInfer/DeepGEMM paths
dsv4_mhc.h, dsv4_quant.h the CPU-side formats the kernels consume
tests/ three GPU tests (dense / attention / MoE) + two pure-CPU unit tests
tools/ the oracles and probes the kernels are validated against
THIRD_PARTY_NOTICES.md licences for the DeepGEMM / FlashInfer / vLLM derived paths

Dropped: c/deepseek_v4.c β€” #165 owns that file β€” plus the c/coli, c/openai_server.py, c/tok.h and docs/api.md edits. Those belong to whichever engine lands, not to a kernel PR. That also removes the file collision @JustVugg warned about: the two PRs each created c/deepseek_v4.c with 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.h includes only <stdint.h>; the .cu includes 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.c streams from disk, backend_cuda.cu holds the hot experts).

DSV4_CUDA_OBJ is empty unless CUDA=1, so a CPU-only build is unaffected.

Makefile

Kernel rules only. deepseek_v4$(EXE) is gone; the .o rules, dsv4-cuda-test and the two header unit tests remain.

Two fixes carried over from the review:

  • VLLM_MHC / DEEPGEMM / FLASHINFER are 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=1 appends instead of replacing. The old override NVCCFLAGS := discarded $(CUDA_GENCODE) β€” silently swapping any CUDA_ARCH for a hardcoded sm_120f β€” and -ccbin, which is how a host whose default g++ is too new for CUDA names a supported host compiler. C++20 now rides on an NVCC_STD variable instead.

Verification

  • colibri, deepseek-v4, cuda-test, bench-omp-grain all still resolve β€” dev's targets are untouched
  • dsv4-cuda-test resolves
  • tests/test_dsv4_mhc and tests/test_dsv4_quant build and pass on CPU (dsv4-mhc: OK, dsv4-quant: OK)
  • NVCCFLAGS checked across CUDA_ARCH = native / portable / sm_90, with and without DEEPGEMM=1: correct standard, architecture preserved, no duplicate -std
  • GPU tests need a device; compile-checked only

Open 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.h exposes today, say so and I will change the header β€” the residency planning is yours and I am not going to duplicate it.

@ZacharyZcR
ZacharyZcR force-pushed the feat/deepseek-v4-flash branch from e535c9b to f6368c0 Compare August 2, 2026 17:30
@JustVugg JustVugg added discussion Proposta / discussione aperta, non un task model-support Supporto a nuovi modelli labels Aug 2, 2026
@ZacharyZcR
ZacharyZcR marked this pull request as ready for review August 2, 2026 22:31
@ZacharyZcR
ZacharyZcR force-pushed the feat/deepseek-v4-flash branch from b9c4160 to 26b42fc Compare August 2, 2026 22:38
ZacharyZcR added a commit to ZacharyZcR/colibri that referenced this pull request Aug 3, 2026
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.
ZacharyZcR added a commit to ZacharyZcR/colibri that referenced this pull request Aug 3, 2026
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.
ZacharyZcR added a commit to ZacharyZcR/colibri that referenced this pull request Aug 3, 2026
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.
@jazir555

jazir555 commented Aug 3, 2026

Copy link
Copy Markdown

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?

@ZacharyZcR

Copy link
Copy Markdown
Contributor Author

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 gate

This 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 cudaMalloc + cudaMemcpy at load, and backend_cuda_dsv4.cu contains no mmap, no pread, and no offload of any kind. If it does not fit in VRAM, it does not start β€” it will not degrade to RAM or disk.

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 run

GLM-5.2 744B, on the normal tiered path β€” this is the configuration the engine is actually built for:

  • dense part resident in RAM at int4: 9.9 GB
  • peak RSS during chat: ~20 GB (auto-capped, fits your 32 GB)
  • routed experts stream from disk: you need ~370 GB of free NVMe
  • your 12 GB card is still useful as the hot-expert VRAM tier

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 .coli_usage learns your workload and pins the hot experts, and a faster NVMe moves it a lot, since cold decode costs ~11 GB of reads per token. But it is not going to feel like a hosted API on that hardware. See docs/benchmarks.md.

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.

@jazir555

jazir555 commented Aug 4, 2026

Copy link
Copy Markdown

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?

@jazir555

jazir555 commented Aug 4, 2026

Copy link
Copy Markdown

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

@rofl0r

rofl0r commented Aug 4, 2026

Copy link
Copy Markdown

the other PR you cross-posted to does what you want, afaict.

@rafpigna

rafpigna commented Aug 4, 2026

Copy link
Copy Markdown

This 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 cudaMalloc + cudaMemcpy at load, and backend_cuda_dsv4.cu contains no mmap, no pread, and no offload of any kind. If it does not fit in VRAM, it does not start β€” it will not degrade to RAM or disk.

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.

@JustVugg

JustVugg commented Aug 4, 2026

Copy link
Copy Markdown
Owner

@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:

One hierarchy, not one memory threshold. VRAM, RAM, and NVMe are placement tiers for the same weights; limited fast memory changes speed, not model semantics.

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: colibri.c streams from disk while backend_cuda.cu holds the hot experts in VRAM.

@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 c/deepseek_v4.c as a new file, with entirely different contents. They also collide on c/Makefile, c/coli, c/openai_server.py and c/tok.h. Whichever merges first makes the other conflict on the spot. If the GPU work stays inside backend_cuda_dsv4.* and the engine file belongs to the streaming implementation, both survive and they compose. If they keep colliding, one of you loses three weeks of work to a rebase for no reason.

The optional build paths

A policy question I want to settle now, because it will come up again with other backends.

VLLM_MHC Γ— DEEPGEMM Γ— FLASHINFER is eight build configurations, and CI cannot compile a single one of them: each needs an external checkout, and one needs Blackwell hardware. Untested build paths rot silently and become maintenance debt for whoever is here in six months.

There is also a concrete defect in one of them:

override NVCCFLAGS := -O3 ... -gencode=arch=compute_120f,code=sm_120f ...

That := replaces NVCCFLAGS, architecture included. With DEEPGEMM=1 the build is sm_120 only β€” RTX 50-series. Any -gencode the user set is discarded. So that path is not "for CUDA users", it is for Blackwell owners, and nobody will notice it breaking because nothing ever builds it. FLASHINFER uses += and does not have this problem.

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 := becomes +=.

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.

@ZacharyZcR

Copy link
Copy Markdown
Contributor Author

@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 PR

Your 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, backend_cuda_dsv4.* holds the hot experts in VRAM.

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:

  • c/deepseek_v4.c is @DrewZt's. I am removing it from this branch rather than negotiating over it β€” it is the file that makes the two PRs collide on sight, and the streaming implementation is the one that should own it. Same for the c/coli, c/openai_server.py and c/tok.h edits: they belong to whichever engine lands, not to a kernel PR.
  • What stays is backend_cuda_dsv4*.{cu,h}, dsv4_mhc.h, dsv4_quant.h, their tests and oracles, and the Makefile rules that build them.
  • The result will not run DeepSeek V4 on its own. That is the point β€” it is a tier for Add DeepSeek V4 target-only CPU inferenceΒ #165 to call, not a second engine.

@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.

:= β†’ += is pushed, and it was worse than one line

c8810c1. You were right about the mechanism and the consequence, and there was a second casualty in it:

override NVCCFLAGS := -O3 -std=c++20 -gencode=arch=compute_120f,code=sm_120f ...

That also discarded -ccbin $(NVCC_CCBIN) β€” the flag a host whose default g++ is too new for CUDA uses to name a supported host compiler. So DEEPGEMM=1 did not merely pin the architecture to Blackwell; on those hosts it broke the build outright, in a path nothing compiles.

C++20 is the only thing DeepGEMM actually needed from that line, so the standard is now a variable an optional backend can raise without rewriting everything around it:

NVCC_STD ?= c++17
NVCCFLAGS ?= -O3 -std=$(NVCC_STD) $(CUDA_GENCODE) -Xcompiler=-Wall,-Wextra
...
override NVCC_STD  := c++20
override NVCCFLAGS += --expt-relaxed-constexpr --expt-extended-lambda -DCOLI_DSV4_DEEPGEMM -I"$(DEEPGEMM_HOME)"

Checked rather than asserted, across CUDA_ARCH = native / portable / sm_90, with and without DEEPGEMM=1: the standard becomes c++20, the architecture stays whatever the user asked for (portable's six gencodes survive), no duplicate -std, default builds byte-identical, and the missing-DEEPGEMM_HOME error still fires.

Unsupported build paths β€” agreed, and stated in the Makefile

Applied as you framed it, and I think uniformly is the right way to have it. The three optional backends now carry this where someone enabling them will actually read it:

VLLM_MHC, DEEPGEMM and FLASHINFER are eight build configurations and CI compiles none of them. […] They gate nothing and block nothing β€” the default build (all three off) is the one CI compiles and the one to report bugs against.

Framing it as a statement about verification rather than quality is yours and it is the correct one. I would rather these be honestly labelled than quietly implied to work.

#773

Closed, with a pointer here. It was this PR's 25 commits plus one, over the same 46 files β€” my mistake for opening both. The extra commit (long-context cache) goes back up on its own after the split, since it is built on a layout that is about to move anyway.


To @jazir555 and @rafpigna directly: your objection is what got this changed, and the answer to "why would someone use this over llama.cpp" should not have needed a maintainer to supply. On the hardware you asked about, GLM-5.2 744B on the tiered path is the honest answer today, and DeepSeek V4 should get to the same bar rather than sit above it.

@ZacharyZcR
ZacharyZcR force-pushed the feat/deepseek-v4-flash branch from c8810c1 to 374d2d6 Compare August 6, 2026 02:31
ZacharyZcR added a commit to ZacharyZcR/colibri that referenced this pull request Aug 6, 2026
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.
@ZacharyZcR ZacharyZcR changed the title feat: add native DeepSeek V4 Flash inference DeepSeek V4: the CUDA kernels, as a tier rather than an engine Aug 6, 2026
@ZacharyZcR

Copy link
Copy Markdown
Contributor Author

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. c/deepseek_v4.c is #165's, and the c/coli, c/openai_server.py, c/tok.h and docs/api.md edits went with it. What remains is backend_cuda_dsv4*, the two format headers, their tests and their oracles. 6,765 lines β†’ 4,940, 25 commits β†’ 1.

That also clears the collision you flagged. Both PRs created c/deepseek_v4.c with different contents; #165 landed first and this branch went CONFLICTING within the hour β€” which is also why its CI stopped running entirely, since GitHub cannot build a merge commit for a conflicted PR. It is mergeable again now.

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 (backend_cuda_dsv4.h includes only <stdint.h>, the .cu only its own header and the CUDA runtime), so nothing here reaches into an engine. #165 calls in, nothing calls out. Same shape as colibri.c + backend_cuda.cu.

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:

  • The three optional backends are marked UNSUPPORTED in place, as a statement about verification: eight configurations, CI compiles none, they gate nothing, and the default build is the one to report bugs against.
  • := β†’ +=. Worth noting it was worse than one line: it also discarded -ccbin $(NVCC_CCBIN), so on a host whose default g++ is too new for CUDA, DEEPGEMM=1 did not merely pin the arch to Blackwell β€” it broke the build, in a path nothing compiles. C++20 now rides on an NVCC_STD variable.

#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 backend_cuda_dsv4.h exposes, I will change the header to match. Residency planning is yours; I only want the kernels to be the thing you call.

Verified: colibri / deepseek-v4 / cuda-test / bench-omp-grain all still resolve untouched, dsv4-cuda-test resolves, and the two CPU unit tests build and pass (dsv4-mhc: OK, dsv4-quant: OK). GPU tests need a device and are compile-checked only.

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.
@ZacharyZcR
ZacharyZcR force-pushed the feat/deepseek-v4-flash branch from 374d2d6 to 280f8b3 Compare August 7, 2026 05:00
@rafpigna

rafpigna commented Aug 8, 2026

Copy link
Copy Markdown

To @jazir555 and @rafpigna directly: your objection is what got this changed, and the answer to "why would someone use this over llama.cpp" should not have needed a maintainer to supply. On the hardware you asked about, GLM-5.2 744B on the tiered path is the honest answer today, and DeepSeek V4 should get to the same bar rather than sit above it.

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!
If you are available to wait 24 hours to have 1 token output, is up to you, not to the engine, nor the model, nor the hardware.
This makes, in my opinion as user, Colibri really different from many other projects that can do basically the same thing. Also llama.cpp can load a model bigger than your vram/ram, but has no optimizations designed for this usage case.

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 :)

@JustVugg

JustVugg commented Aug 8, 2026

Copy link
Copy Markdown
Owner

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 decision

CUDA in colibri has to be ours. The two vLLM-derived files are the part I am declining:

backend_cuda_dsv4_qkv_vllm.cu    +739
backend_cuda_dsv4_mhc_vllm.cu    +535
backend_cuda_dsv4_flashinfer.cu   +71
THIRD_PARTY_NOTICES.md            +60
                                 ─────
                                 1,723 of 4,940  (35%)

The other 3,200 β€” backend_cuda_dsv4.cu at 1,719 lines, dsv4_quant.h, the oracles, the probes, the benchmarks β€” is your work and I want it.

Why, and it is not about licences

The argument for vendoring is real and I am not pretending otherwise: those kernels are battle-tested, they are in-tree rather than a pip install, and writing MLA and MoE CUDA from scratch is months. If the reasoning were only "purity" I would take them.

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:

expert matmul     121.4 s
expert disk       663.4 s

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 way

CUDA 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 quant.h's CPU path β€” including the detail of building the exponent as a bit pattern rather than exp2f, because that reproduces the documented s=0 -> +0 and s=255 -> +inf edges exactly where exp2f would silently diverge.

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 scales[o] multiplication that should not apply to fmt=7 β€” a vendored kernel would not have had that test, and that bug would have shipped.

What I am asking

Split it. Send the DeepSeek V4 CUDA tier without the two _vllm.cu files and the FlashInfer shim, and with whatever your own kernels need to stand alone. If that leaves a gap where the vLLM path was doing the work, say so and we will decide together whether to write it or to leave that path on CPU for now β€” leaving it on CPU is an acceptable answer, given where the time actually goes.

The oracles and probes I want regardless. dsv4_vllm_reference.py and dsv4_vllm_shape_probe.py are fine as validation tools β€” using vLLM to check our output is a different thing from shipping their kernels, and it is a good idea.

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.

ErikTromp pushed a commit to SensAI-PT/aviary-hy3 that referenced this pull request Aug 9, 2026
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.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

discussion Proposta / discussione aperta, non un task model-support Supporto a nuovi modelli

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants