feat(ml): thread primitive + train_hslm windows on N threads, bit-identical at any count (#10876)

[gHashTag]

Closes #10876

What

Lever (a) of #10850: the 66 windows of one tr_train_step_k step run on N worker threads instead of one.

• Compiler (thin Rust): threadlocal var form and thread_spawn / thread_join / thread_cores builtins for --backend ml (bootstrap/src/compiler.rs, +40/-47, dead code deleted; FROZEN_HASH moved).
• Spec proving the primitive: specs/ml/thread_model.t27 (6 tests; parts on N threads sum bit-identically to the fixed-order sum, per-thread state starts at its initializer and does not leak).
• specs/ml/train_hslm.t27: cfg key threads (default = cores, threads=1 is the serial path). Each window writes its own gradient slot in win_g; slots are summed in window order, so the result does not depend on the thread count. nn/hslm state and the b3_* scratch arrays are threadlocal. New test one_two_and_three_threads_give_one_step_the_same_bits_at_speed_2_and_3.

Association change (stated, not hidden)

The old step kept one running sum over windows. The new step sums per-window slots in fixed window order. Floating-point addition is not associative, so:

• at k = 1 window per step the result is bit-equal to the old one;
• at k > 1 it differs from the old binary: state hash 240eb9893ad8eba1 (old) vs a9f594b18f9fc17e (new). Loss agrees to 4 digits (train 6.5916, valid 6.5637) and the .bin hash is the same.

The new value is the same for every thread count, which is the point: determinism no longer depends on the machine.

Lab proof (Railway lab, 48 cores, 8 steps, seed 1729, loaded: load average 55-72 during the runs, so these are floors)

State hash after 8 steps, sha256(ckpt_8.state)[:16]: a9f594b18f9fc17e at threads = 1, 2, 4, 8, 16, 24, 48, at speed 2 and at speed 3. Loss identical at every count.

threads	speed=2 tok/s	speed=3 tok/s
old binary	244.6	484.1
1	238.1 / 245.6	407.3
2	470.8	773.1
4	750.7	1312.7
8	1095.8	1820.7
16	1322.4	1960.9
24	1488.1	1730.1
48	1166.6	1033.8

Speed-up at 16 threads: 5.4x at speed 2, 4.1x vs the old speed-3 binary. Past 24 threads the lab is oversubscribed (load > 48) and throughput falls.

Known costs

• threads=1 at speed=3 is about 16% slower than the old binary (407 vs 484 tok/s): thread-local access on the b3_* scratch. At speed 2 it is level (238-246 vs 245). Not fixed here; a follow-up.
• A Debug build wrote undefined as 0xaa into .data, 1.1 GB for win_g. win_g now starts from a zero initializer (.bss), .data is 178 MB.
• specs/ml/eval_hslm.t27 is BLOCKED (use of undeclared identifier 'b3_q'). Inherited from spec(ml): train_hslm speed=3 -- one matrix over the 27 positions, bit-identical to speed=2 (Closes #10849) #10866 (it fails at this PR's base too), not caused here; follow-up.

Stack

This branch contains the commits of #10857 (issue #10850) until that merges; the diff of this PR alone is the thread work.

Tests on the lab: thread_model 6/6, train_hslm 30/30, t27t 17/17, ngram_hslm 8/8, all with 0 vacuous passes. Seals saved for ml_train_hslm and ml_thread_model.

🤖 Generated with Claude Code

[github-actions]

PR Dashboard

Generated at: 2026-10-19 06:38:10 UTC

• feat(igla): hslm_bwd_ptx kernel 3 -- split-K dW, 4.2-4.9x faster than kernel 2 at K=1782 (Closes #10895) #10914 feat(igla): hslm_bwd_ptx kernel 3 -- split-K dW, 4.2-4.9x faster than kernel 2 at K=1782 (Closes #10895) | Created: 2026-10-19 | Updated: 2026-10-19 | 32 failing / 41 total
• ledger(t27b): rows for five specs master carries unlisted (Closes #10912) #10913 ledger(t27b): rows for six specs master carries unlisted, max_not_pass 129 -> 130 (Closes #10912) | Created: 2026-10-19 | Updated: 2026-10-19 | 27 failing / 35 total
• spec(ml): train_hslm speed=4 -- int8 trit forward, error bounded against speed=2 (Closes #10875) #10908 spec(ml): train_hslm speed=4 -- int8 trit forward, error bounded against speed=2 (Closes #10875) | Created: 2026-10-19 | Updated: 2026-10-19 | 25 failing / 35 total
• spec(ml): eval_hslm mirrors train_hslm's 31 b3_* speed=3 scratch arrays (Closes #10906) #10907 spec(ml): eval_hslm mirrors train_hslm's 31 b3_* speed=3 scratch arrays (Closes #10906) | Created: 2026-10-19 | Updated: 2026-10-19 | 15 failing / 36 total
• analog M4d-b: the delay cell remembers its last edge -- repeated clkc falls within 1.8 % of ngspice PSP103, CK edges within 1 % (Closes #10848) #10898 analog M4d-b: the delay cell remembers its last edge -- repeated clkc falls within 1.8 % of ngspice PSP103, CK edges within 1 % (Closes #10848) | Created: 2026-10-19 | Updated: 2026-10-19 | 15 failing / 35 total
• feat(ml): thread primitive + train_hslm windows on N threads, bit-identical at any count (#10876) #10896 feat(ml): thread primitive + train_hslm windows on N threads, bit-identical at any count (#10876) | Created: 2026-10-19 | Updated: 2026-10-19 | 21 failing / 48 total
• t27b-loop skill: merge-tree against origin/master, FETCH_HEAD and stacked-row traps (Refs #10858) #10894 t27b-loop skill: merge-tree against origin/master, FETCH_HEAD and stacked-row traps (Refs #10858) | Created: 2026-10-19 | Updated: 2026-10-19 | 23 failing / 29 total
• analog: PSP103 drain and gate charge tables (tt, ss), the charge-aware table solver, the 500 fF inverter within 0.02 % of ngspice (Closes #10880) #10893 analog: PSP103 drain and gate charge tables (tt, ss), the charge-aware table solver, the 500 fF inverter within 0.02 % of ngspice (Closes #10880) | Created: 2026-10-19 | Updated: 2026-10-19 | 7 failing / 38 total
• feat(t27b): build record -- skip a build when nothing that feeds the object changed; --force #10890 feat(t27b): build record -- skip a build when nothing that feeds the object changed; --force | Created: 2026-10-19 | Updated: 2026-10-19 | 6 failing / 40 total
• ledger: the five rows master lacks -- sar_adc, sar_comparator_psp, mos_psp103_50c_data, worktrees/reap pass, ngram_hslm blocked; cap equals reality at 129 (Closes #10887) #10889 ledger: the five rows master lacks -- sar_adc, sar_comparator_psp, mos_psp103_50c_data, worktrees/reap pass, ngram_hslm blocked; cap equals reality at 129 (Closes #10887) | Created: 2026-10-19 | Updated: 2026-10-19 | 29 failing / 37 total
• t27t: profile the t27c CPU trainer -- 94.2% of a train_hslm step is matmul, 42.4% outer_add; four levers bounded (Closes #10850) #10857 t27t: profile the t27c CPU trainer -- 94.2% of a train_hslm step is matmul, 42.4% outer_add; four levers bounded (Closes #10850) | Created: 2026-10-19 | Updated: 2026-10-19 | 8 failing / 38 total
• spec(igla): hslm_bwd_ptx -- HSLM backward dX and dW as int8 tensor-core PTX written by t27 (Closes #10847) #10855 spec(igla): hslm_bwd_ptx -- HSLM backward dX and dW as int8 tensor-core PTX written by t27 (Closes #10847) | Created: 2026-10-19 | Updated: 2026-10-19 | 12 failing / 40 total
• igla cell 25 published: the whole-model HSLM trained by t27c --backend ml is on Hugging Face (Closes #10833) #10836 igla cell 25 published: the whole-model HSLM trained by t27c --backend ml is on Hugging Face (Closes #10833) | Created: 2026-10-19 | Updated: 2026-10-19 | 4 failing / 36 total
• t27f: idea factory cover still -- COVER=1 writes the reel's settled title frame, byte for byte (Closes #10828) #10830 t27f: idea factory cover still -- COVER=1 writes the reel's settled title frame, byte for byte (Closes #10828) | Created: 2026-10-19 | Updated: 2026-10-19 | 5 failing / 37 total
• t27f: idea factory pre-flight -- refuse an idea the font cannot draw before writing a frame #10822 t27f: idea factory pre-flight -- refuse an idea the font cannot draw before writing a frame | Created: 2026-10-19 | Updated: 2026-10-19 | 5 failing / 37 total
• t27o: Zig front end, step 1 -- zig_convert.t27 converts tier-1 fns and containers and translates the file's own tests into the oracle (Closes #10765) #10819 t27o: Zig front end, step 1 -- zig_convert.t27 converts tier-1 fns and containers and translates the file's own tests into the oracle (Closes #10765) | Created: 2026-10-19 | Updated: 2026-10-19 | 13 failing / 49 total
• t27f: idea factory batch -- five ideas cover all five presets, and the compiled binary agrees with the spec (Closes #10815) #10816 t27f: idea factory batch -- five ideas cover all five presets, and the compiled binary agrees with the spec (Closes #10815) | Created: 2026-10-19 | Updated: 2026-10-19 | 5 failing / 37 total
• t27f: idea factory picks its style preset from a hash of the idea, PRESET_AUTO (Closes #10803) #10806 t27f: idea factory picks its style preset from a hash of the idea, PRESET_AUTO (Closes #10803) | Created: 2026-10-19 | Updated: 2026-10-19 | 26 failing / 34 total
• spec(port): land the CR-CHAT-03 wave 3 batch -- 23 targets (Refs #8883) #10802 spec(port): land the CR-CHAT-03 wave 3 batch -- 23 targets (Refs #8883) | Created: 2026-10-19 | Updated: 2026-10-19 | 6 failing / 37 total
• feat(t27f): T27 OS home -- the 27-petal triangle menu as one spec #10783 feat(t27f): T27 OS home -- the 27-petal triangle menu as one spec | Created: 2026-10-19 | Updated: 2026-10-19 | 0 failing / 2 total
• queen: key_credit.t27 -- a leased contributor key earns credit only for .t27 that landed (Closes #10589) #10591 queen: key_credit.t27 -- a leased contributor key earns credit only for .t27 that landed (Closes #10589) | Created: 2026-10-18 | Updated: 2026-10-19 | 1 failing / 38 total
• claw: criteria.t27 + claw/criteria check -- commands settle the criteria they can (Closes #10549) #10572 claw: criteria.t27 + claw/criteria check -- commands settle the criteria they can (Closes #10549) | Created: 2026-10-18 | Updated: 2026-10-18 | 0 failing / 30 total
• queen: key_vault.t27 -- one key store, every consumer leases, review lanes first (Closes #10564) #10567 queen: key_vault.t27 -- one key store, every consumer leases, review lanes first (Closes #10564) | Created: 2026-10-18 | Updated: 2026-10-19 | 1 failing / 38 total
• ci: per-PR cancel groups on 9 uncancelled workflows, one provable paths filter (Closes #10552) #10557 ci: per-PR cancel groups on 9 uncancelled workflows, one provable paths filter (Closes #10552) | Created: 2026-10-18 | Updated: 2026-10-19 | 3 failing / 42 total
• claw: reel_still.t27 -- the TRI CLAW /mcp tool that renders one particle still #10548 claw: reel_still.t27 -- the TRI CLAW /mcp tool that renders one particle still | Created: 2026-10-18 | Updated: 2026-10-19 | 0 failing / 28 total
• fix(t27c): seal --verify reads the committed gen file; compile refuses a hand-written output (draft, needs owner-approved-foreign) #10481 fix(t27c): seal --verify reads the committed gen file; compile refuses a hand-written output (draft, needs owner-approved-foreign) | Created: 2026-10-18 | Updated: 2026-10-18 | 2 failing / 46 total
• fix(seal): refuse a seal file another spec owns, re-check the hash after the write, land it by rename #10464 fix(seal): refuse a seal file another spec owns, re-check the hash after the write, land it by rename | Created: 2026-10-18 | Updated: 2026-10-18 | 7 failing / 48 total
• igla ptx v2: cp.async ring, 128x128xBK64 tile, prmt-decoded W; within 1.33x of CUTLASS s8 (Refs #10404) #10456 igla ptx v2: cp.async ring, 128x128xBK64 tile, prmt-decoded W; within 1.33x of CUTLASS s8 (Refs #10404) | Created: 2026-10-18 | Updated: 2026-10-18 | 6 failing / 46 total
• check: refuse a module that declares a top-level name twice (Closes #10421) #10442 check: refuse a module that declares a top-level name twice (Closes #10421) | Created: 2026-10-18 | Updated: 2026-10-18 | 5 failing / 44 total
• demo: claw/sarif alert on a spec that does not parse (DO NOT MERGE, Refs #10356) #10381 demo: claw/sarif alert on a spec that does not parse (DO NOT MERGE, Refs #10356) | Created: 2026-10-18 | Updated: 2026-10-18 | 6 failing / 33 total
• t27b: repair crypto and rust_convert; distinguish PR ledger lag #9637 t27b: repair crypto and rust_convert; distinguish PR ledger lag | Created: 2026-10-15 | Updated: 2026-10-15 | 3 failing / 33 total
• spec(port): share tree helpers across eight gen_w ports (Closes #9633) #9636 spec(port): share tree helpers across eight gen_w ports (Closes #9633) | Created: 2026-10-15 | Updated: 2026-10-17 | 3 failing / 33 total
• Port fpga/verilog/gft_xorpercep_jtag.v (Verilog, 1 module) to specs/port/fpga/verilog/gft_xorpercep_jtag.t27 #9549 Port fpga/verilog/gft_xorpercep_jtag.v (Verilog, 1 module) to specs/port/fpga/verilog/gft_xorpercep_jtag.t27 | Created: 2026-10-14 | Updated: 2026-10-14 | 4 failing / 29 total
• Port gHashTag/trios:crates/trios-git/src/stage.rs (Rust, 2 functions) to specs/port/trios/crates/trios-git/src/stage.t27 #9548 Port gHashTag/trios:crates/trios-git/src/stage.rs (Rust, 2 functions) to specs/port/trios/crates/trios-git/src/stage.t27 | Created: 2026-10-14 | Updated: 2026-10-14 | 4 failing / 29 total
• Port gHashTag/trios:crates/trios-rainbow-bridge/src/bridge.rs (Rust, 5 functions) to specs/port/trios/crates/trios-rainb #9485 Port gHashTag/trios:crates/trios-rainbow-bridge/src/bridge.rs (Rust, 5 functions) to specs/port/trios/crates/trios-rainb | Created: 2026-10-14 | Updated: 2026-10-14 | 4 failing / 29 total
• Port scripts/gen_w734.py (Python, 3 functions) to specs/port/scripts/gen_w734.t27 #9484 Port scripts/gen_w734.py (Python, 3 functions) to specs/port/scripts/gen_w734.t27 | Created: 2026-10-14 | Updated: 2026-10-14 | 4 failing / 29 total
• Port scripts/gen_w726.py (Python, 3 functions) to specs/port/scripts/gen_w726.t27 #9463 Port scripts/gen_w726.py (Python, 3 functions) to specs/port/scripts/gen_w726.t27 | Created: 2026-10-14 | Updated: 2026-10-14 | 3 failing / 29 total
• Port gHashTag/trinity:src/tri/sacred_alu.zig (Zig, 4 functions) to specs/port/trinity/src/tri/sacred_alu.t27 #9427 Port gHashTag/trinity:src/tri/sacred_alu.zig (Zig, 4 functions) to specs/port/trinity/src/tri/sacred_alu.t27 | Created: 2026-10-13 | Updated: 2026-10-13 | 5 failing / 29 total
• Port gHashTag/trios:crates/trios-phd/src/bin/rag_search.rs (Rust, 3 functions) to specs/port/trios/crates/trios-phd/src/ #9426 Port gHashTag/trios:crates/trios-phd/src/bin/rag_search.rs (Rust, 3 functions) to specs/port/trios/crates/trios-phd/src/ | Created: 2026-10-13 | Updated: 2026-10-13 | 4 failing / 29 total
• Port gHashTag/trinity:src/vibeec/spec_loader.zig (Zig, 2 functions) to specs/port/trinity/src/vibeec/spec_loader.t27 #9407 Port gHashTag/trinity:src/vibeec/spec_loader.zig (Zig, 2 functions) to specs/port/trinity/src/vibeec/spec_loader.t27 | Created: 2026-10-13 | Updated: 2026-10-13 | 4 failing / 29 total
• Port gHashTag/trinity:src/background_agent/config.zig (Zig, 5 functions) to specs/port/trinity/src/background_agent/conf #9345 Port gHashTag/trinity:src/background_agent/config.zig (Zig, 5 functions) to specs/port/trinity/src/background_agent/conf | Created: 2026-10-13 | Updated: 2026-10-13 | 5 failing / 29 total
• Port gHashTag/trios:crates/trios-cli/src/cmd/train.rs (Rust, 6 functions) to specs/port/trios/crates/trios-cli/src/cmd/t #9321 Port gHashTag/trios:crates/trios-cli/src/cmd/train.rs (Rust, 6 functions) to specs/port/trios/crates/trios-cli/src/cmd/t | Created: 2026-10-13 | Updated: 2026-10-13 | 4 failing / 29 total
• Port gHashTag/trinity:fpga/openxc7-synth/spi_flash_master.v (Verilog, 1 module) to specs/port/trinity/fpga/openxc7-synth #9263 Port gHashTag/trinity:fpga/openxc7-synth/spi_flash_master.v (Verilog, 1 module) to specs/port/trinity/fpga/openxc7-synth | Created: 2026-10-12 | Updated: 2026-10-12 | 7 failing / 29 total
• spec fix: 4 specs in specs/tri/net the reference blocks -- rewrite them in t27 (1 of 1) #9214 spec fix: 4 specs in specs/tri/net the reference blocks -- rewrite them in t27 (1 of 1) | Created: 2026-10-12 | Updated: 2026-10-12 | 5 failing / 29 total
• Port gHashTag/trios:crates/trios-agent-harness/rings/AH-00/src/lib.rs (Rust, 3 functions) to specs/port/trios/crates/tri #9200 Port gHashTag/trios:crates/trios-agent-harness/rings/AH-00/src/lib.rs (Rust, 3 functions) to specs/port/trios/crates/tri | Created: 2026-10-12 | Updated: 2026-10-12 | 6 failing / 29 total
• Port bootstrap/src/compiler.rs lines 30977-31050 (Rust, 6 functions) to specs/port/bootstrap/src/compiler/LinkSegment_te #9178 Port bootstrap/src/compiler.rs lines 30977-31050 (Rust, 6 functions) to specs/port/bootstrap/src/compiler/LinkSegment_te | Created: 2026-10-11 | Updated: 2026-10-12 | 6 failing / 29 total
• Port gHashTag/trinity:src/storm/links/vibee.zig (Zig, 1 function) to specs/port/trinity/src/storm/links/vibee.t27 #9177 Port gHashTag/trinity:src/storm/links/vibee.zig (Zig, 1 function) to specs/port/trinity/src/storm/links/vibee.t27 | Created: 2026-10-11 | Updated: 2026-10-12 | 5 failing / 32 total
• Port scripts/gen_w706.py (Python, 3 functions) to specs/port/scripts/gen_w706.t27 #9150 Port scripts/gen_w706.py (Python, 3 functions) to specs/port/scripts/gen_w706.t27 | Created: 2026-10-11 | Updated: 2026-10-12 | 7 failing / 30 total
• spec fix: 1 specs in specs/queen the reference blocks -- rewrite them in t27 (1 of 1) #8879 spec fix: 1 specs in specs/queen the reference blocks -- rewrite them in t27 (1 of 1) | Created: 2026-10-11 | Updated: 2026-10-18 | 4 failing / 30 total
• Port gHashTag/trinity:src/vibeec/updated_codex.zig (Zig, 5 functions) to specs/port/trinity/src/vibeec/updated_codex.t27 #8874 Port gHashTag/trinity:src/vibeec/updated_codex.zig (Zig, 5 functions) to specs/port/trinity/src/vibeec/updated_codex.t27 | Created: 2026-10-11 | Updated: 2026-10-12 | 6 failing / 30 total

Summary

Status	Count
Total Open PRs	50
PRs with Failing Checks	46
PRs with All Checks Green	4
READY	3
FAILING	46
PENDING	0
NO CHECKS YET	0

These columns do not partition: 3 + 46 + 0 + 0 = 49, and there are 50 open PRs. A PR is being counted twice or not at all.

Seal Status

• ⚠️ STALE -- sha256(compiler.rs)=509db03c8f4d != manifest seal=87e5cbd3ad94.
  The committed NMSE numbers were certified against an older compiler.rs.
  Run scripts/reseal-check.sh locally for the two-step reseal command (advisory; not a merge gate).