Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
324 commits
Select commit Hold shift + click to select a range
ed44414
results(s54): epilogue on rtx-4090 — differential 206/206, ON vs off-…
LessComplexity Oct 21, 2026
1a64001
results(s54): l2-raster confirm/run2 sessions (A100 x2, RTX 5090, H10…
LessComplexity Oct 21, 2026
b232186
perf(gpu): streams single-wave gate, pin only repeated crossings; ses…
LessComplexity Oct 21, 2026
4f1a98f
results(s54): arch-select on five parts; Ampere open rows closed, H10…
LessComplexity Oct 21, 2026
99df048
perf(gpu): streams only over device-only sites; confirm sessions 3090…
LessComplexity Oct 21, 2026
d09050c
fix(gpu): raster — re-derive the block origin at the epilogue
LessComplexity Oct 21, 2026
973b215
results(s54): epilogue on h100-pcie — differential 206/206; a contend…
LessComplexity Oct 21, 2026
5ff4fa0
results(s54): R5 on five parts, two boxes each — retain-resident off …
LessComplexity Oct 21, 2026
d3a30e3
results(s54): epilogue h100-pcie confirm (different box, host 670790)…
LessComplexity Oct 21, 2026
397b7a6
revert(gpu): raster epilogue re-derivation — measured worse on the 50…
LessComplexity Oct 21, 2026
9a1bbf4
feat(gpu): l2-raster per-part rows — raster OFF on the 3090, gate OFF…
LessComplexity Oct 21, 2026
f89429f
results(s54): rowfold v3 on 3090/5090/H100 (two sessions), A100 (one)…
LessComplexity Oct 21, 2026
b6ad59e
perf(gpu): ship streams on A100 and RTX 5090; pinned and pdl off — fo…
LessComplexity Oct 21, 2026
30b2bc9
feat(gpu): 5090 gate row on the contract faces only; final-table 5090…
LessComplexity Oct 21, 2026
bc98b79
results(s54): l2-raster tables; recovered 3090 run4; vast_run keeps f…
LessComplexity Oct 21, 2026
031daad
results(s54): epilogue decided ON on five parts; tables, ptxas regist…
LessComplexity Oct 21, 2026
18d620d
results(s54): devmap + cell-law on five parts — ON everywhere, differ…
LessComplexity Oct 21, 2026
90e3254
feat(gpu): rowfold per-arch table; results on five parts, two session…
LessComplexity Oct 21, 2026
88f6333
docs(gpu): rowfold rung docs carry the final per-part outcome; fusion…
LessComplexity Oct 21, 2026
ccef7c5
docs(nvptx): plan-s54 4.7 — what l2-raster measured on five parts
LessComplexity Oct 21, 2026
b1c683a
results(s54): integrated 3090 / H100 sessions of the arch-select tabl…
LessComplexity Oct 21, 2026
1e369ab
merge s54/arena into s54/arch-select (R6: one device arena from death…
LessComplexity Oct 21, 2026
1775d02
merge s54/l2-raster into s54/arch-select (grid raster, L2 gate, L2 wi…
LessComplexity Oct 21, 2026
3562187
merge s54/retain-resident into s54/arch-select (R5 on the resident bu…
LessComplexity Oct 21, 2026
0a55d1e
merge s54/rowfold into s54/arch-select (row folds as device sites, re…
LessComplexity Oct 21, 2026
ed2009f
merge s54/epilogue into s54/arch-select (R1: per-cell consumers in th…
LessComplexity Oct 21, 2026
2f67cff
merge s54/devmap into s54/arch-select (R2: per-cell maps as device ke…
LessComplexity Oct 21, 2026
d5d2556
merge s54/streams into s54/arch-select (R7: chain streams, pinned cop…
LessComplexity Oct 21, 2026
dbc2a00
results(s54): the A100 table verify; Phase-A verify sessions named ar…
LessComplexity Oct 21, 2026
ead399f
results(s54): the integrated build's static fusion report, five parts…
LessComplexity Oct 21, 2026
b995cf9
fix(gpu): an epilogue member is born at its site in the arena
LessComplexity Oct 21, 2026
b8399aa
wip(s54): partial integrated prefixes + plan edits left when the work…
LessComplexity Oct 21, 2026
4eed06f
docs(s54): session log, next-session handoff, research notes
LessComplexity Oct 21, 2026
2c38f21
docs(s54): full handoff — per-arch rung table, verifier conflicts, lo…
LessComplexity Oct 21, 2026
a613f53
results(s54): l2-raster H100 PCIe run3 (third box): l2-window contrac…
LessComplexity Oct 21, 2026
4db1f48
results(s54): l2-raster RTX 4090 run1 (third box): gate kept, l2-wind…
LessComplexity Oct 21, 2026
44972af
docs(s55): record consolidated worktree handoff
LessComplexity Oct 21, 2026
e8d069c
fix: restore validation and guarded folds
LessComplexity Oct 21, 2026
4c93c8a
test: gate GPU smoke and retain run failures
LessComplexity Oct 21, 2026
823f64a
fix: reject failed GPU timing processes
LessComplexity Oct 21, 2026
b4152c5
perf: resolve disputed GPU selections
LessComplexity Oct 21, 2026
0ca0507
docs(gpu): record completed frozen validation
LessComplexity Oct 21, 2026
627a17c
fix(rewrite): rebuild acyclic loop route work
LessComplexity Oct 21, 2026
a14feda
docs(gpu): preserve ordering proof and handoff
LessComplexity Oct 21, 2026
3a0d435
fix(gpu): order grouped member consumers
LessComplexity Oct 21, 2026
ce0c201
test(gpu): preserve parallel validation
LessComplexity Oct 21, 2026
bc7af76
test(gpu): record mixed small-M results
LessComplexity Oct 22, 2026
29ad618
docs(gpu): preserve parallel validation prep
LessComplexity Oct 22, 2026
dece72f
chore(perf): organize reports and raw evidence
LessComplexity Oct 22, 2026
d80d8f5
perf: record competitive library comparisons
LessComplexity Oct 22, 2026
50a57e5
fix(gpu): select policy for emitted arithmetic
LessComplexity Oct 22, 2026
293237a
docs: record reviewed observer repair and handoff
LessComplexity Oct 22, 2026
4e1102c
docs: record executed library kernel evidence
LessComplexity Oct 22, 2026
8f1ae02
feat(cli): mapal build|run|dump-ir|test
LessComplexity Oct 22, 2026
b29061f
bench(gpu): runtime measurement harness — request clock, retention lo…
LessComplexity Oct 22, 2026
11c56ef
feat(gpu): cell-fill and asyncstage rungs
LessComplexity Oct 22, 2026
d380843
feat(gpu): ordered tile geometry rungs from shape + device facts
LessComplexity Oct 22, 2026
6c3162c
feat(gpu): opt-in reassoc face, deterministic split-K
LessComplexity Oct 22, 2026
56f90f3
bench(gpu): S62 runtime raw results, RTX 4090 and RTX 5090 session 1
LessComplexity Oct 22, 2026
ff052fc
bench(gpu): S62 runtime raw results, RTX 4090 and RTX 5090 confirm se…
LessComplexity Oct 22, 2026
ba7d850
bench(gpu): submission stage 2, generated vs direct launch enqueue
LessComplexity Oct 22, 2026
9bb2f07
bench(gpu): S62 runtime raw results, stage-2 submission rows and RTX …
LessComplexity Oct 22, 2026
fbf2c75
bench(gpu): S62 runtime raw results, H100 PCIe session 1, H100 SXM co…
LessComplexity Oct 22, 2026
79e5869
bench(gpu): S62 runtime raw results cut by the 19:25 ssh drop
LessComplexity Oct 22, 2026
dad4a56
bench(gpu): S62 runtime raw results, the 19:27 completion sessions cu…
LessComplexity Oct 23, 2026
daeef18
feat(gpu): host CPU profile in GPU builds, per-site placement
LessComplexity Oct 23, 2026
be89bc7
perf(gpu): asyncstage measured on four parts
LessComplexity Oct 23, 2026
f8d7272
perf(gpu): record S62 hostopt/place runs on four parts
LessComplexity Oct 23, 2026
e5d6337
bench(gpu): S62 runtime raw results, the 20:00 completion sessions an…
LessComplexity Oct 23, 2026
2733746
bench(gpu): S62 runtime raw results, retained vs regenerated on one b…
LessComplexity Oct 23, 2026
2b25583
perf(gpu): S62 runtime verdict — allocation latency, submission, rete…
LessComplexity Oct 23, 2026
cc49cae
fix(gpu): gate fill-rows and panel-regs on geometry facts
LessComplexity Oct 23, 2026
c862247
perf(gpu): ship the S62 geometry rungs per part from the v2 A/B
LessComplexity Oct 23, 2026
5dff8df
test(gpu): 4090 64^3 snapshots back to the S53 launch; cite traced pr…
LessComplexity Oct 23, 2026
580f672
perf(gpu): S62 geometry raw results v1-v4 on 5090/4090/H100 PCIe/3090
LessComplexity Oct 23, 2026
40f858c
Merge branch 's62/geometry' into s62-integrated
LessComplexity Oct 23, 2026
ee5770d
Merge branch 's62/asyncstage' into s62-integrated
LessComplexity Oct 23, 2026
c0eaa84
Merge branch 's62/splitk' into s62-integrated
LessComplexity Oct 23, 2026
b94177c
Merge branch 's62/cpugpu' into s62-integrated
LessComplexity Oct 23, 2026
52c57e0
Merge branch 's62/runtime' into s62-integrated
LessComplexity Oct 23, 2026
0601a93
Merge branch 's62/cli' into s62-integrated
LessComplexity Oct 23, 2026
9a9cc20
fix(gpu): cell-fill split partials land in their workspace slab
LessComplexity Oct 23, 2026
6a0e3a0
perf(gpu): S62 integration measured on four parts, per-part rung rows
LessComplexity Oct 23, 2026
b762272
perf(gpu): RTX 3090 executed-code and usage evidence, S62 integrated
LessComplexity Oct 23, 2026
999b89e
perf(gpu): RTX 5090 executed-code and usage evidence, S62 integrated
LessComplexity Oct 23, 2026
98a782a
perf(gpu): RTX 4090 executed-code and usage evidence, S62 integrated
LessComplexity Oct 23, 2026
569b5e3
perf(gpu): H100 PCIe executed-code and usage evidence, S62 integrated
LessComplexity Oct 23, 2026
3ec494e
docs: reconcile S62 integration, gap map and handoff
LessComplexity Oct 23, 2026
5c06a78
feat(float): the default float face — fma, order free, bounded (ADR-0…
LessComplexity Oct 23, 2026
c0103c5
perf(gpu): S63 default-face rungs measured on four parts, OFF rows
LessComplexity Oct 23, 2026
068852d
wip(evidence): S64 intent registry, checker, plan/host evidence lines
LessComplexity Oct 23, 2026
b2645e4
feat(evidence): checker fixes from the first static pass, box arm wiring
LessComplexity Oct 23, 2026
b88e209
perf(evidence): static SASS / ptxas / remarks for the suite and the r…
LessComplexity Oct 23, 2026
59a90b3
perf(evidence): RTX 4090 box run — suite, cells, MLPs, re-judge arms,…
LessComplexity Oct 23, 2026
d90eff1
perf(evidence): RTX 5090 box run — suite, cells, MLPs, re-judge arms,…
LessComplexity Oct 23, 2026
8d758a9
perf(evidence): H100 PCIe box run — suite, cells, MLPs, re-judge arms…
LessComplexity Oct 23, 2026
a1944a8
perf(evidence): RTX 3090 box run — suite, cells, MLPs, re-judge arms,…
LessComplexity Oct 23, 2026
7b6a889
perf(gpu): warp-fill class split on H100 / 3090 (Under -> Partial), S…
LessComplexity Oct 23, 2026
38789b4
docs(evidence): S64 end check, report and handoff; 3090 warp-fill spl…
LessComplexity Oct 23, 2026
de7226f
wip(s65): loop state by reference via byref.rs (reference; to be fold…
LessComplexity Oct 23, 2026
c66c50e
feat(s65): rule R in the existing storage layer; byref.rs deleted
LessComplexity Oct 24, 2026
a7dfff7
feat(s65): compact host allocator; device leg of the property test
LessComplexity Oct 24, 2026
2c5c0e8
bench(s65): static evidence on four archs, RTX 4090 box results, move…
LessComplexity Oct 24, 2026
183032e
bench(s65): RTX 5090 and RTX 3090 box results; movement table grouped…
LessComplexity Oct 24, 2026
a972f76
bench(s65): H100 PCIe box results; instances after the run
LessComplexity Oct 24, 2026
9ca6c6b
perf(gpu): retain-resident ON on the 4090 and 5090 (S65 rows removed)
LessComplexity Oct 24, 2026
78742f4
fix(evidence): the null arm's runtime movement is predicted by the sh…
LessComplexity Oct 24, 2026
c8aae2f
docs(s65): evidence README section and next-session handoff
LessComplexity Oct 24, 2026
249f199
bench: commit the generated algorithm programs under benches/algorithms
LessComplexity Oct 24, 2026
8b499e6
wip(natural-forms): first natural-form generators, natural_forms test…
LessComplexity Oct 24, 2026
30f661f
ir: footprint_plan — every consumer's read edges as one query
LessComplexity Oct 24, 2026
52cb157
natural-forms: piecewise map law over a stepped cell set; column writ…
LessComplexity Oct 24, 2026
c94929a
llvm: generic fusion rule (fuse), window rung, strided 2-D sites
LessComplexity Oct 24, 2026
6b04916
bench(natural-forms): idiomatic decode and softmax generators, regene…
LessComplexity Oct 24, 2026
802fb15
ir/llvm: window row folds; bench: cross-algorithm chain workloads
LessComplexity Oct 24, 2026
e3a5eca
llvm: window row folds behind the window rung
LessComplexity Oct 24, 2026
9c3b969
test: default-face on == off in the chain property test; rustfmt the …
LessComplexity Oct 24, 2026
467c865
bench: the chain workloads in the benches/algorithms snapshot (drift-…
LessComplexity Oct 24, 2026
4186713
natural-forms: twin property test over 256 programs; range-map rung; …
LessComplexity Oct 24, 2026
e203c5d
test(natural-forms): the twin sweep's reach floor at its default size…
LessComplexity Oct 24, 2026
89d5262
natural-forms: a delta wider than its array is the whole upload
LessComplexity Oct 24, 2026
ef7a5f6
testgen(natural-forms): twin range writes may run into the next row
LessComplexity Oct 24, 2026
26f2f68
test(natural-forms): the range-map edge programs at 16-wide rows (the…
LessComplexity Oct 24, 2026
14779ad
natural-forms: the emitter asks each map's piecewise law once
LessComplexity Oct 24, 2026
bade438
merge natural-forms into generic-fusion
LessComplexity Oct 24, 2026
8d1f66b
fix(fuse): a unit kernel carries its cells' nested bodies
LessComplexity Oct 24, 2026
2924d14
bench: regenerate the snapshot — attention's exp in natural form
LessComplexity Oct 24, 2026
a55ead9
evidence: window and range-map intent rows
LessComplexity Oct 24, 2026
fbdec8b
fix(by-ref): an in-place range map owns its bytes, as an Update does
LessComplexity Oct 24, 2026
8dac5b5
evidence(fusion-natural-forms): static evidence of the integrated build
LessComplexity Oct 24, 2026
12727a4
evidence(fusion-natural-forms): CPU-only builds identical to both parts'
LessComplexity Oct 24, 2026
d2dd44c
test(natural-forms): a fold over a range stays a device cell, in a un…
LessComplexity Oct 24, 2026
47a5048
docs(evidence): fusion-natural-forms test results at d2dd44c
LessComplexity Oct 24, 2026
dc0372e
bench: library legs for the generic-fusion chains
LessComplexity Oct 24, 2026
265fe69
fix(bench): attention's p buffer was the operand p
LessComplexity Oct 24, 2026
edf96e9
evidence(fusion-natural-forms): H100 PCIe and RTX 5090 box runs
LessComplexity Oct 24, 2026
8350361
evidence(fusion-natural-forms): attention re-take, H100 PCIe and RTX …
LessComplexity Oct 24, 2026
b8796e8
evidence(fusion-natural-forms): RTX 4090 box run
LessComplexity Oct 24, 2026
f007a6f
evidence(fusion-natural-forms): RTX 3090 box run
LessComplexity Oct 24, 2026
cfb1458
evidence(fusion-natural-forms): RTX 5090 confirm box (Ryzen 9 7945HX …
LessComplexity Oct 24, 2026
7400a00
evidence(fusion-natural-forms): the 5090 confirm box throttled mid-run
LessComplexity Oct 24, 2026
cbe0a66
perf(gpu): fuse ships OFF on all four measured parts
LessComplexity Oct 24, 2026
e2b0e68
evidence(fusion-natural-forms): RTX 3090 confirm box (EPYC 7B12, 250 W)
LessComplexity Oct 24, 2026
678caee
evidence(fusion-natural-forms): the slow 3090 confirm box to confirm-…
LessComplexity Oct 24, 2026
4898e15
evidence(fusion-natural-forms): RTX 4090 confirm box (Ryzen 9 5950X, …
LessComplexity Oct 24, 2026
d071cbc
evidence(fusion-natural-forms): RTX 5090 confirm box (Ryzen 9 7950X, …
LessComplexity Oct 24, 2026
c10c1e4
evidence(fusion-natural-forms): RTX 3090 confirm box on S65's own hos…
LessComplexity Oct 24, 2026
a7fa76c
evidence(fusion-natural-forms): H100 PCIe confirm box on S65's own ho…
LessComplexity Oct 24, 2026
bd546e9
perf(gpu): the fuse rows cite the second hosts; sweep tables
LessComplexity Oct 24, 2026
9a3477a
docs: the fusion-natural-forms sweep — evidence README, next-session,…
LessComplexity Oct 24, 2026
285f61a
perf(ir): facts live on the graph
LessComplexity Oct 24, 2026
349767a
perf(gpu): a kernel no launch site loads is not emitted
LessComplexity Oct 24, 2026
1a881e8
bench: causal attention decode, prefill+decode, multi-channel conv
LessComplexity Oct 24, 2026
9c1e63a
bench(gpu): facts-on-graph sweep on 4 archs, default face
LessComplexity Oct 24, 2026
072f726
fix(bench): default face gated by its bound, never bit for bit
LessComplexity Oct 24, 2026
531eb9c
refactor(gpu): rung fuse is flow
LessComplexity Oct 24, 2026
87d9739
perf(gpu): a flow unit's row is a block of lanes
LessComplexity Oct 24, 2026
c470546
test(gpu): default-face flow rows are lanes with tree folds
LessComplexity Oct 24, 2026
437dcd5
fix(gpu): a retained output's move intent is its steady trip
LessComplexity Oct 24, 2026
81857cc
bench(gpu): torch legs for pj/twin/epi, a stem without a torch leg is…
LessComplexity Oct 24, 2026
5eaad72
bench(gpu): flow sweep on 4 archs (flow ON/OFF, retain-resident h100 …
LessComplexity Oct 25, 2026
04a5cf3
docs(evidence): why-slow — per losing workload x arch gap split, caus…
LessComplexity Oct 25, 2026
7a400a4
perf(gpu): flow ships ON on all four parts
LessComplexity Oct 25, 2026
adc4409
docs: handoff to cpu-gpu-boundary; commit the agent rules (CLAUDE.md,…
LessComplexity Oct 25, 2026
5c8093a
docs: drift — renamed split_count and in_place_source rows
LessComplexity Oct 25, 2026
64de5f0
perf(gpu-rt): kernels, arena and uploads leave the timed path
LessComplexity Oct 25, 2026
c8a49d6
perf(gpu): uniform and nest-cell placement, tree-gate
LessComplexity Oct 25, 2026
317bdf1
merge cpu-gpu-boundary-runtime into placement
LessComplexity Oct 25, 2026
8a4142f
fix(cpu-gpu-boundary): review findings of both builders
LessComplexity Oct 25, 2026
e8c2050
fix(evidence): entry-arena OFF check reads ON-only text (arena base l…
LessComplexity Oct 25, 2026
3190d8c
fix(gpu-bench): atomic split launch line; entry-arena ON applies wher…
LessComplexity Oct 25, 2026
16974f9
fix(gpu-bench): differential ignores the clock line; TF32 atomic spli…
LessComplexity Oct 25, 2026
8d6b215
fix(evidence): a slot the entry arena seeds never allocates (runtime …
LessComplexity Oct 25, 2026
1169548
results(cpu-gpu-boundary): rtx-4090 smoke — values, ABBA arms, evidence
LessComplexity Oct 25, 2026
15f52cf
bench(gpu): integrated_bundle drops a rung arm by emission before com…
LessComplexity Oct 25, 2026
73739a7
bench(gpu): cpu-gpu-boundary sweep on 4 archs (value gates, ON/OFF of…
LessComplexity Oct 25, 2026
9bb981c
docs(evidence): cpu-gpu-boundary end check, rows, refreshed gap table
LessComplexity Oct 25, 2026
8e41f94
docs(readme): GPU workload ratios on 4 archs
LessComplexity Oct 25, 2026
8ef1de8
perf(gpu): rung rows from the cpu-gpu-boundary sweep
LessComplexity Oct 25, 2026
f7925b4
docs: cpu-gpu-boundary session log and handoff; fix swapped ON/OFF in…
LessComplexity Oct 25, 2026
ab82876
docs(adr): ADR-0041 candidate — compute precision as a permission on …
LessComplexity Oct 25, 2026
8e876a9
docs: next-session item 1 — prefill/decode emerge from the loop recur…
LessComplexity Oct 25, 2026
0736051
docs(adr): ADR-0041 accepted for implementation
LessComplexity Oct 25, 2026
71bf770
docs(adr): ADR-0041 surface — a type on the compute step (fold: tf32 …
LessComplexity Oct 25, 2026
7a7f500
docs(adr): ADR-0041 — ': compute' on steps, blocks and functions
LessComplexity Oct 25, 2026
29f0a54
docs(adr): ADR-0041 — function compute type follows the output (fn f(…
LessComplexity Oct 25, 2026
0d20476
test(llvm): bound the differential harness to the core budget; time-b…
LessComplexity Oct 25, 2026
7c25436
feat(ir,gpu): mixed-radix index fact; convmc as an implicit-GEMM tile
LessComplexity Oct 25, 2026
1390658
feat: ADR-0041 compute precision as `: compute` on bodies; tf32-tile …
LessComplexity Oct 25, 2026
70441f0
build: optimize mapal-interp in dev/test builds (the oracle), checks …
LessComplexity Oct 25, 2026
dbfbfd5
docs: next-session — test-speed item with measurements
LessComplexity Oct 25, 2026
85dba44
perf(gpu): prefill and decode emerge from the loop recurrence
LessComplexity Oct 25, 2026
4434029
perf(gpu): nest cells and row folds fold over a warp's lanes
LessComplexity Oct 25, 2026
bd5f7f2
Merge branch 'recurrence-precision-tiles' into recurrence-precision
LessComplexity Oct 25, 2026
87e1295
Merge branch 'cpu-gpu-boundary-latest' into recurrence-precision
LessComplexity Oct 25, 2026
6ab43e3
merge: resolve the mapal-ir re-export conflicts (TripSplit with Preci…
LessComplexity Oct 25, 2026
578e06d
fix(review): radix sites — one output, tensor branch, admission arms;…
LessComplexity Oct 25, 2026
f619fe2
fix(review): trip-batch keeps traps, batches invariants; range-cell g…
LessComplexity Oct 25, 2026
674861f
docs(ir): the lanes test row names lanes_on_unseen_programs (drift)
LessComplexity Oct 25, 2026
436836a
fix(gpu): tensor cores and fused radix steps only where the order is …
LessComplexity Oct 26, 2026
e47f685
fix(gpu): a +0.0 fill carry init is zeroed on the device
LessComplexity Oct 26, 2026
c8215b5
test(review): recurrence generator shapes; lanes gated on the interpr…
LessComplexity Oct 26, 2026
c99ac7f
fix(gates): M through relu and bias chains; trip-batch in the movemen…
LessComplexity Oct 26, 2026
c6cd2d1
results(recurrence-precision): rtx-4090 smoke — values, ABBA arms, ev…
LessComplexity Oct 26, 2026
7d97152
bench(torch): TF32-on legs (matmul + cuDNN allow_tf32) for tf32 / con…
LessComplexity Oct 26, 2026
006cf8c
results(recurrence-precision): sweep step 1 — emission diff vs cpu-gp…
LessComplexity Oct 26, 2026
87bc83c
bench(torch): torch legs for pt_gen (prefill math), the _tf32 program…
LessComplexity Oct 26, 2026
784e69d
results(recurrence-precision): narrowed sweep on 4 archs — changed pr…
LessComplexity Oct 26, 2026
a06c512
docs(recurrence-precision): end check, gap table, what moved, open
LessComplexity Oct 26, 2026
d2a9ab2
docs(readme): GPU results after recurrence-precision
LessComplexity Oct 26, 2026
a2f817d
rungs(rtx-4090): fold-lanes ships OFF
LessComplexity Oct 26, 2026
1950c54
docs: recurrence-precision session log and handoff
LessComplexity Oct 26, 2026
1613a13
docs(adr): ADR-0034 accepted for implementation (sprint site-geometry)
LessComplexity Oct 26, 2026
874944f
geom(B1): per-target geometry table with --geom overrides
LessComplexity Oct 26, 2026
be93a32
bench(changed): emission diff per arch x face vs a base ref, arm tags…
LessComplexity Oct 26, 2026
22b8170
bench(bundle): changed-only bundles and a preflight before any rental
LessComplexity Oct 26, 2026
f30cc05
bench(search): ADR-0034 D2 search driver over benches/gpu/knobs.json
LessComplexity Oct 26, 2026
107017b
fix(preflight): value_gate.py only where a probe answers to D2's bound
LessComplexity Oct 26, 2026
785f2ad
test(device-loops): run the programs over the job slots; touched-file…
LessComplexity Oct 26, 2026
f9c4790
geom(B2): fill-wave and row-split — one fill decision for tile sites
LessComplexity Oct 26, 2026
4eddb74
geom(R1): fit-first — reshape a tall thread grid before shrinking it
LessComplexity Oct 26, 2026
0bffbd2
geom(B4): GEMM tile candidates as table entries
LessComplexity Oct 26, 2026
1f932b9
geom(B3): tf32-radix — the mma.sync tile for a tf32 radix site
LessComplexity Oct 26, 2026
bf6f103
Merge branch 'site-geometry-measure' into site-geometry
LessComplexity Oct 26, 2026
f504761
fix(geom): review findings on site-geometry-sites
LessComplexity Oct 26, 2026
ff750a8
fix(measure): review findings on site-geometry-measure
LessComplexity Oct 26, 2026
e275f17
fix(changed): runtime stamp over what a row's time runs through (runt…
LessComplexity Oct 26, 2026
434c571
fix(gen_tiles): K through a map's output (a fold over a computed array)
LessComplexity Oct 26, 2026
8e9b432
fix(evidence): atomic adds with memory-order qualifiers (newer llc); …
LessComplexity Oct 26, 2026
6a37cd9
fix(gpu): an epilogue group's ordered launch keeps its own tile besid…
LessComplexity Oct 26, 2026
e07a812
fix(bundle): cells keep their reassociation faces under SUITE_ARMS=0 …
LessComplexity Oct 26, 2026
a9c16a9
bench(site-geometry): rtx-4090 smoke — value gates, ON/OFF of the new…
LessComplexity Oct 26, 2026
3526d41
wip(site-geometry): sweep evidence collected before a checkpoint (lap…
LessComplexity Oct 26, 2026
f794649
docs(rules): a mixed ON/OFF result is a missing fact, not an OFF row
LessComplexity Oct 26, 2026
9d550f5
bench(site-geometry): sweep h100 value-gated re-run, search picks, re…
LessComplexity Oct 26, 2026
037b855
bench(site-geometry): h100 search run, picks with shipped values, report
LessComplexity Oct 26, 2026
9583482
perf(gpu-rungs): site-geometry judge — tf32-radix ON on 5090/3090, ro…
LessComplexity Oct 26, 2026
0035aad
fix(box_run): the cuBLAS timing section's _exact face reads the _x mo…
LessComplexity Oct 26, 2026
1f1869e
perf(gpu-rungs): keep tf32-radix OFF on the 5090 — the fitted FFMA ti…
LessComplexity Oct 26, 2026
750c97f
docs(evidence): site-geometry end check, gap table, missing facts
LessComplexity Oct 26, 2026
a25a029
docs(readme): GPU workloads table from the site-geometry sweep
LessComplexity Oct 26, 2026
ba4f295
merge: recurrence-precision-latest (CLAUDE.md mixed-result rule) into…
LessComplexity Oct 26, 2026
991e59a
results(cache): measurement cache from the site-geometry sweep (keyed…
LessComplexity Oct 26, 2026
e3d4710
docs: site-geometry session log and handoff
LessComplexity Oct 26, 2026
1dbe09e
docs: next-session — fix the handoff text (site-geometry standing and…
LessComplexity Oct 26, 2026
3fc0750
docs: summary.md and handoff for the pull request to main
sapiritur Oct 29, 2026
d4e34c6
feat(fmt): mapal fmt source formatter (syntax, CLI, editors) and docs
sapiritur Oct 29, 2026
3966475
evidence: raw run output moves to a GitHub release; manifest, index, …
sapiritur Oct 29, 2026
7df88ff
style: cargo fmt --all (CI formatting step)
sapiritur Nov 1, 2026
a37c984
test(retain_delta): contract oracle fuses on x86-64 (-mfma)
sapiritur Nov 1, 2026
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
The table of contents is too big for display.
Diff view
Diff view
  •  
  •  
  •  
The diff you're trying to view is too large. We only load the first 3000 changed files.
6 changes: 6 additions & 0 deletions .editorconfig
Original file line number Diff line number Diff line change
@@ -0,0 +1,6 @@
[*.mapal]
indent_style = space
indent_size = 4
end_of_line = lf
insert_final_newline = true
trim_trailing_whitespace = true
9 changes: 9 additions & 0 deletions .gitignore
Original file line number Diff line number Diff line change
@@ -1,4 +1,6 @@
/target
/graphify-out/
/.benchmark-artifacts/
__pycache__/
.DS_Store

Expand All @@ -9,3 +11,10 @@ benches/matmul/*.cu

# built artifact — regenerate with package-vsix.py
editors/vscode/*.vsix

# Raw run output goes to a GitHub release (benches/gpu/archive_raw.sh), never into git.
benches/results/evidence/**/*.tgz
benches/results/evidence/**/*.out
benches/results/evidence/**/*.log
benches/results/evidence/**/results-*
benches/results/evidence/**/run-*
1 change: 1 addition & 0 deletions AGENTS.md
108 changes: 108 additions & 0 deletions CLAUDE.md
Original file line number Diff line number Diff line change
@@ -0,0 +1,108 @@
# Mapal — rules for every agent and workflow

## Source formatting

- Format every Mapal source file you create or edit with `mapal fmt <files...>`;
verify it with `mapal fmt --check <files...>` before handing off. If the CLI is
not installed, use `cargo run -p mapal-cli -- fmt <files...>` (add `--check` to
verify). Pass only files touched by the task; do not reformat unrelated sources.
- The formatter owns layout: four spaces, multiline brace bodies, consistent
spacing. Keep comments and descriptive names readable; avoid hand-aligned padding
or compacting generated programs onto long lines. Format generated source too.
- Deliberately invalid syntax fixtures are exempt because formatting rejects invalid
input. Explain that exemption in the test; do not bypass checks for valid programs.
- For code discovery, avoid recursive indexing of archived worktrees, vendored
libraries or generated bundles. Use bounded reads of known files when indexing
stalls; do not launch concurrent index builds for the same checkout.

## Evidence pass (mandatory for any optimization or codegen change)

No keep/remove/gate decision stands on timing alone.

1. **Intent, before measuring.** Each rung declares, machine-checkably:
- what LLVM/NVPTX must derive from the emitted IR (pass/transform: vectorize, unroll, SROA of accumulators, fma
contraction, LSR, cp.async lowering, …), and
- the PTX/SASS features that must appear or disappear (instruction mix, vector widths, LDGSTS/cp.async, FFMA per
load, unroll depth, accumulators/thread, registers, spills, smem, occupancy, launches), and
- the data movement the facts require per program and per loop trip (host↔device copies, allocations, residency);
any copy or allocation the dataflow facts do not require is a finding. Placement, reuse and retention (caches)
are deduced from the graph for every object, loop-carried state included — never a by-value default.
2. **Realized.** For every measured kernel and arch collect LLVM optimization remarks (`-pass-remarks-output`), PTX,
SASS (`cuobjdump -sass` / `nvdisasm`) and `ptxas -v`; check against the intent. A mismatch is a finding, never a pass.
Library kernels (cuBLAS/cuBLASLt/CUTLASS) are captured from the actual launch, not assumed.
3. **Explain.** Every timing difference that drives a decision cites the SASS/usage difference that explains it.
Unexplained ⇒ the decision is "open", not keep/remove.
4. **Breadth.** Judge on the multi-workload suite, not matmul alone: GEMM/MLP, attention/softmax decode, reductions,
FIR/conv/window, elementwise chains, fan-in/fan-out, loops — CPU-only, GPU-only and CPU+GPU — against the best
library (cuBLAS/cuBLASLt/CUTLASS) or PyTorch where no library exists.
Tooling: `benches/gpu/intent.py` (intents), `evidence.py` (checker), `rejudge.py`; a GPU run without its
`evidence.tsv` is not a result.
Raw run output (profile archives, box logs, raw result text) never goes into git: `benches/gpu/archive_raw.sh`
uploads it to a GitHub release and appends `benches/results/evidence/RAW-manifest.tsv`; commit only the evidence
tables, reports and the manifest (`benches/results/evidence/RAW.md`).
5. **End check.** The final integrated build gets the same pass: a table of intended vs realized per optimization per
arch, and how it turned out in time.

## Facts catalog — every decision is based on these (`crates/mapal-ir/src/algo.rs`)

| query | what it proves |
| --- | --- |
| `emission_plan` / `class` | minimal emission: which objects materialize; products dissolved (views, no bytes); consumer counts |
| `last_use_plan` → `position`, `death`, `escapes`, `carried_by`, `dead_after` | death points; escape into output; `carried_by` = value crosses a `LoopBack` (two-iteration liveness) |
| `loop_plan`, `loop_structure`, `sccs`, `topo_order` | canonical loops: init, merge, decide/advance order, exits |
| `write_delta` | the contiguous range of a carried array each trip writes, through `Update`s and piecewise maps (partial update → move/retain only the delta) |
| `path_plan` (+ `is_single_path`) | independent parallel paths / tasks |
| `tile_plan`, `tile_plan_strided`, `move_plan`, `row_fold`, `cell_plan` | affine map/fold sites, data-move sites, row folds, per-cell maps and what each cell reads |
| `elem_plan`, `map_law` (+ `Outside`) | each array's `out[i]` as a closed function of the index; one map's law, piecewise when every cell off a row / range / column / cell is a capture's own (in-place range map) |
| `radix_plan`, `tile_plan_radix` | an integer as digits `(x ÷ s) % r` of the inputs (constant div/mod splits); map/fold sites whose reads are affine per digit as tile sites over `(i, lane, k)` (implicit GEMM) |
| `precision_plan` (`set`, `round_at`, `cross`, `bound`) | ADR-0041: the precisions each float product admits (innermost `: compute`); values rounded once at the producer; edges a rounded copy must not cross; each tile site's a-priori bound (D3 / ADR-0040 D2) |
| `bounds_proof`, `proven` | in-bounds indices; droppable trap guards |
| `guard_plan`, `gated` | per-Phi arm exclusive work; whether an arm must be gated |
| `trip_split` | per counted loop: the trip intervals guards on the counter cut, each interval's recurrent state slots, the work outside them, batched row writes, whether the first interval's trips but the last can be skipped |
| `whole_fold`, `fold_tree_step` (`CellSite::tree`) | a fold of the function itself over one affine read (a one-row row fold); a fold step regroupable as lane partials + tree where the order is free |
| backend: `ResidentPlan`, `arena`, frame slots (`prepare_storage`), `TargetProfile` | device residency, lifetime-colored buffers, host slots/aliases, hardware facts |

New logic starts from this table; if a needed fact is missing, add it to `mapal-ir` (one query, tested) — never
re-derive it inside a backend pass.

**Facts live on the graph.** The IR is a dataflow execution graph; every fact above is stored on the sealed graph
once computed, and its dependencies on other facts form a dataflow graph of their own. A query reads the stored fact;
only a fact that is not there yet is computed (from its stored inputs). Never recompute a fact per call site.

## Standing rules

- Optimizations derive from Mapal dataflow facts (tile_plan, elem_plan, loop_plan, last_use_plan, path_plan,
bounds_proof, use counts, death points) + TargetProfile hardware facts, emitted so LLVM can derive the schedule.
No workload-named special cases, and no parallel module that re-derives what an IR fact or existing allocator
already models (e.g. a loop-only lifetime allocator next to `LastUsePlan` coloring) — extend the existing
machinery so the case falls out. Reuse the plans/facts at hand; add a new fact only when none carries it.
- **A mixed result is a missing fact** (Sapir, 2026-10-26). An optimization that wins on some programs and loses on
others beyond null means its condition lumps different situations together (over-abstraction). Find what separates
winners from losers in the SASS/usage evidence (layout, extent vs SMs, reuse, occupancy, …), add it as an IR or
`TargetProfile` fact, and make the rule read it, so each case gets its own decision. A `GPU_RUNGS_OFF` row is only a
stopgap and is always recorded as an open "missing fact" item.
- **Generic, proven on unseen cases.** Never target a benchmark. Every rule is validated by a property test over
programs from the random generator (`crates/mapal-rewrite/tests/testgen/`, extended as needed) that the change was
not designed on: values vs the interpreter / bound, plus movement and kernel intents. Benchmarks are acceptance
only; a change that helps only the suite is rejected.
- **No "fusion" concept.** Compilation follows data usage in the dataflow graph: which work runs together, what stays
in registers/shared memory and what is stored all follow from producer/consumer footprints (`footprint_plan`), death
points, independence (`cell_plan`, `path_plan`) and `TargetProfile`. Kernel boundaries are a result, never a pass.
The rung is `flow` (formerly `fuse`). Work units must keep the parallelism the graph proves (independent cells across
threads, folds as tree reductions) and compute a value once when it has several readers.
- **Benchmarks are written in idiomatic Mapal** (map / fold / loop / guards / captures), never unrolled or shaped to
what the compiler already recognizes. If the natural form compiles worse, the facts are generalized — the program
is never rewritten to suit them. `benches/algorithms/*.mapal` is the committed, readable snapshot. Every rung has an OFF arm; per-arch defaults only via cited `GPU_RUNGS_OFF` rows.
- Default float face: full precision, a-priori error bound, order unspecified, fastest for the native CPU+GPU
(ADR-0040). Exact/contract/fixed-order are opt-in; value gates on the default face check the bound.
- GPU archs: RTX 4090 sm_89, RTX 5090 sm_120, H100 sm_90a, RTX 3090 sm_86, rented on vast.ai on demand
(`benches/gpu/vast_part.sh`/`vast_run.sh`). `vastai destroy instance <id>` needs `-y`; never touch instances you
did not create — any instance labeled `stocklab-*` belongs to another project; never print keys or provider logs. Same-box ABBA A/B, value gates first, report null variation.
- This is a laptop (14 cores): leave 2 cores free for heavy local work (`CARGO_BUILD_JOBS=12 RUST_TEST_THREADS=12`).
Background QoS (`taskpolicy -b`, efficiency cores only, ~3× slower) only when the user asks for a quiet machine.
- Token budget is tight: wait on a GPU box or long test with ONE blocking background command and read its summary
once (grep/tail, never full logs); no tick-by-tick polling; keep returns to short tables, details in the repo.
- Parallel agents share one scratchpad: each keeps its files in its own prefixed subdir.
- Never push. Results only in README (tables + link).
- Name action items, branches, result dirs and session logs by topic (`deduced-residency`, `generic-fusion`), never
by session number (`s65`, `S66`).
24 changes: 24 additions & 0 deletions CONTRIBUTING.md
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,7 @@ measured, not argued with.
- [Proving a correctness change](#proving-a-correctness-change)
- [Proving a performance change](#proving-a-performance-change)
- [Tests](#tests)
- [Source formatting](#source-formatting)
- [Docs you update in the same change](#docs-you-update-in-the-same-change)
- [Setup and the gates](#setup-and-the-gates)
- [Commits, PRs, licensing](#commits-prs-licensing)
Expand Down Expand Up @@ -189,6 +190,29 @@ Editor grammars have their own suite (`sh editors/test.sh`, 61 assertions) and o
precedence rules between Vim and TextMate — a rule added at one end of one goes at the other
end of the other.

## Source formatting

Use the same formatter for hand-written and generated Mapal source:

```sh
cargo build -p mapal-cli
target/debug/mapal fmt path/to/program.mapal
target/debug/mapal fmt --check path/to/program.mapal
target/debug/mapal fmt --stdout path/to/program.mapal
```

Pass only the files you changed. `fmt` writes explicit paths in place; `--check`
never writes and exits 1 if formatting is needed; `--stdout` previews one file.
`mapal fmt -` reads stdin and writes formatted source to stdout for editor tools.
Invalid input produces diagnostics without replacing the source. Deliberately
invalid parser fixtures are exempt from formatting checks.

The style is four spaces, expanded brace bodies, separate statements and record
fields, and consistent spacing. Comments and literal spelling are preserved.
Use descriptive names and short, useful comments: layout alone does not make code
readable. The [Neovim](editors/nvim/README.md) and
[VS Code](editors/vscode/README.md) plugins use this same CLI.

## Docs you update in the same change

Not after. In the same PR (FRAMEWORK §6.3):
Expand Down
9 changes: 9 additions & 0 deletions Cargo.lock

Some generated files are not rendered by default. Learn more about how customized files appear on GitHub.

6 changes: 6 additions & 0 deletions Cargo.toml
Original file line number Diff line number Diff line change
Expand Up @@ -17,3 +17,9 @@ members = [
[workspace.package]
version = "0.0.1"
edition = "2024"

# The interpreter is the tests' oracle (budgets up to 400M steps): optimized in test builds,
# checks kept (debug-assertions / overflow-checks come from the dev profile). Measured:
# differential suite 777 -> 683 s; rebuild cost 2 s.
[profile.dev.package.mapal-interp]
opt-level = 3
Loading
Loading