Skip to content

qwen36: every layer as one chain on the CUDA device (COLI_CUDA_CHAIN=1), shaped as vk_chain.h -- 30 -> 18.7 ms a token on an RTX 3070 - #1992

Merged
JustVugg merged 10 commits into
JustVugg:devfrom
kreuzzelg:qwen36-cuda-chain
Oct 12, 2026
Merged

JustVugg merged 10 commits into
JustVugg:devfrom
kreuzzelg:qwen36-cuda-chain

Conversation

@kreuzzelg

Copy link
Copy Markdown
Contributor

The CUDA chain you asked for in #1784, after qwen38's per-layer step (#1870): a whole layer per frame with the residual on the device, the host copy canonical, coli_cuda_dn_step's kernels as the DeltaNet ops, and the ops shaped like vk_chain.h's so an engine wires both chains the same way.

What it is

c/cuda_chain.h / c/cuda_chain.cu -- the backend's twin of vk_chain.c. Buffers CC_DEV, CC_UP, CC_DOWN (the host-side ones page-locked and mapped: the device reads or writes them in place, the host through cc_ptr, as your host-visible buffers); frames cc_begin / cc_submit(wait) on one non-blocking stream per device with an event per frame (the stream is the barrier you record by hand); the ops cc_matmul (the resident tensors the tier already uploads, on the chain's stream), cc_norm, cc_rope, cc_attn, cc_dnconv, cc_dnrec, cc_ew -- each the arithmetic of shaders/chain_*.comp ported line for line, the parameter structs field for field the Vkc ones (CcNorm = VkcNorm, ...). A CUDA error marks the device lost to the chain. The backend exports the ops as one table, coli_cuda_chain_ops, versioned by its size: one optional symbol for the Windows loader (the tripwire counts one more optional export), and a DLL that predates it simply has no chain.

c/qwen36_cuda_chain.h -- qwen36_chain.h's frames and contract on that table. Per layer, frame A1: the previous layer's MoE output joining the residual in the CPU's order, the input norm, the gated attention over a KV mirror on the device (q/k norms, RoPE from the CPU's own cos/sin table, the new rows into the device cache and down to the host's, the output gate, o_proj) or the Gated DeltaNet (the tier's fused qkv ++ z projection, b|a, the convolution with its ring and the recurrence with its state on the device, the gated norm, out_proj), the residual add, the post norm, the router logits; waited for. The host routes and runs the routed experts through the tier (moe_ex routed only; that path computed the shared expert regardless of routed_only under the CUDA tier, fixed). Frame A2, not waited for: the shared expert and its gate on the chain's stream while the tier's groups run on the backend's. The last frame: the final norm and lm_head on the last row, or every row of a verify.

State: the host's KV cache and DeltaNet state stay canonical (watermarks per layer, the state pulled before a host read -- a pinned snapshot, a CPU step -- and pushed after a host write -- a reset as zeros on the device, a restored snapshot as an upload); a speculative verify copies the state after each row into a device slot, the convolution and the recurrence split into one dispatch per copy (the same bits), and a rejected draft swaps the slot in (q36cc_rollback) as the CPU swaps its own; a lost device zeroes the state, rebuilds it on the CPU from the prefix record and runs there from then on. COLI_CUDA_CHAIN_FAULT=n fakes the loss.

What it needs, and when it declines: every projection, the shared expert's three matrices and lm_head on one card (COLI_CUDA_CHAIN=1 offers the shared expert to the placer as Q36_OFFER_SHEXP=1 does; COLI_PLACE=auto takes the whole 35B trunk on an 8 GB card); it names the first matrix on the CPU otherwise, and stays off under a qpack container, PILOT, or KV_SLOTS > 1. While it runs, Q36_DN_GPU's per-layer step is not set up: the chain owns the state. The tier exposes the tensors behind its handles (qt_dense_tensor, qt_dnproj_tensor, qt_lmhead_tensor).

A decode GEMV (gemv_i8_rows, backend_cuda.cu): quant_matmul spends a block of 256 threads on one output row and reads the trunk at about a quarter of the bus (117-125 GB/s on the 3070); a warp per row with 16 bytes a lane and the activation in shared memory reads it at 330-418 GB/s (dnproj 0.215 -> 0.063 ms, lm_head 4.07 -> 1.22 ms). coli_cuda_tensor_gemm_async takes it for int8 rows, up to 8 activation rows each on its own (a verify's rows get a decode step's bits, as q36_spec_step asks; the op test holds three rows at once bit for bit to three one-row calls), I a multiple of 16; everything else keeps the block kernel. Only the chain calls it: coli_cuda_matmul's path is untouched.

Measured (RTX 3070 alone, Qwen3.6-35B-A3B int4, COLI_PLACE=auto, tier at 81-82 % hits, 49-token prompt, 200 new tokens, two runs each)

