NVIDIA FP4-Local checkpoints on DGX Spark (TP=1, TP=2): format, model and speed plan in rounds #565

개요

이 이슈는 따로 관리하던 체크리스트 이슈 세 개(#493, #494, #495)를 하나로 합친 플랜입니다. 기존 플랜을 3개 이슈로 나눠놨더니 제가 라운드 처리하기가 힘들어서 통합했습니다!!

전체 작업은 NVIDIA가 내놓은 FP4-Local 체크포인트를 bloomery가 형식 그대로 읽어 DGX Spark에서 빠르게 돌리는 일입니다.
1노드에서는 Qwen3.8-Flash-Next-FP4-Local을 돌리고 2노드에서는 GLM-5.3-Flash-FP4-Local을 tensor parallel(TP=2)로 돌립니다.

크게 3라운드로 나눠놨습니당.

• 1라운드는 베이스를 닦고 있습니다. 여기에 가장 많은 시간을 부으려고 합니당. 지금 다루는 두 모델뿐 아니라 앞으로 나올 모델도 그대로 올라탈 수 있게 가중치 형식과 공용 층을 먼저 깔고 있습니다.
  형식마다 읽기와 커널, TP 분할 단위를 한 번만 만들고 모델은 텐서 이름과 역할, shape만 알려 줍니다.

• 2라운드에서는 이 기반 위에 Qwen3.8과 GLM의 FP4-Local 모델을 올리고 두 모델의 MTP draft도 함께 붙입니다.

• 3라운드는 속도 최적화입니다. 2라운드에서 잰 수치를 기준으로 삼아 PR마다 main과 A/B로 비교합니다.

라운드 안의 PR이 모두 머지되면 main을 한 번 재고 그 라운드를 닫습니다.
이때 잰 수치가 다음 라운드의 기준이 됩니다.

형식과 공용 층을 먼저 깔아 두면 모델이 바뀌어도 형식 코드를 다시 만들 필요가 없습니다.
새 모델을 붙일 때 손댈 곳은 모델의 이름 표와 shape 정도로 줄어듭니다.
라운드가 끝날 때마다 main 수치를 이 이슈에 댓글로 남기니 단계마다 무엇이 얼마나 달라졌는지도 수치로 바로 확인할 수 있습니다.

라운드와 함께 두 갈래를 따로 둡니다. 7절 최적화는 체크포인트 폴더 하나만 받아 flag 없이 띄워도 가장 빠른 설정으로 뜨게 하는 일이고, 8절 사용성은 설치, 튜닝, 오류 안내를 다룹니다.

1. Goal

bloomery serves NVIDIA's FP4-Local checkpoints on DGX Spark with every part in the format the checkpoint ships, and the output matches a pinned reference:

• One DGX Spark (TP=1): nvidia/Qwen3.8-Flash-Next-FP4-Local, revision b9ce18a6.
• Two DGX Sparks (TP=2, tensors split across both nodes): nvidia/GLM-5.3-Flash-FP4-Local, revision 76348ddb.

There is no speed target. Each round is measured against the round before it, and section 2 gives the numbers this work starts from.

This issue replaces #493, #494 and #495. Items are ticked as they finish, and code items link their pull request.

2. Starting point on main

These numbers were measured on main before this work started.

One DGX Spark, main eb1b521f, Qwen3.8-Flash-Next from unsloth UD-Q4_K_XL 766911a6, one request, draft off, median of three rounds. Each cell is decode / prefill in tok/s:

• MTP with the shared Q8_0 draft, 1,024 tokens a request, against draft off [derived]: 1.52 to 1.66 times on code, 1.18 to 1.28 times on greedy English prose.
• A 4-row MTP verify costs 68 to 70 ms against a 33.6 ms plain step, 2.0 to 2.1 times [derived].
• The upload phase of loading takes 29.05 to 30.52 s. The drive reads 10.034 GB/s with eight direct readers.

Two DGX Sparks, main 7de31349 on the expert tier, where node 0 runs all of the attention alone, GLM-5.3-Flash from unsloth UD-Q4_K_XL a38483c8, prompt fed as steps, one request, 128 decoded tokens, draft off. With one observation each, this is a reference point, not a comparison:

Current main cannot open either FP4-Local folder. Until a round serves a checkpoint on main, its speed is compared with these GGUF numbers and labelled a comparison across files.

3. How results are measured

Each round's main numbers are posted as a comment on this issue. They are medians and percentiles over 20 requests a row. ISL/OSL is input and output tokens a request and C the requests in flight at once. Output tok/s is the total over all requests, output tok/s/user one request's rate and prefill tok/s one request's prompt rate. In the Text column, code is the Python function text, en the English prose and ko the Korean prose. A value in parentheses is the change in percent against the previous round, or against section 2 in round 1. Load time is reported apart.

A speed pull request names, before any code, the metrics it should improve and the metrics it must not worsen. Main and the head then run alternately for three rounds in one lease. A metric is better when all three paired rounds favour the head and the median gap is larger than main's own spread over its three runs. The change stays when a listed metric is better, none is worse and the greedy tokens match main; otherwise the code is dropped and the measurement stays on record. Fewer bytes moved, fewer kernel launches, less memory or less CPU time at the same tok/s also counts.

Prompts come from four pinned public texts: code as functions (Python), JavaScript, English prose and Korean prose. A pull request uses the function code and the English prose, and all four when it changes MTP, the tokenizer or prompt handling.

4. Round 1: TP=1

Weight formats and shared layers that any model uses, and Qwen3.8-Flash-Next on one DGX Spark (TP=1) from NVIDIA's NVFP4 and FP4-Local checkpoints, with its MTP drafts and its speed work. Each format's reader, kernels, batched paths and TP cut unit are written once; a model supplies only tensor names, roles and shapes. Speed work targets FP4-Local's prompt processing, time to first token, decode and MTP decode. NVFP4 follows through the routed expert, GDN, attention, hyper-connection, scheduler and concurrent-request paths.

• Safetensors model folder as one tensor source, and the quantization config read into one format per tensor (placement, model, gpu-gates: --place tp plans GLM-5.3-Flash's Hugging Face checkpoint over N ranks, each in its own node's pool #519)
• ModelOpt NVFP4 shards open with their U8 weights and F8_E4M3 scales held as stored (hf: a ModelOpt NVFP4 shard opens, its U8 weights and F8_E4M3 scales held as stored bytes #489)
• PLE table in F8_E4M3 rows with one tensor scale (gguf, engram, model, gpu: Qwen3.8's PLE table in F8_E4M3 rows with one tensor scale, beside IQ4_NL #490)
• Routed NVFP4 experts on the card in decode (card leg: a routed NVFP4 stack runs on the card, a gate and up and down _sel pair with a scale an expert #517)
• MXFP8 weights as a card format: a gemv and a wide prompt GEMM (gguf, gpu, placement: MXFP8 weights (E4M3 blocks of 32 with an E8M0 scale) in a gemv and the F32 tile, bloomery's own types named with rows #521)
• Prepare a GGUF split from a ModelOpt checkpoint folder (hf, model: prepare a GGUF split from a ModelOpt Qwen3.8 checkpoint folder #540)
• Compressed NVFP4 scales (nvfp4_csf, csf_32x2_v1) decoded at load, with FP4-Local's MXFP8 weights and NVFP4 PLE rows read as stored (hf, model: prepare reads Qwen3.8-Flash-Next-FP4-Local, its compressed expert scales decoded, its PLE table in NVFP4 rows and its MXFP8 weights as stored #543)
• The prompt GEMM admits NVFP4 routed experts (gemm, placement: the prompt GEMM admits NVFP4, its store scaling each slot by its expert's tensor scale #545)
• Qwen3.8 served from NVIDIA's NVFP4 checkpoint on one card, against its pinned reference (gpu, model, refset: Qwen3.8 served from the NVIDIA NVFP4 checkpoint on one card, every routed expert on the card, against its pinned oracle #561)
• TP placement over N ranks, each rank in its own node's memory (placement, model, gpu-gates: --place tp plans GLM-5.3-Flash's Hugging Face checkpoint over N ranks, each in its own node's pool #519)
• All-gather of partial rows among N ranks over the existing link (gpu, gpu-gates: the link's all-gather of f32 row spans among N ranks, and link_probe --gather #501)
• FP8 weights with 128 by 128 block scales as a card format: a gemv and a wide GEMM (quant_config: an FP8 weight of one scale a 128 by 128 block is held as FP8_B128, an f32 scale a 128 values of a row #562)
• The TP rank body, GLM first: each rank's share of the model at rank shapes, checked on one card over the loopback link (gate_glm5next_tp: the rank body on the fixture at one, two and four ranks over a loopback #589)
• BF16 weights as shipped on a dense site and its wide GEMM, beside Q8_0 (gpu, placement, model: BF16 dense weights in their words, picked per tensor (Qwen3.8) #592)
• A checkpoint prepared with every part in its shipped format by default: BF16 stays BF16; the Q8_0 rewrite stays as a named opt-in lever with its measured KL (prepare, levers: a checkpoint's BF16 matrices stay BF16 by default; their Q8_0 rewrite behind BLOOMERY_PREPARE_Q8 #607)
• Per-rank slices from a safetensors folder: scales decoded first, each format cut on its own block boundary (gguf, placement, hf, model, model-hf, gpu, gpu-glm5next: a TP rank reads its parts from a GGUF file or a checkpoint folder #618)
• The other FP8 E4M3 scale layouts on the same FP8 kernels: one scale per output channel and one per tensor (hf: FP8 weights of one scale a row, a tensor or a 32 by 32 E8M0 block read onto FP8_B128 and MXFP8 #579)
• One card type table: card columns on each type's QROWS row (crates/gguf/src/qtype.rs), which the card kernels and admission read instead of their own lists (gguf, placement, gpu: a type's card columns in its QROWS row, read by every card operation and layout #588)
• bloomery's private value types named in GgmlType with their row sizes, so they pass the same checks as GGUF types (gguf, gpu, placement: MXFP8 weights (E4M3 blocks of 32 with an E8M0 scale) in a gemv and the F32 tile, bloomery's own types named with rows #521, hf, model: prepare reads Qwen3.8-Flash-Next-FP4-Local, its compressed expert scales decoded, its PLE table in NVFP4 rows and its MXFP8 weights as stored #543)
• A safetensors checkpoint folder served directly on one DGX Spark, NVFP4 and FP4-Local as downloaded, with no prepare step (bloomery-serve, hf, gpu, model-hf: -m <folder> and --hf serve a Qwen3.8 NVFP4 checkpoint folder as downloaded, with no prepare step #635)
• MTP draft state after a prime: fix the draft-state defect before the MTP drafts ship (app, gpu, gpu-gates: a window no verify takes leaves no stale draft state #631)
• Qwen3.8 FP4-Local's MXFP8 attention, indexer and shared-expert weights accepted by the Qwen3.8 placement, so the folder loads as shipped (qwen38: read dense coverage from the by-name sites #649)
• Qwen3.8-Flash-Next-FP4-Local served on one DGX Spark, every part as shipped, against its pinned reference (qwen38: read NVFP4 PLE planes from a checkpoint folder #660)
• Prompt GEMM on tensor cores for MXFP8, FP8 and BF16 weights, with the weights read in their stored format (gemm: BF16 tensor cores for GB10 dense prompt products #668)
• BF16 dense calls with 2 to 8 rows as one multi-column gemv, for several requests and MTP verify (gpu: BF16 sites run small batches through the multi-column gemv #652)
• Qwen3.8 hyper-connection prompt projections on the shared BF16 tensor-core path (gpu: HC prompt projections use the shared BF16 tensor-core dispatch #693)
• The dense BF16 prompt tensor-core GEMM runs a pipelined 128 by 128 tile (gemm: pipeline the dense MMA tile on BF16 stages read by ldmatrix #696)
• A new request zeroes its recurrent state on the card instead of copying 469 MB of host zeros: 16.807 ms to 2.416 ms on a DGX Spark (gpu: zero recurrent stores on the device at reset #643)
• The MTP draft width weighs the batch, so drafting is not slower than draft off with several requests (runtime: learn batch acceptance from the current request #690)
• MTP draft from the FP4-Local checkpoint: W4A16_NVFP4 experts and BF16 dense parts (qwen38: load the checkpoint folder MTP draft in its stored formats #682)
• An unset MTP draft stays off on a server with two or more slots, so concurrent sampled requests keep their throughput (qwen38: keep the unset MTP draft off under --parallel above 1 #695)
• Several slots' sampled steps on a drafted load run together in one batch (qwen38: a drafted load runs a round of several slots' plain steps as one pass #703)
• MTP draft from the NVFP4 checkpoint: FP8 block experts and BF16 dense parts (qwen38: load the checkpoint folder MTP draft in its stored formats #682)
• Shared expert on a side stream beside the routed experts, on Qwen3.8 (qwen38: run the shared expert beside the card leg #686)
• GDN projections on a side stream beside the routed experts, on Qwen3.8 (qwen38: run the delta rule's z and the attention's q on the side stream #697)
• MTP verify reads each routed expert once for all its rows: an NVFP4 run of several tokens dots each selected expert's rows for all its slots in one warp (gpu: an NVFP4 run of several tokens dots each selected expert's rows for all its slots in one warp #808)
• NVFP4 prompt GEMM on the GB10's FP4 tensor cores, with its measured KL (4-bit activations); it shipped as the default path, inside the existing oracle bands (gpu: run routed NVFP4 prompt GEMM on FP4 tensor cores #687)
• Prompt GEMM epilogues fold in F32 and recompute a tile exactly only when it leaves range (gemm: fold the GEMM epilogues in F32 and redo a tile exactly only when it leaves range #700)
• GDN delta rule on tensor cores in prefill, with the Qwen3.8 prompt router, hyper-connection bottleneck and gate projection on BF16 tensor cores and the attention gate folded into the flash store (gpu: run Qwen3.8 prompt layers on tensor cores and fold the attention gate into the flash store #812); the chunked delta products keep both BF16 parts of each operand, so the greedy token at prompt position 4,128 matches the f32 reference again (gpu: retain both BF16 parts in the chunked delta products #816)
• Qwen3.8's prompt card sum folded into the combine on layers without a host leg (gpu: fold Qwen3.8's prompt card sum into the combine on layers without a host leg #701)
• No host union on layers without host experts: skip it in prefill (gpu, model, refset: Qwen3.8 served from the NVIDIA NVFP4 checkpoint on one card, every routed expert on the card, against its pinned oracle #561)
• MTP draft stops at the chosen width instead of walking its full width (width: a chain draft walks only the width the pass plans to verify #705)
• A request no longer waits behind a synchronous cache save when the staging pool is fragmented (serve: keep a finished request's cache save off the next request's path #706)
• Qwen3.8 prompt routes a whole ubatch in one run, so each routed expert is read once (qwen38: route a whole ubatch in one run so each routed expert is read once #740)
• Multi-column F32 and BF16 row walks load two columns at once, so NVFP4 MTP verify rows cost less (gpu: load a pair of columns' chunks of a multi-column F32 row walk before their multiply-adds #707)
• Qwen3.8 router and shared gate read as the checkpoint's BF16 instead of widened F32 (qwen38: keep the router in its BF16 words #732)
• Plain MTP slots walk their waiting draft rows in one store walk, and a round without slot 0 keeps slot 0's rows in place (mtp: walk every plain slot's waiting draft rows in one store walk #735)
• A request with ignore_eos drafts with MTP like any other: banned ids are masked in the drafted rows instead of turning the draft off (serve: a request that bans an id draws in the draft's passes #719)
• Qwen3.8 prompt attention selects and attends a ubatch's selecting rows in one launch each (gpu: select and attend a ubatch's selecting rows in one launch each #753)
• The dense BF16 prompt MMA loops its copied stages apart and widens its raster (gemm: loop the copied dense MMA stages apart and widen the BF16 raster #754)
• Request rows read back into host memory the host wrote first: a 32 MiB readback 12.925 ms to 1.761 ms on a GB10 (gpu, os: read request rows back into host memory the host wrote first #755)
• The fused hyper-connection mix reads whole sectors and prefetches its streams into L2: HC mix 273.0 ms to 145.0 ms per 4,096-token prompt (hc_gated: read the fused HC mix in whole sectors and prefetch its streams into L2 #756)
• GB10 builds select sm_121a by default, so the routed prompt GEMM for NVFP4 and FP4-Local runs in FP4 on the tensor cores (gpu: select native FP4 prompt builds on GB10 #758)
• Qwen3.8 prompt HC reads the normed streams once for an inject site's down and inject on GB10: down and inject GPU time 177.9 ms to about 100.6 ms per 4,096-token prompt on FP4-Local (hc_gated: read the normed streams once for an inject site's down and inject on GB10 #759)
• Routed NVFP4 expert GEMMs on the W4A8 path copy weights ahead and fold the SwiGLU quantization into the up GEMM: NVFP4 4K TTFT -5.8% with sm_121 builds (gemm: speed up the routed NVFP4 expert GEMMs on long prompts #760)
• The BF16 dense prompt MMA copies activations rounded once instead of rounding them in every tile: BF16 dense GPU time 503.9 ms to 418.7 ms per 4,096-token NVFP4 prompt, TTFT about 92 ms lower (gemm: copy the BF16 prompt sites' activations rounded once #761)
• The routers' prompt logits run on the shared F32 tile with two lane sums a thread: router logits kernel 1.022 ms to 0.708 ms a launch on a 4,096-token prompt (gpu: sum the F32 tile two lanes a thread and run the routers' prompt logits on it #763)
• A Qwen3.8 prompt call runs one cut across its checkpoints and reads its PLE rows under the layers (qwen38: a prompt call runs one cut across its checkpoints and reads its PLE rows under the layers #724)
• Qwen3.8's MTP store walk projects and appends the keys and values alone (gpu, gpu-gates: Qwen3.8's MTP store walk projects and appends the keys and values alone #765)
• The routed NVFP4 gate, up and SwiGLU run as one native FP4 launch (gpu: run the routed NVFP4 gate, up and SwiGLU as one native launch #766)
• MXFP8 prompt sites run W8A8 on activations prepared once (gemm: run MXFP8 prompt sites W8A8 on activations prepared once #769)
• A checkpoint folder's parts load in parallel within the host load budget (gpu, model-hf, placement: read a checkpoint folder's parts in parallel within the host load budget #787), sharing the serial packing reserve across a part read run (gpu: share the serial packing reserve across a part read run #809)
• Weight and activation formats on the GB10 tensor cores: W4A4 NVFP4 prompt GEMM (gpu: run routed NVFP4 prompt GEMM on FP4 tensor cores #687), W4A8 routed NVFP4 experts (gemm: speed up the routed NVFP4 expert GEMMs on long prompts #760), W8A8 MXFP8 prompt sites (gemm: run MXFP8 prompt sites W8A8 on activations prepared once #769), W4A16 NVFP4 MTP draft experts (qwen38: load the checkpoint folder MTP draft in its stored formats #682)

Remaining Round 1 items

• The last prompt token runs as a captured graph step instead of a one-row eager GEMM, and PLE rows are read in split parts (serve: skip unused Qwen3.8 prompt heads and read PLE rows in parts #817)
• The MXFP8 gate projection finishes with the shared head norm and writes its output form in the norm epilogue (gpu: finish the MXFP8 gate projection and its output form in one epilogue #819)
• The FP4-Local routed experts on the card route keep their down rows in BF16 and the block sum widens: prefill 2,198.9 to 2,305.7 tok/s (+4.9%) on a 4,096-token prompt, draft off (qwen38: a load whose every layer runs hostless over native NVFP4 downs keeps its card route's down rows in BF16, the product rounding each to nearest even and the block's card sum widening it #914)
• Qwen3.8 prompt HC streams and normed rows walk as BF16 values, each rounded once: prefill 2,296.0 to 2,481.5 tok/s (+8.1%) and TTFT p95 1,790.2 to 1,657.1 ms on a 4,096-token FP4-Local prompt (qwen38: a ubatch walk carries its HC streams and normed rows as BF16 values, each rounded to nearest even once #920)
• Prefill kernels folded into their producers: activation quantize, accumulate, SwiGLU and shared-expert add (gpu: fold Qwen3.8's prompt card sum into the combine on layers without a host leg #701, gemm: speed up the routed NVFP4 expert GEMMs on long prompts #760, gpu: run the routed NVFP4 gate, up and SwiGLU as one native launch #766, gemm: run MXFP8 prompt sites W8A8 on activations prepared once #769, gpu: fuse the hostless Qwen3.8 card sum with the shared add #974, gpu: the fused NVFP4 gate and up writes its SwiGLU in native FP4, and qwen38 counts a folder's NVFP4 rows in that form #1040)
• The routed NVFP4 prompt GEMM copies its operands through a three-stage pipeline in 128-value K steps (gpu: speed up Qwen3.8 prompt prefill in the routed NVFP4 walk and the HC mix #822)
• An MTP verify launches before its PLE rows arrive: the draft ids' PLE rows are read on the step thread and wait at their PLE site, not before the launch (qwen38: launch a verify before its PLE rows arrive #877)
• MTP verify decodes E4M3 in pairs, stages four-column MXFP8 activations once a block, and replays its commit in one launch (gpu: Qwen3.8's verify decodes E4M3 in pairs, stages four-column MXFP8 activations and replays its commit in one launch #884)
• Hyper-connection kernels fuse norm and mix (hc_gated: read the fused HC mix in whole sectors and prefetch its streams into L2 #756, hc_gated: read the normed streams once for an inject site's down and inject on GB10 #759, gpu: walk the native NVFP4 route tile, the HC norm and the prefill flash's query staging with fewer loads #847, gpu: the HC up-mix's passes split the warp's columns, so a mix step reads a whole sector of each stream and its gate stores fall in 32 banks #988)
• The hyper-connection multi-row path reads at the single-row rate
• MTP verify gemv entries on GB10 prefetch their weight rows into L2
• Prompt cache checkpoints at message ends with the full state, read back at the copy rate when the next message reuses them, on one DGX Spark (serve: keep the Qwen3 shared system prompt at message starts #694, gpu: Qwen3.8 slots share a prompt prefix's KV pages and copy them on write #1012, serve: Qwen3.8 slots save to a file and restore from it, the MTP draft's side with them #1027)

Dropped from round 1, not planned now:

• Routed NVFP4 scales kept compressed on the card and decoded inside the expert kernels, for both models
• Long-context decode reads the KV cache at the memory rate: decode attention at one query row, so 64K and 128K decode stay close to the 4K rate (an FP8 KV cache only if a measurement shows it helps)
• Programmatic dependent launch in the decode graph (lowest priority, only if idle time is above two percent)

5. Round 2: TP=2

GLM-5.3-Flash-FP4-Local on two DGX Sparks (TP=2) and its MTP layer, on the round 1 base.

• GLM card paths for the FP4-Local formats: NVFP4 experts, MXFP8, FP8 block and BF16 projections
• GLM-5.3-Flash-FP4-Local loaded on TP=2 from the folder and checked against its reference
• TP=2 decode on two DGX Sparks
• TP=2 prompt batches
• The GLM server on TP=2
• GLM beyond 16,384 positions: run the full context on two Sparks
• NextN (MTP) layer of the GLM FP4-Local checkpoint on TP=2, MTP on against MTP off

6. Round 3: optimization and speed

Optimization against the round 2 numbers, one pull request at a time, each with its own A/B under section 3. The round 2 measurement sets the order. Candidates:

• GLM verify past two rows
• MXFP4 routed prompt GEMM on the GB10's block-scaled tensor cores
• GLM FP4-Local's BF16 projections written as Q8_0 per rank slice, as an opt-in lever with its measured KL; the default stays BF16
• Early send of TP partial rows: start the network send while the producing GEMM still runs
• TP=2 gather wait in one shot on the GB10
• L2 prefetch of the next weights while a gather waits
• Prompt-size gathers striped over both ports of the cable
• Prompt cache on TP=2, each node saving its own half of the state

7. Optimization: the default setup

One folder and one command: bloomery-serve -m <FP4-Local folder> starts the fastest setup the checkpoint allows, with nothing else to download. A flag the user gives always wins, and the server says at start what it chose.

• Qwen3.8 FP4-Local starts at its fastest from the folder alone: the checkpoint's own MTP draft runs at least as fast as the separate Q8_0 GGUF draft, the draft is opened once at load, a context that fits is chosen when --ctx-size is not given, drafting is on where it is faster, and a startup line names each choice and the flag that changes it (qwen38: a checkpoint folder's plan holds its load's least part reads #898)
• A q8_0 KV cache for Qwen3.8 with llama-server's --cache-type-k q8_0 meaning, chosen by the FP4-Local setup when no flag is given, behind an accuracy gate (qwen38: a q8_0 KV cache, set by --cache-type-k and the default on a checkpoint folder #956)

8. Usability

A DGX Spark user installs bloomery and runs a model without building it, tunes it with the flags llama-server users know, and gets an error that names the cause.

• A Linux aarch64 GB10 release archive and installer, so a DGX Spark installs without a toolchain (publishing the release stays with the maintainers) (release, install: a GB10 (DGX Spark) installs a prebuilt linux-aarch64-cuda-sm121a archive #878)
• The Dockerfile builds a linux/arm64 image from the release's GB10 archive (Dockerfile: linux/arm64 takes the release's GB10 archive (SHA256_ARM64), linux/amd64 the x86-64 one as before #882)
• The image workflow publishes the linux/arm64 image when the release carries the GB10 archive (Docker image: linux/arm64 from the release's GB10 archive where the release carries one #883)
• Qwen3.8 takes --ubatch-size/-ub like the other models (qwen38: the seat takes -ub/--ubatch-size over BLOOMERY_QWEN3_UBATCH, one size for its plans and its load #881)
• Help and option errors come before any --hf download or GPU start (bloomery-serve: the seat's flags and levers are checked before the driver is asked and the model fetched #912)
• --levers shows whether each lever that is set applies to this binary and model (levers: --levers says whether each lever applies to the binary and seat; a lever read in place is checked by its kind at main, and one the binary does not read is named as not applied #890)
• The FP4-Local server's time from start to ready is measured and its largest parts cut: a cold start to health 76.34 s to 49.38 s (gpu: a parts load stages each part as it and the parts before it are read, so a run's staging overlaps its reads #913), then 46.13 s to 36.75 s (gpu: a parts load reads each part once the parts before it not yet staged leave it room, so the drive keeps reading across the whole load #919) and 36.94 s to 27.69 s (model-hf: an NVFP4 matrix read whole bounds its read by what the read holds #923)
• The user docs (README, serving, BUILD, the release README) state the defaults the code has: slots, context per slot, the flag table, environment precedence and unknown request fields (docs: one settings page: each seat's flags, which setting wins, what a request field does; slots and per-slot context on every entry page; BUILD.md's card pick, servers and GLM model flag match the code #949)
• --hf takes a revision (hf: --hf takes a revision, <repo>@<rev>[:<quant>]; its listing, files and cache directory are that revision's #963)

9. Rules and related work

Round rules

PRs inside a round merge in a fixed order, and PRs that share files rebase once onto the one before. A round ends when every PR in it is merged and main is measured once. That measurement is the next round's reference, so the next round starts after it.

Speed PRs merged outside the rounds

• Lookup (n-gram) draft on Qwen3.8 (serve seats, generate_qwen3moe: the n-gram lookup draft on Qwen3.8 through the seats' one lookup owner #546): MTP serve decode on one DGX Spark, median of 3, 28.15 tok/s on main against 28.747 tok/s, +2.12 %.
• Dense Q8_0 prompt GEMM at the int8 tensor-core rate on GB10 (gpu: gemm_q8_0p_slab reads the Q8_0 planes a row's step at a time for a full card of tiles that share their slabs #549): prefill at 2,048 prompt tokens on one DGX Spark, median of 4, 1188.63 tok/s on main against 1296.56 tok/s, +9.1 %.
• The m-column F32 walk of the MTP verify (f32 gemv: the m-column walk keeps eight chunks' loads in flight a column #511): MTP decode at 4,096 tokens +7.2 % on code and +7.8 % on English prose against main e8abdc89, with decoded ids equal to main's.

Accuracy and related work

• Formats: every decoder is compared byte for byte with a pinned Python reference written from the format's published reference code, and a structural error is refused by name.
• Models: the reference is a Hugging Face Transformers forward of each checkpoint in f32, with every weight dequantized by its own rule and input_scale ignored. The served model is judged by the existing numerical gates with unchanged bands.
• TP: all ranks are bit-equal at every step, and the TP body run on one node matches the existing body bit for bit.
• A pull request that changes shared code shows with a test that the existing card, host and NVMe paths behave as before.

10. Other pull requests from this work

Fixes and tool changes that came up while doing the work above. None of them is a plan item, and each one stands on its own.

Engine:

• An --hf vision projector outside the plan stays only when its media plan is free (vision: an --hf projector off the plan stays only when its media plan is free #485)
• A tier card on another host plans in its own node's memory pool (placement, gpu-gates: a tier card on another host plans in its own node's pool #503)
• Qwen3.8's placement is workstation::placed_on with the route scratch on top (model: Qwen3.8's machine is workstation::placed_on with the route scratch on top #512)
• Windows over raw memory give back their context handle on every path (gpu, gpu-deepseek41, gpu-gates: windows over raw memory give back their context handle on every path #513)
• A misaligned vector view is refused at launch (gpu: refuse misaligned vector views at launch #520)
• Q8 quantizers refuse non-finite values (quant: refuse non-finite values in Q8 quantizers #597)
• A dense call on a resident weight goes through one site by name, and the card type table picks its launch (gpu: a dense call on a resident weight goes through one site by name, and the card type table picks its launch #639)
• The server's start screen: the wordmark and furnace over a placement panel and a serving panel (serve screen: the wordmark and furnace in half-block pixels over a placement panel and a serving panel #644)
• The expert stream's split pick depends only on the prompt and the plan, so greedy tokens repeat across restarts (gpu: a split pick is a function of the prompt: the expert stream prices its lane by where the probe's experts sit (lane_rate), its probe's timing printed beside it, and an admit's return by the walk's last card layer alone; gate_nvtier's sittings clause loads the paged plan twice, once beside a host copy load #647)
• An opt-in NVMe expert arena victim order by routing count, BLOOMERY_NVTIER_EVICT=heat: 6.58 % fewer decode NVMe fills per step on an RTX 5090 (levers, gpu, model, gpu-gates: BLOOMERY_NVTIER_EVICT=heat takes the NVMe arena's coldest routing count (tierengine S6) #650)
• GLM-5.3-Flash's card leg runs its routed stacks through the shared arm table, with NVFP4 in its planes and both down forms, and fully card-resident layers run without host views (placement, model, gpu, gpu-glm5next, gpu-gates: glm5next's card leg runs its routed stacks by its arms, NVFP4 in its planes and the 32-value downs included; the leg's buffers hold the forms its layers' downs read #651)
• The first Qwen3.8 request makes no pinned allocations, and per-layer embedding page advice is merged into page runs (engram: coalesce row advice into distinct page runs #661)
• Independent BF16 output heads are normalized to MXFP8 (gpu: normalize independent BF16 output heads to MXFP8 #665)
• Q8 activation bytes are counted by their one owner in fit and scratch accounting (gpu: reuse the Q8Act byte owners in fit and scratch accounting #667)
• NeoX rope appends refuse two writers on the same cache row (gpu: require distinct cache writers for NeoX appends #669)
• The server flushes each completed prompt before it starts the next call (serve: flush each completed prompt before the next call #671)
• The long-prompt yield tests keep their own doc comments, and the serve worker pumps a request in one place (serve: hold a request's next call through one owner #689)
• The Qwen3.8 MTP draft store reads only final hidden rows (gpu, gpu-gates: Qwen3.8's MTP draft store reads only final hidden rows #691)
• Qwen3 keeps a shared system prompt for the next request by marking message starts as reuse points (serve: keep the Qwen3 shared system prompt at message starts #694)
• prepare publishes a split only after every shard is written and checked, so an interrupted or concurrent run leaves the old split whole (prepare: publish a validated split without replacing existing shards #708)
• F8 and E4M3 rows whose scaled values overflow F32 are refused like NaN bytes (gguf: refuse F8 rows whose scaled values are not finite #710)
• gemm: replay only the dense MMA cells of an exceptional row or column (gemm: replay only the dense MMA cells of an exceptional row or column #711)
• qwen38: serve paged plans at their requested slot count (qwen38: serve paged plans at their requested slot count #713)
• gpu: count named file formats in weight censuses (gpu: count named file formats in weight censuses #715)
• placement: cap the NVMe tier room at the unified pool (placement: cap the NVMe tier room at the unified pool #716)
• serve: show the start screen's tier row from the load record and share the count with the bar (serve: show the start screen's tier row from the load record and share the count with the bar #721)
• model-hf: a GLM folder's target load no longer plans the NextN tensors it does not use (model-hf: a GLM folder's target load no longer plans the NextN tensors it does not use #727)
• serve: send plain steps before the round's drafted passes (serve: send plain steps before the round's drafted passes #731)
• serve: name an empty numeric field of the start screen as unread (serve: name an empty numeric field of the start screen as unread #733)
• serve, gpu: keep prompt-cache copies and reusable prefixes within their contracts (serve, gpu: keep prompt-cache copies and reusable prefixes within their contracts #737)
• hf: refuse unsupported checkpoint sources before fetching files (hf: refuse unsupported checkpoint sources before fetching files #738)
• gpu, glm5next: a route trace's load views only the layers the host computes (gpu, glm5next: a route trace's load views only the layers the host computes #739)
• runtime: retain batch costs by slot count and pass packing (runtime: retain batch costs by slot count and pass packing #741)
• serve: bound response and decoder waits and preserve half-closed clients (serve: bound response and decoder waits and preserve half-closed clients #742)
• serve screen: reflow the start screen at the terminal width (serve screen: reflow the start screen at the terminal width #743)
• runtime, gpu, app, glm5next: bound slot passes by rows and sequence bindings (runtime, gpu, app, glm5next: bound slot passes by rows and sequence bindings #744)
• serve: observe drafted overlap in engine-call order (serve: observe drafted overlap in engine-call order #745)
• gpu, model, gpu-gates: keep host tier reads protected through copies and page drops (gpu, model, gpu-gates: keep host tier reads protected through copies and page drops #746)
• serve: keep driver calls away from the process exit's teardown (serve: keep driver calls away from the process exit's teardown #747)
• qwen38: drafted-plan headroom counts the stage card's reserves (qwen38: drafted-plan headroom counts the stage card's reserves #793)
• placement: validate card layouts only for uploads (placement: validate card layouts only for uploads #801)
• gpu: group native FP4 launcher arguments (gpu: group native FP4 launcher arguments #815)
• placement: Qwen3.8's restated stage counts the rope rows its load copies to the card (placement: Qwen3.8's restated stage counts the rope rows its load copies to the card #821)
• threads: a mask that leaves the L3 groups unequal pins every allowed cpu before any shares one (threads: a mask that leaves the L3 groups unequal pins every allowed cpu before any shares one #824)
• threads: the default pool takes every allowed physical core before an SMT sibling (threads: the default pool takes every allowed physical core before an SMT sibling #826)
• qwen38: read the next step's PLE rows ahead and keep the command receiver awake between steps (qwen38: read the next step's PLE rows ahead and keep the command receiver awake between steps #831)
• width: the MTP width gate rates recent windows, and a width it left comes back once it pays (width: the MTP width gate rates recent windows, and a width it left comes back once it pays #846)
• gpu: walk the native NVFP4 route tile, the HC norm and the prefill flash's query staging with fewer loads (gpu: walk the native NVFP4 route tile, the HC norm and the prefill flash's query staging with fewer loads #847)
• gpu: shrink a part read run to the live host reading instead of refusing (gpu: shrink a part read run to the live host reading instead of refusing #870)
• placement, gpu: a host-room refusal names the reading that gave the room, so a BLOOMERY_HOST_ROOM value reads as given (placement, gpu: a host-room refusal names the reading that gave the room, so a BLOOMERY_HOST_ROOM value reads as given #885)
• model: count the Qwen3.8 ubatch arena's rows by the arena's own allocation (model: count the Qwen3.8 ubatch arena's rows by the arena's own allocation #889)
• serve: the engine's last call prints each slot's last mtp width record before the tier nodes close (serve: the engine's last call prints each slot's last mtp width record before the tier nodes close #891)
• placement: an unset context takes the largest step under the model's limit that the machine holds (placement: an unset context takes the largest step under the model's limit that the machine holds #906)

Docs:

• The prefetch module docs count one WILLNEED call per merged page run (docs(engram): describe WILLNEED cost per merged page run in prefetch #679)
• docs: a fixture variant's card budget is planned on its own file; a stale one is made again by fixture-requant (docs: a fixture variant's card budget is planned on its own file; a stale one is made again by fixture-requant #800)
• docs: dgx-spark.md states the port at 0.3.10: the sm_121a build, Qwen3.8 FP4-Local on one Spark, tensor parallel's plan and rank body; the documents that cite it follow (docs: dgx-spark.md states the port at 0.3.10: the sm_121a build, Qwen3.8 FP4-Local on one Spark, tensor parallel's plan and rank body; the documents that cite it follow #873)
• docs: README opens with install and a first request; a GB10 installs the release's archive or arm64 image from the release after 0.3.10, from source until then; the docs it points to match (docs: README opens with install and a first request; a GB10 installs the release's archive or arm64 image from the release after 0.3.10, from source until then; the docs it points to match #892)
• docs: re-measure the DGX Spark and RTX 5090 numbers at 0.3.10 (docs: re-measure the DGX Spark and RTX 5090 numbers at 0.3.10 #901)

Gates and tools:

• The nsys-ds41 self-test's stub tree carries gate-bound.sh (nsys-ds41: the self-test's stub tree carries gate-bound.sh #482)
• lds-scan fails when a line count fails (lds-scan: a failed line count fails the scan #483)
• ptx-spill-check fails when a count or a comparison cannot be read (ptx-spill-check: a failed count or comparison read fails the binary #486)
• The Qwen3-VL tower and V4.1 vision gates refuse a set without its full-tap image by name (gpu-gates: the Qwen3-VL tower and V4.1 vision gates refuse a set without their full-tap image by name #515)
• mac-static keeps each full run's results in its own directory (mac-static: keep each full run's results in its own directory #522)
• Test isolation: sj-vs-clef gets its own data directory, and gate-bound's FIFO self-test and executable fixtures are written in a child (tests: sj-vs-clef's own data directory, gate-bound's FIFO self-test and executable fixtures written in a child, so main's PR check passes #557)
• tools/box.sh checks the cargo-oxide pin against the executable Cargo runs (tools/box.sh: check the cargo-oxide pin against the executable Cargo runs #560)
• Stale recorded fixture budgets are refused before planning (gpu-gates: refuse stale recorded fixture budgets before planning #606)
• gate_nvtier loads the all-RAM arm before the paged arm and declares the arm's real host need (gate_nvtier: load the all-RAM arm before the paged arm, and declare its real host need #633)
• The engram gate checks prefetch advice on its own, apart from the row reads (engram: verify prefetch advice independently of row reads #670)
• contrib-check t0 runs the Python self-tests once, inside check-recipes (contrib-check: run the Python self-tests once, inside check-recipes #681)
• qdot imports its x86_64-only quant helpers under the x86_64 cfg, so an aarch64 just lint has no unused-import warning (lint: qdot's activation helpers imported only where the AVX2 paths use them #688)
• The nightly runs the recipes self-test once, through check-recipes.sh (nightly: drop the recipes self-test unit that check-recipes already runs #692)
• The GEMM gates' host transcriptions carry the kernels' exceptional range, so a correct finite result is no longer judged wrong (gpu-gates: carry the GEMM contracts' exceptional range through the host transcriptions #704)
• Every PTX shape table is checked for missing kernels, and a filtered spill check compares only the pins inside its filter (tools: check every PTX shape table for missing entries and filtered pins #709)
• serve: fix test helpers that miss the Cargo binary, leak children and corrupt UTF-8 (serve: fix test helpers that miss the Cargo binary, leak children and corrupt UTF-8 #714)
• tools: cover a failed lane cleanup in the batch self-test (tools: cover a failed lane cleanup in the batch self-test #717)
• tests: isolate CCD pools and quantization refusals (tests: isolate CCD pools and quantization refusals #718)
• gates: preserve the requested host room in Qwen node measurements (gates: preserve the requested host room in Qwen node measurements #720)
• gpu-gates: preserve whole GLM pass wall in throughput (gpu-gates: preserve whole GLM pass wall in throughput #723)
• tools: resolve replay lists before entering arm directories (tools: resolve replay lists before entering arm directories #725)
• gpu-gates: pin the MXFP8 select refusal by kind and tensor (gpu-gates: pin the MXFP8 select refusal by kind and tensor #728)
• tools: count a comparison binary's faults from its post-load prompt (tools: count a comparison binary's faults from its post-load prompt #729)
• gpu: cover pageable snapshot copies and resumed token bits (gpu: cover pageable snapshot copies and resumed token bits #730)
• tools: resolve model profiles and keep failed-arm logs unique (tools: resolve model profiles and keep failed-arm logs unique #736)
• serve: wait for the HTTP reset before advancing the gone-client test (serve: wait for the HTTP reset before advancing the gone-client test #751)
• tools: share t0 test setup and optimize placement tests (tools: share t0 test setup and optimize placement tests #752)
• gpu-gates: the MTP gate's store-behind cuts keep the point they cut to (gpu-gates: the MTP gate's store-behind cuts keep the point they cut to #776)
• recipes: compare imported analysis code by code equality, not marshal output (recipes: compare imported analysis code by code equality, not marshal output #778)
• gpu-gate self-test: the lib-test stub exits at once on TERM (gpu-gate self-test: the lib-test stub exits at once on TERM #780)
• static-sched: time jobs before spawning their processes (static-sched: time jobs before spawning their processes #784)
• gpu-gates, model: plan the fixture tier's card budget beside what its card sets aside (gpu-gates, model: plan the fixture tier's card budget beside what its card sets aside #785)
• recipes: fork independent key self-test sections (recipes: fork independent key self-test sections #789)
• tests, tokenizer and probe: close ten open review follow-ups (tests, tokenizer and probe: close ten open review follow-ups #790)
• recipes: fork global-input key self-tests (recipes: fork global-input key self-tests #794)
• qdot: NVFP4_OFFSET is x86_64-only like its readers, so other targets build without dead_code (qdot: NVFP4_OFFSET is x86_64-only like its readers, so other targets build without dead_code #798)
• tools/lock-back.sh: do nothing under BLOOMERY_BOX_LOCAL=1 (tools/lock-back.sh: do nothing under BLOOMERY_BOX_LOCAL=1 #804)
• gpu, placement: test CPU tier heat and larger host pages (gpu, placement: test CPU tier heat and larger host pages #814)
• ptx-shapes: gemm_iq2s's cc 12.1 pins from the GB10 scan (ptx-shapes: gemm_iq2s's cc 12.1 pins from the GB10 scan #829)
• ptx-shapes: the gemm_nvfp4p walk-less scan holds 32 spill and 16 JIT local bytes (ptx-shapes: the gemm_nvfp4p walk-less scan holds 32 spill and 16 JIT local bytes #833)
• ptx-shapes: the eight cc 12.1 entries the 9df9e7ce backend moved, at the GB10 scan's spill and JIT local (ptx-shapes: the eight cc 12.1 entries the 9df9e7ce backend moved, at the GB10 scan's spill and JIT local #839)
• ptx-shapes: the twelve sm_86 cc 12.1 entries the 9df9e7ce backend moved, at the GB10 scan's spill and JIT local (ptx-shapes: the twelve sm_86 cc 12.1 entries the 9df9e7ce backend moved, at the GB10 scan's spill and JIT local #845)
• ptx-shapes: gqa_prefill_flash_g7_q8 at the GB10 scan's spill and JIT local on sm_121a and sm_121 (ptx-shapes: gqa_prefill_flash_g7_q8 at the GB10 scan's spill and JIT local on sm_121a and sm_121 #871)
• recipes: reuse stable source hashes during reach validation (recipes: reuse stable source hashes during reach validation #875)
• serve: d opens a live view of slots, placement, requests, records and the connect lines (serve: d opens a live view of slots, placement, requests, records and the connect lines #939)
• placement: a unified card refusal names the host room (placement: a unified card refusal names the host room #944)
• gpu: carry the host reading in a tier node's census (gpu: carry the host reading in a tier node's census #946)
• bloomery-serve: the qwen3 seat holds its stage card before it reads the room, so its plan is the CLI's split (bloomery-serve: the qwen3 seat holds its stage card before it reads the room, so its plan is the CLI's split #952)
• gpu: gdn_delta_prepared keeps its state transposed, so each product's accumulator is the next one's A operand (gpu: gdn_delta_prepared keeps its state transposed, so each product's accumulator is the next one's A operand #957)
• bloomery-serve, model-hf: a --hf repo's architecture picks the seats that check the arguments before the fetch (bloomery-serve, model-hf: a --hf repo's architecture picks the seats that check the arguments before the fetch #958)
• serve: n on /v1/chat/completions and /v1/completions answers that many choices of the prompt, queued together, choice i seeded with the request's seed plus i (serve: n on /v1/chat/completions and /v1/completions answers that many choices of the prompt, queued together, choice i seeded with the request's seed plus i #959)
• model: the qwen4exp MTP load refuses a self-contained draft by its own matrices before reading its layer (model: the qwen4exp MTP load refuses a self-contained draft by its own matrices before reading its layer #962)
• serve: a sampled request drafts beside other busy slots, its window a row of the round's one pass with its sampler and history; Qwen3.8's unset MTP draft drafts on any slot count (serve: a sampled request drafts beside other busy slots, its window a row of the round's one pass with its sampler and history; Qwen3.8's unset MTP draft drafts on any slot count #964)
• serve: logit_bias is read, and plain steps and MTP passes draw from the same biased rows (serve: logit_bias is read, and plain steps and MTP passes draw from the same biased rows #966)
• serve: /v1/responses keeps its responses: store, previous_response_id and background, in a bounded store (serve: /v1/responses keeps its responses: store, previous_response_id and background, in a bounded store #968)
• qwen38: a selecting layer keeps its raw indexer keys in a ring of eleven rows, not a plane of the context (qwen38: a selecting layer keeps its raw indexer keys in a ring of eleven rows, not a plane of the context #969)
• serve: structured output takes json_object, arrays, anyOf, $ref, integers and numbers, and samples inside the mask (serve: structured output takes json_object, arrays, anyOf, $ref, integers and numbers, and samples inside the mask #973)
• gpu: fuse the hostless Qwen3.8 card sum with the shared add (gpu: fuse the hostless Qwen3.8 card sum with the shared add #974)
• width: a losing chain probe runs one pass, and the diagnostics carry the gate rates (width: a losing chain probe runs one pass, and the diagnostics carry the gate rates #975)
• gpu: raise faults for overflowing QSA sums and invalid f16 cache values (gpu: raise faults for overflowing QSA sums and invalid f16 cache values #976)
• readers: refuse duplicate checkpoint keys and overlapping GGUF tensor ranges (readers: refuse duplicate checkpoint keys and overlapping GGUF tensor ranges #977)
• app: compare all and selected logits through the same host read boundary (app: compare all and selected logits through the same host read boundary #978)
• model: qwen4exp placement sizes saved states and PLE readers for the load, and refuses a missing plan callback (model: qwen4exp placement sizes saved states and PLE readers for the load, and refuses a missing plan callback #979)
• qwen38: validate completed target rows and own draft input (qwen38: validate completed target rows and own draft input #980)
• hf: normalize NVFP4 expert stacks from bounded input batches (hf: normalize NVFP4 expert stacks from bounded input batches #981)
• serve: --config FILE reads a start's flags from one JSON object, under the command line's (serve: --config FILE reads a start's flags from one JSON object, under the command line's #982)
• serve: keep saturated guard marks short and stable (serve: keep saturated guard marks short and stable #983)
• gemm: retain checked integer conversion for native activation launches (gemm: retain checked integer conversion for native activation launches #984)
• hf: serialize cached fetches and resolve checkpoint shards from the index (hf: serialize cached fetches and resolve checkpoint shards from the index #985)
• gpu: the HC up-mix's passes split the warp's columns, so a mix step reads a whole sector of each stream and its gate stores fall in 32 banks (gpu: the HC up-mix's passes split the warp's columns, so a mix step reads a whole sector of each stream and its gate stores fall in 32 banks #988)
• hf: give each listing its own header file (hf: give each listing its own header file #989)
• serve: the spawned mock's gates wait until it serves before their first request (serve: the spawned mock's gates wait until it serves before their first request #990)
• check-catalog-docs: the self-test edits README's docs/agents.md pointer at its current count (check-catalog-docs: the self-test edits README's docs/agents.md pointer at its current count #992)
• gpu: Qwen3.8's selecting stores and draft store back only the rows a sequence reaches (gpu: Qwen3.8's selecting stores and draft store back only the rows a sequence reaches #993)
• serve: one model by name (GET /v1/models/{model}), Ollama's /api/tags and /api/show, and a 404 that names the request (serve: one model by name (GET /v1/models/{model}), Ollama's /api/tags and /api/show, and a 404 that names the request #994)
• serve: an error about one request field names it in param (serve: an error about one request field names it in param #998)
• gpu: Qwen3.8's ubatch walk allocates its BF16 streams and normed streams at their width (gpu: Qwen3.8's ubatch walk allocates its BF16 streams and normed streams at their width #1001)
• gpu-gates: the qwen4exp e2e fixture budget carries the load's store page rounding (gpu-gates: the qwen4exp e2e fixture budget carries the load's store page rounding #1002)
• serve: a key file line sets its key's scope and its in-flight and rate limits (403, 429 with Retry-After) (serve: a key file line sets its key's scope and its in-flight and rate limits (403, 429 with Retry-After) #1003)
• lease: the perf stub matches the symbol argument, so a TMPDIR holding next no longer skips its held calls (lease: the perf stub matches the symbol argument, so a TMPDIR holding next no longer skips its held calls #1004)
• gemm: a warp's code copy of the native NVFP4 walk reads eight rows' whole steps (gemm: a warp's code copy of the native NVFP4 walk reads eight rows' whole steps #1005)
• serve: /metrics splits decode rounds by width, the engine's waits by phase and the draft's counts by depth (serve: /metrics splits decode rounds by width, the engine's waits by phase and the draft's counts by depth #1007)
• serve: reasoning_loop ends a think span that repeats a cycle, with finish_reason reasoning_loop (serve: reasoning_loop ends a think span that repeats a cycle, with finish_reason reasoning_loop #1009)
• serve: advance GLM and DSML tool calls from their cursor, scanning and streaming only the new text (serve: advance GLM and DSML tool calls from their cursor, scanning and streaming only the new text #1010)
• serve seats: --slot-ctx names one request's context apart from the count, and qwen38 takes --parallel auto (serve seats: --slot-ctx names one request's context apart from the count, and qwen38 takes --parallel auto #1011)
• gpu: Qwen3.8 slots share a prompt prefix's KV pages and copy them on write (gpu: Qwen3.8 slots share a prompt prefix's KV pages and copy them on write #1012)
• qwen38: an image load drafts with MTP, its walks on each sequence's rotation table and a span's rows (qwen38: an image load drafts with MTP, its walks on each sequence's rotation table and a span's rows #1015)
• gate_qwen38_serve: the props clause plans the card and host room the server's --place gate plan record names (a GB10 plans its unified pool) (gate_qwen38_serve: the props clause plans the card and host room the server's --place gate plan record names (a GB10 plans its unified pool) #1018)
• serve: a think or Qwen call tag is markup only where its special token was generated, so the same text in ordinary tokens stays content (serve: a think or Qwen call tag is markup only where its special token was generated, so the same text in ordinary tokens stays content #1020)
• docs: preserve source revision when refreshing fixture budgets (docs: preserve source revision when refreshing fixture budgets #1021)
• serve: a forced tool call follows the model's own close of an open think span (serve: a forced tool call follows the model's own close of an open think span #1023)
• gpu-gate.sh: the PIN_3090 refusal of the Qwen3.8 serve plain arm applies to the real tier only, so the fixture tier runs the plain arm on GB10 and RTX 5090 (gpu-gate.sh: the PIN_3090 refusal of the Qwen3.8 serve plain arm applies to the real tier only, so the fixture tier runs the plain arm on GB10 and RTX 5090 #1024)
• serve: a prompt's last call runs its last id as the call's last row on the Qwen3.8 seat (serve: a prompt's last call runs its last id as the call's last row on the Qwen3.8 seat #1025)
• ptx-shapes: re-pin the sm_121a rows the HC up-mix split and the MXFP8 TMA tile moved (ptx-shapes: re-pin the sm_121a rows the HC up-mix split and the MXFP8 TMA tile moved #1026)
• serve: Qwen3.8 slots save to a file and restore from it, the MTP draft's side with them (serve: Qwen3.8 slots save to a file and restore from it, the MTP draft's side with them #1027)
• gpu: Qwen3.8's decode selected flash merges under the out gate, one launch fewer a selecting layer (gpu: Qwen3.8's decode selected flash merges under the out gate, one launch fewer a selecting layer #1028)
• qwen38: a pass of several slots runs its MTP chains as one walk a depth over every slot's rows (qwen38: a pass of several slots runs its MTP chains as one walk a depth over every slot's rows #1029)
• gpu: a placed load reads the parts its host bytes planned instead of planning them again (gpu: a placed load reads the parts its host bytes planned instead of planning them again #1030)
• placement: plan each census card at its own figures and allocation granule (placement: plan each census card at its own figures and allocation granule #1031)
• qwen38: the plan counts a native FP4 card route by the bytes it holds (qwen38: the plan counts a native FP4 card route by the bytes it holds #1033)
• gpu-gate: the cancellation self-tests signal once the timeout sleeps past its fork (gpu-gate: the cancellation self-tests signal once the timeout sleeps past its fork #1035)
• hf: a large file downloads over up to 16 ranges at once, each resumed where it stopped (hf: a large file downloads over up to 16 ranges at once, each resumed where it stopped #1037)
• gpu: Qwen3.8's ubatch walk stores a delta layer's q, k and v projection as BF16, and runs β and α as one launch (gpu: Qwen3.8's ubatch walk stores a delta layer's q, k and v projection as BF16, and runs β and α as one launch #1039)
• gpu: the fused NVFP4 gate and up writes its SwiGLU in native FP4, and qwen38 counts a folder's NVFP4 rows in that form (gpu: the fused NVFP4 gate and up writes its SwiGLU in native FP4, and qwen38 counts a folder's NVFP4 rows in that form #1040)
• hf: refuse listed names that end in a cache sidecar's suffix (hf: refuse listed names that end in a cache sidecar's suffix #1048)
• serve: a request body sent gzip or deflate is read decoded; another content coding is a 415 (serve: a request body sent gzip or deflate is read decoded; another content coding is a 415 #1049)
• docs: Qwen3.8 on one DGX Spark: both folders' slots, context, MTP draft and memory (docs: Qwen3.8 on one DGX Spark: both folders' slots, context, MTP draft and memory #1053)
• model: coverage lists minimax-m3's own rows: the group-16 flash, the partial NeoX rope, the block-max selector, the 128-expert top-4 router and the swigluoai blocks (model: coverage lists minimax-m3's own rows: the group-16 flash, the partial NeoX rope, the block-max selector, the 128-expert top-4 router and the swigluoai blocks #1054)
• parts, ctx_ask: a folder load's refusal of a paged plan is short on the host, so auto and the context fit plan below it (parts, ctx_ask: a folder load's refusal of a paged plan is short on the host, so auto and the context fit plan below it #1062)
• gpu, qwen35moe, qwen38: a pass of several slots runs its sequence-bound launches once over every slot's rows (gpu, qwen35moe, qwen38: a pass of several slots runs its sequence-bound launches once over every slot's rows #1065)
• asm-census: refuse overwritten loop counters and spilled flags (asm-census: refuse overwritten loop counters and spilled flags #1066)
• serve, app, gates: a drafted pass whose pick reads only a row's best logits takes each verified row from the card's best read (serve, app, gates: a drafted pass whose pick reads only a row's best logits takes each verified row from the card's best read #1067)
• gpu: a row table's rows go to the card as read and decode there, in one owner for every row format (gpu: a row table's rows go to the card as read and decode there, in one owner for every row format #1068)
• serve, qwen38: a long prompt yields to busy slots on the positions' quantum grid, which the Qwen3.8 seat names as its checkpoint grid (serve, qwen38: a long prompt yields to busy slots on the positions' quantum grid, which the Qwen3.8 seat names as its checkpoint grid #1069)
• generate, beside, serve seats: the draft yield is one rule both seats call, and a drafted plan's refusal other than no room for the context stops the load by name (generate, beside, serve seats: the draft yield is one rule both seats call, and a drafted plan's refusal other than no room for the context stops the load by name #1070)
• gpu, chain, runtime: Qwen3.6's MTP store walk ends at the cache append, and a closed width gate judges a request on its own shadows (gpu, chain, runtime: Qwen3.6's MTP store walk ends at the cache append, and a closed width gate judges a request on its own shadows #1073)
• docs(serve): identify the continuation reference source (docs(serve): identify the continuation reference source #1074)
• model, placement, start: a family's placement refusal is its error's source, so the start keeps its no-fit class (model, placement, start: a family's placement refusal is its error's source, so the start keeps its no-fit class #1077)
• Restrict prompt retention to shared arenas (Restrict prompt retention to shared arenas #1079)
• gpu, placement, qwen38: every dense site of the wide walk runs its wide GEMM, its rows holding only the forms their sites read (gpu, placement, qwen38: every dense site of the wide walk runs its wide GEMM, its rows holding only the forms their sites read #1080)
• start: a host-side placement refusal names the bytes it asked for in the plan's refusal line (start: a host-side placement refusal names the bytes it asked for in the plan's refusal line #1081)
• host::costs: the expert stream and the split walk price from the detected device and this load (host::costs: the expert stream and the split walk price from the detected device and this load #1082)
• app, serve: one round of any draft's windows verified together; Qwen3.8's lookup windows ride one pass, and whether a load verifies several slots' windows is the body's fact, not the draft's name (app, serve: one round of any draft's windows verified together; Qwen3.8's lookup windows ride one pass, and whether a load verifies several slots' windows is the body's fact, not the draft's name #1085)
• runtime, gpu, app, gates: one prompt call plan for every seat: a part before the last runs no head, the last call folds its last id where the body allows it, and the unit cut and its read rows live in the shared owner (runtime, gpu, app, gates: one prompt call plan for every seat: a part before the last runs no head, the last call folds its last id where the body allows it, and the unit cut and its read rows live in the shared owner #1086)
• host::xstream, placement::workstation, seats: a load that streams reserves its lane's pinned staging on the host before its plan, the lane pins at most a half of slots, and the tiered open plans on the run's machine (host::xstream, placement::workstation, seats: a load that streams reserves its lane's pinned staging on the host before its plan, the lane pins at most a half of slots, and the tiered open plans on the run's machine #1089)
• serve_client, serve gates: ids_of refuses a response whose token ids are not all u32 ids (serve_client, serve gates: ids_of refuses a response whose token ids are not all u32 ids #1090)
• gate_qwen38_serve, gate_ds41_serve: a completion agrees with its reference by serve_client::agree alone (gate_qwen38_serve, gate_ds41_serve: a completion agrees with its reference by serve_client::agree alone #1092)
• model::arch::qwen35moe::fixture: the recorded card budget plans with the load's own activation stack, so q38dense holds 256 routed experts a layer again (model::arch::qwen35moe::fixture: the recorded card budget plans with the load's own activation stack, so q38dense holds 256 routed experts a layer again #1093)
• runtime::width: a slow start no longer fixes the width chooser's costs (runtime::width: a slow start no longer fixes the width chooser's costs #1095)
• app, gpu, qwen38: a load captures its rounds' passes and each slot's step and draft walks before it serves (app, gpu, qwen38: a load captures its rounds' passes and each slot's step and draft walks before it serves #1096)
• recipes: compile structural schema validators once (recipes: compile structural schema validators once #1097)
• fix(metal): subtract delta tap distance before reducing the ring (fix(metal): subtract delta tap distance before reducing the ring #1098)
• Preserve path errors in snapshot resolution (Preserve path errors in snapshot resolution #1099)
• gpu: one SwiGLU epilogue writes the down's q8 blocks of 32 from the up GEMM of either family: Q4_K over a card route, Q8_0 planes over a dense block, NVFP4 planes as before (gpu: one SwiGLU epilogue writes the down's q8 blocks of 32 from the up GEMM of either family: Q4_K over a card route, Q8_0 planes over a dense block, NVFP4 planes as before #1100)
• hf, model-hf: an FP8 or unquantized stack normalizes from bounded input batches into one output reserved at its length, as an NVFP4 stack does (hf, model-hf: an FP8 or unquantized stack normalizes from bounded input batches into one output reserved at its length, as an NVFP4 stack does #1101)
• gpu::draft: one executor for a draft's walks of several slots, with Qwen3.8's MTP and GLM's NextN as its programs (gpu::draft: one executor for a draft's walks of several slots, with Qwen3.8's MTP and GLM's NextN as its programs #1102)
• model-hf::folder: an FP4-Local folder's compressed block scales are read back once, so the host reads its experts in place (model-hf::folder: an FP4-Local folder's compressed block scales are read back once, so the host reads its experts in place #1103)
• gpu, qwen38: a pass of one row a slot steps each delta state in place, so its commit replays nothing (gpu, qwen38: a pass of one row a slot steps each delta state in place, so its commit replays nothing #1104)
• app::mtp: a pass of several slots' drafts calls the shared draft executor, with no per-body override (app::mtp: a pass of several slots' drafts calls the shared draft executor, with no per-body override #1105)
• q8f32: an MXFP8 gemv of five to eight columns runs the staged walk over an eight-column stage (q8f32: an MXFP8 gemv of five to eight columns runs the staged walk over an eight-column stage #1106)
• fix(metal): reject every malformed GEMM shape component (fix(metal): reject every malformed GEMM shape component #1109)
• gpu, placement: every family's prompt ring pass reads the ring's half through one common view, so an NVFP4 folder splits, and a plan with no host expert reserves no stream staging (gpu, placement: every family's prompt ring pass reads the ring's half through one common view, so an NVFP4 folder splits, and a plan with no host expert reserves no stream staging #1111)
• serve readers: read SSE frames whole in every gate and check the soak's paired chat answer (serve readers: read SSE frames whole in every gate and check the soak's paired chat answer #1114)
• fix(metal): reject missing values after repeated probe flags (fix(metal): reject missing values after repeated probe flags #1115)
• gates, qwen38, fixture: the NVFP4 gate's clauses run on a four-layer cut of the NVFP4 and FP4-Local folders in the fixture tier (gates, qwen38, fixture: the NVFP4 gate's clauses run on a four-layer cut of the NVFP4 and FP4-Local folders in the fixture tier #1116)
• mm31 oracle: a resume refuses the state of another cut generation (mm31 oracle: a resume refuses the state of another cut generation #1118)
• gpu, qwen38, qwen3: the KV cache's unset format is one rule (gpu::kv_format), f16 for a GGUF file and a checkpoint folder alike, and the load names what chose it (gpu, qwen38, qwen3: the KV cache's unset format is one rule (gpu::kv_format), f16 for a GGUF file and a checkpoint folder alike, and the load names what chose it #1119)
• ref, qwen4exp: the NVFP4 oracle reads a cut only under the # fixture line of its first shard (ref, qwen4exp: the NVFP4 oracle reads a cut only under the # fixture line of its first shard #1122)
• gates, placement: --parallel auto counts the slots in one shared fit on every seat of resident slots (gates, placement: --parallel auto counts the slots in one shared fit on every seat of resident slots #1123)
• gate-batch: a local batch runs one lane and packs against the machine's own memory (gate-batch: a local batch runs one lane and packs against the machine's own memory #1125)
• gpu, qwen38: the card route's down form is read from the whole slot map, so a load with an expert tier rounds its down rows to BF16 as its union's one-card load does, and the tier join reads the tier's rows rounded with them (gpu, qwen38: the card route's down form is read from the whole slot map, so a load with an expert tier rounds its down rows to BF16 as its union's one-card load does, and the tier join reads the tier's rows rounded with them #1126)
• gpu, gpu-gates, runtime: Qwen3.5/3.6 runs its GDN prompt, verify and closing call on the shared delta store (gpu, gpu-gates, runtime: Qwen3.5/3.6 runs its GDN prompt, verify and closing call on the shared delta store #1127)
• gpu-gates: read the Qwen3.8 greedy draft and cache counts strictly (gpu-gates: Qwen3.8 greedy reader hides malformed draft counters #1129) (gpu-gates: read the Qwen3.8 greedy draft and cache counts strictly (#1129) #1130)
• docs: the Spark and RTX 5090 tables hold one value a cell at main e7180ea, the Spark page's plan and speed tables follow, and stale citations point at the current code (docs: the Spark and RTX 5090 tables hold one value a cell at main e7180ea9, the Spark page's plan and speed tables follow, and stale citations point at the current code #1132)
• model: the qwen35moe trunk and two-3090 plan pins count each GDN layer's carry and lane stamp (model: the qwen35moe trunk and two-3090 plan pins count each GDN layer's carry and lane stamp #1134)
• qwen35moe e2e gate: the last-step test's unplaced CLI run drops the fixture tier's placed levers, which it would refuse by name (qwen35moe e2e gate: the last-step test's unplaced CLI run drops the fixture tier's placed levers, which it would refuse by name #1136)
• gpu: the GDN prompt delta's products take each f32 operand as three BF16 parts, so a chunk keeps the one-token rule's precision (gpu: the GDN prompt delta's products take each f32 operand as three BF16 parts, so a chunk keeps the one-token rule's precision #1140)
• gpu: the GDN prompt delta seeds only its high parts' term, so the smaller terms keep their own alignment (gpu: the GDN prompt delta seeds only its high parts' term, so the smaller terms keep their own alignment #1145)
• qwen38 serve and qwen4exp MTP gates: the floor band and the deep plain walk rest on premises that hold (qwen38 serve and qwen4exp MTP gates: the floor band and the deep plain walk rest on premises that hold #1147)

AI help

AI 코딩 에이전트가 문서를 작성한 후 이슈를 올렸습니다.