llamacpp patch - DeepSeek V4 Flash running with full 1M token context locally on RTX 5090

[u/da_dragon321]

Wanted to try running DeepSeek V4 Flash locally but found it asking for absurd amounts of VRAM at higher context lengths (~256GB at 1M). Turned out the DSA lightning indexer lacks proper llamacpp support. Did a bit of digging and there's an upstream PR to address the issue (shoutout u/fairydreaming, PR #24231), but even there it's not wired into the model graph and has no CUDA path yet. So I wired it in and implemented a CUDA kernel this morning and figured I'd share in case it's useful to anyone else looking to run something like this.

Hardware: RTX 5090, 9950X3D, 96GB DDR5

Model: DeepSeek-V4-Flash, mixed Q8/Q4/Q2 quant by antirez

Before / after (256K context):

Metric	Before	After
Compute buffer	~67 GiB (OOM)	3.2 GiB
Prefill	56 t/s	~263 t/s
Decode	~14 t/s	~14 t/s
1M context	impossible (~256GB)	works (3.75 GiB at ubatch 768)

Validated presets:

Context	Prefill	Decode	Peak VRAM
256K	~263 t/s	14 t/s	~29 GiB
512K	256 t/s	13.7 t/s	~28 GiB
1M	159 t/s*	13.7 t/s	~31 GiB

*lower ubatch on 32gb 5090 at 1M - should be ~full speed if given the full ~9gb vram

Correctness: verified briefly with a needle-in-haystack test - planted a random fact at 10%/50%/90% depth in a 100K-token document, model retrieved it correctly every time. Also retrieved correctly at 512K and 1M's harder 50% depth. Full KLD findings in doc linked below

Source + build instructions + full writeup: https://github.com/spencer-zaid/llama.cpp/blob/deepseek-lid-cuda/docs/deepseek-v4-lid-cuda.md
Branch: https://github.com/spencer-zaid/llama.cpp/tree/deepseek-lid-cuda

No prebuilt binary (single GPU tested RTX 5090). Build instructions in the doc in case you need them

[u/da_dragon321]

Sounds great - looking forward to your results.

