Skip to content

GPU boundary, recurrence-derived prefill/decode, site geometry, and mapal fmt - #1

Open
LessComplexity wants to merge 322 commits into
mainfrom
site-geometry-latest
Open

GPU boundary, recurrence-derived prefill/decode, site geometry, and mapal fmt#1
LessComplexity wants to merge 322 commits into
mainfrom
site-geometry-latest

Conversation

@LessComplexity

@LessComplexity LessComplexity commented Oct 29, 2026 •
edited
Loading

Copy link
Copy Markdown
Owner

Summary

Three optimization sprints (cpu-gpu-boundary, recurrence-precision, site-geometry), the measure-only-what-changed tooling, and the mapal fmt source formatter. Plain summary with results: summary.md.

Ratios are Mapal time ÷ the fastest competitor on the same program and GPU (cuBLAS / cuBLASLt / CUTLASS / cuDNN via PyTorch / hand-written CUDA / PyTorch / torch.compile); below 1 = Mapal faster. GPUs: RTX 4090 / RTX 5090 / H100 / RTX 3090.

workload before now
causal attention 2.87 / 2.74 / 5.62 / 2.89 0.18 / 0.26 / 0.25 / 0.14
attention prefill 5.57 / 5.15 / 8.62 / 5.51 0.18 / 0.22 / 0.26 / 0.15
gen (prefill + decode in one natural loop) — 0.17 / 0.24 / 0.25 / 0.15
decode 1.96 / 1.34 / 3.58 / 1.89 1.01 / 0.82 / 1.51 / 1.00
multi-channel convolution ~2000–5000 (ran on the CPU) 1.38 / 0.89 / 1.18 / 1.43
elementwise chain 2.59 / 5.33 / 4.75 / 2.27 0.46 / 0.45 / 1.01 / 0.61

What changed

  • CPU↔GPU boundary: work stays on the GPU (loop scalars as kernel arguments, maps with inner sums as device cells, kernels resolved and memory allocated before the clock).
  • Recurrence: prefill and decode are derived from the loop's dataflow (trip_split), not hardcoded; many threads per cell; convolution as an implicit matrix multiply (radix_plan); compute types in the language, fold: tf32 | f32 { … } (ADR-0041).
  • Site geometry: per-GPU geometry table with --geom= overrides (ADR-0034), measure-only-changed tooling (benches/gpu/changed.py, result cache, pre-rent preflight, search driver); mixed sizing rules recorded as missing facts (CLAUDE.md rule "a mixed result is a missing fact").
  • Tests: differential harness bounded to the core budget and emulated runs time-boxed (the laptop hangs); the interpreter oracle builds optimized in tests.
  • Formatter: mapal fmt [--check], editor format on save, docs.

Test plan

  • Full workspace tests on the integrated branch: 1296 passed, 0 failed (site-geometry judge)
  • CPU-only emission byte-identical to the previous release branch at each integrate step
  • GPU value gates on all four GPUs for every changed program (sweep evidence under benches/results/evidence/)
  • Formatter: mapal-syntax 215 passed, mapal-cli format 8 passed; reformatted bench programs give identical interpreter output
  • Review the per-GPU stopgap rows listed in docs/performance/evidence/site-geometry.md (next sprint replaces them with facts)

Raw run output (profile archives, box logs, raw results; 277 files) is not in git: it is in the release evidence-raw-2026-10-29, indexed by benches/results/evidence/RAW-manifest.tsv (path, size, sha256, source commit). The branch history was rewritten to drop it (701 → 311 MB of new files).

Next steps and handoff: docs/next-session.md, session logs in docs/sessions/.

🤖 Generated with Claude Code

https://claude.ai/code/session_018Mdxerz1UZ8nCSLz1kHkJc

LessComplexity and others added 30 commits October 21, 2026 06:50
@LessComplexity @claude
results(s54): epilogue on rtx-3090, a100 (SXM4), rtx-5090 — different…
1ba0df7 
…ial 206/206 each, ON vs off-epilogue on 9 probes

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
fix(gpu): column-major store rasters the grid down the rows
595eef2 
The revised build (2d70913, rowfold-<part>.txt + -confirm.txt) still paid for
the column-major store on the tile kernel: H100 PCIe 8.80 vs 7.29 ms exact,
4.50 vs 4.15 ms contract at 4096 (traced). The lane order was not the cost: a
wave of blocks with one row block wrote 4096 column pieces of 512 B, 16 KB
apart. Under the column-major store the grid's fastest axis now runs down the
rows (ctaid.x = row block, the host swaps gx/gy), the warp rule one level up.
CPU emulation 120/120 (rect 128x384, 512, f64, five parts, both faces, arms).

