Skip to content

qwen36: the DeltaNet layer runs on the card (Q36_DN_GPU=1); expert homes per layer range (QT_HOME=layer); shexp offered on request - #1784

Merged
JustVugg merged 18 commits into
JustVugg:devfrom
kreuzzelg:qwen36-kda-gpu
Oct 3, 2026
Merged

JustVugg merged 18 commits into
JustVugg:devfrom
kreuzzelg:qwen36-kda-gpu

Conversation

@kreuzzelg

@kreuzzelg kreuzzelg commented Sep 28, 2026 •

Copy link
Copy Markdown
Contributor

Three changes that the same measurement asked for, in the order they were found. Reading why Ollama 0.34.4 decodes this model at 36 tok/s against our 25 (both cards, same box): llama.cpp keeps the whole gated-delta layer on the card and splits the model by layers. Our trunk-in-VRAM token spent 8 of 39 ms on DeltaNet host round trips, our two-card form joined every layer across the cards, and -- found on the way -- the shared expert on the card cost more than on the CPU.

1. The DeltaNet layer on the card (Q36_DN_GPU=1, opt-in)

Backend (coli_cuda_dn_*, five entry points, optional in the DLL loader): one object per DeltaNet layer holds the short-conv ring, the recurrent state [vh x kdim x vdim], the conv and norm weights and one token's scratch in VRAM. coli_cuda_dn_step runs a decode token end to end: x up, in_proj GEMV, causal conv + SiLU with the ring advancing, per value head the L2 norms (double, eps inside the sqrt as on the CPU), the decay, the delta against the decayed state, the write-back, out = q^T S over the updated state, the gated RMSNorm with silu(z), out_proj GEMV, out down. The gates (a/b, 32 outputs) stay on the CPU and travel as kernel parameters. One upload, one download of 8 KB per layer; the state never crosses the bus. Same arithmetic as deltanet(), so the CPU path stays the reference; the fake backend carries a host-side copy of the steps and can be told to fail.

Tier (qt_dn_gpu_*): where the placer put a layer's in_proj and out_proj on one device, the layer joins them there. The host copy of the state stays the engine's, moved with set/get_state; get_state works after a failed step, so the engine continues from what the card holds.

Engine: a decode token computes the two gates and hands x to qt_dn_gpu_step. The host arrays remain canonical, with two flags per layer: the device copy is fresh (nothing to upload before a GPU step) or ahead (download before any CPU use). A prompt (S > 1 keeps the batched CPU path) pulls first, a --pin snapshot pulls all layers, reset_recurrent and a pin restore invalidate the device copy, a failing GPU step pulls what the card holds and continues on the CPU. Under auto a layer's dnout now goes to the device its dnproj went to, so every layer qualifies.

RTX 3070 alone, trunk in VRAM, 200 tokens Q36_DN_GPU=0 Q36_DN_GPU=1
DeltaNet ms/token 16.2 10.8
token ms 39.4 / 39.3 33.9 / 33.3
decode tok/s (1000 / token ms) 25.4 30.0
tok/s as the engine prints it (TTFT included, 200 tokens) 21.2 / 21.3 24.1 / 24.4

With the shared expert on the CPU (part 3) the same card reaches 31.8-32.1 ms, 31.1-31.4 tok/s decode (25.6 on the engine line), from 25.4 before this branch: +23 %.

The 10.8 ms that remain are the two GEMVs (16 and 8 MB per layer) and the two copies. Two GPU runs are byte-identical to each other; against the CPU path the GEMVs sum in a different order, as for every placed component.

2. Two cards as a pipeline (QT_HOME=layer, opt-in)

By default expert eid lives on device eid % n_gpus in every layer, so each layer's group is issued to both cards and qt_take waits for the slower one: take 5.2 ms/token on this 3070 + Quadro RTX 4000 pair against 0.2 on one card, and two cards at 99 % residency no faster than the 3070 alone. QT_HOME=layer homes every expert of a layer on one device, chosen per layer range by allowance divided by the card's dense-GEMV time (a 64 MB int8 probe at startup; the log names how much slower the slow card is per dense byte; QT_LAYER_SPLIT=<n> overrides), puts the layer's trunk components on the same device and lm_head on the last layer's, and issues a group to exactly one device -- llama.cpp's layer split. Budgets, warmstart, LFRU and the swap path follow the homes (home2(layer, eid)); a hand-written list's reservation does not remove a card in this mode. One card: unchanged. COLI_PLACE also gained experts=all for the index mode: the experts stay on every card even when a list reserves one for the trunk.