TTFT tok/s (the run, TTFT included) ms per decode token decode tok/s device wait / token routed experts on the host / token
Q36_DN_GPU=1 (#1784, the per-layer step) 1.22-1.26 s 26.2 / 27.6 30-32 32-33 12.3-13.5
COLI_CUDA_CHAIN=1, the block GEMV 0.98-1.00 s 30.8 / 31.1 27.6 36 14.7 ms 9.9
COLI_CUDA_CHAIN=1, the warp GEMV 0.98 s 42.7 / 42.8 18.7 53 5.7 ms 9.9
Ollama 0.34.4, the same card (28 Sept) 26.2 38.1

The chain alone removed the CPU's parts (the attention core 3.9 ms, the shared expert 4.1 ms, about seventy bus round trips) and left the token as the device's time: 14.7 ms, the trunk's 1.9 GB of int8 rows at 120 GB/s. With the warp GEMV the device's part is 5.7 ms, the token 18.7 ms: -38 % against the per-layer step, and the first time this engine decodes faster than Ollama on this card (+39 %). What is left is the routed experts on the host (9.9 ms) and the routing, launches and forty waits (about 3 ms). The chain is deterministic run to run; its text opens as the per-layer path's and diverges after about a hundred characters, as two runs of the per-layer path diverge from each other.

Opt-in for now, as Q36_DN_GPU was; the default is yours to flip.

Verified

  • tests/test_qwen36_cuda_chain.c (in CI with the engine tests): an in-memory two-layer model (a DeltaNet layer, a gated-attention layer, four routed experts and a gated shared expert per layer, lm_head) through step_ex on the fake tier, the chain's ops from a host-side table in tests/qwen36_fake_cuda.h with the kernels' arithmetic -- 31 checks: the CPU path's logits for a prompt and decode tokens (max diff 2e-5 of 20), the host state untouched while the chain runs and read back as the CPU's, the K/V rows in the host cache, a reset and a CPU step in between, a verify's three rows and the rollback of a rejected draft, a frame that fails (the state rebuilt, the step and the rest on the CPU, the same logits), a backend without the table, KV_SLOTS > 1.
  • tests/test_cuda_chain.cu (make cuda-chain-check CUDA=1; gpu-compile builds it, so the HIP job compiles it too): every op on the 3070 against the references -- buffers and frames, int8 and f32 matmuls at offsets (both GEMV paths), the norms, RoPE, causal GQA with the gate and a selection list, the convolution in both orders with its ring and snapshot, the recurrence at KD 8 and 128 with both gates, every element-wise op, a failed launch marking the device lost. All pass.
  • Builds: CPU-only, VK=1, CUDA=1; test_makefile_vk_obj.py and test_backend_loader.py green; a cuda_chain.cu step in the CUDA syntax check.
  • Not run on the real container: --pin, the prompt cache and prompt-lookup drafts under the chain (the code paths pull or invalidate as the Vulkan chain's; the fake test covers each), and a second card (one device per chain for now; the layer pipeline of QT_HOME=layer would need a chain per card, as your COLI_VK_DEV2).

Thanks to Claude Code for the help along the way.

kreuzzelg and others added 6 commits October 6, 2026 21:42
…s shape

A layer issued on one stream: buffers (CC_DEV, and CC_UP / CC_DOWN as
page-locked host memory the device reads or writes in place), frames
(cc_begin / cc_submit on a non-blocking stream per device, an event per
frame), and the ops cc_matmul (the backend's GEMV / GEMM kernel over the
tier's resident tensors, on the chain's stream), cc_norm, cc_rope, cc_attn,
cc_dnconv, cc_dnrec and cc_ew, each the arithmetic of the shader of the
same name in shaders/chain_*.comp ported line for line, the parameter
structs field for field the Vkc ones. A CUDA error marks the device lost
to the chain. Exported as one table (coli_cuda_chain_ops): one optional
loader symbol, versioned by its size, so a DLL without it has no chain.
backend_cuda.cu gains coli_cuda_tensor_shape and
coli_cuda_tensor_gemm_async for it; the HIP shims the mapped buffers need.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
qwen36_chain.h's frames and state contract on cuda_chain.h: per layer,
frame A1 (the previous MoE output joining the residual, the input norm,
the gated attention over a KV mirror on the device or the DeltaNet with
its state on the device, the residual add, the post norm, the router
logits), the host's routed experts through the tier, frame A2 (the shared
expert, not waited for), the last frame the final norm and lm_head. The
host's KV cache and DeltaNet state stay canonical (watermarks, pull
before a host read, push after a host write, a verify's copies per row,
the rollback of a rejected draft); a lost device rebuilds the state on
the CPU from the prefix record. The tier exposes the tensors behind its
handles (qt_dense_tensor and friends); COLI_CUDA_CHAIN=1 offers the
shared expert to the placer, and moe_ex's routed-only form leaves the
shared expert to the caller under the CUDA tier too.

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

tests/test_qwen36_cuda_chain.c: an in-memory two-layer model through
step_ex on the fake tier, the chain's ops from the fake's host-side table
(the same arithmetic as the kernels): the CPU path's logits for a prompt,
decode tokens, a reset, a CPU step in between, a verify and its rollback;
the state read back as the CPU's; a failed frame rebuilding on the CPU; a
backend without the table. tests/test_cuda_chain.cu: every op on a real
device against the references (make cuda-chain-check; gpu-compile builds
it). cuda_chain.cu in the CUDA and HIP DLL rules and the CI syntax check.
docs: ENVIRONMENT rows, the chain in qwen36-cuda-tier.md and cuda.md.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
… chain's decode accounting