Results: rowfold ON is stable+ on 3090, 5090, A100 in two sessions each
(rowsum 2.1-9.2x, softmax_rows 1.03-1.27x after the nocapture fix).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
results(s54): epilogue on rtx-4090 — differential 206/206, ON vs off-…
ed44414 
…epilogue on 9 probes

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
results(s54): l2-raster confirm/run2 sessions (A100 x2, RTX 5090, H10…
1a64001 
…0 PCIe); table.py traced kernel column

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
perf(gpu): streams single-wave gate, pin only repeated crossings; ses…
b232186 
…sion 1 on four parts

Session 1 (rtx-4090, rtx-3090, rtx-5090, h100-pcie; differential 52/52,
52/52, 74/74, 74/74; trap-before-site all pass):
- streams opened by the last-wave form at 4096 (3090, 5090) and
  2048/4096 (H100): one stream 8-37% faster. The gate now opens only grids
  that fit one wave together (b_a/w_a + b_b/w_b <= 1).
- pinned single crossings lose on every part: registering 0.06-0.09 ms/MB
  (0.3-0.5 at 1 MB) vs 0.002-0.08 ms/MB saved per crossing (box copy_rate).
  Only a copy inside a loop is pinned now; the RTX 5090 loop measured
  +17-19% with pinning.
- runtime: one device wait before the frees (was one per free), events
  reused per recording stream.
- twin probe: indep without its join (corner reads only), the device-only
  shape streams can pay on.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
results(s54): arch-select on five parts; Ampere open rows closed, H10…
4f1a98f 
…0 warp-fill

Five sessions of the shipped table (arch-select-<part>.txt), each paired with the
S54 session on the same table (3090 -verify2/-verify, A100 -verify3, 4090 base and
-confirm, H100 -verify, 5090 -verify and base/confirm for the FMA classes).
Differential 273-294 pass / 0 fail per part, trap-before-site 11-12 / 0.

Rows (rungs.py: a decided row flips only on a stable sign in every session):
- rtx-3090: warp-fill and unroll OFF under-filled contract; min-blocks ON on the
  FMA tile, OFF on TF32 (Face::Tf32); warp and idx32 ON under-filled contract.
- a100: idx32 ON filled contract; wide-fma stays OFF (with idx32 ON it costs);
  ldmatrix OFF under-filled TF32.
- h100-pcie: warp-fill OFF under-filled contract (disjoint, both sessions).
- rtx-4090, rtx-5090: no change (no stable sign across 3-4 sessions).

rungs.py: null-corrected kernel arms; flip rule vs plan rule (--new); mixed held.
box_run.sh: one discarded warm-up bench (the first bench ran 3.6% off its repeat).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
perf(gpu): streams only over device-only sites; confirm sessions 3090…
99df048 
…/5090, A100 session 1

