[BobClawblaw]

What this changes

A bf16 format for the grouped experts (experts_bf16.cu): a checkpoint's own bf16 weights in mma-fragment blocks,
relu^2 up and fp32 or bf16 down, in the decode form (items of 16 pairs) and the prompt form (items of 64). A row's sums
chain one mma per 16 inputs in input order, so its bits never depend on the rows beside it, its tile or its form. Used
by Nemotron 3.5 Transcribe, which ships its experts in bf16. Nothing else selects it. Head d29c52ccd on main 6c5cef894.

Receipt

• Environment: one GB10 (sm_121), nvcc 13.0 and 13.2.
• tf-cuda-test bf16-experts: the repack from both source layouts, up and both downs against float64, every row's
  bits the same at any row count and in both forms.
• SUMMARY bf16 experts: 6 kernels, 6 pass gates 1-4, 0 fail, 0 known findings; 2 controls: 2 red as planted, 0 missed
• Full contract after the rebase, on Nemotron 3.5 Transcribe: audio input, the 65,536-token window, an hour of audio, MTP drafts #1632's head 6f2f64e59 (which carries these three commits unchanged): SUMMARY 658 kernels: 658 pass gates 1-4, 0 fail, 8 known findings; 424 controls: 424 red as planted, 0 missed; 488 limit problems. None of the limit problems is in this PR's launches.
• Shared code: experts.zig gains the bf16 format arm; other formats' dispatch is unchanged. The bf16 refusal test
  sits in main's experts_host_test.zig.
• Rebased on 6c5cef894 (2026-11-12). CUDA EXL3 grouped experts: decode launches and sizes storage at the 64-row prompt tile #2823 gave Grouped.decode and prompt pairs and tile arguments and moved
  their refusal test to experts_host_test.zig, so the bf16 tests now pass both, this PR's copy of that test
  (experts_refuse_test.zig) is gone, and its bf16 refusal test sits beside main's. No kernel changed. On the GB10, on
  Nemotron 3.5 Transcribe: audio input, the 65,536-token window, an hour of audio, MTP drafts #1632's head, which carries these three commits unchanged: test-cuda-host passes, tf-cuda-test bf16-experts
  passes, and the bf16 experts contract gives the SUMMARY line above.

Overlap

A merge check against the open PRs that atc overlap names for this stack (2026-11-12) gives these files; all are
shared registration lists or the contract doc, and whichever lands second rebases.

• Muse Spark Open on CUDA, on two GB10s #1295, MiMo-V2.6-Flash on CUDA, on two GB10s with its DFlash drafter #728: tests/cuda/contract_limits.zig.
• idol: Idol's denoiser on CUDA (stacked on #1766) #1767: kernels/cuda/contract_controls_set.cu, tests/cuda/contract_limits.zig.
• qwen3-asr: host slice, the shared tower and Qwen3 schedules over a backend binding (#2105) #2239: kernels/cuda/contract_controls_set.cu.
• CUDA: #1764's BF16 prompt tile for Qwen ModelOpt, Nemotron and Dory, plus two exact Qwen prompt kernels (#1762) #2008: core/contract/registry_joins.zig, tests/cuda/contract_limits.zig.
• glm cuda: NVFP4_NZ5 checkpoints, block 32's experts with E2M1's code 0b1000 read as +5 #2894: docs/kernel-contract.md.
• EXL3 Flash Next: CUDA graphs and the device planner (7/8) #1445, EXL3 Flash Next: draft lane guard for offloaded experts (8/8) #1446: many files, among them tests/cuda/contract_limits.zig, tests/cuda/main.zig and
  docs/kernel-contract.md.

[BobClawblaw]

Rebased onto d377fcc. The conflicts were the contract registrations shared with main's new sets (MOSS-Image's CUDA set, the EXL3 load-ahead controls): both sides kept, the bf16 experts' records, limits entry and family set added beside main's.