Measured (two cards, Q36_DN_GPU=1, shared expert on the CPU): the join is gone (take 0.17 ms), the MoE phase falls from 12.4-13.4 to 9.4 ms at 99.7 % residency, and the pipeline ends level with the 3070 alone -- 33.6 ms = 29.8 tok/s at a hand-written 28/12 split against 31.8-33.1 ms = 30.2-31.4 tok/s -- because lm_head (+2.8 ms) and twelve layers (+1.6 ms) now run on the slower card. The probe at startup measures that card 154 % slower per dense byte on a 64 MB GEMV and picks 28/12 by itself; under auto that lands at 36.0 ms = 27.8 tok/s (hit rate 97 %, more trunk than the hand list put on the fast card), 12 % behind the 3070 alone. Two equal cards would keep the MoE gain without that price; cold, two cards win either way (45.3 against 52.3 ms; 22.1 against 19.1 tok/s). A split by allowance alone was the wrong one on unequal cards (19/21 gave 38.5 ms = 26.0 tok/s, 28/12 gave 33.6 ms = 29.8 tok/s), which is why the default weighs by the probed GEMV time. Ollama 0.34.4 on the same pair has the same shape, for the record: 38.1 tok/s on the 3070 alone, 36.6-37.5 on both cards, 33.0 on the Quadro alone -- on this pair the second card does not pay for anyone.

3. The shared expert stays on the CPU (Q36_OFFER_SHEXP=1 to offer it)

On the card shexp is 120 synchronous small GEMVs per token that sit between qt_issue and qt_take, where on the CPU the shared expert hides behind the expert group. Measured with the DeltaNet layer on the card: shared 3.7-4.0 ms/token on the CPU against 6.3 on the same 3070 and 11-12.7 ms with layers on the slower card; the token went 33.3 -> 31.8 ms (30.0 -> 31.4 tok/s) on one card and 38.6 -> 33.6 ms (25.9 -> 29.8 tok/s) as a pipeline once shexp left the card. #1657 saw the +3.9 ms on one card already; with a second card it doubles. So the engine offers shexp only on request; a hand-written list naming it is still obeyed.

Two defaults change on CUDA builds

Both are opt-out/opt-in switches the measurement asked for; the CPU path and COLI_PLACE=off are untouched.