Two sessions per part (v1 deferred the chained copy-outs to their host
reader; v2 added one device wait per free run and events reused) measured
every probe with a host join (indep, diamond, mlp2 at 512) 2-38% slower on
streams than on one stream, on all five parts: the host sits between the
chains and waits for a transfer five to ten times one kernel. Streams now
open only when every straight-line site's output stays on the device; the
copy-out deferral is gone. twin (indep without its join) on the RTX 5090:
+10-18% by minima with streams (one session). px_fan_128: the
differential of two chains and a cross-chain wait under this rule.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
fix(gpu): raster — re-derive the block origin at the epilogue
d09050c 
ctaid reads are rematerialised where they are used; the remapped origin was
carried across the panel loop instead (3090 exact at 8192 spilled 24 bytes at
the 128-register cap, l2-raster-rtx-3090-run2.txt). The epilogue now reads ctaid
again through an asm sideeffect move (CSE would fold a readnone read back into
the head's) and remaps it there. Row-major text unchanged.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
results(s54): epilogue on h100-pcie — differential 206/206; a contend…
973b215 
…ed host (process walls 0.9-1.5 s, iter outliers to 60 ms)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
results(s54): R5 on five parts, two boxes each — retain-resident off …
5ff4fa0 
…on Ada and Blackwell

decode_1024_256, every arm interleaved, both faces pooled, the loop
(traced wall − load) as the decision metric (context creation varies
±200 ms per process on the H100s and no arm changes it), process wall
beside it. Differential 71/71 (session 1) and 38/38 (session 2) on
3090, A100, 4090, H100 PCIe, 5090.

  rung              3090        A100        4090        H100        5090
  retain-resident   1.05/1.04   1.09/1.04   1.03/0.98   1.06/1.09   1.01/1.02
  delta-h2d         1.14/1.21   1.14/1.10   1.17/1.20   1.40/1.10   1.24/1.28
  lane-stride       1.57/1.45   1.32/1.42   1.61/1.59   1.40/1.30   1.58/1.58
  all three (base)  2.21/2.29   1.78/1.99   2.20/2.35   2.70/1.38   2.31/2.47

GPU_RUNGS_OFF: retain-resident off on the 4090 (sign changes) and the
5090 (≤ 1%): the whole K/V site costs what its row does there. delta-h2d
and lane-stride ship on all five.

Decode wall 0.51 (3090, 4090, 5090), 0.86–1.01 (A100), 0.96 (H100) s,
half the s54-base placement in the same run — and still 8–28× slower in
the loop than the eager cuBLAS loop with or without a hand-written cache
(21–86 ms). Mapal's per-step device work is now four launches and three
1 KB copies; the rest is host time, where the loop state copies the
carried 1 MB x four times a step.

Harness: box_run five decode reps; vast_run gives up a rental created
stopped (its GPU taken by another renter) at once.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
results(s54): epilogue h100-pcie confirm (different box, host 670790)…
d3a30e3 
… — 206/206; probe_arms.py (a placement rung on the probes, vs the base session)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
revert(gpu): raster epilogue re-derivation — measured worse on the 50…
397b7a6 
…90 FMA tile

The epilogue re-read of ctaid (d09050c) cleared the exact tile's spill but gave
the RTX 5090 FMA tile 16 bytes of spill: ungated raster there went from
0.980/0.959 (8192, run2/run3) to 1.031/1.008 OFF/ON (run4/run5), and TF32 at
4096 to 1.018*/1.021*. The head-only remap (3901fc2) is what ships; its PTX is
byte-identical to the bundle the run2/run3 sessions measured. Results of the
post-fix-2 sessions stay as measured (l2-raster-*-run4/5.txt).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
feat(gpu): l2-raster per-part rows — raster OFF on the 3090, gate OFF…
9a1bbf4 
… on the 5090, l2-window OFF but H100 contract

GPU_RUNGS_OFF rows from the l2-raster sessions (two boxes each for every row
that ships or withdraws an optimization):
- rtx-3090 raster OFF: the derived grouping at 8192 measured no stable sign
  (exact 0.986/1.029, contract 0.992/1.017, TF32 0.991/0.956 OFF/ON).
- rtx-5090 l2-gate OFF: the ungated grouping (h = 16) measured 2-5% faster at
  8192 on all three faces in both sessions; no effect at 4096.
- l2-window OFF on the 3090 (no probe fits its L2), A100, 4090 (it costs up to
  7% on contract), 5090, and the H100 exact face; ON on H100 contract
  (mlp2_2048_c traced kernel 1.012/1.016).
Snapshots are back to s54-base's text (the H100 exact face and the 4090 ship no
window); the l2-window placement test forces the rung on.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
results(s54): rowfold v3 on 3090/5090/H100 (two sessions), A100 (one)…
f89429f 
…; layout OFF rows

rowfold-<part>.txt / -confirm.txt: the raster-swap build (595eef2). rowfold
ships ON (every part, iter rule, two sessions where done): rowsum 2.3-9.2x the
host fold, softmax_rows 1.0-1.3x. rowfold-layout ships OFF on 3090, 5090,
H100 PCIe (device clock: column-major tile store + coalesced fold vs
row-major + strided fold, 4096 losing in both sessions; H100 8-22%).
probe_rungs.py also states the rule on the traced kernel sums.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
perf(gpu): ship streams on A100 and RTX 5090; pinned and pdl off — fo…
b6ad59e 
…ur sessions, five parts

The gate is the last-wave form again, now over device-only sites only: on
the A100 at 2048 two grids of ~1.2 waves run in 3 waves instead of 4,
1.27-1.43x faster in three sessions (twin, two independent sites). The
first build's losses with this form were all host-joined programs, which
no longer open streams.

