Repository navigation
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
Conversation
…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>
|
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,
A question for after this one: would you take the same chain to Qwen3.8-Flash-Next (
So it would mostly be a |
…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>
Independent RTX 3070 validation — Qwen3.6 CUDA chain (PR #1992)Independent real-model validation of PR #1992 at commit Scope: This is a four-GPU validation, not an exact reproduction of the original single-GPU 18.7 ms/token result. Hardware and configuration
Placement:
Runtime validation
Numerical check: Default A/B final logits differ because the CPU router rounds activations to int8, while the GPU router uses float activations. With Balanced A/B benchmarkFixed 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.
Results:
Memory and stability
Measurement notes
Comparison with the original benchmarkThe upstream single-GPU result of 30–32 → 18.7 ms/token was not directly reproduced. This test differs in several ways:
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. ConclusionThe 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 |
|
@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 Yes to Qwen3.8, and it is done on top of this branch ( 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):
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 @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 ( Thanks to Claude Code for the help along the way. |
|
@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. |
|
@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:
With the same four GPUs reordered so the dense chain runs on the x16 GPU:
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 |
|
@kreuzzelg — I tested your The existing Validation on macOS, CPU fake CUDA backend:
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 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 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; |
|
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
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 |
) 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>
|
@mfethe1, yes, that is useful, and it is in: applied to @JustVugg, a question on scope rather than code: the Qwen3.8 chain is finished on |
|
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. |
) 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>

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 likevk_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 ofvk_chain.c. BuffersCC_DEV,CC_UP,CC_DOWN(the host-side ones page-locked and mapped: the device reads or writes them in place, the host throughcc_ptr, as your host-visible buffers); framescc_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 opscc_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 ofshaders/chain_*.compported line for line, the parameter structs field for field theVkcones (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_exrouted only; that path computed the shared expert regardless ofrouted_onlyunder 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=nfakes 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=1offers the shared expert to the placer asQ36_OFFER_SHEXP=1does;COLI_PLACE=autotakes 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, orKV_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_matmulspends 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_asynctakes it for int8 rows, up to 8 activation rows each on its own (a verify's rows get a decode step's bits, asq36_spec_stepasks; 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)Q36_DN_GPU=1(#1784, the per-layer step)COLI_CUDA_CHAIN=1, the block GEMVCOLI_CUDA_CHAIN=1, the warp GEMVThe 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_GPUwas; 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) throughstep_exon the fake tier, the chain's ops from a host-side table intests/qwen36_fake_cuda.hwith 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-compilebuilds 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.VK=1,CUDA=1;test_makefile_vk_obj.pyandtest_backend_loader.pygreen; acuda_chain.custep in the CUDA syntax check.--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 ofQT_HOME=layerwould need a chain per card, as yourCOLI_VK_DEV2).Thanks to Claude Code for the help along the way.