[jayleaton]

Stacked on #1766 (Idol's denoiser on Metal). This branch holds #1766's commits too; its own change is the five commits after e34dab7ad (#1766's 9aaf5d318, replayed on main d367f0fb7): 3d8d9950c (the CUDA denoiser), 2e7d31d1a (recipe), 15d7a9bf0 (harder attention sums inputs), 0e0394fa6 (the notice) and f710b4a77 (shared common/ reduction and mma helpers, for cuda_common_check.py). Review after #1766, or diff against its branch.

What this changes

Idol's denoiser on the CUDA runtime, beside the Metal family and driven by the same case driver and gates (zig/tests/idol_run.zig; tf-idol on CUDA, step idol-cuda). For #1757.

Run on a GB10. It compiles for sm_120 and sm_121 with nvcc 13.0, and the CUDA host tests pass in a Linux box. On a GB10 (ashhart's rerun at 776a3335d), all 13 kernels pass the kernel contract's gates 1-4, and all three controls go red as planted. The restack onto ca429e95a moved the controls into their own header, so idol-contract needs one more GB10 run on f5b88dd98 (or 36828402d, whose SASS is byte-identical). On sm_120, Nipale-ai's RTX 5090 run (CUDA 13.3, nvcc 13.3.73) on f5b88dd98 passes the kernel contract (13 of 13 kernels, 3 of 3 controls red) and the synthetic case; its 13.3 limit numbers differ from the 13.0 records but stay within their bounds. Idol's CUDA sources are byte-identical at f710b4a77.

What it adds

• zig/kernels/cuda/idol.cu (built with the glue flags: --fmad=false, IEEE division and square root) and idol_attention.cuh: the norms, RePo's angles, q/k rotation and F.normalize (a NaN norm stays NaN, as torch's clamp keeps it), the gated residual, bias with GELU, the embedding, the condition sum and the fp32 bias. Attention is Idol's own head-64 kernel with vision_kernel's row arithmetic: fp16 q, k, v and P (rounded to nearest), fp32 scores, softmax and sums. No shared attention template gains an instance, so no other family's SASS changes.
• zig/src/families/idol/cuda*.zig: weights, the model and launchers. Model.load admits sm_121 and sm_120 and refuses any other GPU by name, with a log line.
• Projections on cuBLASLt with CUBLAS_COMPUTE_32F: fp16 block linears into fp32 sums (bias and one fp16 rounding in idol.cu), and true fp32 (fp32 inputs, no TF32) for RePo's maps, the condition embedding and the output, so CUDA needs no hi/lo split. A host test pins the compute type's value.
• Kernel contract: tf-cuda-test idol-contract, gates 1-4 over every idol.cu kernel and the attention, with three controls that each drop one fp16 rounding; it is in the full contract's family sets and contract_limits.zig. Gate-5 records (idol_cuda_records.zig) are nvcc 13.0's for sm_80, 86, 89, 120 and 121, and match the repo's parser on all 50 entries.
• docs/recipes/idol-cuda.md, and a changelog line.

Shared code

• cublaslt.zig: Linear.initTyped takes the input and output types; Linear.init is now initTyped(.bf16, ...), so every existing caller builds the same plan as before.
• contract_attn.zig: Ref.out lets the attention check read fp16 outputs; it defaults to bf16, the only value before.
• host.zig / weights.zig: the fp32 host tensors' struct moves into host.zig so both backends load it; Metal is unchanged.
• zig/tests/idol.zig: the case driver moves into idol_run.zig, shared with CUDA. On Metal it gives the same gates and numbers as idol: Idol's denoiser as a native Metal family #1766's tool on the same case.
• Registration beside MOSS-Image's: the records through idol_cuda_sets.zig in registry_joins.zig, contract_limits.zig, main.zig and cuda.zig. Idol's set goes after joint_schema in contract_family_sets.zig, and the host tests sit beside MiniMax's in cuda_family_host.zig. The three controls are in their own contract_idol_controls.cuh, included from contract_controls_set.cu. registry.zig and contract_controls.cu are untouched.
• THIRD_PARTY_NOTICES.md: idol.cu and idol_attention.cuh join idol: Idol's denoiser as a native Metal family #1766's idol-sdk entry.

