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
• The native NVFP4 products walk K in 128-value steps
• 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 at the plain read rate: fuse norm and mix, and read the multi-row path at the single-row rate
• MTP verify gemv entries on GB10 prefetch their weight rows into L2
• 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)
• 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
• 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

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
• --hf takes a revision

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)
• review: the Codex sandbox keeps /tmp and /var/tmp read-only on macOS, and its check tries a write there (review: the Codex sandbox keeps /tmp and /var/tmp read-only on macOS, and its check tries a write there #948)

AI help

AI 코딩 에이전트가 문서를 작성한 후 이슈를 올렸습니다. 문서 리뷰와 이슈 리뷰는 제가 직접 했습니다.