feat(igla): hslm_bwd_ptx kernel 3 -- split-K dW, 4.2-4.9x faster than kernel 2 at K=1782 (Closes #10895)

[gHashTag]

Closes #10895

Follows #10855 (merged, fc611e7): kernel 3 builds on kernels 1 and 2. Rebased onto master after the squash; the diff is only the split-K change (spec, its regenerated Rust and seal, and driver.rs, net 0 lines).

CI note: t27b-native-ratchet, Corpus ratchet (expected-failure ledger), duplicate-bodies, the orphan ceiling and the types Range check were already red when #10855 merged (non-required). This PR does not bless the ledgers: that needs a full t27b lab corpus run and would also bless three unrelated specs (sar_adc, ngram_hslm, reap) that sit UNLISTED today, so it belongs in a ledger-steward change.

What

specs/igla/coder/hslm_bwd_ptx.t27 gains kernel 3, hslm_dw_splitk_s8: kernel 2 (dW = X^T x dY) split along k.

• grid z = slices; slice z sums k in [zper32, (z+1)per32), per = ceil(ceil(k/32)/nz) read from %nctaid.z in the kernel.
• Each slice ends with red.global.add.s32 of its i32 partials into C. The host zeroes C before each launch (zero_first()), so the result equals kernel 2's bit for bit.
• grid_z() aims at SPLIT_TARGET_BLOCKS = 64 blocks, never more slices than k-steps, and drops slices the rounding would leave empty (no idle block).
• A warp covers one 16-byte chunk and 32 k rows (kernel 2: 8 lanes per row), so the transposed byte stores spread over 8 banks. Measured effect alone: small (0.034 -> 0.0325 ms at the old target).
• Kernels 0, 1, 2, their integer CPU references and their text pins are unchanged.

Measured (RTX 4090, RunPod, t27c ptx-run --iters 500 --stride 7, mismatches=0 everywhere; a stride-1 run at 256x736x1782 also reports mismatches=0)

m x n x k	kernel 2 ms	kernel 3 ms	CUTLASS s8 ms	kernel 3 / CUTLASS
256 x 256 x 1782	0.0747	0.0153	0.0066	2.3x
256 x 736 x 1782	0.0752	0.0180	0.0068	2.6x
736 x 256 x 1782	0.0746	0.0180	0.0068	2.6x
256 x 256 x 27	0.0037	0.0072	-	-

Kernel 3 is 4.2-4.9x faster than kernel 2 at K=1782 (was 11x behind CUTLASS, now 2.3-2.6x).

The goal was not reached

The issue asked for within about 1.5x of CUTLASS s8. It is at 2.3-2.6x. Honest notes:

• The ptx-run time is wall time per call and for kernel 3 includes the zero-fill launch. At K=27 (one step, nothing to compute) kernel 3 costs 0.0072 ms against 0.0037 ms for kernel 2: that is two launches against one, about 3.7 us each from the CPU. The CUTLASS figures are GPU-event times without that floor, so the table is not apples to apples; about 7 us of the 15-18 us is launch cost.
• What remains on the GPU is the red traffic (every slice adds the full 128x128 tile into C) and the 128x128 tile walk with u8 transposed shared stores. The sweep of SPLIT_TARGET_BLOCKS (8..256) is U-shaped with its floor at 48..64: fewer slices starve the SMs, more slices pay in reds.
• Next steps if the owner wants the last 1.5x: fuse the zero fill (e.g. a graph or a cooperative first-slice store), or a second-pass reduction instead of reds. Both are separate changes.

Foreign-line accounting (shrink-only budget)

bootstrap/src/host/driver.rs: net 0 lines. The cuMemsetD32_v2 declaration joins an existing extern line; the launch closure zeroes C when zero_first(kernel) and passes grid_z(...) as the launch's z dimension, in place. bootstrap/gen/rust/... and the seal are regenerated on the Railway lab (test-report 26/26, 0 vacuous; gen-rust; seal --save).

Test plan

• t27c test-report specs/igla/coder/hslm_bwd_ptx.t27 on the lab: 26/26, 0 vacuous
• gen-rust and seal regenerated on the lab, hashes verified after transfer
• t27c ptx-run --kernel 3 on RTX 4090: mismatches=0 for 256x256, 256x736, 736x256 at K=1782 and 256x256 at K=27

🤖 Generated with Claude Code

[github-actions]

PR Dashboard

Generated at: 2026-10-19 07:15:23 UTC

• feat(t27os): PTY server that owns the agent terminals and restores them (Refs #10782) #10934 feat(t27os): PTY server that owns the agent terminals and restores them (Refs #10782) | Created: 2026-10-19 | Updated: 2026-10-19 | 33 failing / 38 total
• model routing: Sonnet 5.5 by default for every agent, the orchestrator included; Opus only on evidence (Closes #10921) #10929 model routing: Sonnet 5.5 by default for every agent, the orchestrator included; Opus only on evidence (Closes #10921) | Created: 2026-10-19 | Updated: 2026-10-19 | 16 failing / 35 total
• fpga(ax7203): the SAR converter of the golden deck on the AX7203 -- 31 conversions, UART receipts equal to the spec's codes, a one-literal negative control (Closes #10910) #10919 fpga(ax7203): the SAR converter of the golden deck on the AX7203 -- 31 conversions, UART receipts equal to the spec's codes, a one-literal negative control (Closes #10910) | Created: 2026-10-19 | Updated: 2026-10-19 | 8 failing / 42 total
• t27o: Zig front end step 2 -- --check on .zig, sweep lists .zig, imported consts, std.math.pow (Closes #10823) #10918 t27o: Zig front end step 2 -- --check on .zig, sweep lists .zig, imported consts, std.math.pow (Closes #10823) | Created: 2026-10-19 | Updated: 2026-10-19 | 9 failing / 42 total
• 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 | 29 failing / 37 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 | 3 failing / 36 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 | 13 failing / 50 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 | 6 failing / 39 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
• 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 | 19 failing / 39 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 | 4 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 | 4 failing / 37 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 | 4 failing / 37 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 | 5 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 | 1 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
• Port gHashTag/trios:crates/trios-chat/rings/CR-CHAT-02/src/ratchet_replay_window_guard.rs (Rust, 2 functions) to specs/p #8873 Port gHashTag/trios:crates/trios-chat/rings/CR-CHAT-02/src/ratchet_replay_window_guard.rs (Rust, 2 functions) to specs/p | Created: 2026-10-11 | Updated: 2026-10-12 | 9 failing / 31 total
• Port scripts/tri-doc-sync.py (Python, 1 function) to specs/port/scripts/tri-doc-sync.t27 #8850 Port scripts/tri-doc-sync.py (Python, 1 function) to specs/port/scripts/tri-doc-sync.t27 | Created: 2026-10-10 | Updated: 2026-10-12 | 7 failing / 33 total
• gen-verilog: an entry parameter named clk/rst_n/en is declared once; a void body's last call stays a statement (Closes #5649, Closes #5650, Closes #5651) #8844 gen-verilog: an entry parameter named clk/rst_n/en is declared once; a void body's last call stays a statement (Closes #5649, Closes #5650, Closes #5651) | Created: 2026-10-10 | Updated: 2026-10-19 | 7 failing / 47 total
• Port gHashTag/trios:crates/trios-mesh-node/src/persist.rs (Rust, 5 functions) to specs/port/trios/crates/trios-mesh-node #8834 Port gHashTag/trios:crates/trios-mesh-node/src/persist.rs (Rust, 5 functions) to specs/port/trios/crates/trios-mesh-node | Created: 2026-10-10 | Updated: 2026-10-12 | 5 failing / 32 total
• Port gHashTag/trios:crates/trios-chat/rings/CR-CHAT-06/src/agent_context_window_budget_guard.rs (Rust, 1 function) to sp #8822 Port gHashTag/trios:crates/trios-chat/rings/CR-CHAT-06/src/agent_context_window_budget_guard.rs (Rust, 1 function) to sp | Created: 2026-10-10 | Updated: 2026-10-12 | 6 failing / 32 total

Summary

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

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

Seal Status

• ⚠️ STALE -- sha256(compiler.rs)=88b8cb73ceda != 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).