quant_matmul spends a block of 256 threads on one output row and reads the
trunk at about a quarter of the bus (117-125 GB/s on an RTX 3070): on the
chain's decode token that was 14.7 ms of device time for 1.9 GB of rows.
gemv_i8_rows gives a warp to each row, 16 bytes a lane per step, the
activation in shared memory, a shuffle reduction: 330-418 GB/s on the same
shapes (dnproj 0.215 -> 0.063 ms, lm_head 4.07 -> 1.22 ms), 5.7 ms a
token. coli_cuda_tensor_gemm_async takes it for fmt 1, one row, I a
multiple of 16; the block kernel stays for the rest (the op test covers
a width it leaves alone). The chain's report now gives the decode token's
own waits and expert time; the op test compares whole buffers.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
…decode step's bits)

A speculative verify hands the chain 2 to 6 rows; with those on the block
kernel its rows got other bits than a decode step's, which the engines
ask not to happen (q36_spec_step: an accepted draft equals what decode
would have said). gemv_i8_rows now runs up to 8 activation rows as
blockIdx.y, every row the one-row call's sums; the op test holds three
rows at once bit for bit to three one-row calls.

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

JustVugg commented Oct 8, 2026

Copy link
Copy Markdown
Owner

Thanks, this is a big step: 18.7 ms a token, measured on the card, with the chain shaped like the Vulkan one.

The three red checks come from the PR itself, not from flakes. Each fix is small, and I checked them locally (the ASan run, test_makefile_deps.py, test_backend_loader_header_parity.py):

  1. macOS: tests/test_qwen36_cuda_chain exits 1 in section C. step_ex(&ref, ids + 6, 3, 6, 3) reads ids[6..8], but ids has 8 entries, so ids[8] is past the end. ASan reports it as a stack-buffer-overflow at qwen36.c:4254, called from test_qwen36_cuda_chain.c:212. With gcc -O3 on Linux the stray value happens to be a valid token id. On macOS it is not, so the engine refuses it with "token id ... out of range" and exits 1, with no FAIL line. A ninth id fixes it (ASan clean, 31 checks pass):
    const int ids[9] = {3, 17, 42, 9, 21, 60, 5, 33, 11};
  2. Python tests: test_no_generated_rule_lists_headers_by_hand. The test binary's rule lists its headers, which -MMD already tracks:
    tests/test_qwen36_cuda_chain$(EXE): tests/test_qwen36_cuda_chain.c qwen36.c qwen36_tier.c $(VK_OBJ)
  3. Python tests: test_every_header_export_is_resolved_by_the_loader. coli_cuda_tensor_shape and coli_cuda_tensor_gemm_async are declared COLI_CUDA_DLLEXPORT in backend_cuda.h, but the loader never resolves them. Only cuda_chain.cu calls them, inside the same DLL, so they need no export. Declaring them without COLI_CUDA_DLLEXPORT is enough, and the chain stays one optional symbol.
  4. Windows: test_abi_is_derived_from_the_loader_source (16 != 15). This is the tripwire you mention, and it only runs on Windows, so it is skipped on Linux. chain_ops is one more optional export: 15 -> 16 optional, 62 -> 63 exports, plus +chain_ops in the comment next to the count.

A question for after this one: would you take the same chain to Qwen3.8-Flash-Next (qwen38.c)? The pieces are already there:

So it would mostly be a qwen38_cuda_chain.h on your ops table. One difference: Qwen3.8's trunk is heavier than its experts (8 GiB in bf16, which is why it runs in int8), so on an 8 GB card it leaves less room for the tier. We have no NVIDIA card to measure on, and your 3070 numbers are what would tell us whether it pays there too.

kreuzzelg and others added 3 commits October 8, 2026 20:50
…kend calls kept inside the backend; the test's id array

The Windows loader test expects 16 optional exports now (chain_ops).
coli_cuda_tensor_shape and coli_cuda_tensor_gemm_async are cuda_chain.cu's
alone (the same DLL): declared there, not in backend_cuda.h, so the header
parity test does not ask the loader for them. The engine test's rule gets
its headers from -MMD like every other. The test's verify read three ids
from index 6 of an array of eight: under clang on macOS the ninth was
garbage and step_ex refused the token (ASan: stack-buffer-overflow); the
array is twelve long now, and the test runs clean under ASan and UBSan.

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

The cache is malloc'd; past the written positions it holds whatever malloc
gave, zeros on Linux and not on Windows (UCRT64 compared the two models'
garbage). The check covers the four positions of the prompt, per head.

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

m4t5k4 commented Oct 8, 2026

Copy link
Copy Markdown

