Repository navigation
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
Conversation
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>
|
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 Two smaller things: two defaults change on CUDA builds (the shared expert is no longer offered to the card unless |
5f0a73a to
f7db740
Compare
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>
|
Thank you for running the segment adapters, that gate is exactly the one the fake could not see.
#1769 is rebased onto the same dev as well. |
f7db740 to
d992beb
Compare
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>
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>
5512cf0 to
aaefaf9
Compare
|
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 ( Reading That changes the question at the end of the description, so let me ask it concretely:
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; |
|
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:
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. |
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_stepruns 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 Sover the updated state, the gated RMSNorm withsilu(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 asdeltanet(), 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 withset/get_state;get_stateworks 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 > 1keeps the batched CPU path) pulls first, a--pinsnapshot pulls all layers,reset_recurrentand a pin restore invalidate the device copy, a failing GPU step pulls what the card holds and continues on the CPU. Underautoa layer'sdnoutnow goes to the device itsdnprojwent to, so every layer qualifies.Q36_DN_GPU=0Q36_DN_GPU=1With 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
eidlives on deviceeid % n_gpusin every layer, so each layer's group is issued to both cards andqt_takewaits for the slower one:take5.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=layerhomes 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_PLACEalso gainedexperts=allfor 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 (take0.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; underautothat 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=1to offer it)On the card
shexpis 120 synchronous small GEMVs per token that sit betweenqt_issueandqt_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 onceshexpleft the card. #1657 saw the +3.9 ms on one card already; with a second card it doubles. So the engine offersshexponly 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=offare untouched.Q36_DN_GPU=1, 200 tokens)Q36_OFFER_SHEXP=1auto, a layer'sdnoutgoes to the device with the most roomdnprojwent to, when that device has the roomauto: 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_PLACEnamingshexpis still obeyed when the offer is made;COLI_PLACE=offstill 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.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
downin 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 throughdeltanet()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.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).auto).--pinsnapshots and prefix reuse underQ36_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.hsays 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.