GPU_RUNGS_OFF (S54 R7): streams OFF on the RTX 3090 (twin n=512 no effect
in two sessions), RTX 4090 (0.94-0.99x) and H100 PCIe (not established:
one noisy PCIe box, one SXM session inside the noise floor); ON on the A100
(n=512 1.04-1.11x, three sessions) and RTX 5090 (n=512 1.01-1.11x, three
sessions; n=4096 1.04-1.05x). pinned OFF on all five: one crossing never
repays registration, and in the loop the sign follows the host's pageable
rate (RTX 5090 +17-19% on two boxes, -5% on one at pageable = pinned).
pdl OFF on H100 and RTX 5090: no stable sign.

Differential N/N on every session (52/52 ... 74/74, last 32/32 and 58/58 on
the shipped table), trap-before-site all pass. Shipped host and device text
on the RTX 3090, RTX 4090, H100 PCIe and H100 is byte-identical to s54-base
(360/360); every OFF arm is (450/450). Results: streams-tables.md.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
feat(gpu): 5090 gate row on the contract faces only; final-table 5090…
30b2bc9 
… session

The final-table session (third 5090 box) measured the shipped grouping at 8192
TF32 1.014x (cuBLAS-normalized), FMA 1.003x, exact 0.999x; with run2/run3 TF32 is
faster in all three sessions, exact is not, so the gate stays on the exact face.
The multi-computation probes at 4096 (every pt_*_4096 now groups on the 5090)
measured traced kernel time 0.996-1.008 OFF/ON: no effect. wide_8192x1024x1024
puts the 5090's grouped order in the box differential (46/46).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
results(s54): l2-raster tables; recovered 3090 run4; vast_run keeps f…
bc98b79 
…etched results

