docs/open-gaps.md
Open gaps
| Track | Status | Next gate |
|---|---|---|
Device expert slices, CUDA against CPU (ENG-EXPERT-STREAM-DEVICE, #1299
) | G0-SPEED VOID behind G0-CORRECT FAIL; no ratio is accepted | Resolve the block-0 router-input divergence, then run a clock-controlled repeat |
Surface coverage (ARCH-ONE-SURFACE) | CORRECTNESS COMPLETE: #139 restores DSR 32 (kcuda=0) via registry/name resolution; ABI-v14 selection unchanged; no speed claim | Selector 2/2·11 plus execution-bound CMake/File-API/CTest + CI/preflight + manifest-integrity guard 52/52; CPU platform/loader/C-ABI tests green; CUDA A/B remains residual |
| 35B prefill TTFT | 0.920x c1 / 0.849x c4 against a correctly FUSED denominator at the pin; the 0.93x-0.98x reading came from an UNFUSED oracle (#414) and is void | Re-attribute the residual against the fused denominator, then close |
| c16 and above, both models (#577) | VOID, not a number. 93/96 against the pin’s 96/96, the three missing being the SLOWEST: our SSE keepalive was ON. Recurred on Qwen3.8-27B at 5/6 and 36/48 (#931 ) | Keepalive now default OFF and rate harnesses refuse failed != 0; re-run c16/c32 on a binary carrying both |
| Clock-controlled pin grid (#520, #414, #543) | First defensible series LANDED at 1024/128, n=3, interleaved, one boot, flat 2184 MHz. Only c1, c4 and a partial c16 exist | c2, c8 and c32 at the pin under clock control, so the sweep is a sweep |
| 35B low-batch MoE decode | CLOSED at low batch (c1 0.975x, c4 wins); c16 0.93x. VT_ASYNC_DEVICE_MIRROR default ON for correctness. VT_ASYNC_EXECUTOR Option A (H2D out of capture) A/B’d speed-NEUTRAL | c16 lever is prefill glue (task #61), not the decode drain. test_qwen36_async_serving GREEN |
| CPU keep-quant MoE decode | No number owed: correctness-only P0. The grouped keep-quant GEMM read activations as f32 whatever their dtype, so CPU MoE decode emitted token-0 garbage from b4f5610a (2026-07-31) | Speed unmeasured and unclaimed; test_ops_quant_dot GREEN (150224 assertions) |
Accepted-and-inert serve args (SERVE-RECIPE-ARGS, #606) | No number owed: argument parsing only, so nothing to time and no oracle leg. Correctness gate 4 cases / 58 asserts GREEN, RED-first, mutation-proven | None. A speed axis would be fabricated; closes on review plus the operator gate rerun |
GDN core/z at the model dtype, MoE arms (GDN-MOE-BF16-OUT, #1168
) | No number, none claimed. GATING, CPU tier only. A token gate cannot see this axis: f32 is the MORE precise deviation (spec
) | 35B correctness first, then 315/315, 235/235 dense inertness, the VT_GDN_OUT_BF16 0-versus-1 A/B per leaf, and nsys for the memory format. All PENDING a GPU host |
NemotronH model row (MODEL-TEXT-nemotron-h-..., #517
) | No number, none claimed. The row’s FALSE KERNEL-SSM-MAMBA block is corrected (#1074); it STAYS INVENTORIED, because the A3 96/96 belongs to 0ea5d249f, not to current main | A re-run of the A3 gate against main. Only a pass measured there moves this row, or gives a grid a denominator |
| DeepSeek-V2-Lite MLA | Attributed miss, ACTIVE | Throughput at every concurrency |
| Qwen3.5 text-only arms (#490) | No number; run gates OWED, both PARTIAL. The loader half is CLOSED (#740, #864 DONE), so what blocks these is hardware, not a refusal | No fitting ckpt for either causal-LM arm: no denominator. Qwen3.8-2.4T-A95B is ~4.8 TB vs 128 GB; its load plan resolves, which is not a token |
| Darwin Qwen3.5 build repair (#1054, 2026-08-16) | NOT APPLICABLE. Removing a redundant namespace-scope lambda capture that Apple Clang rejects under -Werror changes no generated refusal text, no model math and no runtime path | None. The binding gate is the Apple Clang build |
| Qwen3.6-35B-A3B published BF16 (#740, #864) | No number, and none was owed: the 2026-08-15 gate measured TOKENS. Correctness MET vs the pinned oracle: 6/7 prompts STRICT 16/16, the 7th an exact tie (#910); SACRED 3/3 byte-identical | A throughput / latency / memory grid on this checkpoint. Nothing is measured, so nothing is claimed |
| MoE vision tower image + video (#891) | NOT gated, no number. The 333 model.visual.* tensors load and the tower computes, on sm_110 FALLBACK attention, not the shipped GB10 path. The token-exact mm gates never ran | Both modality gates on GB10 through the shipped fast path, then a per-modality speed grid |
| Dense image/video after the #891 merge (#908) | UNVERIFIED, network-blocked. Dense TEXT is 235/235 at 2f2bce926, a true before/after (binary md5 db889909d4… vs 49ded1ece8…, 500 TUs recompiled). The modality arms were not re-run | Re-run the dense image and video gates once the fixtures are reachable |
| Qwen3.5 upstream throughput levers (roadmap C10) | NOT MEASURED. vLLM’s 2026-08-06 25K tok/s/GPU is a GB200/NVLink72 disaggregated cluster result, not comparable to one GB10, and is NOT adopted as our bar | Advance the parity pin past 555967922 so the referenced PRs exist, re-capture goldens at zero drift, then port the GDN prefill kernel |
| DeepSeek-V4-Flash | Parity with ds4 (0.997x) | Optional beat-path: f16 tensor-core DSA/router (near-tie class) |
| DeepSeek-V4-Flash vs vLLM | Infeasible on one Spark | 2x GB10 with TP2 over the NCCL seam |
| Tensor parallelism (task #287 spike, 2026-08-08) | No number owed (benchmark_binding=false): records-only scope at pin 555967922 (spec
). TP-W1 landed (group table, 6/6); TP-W2..W7 gate on token-exactness, not speed | Perf gate is TP-W6: at or above vLLM TP=2 on every axis, same 2-GPU box (PENDING-HW) |
| Multimodal image, audio, video | Correctness gated, speed unmeasured | Per-modality speed grids |
Voxtral enc attn (VT_WHISPER_ENC_FA2, #432) | FA-2 hd-64: 731.7 -> 133.0 ms (5.50x); enc fwd (not TTFT) vs pin TTFT 46.02: 15.90x -> 2.89x | OPT-IN: 3 near-tie diverg (band PASS) vs 0. Dev call |
/v1/videos OpenAI + ONE-SURFACE ROW 2 | No number owed: ABI-v12 device selection is backend-dispatch plumbing; generation math and speed paths are unchanged | DSR 34→32; baseline/allowlist unchanged; 25/25 checker mutations; CPU fold 6/137, including one-queue/device-provenance mutations |
| Qwen3-dense decode CUDA-graph | Token-exact pass, ~4.3% e2e directional | Steady-state per-step tok/s |
| Kimi-Linear-48B-A3B (KDA+MLA+MoE) | RUNNER FOLD LANDS (ROW 7, §21, #122): engine==CLI 128/128 byte-identical; vs golden 122/128 (near-tie profile); FA2 MLA default-ON; SACRED green. Server 19.0 tok/s wall; CLI 18.93 reproduced | vLLM ~21 (#111 floor; in-session re-measure ABORTED by GB10 reboot at util 0.82, §21): ~0.90×, >= vLLM NOT met; residual = KDA host islands + grouped MoE + decode graph |
| vLLM 0.26 re-benchmark | Pending | Re-run the binding grids on the advanced pin |
| MiniMax-H3 FP4 speed (W-FP4a) | Measured GB10 (row/H3-FP4-GPU-E2E). Marlin W4A16 byte-exact vs bf16; fp4 a memory win, 0.8x bf16/forward. Real-ckpt fp4-resident e2e RUNS (mp4/wav) | fp4 speed CLOSED. bf16-vs-quant A/B: ENCODER half MEASURED (§8.15), DiT half NOT (no bf16 render exists). Detail: benchmark-record + spec §8 |
| LTX-2.5 axes | Speed PENDING (vllm-omni#6066 has no native 2.5), binding oracle too. SIZE: 704x448/25f and 448x256/25f both COMPLETE on GB10 (4231 s, 3085 s); one run each, contended box, no oracle, no ceiling (#1088) | NOT the VAE decode (#1041/#1009): 39-100% of the ~1731 s phase is the caption projection (#1208), x86 671.8→78.4 s; LoRA fusion 17.78→0.124 s (#1202), add-back 59% (#1254). #1286 REFUTED (#1317); alloc bound 464 MiB |
| LTX-2.5 video VAE decode device arm | Speed PENDING, and deliberately unmeasured: the row that landed vt::Conv3d (#1007) was dispatched with the GPU lease withheld and made no ratio claim | Correctness only: test_ops_conv3d 4/4 byte-exact vs an independent reference, goldens unmoved at 44/44, non-CPU dispatch byte-identical on a fake backend. Owed: a CUDA run (#1452); the non-conv stages (#1451) |
LTX-2.5 FULL 21.004B (one_stage, bf16 dev DiT) | FIRST render on the full model: 768x448/25f on GB10 in 2990 s, 25 frames + 1.01 s 48 kHz stereo, verifier PASS both arms. Binary 0a43a750 from 7b9e207b1, run 20260819T150230Z | REFUSED 1024x576, same lease: 162.0 s/fwd x 60 = 10803 s vs the rung’s 7153 s LEASE BUDGET (#1375). STG/modality were OFF (#1092); that refusal is LIFTED, full fidelity is 4 fwd/step. Detail: benchmark-record |
LTX-2.5 DiT self-attention kernel (LTX25-DIT-ATTN-FLASH, #1549) | A/B CLOSED, one binary and one lease at 768x448/49f (2352 tok), n=119 per arm, no stack sampler: naive 45.547 s vs flash 6.404 s per DiT forward = 7.112x. A third render repeats flash to 0.17% on the median. This replaces the old cross-run ~6.0x (6.03-6.23x) range and its single-arm 7.680 s figure | PIXEL GATE TAKEN AND FAILED (#1612 → #1743): the two arms render VISIBLY DIFFERENTLY — mean |delta| 6.414 vs a <= 1.0 bound, worst-frame PSNR 22.269 dB vs >= 40, worst SSIM 0.881 vs >= 0.99, 71% of the render’s own motion step vs a <= 10% bound; audio 29.368 dB and r 0.933. The control render is BIT-IDENTICAL, so the delta is entirely the kernel’s, and the op matches host to 8.94e-08 against a 2e-5 tolerance — 30 sampler steps at 4 DiT forwards each amplify it. Two NAIVE renders across builds diverge more still (9.452). Cap-raise REVERTED, head_dim evidence WITHDRAWN (#1578). See benchmark-record |
MiniMax-Music3 (MiniMaxMusic3ForConditionalGeneration) | Every axis vs the reference stays PENDING. A PARTIAL device arm now exists (#672): the 8.6B LM and the 2.4B fp32 DiT run on the accelerator, so the rows below are internal two-arm numbers and NOT parity ratios | Denominator: SGLang-Omni 748a0b43 in its production configuration (both CUDA graphs, compiled DIT and DAV, batched seeded sampling) |
| MiniMax-Music3 device arm, Jetson Thor sm_110 (#672) | --device 1 vs --device 0, same request/seed, idle box: 2 AR frames 846.6 vs 835.1 s (1.014x SLOWER); 10 frames 1430.4 vs 1512.1 s (0.946x). Fit: -11.65 s/frame, +34.8 s fixed | A third duration (the fit has no residual), and moving the depth decoder + DiT + vocoder, which are 5 of 6 stages and still host scalar loops |
MiniMax-Music3 DiT device arm, thor:gpu0 sm_110 (#672) | Per DiT forward at the capture’s geometry, same binary/weights/inputs, idle box: 204.955 s host vs 0.186-0.187 s device, 1094-1102x (1201x fitted). Staged ONCE (0.61 s; loop intercept 0.063 s). Whole process 3.5-4.5x | e2e song pair NOT runnable (host DiT alone ~37.6 h at 30 steps). Depth decoder/condition mix (bf16-storage), vocoder (no ConvTranspose1d) still host. Detail: benchmark-record |
| MiniMax-Music3 depth stage, x86-64 20-core (#672) | Thor sm_110, real checkpoint, 3 pairs (7b22b5b0): depth_forward 348.27 -> 78.32 s, 4.45x (1414 -> 808 calls); wall 446.33 -> 163.00 s, 2.74x. E2E audio byte-identical, 6 runs. x86 synthetic: 3.50x | Wall 2.74x not 4.45x: vocoder.decode_window a fixed 53.6 s, 32.9%, unowned. lm_decode_step fell 2.65 s unexplained; Thor vs x86 a HYPOTHESIS. Arms proved by CALL COUNT, not sha256 (#1516). Detail: benchmark-record |
| MiniMax-Music3 vocoder DECODE WINDOW (#672, #1334 ) | Thor sm_110 rc lease, host arm, two trees, two binary sha256, 3 alternated rounds: 1.364x at 20 latents to 1.439x at 344; second lease, 1 thread 2.157x | Never the conv kernel alone (5.2-5.8x at -O2). Do NOT multiply onto an e2e bucket: e2e reads 54.091 s where this reads 97.4463 s at 344, unexplained (#1512). CUDA memcmp holds. Detail: benchmark-record |
| MiniMax-Music3 CPU host kernels, x86-64 20-core (#672) | KERNEL A/B at the vocoder’s real geometry, min of 5 interleaved rounds: convolution chain 13.36 -> 1.25 s, 10.7x; Conv1d 12.03x, LinearNoBias 10.88x. Output fingerprints IDENTICAL on both arms | e2e pair VOID (cold CIFS cache; a foreign ctest at load 76.6) and re-running. Stages 0/1 only ~2x: the pivot trades WEIGHT locality for accumulator locality. Detail: benchmark-record |
MiniMax-Music3 END TO END on main, thor:gpu0 sm_110 (#672, #1512
) | Two clones d0598a255 -> a50c57d69, one lease, staged ckpt, medians of 3 alternated rounds. 4 s/4 steps wall 166.04 -> 48.07 s, 3.45x; ar.depth_forward 80.73 -> 4.27 s, 18.90x. 20 s/30 steps 5.49x | 4.53x with both loads excluded (the 20 s pair also crosses CIFS -> local). The DiT is now 62.24 % at 370.556 s/120 calls, unchanged, owned by no row. lm_decode_step fell 1.49x unexplained. Detail: benchmark-record |
MiniMax-Music3 DiT intra-forward split, thor:gpu0 sm_110 (#672, #1542
) | One binary, VLLM_CPP_MUSIC3_DIT_SPANS the only variable, idle box, staged ckpt, seq 690: vt::AttentionCross is 43.9% of the forward for 4.0% of its flops, 0.204 vs the GEMMs’ 3.98 TFLOP/s; 27.6% of the run | Instrument only; nothing got faster. GEMMs sit at 71.5-79.0% of the fp32 peak, so the lever is a kernel (#1555 ), not precision. Detail: benchmark-record |
MiniMax-Music3 DiT attention kernel, thor:gpu0 sm_110 (#672, #1555
) | ONE binary, arms by VT_OP_PROVIDER_DISABLE, engine-announced provider as the control, staged ckpt: at 20 s/30 steps wall 595.496 -> 449.969 s, 1.3234x, denoise.dit_device 1.6461x, dit.attn 8.607x | Same 1.6461x at 2 steps over 3 alternated pairs, spread 0.080%. All 6 mutations RED. ncu UNAVAILABLE here, so the mechanism is ABLATION: occupancy and bandwidth REFUTED. Detail: benchmark-record |
| MiniMax-Music3 vocoder decode-window SCALING (#672, #1664 ) | Thor sm_110 lease, ONE boot, six distinct binary sha256, conv_out user/wall 1.00 -> 13.06 as the behavioural control. THE SHIPPED ARM, SWEPT: 2.85x -> 11.54x of 14 threads, 3 alternated rounds; 4.08x against the baseline, PAIRED over 7 alternated rounds in one job. vocoder.snake 11.8x | The 11.48x/4.11x pair was arm C’s curve and a ratio composed across jobs; both are now measured and #1683 is CLOSED. #1770 is CLOSED too: the blocking condition was re-measured against a NULL arm over 31 paired rounds at three lengths, read 1.0086x/0.9721x/0.9700x on its own shapes – every one inside its own null – and was REMOVED. The 11.54x and 4.08x stand; the condition decided 4 shapes of 11 and was worth 0.36 ms of a 3.5 s window |
MiniMax-Music3 vocoder CUDA arm, thor:gpu0 (#672, #1474
) | Same binary, VLLM_CPP_VOCODER_DEVICE the only variable, idle box, one lease. 344 latents host 15.078 vs CUDA 19.288 s = 0.782x; ~689x4 122.169 vs 150.060 = 0.814x. f32 bought the device arm 2.78x | It LOSES: #1356 + #1474 moved the HOST kernel 3.59x against the device arm’s 2.78x, so §13.10’s f64 suspicion was right and is no longer the blocker. WAV byte-identical across arms, 12 runs. Detail: benchmark-record |
MiniMax-H3 render coherence (row/H3-RENDER-CLOSE #77) | CLOSED: a COHERENT scene on GB10. #70/#74 white was wrong-PARTITION usage (t2va on the ref2va ckpt); t2va on the FL2VA GGUF renders a prompt-matched orange cat (adj-cos 0.95 vs 0.06, no patch-grid) | Verified first: t2va inputs byte-exact vs upstream; CUDA device==host at seq 1920. Follow-up H3-TASK-PARTITION-GUARD: the task/partition mismatch now RAISES 1:1 with _resolve_task (spec §8.6-8.7) |
MiniMax-H3 image conditioning (row/H3-CONDITIONED-E2E, row/H3-VISION-SCATTER, row/H3-REF2VA-ASSEMBLY) | fl2va COHERENT; ref2va assembly bug FIXED+gated. vision→cond scatter gated; ref2va block-dim double-division fixed + RED-first gated (128 vs 512) + a permanent ref2va DiT-forward rung (§8.10) | grid RE-ATTRIBUTED: with the fix ref2va grids in fp4 AND bf16, and t2va with no refs on the ref2va NVFP4 also grids while FL2VA-GGUF renders, so it is the NVFP4 checkpoint/loader, NOT assembly/fp4 (§8.10) |
| MiniMax-H3 Thor render speed (sm_110, no FA2) | 34.6 s/step at 864x480/124f/50 steps on Q4_K_M, 16.6x off 574.5 (render ~28 min, was ~8 h). Landed: warp-per-query, chunked warp reduce-scatter (1.76x), bf16 mma.sync (9.82x) | Shared-memory K/V tiling (23% SLOWER) and register Q-blocking (-0.8%) both measured and REVERTED: memory traffic is not the bound (one head’s K+V is 3.9 MB against 32 MB of L2) |
| MiniMax-H3 render duration (audio halving) | CLOSED: a 124-frame render silently muxed as 61. The decoded audio ran half the video’s duration and the muxer passes -shortest. Every structural check passed: shapes were self-consistent, just halved | Gated on the duration invariant (latent steps / 40 Hz equals the video duration), not on shape self-consistency, which a halved pipeline satisfies |
| MiniMax-H3 quantization floor | Use Q4_K_M, not Q3_K_M. H3’s split-half RoPE produces channel-wise magnitude outliers 3-bit cannot hold; a controlled A/B (same prompt, seed, code) turned a murky lattice-covered silhouette into a photoreal close-up | Per-tensor mixed precision, if a smaller footprint is ever owed |
MiniMax-H3 image conditioning (row/H3-CONDITIONED-E2E, row/H3-VISION-SCATTER, row/H3-REF2VA-ASSEMBLY, row/H3-NVFP4-LOADER-DIFF, row/H3-NVFP4-STREAM-DIFF) | fl2va COHERENT; ref2va NVFP4 nibble loader bug FIXED (byte-verified); grid residual DIAGNOSED §8.12 | Activation diff + fingerprints (PR #95): NO load-path defect; all weights/islands/RoPE quant-noise-close to the coherent GGUF; grid = community-NVFP4 quant fidelity, not a loader fix. See benchmark-record |
MiniMax-H3 encoder quantization (H3-ENC-BF16-COND-DIFF) | Measured Thor (d1085374). Q4_K_M vs bf16 encoder, same 233-token prompt, same forward: rel RMS 0.0340 (0.0685 excl. sink), per-token cosine mean 0.99745 / min 0.909, median rotation 3.5° | NOT a scale change (best rescale 0.0340->0.0328). Same energy as a ONE-WORD prompt edit but DIFFUSE: 232/233 tokens rotate vs 172/233 untouched. Render A/B owed. Detail: benchmark-record |
| MXFP4 Qwen3-8B (W4A16 Marlin) | KERNEL-MARLIN-DENSE-EXEC x3 (dense-ON default): c1 1.020, c2/c4/c8 0.962/0.966/0.969, GPU mem 2.63x less (beats #51 1.005/0.925/0.939/0.953 EVERY axis); #44 3/3, 32B-NVFP4A16 6/6; -Werror test-guard fixes x2 | VT_MARLIN_DENSE default-ON (+951us). FLASH-PTXAS #82: cuModule A/B ties our+vLLM PTX across ptxas 13.0/13.2/driver-JIT (~144us); +10us is engine CONTEXT not codegen, no ptxas lever/flip (retires #75) |
Vulkan load memory (BACKEND-VULKAN-LOADMEM) | The load held the model TWICE. 27B bf16, GB10, VT_ADOPT_DEVICE_BYTES A/B: VmRSS 100.759 -> 53.413 GiB, MemAvailable floor 13.85 -> 47.3 of 119.6. Device bytes identical. Detail | Load-phase peak (host build), and the page cache that tracks copied bytes 1:1 |
Memory footprint vs declared workload (ROAD-V1-MEM, #83) | Never measured, and not measurable today: there is no auto-sizing to compare against, because the KV pool is a hand-typed --num-blocks, so “what the run actually needed” has no number | Once M1’s MemoryBudget lands: predicted-vs-actual bytes per allocation class, then peak footprint ours-auto vs vLLM at its 0.9 default on the same model and config |
Startup latency (cold to first /health) | 36.51 s vs vLLM 0.25.0’s 221.51 s = 6.07x (medians of 3, 27B-NVFP4, GB10). PROVISIONAL: 3 of 6 legs contended, repeat killed by a host reboot. Detail | Uncontended 3-rep re-run on a quiet box |
Speculation depth (ROAD-V1-D3-SPEC-K, #81) | PENDING. 27B k=4 acceptance 0.875/0.750/0.618/0.507 vs control 0.000. One-window re-run 2026-08-17: void lifted, 6/6 clear. vLLM leg blocked, the oracle OOM-reboots this box (spec ) | the vLLM leg (our-ON == vLLM-ON) on the 27B and the whole 35B lane, then the c1/c>1 A/B at matched k + the prose vs code acceptance-vs-depth curve any depth policy needs |
Vulkan vs llama.cpp Vulkan (BENCH-VK-LLAMA
) | 25 NATIVE (+8 GDN). 27B prefill 21.5x, a SELF-ratio not a llama.cpp one; decode 4.36 vs 4.35, MET, denominator SUPERSEDED (7 clean legs). Smart barriers skip 19.8%/tok, GPU -1.09 ms; e2e 8/12, unresolved. OFF. | VK-C coopmat A/B on Thor (VT_VULKAN_COOPMAT=0 A/Bs it): 11.1x-32.9x vs our UNTILED scalar kernel, not a competent GEMM. VK-E: llama.cpp -DGGML_VULKAN=ON at 237ad9b96 on dgx, SUPERSEDED, same GGUF, 3 columns |
ROCm (BACKEND-GATE-ROCM-VLLM / -SGLANG) | PENDING: no binding throughput number. The directional d=128 decode row below is not one. Runtime-green on 5 gfx archs. Gemma-3 is 48/48 exact vs two vLLM-ROCm oracles; Qwen3.5-0.8B correctness remains open | Same model, quantization, request shape and cache policy vs pinned vLLM-ROCm on one idle AMD host. Add equivalent SGLang; close correctness first (#41 ) |
ROCm d=128 decode arm (BACKEND-ROCM, #382
) | DIRECTIONAL, not binding. gfx1200, both sides in the pinned oracle container, 1024/128 c1, 8 prompts, 3 reps: TPOT 42.40 -> 11.67 ms with VT_ATTN_DECODE_D128=1; vLLM 555967922 6.68 ms, so 6.35x -> 1.75x | Harnesses differ (oracle over HTTP, ours in-process) and no same-tool per-call trace exists, so #488
stays open. Owed: decode-windowed rocprofv3 both sides |
| Gemma4 prefill-peer Finish-barrier cost (#1047 item 3) | Attribution GREEN at T=2029, not a product ship number. A (wait-only BEFORE) 1122.10 tok/s vs B (AFTER) 1094.24; +2.55% / 46.05 ms/req. Wait deletion is not authorized to land. | T=19 not run. Overlap/async retirement unmeasured. Detail: benchmark-record |
Tenstorrent Blackhole (BACKEND-TENSTORRENT) | NOT APPLICABLE (speed). Correctness: OPT-125m STRICT 6/6 e2e on real hardware. Qwen3-0.6B has a device-specific golden and short 4-token warm smoke (~0.28 tok/s), not a completed speed run | Full 16x16 Qwen3 gate, then device-resident tensors + ttnn::sdpa_decode before any performance comparison. Spec |
Mistral-7B-v0.3 on Tenstorrent (BACKEND-TENSTORRENT-MISTRAL) | PENDING (speed gate), figure reproduced: warm eager median 9.817 tok/s (b1, 32-tok greedy, n=15 under one lock, 2026-08-26, post-W2c 21fe11cf1), replacing the single-run 4.26 anecdote. No vLLM ratio exists or can (no TT backend); clock-unattributed | A binding speed gate still needs the same-binary paired method with a sampled board window; Record , spec |
Host-free decode graph (BACKEND-TENSTORRENT-HOST-FREE-FORWARD) | INVERTED, CONFIRMED REPRODUCIBLE AT CLOCK PARITY, AND DECIDED: stand-pat, #2003
closed. At a1d04bbc9 + the #2005
W2 harness (2026-08-30): default eager median 10.998 tok/s vs median 14.299 for VT_TT_HOST_FREE_DECODE=0, ratio 1.300 median / 1.294 mean (Qwen3-0.6B b1, order-alternated pairs x3, judge PASS). Every busy slice of both arms sits at exactly {1350} MHz with the cap read from FIRMWARE (limits.asic_fmax via tt-smi smbus AICLK_LIMIT_MAX), cross-checked by the new in-process pyluwen sampler at 2.7x the subprocess cadence — superseding the 2026-08-26 1.254 at 21fe11cf1 and retiring the claimed-max UNVERIFIED-pin owed. The default stays host-free: a faster opt-out is a measurement, not a corrected mechanism, so no env flip | The corrected mechanism for the opt-out arm’s ~2.5x move is the owned residual in the host-free spec
## Owed: the EnsureDevice2D hypothesis is REJECTED (the per-op delta b86e3705f..21fe11cf1 shows 353511e72/101b415d7 only ADD staging work). Captured-arm hang #1625
, #1627 TT async readback. Record |
Qwen3.5-0.8B on Tenstorrent (BACKEND-TENSTORRENT-QWEN35, #1715
) | FIRST TT numbers for the family, and they are ~100x under a 7B dense on the same board: warm 0.089 tok/s (~11.2 s/token), reproduced exactly across runs, arms and a 96-token generation (P150, single lock hold, production entry point). The host-free opt-out arm is IDENTICAL (0.080/0.089/0.089 vs 0.080/0.090/0.089), so #1604/#2003 polarity is irrelevant here. Output coherent. Unattributed by sampler; minute-scale duty cycles make cap-pegging near certain, which only sharpens the gap. Leading hypothesis: a per-layer host/reference fallback inside the hybrid-GDN stack (~400 ms/layer-token over ~28 layers) | Profile ONE eager decode step and name the dominant op BEFORE registering more kernels blindly. DONE 2026-08-27: the wall is host-side staging around the TT GEMM — vt::Tensor::Numel() 27% of samples, repeated MetalContext/UMD chip discovery ~12%, EnsureDevice2D→MatmulBTKernel 24.1% of the call graph feeding DenseMlpBlock (9.4%) and the f32 logits GEMM (9.3%); no device kernel ranks. CPU fallback control runs the same leg ~73x faster (record
, evidence docs/bench-evidence/tt-qwen35-eager-profile-20260827.log). W4 (#2107) landed 2026-08-27, levers 1+2: bulk bf16 staging (one from_span over the master’s own bytes — Numel() 27.09%→1.76%, the EnsureDevice2D→MatmulBTKernel 24.1% chain gone) + single-locked slot resolution; eager leg 28.934→16.961 s, 0.104→0.177 tok/s (+70%). The residual lead is tt-metal’s per-upload context/UMD/CQ work inside from_span→to_device (~23% of the smaller pie) and CPU-threadpool spin (19.2%). W5 (#2244) landed 2026-08-29: allocation-free per-slot persistent uploads (residual allocs 0.02%), wall honestly unmoved (−0.14%) — the residual re-attributed to the per-CQ-op tt-metal stack. W6 (#2273) closed 2026-08-29: batch per-layer staging named INEXPRESSIBLE on the pinned tt-metal (every write primitive targets one MeshBuffer; the staging fan-in is causally interleaved). W7 (#2282) landed 2026-08-30: the staging-write elimination is a measured NULL — staging counters byte-identical before/after on both production workloads (e2e ambient 0.114 writes/step in both arms; wall 1385.5 s vs 1385.9 s mean), wall unchanged; the residency state does not manufacture the writes (they are the causally required ones W6 identified), and the attempt’s gate failure caught a real latent defect (the reservation arm served stale bytes over a live device shadow), fixed by the spend/refuse rule |
Prompt logprobs (SAMPLE-PROMPT-LOGPROBS, #223) | NO number measured, claimed or owed. Correctness-only, CPU. Upstream ships this path explicitly unoptimized (gpu_model_runner.py:5622-5623); a step where no request asks is unchanged | Floor if one is ever wanted: vLLM’s own prompt_logprobs=k, same model and prompt |
logprobs_mode (SAMPLE-LOGPROB-TOKEN-IDS, #238) | NO number measured, claimed or owed. Correctness-only, CPU. One [n, vocab] device->host copy per step when a processed mode is engaged, nothing when not | Nothing to close: observation modes, not a path vLLM optimizes either |
logprob_token_ids scoring (#264) | No number owed: correctness-only, CPU-gated; inert unless set | Owed once the OpenAI field is wired |
SGLang floor arms (BACKEND-GATE-CUDA-SGLANG) | PARTIAL, not “never ran”. 27B-NVFP4 c8/c16, 3 reps, 0 errors, 2026-07-28. This cell read “Never ran” until 2026-08-16 (#979 ) | c1/c2/c4, the in-series vLLM arm, 35B, the token-ID cross-check, paired traces, and every Qwen3.8-27B point. The image path stays forbidden, and the oracle is reachable again by a different one. On 2026-08-23 two rc jobs on dgx:gpu0 installed the pinned SGLang from PyPI wheels into a virtual environment, asserted the installed tree against a committed 3338-file manifest, served Qwen3.8-27B bf16 to readiness in 454 s and completed one c1 leg with 6/6 requests, 0 errors and exactly 768 output tokens. The oracle is gateable = yes again on that route. That is reachability, not a floor: the leg is an absolute on an unpinned clock that drifted 7.59% against the 5% ceiling, no vllm.cpp arm ran beside it, and the arms listed here are still unrun (#1265
) |
llama.cpp on a CURRENT CUDA card (BACKEND-GATE-CUDA-LLAMACPP) | No run. Row filed 2026-08-16 (#979 ); this arm had no owner before | Our Q4_K_M arm (#821 ), then a same-file same-box quant-matched run |
Qwen3.8-27B three-way at each engine’s best (BENCH-QWEN38-27B-SOTA) | No run, filed 2026-08-21 (#1574 ). The competitors’ own numbers on their own checkpoint are quoted in the spec as THEIR claim | fp8 KV on CUDA gates every cell (KV-FP8 W2-W4 read later), then QUANT-QWEN38-27B-NVFP4-ARM W5. DFlash2’s mechanism landed 2026-08-20. Spec |
Embeddings on the ONE surface (ROW 6, LlamaModel + vllm_embed + /v1/embeddings) | NO number measured, claimed or owed. Correctness-gated only, CPU: the 2026-08-08 fold (engine path == direct registry path, f64 LAST+normalize reference on the committed fixture) is plumbing, no speed claim | A REAL embedding checkpoint (e5-mistral class) + a same-box vllm.LLM(task="embed") oracle; only then does an embed-throughput bar exist |
| Parakeet/FastConformer ASR (P1-P4 + ONE-SURFACE fold ROW 1) | NO number measured, claimed or owed. Correctness-gated only, CPU f32; the 2026-08-07 surface fold (vllm_transcribe, /v1/audio/transcriptions) is transcript-byte-identical plumbing, no speed claim. | Floor is parakeet.cpp, same clip and box; needs a CUDA provider and a pretrained checkpoint |
| Async serving correctness (#323) | FIXED: the decode graph replayed stale HOST token ids, degenerating concurrent requests past slot 0 (classic-dense, graph on = default). Graph declines while the mirror is live; async 7/7, SACRED 184/184 | Graph to read ids at REPLAY |
Ampere consumer (sm_86, RTX 3090 class) | No number owed; no such board here. 2026-08-06 build-verify: 7/7 FA2 TUs 0-warn, real sm_86 SASS. Detail | External RTX 3090 report. Floor is llama.cpp on that card (GGUF, not our Blackwell-only NVFP4 grid) |
Pre-Ampere breadth (Turing sm_75 / Volta sm_70 / Pascal) | No number owed; nothing runs on these arches. 2026-08-06 sm_75: 20/20 TUs PASS (0 err/warn), WMMA bodies + all 3 selectors arch-gated; GB10 SASS byte-identical. Detail | Full-library LINK at sm_75 + cuobjdump SASS, then a build-supported row. The fp16 fattn port is speed-only now; its floor when a card exists is llama.cpp on that card |
Qwen3.8-27B Q4_K_M GGUF vs llama.cpp b10451 (QUANT-QWEN38-27B-GGUF-ARM, #821
) | TOKEN GATE RAN AND FAILED, so NO speed or memory number is admissible (2026-08-23, thor:gpu0, jobs 64f66cda / 8e0d8e54). Same GGUF file both sides, greedy, 48 tokens, MTP OFF so both decode the same 851 tensors. Tokenizer EXACT 6/6; generation DIVERGES 5/6 (first index 7/34/20/-/14/32, prompt 3 exact 48/48). Teacher-forced over 288 steps our token is the oracle’s rank 1 on 282, rank 2 on 6, NEVER rank 3+, losing 0.027-0.178 logits on absolutes of 15.9-22.6 - a PRECISION difference, not a wiring defect. Near-tie band not reached for: the oracle is deterministic. Resident bytes measured only to refuse a dequant hypothesis, ours 24.997 GiB vs the oracle’s 30.917 GiB (evidence
) | A logit vector off a production path, which this tree does not expose, then a per-layer bisection against llama.cpp separating the 48 GDN layers from the 16 full-attention ones. Only then a speed grid |
Vision-tower skip host RSS (ENG-MM-INPUT-PIPELINE L3, #607
, #1358
) | PARTLY MEASURED, remeasured 2026-08-28. qwen3-vl MET half 1 on both pairs twice: 1.542 GiB freed at load on thor:gpu0 at 41ab550b9 (2026-08-24), then 0.770 GiB on dgx:gpu0 at 525d2b991 (2026-08-28) once #1359
’s Qwen3-VL half had landed — both --device cpu, Qwen3-VL-4B-Instruct (result
). The 0.499x fall between them is the predicted halving and is CORRECT, not a regression: the skip now frees only the tower the checkpoint ships, our widening having gone. STILL OPEN: half 2, the default arm within 2% of the pre-L3 edbc47ce0 binary, which neither run asserted; muse-glimmer-30b on BOTH halves against its own 90%-of-7.161-GiB threshold; and any GPU/VRAM arm, since both measured builds are CPU-only | Half 2 for qwen3-vl, then muse-glimmer-30b on a device with ~56 G of worker-local disk (thor could not spare it, dgx has 2.3 T). muse-glimmer-30b still widens and its own threshold is unchanged, blocked on #2166 |