Independent RTX 3070 validation — Qwen3.6 CUDA chain (PR #1992)

Independent real-model validation of PR #1992 at commit eaccb95ef180ea57249824170b02992c5c20026a: four-GPU smoke tests, 200-token inference and balanced A/B benchmarking passed without model swapping.

Scope: This is a four-GPU validation, not an exact reproduction of the original single-GPU 18.7 ms/token result.

Hardware and configuration

  • Model: Qwen3.6-35B-A3B, validated per-row int4 checkpoint.
  • GPUs: Four NVIDIA RTX 3070 8GB cards, physical indices 0, 1, 2 and 4, isolated by UUID and verified as CUDA-local ordinals 0–3.
  • CPU: Intel Core i5-6600K, four cores.
  • RAM: 32 GB installed (~30 GiB usable).
  • CUDA: Toolkit 12.4.
  • Power limits: 120 W per GPU, unchanged.
  • Expert cache: 256 experts/layer, default warmstart.

Placement:

COLI_GPUS=0,1,2,3
COLI_CUDA=1
CUDA_EXPERT_GB=6.1
COLI_PLACE=lmhead=0,dnproj=0,dnout=0,attnproj=0,shexp=0,experts=all
Q36_OFFER_SHEXP=1

QT_HOME and CUDA_RELEASE_HOST were unset. The latter is not a Qwen-engine setting.

Runtime validation

  • Complete 40-layer dense CUDA chain and LM head resident on local GPU 0.
  • 190 additional trunk/shared matrices and 30 DeltaNet input projections placed on GPU.
  • 10,240/10,240 routed experts resident in VRAM across four GPUs.
  • 100% decode VRAM expert hits.
  • Warmstart released redundant int8 host copies of resident experts while retaining packed int4 sources.
  • No CPU expert misses, expert swaps or backend fallback.
  • Eight-token smoke tests completed successfully.
  • All completed 200-token runs produced byte-identical generated text and finite final logits.

Numerical check: Default A/B final logits differ because the CPU router rounds activations to int8, while the GPU router uses float activations. With COLI_DENSE_IDOT=0 in a separate diagnostic, final-logit cosine similarity reached 0.999999999999657, with maximum absolute difference 1.57e-5. This diagnostic was excluded from performance comparisons.

Balanced A/B benchmark

Fixed 49-token prompt, 200 new tokens, run order A1 → B1 → B2 → A2.

Both modes used identical model weights, binaries, GPU devices, budgets, placement, pinned CPU threads and fresh-process/full-expert warmstart policy.

  • A: Q36_DN_GPU=1, COLI_CUDA_CHAIN=0
  • B: COLI_CUDA_CHAIN=1, Q36_DN_GPU unset
Run Mode Decode ms/step Decode tok/s TTFT (s) Generation tok/s Dense-chain wait ms/step Host routed interval ms/step Expert hits
A1 A 42.1 23.75 1.68 19.75 N/A N/A 100%
B1 B 24.3 41.15 1.13 33.24 12.83 9.35 100%
B2 B 24.8 40.32 1.39 31.23 13.41 9.31 100%
A2 A 40.8 24.51 1.86 19.90 N/A N/A 100%

Results:

  • Mean A decode latency: 41.45 ms/step.
  • Mean B decode latency: 24.55 ms/step.
  • 40.8% lower decode latency with the CUDA chain.
  • Approximately 1.69× higher reciprocal decode throughput.
  • A separate full CUDA-chain run measured 23.6 ms/step, 1.08 s TTFT and 34.17 generation tok/s.

Memory and stability

  • Peak RSS during successful runs: 21.88 GiB.
  • Minimum sampled available system RAM during A/B: 5.15 GiB.
  • Zero process swap and zero system swap-out across the four benchmark runs.
  • Maximum sampled VRAM usage:
    • GPU 0: 5,991 MiB.
    • Other three GPUs: 4,037 MiB each.
  • No backend fallback.
  • All selected GPUs returned to idle after teardown.
  • Safety guards remained active throughout the runs.

Measurement notes

  • Decode latency figures are rounded aggregate model-forward times across 199 decode steps.
  • Reciprocal decode throughput excludes outer argmax and streaming bookkeeping.
  • Generation throughput includes TTFT but excludes load/warmstart.
  • No latency percentiles are inferred.
  • B's host routed interval includes GPU group waits, not CPU expert arithmetic.
  • A's cpu-miss-labeled timing region includes shared-expert GPU work and is not directly equivalent to B's dense-frame wait metric.
  • Historical system swap occupancy should not be confused with new swapping during the tests.

Comparison with the original benchmark

The upstream single-GPU result of 30–32 → 18.7 ms/token was not directly reproduced.

This test differs in several ways:

  • Four RTX 3070s instead of one.
  • Full expert residency with 100% decode VRAM hits.
  • Different CPU and shared-expert placement.
  • A separately selected 49-token prompt, because the original prompt text was unavailable.

The balanced comparison holds expert residency constant between A and B and demonstrates the improvement from switching execution paths in this configuration. It does not isolate a kernel-only effect.

Conclusion

The results provide independent real-checkpoint validation of the shared NVIDIA CUDA chain and multi-GPU expert-group execution on RTX 3070 hardware.

No production source changes, commits or pushes were made. The tested commit and existing test executables were preserved.

Local benchmark artifacts are stored under /tmp/colibri-pr1992-4gpu/.
1790791519772

@kreuzzelg

kreuzzelg commented Oct 9, 2026 •

Copy link
Copy Markdown
Contributor Author

@JustVugg: Thanks. The four are in (4eea949, 083d5e5, eaccb95): the ninth id (twelve now; the test also compares the K/V rows it wrote rather than the whole malloc'd cache, which UCRT64 caught), the rule without its headers, the two calls declared inside cuda_chain.cu alone, and the loader tripwire at 16 optional / 63 exports. CI is green on the last push.

Yes to Qwen3.8, and it is done on top of this branch (kreuzzelg:qwen38-cuda-chain, one commit plus its numbers; I will open it as its own PR once this one is in): qwen38_cuda_chain.h on the same table, qwen38_chain.h's frames and contract -- the four streams on the device through the gated residuals (HC ops), the PLE layer (the n-gram rows from the host, the projections, gate and dilated convolution on the device), QSA (the index keys into the device cache, each completed block's pooled key once, the indexer's top-k per row, the attention over the listed positions), the DeltaNet with the sigmoid gate, the final mixer and lm_head, a verify's per-row copies of the DeltaNet state, the conv rings and the PLE ring, the rollback, a lost device rebuilt from the prefix record. The table gained cc_qsa and cc_ple (chain_qsa.comp and chain_ple.comp ported line for line) and the recurrence at KD 4 for the tiny fixture. The matrices are the tier's where the placer put them; the rest (the gated residuals' low-rank pairs, b/a, the PLE projections, the final mixer: what Q38_TRUNK_MIN_KB never offers) the chain uploads itself from the same bytes, int8 rows where the CPU holds them, else BF16 as f32.

Measured on the 3070 (FP8 checkpoint, the trunk's 553 matrices as int8 in VRAM, 16 % expert hits behind it, 100 tokens, two runs each):

TTFT tok/s ms per token resident-mm / forward DeltaNet QSA PLE lm_head shared
Q38_DN_GPU=1 (#1870) 19.0 s 1.22 / 1.26 612-632 172 ms 161 23 12 26 23
COLI_CUDA_CHAIN=1 17.7-17.9 s 1.41 / 1.43 526-535 74 ms in the chain in the chain in the chain 10 0.4

So the chain takes the whole dense part at once: -15 % per token and 1.2 s off the first token. The token stays the experts' -- about 1 s of expert reads per forward (24,576 FP8 experts behind a 4 GiB trunk on 8 GB) and 290 ms of routed experts -- so on this card the trunk's weight is what you said it would be: it pays, but it cannot carry the model past its expert reads. On a card with room for the experts the dense share would be the visible one. VRAM: the chain adds 174 MiB of its own uploads and 381 MiB of state (the K/V, index-key and pooled-key mirrors, the DeltaNet state, the PLE ring) beside the tier.

Verified: the tiny FP8 fixture's oracle on the fake tier with the chain on (8/8 tokens, cosine 0.99998), a frame failing mid-run (rebuilt on the CPU, within the limits), and the per-matrix path, from qwen38-cuda-chain-check beside qwen38-dn-gpu-check; the QSA block keys and selections (ties as the CPU sorts, the causal case) and the PLE gate and convolution on the 3070 in test_cuda_chain. One thing of yours I found on the way: q38_forward's final lm_head was bound to chain_logit under COLI_VULKAN only, so a second chain backend recomputed the logits from the CPU's mixed; it is bound for either now.

@m4t5k4, thank you for the four-card run, that is the first independent real-checkpoint measurement of the chain. Two things in your table are worth more than the headline: with every expert resident your host routed interval is 9.35 ms a step, and mine on one card at 81 % hits is 9.9 -- so the routed-expert cost is the tier's issue/take path itself, not the CPU misses, which makes it the next lever after this PR (I had guessed the misses). And your chain waits 12.8-13.4 ms a step on the device where mine waits 5.7 for the same matrices: the chain's host-mapped buffers (the routed sum up, the rows and logits down, 40 times a token) cross the bus on every frame, so I would expect that gap to be the lanes the cards sit on (x4 or a riser on a four-card board?) and the 120 W cap rather than the kernels. A mining-style board puts each card on a PCIe x1 riser, about 1 GB/s, and every one of those small crossings then costs its latency over a narrow link, where on x16 they vanish. If you can say what the lanes are, that tells us whether the UP/DOWN buffers should become device-side with one copy per frame instead of mapped memory. And if you have the time for one more run, the one that would settle it: a single 3070 in an x16 slot, the same A/B (COLI_GPUS=0, the trunk and whatever experts fit on it), and its [chain] line here. On a single card at x16 the chain waits 5.7 ms a step for me; your number from a second machine would tell us whether the rest of the gap is the lanes or the kernels.

Thanks to Claude Code for the help along the way.

@m4t5k4

m4t5k4 commented Oct 9, 2026

Copy link
Copy Markdown

@kreuzzelg Thanks for looking into my 4× RTX 3070 results!

My inference server is built from hardware left over from the Ethereum PoW mining days: an Intel i5, an ASRock Q270 mining motherboard, and RTX 3070s connected through PCIe x1 risers.

The motherboard has no M.2 NVMe slot, just one PCIe x16 slot and twelve PCIe 3.0 x1 slots. All four GPUs used in the benchmark are on x1 risers, while the x16 slot holds a PCIe-to-NVMe adapter with my 2 TB SSD.

Tomorrow I'll move one RTX 3070 directly into the x16 slot. I've already copied the Qwen3.6 checkpoint to a SATA SSD so I can remove the NVMe adapter.

I also had the GPUs capped at 120 W on Codex's recommendation, since we expected inference to use less power. I'll test a higher power limit as well.

I'll rerun the same A/B benchmarks, check the actual PCIe link width and speed, monitor GPU power and clocks, and share the chain results.

@m4t5k4

m4t5k4 commented Oct 11, 2026 •

Copy link
Copy Markdown

@kreuzzelg Quick update on the PCIe investigation!

I temporarily moved one RTX 3070 into the CPU-connected PCIe 3.0 x16 slot, replacing my NVMe adapter, and reran the Qwen3.6 CUDA-chain benchmark on 4 GPUs using the SATA checkpoint.

With the dense CUDA chain on an x1 GPU:

  • 23.4 ms/token, chain wait 12.69 ms/token

With the same four GPUs reordered so the dense chain runs on the x16 GPU:

  • 18.2 ms/token, chain wait 6.83 ms/token

That's 22.2% lower decode latency and 46.2% lower chain wait. All 200 tokens and final logits matched exactly, with 100% decode expert-cache hits.

These are single diagnostic runs, and GPU placement changed as well as PCIe connectivity, so it's not a pure PCIe isolation test. But the results strongly support your suspicion about the x1 risers.

I'm now looking into which host/device crossings contribute most to the remaining wait time. Thanks again for pointing me in this direction!

Update on power limit:

Metriek 120 W 170 W
Decode 18,2 ms/token 17,7 ms/token
Chain-wait 6,83 ms/token 6,32 ms/token
Routed experts 9,16 ms/token 9,29 ms/token
Output en logits Correct Correct

@mfethe1

mfethe1 commented Oct 11, 2026

Copy link
Copy Markdown
Contributor

@kreuzzelg — I tested your qwen38-cuda-chain branch at 37fe0589ccfc50628f344941f77e59055f65238f and have a small test-only patch to offer for its eventual PR.

The existing qwen38-cuda-chain-check covers chain-on, a frame failing mid-run with CPU recovery, and chain-off, with Q38_MTP=0. This patch adds a generated tiny FP8-experts MTP forced-rejection case, checking the oracle, chain frames, and DeltaNet snapshot writes. No production code changes: three files, 23 additions / two deletions.

Validation on macOS, CPU fake CUDA backend:

  • make check: exit 0; Python portion ran 2,417 tests, with 270 skipped.
  • make -C c qwen38-cuda-chain-check: exit 0; all four cases matched 8/8 reference tokens. The new rejection case had numeric cosine 0.99998191, max absolute error 0.00078176 against the test's 0.01 bound.
  • Negative control: in a disposable copy, removed the recurrent/convolution buffer swap-back in q38cc_rollback and rebuilt. The original three cases still passed; the new case failed with 6/8 tokens and numeric cosine -0.18893308, executable exit 1. Restoring the code and rebuilding returned it to passing.

The new case reports 0/11 drafts accepted and executes snapshot-enabled chain operations. This is all-draft rejection/rollback coverage, not mixed acceptance, real CUDA-device validation, or a performance measurement. I have not run cuda-test on Linux hardware. Dependencies for the real tiny-fixture generator were Python 3.14.7, Torch 2.14.1 and Transformers 5.16.1; fixture generation ran offline, without downloading a checkpoint.

The patch below applies cleanly to that exact branch revision. Would this be useful in your Qwen3.8 PR? Happy to follow your preferred handoff/target rather than open a competing PR against dev before your branch lands. If a newer/unpublished test already covers this, I'll avoid duplicating it.

Disclosure: AI-assisted test development and review; the checks and mutation were executed and independently rerun by the coordinating agent. No claim that this reveals a current production bug.

Test-only patch against 37fe058 (forced-rejection regression)
diff --git a/c/Makefile b/c/Makefile
index 0495241..7eea307 100644
--- a/c/Makefile
+++ b/c/Makefile
@@ -1501,7 +1501,7 @@ qwen38-tiny-check: qwen38-tiny-generate qwen38$(EXE)
 # plus BF16 128x128 block scales, gate/up and down in separate shards. The
 # native slab path and the expanded path must both reproduce the reference;
 # then the engine runs on the fake CUDA tier (no toolkit, no GPU).
-.PHONY: qwen38-tiny-fp8-generate qwen38-tiny-fp8-check qwen38-tier-engine-check qwen38-cuda-chain-check
+.PHONY: qwen38-tiny-fp8-generate qwen38-tiny-fp8-mtp-generate qwen38-tiny-fp8-check qwen38-tier-engine-check qwen38-cuda-chain-check
 qwen38-tiny-fp8-generate:
 	$(PYTHON) tools/make_qwen38_tiny.py --out ./qwen38_tiny_fp8 --fp8-experts
 
@@ -1518,9 +1518,12 @@ qwen38-tier-engine-check: qwen38-tiny-fp8-generate tests/test_qwen38_tier_engine
 qwen38-dn-gpu-check: qwen38-tiny-fp8-generate tests/test_qwen38_dn_gpu$(EXE)
 	./tests/test_qwen38_dn_gpu$(EXE)
 
-qwen38-cuda-chain-check: qwen38-tiny-fp8-generate tests/test_qwen38_cuda_chain$(EXE)
+qwen38-cuda-chain-check: qwen38-tiny-fp8-generate qwen38-tiny-fp8-mtp-generate tests/test_qwen38_cuda_chain$(EXE)
 	./tests/test_qwen38_cuda_chain$(EXE)
 
+qwen38-tiny-fp8-mtp-generate:
+	$(PYTHON) tools/make_qwen38_tiny.py --out ./qwen38_tiny_fp8_mtp --fp8-experts --mtp
+
 # The routed experts as int4-g64: tools/convert_qwen38_experts_int4.py writes
 # the experts-int4g64/ sidecar next to the FP8 shards and qwen38 picks it up.
 # The fixture is the FP8 one with its experts scaled by 3 (at the default
diff --git a/c/tests/qwen36_fake_cuda.h b/c/tests/qwen36_fake_cuda.h
index 5783288..28e1b8a 100644
--- a/c/tests/qwen36_fake_cuda.h
+++ b/c/tests/qwen36_fake_cuda.h
@@ -154,6 +154,7 @@ int coli_cuda_matmul(ColiCudaTensor **tensor, float *y, const float *x, const vo
 struct ColiCudaDn { int device, vh, vk, kdim, vdim, conv_dim, convk, hidden; float eps; int gate_sigmoid;
                     float *ring, *rec, *conv_w, *norm_w; };
 static int fake_dn_steps, fake_dn_fail;
+static int fake_dn_snapshots; /* chain MTP tests: prove verify snapshots were written */
 static void fake_gemv_i8(float *y, const float *x, const ColiCudaTensor *t) {
     const int8_t *q = (const int8_t *)t->w;
     for (int o = 0; o < t->O; o++) { float a = 0.f; const int8_t *w = q + (size_t)o * t->I;
@@ -371,6 +372,7 @@ static int fcc_attn(CcBuf *q, CcBuf *kc, CcBuf *vc, CcBuf *o, CcBuf *gate, CcBuf
 static int fcc_dnconv(CcBuf *in, CcBuf *w, CcBuf *ring, CcBuf *out, CcBuf *snap, const CcDnConv *p) {
     if (!fcc_ready() || !in || !w || !ring || !out || !p || (p->snap_row >= 0 && !snap)) return 0;
     fake_chain_ops++;
+    if (p->snap_row >= 0) fake_dn_snapshots++;
     int nh = p->CK - 1;
     for (int c = 0; c < p->CD; c++) {
         float hist[8] = {0};
@@ -393,6 +395,7 @@ static float fcc_softplus(float x) { if (x > 20.f) return x; float e = expf(x);
 static int fcc_dnrec(int KD, CcBuf *cv, CcBuf *ab, CcBuf *z, CcBuf *st, CcBuf *prm, CcBuf *y, CcBuf *snap, const CcDnRec *p) {
     if (!fcc_ready() || !cv || !ab || !z || !st || !prm || !y || !p || (p->snap_row >= 0 && !snap) || KD > 256 || p->VD > 128) return 0;
     fake_chain_ops++;
+    if (p->snap_row >= 0) fake_dn_snapshots++;
     for (int h = 0; h < p->VH; h++) {
         int kh = h / (p->VH / p->KH);
         float *S = st->d + p->st_off + h * KD * p->VD;
diff --git a/c/tests/test_qwen38_cuda_chain.c b/c/tests/test_qwen38_cuda_chain.c
index d7b7e9d..9f7146b 100644
--- a/c/tests/test_qwen38_cuda_chain.c
+++ b/c/tests/test_qwen38_cuda_chain.c
@@ -67,6 +67,21 @@ int main(void) {
     tk(fake_chain_frames == 0, "no frame reached the card");
     qt_shutdown();
 
+    printf(" MTP draft rejection with the chain on the fake card\n");
+    setenv("SNAP", "./qwen38_tiny_fp8_mtp", 1);
+    setenv("Q38_MTP", "1", 1);
+    setenv("Q38_MTP_FORCE", "reject", 1);
+    argv[3] = (char *)"./qwen38_tiny_fp8_mtp/ref.json";
+    setenv("COLI_CUDA_CHAIN", "1", 1);
+    fake_chain_frames = 0; fake_chain_fail_at = 0; fake_dn_snapshots = 0;
+    rc = qwen38_main_unused(4, argv);
+    tk(rc == 0, "rejected MTP drafts roll back and remain within the oracle's limits");
+    tk(fake_chain_frames > 8, "rejected MTP verification ran through chain frames");
+    tk(fake_dn_snapshots > 0, "MTP verification wrote DeltaNet chain snapshots");
+    unsetenv("Q38_MTP_FORCE");
+    setenv("Q38_MTP", "0", 1);
+    qt_shutdown();
+
     if (t_fails) { printf("test_qwen38_cuda_chain: %d failure(s)\n", t_fails); return 1; }
     printf("OK test_qwen38_cuda_chain: Qwen3.8's chain on the card gives the oracle's tokens\n");
     return 0;

@kreuzzelg

Copy link
Copy Markdown
Contributor Author

A supplementary data point, since Ollama shipped 0.40.3 in the meantime: the same prompt and the same card, Ollama against this PR's chain. Ollama's per-card numbers are its eval rate (decode only, --verbose, --think=false, one cold run and two warm, the single-card instance under CUDA_VISIBLE_DEVICES=0 OLLAMA_VULKAN=0); Colibrì's are the 06 Oct runs from the table above (same build as this branch but for the tests).

RTX 3070 8 GB, Qwen3.6-35B-A3B, 200 new tokens decode tok/s (warm) ms per token prompt VRAM
Ollama 0.40.3, Q4_K_M, the 3070 alone 38.1 / 38.4 26.2 61 tokens at 34-40 tok/s 7.0 GB
Ollama 0.40.3, Q4_K_M, 3070 + Quadro RTX 4000 (its layer split) 37.1 / 37.0 27.0 37.9 tok/s 4.7 + 7.3 GB
Ollama 0.34.4, the 3070 alone (28 Sept) 38.1 / 38.2 26.2 38-39 tok/s 7.1 GB
Colibrì Q36_DN_GPU=1 (#1784, the per-layer step) 32-33 30-32 49 tokens, TTFT 1.22-1.26 s
Colibrì COLI_CUDA_CHAIN=1 (this PR) 53 18.7 49 tokens, TTFT 0.98 s ~7.2 GB

Ollama did not move between 0.34.4 and 0.40.3 on this box (and its second card still costs it about a token a second, as on 28 Sept). The chain puts this engine 38 % ahead of Ollama's decode on the same card, where the per-layer step was 15 % behind it. The quantizations differ, as before: Ollama's Q4_K_M keeps down and output at Q6_K, this container is per-row int4 with the int8 trunk; the comparison is of the two engines as they ship, not of the kernels alone. Raw logs in my workspace if anyone wants them.

kreuzzelg added a commit to kreuzzelg/colibri that referenced this pull request Oct 12, 2026
)

The fourth case of test_qwen38_cuda_chain: the tiny FP8 fixture with its
MTP head (qwen38-tiny-fp8-mtp-generate), Q38_MTP_FORCE=reject, the chain
on: every draft rejected, every verify rolling the device's copies back
(q38cc_rollback), the run within the oracle's limits (8/8 tokens, cosine
0.99998), and the fake counting the DeltaNet snapshot writes. Dropping the
rollback's buffer swap fails it (6/8 tokens): the case guards the
contract the first three did not reach. Offered on JustVugg#1992 against this
branch's 37fe058; applied as is.

Co-authored-by: mfethe1 <mfethe1@users.noreply.github.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
@kreuzzelg

Copy link
Copy Markdown
Contributor Author

@mfethe1, yes, that is useful, and it is in: applied to qwen38-cuda-chain as it was (now de67631), with you as co-author, after running it here -- the four cases pass (the new one 8/8 tokens, cosine 0.99998), and it covers exactly what the first three did not reach, the verify's device copies and q38cc_rollback, which I had listed as untested. It will go into the Qwen3.8 PR with the rest once #1992 is in; no need for a PR of your own. One note for anyone generating the fixture locally: make_qwen38_tiny.py pins transformers 5.16.1, the oracle requirements file installs newer, so the pin has to be met by hand (the CI job does).

@JustVugg, a question on scope rather than code: the Qwen3.8 chain is finished on kreuzzelg:qwen38-cuda-chain, three commits on top of this branch (+1418 / -17 in 12 files: cc_qsa, cc_ple and the KD-4 recurrence in the table, qwen38_cuda_chain.h, the fixture test with mfethe1's rejection case, the 3070 numbers in docs/qwen38.md), fast-forwardable onto this one. This PR stands at +2561 / -21 in 18 files, about twice #1871; with Qwen3.8 it would be four times. I read your "after this one" as a second PR and have kept it that way, so this one stays the commit m4t5k4 and mfethe1 measured -- but if you would rather review the two chains as one, I push the three commits here today. Your call; nothing moves until you say.

@JustVugg

Copy link
Copy Markdown
Owner

Keep them separate, please: #1992 as it stands (the commit m4t5k4 and mfethe1 measured), and the Qwen3.8 chain as the next PR on top of it once this one is in. Two reviews of this size are easier to get right than one of four times #1871. Thanks for folding in mfethe1's rejection case, that was exactly the gap in the first three. I'll review #1992 now.

@JustVugg
JustVugg merged commit d6733c5 into JustVugg:dev Oct 12, 2026
114 checks passed
kreuzzelg added a commit to kreuzzelg/colibri that referenced this pull request Oct 14, 2026
)

The fourth case of test_qwen38_cuda_chain: the tiny FP8 fixture with its
MTP head (qwen38-tiny-fp8-mtp-generate), Q38_MTP_FORCE=reject, the chain
on: every draft rejected, every verify rolling the device's copies back
(q38cc_rollback), the run within the oracle's limits (8/8 tokens, cosine
0.99998), and the fake counting the DeltaNet snapshot writes. Dropping the
rollback's buffer swap fails it (6/8 tokens): the case guards the
contract the first three did not reach. Offered on JustVugg#1992 against this
branch's 37fe058; applied as is.

Co-authored-by: mfethe1 <mfethe1@users.noreply.github.com>
Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
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.

4 participants