l2-raster-tables.md: the shipped table, the derivation against the measurement,
grouped vs row-major in every session, l2-window per part, kernels and chain
probes vs the libraries, sessions. vast_run.sh: the post-profile re-fetch no
longer overwrites a fetched results file with an empty one when the box drops
(it cost the 3090 run4 file; recovered from the box's stdout).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
results(s54): epilogue decided ON on five parts; tables, ptxas regist…
031daad 
…ers, plan §4.1 measured

ON vs off-epilogue (own clock, medians, 9 probes x 2048/4096 x both faces):
3090 1.55-3.24, A100 1.20-4.72, 4090 1.73-4.85, 5090 1.69-7.00, H100 PCIe
0.26-3.84 and 1.10-20.44 on two contended hosts (the 0.26 row: host stalls in
3 of 5 ON reps, minima 2.38); every ratio of minima >= 1.29. No GPU_RUNGS_OFF
row: the rung ships ON everywhere (filled class). Differential 206/206 on every
box. H100 exact `epi`: 132 vs 128 registers, one resident block, kernel
13-21% slower ON (both sessions). Mapal still loses the whole program to the
best library on every probe (contract 4096: 0.44-0.94, loop 0.07-0.42):
cuMemFree/alloc in the clock and the loop's host state copies remain (item 6).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
results(s54): devmap + cell-law on five parts — ON everywhere, differ…
18d620d 
…ential 72/72

devmap-<part>.txt (rtx-3090, a100 SXM4-80GB, rtx-4090, h100-pcie, rtx-5090) and
-confirm.txt (A100 SXM4-40GB, H100 PCIe, other boxes): arms on / off-devmap /
off-cell-law interleaved, both faces, 2048 and 4096, library chains gated 54/54.
Differential 72/72 per part, 16/16 per confirm box.

devmap off/on (Mapal's clock): site probes 1.30-4.76x (H100 to 15.5x), loop
1.76-2.53x, ew 4.4-121x. cell-law on ew at 4096: 1.11-9.32x (confirm 1.31-2.20),
residual no effect. ew's one kernel vs the library's fused kernel at 4096:
0.93-3.66x (A100/H100 compute-bound on two IEEE divides per cell). Losses vs the
best library chain at 4096 contract: 0.35-0.91 on site probes, loop 0.08-0.32
(cuMemAlloc/Free in the clock; five 64 MB host memcpys per loop step on the
device-carried state). Both rungs ship ON on all five parts: no GPU_RUNGS_OFF
row; the evidence is on the GpuRungs fields, pinned in
gpu_rungs_are_the_measured_table. devmap-tables.md has every row.

Also: cell_grid unit test; ARCHITECTURE R2 as built, cell-law in the rung list.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
feat(gpu): rowfold per-arch table; results on five parts, two session…
90e3254 
…s each

rowfold ships ON on every part (own clock, two boxes each): rowsum
2.3-9.9x the host fold, softmax_rows 1.0-1.3x. rowfold-layout ships only
where the device clock measured it faster in both sessions: the 4090's
contract face (4096 2.507 vs 2.564, 2.605 vs 2.640 ms) and the A100's exact
face (15.479 vs 15.568, 19.911 vs 20.027 ms); OFF elsewhere (the column-major
tile store costs 0-6% on GDDR parts, 7-22% on H100; the strided fold it
replaces costs only 0.01-0.1 ms at 4096). GPU_RUNGS_OFF rows cite the files;
rowfold_rungs_are_the_measured_table pins them.

4090 and A100 confirm sessions ran on hosts with 5-10 Gbps links
(vast_run.sh START_TRIES: slow image pulls left five 4090 hosts loading).
rowfold-tables.md: the tables. plan-s54 4.4 outcome, ARCHITECTURE R0'.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
docs(gpu): rowfold rung docs carry the final per-part outcome; fusion…
88f6333 
… report at the final build

680/680 OFF/base emissions still byte-identical (five profiles, every face).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
docs(nvptx): plan-s54 4.7 — what l2-raster measured on five parts
ccef7c5 
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
results(s54): integrated 3090 / H100 sessions of the arch-select tabl…
b1c683a 
…e; 3090 under-filled FMA takes one step

integrated-rtx-3090.txt (new table vs the old as arm `old`, same box): exact n=1024
1.15, contract n=2048/4096/8192 1.06/1.01/1.07, TF32 n=1024 1.13 (= cuBLAS) — but
contract n=1024 0.94: the four single flips that each paid on warp-fill's tile lose
together on the S52 tile (on it idx32 ON 0.944, unroll OFF 0.931). The 3090
under-filled FMA class ships warp-fill OFF alone; idx32 and warp rows restored,
unroll stays ON. integrated-h100-pcie.txt: contract n=1024 1.047, elsewhere the same
text. Differential 292/292 on both, trap-before-site 14/14.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
merge s54/arena into s54/arch-select (R6: one device arena from death…
1e369ab 
… points)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
merge s54/l2-raster into s54/arch-select (grid raster, L2 gate, L2 wi…
1775d02 
…ndow)

Unions: GpuRungs fields / NAMES / rung(), resident.rs (arena + l2_window), the
resident tests, bundle.sh arm lists (HOST_RUNGS gains arena and arch-rungs),
box_run.sh (probe_arms file, else every built off arm), fusion_report.py columns.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
merge s54/retain-resident into s54/arch-select (R5 on the resident bu…
3562187 
…ffer, delta h2d, lane stride)

resident_plan takes the rung set (devgen, arena, R5); the arena slot follows the
retain filter; ResidentPlan carries both arena and delta; rows and tests unioned.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
merge s54/rowfold into s54/arch-select (row folds as device sites, re…
0a55d1e 
…ader-chosen store order)

The column-major store's block ids live in block_ids (raster skips it: that order is
its own); emit_site_resident takes the slot list (rowfold) and the L2 window
(l2-raster); l2_window's next reader reads DevSite slots; the rowfold test counts
arena members as allocations.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
merge s54/epilogue into s54/arch-select (R1: per-cell consumers in th…
ed2009f 
…e site's store)

Integration: kernel_text takes both the column-major store and the epilogue; the
device unit skips grouped sites' plain kernels; emit_site_resident launches
launchv for a group (no L2 window on it yet), call's launch otherwise; epilogues
skip row-fold sites; carries count member outputs as device-made; the arena keeps
a host slot for a stored member a host reader copies out; launch_l2 and launchv
both live in mapal_gpu.c; fusion_report reads EMIT_FLAGS and FLAGS.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
merge s54/devmap into s54/arch-select (R2: per-cell maps as device ke…
2f67cff 
…rnels)

Integration: the plan's device outputs are sites', cells' and stored epilogue
members'; cells skip epilogue members (R1 first) and read DevSite slots; slot_use
and uses see both; carries and produced_on_device count all three; the arena is
laid out after the cells; one launchv declaration (it was emitted twice); the
fusion report counts cell kernels, L2 windows, raster groups and the arena.
Each item's tests pin its own rule (devmap: epilogue and arena off; epilogue:
devmap off; rowfold: epilogue and devmap off).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
@LessComplexity @claude
merge s54/streams into s54/arch-select (R7: chain streams, pinned cop…
d5d2556 
…ies, pdl)

Integration: a site launch takes launchv on its chain's stream when one is open
(and for an epilogue group, waiting on every operand's stream, the group's
complete operands included); copy-outs go through the stream-aware d2h and pin;
chains() runs over tile DevSites and refuses when a launch the chains do not order
(a fold, an unchained site, a cell kernel) reads a chained output; the arena and
the L2 window stand down in a function that opens streams; one launchv in the
runtime (the stream-aware one) and one declaration; the streams init flag follows
a launch that names a stream, not any launchv. Each item's tests pin its own rule.
rungs.py: S54's host-read rungs key fill Any. cargo test --workspace --release:
87 suites, 1171 passed.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01MS7jWhiYwxiwxQkzPsqJeE
LessComplexity and others added 29 commits October 26, 2026 02:39
@LessComplexity @claude
geom(B2): fill-wave and row-split — one fill decision for tile sites
f9c4790 
fill-wave: a tile grid with fewer warps than Gpu::fill_warps cuts its
free fold into the fewest partitions that reach it, counted at the
tile's own residency (not grid-share's cap at the unsplit grid); the
split kernel's panels are sized for the split grid (C9 attn P.V, C10
GEMV, H100 rt_mlp2_1024). row-split: the retained one-row pass takes
the chooser (atomic / stream-K, workspace on a free order) instead of
one 256-thread block (D1). OFF arms byte-identical to the base on 191
programs x 4 GPUs x 3 faces; ON changes the default face only (91
programs). Property test site_geometry.rs over testgen geom programs
(emulated ON = OFF); intent rows fill-wave, row-split.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013YjipM6wchPQHbhsabHXJk
@LessComplexity @claude
geom(R1): fit-first — reshape a tall thread grid before shrinking it
4eddb74 
An FMA tile whose thread grid outgrows the rows (ty*tm > rows) takes
row-fit's reshape (ty/2, tx*2, threads kept) when that alone gives every
SM a block, before the fill-the-SMs step halves its threads: convmc on
the 4090 784 x 64 -> 196 x 256 threads (same tm x tn, same warps). OFF
arm byte-identical to B2; ON changes 5 programs (convmc, convmc_tf32,
m16 x 18432 skinny). Property test over testgen radix + wide products
(emulated ON = OFF, 19 moved); splitk.rs pins the S62/S63 chooser with
fill-wave / fit-first off; intent row fit-first.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013YjipM6wchPQHbhsabHXJk
@LessComplexity @claude
geom(B4): GEMM tile candidates as table entries
0bffbd2 
Geom gains the FMA tile's tm, tn, warps, bk, stages and the wgmma tile's
wg_n, wg_groups (0 = derived). gpu_tile_forced rebuilds the derived tile
around the forced values (block tile square-most, residency from held
registers, the deepest panel the share holds for the forced stages) and
falls back to the derived tile where the site cannot take it. knobs.json
lists per-arch candidates: H100 8x16 + cp.async stages (C6a), 5090 8x16
(C6b tile), K=1024 panel depths / warps (C6c), H100 wgmma n / groups (T1).
Default emission byte-identical to the previous commit. Property test:
one knobs.json candidate per generated product, emulated = table row
(22 moved a kernel, 21 valued).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013YjipM6wchPQHbhsabHXJk
@LessComplexity @claude
geom(B3): tf32-radix — the mma.sync tile for a tf32 radix site
1f932b9 
A mixed-radix site whose products admit tf32 takes the mma.sync tile:
digit-affine operands copied element by element (cp.async 4 B, the k
digits added per panel), the digit-affine output stored cell by cell;
wgmma (swizzled 16 B panels) is not used for radix sites. convmc_tf32:
FFMA tile -> kind=mma on all four parts; PTX assembles with llc for
sm_89/90a/120/86 (values are the box's gate: the emulator runs no mma).
Also: fill-wave keeps the stream-K ladder (stream-K weighed against the
S63 count), so H100 sites that took stream-K keep it. All four new
rungs OFF = base byte for byte (191 programs x 4 GPUs x 3 faces).
Property test tf32_radix_counted_on_unseen_programs (14 of 128 generated
convolutions reach the tensor tile, all assemble); intent row tf32-radix.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013YjipM6wchPQHbhsabHXJk
@LessComplexity
Merge branch 'site-geometry-measure' into site-geometry
bf6f103
@LessComplexity @claude
fix(geom): review findings on site-geometry-sites
f504761 
- forced tile knobs (tm/tn/warps/bk/stages) replace the derivation step they
  name inside gpu_tile_split / gpu_tile_async; gpu_tile_forced deleted, one
  residency helper (tile_capacity, cell, mma, wgmma, FMA tile)
- forced wg_n / wg_groups the site cannot take fall back to the derived wgmma
  tile, never another kernel class
- fit-first reshapes only where its block tile divides the site
- fill-wave keeps the re-sized split launch only when bk' | k/S
- the fill key (gpu_tile) reads the seed row: a --geom arm moves the launch,
  never which GPU_RUNGS_OFF rows apply; naive launch takes its own grid
- tf32-radix ships OFF (cited row) on parts no box has value-gated
- plan lines list row-split / fill-wave / fit-first / tf32-radix in rungs_on,
  split lines carry fill_warps

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013YjipM6wchPQHbhsabHXJk
@LessComplexity @claude
fix(measure): review findings on site-geometry-measure
ff750a8 
- runtime/harness stamp (crates/mapal-rt + box harness): a change re-measures
  every row; Mapal cache rows keyed by stamp|nvcc|driver
- default-face-only changes of differential programs are value-gated on the
  box (bundle.sh DEFAULT -> <prog>_d, box_run.sh D2 bound with tiles/<p>.kmax;
  gen_tiles.py writes K for every loop-free program); no K -> changed_bundle
  stops (UNCHECKED)
- diff and changed_bundle rebuild the emitter / runtimes first
- search pick counts tile_bench GATE-FAILED candidates as gate failures
- box_run drops only the failing arm and times the rest
- CPU-only changes go to one box (--cpu-box / CPU_BOX)
- parts with GPU_RUNGS_OFF rows probe --gpu-on=<r> arms
- changed_bundle passes SUITE_ARMS=0 (arms from arms.tsv, cpu where changed)
- cell library cache key (cell= carries mm_); preflight checks non-empty
- evidence/intent: new rungs in KERNEL_RUNGS; row-split / fit-first /
  fill-wave intents keyed on plan facts (fill_warps, retained row)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013YjipM6wchPQHbhsabHXJk
@LessComplexity @claude
fix(changed): runtime stamp over what a row's time runs through (runt…
e275f17 
…ime, timing harnesses)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013YjipM6wchPQHbhsabHXJk
@LessComplexity @claude
fix(gen_tiles): K through a map's output (a fold over a computed array)
434c571 
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013YjipM6wchPQHbhsabHXJk
@LessComplexity @claude
fix(evidence): atomic adds with memory-order qualifiers (newer llc); …
8e9b432 
…smoke plumbing (TORCH passthrough, --rungs on-arms)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013YjipM6wchPQHbhsabHXJk
@LessComplexity @claude
fix(gpu): an epilogue group's ordered launch keeps its own tile besid…
6a37cd9 
…e a fill-wave split

