Qwen3.8-2.4T on --device cuda: the model loads, then the first forward dies on cudaMalloc: out of memory

[localai-bot]

What happened

Qwen3.8-2.4T-A95B UD-Q1_0 (370 GiB, 10 shards) loads to a serving state on --device cuda on a DGX Spark (GB10, 119 GB unified memory) and then dies on the first request:

=== READY after ~1560s ===          (26 min, so the load itself succeeded)
=== measuring (4 tokens) ===
TOTAL 0 tokens in 402.8s
TEXT: ''
== CUDA ==  CUDA Version 13.0.1
engine-fatal: EngineCore busy loop threw: vt cuda: cudaMalloc: out of memory
async-llm: output handler saw engine death: EngineCore encountered an issue [vt cuda: cudaMalloc: out of memory]
sse: stream aborted mid-flight: ...

The failure is in the EngineCore busy loop, not the loader — so weight residency is fine and it is the forward that cannot allocate.

Why this matters

The same checkpoint serves correctly on --device cpu on the same box, same binary, same prompt: TTFT 667.0 s, steady decode 44.2 s/token, coherent output. So the capability exists on one device path and not the other, and the GPU path is where the compute headroom is.

The plausible mechanism, not yet confirmed: the disk-borrow residency that makes this fit is a CPU-path capability. HostExpertSlotStore is host memory by construction, and the expert-streaming lane is CPU-gated. On the CUDA path the forward presumably needs expert slices (and activations) in device allocations, and cudaMalloc draws from the same unified 119 GB pool that already holds the dense remainder plus page cache. 2790 slices/token at 2.38 MB is ~6.5 GB per token of expert bytes that would have to be staged. Someone should confirm which allocation actually fails and its size rather than trust that reasoning — the log line does not name it.

Related: keep-quant device slice needs ResidentWeight is a known constraint in this area.

Reproduction

Under an rc lease on dgx:gpu0 (rc hold dgx:gpu0 --ttl 3h, then ssh — the checkpoint is on the host/NAS and rc run gives the container only):

docker run -d --name q1cuda --gpus all \
  -v /home/mudler/iq1bench:/src -v /home/mudler/ckpt:/ckpt:ro -w /src \
  -e VT_GGUF_PREFAULT=0 \
  vllmcpp-build:gb10 ./build-srv/examples/vllm-server \
  --model /ckpt/qwen3.8-q1_0/Qwen3.8-2.4T-A95B-UD-Q1_0-00001-of-00010.gguf \
  --device cuda --max-num-seqs 1 --max-model-len 512 --port 8899

Checkpoint: unsloth/Qwen3.8-2.4T-A95B-GGUF at revision 567d3e6ac26c5474b18311e619c04350fb9a5556, also staged at /usr/local/nas_share/checkpoints/qwen3.8-2.4t-a95b-ud-q1_0.

Readiness must be probed with python3, not curl — that image has no curl, and a suppressed curl exit 127 reads as "not ready" forever.

What would close this

Either the forward allocates within the pool on a larger-than-memory GGUF, or the engine refuses --device cuda for this configuration by name at load time with a message saying why. Loading for 26 minutes and then dying mid-stream is the worst of the three outcomes, and per policy an unimplemented arm should refuse with a message naming the missing part rather than fail late.

Not a blocker for the CPU result

The measured CPU numbers stand and are unaffected. Recording this so the --device cuda gap is visible debt rather than something rediscovered by the next person who tries the obvious thing.

[localai-bot]

Phase 1: the allocation, named and sized

The log line could not name it. CudaBackend::Alloc is Check(cudaMalloc(&p, bytes), "cudaMalloc") (src/vt/cuda/cuda_backend.cu:77-81) and Check composes "vt cuda: " + what + ": " + cudaGetErrorString(err) (:48-52), where what is the compile-time literal. bytes is in scope and discarded.

The allocation is d.b.Alloc(nb) in ResidentWeight, src/vllm/model_executor/models/qwen3_5.cpp:1010-1011, with nb = w.bytes.size() — for a routed-expert weight that is the whole STACKED [E*N,K] keep-quant tower: every expert of one matrix of one layer, in one contiguous cudaMalloc.

Both switch positions reach that same line, which is why no knob avoids it:

Configuration	Path
default (grouping on)	MoeBlock:6615,6616,6620 → KqGrouped:5694 → ResidentWeight
VT_MOE_EXPERT_STREAM=1 (disables grouping, :5670-5676)	ExpertMlpKq:5638 → MatmulF32Slice:5611 → KqExpertSlice:5595 → KqResidentSlice:5112 → ResidentWeight

KqExpertSlice's slot arm is guarded by is_cpu() (:5578), so on a device platform it falls through before the store is even constructed. Ruled out by reading: cuda_moe.cu and cuda_glue.cu contain no allocation at all, and BuildMoeMarlinResident (:6010-6062) is not on this path — MoeBlock:6555 needs expert_*_fp4, and a GGUF populates expert_*_kq.

The size, from the checkpoint

Both GGUF tensor tables re-censused at revision 567d3e6ac26c5474b18311e619c04350fb9a5556 over all ten shards by HTTP range request, no tensor data downloaded: 1702 records parsed against the 1702 declared in split.tensors.count.

Tower	Bytes		Count
IQ1_XXXS ffn_{gate,up,down}_exps	1,275,068,416	1.1875 GiB	276
Q2_K, block 92 (the nextn MTP block)	2,818,572,288	2.6250 GiB	3
all *_exps	360,374,599,680	335.62 GiB	279

Whole file 397,245,341,184 bytes = 369.96 GiB; expert_count=512, embedding_length=8192, expert_feed_forward_length=2048, block_count=93. Cross-check: 1,275,068,416 / 512 = 2,490,368, exactly the slot_bytes=2490368 the W4 banner printed.

The budget, and the instrument

nvidia-smi --query-gpu=memory.total,memory.free,memory.used answers [N/A], [N/A], [N/A] on this box, and the rc fleet label records vram=[N/A]M. cudaMemGetInfo answers honestly. Measured on dgx:gpu0 under an rc hold, through libcudart.so.13 in vllmcpp-build:gb10:

cudaGetDeviceCount rc = 0 count = 1
cudaMemGetInfo     rc = 0
free  = 122059919360 (113.677 GiB)
total = 128452956160 (119.631 GiB)
attr Integrated           rc=0 value=1
attr UnifiedAddressing    rc=0 value=1

total is EXACTLY /proc/meminfo MemTotal (125442340 kB) times 1024, and equals the fleet label's own mem_total_bytes=128452956160.

Why the loader survives and the forward does not

335.62 GiB of tower staging is 2.8x the whole machine. The load survives because a borrowed tower costs ZERO anonymous bytes; staging converts each borrow into a real allocation, so with ~62 GiB of dense remainder already resident the pool is exhausted after roughly (119.631 - 62) / 1.1875 = 48 towers — partway through layer 16 of 93, on the first forward. Hence 26 minutes to READY, then death mid-stream.

So the hypothesis in the issue body has the mechanism right and the quantity wrong: it is not ~6.5 GB of per-token staging that fails, it is the whole 335.62 GiB tower set being made device-resident once. Per-token slicing never gets a chance to run.

Phase 2 and Phase 3

Phase 2, the load-time refusal, is implemented and opening as a pull request. Phase 3, the device-side slot store, is scoped and filed as #1124 rather than started here; #1126 covers the missing CudaBackend::DeviceMemoryInfo override that this measurement found, and #1127 moves the budget knob into ENG-RESIDENCY-CONFIG's config surface once #1119 lands.