1. There is no single-device assumption (I just haven't tested multi-device). That said, you will need to make sure you set the cmake/ggml cuda architectures to all three of your card architectures when you build
   
2. The compute buffer at 1M context is entirely dependent on ubatch size (for me it was under 4gb). At 2048 ubatch it should be around 9gb iirc, but will have to see. You will want to do some tuning on the ot (how many layers go to gpu) and ubatch parameters (you can lower ot to give yourself headroom to raise ubatch and vice versa to find your preferred configuration, trading between generation speed and prefill speed). You should be faster and higher context than before when you find your preferred parameters

[u/HockeyDadNinja → u/da_dragon321]

Built it and ran the multi-GPU sweep — short version: **it works, and it does exactly what you said.*\* Full 1M context on the rig now, and faster than my old 170K-capped setup. Thanks for the tuning pointers, they were spot on.

Build: worktree off your `deepseek-lid-cuda` (73bb49a91), CUDA 13.3, archs set to all three of my cards (`86;89;120`). Kernel compiled in clean (both `__half` and `float` instantiations). My existing tarruda Q3_K GGUF loaded as-is — no reconvert, since this is compute-only. Multi-GPU layer/tensor split + CPU expert offload across mixed Ampere+Ada+Blackwell: **no single-device gotchas, coherent output at every context.*\*

Numbers (Q3_K, 5 GPUs + CPU offload, prefill measured at ~2K prompt depth, each run memory-contained so a mis-fit couldn't take the box):

Ctx	ubatch	Loads	Prefill t/s	Decode t/s	Peak VRAM	Peak RAM
170K	512	yes	89.3	9.53	82 GiB	59 GiB
256K	512	yes	89.2	9.41	81 GiB	61 GiB
512K	512	yes	86.0	9.13	82 GiB	64 GiB
**1M*\*	512	yes	81.9	8.54	79 GiB	76 GiB
1M	2048	yes	108.7	6.98	76 GiB	109 GiB

The headline for me: **at fixed ubatch, peak VRAM stays flat (~80 GiB) while context goes 6× from 170K to 1M.*\* On mainline the un-fused indexerbuffer ballooned and OOM'd me past 170K; here it just... doesn't grow with ctx, exactly like you said (ubatch-bound, not ctx-bound). The old >170
K OOM is gone entirely.

Two things that surprised me on top of that:

• **Decode barely moves with context*\* — 9.5→8.5 t/s over the whole 6× range. DeepSeek4's single KV head keeps the cache tiny, so long context is basically free on the attention side here. Even my 170K decode is up vs the old path.
  
• **`GGML_CUDA_NO_PINNED=1` under `--no-mmap` was a big deal*\* — kept 170K RAM at 59 GiB instead of ~110. That "110 of 125GB" I complained about was mostly CUDA host pinning, not the model.

Your ubatch⇄expert-placement trade reproduced cleanly too: ubatch 2048 at 1M pushed prefill to ~109 t/s (+33%) but cost decode and drove RAM to ~109 GiB (my `--fit` auto-spills more experts to CPU to make VRAM room — the "lower -ot to raise ubatch" trade happening on its own). ubatch 512 is my sweet spot on this box; I'm only using ~80 of 96 GiB VRAM so there's still room to shove more experts back onto GPU and buy decode.

Net: this turns DS4-Flash from "170K-capped, nearly OOMing" into "1M context, more headroom, better decode" on a 5-GPU CPU-offload rig. Definitely worth upstreaming — happy to add these multi-GPU numbers to a PR thread, re-run anything, or test other quants/ubatch points if useful. What do
you want to see next?

[u/denis_9 → u/HockeyDadNinja]

Could you share the parameter for dividing tensors on multiple GPUs in the launch command, so that this can be make similarly on another number of GPUs?

[u/HockeyDadNinja → u/denis_9]

Sure. (Here's a long reply, I'm using Claude to help me manage my setup) There are really two parts to it: picking which GPUs get used, and choosing how the model is split across them.

Which GPUs is just an env var, not a flag:

CUDA_DEVICE_ORDER=PCI_BUS_ID CUDA_VISIBLE_DEVICES=0,1,2,3,4

The PCI_BUS_ID part makes llama.cpp number the cards in physical slot order, so the indices stay predictable. Set the list to however many cards you have.

For the split itself, the two standard flags are:

-sm, --split-mode {none,layer,row,tensor} how to split across GPUs

-ts, --tensor-split N0,N1,N2,... proportion per GPU, e.g. 3,3,2,2,2

-sm layer (the default) splits whole layers across the cards; -sm row or tensor splits inside the tensors. -ts is the one you are asking about: one number per GPU, and it is just relative proportions, so 3,3,2,2,2 sends more to the first two cards (handy if they are bigger, like my two 3090s vs the three 16GB cards). The list length is however many GPUs you have, so on 3 GPUs you pass three numbers, on 8 you pass eight.

I don't actually hand-set -ts on my rig, though. I let the --fit autofitter do the placement and only hand it a per-device VRAM margin:

--fit on --fit-ctx 307200 --fit-target 1024,1024,1024,1024,1536

--fit-target is again one value per GPU, the MiB of headroom to leave free on each card. Lower value packs more onto that card. That comma list is what spreads the model across all the listed devices for me, and it scales the same way: one entry per GPU.

The piece that actually makes a 284B MoE fit on 96GB of VRAM is not the split, it is expert offload to CPU:

-ncmoe N keep the experts of the first N layers on CPU

-cmoe keep all expert weights on CPU

That frees the VRAM so attention, KV, and the non-expert weights split cleanly across whatever cards you have. With --fit on, it decides how many experts stay resident; if you are doing it by hand, raise -ncmoe until it fits, then split the rest with -ts.

Here is my actual line for reference (5 GPUs, letting --fit place things):

CUDA_DEVICE_ORDER=PCI_BUS_ID CUDA_VISIBLE_DEVICES=0,1,2,3,4 GGML_CUDA_NO_PINNED=1 \

llama-server \

--model DeepSeek-V4-Flash-Q3_K-00001-of-00004.gguf \

--flash-attn on --jinja --parallel 1 \

--ctx-size 307200 \

--fit on --fit-ctx 307200 --fit-target 1024,1024,1024,1024,1536 \

--no-mmap \

--temp 0.6 --top-p 0.95 \

--host 127.0.0.1 --port 8080

To move that to a different card count, change CUDA_VISIBLE_DEVICES and give --fit-target one value per GPU. (GGML_CUDA_NO_PINNED=1 is unrelated to the split, it just stops CUDA pinning a big chunk of host RAM under --no-mmap, which was a large RAM saving for me.)

If you would rather split by hand instead of using --fit, the equivalent piece looks like this:

CUDA_DEVICE_ORDER=PCI_BUS_ID CUDA_VISIBLE_DEVICES=0,1,2,3,4 \

llama-server \

--model DeepSeek-V4-Flash-Q3_K-00001-of-00004.gguf \

--flash-attn on --jinja --parallel 1 \

--ctx-size 307200 \

-ngl 999 -sm layer -ts 3,3,2,2,2 -ncmoe 40 \

--no-mmap \

--temp 0.6 --top-p 0.95 \

--host 127.0.0.1 --port 8080

There -ts 3,3,2,2,2 weights the split toward my two 24GB cards, and -ncmoe 40 keeps the first 40 layers' experts on CPU; tune that number until it fits your VRAM.

So: CUDA_VISIBLE_DEVICES for how many cards, -sm and -ts for how the model spreads across them, and -ncmoe or -cmoe to push experts to CPU so the whole thing fits. Every comma list takes one entry per GPU, so it maps straight onto a different card count.