The smoke box (rtx-4090): deep_rmlp*_d default face hit an illegal memory
access — launch_in re-sized the tile for the split combine, and the ordered
(epilogue) kernel was launched with the split tile's shared memory (43008 B
for a kernel sized for 86016 B). The ordered path now takes the unsplit
launch. Regression: site_geometry epilogue_launch_keeps_its_own_tile_beside_a_split
(every launch's smem equals its kernel's plan); testgen products feed a
bias + relu epilogue a third of the time. Default-face differential gate:
magnitude oracle (D2 absolute) for sums-of-products programs.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013YjipM6wchPQHbhsabHXJk
@LessComplexity @claude
fix(bundle): cells keep their reassociation faces under SUITE_ARMS=0 …
e07a812 
…(box_run's split-aware cell section keys on _r)

The smoke box timed the cells in the single-arm section, which reads an
atomic split kernel's launch as an ordered one (values wrong, no time).

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013YjipM6wchPQHbhsabHXJk
@LessComplexity @claude
bench(site-geometry): rtx-4090 smoke — value gates, ON/OFF of the new…
a9c16a9 
… rungs

smoke-run1: default-face differential caught the epilogue launch bug (fixed
6a37cd9). smoke: differential 51/51 (13 default faces within D2), 30 probe
programs value-gated, probe ON/OFF per rung (onoff.tsv). smoke-cells: 13
cells via tile_bench --splitk, 0 gate failures, 65 split programs within D2.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_013YjipM6wchPQHbhsabHXJk
@LessComplexity @claude
wip(site-geometry): sweep evidence collected before a checkpoint (lap…
3526d41 
…top closed) — 4 archs, pass2 results, evidence, ncu/traces; summary not written yet

Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01LMRRJUctgafB7GWyDbYrmt
@LessComplexity @claude
docs(rules): a mixed ON/OFF result is a missing fact, not an OFF row
f794649 
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018uFHDgKZReAjnzRpYcT1NY
@LessComplexity
bench(site-geometry): sweep h100 value-gated re-run, search picks, re…
9d550f5 
…port

Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>

Claude-Session: https://claude.ai/code/session_018uFHDgKZReAjnzRpYcT1NY
@LessComplexity
bench(site-geometry): h100 search run, picks with shipped values, report
037b855 
Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>

Claude-Session: https://claude.ai/code/session_018uFHDgKZReAjnzRpYcT1NY
@LessComplexity @claude
perf(gpu-rungs): site-geometry judge — tf32-radix ON on 5090/3090, ro…
9583482 
…w-split and fit-first stopgaps

tf32-radix: the forced-ON arm is within its bound on the 5090 and 3090 boxes and
1.22x / 1.26x faster (mma.sync replaces the FFMA radix tile): rows lifted.
row-split OFF on 4090/5090/3090: there the one-row pass takes the atomic combine
(naive atom.add kernel, 2 more zero launches per trip) and loses 14-26 %; H100's
workspace split wins 1.41-1.56 (open G2). fit-first OFF on 5090/H100/3090, where
it lost every measured program (open G4). Emission: 49 rows move, each equal to
its measured arm.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018uFHDgKZReAjnzRpYcT1NY
@LessComplexity @claude
fix(box_run): the cuBLAS timing section's _exact face reads the _x mo…
0035aad 
…dule

Since ADR-0040 the unsuffixed mm_<cell> is the default face (order free), so the
in-order check of that section flagged every cell as MAPAL VALUES WRONG (101 /
89 / 57 / 39 lines in the site-geometry sweep, on the shipped arm too) while the
gated section's _x legs are bit-exact.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018uFHDgKZReAjnzRpYcT1NY
@LessComplexity @claude
perf(gpu-rungs): keep tf32-radix OFF on the 5090 — the fitted FFMA ti…
1f1869e 
…le is faster

Composed in the same 5090 run: the mma.sync tile 0.07003 ms vs the FFMA radix
tile with fit-first off (now shipped there) 0.06797 ms (null 1.0 %). The 3090
lift stands (0.1327 vs 0.1477 ms). Emission: 49 rows move, each equal to its
measured arm.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018uFHDgKZReAjnzRpYcT1NY
@LessComplexity @claude
docs(evidence): site-geometry end check, gap table, missing facts
750c97f 
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018uFHDgKZReAjnzRpYcT1NY
@LessComplexity @claude
docs(readme): GPU workloads table from the site-geometry sweep
a25a029 
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018uFHDgKZReAjnzRpYcT1NY
@LessComplexity @claude
merge: recurrence-precision-latest (CLAUDE.md mixed-result rule) into…
ba4f295 
… site-geometry-latest

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018uFHDgKZReAjnzRpYcT1NY
@LessComplexity @claude
results(cache): measurement cache from the site-geometry sweep (keyed…
991e59a 
… by emission hash; lets the next sweep carry unchanged rows)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018uFHDgKZReAjnzRpYcT1NY
@LessComplexity @claude
docs: site-geometry session log and handoff
e3d4710 
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018uFHDgKZReAjnzRpYcT1NY
@LessComplexity @claude
docs: next-session — fix the handoff text (site-geometry standing and…
1dbe09e 
… missing facts)

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018uFHDgKZReAjnzRpYcT1NY
@sapiritur @claude
docs: summary.md and handoff for the pull request to main
3fc0750 
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018Mdxerz1UZ8nCSLz1kHkJc
@sapiritur @claude
feat(fmt): mapal fmt source formatter (syntax, CLI, editors) and docs
d4e34c6 
From the 2026-10-28 source-formatter session (docs/sessions/2026-10-28-source-formatter*.md):
the mapal-syntax formatter, `mapal fmt [--check]`, nvim / VS Code format on save,
the OpenSpec source-formatting spec, the CLAUDE.md formatting rule, README and
contributor docs, docs/architecture-explorer.html; two bench programs
reformatted (interpreter output byte-identical, fmt --check clean).
Verified here: mapal-syntax 215 passed / 0 failed, mapal-cli format 8 / 0,
workspace tests compile.

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018Mdxerz1UZ8nCSLz1kHkJc
@sapiritur @claude
evidence: raw run output moves to a GitHub release; manifest, index, …
3966475 
…archive script, ignore rules

Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_018Mdxerz1UZ8nCSLz1kHkJc
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

None yet

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

2 participants

@LessComplexity @sapiritur