Receipt

• Environment: this branch on idol: Idol's denoiser as a native Metal family #1766's commits (9aaf5d318, replayed as e34dab7ad) on main d367f0fb7; a Linux box with Zig 0.17.0 for builds, and with nvcc 13.0 for the earlier fatbin runs below; Mac Studio M5 Ultra for the Metal side.
• CUDA, in the Linux box (no GPU), before the restack:
  • zig build test-cuda-host: 278 of 278 steps. On the first run one test failed (3,097 of 3,106 passed, 8 skipped); the rerun ran the failed binary again and it passed. I did not identify which test that was; nothing in this branch changed between the runs.
  • zig build fatbins -Dnvcc=... -Dsm=121,120: 218 of 218 steps; idol.fatbin holds sm_121 and sm_120.
  • zig build for Linux: tf-idol, tf-cuda-test and tensorfold build (before the rebase onto MOSS-Image's CUDA port; the box run after it built test-cuda-host and the fatbins).
  • contract_controls.cu no longer changes: main closes tf_ctl_pocket_in_fmax, so the restack dropped f4a07e29e, and Idol's controls moved into contract_idol_controls.cuh.
• Metal, on the M5 Ultra: zig build test test-golden tf-idol -Dcpu=apple_m1: 593 of 593 steps, 5,254 of 5,283 tests pass, 29 skipped, 0 fail. tf-idol on the SDK's avocado case through the shared driver: all 101 gates pass, with the same run worst row (0.04738), final-sample cosine (0.999993) and 50-step time (27.53 s) as idol: Idol's denoiser as a native Metal family #1766's tool, and identical repeats.
• On f710b4a77 (the rebase onto d367f0fb7; idol's sources unchanged), in Linux (no GPU): cuda_common_check.py 0 violations, cuda_fragment_check.py 0 misses, the checker unittest 9 of 9, check_cuda_sources.py 12 of 12 on idol and the controls set, zig fmt --check clean, pr_check.py --pr 1767 0 errors, lean_check.py 0 new against main, git diff --check clean.
• On f5b88dd98, in Linux (no GPU): zig fmt --check build.zig zig clean; lean_check.py 345 problems, the same as main; test-kernel-contract, test-idol-cuda-host and compile-cuda-tests 48 of 48 steps, 251 of 251 tests; cuda_common_check.py 0 violations, cuda_fragment_check.py 0 misses, the checker unittest 9 of 9; zig build test 694 of 694 steps, 7,269 of 7,313 tests (44 skipped), exit 0. nvcc 13.0.88 (container, no GPU) gives SASS and resource usage byte-identical to 36828402d for the idol and controls images on all five SMs.
• CUDA, on a GB10 (ashhart's rerun at 776a3335d): SUMMARY idol: 13 kernels, 13 pass gates 1-4, 0 fail; 3 controls: 3 red as planted, 0 missed. The attention sums case gives exact sums at k 12294, where fp16 sums would miss by 248135x and bf16 by 1723x. The first run, at d153b6cf0, passed 12 of 13: that case's inputs were too easy for fp16 to miss by 100x, and f35cdd6f1 made them harder. The controls moved in the restack, so this needs a rerun on f5b88dd98 (or 36828402d; same SASS).
• Not run, and why: the GB10 idol-contract rerun on f5b88dd98 (or 36828402d), which is ashhart's. Gate 5 against a live driver, kernel-contract-limits, the cuBLASLt plans' algorithm search, admission on a device, weight loading, and the case's parity, repeats, timing and memory, which need a GPU run beyond the contract. nvcc 13.2 and 13.3 records are not measured yet (those compilers fall back to 13.0's).

Checklist

• Every token still goes through the lane rounds. No serial path, and nothing that needs drafts off. (No token path.)
• Drafted output equals "draft": false, a resumed prompt equals a fresh one, and concurrent equals solo. (n/a.)
• No precision traded for speed. If the bits change, the description says which sums change and why.
• A new or changed GPU kernel is registered with the kernel contract (docs/kernel-contract.md); each backend's SUMMARY line is in the receipt. (Registered; the GB10 SUMMARY line, 13 of 13, is above.)
• Prompt processing is no slower than the last release. (Not touched.)
• The description says which platforms I ran, Metal M1 to M5 and CUDA, and which I could not.
• New tests fail before the change, pass after it, and skip cleanly without their dependency. (Host tests yes; the device contract passes on GB10; the device case run is still pending.)
• Comments and docstrings are one line. No measurements or history in the source.
• No personal data, machine names, internal hosts or local paths. No AI attribution lines.

[ashhart]

Thanks @jayleaton. The shared changes read as no-ops for existing callers: Linear.init goes through initTyped(.bf16, ...) and Ref.out defaults to bf16.

Two things before it can land:

• On current main, zig/src/core/contract/registry.zig and zig/build/cuda.zig are both at exactly 600 lines. This PR adds one line to each (the _ = @import("idol_cuda_records.zig") test line and the idol_cuda.zig targets hook), which takes both past the lean limit. Could you fold them into an existing line or move them into Idol's own files, as idol: tf-idol-generate end to end on CUDA, and one GB10 qualification script (stacked on #1787, #1858-#1861; not yet run on a GPU) #1862 does for the component records?
• Device receipts: the idol-contract SUMMARY line, the case's parity, repeats, timing and memory from a GB10. If you don't get time on one, we can run it on ours and post the results here.

Please also rebase onto #1766's new head once it's on main.

[jayleaton]

Thanks @ashhart. The line-limit fix is pushed:

• registry.zig: the cuda list now takes idol_cuda_sets.zig (as in idol: tf-idol-generate end to end on CUDA, and one GB10 qualification script (stacked on #1787, #1858-#1861; not yet run on a GPU) #1862), which joins Idol's CUDA records and holds their tests, so the separate test import is gone.
• build/cuda.zig: idol_cuda is one more entry in the existing runner-imports line. Its targets builds tf-idol and returns the launchers.

Both files gain no lines, so they stay at 600 once this is on main. The later Idol PRs will add to idol_cuda_sets.zig instead of registry.zig.

Host gates pass here (Linux, Zig 0.17.0): zig fmt, test-idol-cuda-host, compile-cuda-tests, test-contract-images-host, test-kernel-contract, zig build test and check-release-paths, with no new lean findings.

We don't have GB10 time, so yes please: we'd be glad if you ran the idol-contract receipts on yours (the SUMMARY line, the case's parity, repeats, timing and memory). We'll rebase onto #1766 once it's on main.

[ashhart]

Thanks @jayleaton, both files now stay at 600 lines, and moving Idol's records into its own set file is the right shape for the PRs after this one. We'll run the idol-contract receipts on a GB10 (the SUMMARY line, the case's parity, repeats, timing and memory) and post them here. Rebase onto #1766 once it's on main and this is ready. We have no sm_120 card, so that admission stays unrun unless someone with one can post it.

[ashhart]

@jayleaton, our GB10 run of idol-contract at d153b6cf0 is in:

SUMMARY idol: 13 kernels, 12 pass gates 1-4, 1 fail; 3 controls: 3 red as planted, 0 missed

• Passing: the norms (tf_idol_norm affine and modulated, tf_idol_norm32) pass non-finite, repeat (five runs byte-identical), shapes and sums. For example, tf_idol_norm32 notes that fp16 sums would miss by 453x and bf16 by 4405.5x.
• Controls: the three planted controls (the modulated norm unrounded before scale, gate_add without fp16(g y), and GELU of the unrounded linear) each go red at shapes, as planted.
• The one failure is the sums gate on tf_idol_attention, full<64, 8>, fp16: "its stress inputs make fp16 sums miss by only 0.3x (needs 100x)". The case's inputs aren't hard enough yet to show the kernel's sums do better than fp16 accumulation would. Please strengthen that case's stress inputs, for example large cancelling terms across the full 64 keys, so fp16 sums would miss by 100x or more, and the gate can then show the kernel's own sums meet the bound. Or, if fp16 accumulation is what the kernel is meant to do there, say so and declare it.

The job's general CUDA kernel contract half compared against a base build that didn't complete on our side, so we're not reading anything from those lines. They're ours to rerun, and nothing is owed from you there. With the attention case fixed, this rebases once #1766 lands, as planned.

[jayleaton]

Thanks @ashhart. The attention kernel sums in fp32 (both mma steps accumulate in f32, and so does l), so I strengthened the stress inputs rather than declaring fp16.

Why it missed by only 0.3x: the values were odd integers up to 15. In the strided order each of the 32 fp16 lanes stayed below 2048, where fp16 integers are exact, so the only fp16 error was rounding l (12294 to 12296).

f35cdd6 raises them to odd integers up to 255, still mirrored so they cancel. The lane sums now pass 2048 and round. Every fp32 partial sum stays an integer below 2^24, so fp32 is still exact in any order.

On the host, gate 4's proof at the joint shape (12294 keys) now gives:

• fp16 sums miss by 177,485x (strided; sequential and pairwise overflow to inf)
• bf16 sums miss by 881,106x
• fp32 sums are exact in all three orders

A new host test in test-idol-cuda-host runs this check: it fails on the old values and passes on the new ones. zig fmt, test-idol-cuda-host, compile-cuda-tests and test-kernel-contract pass, and lean_check.py stays at the base count (424).

Could you rerun idol-contract on the GB10? The exact ratio depends on the run's seed, but it should be far above 100x. That run is also the first time the kernel's own output goes through the sums gate.

[ashhart]

/atc exchange 7979a3260563fc76a06d631a0348ff6971edb9a158b0bbbc0c3d970b63c0fa93

[ashhart]

Thanks @jayleaton, that's a clear diagnosis and the right fix. Odd values below 256
(zig/tests/cuda/contract_idol_attn.zig:111) push fp16 lane sums past 2048 while every fp32 partial sum stays an exact
integer, and the 100x bar and the kernel's bound are unchanged. So the gate got harder, not easier, and the host test
at :166 would fail on the old inputs. We'll rerun idol-contract on a GB10 at f35cdd6f1 and post the SUMMARY line
here.

The CUDA case's parity, repeats, timing and memory need the same case files we asked for on #1766 (case, exact and
the blocks panels, with their SHA-256s). Once those are shared we can run both. After #1766 lands, please rebase this
onto main. The conflicts are registration lists, and contract_family_sets.zig has grown on main, so Idol's set goes at
its new end. Please also update the title and changelog line, which still say it hasn't run on a GPU. Main is frozen
for a release right now, so nothing lands until that's out.

• Synthena (Ash's Agent)

[jayleaton]

Thanks @ashhart. The title and changelog line are updated in 776a333. The title drops "not yet run on a GPU". The changelog now says that on a GB10 the kernel contract passed gates 1-4 for 12 of 13 kernels, that the attention sums case got harder inputs and is waiting on your rerun, and that sm_120 is still unrun. The rebase onto main comes after #1766 lands and the release freeze ends, with Idol's set at the new end of contract_family_sets.zig.

[ashhart]

Thanks @jayleaton. The title and changelog line now match where this stands. The idol-contract rerun at f35cdd6f1 on a GB10 is ours to do, not yours; your latest commit changes only the changelog, so nothing is needed from you for it.

For the rebase once #1766 is on main:

1. Please drop f4a07e29e. Main already closes tf_ctl_pocket_in_fmax, so reapplying it would add a second brace.
2. zig/kernels/cuda/contract_controls.cu is at 600 lines on main, and Idol's controls would add 56. Main now builds the controls image from contract_controls_set.cu, which includes each family's controls from its own .cuh (as contract_js_controls.cuh does). Could Idol's go in one there?
3. Idol's set goes at the new end of contract_family_sets.zig, as before.

The CUDA case's parity, repeats, timing and memory still need the case files from #1766 (case, exact and the blocks panels, with their SHA-256s). The description's "Not run on a GPU" paragraph could also be updated.

• Synthena (Ash's Agent)

[ashhart]

Thanks @jayleaton. The idol-contract rerun on a GB10 is in, on this PR's head 776a3335d (the same code as f35cdd6f1, plus your changelog edit):

SUMMARY idol: 13 kernels, 13 pass gates 1-4, 0 fail; 3 controls: 3 red as planted, 0 missed

The attention sums case now passes too, with exact sums at k 12294 where fp16 sums would miss by 248135x and bf16 by 1723x. All three planted controls fail at the shapes gate, as they should. The full kernel contract on the same tree adds no failure line that names an Idol kernel. Its other differences from main come from this branch's older base, and the rebase clears them.

With that, the changelog line can say 13 of 13. The rest stands as before: the case files from #1766 for parity, repeats, timing and memory, then the rebase once #1766 is on main.

• Synthena (Ash's Agent)

[jayleaton]

Thanks @ashhart, good to see 13 of 13. The changelog line now says all 13 kernels pass gates 1-4 on a GB10, all 3 controls go red as planted, and sm_120 has not run (8fc8dc052). The description's receipt now quotes your rerun's SUMMARY line and sums figures.

[ashhart]

Thanks @jayleaton, and for putting 13 of 13 into the changelog line (8fc8dc052); for this code the GB10 contract is settled. What's left is the rebase and the case files, so here is everything in one place:

1. Restack: the six idol: Idol's denoiser as a native Metal family #1766 commits this branch carries are its copies from before the rebase (ending at 94a039249; idol: Idol's denoiser as a native Metal family #1766 is now at 5f86cedf9). Their changes match except the idol_records line, but their context is older main. Please drop them and replay your own commits on main once idol: Idol's denoiser as a native Metal family #1766 lands.
2. Drop f4a07e29e: main already closes tf_ctl_pocket_in_fmax.
3. Move Idol's controls: the lines in f04ce6116 would take zig/kernels/cuda/contract_controls.cu past 600. Please put them in their own .cuh, included from contract_controls_set.cu as contract_js_controls.cuh is.
4. Registration on current main: the CUDA join moved from registry.zig to zig/src/core/contract/registry_joins.zig, so @import("idol_cuda_sets.zig").cuda ++ goes beside MOSS-Image's line (:19) and registry.zig stays untouched; contract_family_sets.zig now ends with joint_schema, so Idol's set goes after it, with an errdefer for joint_schema now that it isn't last; "idol" joins the list at zig/tests/cuda/contract_limits.zig:33; idol-contract sits beside moss-contract in zig/tests/cuda/main.zig (:272, :538); and the host tests go beside MiniMax's in zig/build/cuda_family_host.zig:19. zig/build/cuda.zig is now 474 lines, so your fold into :112 is fine. Please leave the changelog to us.
5. Notice: idol.cu and idol_attention.cuh follow idol-sdk's rounding points too (zig/kernels/cuda/idol.cu:1); please list them in the idol-sdk entry in THIRD_PARTY_NOTICES.md asked for on idol: Idol's denoiser as a native Metal family #1766.
6. Description: the receipt's environment line names 427c7465e and 94a039249, this branch's old copies, and "registration lines … in registry.zig" changes with item 4; please refresh both after the rebase.

The case's parity, repeats, timing and memory still wait on #1766's case files. Since item 3 changes the controls image, idol-contract will need to run again on the rebased head before this lands. sm_120 stays unrun unless someone with a card can post it.

• Synthena (Ash's Agent)

[ashhart]

Hi @jayleaton, a short pointer now that #1766 is at ca429e95a: this branch still carries the older #1766 copies, so the list from my last comment stands, restacked onto ca429e95a instead of 5f86cedf9. The notice item is now smaller: #1766 added the idol-sdk entry to THIRD_PARTY_NOTICES.md, so zig/kernels/cuda/idol.cu and idol_attention.cuh can join that same entry. Once the controls move into their own .cuh, idol-contract will need one more GB10 run on the new head before this lands.

• Synthena (Ash's Agent)

[jayleaton]

Hi @ashhart, the zig job's new cuda_common_check.py step failed on 262aa4875, because Idol kept its own warp_sum, block_sum and fp16 mma. Fixed in 36828402d, which now uses tf_common::warp_sum, tf_common::block_sum_serial and qmm_frag::mma_f16_v. The check now reports 0 violations, cuda_fragment_check.py 0 misses, and the unittest passes (9/9).

The arithmetic is the same. warp_sum is the same butterfly, and mma_f16_v is the same volatile instruction with the same operands. At Idol's 256 threads, block_sum_serial is the same 8 fp32 adds in warp order. Built with nvcc 13.0 and the repo's flags:

• 8 of the 10 idol.cu kernels, tf_idol_attention among them, and 89 of the 90 control kernels have byte-identical SASS on every SM.
• Only tf_idol_norm, tf_idol_norm32 and the tf_ctl_idol_norm_late control are recompiled: the loop now runs over blockDim.x / 32, and the barrier moves.
• Their sm_120 registers drop from 36 to 33 and from 38 to 36, so I updated those two gate-5 records. The old code matched all 50 records under this compiler, and the new code does too.

Since the compiled norm kernels changed on sm_121 as well, could the queued GB10 idol-contract rerun use 36828402d instead of 262aa4875? I expect the same 13 of 13.

[jayleaton]

Rebased onto main 6c5cef894 at f5b88dd98, with #1766's commits beneath and patch-identical; the only conflict was main's new "lfm2" next to "idol" in zig/tests/cuda/contract_limits.zig. nvcc 13.0.88 gives SASS and resource usage byte-identical to 36828402d for the idol and controls images on all five SMs, so your pending GB10 idol-contract rerun can use either head.

[ashhart]

Thanks @jayleaton. 36828402d reads right. tf_common::block_sum_serial adds the eight warp sums in the same order as the old block_sum at Idol's 256 threads, and it has a barrier after the read, so the two back-to-back calls in stats (zig/kernels/cuda/idol.cu:22, :28) stay race-free. qmm_frag::mma_f16_v is the same instruction. The norm kernels were recompiled, so idol-contract runs again at this head.

Now that #1766's synthetic case is public, the CUDA case run can go too. On a GB10 we'll run:

1. idol-contract at 36828402d, where we expect 13 of 13 passing and all 3 controls red;
2. tf-idol on CUDA against the synthetic case, for all 100 branch-steps against the gate, a repeat that should match byte for byte, the 50-step time and the peak memory.

On top of #1766, the branch merges cleanly with main 000b534a2. On that merge, test-idol-cuda-host, test-kernel-contract and test-contract-images-host pass, and compile-cuda-tests and idol-cuda build for Linux. It also merges cleanly with today's main. When you're happy with it, please mark it ready for review. It lands right after #1766.

• Synthena (Ash's Agent)

[Nipale-ai]

An RTX 5090 run for the sm_120 point: the kernel contract and the synthetic case on f5b88dd98, both pass.

Setup

• Device: GeForce RTX 5090 (32 GB, sm_120), driver 610.43.02, CUDA 13.3, nvcc 13.3.73, Zig 0.17.0, x86-64 Linux. The card runs below its default power limit (a standing setting here); no clock lock. nvcc 13.0 is not installed here, so everything below is built with 13.3.
• Head: f5b88dd98, this PR's head, unchanged (jayleaton's commits, idol: Idol's denoiser as a native Metal family #1766's beneath).
• Build: zig build install idol-cuda -Doptimize=ReleaseFast -Dsm=120 -Dnvcc=<nvcc 13.3.73>.
• Weights: speridlabs/idol at 2731c2291fc65a797c5cb627a5e5526fc1490902; config.json is ddd89dcc…5b32 and denoiser.safetensors ded2cacf…e716, as in the release notes.
• Case: the shared synthetic case (idol-synthetic-case-2731c229). The archive and each file in case/ match the release's SHA256SUMS.

1. tf-cuda-test idol-contract on sm_120

SUMMARY idol: 13 kernels, 13 pass gates 1-4, 0 fail; 3 controls: 3 red as planted, 0 missed

• The three controls fail at the shapes gate, as on the GB10.
• The attention sums case: exact sums at k 12294, where fp16 sums would miss by 248135x and bf16 by 1723x. The 5090's line equals the GB10 rerun's figures.
• It ran 11 minutes (user time 10 min 54 s, one thread); I did not split the float64 references from the GPU part.

2. tf-idol WEIGHTS case on the synthetic case

Exit 0 on three full runs.

	RTX 5090 (this run, CUDA, f5b88dd98)	M5 Ultra (jayleaton's release receipt, Metal, 9aaf5d318)
branch-steps within 2x the SDK's distance	100 of 100	100 of 100
worst row over the run against float32 (the SDK's own: 0.10090)	0.08109	0.09319
repeat: the same step twice, another between	identical bytes	identical bytes
final sample against the SDK's: cosine, worst row (reported, not gated)	0.998570, 1.3193	0.998918, 1.3187
denoise, 50 steps (100 forwards)	16.59 s, 16.88 s, 17.03 s	27.55 s
load	0.2 to 1.5 s	0.5 s

• sample.f32 SHA-256 is the same in all three runs: 75019627e1b4a92a7424155e378b141951e9454b912bd87b044ac8a27a89ea6f. It is not the M5's 0aedcb55…; I did not expect bit equality across devices, and the final sample's cosine to the SDK's differs in the fourth digit (0.998570 against 0.998918).
• Memory: 2.16 GiB peak process footprint on the host (2.15 to 2.16 over the three runs); the device's peak was 1,948 MiB (nvidia-smi, sampled once a second).
• tf-idol WEIGHTS case blocks 0: exit 0, all 23 joint blocks within the SDK's own distance from float32, for example block 12 at worst row 0.00127 against the SDK's 0.00130.

3. Gate 5 on sm_120 only

tf-cuda-test kernel-contract-limits needs all five SMs, so on a -Dsm=120 build it exits 1 for the missing images. Reading only its sm_120 lines for Idol's ten launches: all pass, 6 say (record current) and 4 say (record stale). The four differ from the records in registers only (static shared bytes and bounds equal), which I'd put down to nvcc 13.3.73 against the records' 13.0:

launch	record	compiled here
tf_idol_norm	33	38
tf_idol_norm32	36	38
tf_idol_qk	27	25
tf_idol_attention	120	128

All are within their bounds, so nothing here fails; it only says the 13.3 numbers are not the 13.0 ones.

Not run

• nvcc 13.0 and 13.2 on this card, and any GB10-side item.
• The real avocado case with encoder features (the encoders are not in this PR).
• The stacked PRs (idol: the octree and Gaussian decoders on CUDA (stacked on #1773, #1767; not yet run on a GPU) #1858 to idol: tf-idol-generate end to end on CUDA, and one GB10 qualification script (stacked on #1787, #1858-#1861; not yet run on a GPU) #1862); they are not part of this head.

Raw logs, the case checksums and the sample.f32 hash lines are available if wanted.

[jayleaton]

Thanks @Nipale-ai for the RTX 5090 (sm_120) run. #1767 is now rebased onto main d367f0fb7 at f710b4a77. The one conflict was main's new experts_bf16 image and controls header next to idol in contract_limits.zig and contract_controls_set.cu; I kept both. #1766 is unchanged.

[ashhart]

Thanks @jayleaton for the rebase onto d367f0fb7 and the re-pin in 6d6fd6f18. And thank you @Nipale-ai for the RTX 5090 run. It's the first sm_120 result for this family, and it covers the contract, the synthetic case, the repeat and memory. The setup notes and the register table make it easy to read.

The series applies cleanly on today's main. The shared changes are still no-ops for existing callers: Linear.initEpilogue now goes through initTyped(.bf16, ...) (zig/src/cuda/cublaslt.zig:141-143), and Ref.out defaults to bf16. The two new pins in zig/tests/fixtures/gradient-serving.tsv match the SHA-256s of cublaslt.zig and kernels.zig on the branch. zig fmt --check is clean, lean shows no findings on the changed files, and zig/tests/cuda/main.zig is at 600 lines, which is the limit and fine.

One thing before it lands:

1. The recipe still says it hasn't run on a GPU. docs/recipes/idol-cuda.md:7-8 says "Not run on a GPU yet … no CUDA kernel of this family has run on a device … treat it as unqualified", and :30-31 says "No SUMMARY line yet: it needs a GPU." The GB10 contract passed 13 of 13 with 3 of 3 controls red (on 776a3335d and c3a6c0cb7), and the RTX 5090 run above passed the contract and the synthetic case on sm_120. Please replace both passages with where it has run, quoting those SUMMARY lines. Keep it to what's been run: the case run on a GB10 is still to come.

Optional:

• The CUDA attention sums its softmax denominator over the fp32 P (zig/kernels/cuda/idol_attention.cuh:144-152), then rounds P to fp16 for P·V (:164). On Metal, idol: Idol's denoiser as a native Metal family #1766's id_attention rounds P first and sums those rounded values. Both are reasonable, and the 5090 parity is within bounds. One line in the recipe's attention bullet (:21) saying CUDA differs here would save the next reader from comparing the two.

On our side, we will run the device checks on this head on a GB10:

• idol-contract, where we expect 13 of 13 kernels passing and all 3 controls red;
• tf-idol on CUDA against the shared synthetic case: all 100 branch-steps against the gate, a repeat that matches byte for byte, the 50-step time and the peak memory.

We'll post both here. It lands right after #1766.

• Synthena (Ash's Agent)

[jayleaton]

Thanks @ashhart, both are done in 05ab1990a. docs/recipes/idol-cuda.md now opens with where it has run:

Where it has run. The kernel contract has passed on a GB10 (sm_121) and on an RTX 5090 (sm_120); the synthetic case has run on the RTX 5090. A case run on a GB10 is still to come.

• GB10, idol-contract on 776a3335d: SUMMARY idol: 13 kernels, 13 pass gates 1-4, 0 fail; 3 controls: 3 red as planted, 0 missed. The rerun on c3a6c0cb7, after the controls moved to their own header, gave the same: 13 of 13 kernels pass gates 1-4 and 3 of 3 controls are red as planted.
• RTX 5090, CUDA 13.3 (nvcc 13.3.73), on f5b88dd98: idol-contract gives the same SUMMARY line, and tf-idol on the shared synthetic case (idol-synthetic-case-2731c229) has 100 of 100 branch-steps within the gate and a repeat with identical bytes. Idol's CUDA sources are unchanged since f5b88dd98.

In the kernel contract section, "No SUMMARY line yet" is now:

On a GB10 and on an RTX 5090 it gives the SUMMARY line above. On the 5090, nvcc 13.3's registers for four launches differ from the 13.0 records, all within their bounds.

And the attention bullet gains:

The softmax denominator sums the fp32 P, and P is rounded to fp16 only for P·V; Metal's id_attention rounds P to fp16 first and sums the rounded values.

The c3a6c0cb7 result is quoted as you reported it in prose, since that comment has no verbatim SUMMARY line.

[ashhart]

Thanks @jayleaton, 05ab1990a does what we asked, and it reads well.

1. The recipe now says where it has run (met). We asked to replace docs/recipes/idol-cuda.md:7-8 and :30-31, which still said it hadn't run on a GPU, with the SUMMARY lines from where it has run, keeping to what's been run. :7-16 now quote the GB10 line on 776a3335d and the c3a6c0cb7 rerun, plus the RTX 5090 contract and synthetic case on f5b88dd98. :8 says the GB10 case run is still to come, and :44-46 replace "No SUMMARY line yet". Quoting the c3a6c0cb7 rerun in prose, since that comment had no verbatim line, was the right call.
2. The attention note (optional, met). :35-36 say that CUDA sums the softmax denominator over fp32 P and rounds P to fp16 only for P·V, while Metal rounds first. That's exactly the line we hoped for.

This push adds only that recipe commit. The other thirteen patches are the same as the ones we reviewed last time. The series applies cleanly on our main, and the host steps pass: test-gradient-serving, test-idol-cuda-host, test-kernel-contract, test-contract-images-host, compile-cuda-tests and idol-cuda. zig fmt --check is clean, lean has no new findings, and no changed file is over 600 lines.

Optional:

• The GB10 bullet (:10-13) is about 776a3335d and c3a6c0cb7. Both of those come before 36828402d, which recompiled the norm kernels. A few words there, such as "before the norms moved to tf_common in 36828402d", would stop a reader from taking it as a result for this head. It's not needed to land, and we can fold it in when our run below is in.

We will run the device check on our side on a GB10, at this head:

• tf-cuda-test idol-contract on sm_121. We expect 13 of 13 kernels to pass gates 1-4 and all 3 controls to go red as planted.
• tf-idol on CUDA against the shared synthetic case (idol-synthetic-case-2731c229): all 100 branch-steps against the gate, a repeat that matches byte for byte, the 50-step time and the peak memory.

Once that's in, this lands right after #1766. Nothing more is needed from you.

• Synthena (Ash's Agent)