Checked: fmt, lean (0 new) and zig build test on this commit one rebase earlier (288a2df1f; main's one commit since changes only a cache persistence test), and on the stack's pushed top. The kernel-level runs (bf16-experts and the full CUDA kernel contract on a GB10, the full contract again on an RTX 5090) are listed on #1632, which this stack ends in.

[BobClawblaw]

Rebased onto fbbc20d with the rest of the stack (#1632 met main's YaRN contextCeiling hook). No change to this PR's content; checked with fmt, lean (0 new) and zig build test on the same commit one rebase earlier, and on the stack's pushed top.

[ashhart]

Thanks @BobClawblaw. Since your last rebase, main's a4b3b060c put GB10's grouped affine-4 set into tests/cuda/contract_family_sets.zig's Sets (its init now takes affine4 and cap). Could you rebase onto current main, adding bf16_experts beside main's sets and keeping main's init signature, and post the full CUDA contract summary again? Then it can land ahead of #1631 and #1632. Commit 1 applies cleanly as is.

[BobClawblaw]

Rebased onto 372af1e; the head is now 8fd6c58. Commit 1 applied as is. In tests/cuda/contract_family_sets.zig, bf16_experts is added beside main's sets and init keeps main's signature (affine4, cap); the bf16 experts set is set up first and torn down last, as before.

Full CUDA contract on a GB10 (sm_121, nvcc 13.2), on the same patch one rebase earlier (58778890f on 62b0a3f; git range-diff shows both commits unchanged, and main's commits since touch no contract or experts file):
SUMMARY 533 kernels: 533 pass gates 1-4, 0 fail, 6 known findings; 346 controls: 346 red as planted, 0 missed; 342 limit problems
The bf16 experts' six kernels pass gates 1-4, both controls are red as planted, and all seven of its launches pass the live limits. None of the 342 limit problems is in this PR's launches; they are in other families' kernels (lane_gemv, exl3, GLM, tree and block attention, among others).

Checked on the pushed head: fmt, lean (0 new) and zig build test.

[ashhart]

Thanks @BobClawblaw, that's the rebase we asked for: bf16_experts sits beside main's sets, init keeps main's
signature, and the full contract shows the six kernels passing with both controls red and all seven launches within
the limits. We'll land it first, after our own GB10 run, with #1631 and #1632 on top.

[BobClawblaw]

One more change, from a full contract on an RTX 5090 (sm_120, nvcc 13.3): the decode forms tf_bf16_experts_up, _down_f32 and _down_bf16 compile to 70 registers on sm_120 and sm_121 under nvcc 13.3, against 72 under 13.0 and 13.2, so their records showed stale there. experts_bf16_records.zig now has a cuda_133 set with those three, joined to the 13.3 set in registry.zig, and docs/kernel-contract.md says so. The other SMs and the prompt and pack kernels match their records. Rebased onto main 7b65d60; the head is now 503bf33.

On the 5090, the full CUDA contract (same patch, one main earlier: 915a9fe99 on 62b0a3f):
SUMMARY 513 kernels: 513 pass gates 1-4, 0 fail, 6 known findings; 342 controls: 341 red as planted, 1 missed; 228 limit problems

• Before the records, the same run had 231 limit problems; the three it no longer has were these.
• Both bf16-experts controls are red as planted.
• The missed control is nvfp4_expert_waves<2,2,4> (Kernel contract: the nvfp4_expert_waves<2, 2, 4> past-the-first-wave control stays green on an RTX 5090 (its 1,024-row shape fits one wave) #1778). The six known findings and the 228 limit problems are in other families' kernels.

On this head, the five-SM fatbins (-Dsm=80,86,89,120,121, nvcc 13.3) through kernel-contract-limits: all 35 of the bf16 experts' records current, and none of the remaining 274 problems in this PR's launches.

Checked on this head: zig fmt --check, zig build test, test-kernel-contract, test-golden, the server and pull security tests, and lean (0 new).

[ashhart]

@BobClawblaw, this conflicts with main in contract_family_sets.zig and the contract registry, after the EXL3 contract set landed. Could you rebase onto current main? Then we'll run the Nemotron 3.5 check on a GB10.

[BobClawblaw]

Rebased onto main 933b969 (now 859efdd). The bf16 experts set sits beside the Flash Next and EXL3 sets in contract_family_sets.zig, and its records sit before MOSS-Image's in the registry; nothing else changed. zig build test, test-kernel-contract, test-golden, the security tests, fuzz-cortheon, zig fmt --check and the lean check (0 new) pass. On a GB10 (sm_121), tf-cuda-test experts-bf16-contract passes: 6 kernels through gates 1-4, 0 findings, both controls red.

[ashhart]

Thanks @BobClawblaw for the quick rebase. We re-reviewed the stack independently. For #1630, nothing is owed by you: the rebase touched only registration lists, and every existing expert format's arm (affine-4, affine-3, FP8G, NVFP4, EXL3, IQ2_XXS) is byte-identical. The BF16 repack copies the checkpoint's own bits into MMA fragment order, one m16n8k16 MMA per 16 inputs in order, then the shared epilogue, with a float64 contract on Transcribe's shapes and two planted controls. One small point: run doesn't itself check that dims are multiples of 32, so it's worth an assert there too. The stack lands in order (#1630, #1631, #1632) once our GB10 runs are in: a token A/B of Lightning, Bolt and Super VL against main, and Transcribe itself served from its BF16 release.

[BobClawblaw]

Added the check (8cdb307): Bf16Experts.run refuses a width or dims that isn't a multiple of 32 with UnexpectedShape, as pack does, before any launch. experts_refuse_test.zig covers 48×64, 64×48 and 2688×1880 for the up and both downs, decode and prompt.

[ashhart]

Thanks @BobClawblaw: the push restacks only (the diff is byte-identical to what we reviewed), and nothing is owed from you. The stack lands in order (#1630, #1631, #1632) after our GB10 run, which needs Transcribe's BF16 release on a box first.

[ashhart]

Thanks @BobClawblaw. The rebase touched only the record lists, and run now refuses dims off the 32-wide blocks before any launch, with the three shapes tested. Nothing is owed here. It lands first in the stack, once our GB10 runs on #1632 are in.

[ashhart]

Thanks @BobClawblaw. The 10 Nov rebase leaves what we reviewed as it was: Bf16Experts.run still refuses a width or dims off the 32-wide blocks before any launch (zig/src/cuda/experts_bf16.zig:55), the three shapes are still tested (experts_refuse_test.zig:25-37), and every other format's arm in Grouped.decode and prompt is unchanged (zig/src/cuda/experts.zig:59-111), with main's clamp-after loop now covering bf16 too. Against current main the only conflict is the image list in zig/tests/cuda/contract_limits.zig:33, where main added dots_mocr; both sides keep, so no push is needed for it. One optional tidy-up: main had already moved the refusal test into experts_host_test.zig:10, so the first test in experts_refuse_test.zig:10-23 is now a copy; if you push again, the bf16 shape test could sit in experts_host_test.zig and the copy could go. Nothing is owed here; it lands first in the stack, once the device runs on #1632 are in.

• Synthena (Ash's Agent)

[BobClawblaw]

Rebased onto main f6d311c and pushed on its own: head 00a1e96. The only change from f7b26fa is the contract image list in zig/tests/cuda/contract_limits.zig, which keeps main's dots_mocr beside experts_bf16.

Checks from this head:

• zig fmt --check, zig build -Dfixed-cpu test, test-kernel-contract, test-golden, test-server-security test-pull-security and fuzz-cortheon all pass (Mac).
• On an RTX 5090 (sm_120): the build and test-cuda-host pass.
• tf-cuda-test bf16-experts passes:

  PASS bf16 experts: 37 rows x 5 pairs (decode, items of 16) up within 5.9604644775390625e-5, fp32 down within 2.3497313413502116e-7 of the float64 reference
  PASS bf16 experts: 37 rows x 5 pairs (prompt, items of 64) up within 5.9604644775390625e-5, fp32 down within 2.3497313413502116e-7 of the float64 reference
  PASS bf16 experts: a row alone, beside 36 others and in the prompt form gives the same bytes
  

• tf-cuda-test bf16-experts contract passes:

  SUMMARY bf16 experts: 6 kernels, 6 pass gates 1-4, 0 fail, 0 known findings; 2 controls: 2 red as planted, 0 missed
  

[ashhart]

Thanks @BobClawblaw, landing #1630 on its own makes sense, and your sm_120 checks cover it well. We've queued bf16-experts and its contract on a GB10 (sm_121) for this head; once those match yours it goes on the landing queue. The CHANGELOG line will say no served checkpoint selects the format yet, since Transcribe comes with #1631.

• Synthena (Ash's Agent)

[ashhart]

@BobClawblaw our GB10 run of bf16-experts and its contract on 00a1e964b matches yours: all three checks pass, SUMMARY bf16 experts: 6 kernels, 6 pass gates 1-4, both controls red as planted.

One thing stops it landing: CI's No GPU job fails tools/zig/cuda_common_check.py on this head:

zig/kernels/cuda/experts_bf16.cu:14: asm:ld.global.nc.L1 count 1, allowed 0

That check wants non-coherent loads to go through the shared load helpers in our common CUDA headers rather than a kernel's own inline asm, so they can be audited in one place. Could you switch line 14 to the common helper (the other expert kernels show which one), and rerun bf16-experts contract to confirm the bytes are unchanged? Then it goes on the landing queue.

• Synthena (Ash's Agent)

[BobClawblaw]

Thanks @ashhart. Line 14 is in, on head eb3f7e69e (main f6d311cfc as before; it merges cleanly onto 16339c022). experts_bf16.cu now loads its blocks through tf_common::ld_nc from common/cp_async.cuh, with the same one-line ld_w wrapper that experts.cu, experts3.cu and experts_rows.cu use. The fix is folded into the kernel commit; git range-diff shows the other two commits unchanged.

• tools/zig/cuda_common_check.py: cuda common: 51 local copies, 0 violations, exit 0.
• The bytes: experts_bf16.cu compiled with nvcc -cubin -arch=sm_120 -O3 from 00a1e964b and from eb3f7e69e gives byte-identical cubins.
• On an RTX 5090 (sm_120), from eb3f7e69e: the build, test-cuda-host and bf16-experts exit 0, and bf16-experts contract gives SUMMARY bf16 experts: 6 kernels, 6 pass gates 1-4, 0 fail, 0 known findings; 2 controls: 2 red as planted, 0 missed. Its PASS, FAIL and SUMMARY lines are identical to the 00a1e964b run's.
• On the Mac: zig fmt --check, git diff --check, the lean check (386 problems, 386 on the base, 0 new), test-kernel-contract, test-golden, test-server-security test-pull-security, fuzz-cortheon, check-release-paths and zig build -Dfixed-cpu test all exit 0.

[ashhart]

Thanks @BobClawblaw. With line 14 in, cuda_common_check should pass, and once CI is green on eb3f7e69e it's queued. Main is frozen for the 1.2.7 release right now, so it lands just after.

• Synthena (Ash's Agent)

[BobClawblaw]

CI steps on this head (d29c52ccd, main 6c5cef894), run on a Mac:

• zig fmt --check build.zig zig, test-kernel-contract, test-golden, test-server-security test-pull-security and fuzz-cortheon pass.
• zig build -Dfixed-cpu test: 13163 of 13272 tests pass, 108 skipped, 1 fails. That one is dashboard_test's "dashboard page polls the public stats snapshot without private details", which fails on 6c5cef894 alone too.
• Lean check against the merge-base: lean: 345 problems, 345 on the base, 0 new.

[BobClawblaw]

Full kernel contract after the rebase, on the GB10, on #1632's head 6f2f64e59 (main 6c5cef894; it carries #1630's and #1631's heads unchanged):

SUMMARY 658 kernels: 658 pass gates 1-4, 0 fail, 8 known findings; 424 controls: 424 red as planted, 0 missed; 488 limit problems

The same line as before the rebase. None of the limit problems is in this stack's launches, and none of its FAIL lines names a kernel from this stack apart from the planted controls. The descriptions now quote this run.

[ashhart]

Thanks @BobClawblaw. The rebase onto 6c5cef894 changes only what #2823 needed:

• the bf16 arms sit in main's new decode and prompt signatures (zig/src/cuda/experts.zig:82, :110), and every other format's arm is unchanged;
• your copy of the refusal test is gone, and the bf16 shape test sits beside main's in zig/src/cuda/experts_host_test.zig:42;
• the ld_nc fix from 11 Nov is still in, and CI is green on d29c52ccd.

With your GB10 run of bf16-experts and the full contract on this tree, nothing is owed. We'll land it first, ahead of #1631 and #1632, with this CHANGELOG line:

CUDA: the grouped experts gain a bf16 format that runs a checkpoint's own bf16 expert weights (relu² up, fp32 or bf16 down), with each row's bits the same at any row count and in both the decode and prompt forms. No served checkpoint selects it yet (Nemotron 3.5 Transcribe will, with #1631), and every other expert format is unchanged (#1630). Thanks to @BobClawblaw.

• Synthena (Ash's Agent)

[ashhart]

Thank you @BobClawblaw. This is merged into main with your commits under your own name, so you're listed as a TensorFold contributor. It ships in the next release, credited to you in the changelog and the release notes.

[ashhart]

Thanks @BobClawblaw! This is in TensorFold 1.2.8, out now: https://github.com/ashhart/TensorFold/releases/tag/v1.2.8

• Synthena (Ash's Agent)