default before after measured (RTX 3070, 35B int4, Q36_DN_GPU=1, 200 tokens)
the shared expert is offered to the placer yes (since #1657) only with Q36_OFFER_SHEXP=1 one card 33.3 -> 31.8 ms/token (30.0 -> 31.4 tok/s), shared 6.3 -> 3.7-4.0 ms; two-card pipeline 38.6 -> 33.6 ms (25.9 -> 29.8 tok/s), shared 11.0 -> 3.9 ms. Without the DeltaNet layer on the card the one-card half of this was already in #1657 (+3.9 ms/token for shexp on the card).
under auto, a layer's dnout goes to the device with the most room yes to the device its dnproj went to, when that device has the room two cards, auto: 28 of 30 DeltaNet layers qualified for the card -> 30 of 30; 37.7 -> 37.4 ms/token (26.5 -> 26.7 tok/s). On one card nothing changes (one device).

A hand-written COLI_PLACE naming shexp is still obeyed when the offer is made; COLI_PLACE=off still places nothing.

Against Ollama on the same box (for scale, not as a target)

Same prompt (49 tokens, ~200 generated), same day, Ollama 0.34.4 with qwen3.6:35b (Q4_K_M, 23 GB, --verbose --think=false), one cold run and two warm ones per configuration. Colibri from this branch, per-row int4 container, Q36_DN_GPU=1, shared expert on the CPU.

Ollama 0.34.4 (Q4_K_M) Colibri (this branch, int4)
decode tok/s, 3070 alone, warm 38.1 31.1-31.4
decode tok/s, both cards, warm 36.6-37.5 27.8 (pipeline, auto) / 29.8 (hand list)
decode tok/s, Quadro RTX 4000 alone, warm 33.0 not measured
first token, warm (prompt eval / TTFT) 61 tokens at 38 tok/s = 1.6 s 1.5 s
cold start: time to load + first token 39 s (page cache warm) to 2 min 24 s (cold) 9 s load + 1.8 s TTFT
cold decode, first 200 tokens 9-16 tok/s (mmap paging behind the generation) 19.1 tok/s (3070) / 22.1 (two cards)
VRAM 7.1 GB (one card) / 4.5 + 7.2 GB 6.5 GB (one card)
RAM 23 GB model, mmap ~40 GB (every expert resident in RAM)

Where this engine is ahead: the cold start and everything before the heat file is warm -- 9 s to load the container and 1.8 s to the first token against Ollama's 39 s to 2.4 min, and 19-22 tok/s cold against 9-16 while Ollama pages the experts in; the two-card cold number is where the pipeline pays. Where Ollama is ahead: warm decode by 21 % on one card, at a smaller RAM footprint, and its layer split loses nothing when a second card is present (38.1 -> 36.6, ours 31.4 -> 27.8-29.8). Both engines are faster on the 3070 alone than on the pair; the Quadro costs both of them. Two caveats on fairness: Q4_K_M keeps down in Q6_K and measures 2.5 % better perplexity than our per-row int4 (#1370), and Ollama's cold run had a cold page cache twice while ours reads 23 GB of experts into RAM at startup by design.

Verified

  • tests/test_qwen36_dn_gpu.c: one in-memory DeltaNet layer through deltanet() on the fake backend -- GPU outputs equal the CPU outputs; the host copy is left alone while the card advances and reads back the CPU path's state; a two-row CPU prefill after GPU tokens pulls first, the GPU token after it pushes first; a failing step turns the layer off and the CPU continues from the card's state.
  • tests/test_qwen36_tier_layer_home.c: layer homes, trunk on the layer's device, lm_head on the last, budgets and residents per card, a group to exactly one device, QT_LAYER_SPLIT, one card unchanged.
  • tests/test_qwen36_tier_autoplace.c: dnout follows dnproj; experts=all.
  • Existing qwen36 tier tests unchanged and green; CPU-only build carries the stubs; colibri (the loader host) builds; the loader ABI tripwire counts the five optional symbols.
  • tests/test_segment_adapters_real qwen36 <container> 0 4 8 8, plain and under ASan+UBSan: green after the NULLs (it reproduced the use-after-free before).
  • The weighted default split measured on the pair (probe 0.50 against 1.28 ms per 64 MB GEMV, 28/12 chosen, 36.0 ms = 27.8 tok/s under auto).
  • Not exercised on the real container: --pin snapshots and prefix reuse under Q36_DN_GPU=1 (the code paths only read host arrays, pulled first; the fake test covers the order). I would run a serve smoke before merging if you want it.

Where this goes next

c/delta_attention.h says it: the same gated delta step runs the linear-attention layers of Kimi K3 (KDA), GLM-5.3-Flash and Qwen3.8; only the dimensions and the gate differ. The backend object takes those dimensions as parameters (vh <= 64, kdim/vdim <= 256), so the other engines can hand their layers to the same five calls. Kept to qwen36 here so the numbers are the engine's own. Against Ollama 0.34.4 on the 3070 alone (38.1 tok/s decode, Q4_K_M) the gap goes from 50 % to 21 %: 25.4 -> 31.4 tok/s decode on this box.

Thanks to Claude Code for the help along the way.

kreuzzelg added a commit to kreuzzelg/colibri that referenced this pull request Sep 28, 2026
hipcc has no __shfl_down_sync (the file's f8_shfl_down macro already maps
it) and backend_gpu_compat.h mapped cudaMemsetAsync but not cudaMemset.
Found by the HIP syntax check on JustVugg#1784.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
@JustVugg

Copy link
Copy Markdown
Owner

Thanks. On the CPU path this is clean: every qwen36 tiny oracle is token-exact with logits byte-identical to dev (MoE caps 1/2/8, 2.4T geometry, the 27B dense and the 27B vision fixture), and all 31 qwen36 C tests pass.

There is one use-after-free that also hits CPU-only builds. At c/qwen36.c:4639 dn_dev_fresh and dn_host_stale are freed, but line 4640 only resets DN_rec and DN_conv. The next segment run writes into the freed array (line 3038), and teardown frees both again (line 4577). Under ASan tests/test_segment_adapters_real qwen36 ... reports heap-use-after-free; the plain build segfaults 3 times out of 3, dev passes. CI misses it because segment-adapters-real is a manual gate. Setting both pointers to NULL next to the other two fixes it.

Two smaller things: two defaults change on CUDA builds (the shared expert is no longer offered to the card unless Q36_OFFER_SHEXP=1, and under automatic placement a layer's dnout follows its dnproj). Please put that in the description with a before/after number. And one log line is in German. With the NULLs it can go in.

kreuzzelg added a commit to kreuzzelg/colibri that referenced this pull request Oct 1, 2026
hipcc has no __shfl_down_sync (the file's f8_shfl_down macro already maps
it) and backend_gpu_compat.h mapped cudaMemsetAsync but not cudaMemset.
Found by the HIP syntax check on JustVugg#1784.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
kreuzzelg added a commit to kreuzzelg/colibri that referenced this pull request Oct 1, 2026
…the state; one log line in English

JustVugg's review of JustVugg#1784: the segment engine freed dn_dev_fresh and
dn_host_stale next to DN_rec/DN_conv but reset only the latter two, so the
next segment run wrote into freed memory (deltanet marks the device copy
old) and teardown freed both again -- heap-use-after-free under ASan in
tests/test_segment_adapters_real, a segfault in the plain build, CPU-only
builds included. Both pointers are NULL now, like the other two.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
@kreuzzelg

Copy link
Copy Markdown
Contributor Author

Thank you for running the segment adapters, that gate is exactly the one the fake could not see.

  • The two pointers are NULL now next to DN_rec/DN_conv (502c81e1). tests/test_segment_adapters_real qwen36 <container> 0 4 8 8 is green here in the plain build and under ASan+UBSan; before the fix it reproduced your heap-use-after-free.
  • The German log line is English.
  • The two default changes have their own section in the description now, with before/after numbers for each (shexp offered on request: 33.3 -> 31.8 ms/token on one card, 38.6 -> 33.6 as a pipeline; dnout following dnproj under auto: 28 -> 30 layers on the card, 37.7 -> 37.4 ms on two cards).
  • Rebased onto dev 4521832 (the Makefile's new dependency style for the two test rules; otherwise clean), 16 commits, CI running.

#1769 is rebased onto the same dev as well.

kreuzzelg added a commit to kreuzzelg/colibri that referenced this pull request Oct 2, 2026
hipcc has no __shfl_down_sync (the file's f8_shfl_down macro already maps
it) and backend_gpu_compat.h mapped cudaMemsetAsync but not cudaMemset.
Found by the HIP syntax check on JustVugg#1784.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
kreuzzelg added a commit to kreuzzelg/colibri that referenced this pull request Oct 2, 2026
…the state; one log line in English

JustVugg's review of JustVugg#1784: the segment engine freed dn_dev_fresh and
dn_host_stale next to DN_rec/DN_conv but reset only the latter two, so the
next segment run wrote into freed memory (deltanet marks the device copy
old) and teardown freed both again -- heap-use-after-free under ASan in
tests/test_segment_adapters_real, a segfault in the plain build, CPU-only
builds included. Both pointers are NULL now, like the other two.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
kreuzzelg and others added 17 commits October 3, 2026 10:53
One object per DeltaNet layer holds the short-conv ring, the recurrent
state [vh x kdim x vdim], the conv and norm weights and one token's scratch
in VRAM. coli_cuda_dn_step runs a decode token end to end on the card:
x (host) -> in_proj GEMV -> causal conv + SiLU -> per value head the L2
norms (double, eps inside the sqrt), the decay, delta against the decayed
state, the write-back, out = q^T S over the updated state, the gated
RMSNorm with silu(z) -> out_proj GEMV -> out (host). Gates arrive by value.
Only x and out cross the bus; the state never does. Same arithmetic as
qwen36.c's deltanet(), so the CPU path stays the reference.

Optional in the DLL loader (five RESOLVE_OPT); the fake backend carries a
host-side reference of the same steps so an engine test runs without a
card, and can be told to fail.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…s on the card

Where the in_proj (qt_dnproj_init) and the out_proj (a qt_dense handle)
of a DeltaNet layer sit on the same device, qt_dn_gpu_init puts the conv
ring, the recurrent state and the norm/conv weights there too, and
qt_dn_gpu_step runs a decode token through the backend chain. The host
copy of the state stays the engine's, moved with qt_dn_gpu_set/get_state
when the CPU path needs it; get_state works after a failed step as well,
so the engine can continue from what the card holds. A failing step turns
the layer off like every other GPU path here.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
… end to end

With the trunk in VRAM the CPU still spent ~8 of 39 ms/token on the
convolution, the L2 norms, the recurrence and the gated norm of the thirty
DeltaNet layers -- microseconds of arithmetic, paid as host round trips:
the in_proj result came down, the recurrence ran on the CPU, its output
went back up for the out_proj. Now a decode token computes the two gates
on the CPU (two tiny matmuls) and hands x to qt_dn_gpu_step, which returns
the layer's output; one upload, one download per layer.

The host arrays stay canonical. Per layer, dn_dev_fresh says the card
holds the host's state (nothing to upload before a GPU step) and
dn_host_stale says the card advanced past the host copy. A CPU step (a
prefill, S > 1) pulls first and marks the device copy old; a snapshot
pulls all layers; reset_recurrent and a pin restore invalidate the device
copy. A failing GPU step pulls what the card holds and continues on the
CPU. Opt-in for now: Q36_DN_GPU=1, and only layers whose in_proj and
out_proj the placer put on one card.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…inst the CPU path

One in-memory DeltaNet layer, no container, driven through deltanet() on
the fake backend: decode tokens through the GPU path give the CPU outputs;
the host copy is left alone while the card advances and reads back the CPU
path's state; a two-row CPU prefill after GPU tokens pulls first and
continues, the GPU token after it pushes first; a failing step turns the
layer off and the CPU continues from the card's state.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…experts=all

Two placement details the GPU DeltaNet layer needs. The automatic placer
chose the device for each offer by room alone, so a layer's in_proj and
out_proj could land on different cards (28 of 30 layers qualified on two
cards): now dnout goes to the device its dnproj went to when that device
has the room. And a hand-written list that reserves a card for the trunk
hands that card's experts to the others (JustVugg#1361), which rules out "trunk on
the fast card, experts on both"; experts=all keeps every COLI_GPUS card
for the experts, reserved or not. Both in test_qwen36_tier_autoplace.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…counts model layers

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
… as a pipeline

By default expert eid lives on device eid % ndev in every layer, so every
layer's group joins across the cards and the slower card paces each layer:
measured on a 3070 + Quadro RTX 4000 pair, take() 5.2 ms/token against 0.2
on one card, and two cards at 99 % residency no faster than the 3070 alone.
llama.cpp splits the other way, by layers, and the token crosses the cards
once.

QT_HOME=layer does that here: every expert of layer l is homed on one
device, chosen per layer range in proportion to the cards' allowances (or
QT_LAYER_SPLIT=<n> layers on the first card), the automatic placer puts the
layer's trunk components on the same device and lm_head on the last
layer's, and a group is issued to exactly one device. Budgets, warmstart,
LFRU and the swap path go through home2(layer, eid); the hand-written
list's reservation rule does not remove a card in this mode, its layers'
experts live there. One card: unchanged.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…split

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…FER_SHEXP=1)

On the card shexp is 120 synchronous small GEMVs per token that sit
between qt_issue and qt_take, where on the CPU the shared expert hides
behind the expert group. Measured on the 35B with the DeltaNet layer on
the card: shared 3.7-4.0 ms/token on the CPU against 6.3 on the same 3070
and 11-12.7 with layers on a slower second card; the token went 33.3 ->
31.8 ms on one card and 38.6 -> 33.6 as a two-card pipeline once shexp
left the card (JustVugg#1657 saw the one-card half of this). A hand-written
COLI_PLACE naming shexp is obeyed when the offer is made; the trunk_dense
test asks for the offer explicitly.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
… expert on the CPU

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
… speed

The allowance alone put 21 of 40 layers on the slower card of this pair
(38.5 ms/token against 33.6 at a 28/12 split). At startup the tier now
times a 64 MB int8 GEMV on every card (ten calls, best of three rounds)
and weighs a card's share of the layers by allowance / time; the log says
how much slower the slow card is per dense byte, which is what every layer
it holds pays. A backend that answers in no time (the fake) leaves the
weights at the allowances, so the tests stay deterministic; QT_LAYER_SPLIT
still overrides.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…mber for the record

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
hipcc has no __shfl_down_sync (the file's f8_shfl_down macro already maps
it) and backend_gpu_compat.h mapped cudaMemsetAsync but not cudaMemset.
Found by the HIP syntax check on JustVugg#1784.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…the state; one log line in English

JustVugg's review of JustVugg#1784: the segment engine freed dn_dev_fresh and
dn_host_stale next to DN_rec/DN_conv but reset only the latter two, so the
next segment run wrote into freed memory (deltanet marks the device copy
old) and teardown freed both again -- heap-use-after-free under ASan in
tests/test_segment_adapters_real, a segfault in the plain build, CPU-only
builds included. Both pointers are NULL now, like the other two.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…er link that includes qwen36.c

dev's test_makefile_vk_obj asks for it: under VK=1 qwen36.c calls coli_vk_*.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
@kreuzzelg

Copy link
Copy Markdown
Contributor Author

Rebased a third time, now onto dev fe744f4 (across the dense chain, #1849/#1851): the two places where both sides mark their device copy of the DeltaNet state old (reset_recurrent, the pin restore) now carry your q36c_host_wrote line and my dn_gpu_invalidate side by side; otherwise clean. Tests, the Makefile VK_OBJ tripwire and test_segment_adapters_real qwen36 are green here, CUDA build fine. #1769 is on the same dev.

Reading vk_chain.c after the fact: it is the same idea on the other backend, with the same contract -- the host's copy stays canonical, pulled before anything reads it there, pushed after anything writes it -- only broader (the whole layer as one frame, the residual on the device). Since the chain declines under the CUDA tier, the two never meet at run time, and a layer-homed two-card split (QT_HOME=layer) is what a frame per layer would want on two cards anyway.

That changes the question at the end of the description, so let me ask it concretely:

  1. Do you want CUDA to grow a chain of its own over time, mirroring vk_chain.c (whole frame per layer, residual on the device)? Then coli_cuda_dn_step is its first building block and I would rather shape it the way you want the rest to look -- say where.
  2. Or should this stay a per-layer step, with qwen38 as the one follow-up (same DeltaNet, 48 value heads, the CUDA trunk from qwen38: CUDA VRAM tier for Qwen3.8-Flash-Next -- hot experts, and the dense trunk as int8 residents (+43 % on one 8 GB card, +60 % on two) #1424 already there; about a day)? I would not take it to kimi_k3 or glm53 on CUDA: neither has a CUDA dense placement, and your direction for both is the shared Vulkan tier.

One datapoint from trying to measure your chain against this on the same card, since the docs say no discrete GPU has been: on an RTX 3070 with the launch VBIOS (94.04.3A, no Resizable BAR; resource1_resize = 0x1c0, so no BIOS switch) the Vulkan backend comes up, warns "only 246 of 8192 MB VRAM is host-visible", and then places nothing -- 0 matrices resident, the tier failed 3004, the chain "a matrix did not reach the device" -- the NVIDIA driver refuses the allocations past the window rather than spilling to system RAM as RADV did in #523. That is the gap #1338 described (staged uploads into device-local memory); every Ampere card still on its launch VBIOS and every Turing card is in that situation. Happy to open it as its own issue with the log if that is useful.

@JustVugg

JustVugg commented Oct 3, 2026 •

Copy link
Copy Markdown
Owner

Thank you. This goes in as soon as CI is green on the updated branch: I pressed update-branch, so it carries the current dev.

On your question: CUDA keeps going alongside Vulkan, not instead of it. The order we would like:

  1. First keep this a per-layer step and take it to qwen38: the same DeltaNet, 48 value heads, and the CUDA trunk from qwen38: CUDA VRAM tier for Qwen3.8-Flash-Next -- hot experts, and the dense trunk as int8 residents (+43 % on one 8 GB card, +60 % on two) #1424 already there. It is a measured win for about a day of work.
  2. Then a CUDA chain that mirrors vk_chain.c: a whole layer per frame, with the residual on the device. Same contract: the host copy stays canonical, is pulled before any host read and pushed after any host write, and a lost device is rebuilt from the prefix record. coli_cuda_dn_step would be its DeltaNet op. When you get there, please shape the ops like vk_chain.h's (geometry given by the caller, one recorder, the engine owning its state), so an engine can wire both chains the same way.

kimi_k3 and glm53 stay on the Vulkan path, as you say.

The Resizable BAR finding is a real gap, thanks for measuring it. Staged device-local uploads, the approach of #1338 by @crichalchemist, are being written for the current backend: resident weights, the shared tier's expert pool and the chain's buffers. When that PR is up, could you run it on the 3070 with the launch VBIOS? That is the one configuration we cannot test here.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants