Skip to content

Latest commit

 

History

History
186 lines (162 loc) · 175 KB

File metadata and controls

186 lines (162 loc) · 175 KB

Mapal — Global Status

Last updated: 2026-08-10 · Sessions 49–50 (the scheduler was the kernel). Sapir asked for a linear-NN benchmark — NumPy vs Mapal, both machines, at scale, plus two nets at once. It could not be run: every chained bulk program in the tree was emitting fully sequential code, attn_256.mapal included, since S25. path_plan grouped scalar morphisms into undirected components of a "shares an object in any role" graph; every matmul site's capture tuple captures the same k range, so one shared read fused them all into one task, and layer 2's group then reads layer 1's output — so that task sat both before and after the layer-1 map. A cycle. llvm/src/lib.rs:261 filters a cyclic plan out and emits the sequential form, so the symptom is silently lost parallelism, and path_plan's own debug_assert is compiled out of release. The numbers those programs printed were also fiction — pre-fix nn2_2048 self-timed at 1.81 ms = 19,000 GF/s on a 4,100 GF/s part, because the clock fence is task dependency edges and the sequential path has none. Fixed by making component adjacency dataflow plus co-writing, never sibling-reading: two morphisms that merely READ one object are siblings, not a chain. Then the benchmark ran, and it said the concurrency model cost more than the compute at small sizes — at d=128 a 4-layer net spent ~95% of its parallel wall in the runtime. Six more changes, all measured: P0 a packed tile site is two tasks in the outer run, not a wrapper opening two nested runs (36.8 µs each, per site, per invocation, inside the timed region); P0b a dispatch wakes only the lanes it filled (empty dispatch 5.1 → 0.46 µs at 14 lanes); Placement collapsed to one variant — dep-unlocked tasks are seeded, not placed on the producer's lane, which S43 predicted and which is worth 1.49–1.58× on its own; P2 the run graph is sealed once and the completion path is lock-free (execute takes no lock, every slice but the last is one atomic) — this closed a 0.855× regression on the 32-lane i9 and the two cells it hurt most were the two it helped most, 1.733× and 1.433×, which is a diagnosis confirmed by its own repair; P1 the host spins before parking while workers do not — spinning workers are REFUTED, 0.80–0.93× on real shapes at every budget swept, because handoff_floor gave every thread work every wave and a real dispatch leaves most lanes idle; P3 a task that cannot pay for its own dispatch is not dispatched — one-slice tasks run on the calling thread (constant-free), and a roofline estimate max(2nk/flops_per_ns, bytes/bytes_per_ns) against a measured dispatch_ns decides the rest. A pure flop count cannot make that call: conv2d_1024 has HALF the flops of an nn4_256 layer and takes 3.7× longer. Results, M4 Pro, one run, 41 cycles, values bit-identical to NumPy: 4-layer net at d=128 0.3636 → 0.0810 ms (4.49×), d=256 1.82×, d=512 1.38×; against NumPy the tie moved to d=512 (1.02×) and d≥1024 is 1.23–1.47× ahead. Two independent nets are 1.34–1.36× faster than the same flops chained with nothing in the source saying so; NumPy's equivalent ratio is 1.00–1.09 with a matching control, i.e. no overlap. The whole non-matmul ladder moved: 13 of 13 shapes improved (saxpy 1.40×, conv2d_s1024 1.30×, transpose_512 1.29×, conv2d_512 1.26×, conv2d_1024 1.25×, gather 1.13×, reduce 1.10×) — streaming, stencil, permutation, gather, reduction, no matmul among them, which is the evidence the parallelism is in the graph and the runtime rather than in a kernel. matmul: 1024 0.81 → 0.72, 2048 5.17 → 5.03, 4096 unchanged at 38.7 (NumPy 44.1). Three bugs found and fixed, one of them introduced by this session's own P2: retire published Done before unlocking its dependents, so a waiter could beat the unlock and trip mapal_par_run_pinned's assertion — survivable only while the host always parked, and P1's host spin made it fire on conv2d_1024 within seconds (hammered ×300 after the fix). ADR-0032 held throughout: mapal-ir learned no machine fact — its whole diff is the graph rule plus tests, and every scheduling and machine decision lives in backends/llvm and mapal-rt. A profile with no MEASURED dispatch economics carries None and dispatches everything, so no part is made worse by a number nobody took — verified, 0 of 174 cells move under --target=generic; raptorlake is owed its three numbers because the arch box went down mid-session. MAPAL_PAR now means total executors including the host (Pool::new spawns threads - 1, because help_until already claimed lane threads - 1) — measured neutral, but it changes what a number taken at a given MAPAL_PAR means and needs Sapir's sign-off. Gate: 1056 passed / 0 failed, differential byte-equality oracle green throughout — results never changed, only speed; fmt clean. The remaining gap is compute, not scheduling: at d=128 the parallel time is now within 20% of the one-thread time, and mapal's single thread is 0.28–0.35× of NumPy below d=512 (0.87× on the i9). S43's operand-residency instrument already measured +71% at one thread, assembly-verified — 939 × 1.71 = 1606 GF/s against Accelerate's 1546, i.e. the gap and the unclaimed win are the same size — and it is unclaimed because every lever that captures it (kc +6.1% 1t / −25.5% threaded, nc +18.7% 1t / parity) costs the configuration that ships. Three measurement rules earned (25 a microbenchmark that gives every thread work every wave does not model a dispatch; 26 a constant derived from the runtime must be re-derived when the runtime changes; 27 label every table with the run it came from). Everything is UNCOMMITTED on main @ 680ce44, alongside the untouched S48 work. Logs: sessions/2026-08-10-s49-s50-the-scheduler-was-the-kernel.md. Perf: performance/s49-two-nets-and-the-cyclic-plan.md, s49b-the-dispatch-price.md, s50-p0-the-nested-runs.md, s50b-p2-the-shared-lock.md, s50c-p1-spinning-workers-refuted.md, s50d-dont-dispatch-what-cannot-be-shared.md, s50e-the-roofline-threshold.md. Plans: components/runtime/plans/plan-s50-scheduler-handoff.md, components/backend-llvm/plans/plan-s49-linear-nn-two-nets.md. New component: components/runtime/.

Previous: Sessions 44–47 (conflict, not capacity). See sessions/2026-07-31-s44-s47-conflict-not-capacity.md; the S44–S47 roll-up that stood here is preserved there and in performance/s44-conflict-not-capacity.md.

Previous: Session 42 (the constant that cost a session — recorded here late; S42 closed without updating this roll-up). S42's stated P0, k-loop software pipelining, was refuted in the first hour (+0.1–0.2%, overlapping, 3 sizes, both kernel layouts). The genericity work shipped clean: f32_tiles DELETED (ZA holds exactly sizeof(elem) tiles at a given width — an ISA rule, not a recorded fact, so recording it could only make it wrong), --target=native detects SVL/L1D/L2 from sysctl and emits byte-identical IR to the hand-written apple-m4-sme, SME b-addressing derived from recorded facts rather than three coincidental literals, and the A pack reordered row-outer (scalar float loads 51 → 5, worth 3%). KC blocking built and default OFF: at the corrected depth +6.1% at 1 thread / −25.5% threaded at N=4096, and threaded is what ships. The session's process failure is its lesson — one wrong constant (sme_kc = 512 instead of the swept 1024) made KC look like a 1.27× loss and sent SIX investigations down the wrong road, all refuted at the wrong depth and all wasted; the sweep that overturned it was the simplest experiment available and was run last (rule 17). Also: the matrix unit count became a measurement rather than an inference — exactly 2 units, ~2000 GF/s each, ~4100 aggregate (units.c, zero memory traffic) — and the i9 box leg closed S29's open item, finding a step function rather than a curve (any blocking costs ~1.8× at 1t, ~1.3× threaded, at every depth). Its §5e conclusion — "the gap is operand cache residency, worth ~1.79×, which would pass Accelerate" — was partly retracted by S43: the direction survives and was confirmed in the emitter at one thread, but the 1864 ceiling was thermal drift, 1043 was the wrong cell, and the magnitude, the level and "would pass Accelerate" all fall. Log: sessions/2026-07-31-s42-the-constant-that-cost-a-session.md. Perf: performance/s42-sme-roofline.md (carries its own retraction boxes at §0 and §5e). Previous: Session 41 (the NVPTX leg — plan RATIFIED, step 1 of 8 built; and the geometry record finally has gates). Three things landed, none of which moved a single emitted byte. (1) func.rs is now func/ (Sapir's directive, a plan-s41 prerequisite): 7,299 lines — of which impl<'a> FnEmit<'a> alone was 6,567 — became eleven child submodules each with its own impl block (core 820 · frame 232 · drive 672 · ops 743 · tile 492 · window 382 · conv 828 · packed 971 · trio 761 · vec 377 · bulk 428 · mod 745). Visibility preserved exactly, not widened — private methods became pub(super), the surface one file already had; pub(crate) unchanged at 13. Proven a pure refactor rather than asserted: 159/159 A/B emissions byte-identical (53 sources × 3 faces), re-verified after cargo fmt and again after comment fixes; 113 methods in, 113 out. Note the instrument choice — the test suite alone would NOT have caught a silent emission reorder, since goldens exist only where goldens exist; 159 emissions is the wider net. Side effect that matters for the GPU leg: all four mapal_par_* sites now live in ONE file (func/drive.rs). (2) The §2.2 gates — the durable output of this session. Gate A (crates/mapal-ir/tests/consumer_coverage.rs) is ADR-0033's hand-run grep, automated: every backend is Required or Exempt with a reason and an end condition; an unlisted backend fails; a stale exemption fails; a self-check prevents the gate passing vacuously. Gate B (tile_sites_pin.rs, +2) snapshots the record's field values — not just the site counts — and proves the plan is a pure function of the graph. Together they are the answer to the nine-session drift, and they are independent of packaging, which is the correction Sapir made to the plan's first draft. (3) Step 1: Machine::{Cpu, Gpu(Gpu)} on TargetProfile + cuda-ada (sm_89), every pre-S41 profile pinned Cpu, GPU facts reachable only via gpu() -> Option. Gate: 1015 passed, 0 failed (1006 → 1009 → 1015), fmt clean, emission 159/159 identical, tree UNCOMMITTED on main @ aaaa5dd. Two corrections Sapir made, both kept: the plan's first draft claimed forking caused CUDA to stop consuming tile_planretracted, the record says tile_plan landed S25 and the last CUDA session was S23, so work stopping is the cause and packaging is independent; and the kind-1/kind-2 taxonomy was too coarse — every matrix unit (SIMD FMA, SME, Intel AMX, tensor cores) is the same sentence (stage operands → issue a block MAC → accumulate into a resident accumulator → keep it hot across the reduction axis → store once), so the tile nest is shared and only the innermost leaf differs; the one genuinely structural split is cooperation, which is GPU-only. And the session found a second, cheaper leg: ARM SME. Probed rather than assumed (the S38 method): the M4 Pro reports FEAT_SME2=1, LLVM 22.1.8 lowers llvm.aarch64.sme.mopa.nxv4f32fmopa za0.s, …, rdsvl measures SVL = 64 B ⇒ 16×16 f32 ZA tiles × 4, a 16×16 matmul runs 0/256 mismatched, and a hand-written 1024² SME GEMM does 5.0320 ms vs Mapal's 17.5449 (3.49× faster) against numpy's 1.2977 — narrowing the numpy gap from 13.5× to 3.88× while using only 1 of 4 ZA tiles, no B packing and no KC blocking, i.e. the three rungs the backend already implements. Two machine facts recorded that would have cost days: -march=armv9-a+sme2 compiles but SIGILLs (this part has SME without SVE, so armv9-a's implied +sve emits non-streaming SVE — use armv8-a+sme2), and fmopa fuses (92/256 differ vs separate mul+add, 0/256 vs fmaf), so SME is a contract-face realization governed by ADR-0032 D1/D3, not a free speedup — which is exactly where the published flow-fma numbers already live. Artifacts: benches/sme/. Also verified: the Arch box has an RTX 4070 Ti (sm_89) and CUDA 13.3.1 already installed — no vast.ai rental, and GPU/CPU numbers become same-machine comparable. Plan: components/backend-nvptx/plans/plan-s41-the-nvptx-leg.md. Previous: Session 40 (the arm owns the loop — plan-s40-the-arm-owns-the-loop, SHIPPED; the gate is GREEN again). S39's §4a P0 is closed: gating is now stable across LiftLoops, and eval ∘ rewrite = eval holds on all three pinned seeds plus a 1024-case hammer. The root of §4a was the v1 refusal keying a site's semantics on graph shape (SCC presence) — rewriting changes shape while preserving meaning, so no shape-based refusal could ever commute with LiftLoops. The replacement: loops join an arm as units — the unit is exactly the region the flat walk hands the loop driver, machinery is never per-morphism ownable (which is also the precise explanation of S39's failed "delete the refusal" attempt: per-morphism closure cannot complete a cycle, so it gates exit-boundary fragments and starves the driver), a canonical unit joins whole and is represented in the own-list by its LoopEnter handle alone, carrying the unit's transitive trap flag; a loop is heavy by definition. A site inside a loop never joins its own enclosing unit (found by test on first run — from inside, every "external" consumer is vacuously satisfied). Two topologies, one mechanism: loop-inside-arm fires the driver from the Phi (interp run_loop, LLVM emit_loop); arm-inside-loop-body gates through ordinary closure, and the LLVM loop emitter's cones gained the gated-skip interp already had. path_plan folds a gated loop into its Phi's sequential task; ConstFold refuses to fold a Phi whose arm owns a loop (refolds after LiftLoops — lazy, correct, marked). CUDA skipped on Sapir's direction ("cuda can be skipped it will be translated to nvptx later anyways"): the host emitter keeps strict semantics for loop-touching sites via a site filter; its goldens are unmoved and surface Mapal cannot build the shapes (L1406). A 17-agent adversarial review of the first build confirmed SEVEN defects (13 findings, 6 refuted), one proven by execution — an interpreter panic on a valid graph — and all seven are fixed and pinned (mapal-rewrite/tests/guard_loops.rs; full account plan-s40 §6b): the re-close could not unwind a joined unit (handle now re-tested with the join predicate); sink members joined vacuously (a loop writing Return gated); in-body loop-invariant arm work was double-owned; DCE's dead-Phi pin keyed on can_trap and dropped heavy-gated sites (now gated()); ConstFold dropped a resolved guard's arm even when its alias refused an SCC winner (replay panic — the drop now mirrors the alias conditions); and the S39-class instability beyond LiftLoops was REAL — DCE deleting a dead sibling reader flipped a site strict→gated and suppressed a trap, reproduced independently by the 1024-case hammer. That last fix went through two rejected designs (plan-pinning dies in replay, which materializes only read objects; dead-sink ownership in guard_arm is unsound on cross-arm dead cones — the invariant suite's disjointness assert caught it) before landing as DCE pinning tainted dead sinks in the verdict cone: verdicts preserved, pure dead cones still droppable, surface emissions unmoved. Evidence: full workspace gate GREEN — 1006 passed, 0 failed (post-review-round) — fmt clean; LLVM differential 37/37 in 419.6 s with the 1,280-run sweep; CUDA suite 163 green; A/B emission vs 8b40442 — 103 byte-identical, 1 differs (examples/calc.mapal raw, S39's own signed-off change, pinned by its goldens), 0 new emit failures — so S40 moved ZERO surface emissions and no perf run is owed (measurement rule 9/10). Coverage gap recorded: testgen builds only loop-inside-arm; arm-inside-loop-body rests on two hand-built tests (ir + interp). And byte-identity does not cover the COMPILER's own time (Sapir's catch, post-close): a 51-run alternating compile-time A/B against 8b40442 measured the emission sweep at +16.4% (663.2 → 772.1 ms median, distributions non-overlapping — S39+S40 combined, no pure-S39 build exists to split), because guard_plan built units/bounds/trap-fixpoints even for Phi-FREE functions per pass per fixpoint round, and DCE's verdict-cone walks ran on Phi-free graphs; two early-exits (exact no-ops on results, byte-identity re-verified: 103/104, calc only) cut it to +1.7% (651.0 → 662.3 ms), the residual being guard machinery on the functions that actually carry Phis. Report performance/s40-compile-time.md, raw series benches/results-s40/; measurement rule 12 added (A/B the compiler's wall time whenever a deduced query grows or gains a consumer). Work is uncommitted on main together with S39's. Plan: components/ir/plans/plan-s40-the-arm-owns-the-loop.md (§6a records what building added to the plan). Previous: Session 39 (guards gate the flow — plan-s39-guards-are-conditional, SHIPPED). (1 > 0) -> { -true-> 42; -false-> 7 / 0; } printed nothing and exited 101; it now prints 42. Sapir's framing is the whole plan: "not all branches should compute if a branch is not gonna be taken — they should be rendered/compiled yes, but not computed if not taken", and "it's not about select, it is about dataflow, if the flow condition allows it". The root error was a realization promoted to the semantics: category-ir §4.4 justified computing both arms by "both datapaths exist" in hardware and the "branchless-by-default bias for GPU codegen" — true statements about how to emit a guard on one machine, written down as what a guard means (FRAMEWORK §4.2: a TrnLoc promoted to a Trn). Pure is not total — the arm restrictions buy purity, but Div/Mod/Index/Update are pure partial morphisms, so evaluating both arms implements the copair only where BOTH are defined, which is the wrong morphism rather than a different schedule of the right one — and compiled is not computed: both arms stay in the binary, only running is conditional. Mapal already gated one placement and not the other (I4's loop fork fires LoopBack/LoopExit mutually exclusively off one shared Bool, validate()-enforced), so this removes an exception rather than adding a feature. Shipped as mapal_ir::guard_plan, a deduced query beside path_plan/tile_plan/elem_plan, consumed by interp, LLVM and the CUDA host emitter. The trap was only the loud instance: a map { x -> x/0 } in an untaken arm ran too, and with the rewriter on the compiler proved the answer was 99 and still dispatched a 3-task parallel job whose only effect was to trap (DCE pins trapping cones, R4) — 130 lines of IR to 47 after the fix. Five defects the plan did not predict, all found by running rather than reading. (1) Ownership is consumer closure, not liveness — the 1,280-run differential caught it (read before write, testgen case #94): nothing deletes dead code before execution, so a Proj feeding both an arm and a dead Neg looked exclusive, got gated, and the dead Neg read what the unchosen arm never wrote. (2) Subtracting a nested site's work can orphan a morphism whose consumers left with it — re-closed to a fixpoint. (3) can_trap/heavy must be transitive through nested sites in guard_plan and in ConstFold's losing-arm drop; calc's right-folded 5-arm match hides Div/Mod innermost, so without it calc(0, 20, 0) still trapped and the rewritten build diverged from the oracle (R1). The cost cut is measured: an arm's own-list is never empty (it ends with its boundary Pair edge), so the first cut branched even for -true-> x, and in sepia that branch landed inside a per-element map body where it would cost the loop its vectorization — gated() is now legality-then-cost (can-trap MUST gate; heavy — bulk op or call — is worth gating; two scalar arms are left alone). Perf: no change, proven structurally rather than measured — 103 of 104 A/B emissions byte-identical against 8b40442 (= 24f52c9 + a rename-only commit) across every bench shape, every matmul and every example in all three faces, with linked binaries byte-identical and zero new emit failures; the single change is examples/calc.mapal. Report: docs/performance/s39-guards-gate-the-flow.md, raw series in benches/results-s39/. Runtime medians were taken anyway (51 alternating runs × 6 shapes) and spread −5.9%…+1.2% between byte-identical binaries — the strongest form of S38's measurement rule 6, and a reminder that anything under ~6% on an unpinned Mac at sub-millisecond sizes is noise. The class had zero differential coverage (0 trapping arms in 82 guard sites over 320 testgen programs — which is why the bug reached production); Step::PhiTrapArm now generates them, 60 trapping arms in 139 sites. Gate: 992 passed, 2 FAILED — both are one defect, S39's own: gating is not stable across LiftLoops. guard_plan refuses to gate a site whose arm work touches a loop SCC; LiftLoops removes the SCC, so the raw graph runs strict (traps) and the rewritten one gates (does not) — eval ∘ rewrite = eval broken. The gated side is the correct one; the raw graph should gate and refuses to. The refusal cannot simply be deleted (tried: loops break outright, route object built before read), so the fix is to fire arm-owned work through the loop driver rather than refusing the site. IR-only: surface Mapal cannot build it (L1406 rejects -> loop in a Phi arm and lower never puts loop machinery in an arm), so every example, bench shape and matmul is unaffected — the 1,280-run differential RAN and passed, all goldens pass, 0 pending snapshots, fmt clean. Three proptest seeds pinned. Full account: session log §4a. Owed: the CUDA device (kernel) form is still strict — only the host form is gated — and no CUDA change has hardware verification. Previous: Session 38 (CLOSED — log sessions/2026-07-27-s38-trap-order-is-source-order.md) — trap order is now source order, the gate is GREEN for the first time since S33, and the bug turned out to be bigger than the plan said. topo_order's ready worklist became a min-heap on (loc.start, loc.end, insertion index) instead of a FIFO, because insertion order is a property of the compiler and trap order is observable, so it must be a property of the program — the graph does not order two independent trapping ops at all, rewriting creates and destroys objects, and Inline could turn Trapped(IndexOob) into Trapped(DivZero). Approach A rather than a separate selection key, because three sites must agree on which trap is reported (the oracle walks topo_order, sequential LLVM emits in it, the parallel runtime CAS-mins on the same index) and they agree today only because all three derive from it. A′ was priced and REFUTED: 62.2% of raw lowered objects (484/778) change position, only 9 of 36 functions are already in source order, and the deviation is systematic — an object's loc is the operator token, not the sub-expression extent — so "make lowering create objects in source order" churns every golden rather than A's 38. The wider bug is the better demonstration: the same root cause let a trap swallow output written before it, because mapal_par_run_pinned runs its body synchronously on the host thread; same program, both exit 101, PRE prints nothing and POST prints 111\n222\n. Invisible to the 1,280-run sweep by constructionexpect_native maps Trapped to (None, 101) and the stdout assert is if let Some(want), so trapping runs compare the exit code and discard stdout — a class with zero coverage, now pinned by differential_trap_preserves_preceding_output as a verified negative control (PRE fails it). That test must pin a literal, because interp::run derives output from the IoToken's accumulated log and only on Done: interpreted output is a value that dies with an abort while compiled output is a side effect that survives, so the two I/O models diverge exactly on the trap path. Gate 981 passed / 0 failed, differential 37/37 in 403.93 s at -O0/-O2, fmt clean. 38 goldens moved: 37 ordering-only, one intended behaviour change (example_calc, both backends) signed off by Sapir. 22 were adjudicated by subagent panels (26 + 9 agents, every finding independently attacked by a refuter); the round covering the remaining 16 was stopped before producing a verdict, so those rest on the orchestrator's two mechanical sweeps alone — weaker evidence, recorded as such in the session log with the list and the resume command; 0 of 38 changed any tiling marker, guard count, attribute or trap count; and the observable-effect axis is closed by proof — 34 sequences unchanged, and the 3 where a PRINT crosses a mapal_par_check are safe because those modules hold zero mapal_par_trap sites and the only production writer of run.trap is mapal_par_trap, so check_trap's if trap != 0 can never fire. Perf: the "nothing moves" pre-registration is refuted — i9, 3 passes × 101 alternating, both faces, values byte-identical throughout: saxpy 1t +5.3% three times (the third pass a byte-identical rebuild, so it doubles as the ±1% noise-floor control), conv2d 1t −3…−8% but par +3…+7%, mm1024 +2.6% conformance yet flat on the FMA face — running the FMA leg (Sapir's catch) prevented publishing a one-face regression. Mechanism deliberately not isolated: vector-instruction counts are byte-identical pre/post (199/199, 295/295, 583/583) so it is scheduling not codegen, and %Frame member order vs task interleaving were not separated — the earlier "%Frame is the cause" claim was withdrawn as over-stated; S36c's %Frame alias barrier stays refuted. Incidental: --contract is a no-op on 4 of 7 ladder shapes (saxpy/reduce/transpose/gather emit byte-identical IR in both faces — contraction is applied only in tile kernels and those are not tile sites). And the GPU decision was taken on evidence: NVPTX, not CUDA C. An 8-agent audit recommended keeping the CUDA C emitter; Sapir challenged it; the audit's own most load-bearing claims were flagged unverified and a 15-minute llc -march=nvptx64 probe (LLVM 22.1.8) refuted themaddrspace(3).shared, ptx_kernel cc → .visible .entry, and llvm.nvvm.mma.m16n8k16.row.col.f32.f32mma.sync.aligned… with 804 mma/wmma intrinsics available incl. MXFP block-scale. Also corrected: the framing that two emitters over one IR violate FRAMEWORK §3/§5 is wrong — §4.2 sanctions one Trn at two Locs as "different code … the strategy shape", and §5 lists backends among pluggable variants. What survives from the audit: NVPTX does not build the smem rung for you (zero __shared__/__syncthreads/dim3 in backends/cuda/src), two real §5 duplications are language-independent (≈28 character-identical lines of type-erasure remap across the two backends; the mapal-rt ABI declared twice), and llvm/src/func.rs:342 packing_site is a CPU packed-format decision wearing a legality predicate's name (rename packed_layout_admits). CUDA C structurally cannot express ADR-0032 D1's per-region precision lattice, which is why NVPTX was always going to be forced. Also recorded at S38 close (external review, verified): guards are strict — an untaken arm still traps ((1 > 0) -> { -true-> 42; -false-> 7/0 } exits 101, and examples/calc.mapal's own header documents it). Sapir's direction is to fix the cause, not rename to select: a flow with a condition should have the condition execute and determine whether the path is taken; the IR exposes the flow as conditional and the backend honours it. Type-system context for the plan: Ty has products and no coproducts, and a branch is a coproduct — the same gap scan and first-class functions will hit. S39 opens on three Sapir directions: the GPU leg via NVPTX, beating OpenBLAS at one thread (a flat 1.20× behind at 1024/2048/4096, size-invariant ⇒ a micro-kernel deficit, measured on the untuned generic profile), and hardware-specific units in the backend (AMX / tensor cores as a per-Loc capability, never a mapal-ir fact). Previous: Session 37 (CLOSED — log sessions/2026-07-27-s37-elem-plan-and-the-dead-array.md) — the compiler now records what out[i] IS, and the biggest win came from deleting an array nobody read. Sapir reframed S36c's "iota is an index law" from an emitter special case into a graph law: a stage carries a structure, and composing stages composes structures. Written out, the five candidate stages have the same shape of element law and the same target object, which is FRAMEWORK §3's criterion for one object with extra structure — Enumerate needs no constructor at all, being Pair(Index, ·), i.e. enumerate a ≅ zip(iota n, a). elem_plan (mapal-ir/src/algo.rs) is the deduced query: ElemSrc is Index | Broadcast | Load (the cut) | Pair | Apply, three guards (single in-edge, outside every loop SCC, depth ≤ 16), producers recognised by an exact op-tag set because trap-freedom is a documented guarantee of those four tags specifically. It is a query, not a rewrite (Sapir): it records that an intermediate could be skipped and never removes it, which is what leaves the elide-vs-materialise decision with the backend — the answer differs per target (GPU bandwidth vs CPU L2 vs FPGA BRAM) — and keeps S27 rung-3's deliberate packing expressible. Same commit migrates tile_iota_size off literal op-tag matching: it asked op == Operation::Iota, what the node is tagged, where it means to ask what the element is, and tile_site calls it twice per site with a silent fallback to the scalar emitter; landed behind a pin first (tile_sites_pin.rs, 15 sources) because losing a site is a 4.0× cliff with no diagnostic. The pin surfaced its own fact: raw = 0 everywhere — no tile site is recognised before rewrite, so the entire tiled path is downstream of Inline/LiftLoops. The payoff was not the arithmetic. Once every consumer rebuilds the element the array is write-only, and saxpy's zip task was writing 8 MB per run that nothing read, inside the time bracket; %Frame also lost the 4 MB iota, 12 MB total. Baseline main @ c5f48c9, conformance face (verified 0 fmla), interleaved runs, program's own iter ms=: saxpy 1t 0.4769 → 0.0981 ms (4.86×), par 0.1860 → 0.0833 ms (2.23×); transpose/gather/fir/reduce/conv2d flat; matmul flat at 512, 1024, 2048 and 4096 — confirmed independently by emission, where no tiled kernel changed in any of 51 sources (total structural delta 291 lines, all of it GEP+load → trunc). Apply (Map-as-producer) ships as legality only and this backend declines it on measurement — enabling it put two calls inside saxpy's timed loop regenerating the inputs from the index, for 0.72× — which is Table B working as designed. Separately, the per-push differential sweep went 409 s → 15 s (27×) with the cross product intact: the cost was never the compile, it was minting and first-running 1,280 fresh binaries, i.e. macOS code-signature validation in a single-threaded system daemon outside the process tree (which is why the existing fan-out made it slower); merging 32 cases into one translation unit fixes it, and Linux's ld-bound version of the same total too. Two P0s were refuted rather than built: the %Frame alias barrier (S36c's 2.3×) does not exist in emitted code — a struct-field control vectorises with no metadata at all, and across 61 tasks in 7 shapes exactly one reports unsafe dependent memory operations, saxpy's Zip task, whose output nothing reads, while saxpy's timed loop already vectorises; and "halve the differential cross product" was unnecessary once the real cost was found. Open, and the reason the gate is red: open_inline catches Inline turning Trapped(IndexOob) into Trapped(DivZero) — pre-existing on main, seed pinned, and the fix (approach A: topo_order ties break on source position) is built, measured green on the 1,280-run differential, and deliberately reverted because it churns 19 goldens across three crates and reorders emission for programs that were never rewritten; ratified for S38 in components/ir/plans/plan-s38-trap-order-is-source-order.md. The work is on branch s37-elem-plan, 7 commits, not pushed and not merged. Previous: Session 36d (CLOSED — log sessions/2026-07-27-s36d-readme-editors-and-the-fma-question.md) — the S36 block closes. The README now carries current numbers only (Sapir's rule: pre/post-fix belongs in status and session logs), both build faces per table, and a regenerable matmul 4096 row — that row had no source behind it (only stale .ll; gen_flow.py emits the untimed form), so benches/matmul/matmul4096_cap_f32.mapal was written and the row re-measured: 245 ms conformance / 155 ms FMA against C++ 33,439 and NumPy 44.3, i.e. 216× the naive baseline at 4096² and 55× at 1024². The editor tooling broken by the rename is fixed — all three failures were stale build artifacts, not code: the only built vsix was flow-lang, the nvim plugin still required flow.icon after the module moved to mapal/icon.lua, and only FlowIcons.ttf was ever installed. Rebuilt mapal-lang-0.1.0.vsix and MapalIcons.ttf (which reports family MapalIcons internally — the ADR-0037 font item needed a rebuild, not a fix), installed both, rewrote the nvim plugin and kitty's symbol_map. And the FMA question got an answer: it cannot be deduced, because fusing changes the value — it is a permission on ADR-0032 D1's lattice, the same one reassociation rides, and Ty carries neither. The default cannot flip today because contraction is a flag LLVM interprets, so the interpreter cannot predict the fusion set and the 1,280-run byte-equality oracle would have to weaken to a tolerance; the fix is Operation::Fma as a Core op with mul_add in the interp, after which the default can be contract (worth 1.62× on matmul) with byte-equality intact — components/ir/plans/plan-s37-scan-recurrence.md §8, ir suggestion #3. Checked and corrected in passing: Rust does not contract by default (0 FMA in both baselines) — the README had claimed it did. Previous: Session 36c (CLOSED — log sessions/2026-07-27-s36c-the-real-gaps.md) — two published numbers were wrong, and the two largest measured wins in the tree are not where anyone was looking. (1) Every Mapal leg ever published was the CONFORMANCE face: EmitOpts::contract defaults off and no harness passes --contract, so the emitted object carries zero FMA instructions while C++/Rust get FMA from -ffp-contract=fast and NumPy from BLAS. Re-emitted (0 → 28 vfmadd in the same object): the OpenBLAS gap on identical hardware is 1.23×, not 1.85×; matmul 1t 17.42 → 14.82 ms; on the Mac par 3.65 → 2.25 (its residual 3.3× is Accelerate reaching AMX, which is not a compiler comparison). (2) The reduce row was a semantics gap plus a governor artifact — Rust's baseline reassociates and its answer depends on thread_width(), NumPy's is pairwise, and at one thread where all three compute the same left fold Mapal is the fastest (0.3668 vs 0.3821 vs 0.3821); separately, perf shows the i9 kernel costs a constant 2.1 M cycles in every configuration while wall time swings 0.38–2.01 ms with the boost clock, so that box's sub-5 ms cells must be reported in cycles. (3) The non-compute shapes are not blocked by arithmetic intensity. Two self-inflicted memory-layout facts: the %Frame struct destroys LLVM's alias analysis even though build_frame_layout already proves disjointness (2.3× on saxpy 1t), and iota is materialised as an array instead of being an index law so every map over it is an indirection (3.1×). Together they put saxpy at parity with clang -O3 -march=native with zero mapal-ir changes — closing the S35 finding without touching the tile ladder. Map fusion turns out to be already shipped and to fire on nothing; a Scan primitive unifying fold/window/scan is planned in components/ir/plans/plan-s37-scan-recurrence.md. Previous: Session 36b (CLOSED — log sessions/2026-07-27-s36b-cross-machine-validation.md) — the clock-read barrier validated A/B on seven shapes and two machines, 8,400 timed runs (benches/results-s36/). Each shape was emitted twice, by the compiler at 35fb681 and at 896fb3c, from a worktree so both binaries exist at once. The finding is that the 0.01 ms race counter was a proxy calibrated to fir and it missed the worst case: matmul1024 pre-fix reported a par minimum of 0.0209 ms against a 31.22 ms single-threaded median — an apparent 1494× on 14 cores — without tripping it. The test that needs no calibration is the machine: a cell is impossible when 1t_median / par_min exceeds the thread count. On that test, pre-fix 3/7 cells on the M4 Pro, 5/7 on the unpinned i9 and 3/7 on the pinned i9 were impossible; post-fix 0/7 everywhere; post-fix maxima are 8.6× on the Mac, 11.6× on the pinned box (16 threads) and 24.4× on the unpinned box (inside its 32-thread bound, above its 24 physical cores, and ramp-inflated — 3.8× against that cell's 1t min). Two limits stated rather than glossed: reduce and saxpy never showed the defect at all, so their post-fix result is a control not a repair, and n=100 with zero hits bounds the residual rate near 3%/run at 95% rather than zeroing it. MAPAL_PAR=1 is the control: unmoved on the pinned box and the Mac (six of seven pinned cells within 1.5%), with byte-identical output on all seven shapes (value_identity.log). On the UNPINNED box it swings up to 34% between runs — a leg with no workers cannot race, so that is the machine, and it is why the pinned log is the one to quote. Two things the campaign separated from the defect: the i9's powersave governor (a 5× spread visible on the 1t leg, which cannot race — pin the CPU) and post-fix par spread on small kernels (pool dispatch dominating a 50 µs kernel; both ends physically reachable). Also fixed: ladder2_baseline.cpp used std::string without <string> and would not build under gcc 16 — the baseline leg of a benchmark, unbuildable on the measurement box. Previous: Session 36 (CLOSED — log sessions/2026-07-27-s36-clock-read-barrier.md) — S33's remaining P0 is closed: mapal_par_wait let workers run ahead of the clock, so a threaded kernel could finish before the time read meant to bracket it. The fix is plan-s33b §3 as written — a clock read stops being a bare host-spine checkpoint and becomes a pinned DAG node with edges both ways, so the work written after it cannot be dispatched until it fires. It landed entirely in mapal-ir's path_plan: crates/mapal-rt/ has a zero-line diff and the emitter needed no new concept, because the pinned-task machinery built for trap-capable calls already carried it. fir 65 536 at MAPAL_PAR=14: 6/100 readings under 0.01 ms → 0/100, values byte-identical, total wall time unchanged (2.73 → 2.62 ms median) — and the self-timed interval RISES ~12%, which is the correction: work that used to start before t0 is now inside the bracket. Acceptance 2 was amended by measurement rather than declared met — the residual min/median gap is pool wake-up jitter and tracks kernel size (0.57 at fir 65 536, 0.84 at fir 1 048 576, zero sub-0.01 readings at either), so min is trustworthy again where the kernel dominates wake-up and tiny cells stay on medians for a reason unrelated to the clock. Gate 972 passed, 0 failed; both launch-contract guard rails pass unmodified. Previous: Session 35 (CLOSED — log sessions/2026-07-26-s35-shape-ladder-and-the-ast-question.md) — the shape ladder grew four non-compute classes (saxpy, reduce, transpose, gather), each with a C++ and NumPy baseline and identical printed values before timing. Threaded, Mapal takes transpose and gather and loses saxpy; single-threaded it loses all four. The finding: a plain map is not a tile site, so streaming and permutation kernels emit scalar loops — the saxpy task has one fmul and zero q-register loads, while the generation task in the same binary is full-width NEON. That is the first measured bound on how far the graph deduction generalizes past dense compute. Also: the README claimed the compiler translates source "instead of a traditional AST", which is false — crates/mapal-syntax/src/ast.rs is a 384-line recursive tree — and is now corrected to the accurate and stronger framing (the syntax is a serialization of the execution graph; optimization happens on the graph, dataflow-first, not on the control-flow-first IRs everyone else uses). "Skip the AST" was priced rather than argued: lex+parse is 4.4 µs against 140 ms of clang, 0.003%. Ten syntax snapshots invalidated by the S34 rename (recorded byte offsets over rewritten fixtures) were regenerated and verified against source bytes; gate green at 971 passed, 0 failed. Previous: Session 34 (CLOSED — log sessions/2026-07-26-s34-mapal-rename-and-trap-deleting-rewrite.md) — THE PROJECT IS NOW NAMED Mapal; source files are .mapal (ADR-0037). מפל, Hebrew for waterfall: flow falling through stages, and it contains map, the central operator. Forced by evidence rather than taste — flowc on crates.io is someone else's dataflow-language compiler (92k downloads), and Meta's Flow owns the search. Renamed: 8 crates, 1,409 runtime symbols, 202 environment variables, 49 source files, 323 living docs/editor files, 83 TextMate scopes. The extension went .flow.mp.mapal: our own editor suite caught .mp colliding with MetaPost in every Vim install (mfNumeric groups leaking in), and GitHub's Linguist would have mislabeled all 49 files. .mpl is JetBrains MPS, .ml is OCaml, .map is source maps — all rejected on evidence. Immutable by ADR-0037 D3: session logs, ADR-0001…0036, docs/performance/** and recorded result CSVs keep the old name, because a number's provenance includes the binary that produced it. Lowercase flow survives as the name of the language construct (a -> b; is a flow statement). Gate green: full suite, differential 36/36 in 465 s, fmt clean, 61 editor assertions. Previously this session — the P0 is CLOSED and the gate is GREEN. The rewriter's trap deletion was map(id) → id, not fusion proper: functor_laws.rs:is_identity_body judged a map body by its Return writer alone, so a body that returns its parameter and computes a dead trapping Div read as the identity and the entire Map was aliased away — Trapped(DivZero)Done(0). The earlier passes were the enabler, not the culprit: ConstFold's proj∘pack forwarding is what puts a body into identity shape while DCE correctly keeps the impure dead Div (R4). Re-shrunk to a 3-step program (PROPTEST_MAX_SHRINK_ITERS=200000, 0.4 s) which is why the S33 bisect pointed at MapFusion. The law's precondition is an equality of morphisms, so the guard now quantifies over the body's whole morphism set via graph_rewrites::is_pure — now pub(crate), because "may this dead cone go?" (DCE) and "does this body denote id?" (the identity law) are the same question and two lists would drift. All four failing property entry points were this one bug; the pinned seed is retained and passes. Negative-controlled: with the guard reverted, the new hand-written case and all four suites fail again. cargo test --workspace --release green (LLVM differential 36/36 in 462 s, so it ran rather than skipped); cargo fmt clean. Also this session: the public-repo community files — CONTRIBUTING.md (model-first workflow, the measurement rules, the open-ADR pickup list), CODE_OF_CONDUCT.md, SECURITY.md, three issue forms and a PR template. Detail: components/rewrite/plans/plan-s34-identity-map-trap.md. Remaining P0: the mapal_par_wait clock race (measurement-only). Previous: Session 33 (CLOSED — log sessions/2026-07-26-s33-boundary-openblas-parity-open-source.md). The repo is now PUBLIC (github.com/LessComplexity/mapal) with CI, and CI is RED for a real reason: on its first run a randomised property test found that the rewriter deletes a trap that must fire (open_default: original Trapped(DivZero), rewritten Done(I32(3))) — pre-existing at 1daddaa, pinned as a proptest seed so it cannot pass on a lucky draw, and now the project's highest-priority item. Second new P0: mapal_par_wait lets workers run ahead of the clock (3–4% of threaded runs self-time far too low; one live case read 0.0001 ms), which inverts the measurement rule to min for 1t, median for par and makes every par minimum in S28–S32 suspect. A runtime-only fix was built and REVERTED — see plan-s33b-clock-read-barrier.md §4 for why not to retry it. Also this session: the emitted .ll/.cu artifacts (112 files, 4.4 MB) are no longer tracked and regen.sh derives its own worklist; the README was rewritten results-forward with category theory as a first-class section; editor support gained a TextMate grammar, a real logo (SVG + our own single-glyph font at U+F8F0, because the Rust/C++ marks in a file tree are font glyphs), and 61 test assertions across both editors after five reported bugs. Previously (S33 in progress) — the conv2d "per-core gap" is CLOSED, and it was never a kernel defect. mapal_rt_alloc returned a reserved address range rather than memory: a large alloc is served by mmap, which delivers no physical pages, so the first store to each page trapped into the kernel (2 MiB of zeroing apiece under THP) inside whatever () -> time region wrote first. conv2d exposed it because nothing writes its output array until the convolution itself does, while the C++ baseline pre-pays that cost above its own timer (std::vector<float> out(n) value-initializes). Shipped reside (one byte per 4 KiB after each alloc — not a memset; the fault's own zeroing IS the initialization, and a memset would double the traffic, ~15 ms wasted on a 64 MB matmul frame): no emitter change, no mapal-ir change, no .ll moved (the gate was green at that point; it is now red for the unrelated rewriter bug above). Proven before the fix existed by exact differenced counters — flow's kernel 905,100 cycles / 299,221 ref-cycles / IPC 2.25 vs cpp's 1,072,928 / 382,489 / IPC 1.78, i.e. 18% fewer cycles and 22% less real time, with ref-cycles (frequency-invariant) predicting the post-fix window at 0.150 ms against 0.144 measured. The recorded "IPC 3.11 vs 1.57" was process-level, contaminated by Mapal's generation legs (IPC 0.86–1.04) whose instruction counts differ 7× from C++'s. Alternating same-session A/B: conv2d 1.72–1.89×, fir 1.12–1.17×, matmul 1.12→1.04→1.01→1.003× as N grows — the effect scales as output-size ÷ kernel-length, so matmul is immune. conv2d is now 1.21× AHEAD of naive C++ per core on BOTH NEON and AVX2 (was 1.56× behind). All eight previously-eliminated in-kernel hypotheses were correctly refuted — the kernel was never the problem. Second machine, full suite: on the i9-14900F (no matrix coprocessor, numpy → OpenBLAS 0.3.30 on the same AVX2 units) Mapal's generated GEMM reaches PARITY with hand-tuned OpenBLAS — 1024² dead even (1.526 vs 1.506 ms) — while the M4's AMX-backed numpy is 3.3× ahead, so the M4 matmul gap is silicon, not code generation. Stated in full rather than by its best cell: 1t is a flat 1.20× behind at 1024/2048/4096 (146 vs 174 GFLOP/s, both size-invariant ⟹ a steady micro-kernel deficit, not a blocking failure), threaded within ±10% (ahead 1.08× at 2048, behind 1.24× at 512 and 1.06× at 4096). Confound ruled out: OpenBLAS default == OPENBLAS_NUM_THREADS=32, so it does use the whole machine. The scheduler claim was TESTED and REFUTED as stated. Mapal scales 9.2–9.8× vs OpenBLAS's 7.1–8.1×, which invites "our scheduler is better". Controlled experiment on the same box — same binaries, 8 threads in every cell, only CPU uniformity varying (cpu 0–15 are P-core threads at 5.5–5.8 GHz, cpu 16–31 are E-cores at 4.3): 8 E-cores (uniform) flow 5.894 / numpy 5.591 → numpy 5% ahead; 8 P-cores (uniform) flow 2.438 / numpy 1.724 → numpy 41% ahead; 4 P + 4 E (mixed) flow 3.384 / numpy 5.573 → flow 1.65× ahead. On uniform cores OpenBLAS wins; Mapal wins only when cores are mixed. Mechanism visible in numpy's own column: 8 E-cores 5.591 → 4P+4E 5.573, i.e. swapping in four 35%-faster cores bought OpenBLAS nothing (static partitioning waits on the slowest thread), while Mapal went 5.894 → 3.384 = 1.74× (work stealing absorbs the slack). So what Mapal has is heterogeneity tolerance, not a better scheduler — the full-machine parity exists because 16 of 32 threads are E-cores, and on homogeneous server hardware OpenBLAS should lead at both widths. Still valuable (nearly every consumer CPU is hybrid) but NOT a claim about beating BLAS on a server. The 1t gap ran on the untuned generic profile — no hardware excuse in it, and the honest remaining target. On the i9 Mapal wins every shape at every size against every baseline including numpy. The fix's effect is also platform-dependent: conv2d 512 gains 1.89× on macOS's 16 KiB pages and exactly nothing on Linux, where a 1 MB output hides inside a 2 MiB huge page its neighbours already faulted. NEW P0 — a help-first race in mapal_par_wait: the waiting host executes work past the watermark it waits on, so a kernel can finish before the clock meant to bracket it starts. 3–4% of threaded runs self-time far too low (MAPAL_PAR=1: 0/100, so every 1t number is sound). This inverts the standing "quote min, never median" rule — frequency ramp makes SLOW outliers, this race makes FAST ones, so min is maximally vulnerable and worsens with N. Pre-existing (PRE 3/100, POST 3/100 against a HEAD~1 runtime): every par minimum in S28–S32 is suspect and the S32 scheduling verdict needs re-confirming under a median. Perf: performance/matmul/s33.md. Previous: Sessions 31+32 (CLOSED — log sessions/2026-07-26-s31-s32-deduced-blocking-and-scheduling.md). Shipped: TargetProfile (six hand-swept constants become one named table plus arithmetic; generic byte-identical, verified as 66 A/B emissions; the KC gate now closes by derivation), i_reuse-driven row blocking (ci == 0 and ci == cq are ONE predicate at q=0/q=1, so conv is blocked because the record says its read slides — conv2d −25% at 1t, FMA:load 0.80 → 1.20), and per-region slice sizing (the pool receives sizes instead of inventing them; slices cut on the region quantum — matmul512 1.43×, matmul1024 1.41× at the default width). Also fixed a 1.8× slicer defect: it derived a count then equal-divided n, leaving ragged pieces. OPEN diagnosis: conv2d's kernel is 1.55× slower than naive C++ on BOTH NEON and AVX2 — same ratio on unrelated architectures, so it is structural. Cache exonerated on both machines (flow-zen3 has fewer misses than C++ and still loses); IPC is the gap, 3.11 vs 1.57. Eight hypotheses eliminated by measurement — see the log §3 before proposing another. Our tiling is not the problem: 8.6× over Mapal's own untiled path. Previous header: Session 31 (IN PROGRESS) — plan-s31-deduced-blocking items 2/3/4 SHIPPED: conv2d is −25% at one thread (0.5343 → 0.3992), closing cpp-1t from 2.09× to 1.56×, with blocking applied because the RECORD says the read slides (i_reuse: ci == 0 and ci == cq are one predicate at q=0/q=1) and TI taken from the register file, not a literal. FMA:load 0.80 → 1.20 in the disassembly. Two predictions of my own were refuted by measurement and recorded as such: the vector accumulator's 2.9× (delivered ~10% — LLVM was already promoting it) and the ordering built on it (row blocking beat it, ~17%). conv2d's remaining deficit is now majority scheduling — optimum 4 threads vs the default 14. Previously this session — TargetProfile: the emitter's machine facts become data. Sapir's redirect reframed the S31 opener: row blocking must be generic per algorithm with TI detected from the execution graph, not a conv2d special case. Investigating that produced three findings that reordered the work. (1) The geometry half needs nothing new in mapal-irci == 0 (matmul) and ci == cq (conv) are the same reuse predicate at q=0 and q=1, pure arithmetic over recorded TileRead fields. (2) The magnitude half is not a graph question at all — it is the vector register file, which existed nowhere in the repo. So TargetProfile (next-session item 0) had to land first, and it did: crates/backends/llvm/src/profile.rs replaces tile_j_for/TILE_I/TILE_KC/tile_nc_for/HEAP_MIN_BYTES with one named table (generic/apple-m/zen3) plus five derivations; generic reproduces every literal, and tile_i = vec_regs/(2×acc_vecs_per_row) reproduces S26's swept 4 including its recorded TI=8 spill. The KC gate now closes by derivation (16 MB L2 ⟹ kc=4096 ≥ K) — pinned in the strong form that apple-m WITH --kc is byte-equal to generic WITHOUT it. Threshold, not off-switch (Sapir's catch): at K=8192 it reopens, and there the derivation disagrees with S30's measurement (the nest lost at every size, deficit growing in N) — sound as "how deep a panel fits in L2", not as "when the nest pays"; default-OFF keeps it out of shipped builds and the box leg settles it. Value-invariance proven under zen3 (TJ 16→32, TI 4→2, oracle-equal at -O0/-O2). (3) TILE_I was two quantities sharing a number: the FIR window rung blocks lanes over a memory accumulator no register budget bounds — now WINDOW_SUBROWS, named apart. Rule 1 was checked as 66 byte-identical A/B emissions against HEAD, because benches/matmul/regen.sh cannot serve as the gate — the 72 checked-in .ll are stale at HEAD (pre-existing: they predate S30b's time migration, and regen.sh exits 1 on the CUDA leg which rejects time). Deferred with reasons: native (vec_regs is not probeable) and the kc_nest tri-state (its auto would have enabled the nest by default for every K>128 site, breaking rule 1). Still open — the item this unblocks: deduced TI + conv2d row blocking, planned in components/backend-llvm/plans/plan-s31-deduced-blocking.md, whose mem-op accounting says the queue order is backwards (the conv vector accumulator is ~2.9× and row blocking ~1.2×; the reuse factor at TI=4 is 2.0×, not the 3× in suggestion #11 — that is the TI→∞ limit). Also this session: examples/matmul4.mapalmatmul4_loop.mapal (Sapir's rename) followed through into its three test call sites and two snapshots. Previous: Session 30 + 30b (CLOSED — logs sessions/2026-07-25-s30-vector-accumulators.md + -s30b-measurement-and-readme.md). S30b: the full CPU comparison, and the finding that the default thread count is wrong. matmul migrated off MAPAL_PERF (gen_flow_capture.py brackets the kernel with time; new benches/matmul/matmul_ab.sh), so the whole 512→4096 × f32/f64 × all-CPU-legs matrix is compute-only on both sides for the first time: Mapal beats naive threaded C++/Rust 11× at 512² growing to 192× at 4096², and NumPy stays 4× ahead on matmul (matrix coprocessor) while Mapal is 4–16× ahead of NumPy on fir and conv2d. A thread sweep then found the default (every core) is the WRONG width for three of four benchmarks — conv2d is 2.1× faster on 4 threads than on 14, fir peaks at 8, matmul 1024 at 8, only matmul 4096 wants all of them; a C++ control degrades at 14 too (the chip is 10 P + 4 E), but ~2× ahead of us at EVERY width, so conv2d's gap is per-core (TI=1: 24 vector loads per 36 FMAs against matmul's 4 per 32), not scheduling. Sapir's direction, recorded: the thread count is deducible — graph facts (element count, work and bytes per element) × machine facts (core count, P/E split) — and belongs with TargetProfile. Also: README.md + LICENSE (Apache-2.0 with the LLVM exception) for open-sourcing, the README rewritten after an adversarial audit found four framing problems (naive baselines unlabeled, benchmarks from the non-bit-exact --contract face, a CI claim with no CI, and a headline contradicted by its own conv2d row). Previous header: Session 30 (CLOSED — log sessions/2026-07-25-s30-vector-accumulators.md). S30 cashed the S29 diagnosis: tile accumulators are now SSA values, not stack memory. The constant-width main tile carries TI accumulators of type <TJ x elem> across the k loop by phi — no alloca, no GEP, no accumulator load/store inside the loop. Result: the KC leg's str q…,[sp] 92 → 0, its runtime alias checks 8 → 0, and its hot loop is instruction-for-instruction the baseline's; KC-on 1024 f32 59.9 → 21.7 ms, f64 158.4 → 45.2, 2048 488.5 → 176.3, 4096 4097 → 1564. The shipping path is unchanged within noise but is now emitted in register form rather than granted it by an LLVM heuristic — which is the durable part: the S29 regression happened because that grant was withdrawn when unrelated code touched the same slot. With both legs fair, the KC traversal still loses at every size on M4 Pro and the deficit grows with N (+5% → +14%), so kc_nest stays default OFF for a measured reason; the box leg is now a fair test of the order. Gate green (72 suites). Previous header: Session 29 (CLOSED — log sessions/2026-07-25-s29-kc-verdict-time-builtin.md). S29 opened on a broken tree — the previous session was killed mid-experiment — and closed with the workspace green, three builds finished, and two of the session's own headline claims corrected by measurement. (1) Repair: the tree did not compile (5 Operation::TimeMs/ExprKind::Unit match arms unwired) and the KC k-panel nest was frozen mid-experiment (// EXP-A markers, TILE_KC left at a probe value, an orphaned jb0 param). The experiment was FINISHED, not reverted: with the (jc, kc, ic) order partial sums park in out at every panel end, so the accumulator is one j-tile wide and TI×NC would be 32× dead space. (2) The KC verdict — a 3× LOSS (tile_ab.sh matmul1024_cap_f32, MAPAL_PAR=1: fma 59.82 ms with the nest / 19.80 without, the OFF column reproducing S28's 18.9). Diagnosed, and the first explanation was wrong: parking is ~2.5% of the gap; the cause is codegen — the [64 x float] accumulator alloca is register-promoted across the k loop in the jt-outer leg and NOT in the KC leg (92 str q…,[sp] vs 0), so every FMA round-trips the stack. Invariant across every tile constant, so the follow-up is a promotable-accumulator fix, not a re-tune (s29.md §1, suggestions #16). Shipped default-OFF behind EmitOpts::kc_nest — bit-exact either way (differential-enforced), kept because it was designed for box-scale traffic (4096 on zen3) where it is not measured. (3) The time builtin, end to end (Operation::TimeMs : IoToken → (IoToken, f64); () -> time is the first wire-LESS stage; () now parses to ExprKind::Unit instead of P0001): syntax→lower→check→ir→interp→rewrite→llvm→rt, cuda a recorded ✋ cell, no new L-code. Three defects found and fixed while building it, each mutation-verified: a clock read raced the tasks it bracketed (no value producer to wait for ⇒ path_plan now fences every task written entirely before it IN THE SOURCE — source order, because the dataflow graph orders pure work against a clock read not at all); a clock value consumed by a task produced a NEGATIVE elapsed (§4.5 Law 1 data teleport ⇒ the consumer cone stays on the host spine); and a loop-body read was hoisted out of the cycle (two of lower's four effect detectors still tested print alone). (4) Heap lowering: HEAP_MIN_BYTES = 256 KB gates FnEmit::entry_allocmapal_rt_alloc arena with one mapal_rt_free_all() before the entry ret; matmul2048_cap_f32 runs locally for the first time (-1045/51275) — it used to SIGSEGV on the 64 MB macOS stack. Entry-fn-only is the recorded ceiling (BL9). (5) The first honest kernel-only shape numbersbenches/shapes/*.mapal self-time and --perf is retired there: fir wins every column at both sizes (1M fma-par 0.402 vs cpp-mt 1.462 = 3.6×, numpy 6.368 = 15.8×, cpp-1t 11.395 = 28×), conv2d beats cpp-mt at 512 (0.083 vs 0.112) and LOSES 3.4× at 1024 (0.445 vs 0.133) — the TI=1 row-blocking ceiling is now measured, not predicted (suggestions #11/#12/#17). This corrects S28: its "conv kernel ≈0.04 = 3× over cpp-mt" was a subtraction, not a measurement; the real 512 kernel is 0.083/0.107, ahead by 1.3×. Gate: cargo test --workspace --release green (72 suites); every pre-existing llvm golden untouched. Box leg NOT run — the KC verdict and every number above are local (M4 Pro). Previous header: Last updated: 2026-07-24 · Session 28 (CLOSED — log sessions/2026-07-24-s28-shapes-ladder.md; commits pending Sapir). S28: the ladder generalizes to fir & conv2d (Sapir's S27c focus directive) + the S27 box debt paid. (1) mapal-ir k-split record (A1): TileRead.ksplit? : TileRead → TileKSplit{div,cq,cr} — the fold-body Div/Mod pair on the counted axis (shared literal, depth % div == 0) binds derived (k÷div, k%div) walker axes — the map-body (t÷C,t%C) move one level down; rules: XOR raw-k, unused pair ⇒ None (matmul/fir records bit-identical); conv2d_16 records its site (S27c's priced refusal cashed). (2) FIR 1-D window rung (B, zero mapal-ir change): window1d_siteemit_tiled_map_blocked_1d — the rung-2 DUAL (TI=4 blocks over the lane axis, ONE scalar w[k] per k shared across subrows, constant-TJ main, ×2 k-unroll; remainder = TI=1 j-split; non-window 1-D byte-stable). fir both tables WON local (fma-par 0.2133 vs cpp-mt 0.2395 / rust-mt 0.3017 / numpy 0.3932; fma-1t 0.2156 vs cpp-1t 0.9239 = 4.3×) AND box (0.287@16T vs numpy-1t 1.39; 1t 0.786 = 3.4× cpp-1t). (3) conv2d unrolled micro-kernel (A3): conv_siteemit_tiled_map_conv — the (kq,kr) taps fully unrolled at compile-time offsets, ZERO div/mod; non-conv ksplit keeps the untiled fallback (rule-3 guard); packing_site excludes ksplit. conv kernel ≈0.04 par = 3× over cpp-mt local; box par table WON at the leg level (0.742 vs cpp-mt 2.32); the M4 par leg stays OPEN on the gen measurement boundary (MAPAL_PERF brackets the untiled img-gen the baselines exclude — recorded finding, suggestion #14). (4) Box (S27 debt): "balance 0" was an agent misread (credit 15.41 + autobill — corrected per Sapir); box #1 (45692618) VANISHED mid-run (interruptible-class), relaunched on-demand 45712913 (EPYC 7B13 zen3, destroyed ≈$0.45): full matrix + S27/S28 shapes — disasm gates pass (conf 0 vfmadd / fma 128 vfmadd 0 unfused = the S26 finding CLOSED), 2048/4096 flow rows clean, OpenBLAS frontier measured: threaded 9.7×/5.9× ahead @1024/@4096 f32, numpy-1t 2.7× ahead of fma-1t wall (agenda-2 target), GRAIN quantization measured (fir 61T 0.526 → 16T 0.287 — 16 slices = 0.26 waves @61T; suggestion #15). Gate: full workspace green (69/69); matmul .ll byte-identical; 180 ir · 53 llvm (28 differential + 25 golden). Previous header: Session 27+27b+27c (CLOSED + COMMITTED — three logs sessions/2026-07-24-s27*.md; S27c: the local same-machine matrix + shapes baselines + the compute-vs-compute fairness rule (Sapir's catch — flow ahead at every size on the fair basis; conv2d refusal priced 7.9× = the S28 focus); box still owed on balance). Previous header: Session 27+27b (CLOSED — sessions/2026-07-24-s27-fma-packing-fnstrip.md + -s27b-loop-lift-panel-residence.md; box leg BLOCKED on vast.ai balance 0 — everything else done). S27: the three numpy-gap closers + fn-strip, shipped and locally measured. (1) FMA contraction (product face): hypothesis verified in one golden (LLVM≥14 gates contraction on per-instruction IR flags — plain IR + -ffp-contract=fast → 0 fused; flagged → 34 fmla.4s, 0 unfused) → EmitOpts::contract default-OFF, tile-kernel fmul contract/fadd contract; conformance face bit-exact untouched (S24b's CPU twin); _fma.ll artifact twins + tile_ab fma leg (rel-tol vs max(|e|,1) + disasm asserts). (2) BLAS rung 3 packing: b read through a packed j-tile-major 64-aligned panel; parallel flavor packs in a run-once wrapper task before a nested global-pool slice dispatch (help-first finish ⇒ width-1 sound); Seq-task/loop sites pack inline per iteration (orchestrator review caught codex's uninitialized-buffer hole on loop-membered sites — loop-carried-b regression pinned, pre-fix garbage); --no-pack A/B control. (3) Micro-kernel finishing: per-width TJ (f32 16 / f64 8), k-unroll ×2, packed-line prefetch. Local 1t @1024: f32 tile 32.8 → fma 19.3 ms = 1.81× vs S26; f64 64.0 → 38.6 = 1.73×. (4) fn-strip WIRED (Sapir directive): PassId::Inline first in default rewrite(); loop-bearing callees never inlined (nested-SCC guard); cap 64→256; map-body Calls strip — a map body calling a helper fn TILES (differential_tiled_matmul_via_helper_fn). S27b (same-day continuation, Sapir's close-review directives): (5) panel residence — jt-outer nest for packed sites (the k-panel deferral corrected: packing fixed access pattern, not volume), per-thread b-traffic ÷4 @1024 / ÷16 @4096, ZERO acc spill, byte-exact; KC-split stays box-gated. (6) loop→map/fold lifting SHIPPED (ratified in-session: "even loop naive implementations enjoy the perf boost automatically"): PassId::LiftLoops after Inline — R-LF/R-LM guarded-trace lifts off loop_plan facts, K≥1 (adjudicated in-flight — Core has no empty arrays), every rejection pinned; matmul4 loop form lifts → inlines → tiles, -275/3748 byte-exact at -O0/-O2 × {default, MAPAL_PAR=1} — THE S26 NON-TILING PIN INVERTED; fir's loop fold-lifts too; loop-form bench legs (matmul16..128) regenerated TILED — the N⁴-wall legs retire. (7) suggestions #10: block_plan backend-generic schedule query (Sapir direction, gated on cuda as second consumer). Harness to 4096 + fma legs; s27_box.sh prepped (S28 opener). Full workspace green orchestrator-run (72 suites, fmt clean): ir 176 · llvm 47 (23 differential incl. 1280-run + lift acceptance · 24 golden) · rewrite 68 · rt 16. Previous: Session 26+S26b (CLOSED — sessions/2026-07-23-s26-register-blocking.md + 2026-07-23-s26b-par-on-par-reframe.md; commits pending Sapir's confirm). S26: BLAS rung 2 — TI=4 register blocking + the fixed-TJ main/remainder split, shipped end-to-end in one session: backend-llvm func.rs only (TILE_I=4/TILE_J=16, gate rows>1 && b.ci==0; mapal-ir untouched; per-cell op order preserved ⇒ byte-exact R1, differential-gated). Local A/B (MAPAL_PAR=1, min-of-3): tile vs no-tile 512 f32 12.8× · 1024 f32 23.4× · 1024 f64 12.3× · attn_256 12.1×; disasm 34 fmul.4s+34 fadd.4s, zero scalar; TI sweep 2/4/8 → 4 (8 spills). Box (EPYC 7B12 zen2, cgroup-v1 quota 61.44 → pool 62, clang-18 via llvm.sh): flow vs chapel-multicore f32@1024 15.9 vs 117.4 (7.4×) · f64 23.3 vs 118.3 (5.1×) · N=256 flips flow 1.4× — chapel loses every cell ≥256; numpy f64 gap 13.8× → 7.4× @1024; 1t f32@1024 84.9 vs S25's 568.3 (6.7×), flow 1t vs cpp 1t 89×; par f32@1024 flat vs S25 (memory/startup floor — rung 3 owns it). Disasm: full-width AVX2 ymm, 0 xmm (the split worked) but vfmadd absent under -ffp-contract=fast (~2× FLOP density left — recorded finding). Sapir standing directive: comparisons same-machine + machine specs stamped ON every results CSV (runner.py). matmul4 loop form answered: does NOT tile (graph-shape detector, byte-identical emission; its cap-form twin tiles — loop→map lifting is a rewrite-level rung, Sapir's call). 906 green (ir 176 · llvm 42 — differential 19 incl. 2 new tiled cases + goldens 23 re-pinned · rt 16 · others 672); box destroyed (≈$0.12). S26b (Sapir framing directive — 1t-on-1t / par-on-par verdicts only, mt-flow-vs-1t-baseline rows dead): quota-aware threaded cpp/rust baselines (cpp_mt/rust_mt, cgroup rule = mapal-rt's own) + numpy-1t/chapel-1t legs, one mini-box (destroyed ≈$0.036, every out= byte-equal to s26.csv). Par-on-par: flow beats EVERY threaded naive-class baseline — cpp-mt 3.1–10.9×, rust-mt 3.0–9.5×, chapel-mc 3.1–9.6× (f32@1024: 12.7 vs 138.1 = 10.9×; the deleted 475× row's honest form); numpy-threaded still ahead 3.9× f32 like-for-like (flag: numpy_bench.py runs fp32 — the standing f64 pairing overstates flow's gap; Sapir's call); numpy-1t beats flow-1t 3.3× f32@1024 (the pure kernel gap — rung 3's target). Previous: S25 CLOSED — tile emission v1 — the BLAS ladder's rung 1, shipped end-to-end in one session (Sapir directive "continue until 1+2 implemented and tested" + the mid-session shapes/verify directive): mapal_ir::tile_plan (affine-triple recognition, 2-D + 1-D lane modes, proof-gated) → the backend-llvm TILE_J=16 register micro-kernel (per-cell chain order exact ⇒ stdout byte-equal at any thread count/opt level — R1 by construction, differential-gated) + cgroup-quota pool width + the MAPAL_PERF llvm compute timer. Box (EPYC 7702P, 62-core quota): flow 3–8.6× AHEAD of chapel-multicore at 512/1024 both widths; numpy gap 130× → 13.8× f64@1024. Shape corpus: attn tiles (2 chained sites, 4.6× local), fir tiles (large-K 2.2×; small-K clang self-SLPs), conv2d/rowmajor refusals exact. Suggestion #9 found already-shipped (S20c) — s24 attribution corrected. 904+ green; box ≈$0.55. Previous: S24 — the parallel orchestrator: Sapir's parallel-first directive shipped end-to-end in one session (ratified plan → mapal-ir path_plan task-DAG query → mapal-rt work-stealing scheduler → backend-llvm parallel mapal_main with the speculate-and-order trap protocol), R-PAR proven live (output byte-equal to the oracle at any thread count, -O0/-O2), and measured on a fresh box: N=1024 f32 flow-llvm 184.0 ms vs chapel-multicore 192.7 ms — flow AHEAD, 19.1× over its own single thread, 18× over single-thread C++ — the S23 60× gap closed (docs/performance/matmul/s24.md). Two codex-round review finds fixed pre-merge (pinned pool race; effectful-loop seed hoist) + the v3 protocol derivation when review killed the draft's abandon semantics. Workspace green; box destroyed (≈$0.25). Previous: S23 — the S22 mandate's tail delivered to its done-bar: WP-D hoisting (64f1f50, codex-implemented per the restored delegation split — root cause of the "network-dead" codex was a STDIN wait, </dev/null fixes it), WP-E llvm assessment (deferred with sroa measurements, suggestions #10, 1daef83), and the FULL S23 performance matrix on a fresh 4090 (7b7680c, 84 rows, all legs incl. chapel, N=4→4096). The box differential — the FIRST hardware run of the S22 minimal-emission emitters — caught a real S22 bug pre-sweep (in-twin Fold scalar results Inline-classed → value dropped; fixed + pinned acdb319, and emit_sweep.rs now runs the 320-draw emission sweep locally without nvcc). Differential 15/15 green at 16-core pinning (48-vCPU fan-out starves CUDA context init — standing gotcha). Headline number: the GEMM kernel alone is 1.60× from naive-CUDA at N=512 f32 (0.125 vs 0.078 ms). Workspace 853 green (200 syntax · 153 ir · 155 lower · 29 check · 62 interp · 63 rewrite · 27 llvm + 1 ignored perf · 1 mapal-rt · 163 cuda); fmt clean; tree committed; box destroyed (≈$0.42).** Previous: S22 (closed — sessions/2026-07-22-s22-minimal-emission.md). Milestones: M1 ✅ · M2 ✅ · M3 ✅.

Components

Component Status Tests One-line state Docs
syntax tested 201 ✅ P1 complete: lexer + parser (ADR-0005/0009/0010/0011/0012); golden trees for all 9 examples (+zip_demo/vector_add S09, +seq_demo S11), zero diags; P-code out-of-Core rejection; ADR-0019 seq statement block (SeqBlock node; P0117 for dropped fanout-block statements); bench recorded. S20: P0108 carve (ADR-0029 stage 2a). S22: ADR-0031 — the carve is REMOVED; P0108 rejects every call expression uniformly (iota/fill get the arrow-form teaching message); 200 ✅. status
ir tested 189 ✅ S39: guard_plan — per Phi-position guard, the condition and each arm's exclusive work; an arm that is not taken does not run. Ownership is consumer closure (not liveness — a dead consumer still reads its operand), re-closed after nested-site subtraction, with can_trap/heavy transitive through nested sites. GuardSite::gated() = legality (can-trap MUST gate) then cost (heavy = bulk op or call). S36: a clock read is a DAG node — path_plan makes TimeMs its own pinned task with edges BOTH ways off source position (max_loc < ⇒ dep, min_loc > ⇒ dependent), so the bracketed work can no longer start before the read that opens the bracket; fir 65 536 par went 6/100 sub-0.01 ms readings → 0/100, values byte-identical. Core graph IR per ADR-0013/LC-4: edge-only dataflow, sealed builder + independent validate, inline-trace loops, IO token, SCC/topo, linted Mermaid; bench recorded. S05: empty-struct hole fixed. S07: Print{newline}/println (ADR-0015, +1). S09: §5.1 typing-table golden oracle (+1); Zip/Enumerate pure collection ops (ADR-0018). S20: last_use_plan (death/escapes/carried_by/dead_after) + bounds_proof (interval lattice; S20c capture-range recursion) deduced queries; Iota/Fill Core ops (ADR-0029 stage 1). S21: Operation::Widen (33-variant set) — four-edge lattice validate-enforced, InvalidWiden twins, family-Dest builder. S22: emission_plan deduced query (minimal-emission WP-A — Dissolved/Inline/Named; Pair-built-only dissolution after the review-caught R-NODUP count-drop) + FnBuilder::ty_of pub (ADR-0031). S24: path_plan deduced query — the execution graph's task DAG (paths/deps/ranks/transitive-trap sites/threshold checkpoints; token+effectful-loops stay host), the backend-independent parallel-orchestrator source. S25: tile_plan deduced query — bit-exact-interleavable map{fold} sites (affine-triple reads, 2-D + 1-D lane modes, proof-gated legality; #9 verified shipped-at-S20c + pinned). S28: TileRead.ksplit? : TileRead → TileKSplit — the fold-body (k÷div, k%div) derived-var split recorded (the map-body move one level down; XOR raw-k, depth % div == 0; unused pair ⇒ None, matmul/fir records bit-identical); conv2d_16 site recognized (was refused); 180 ✅. status
backend-nvptx planned No crate; grounded by a ratified plan, not by code. The GPU tile kernel is the SAME Trn at a second Loc — no mapal-ir edit, no new Operation. Step 1 of 8 built (the Machine class + cuda-ada profile, in backends/llvm); the two §2.2 gates built. Steps 2–4 need NO hardware (llc -march=nvptx64 is local); steps 6–8 use the owned RTX 4070 Ti. __syncthreads is modelled as the morphism that completes the Gmem→Smem transmission, so a missing barrier is Coherence Law 1 failing rather than a forgotten idiom. Packaging (own crate vs Machine discriminator) deliberately UNDECIDED — §2 is a default, §8.5 judges it on built code. status
lower tested 161 ✅ P2 second half complete: full Core surface lowers to sealed validate-empty IR; all 9 examples golden (+countdown/effectful-call, +seq_demo S11); 51 L-codes with rejection matrix; literal-width unification; token laws; zip/enumerate collection builtins (ADR-0018); ADR-0019 seq statement block (SeqBlock emit — no IR footprint, ordering is the token thread; L1611 no-tail); bench recorded. S20: iota/fill lowering — L1612/L1613, ExprKind::Call WTy::Array synthesis, static_count_arg (+5 rejections). S21: widen_i64/widen_f32/widen_f64 builtin family (is_pure_builtin + widen_target, L1614 teaching diagnostic, L1009 reserved). S22: ADR-0031 — n -> iota / (x, n) -> fill join the stage family over builder iota/fill_from; L1612/L1613 reworded (oversize → width-owned L1202); bound-literal counts legal (pinned); 155 ✅. status
check tested 30 ✅ P3 complete (S10): check(source, program, ir) — T0101 Return exclusivity (strict CK3) + T0201 E2 effect legality; ADR-0019 rebased the effects walk to key on node kind (Fanout opens the illegal-effect context, SeqBlock sticky), OQ-C1 closed (effectful seq in a branch is a composite rejected by E2; CK5 pin→theorem); typing discharged at validate boundary; E3 vacuous-by-proof (no heap ops), reopen trigger pinned; 10 examples clean incl. calc + seq_demo. status
interp tested 63 ✅ M1 oracle live (S08): fueled evaluator; guard-first loop driver (ADR-0016); nine example goldens (+seq_demo S11) + countdown + 55/fir value contracts + zip/enumerate contracts (ADR-0018) + seq-wrapped reassign contract (ADR-0019) + traps + fueled divergence + determinism; bench recorded. S20: Iota/Fill oracle arms + tests/iota_fill.rs (3 contracts + 2 e2e). S21: Widen arm — Rust as (round-ties-even ≡ sitofp/fpext/C casts), boundary-value contract; 62 ✅. status
rewrite tested 70 ✅ S34: the P0 fixedmap(id) → id judged a body by its Return writer alone and deleted map bodies whose dead Div must trap; is_identity_body now quantifies over the whole body via the crate-shared is_pure. One bug, four property entry points; seed retained, gate green. P4 complete (S12): shared plan+replay rewriter, fixpoint driver, oracle-exact const fold, CSE, conservative DCE, Map fusion, R1 properties, goldens, bench. S27: Inline default-first, loop-bearing-callee guard, cap 256, body Call stripping. S27b: LiftLoops immediately after Inline; R-LF/R-LM consume loop_plan, require K>=1, and synthesize captured Fold/Map bodies. Matmul4 now rewrites to zero Calls/loops with Map-with-Fold and reaches tiled LLVM; focused rejections, six-pass R1 battery, generated lift shapes, and 1,280 differential are green. status
backend-llvm tested 78 ✅ S36: the clock read moved into a pinned task and this backend needed no new code — the host spine emits the existing pin{id}_entries wait → checkrun_pinned trio at the read's topo position; total wall time unchanged, the self-timed interval rises ~12% because work that used to precede t0 is now inside the bracket. P5/M2 (S13): full textual-LLVM emitter emit(&CategoryIr) -> Result<String, EmitError> (ADR-0020) — ty/module/func/loops (ADR-0016 guard-first CFG via the shared mapal_ir::loop_plan predicate); mapal-rt runtime seam (7 print externs + mapal_trap exit-101, render-parity vs interp). 7 golden_ll (incl. two-loops + exit-only-payload-once pin) + 9 differential (320-case closed testgen sweep raw+rewritten, u8 ABI — caught the zeroext-before-type invalid-LLVM bug, matmul loop-driven Update, traps exit-101, by-ref call-args row S20); nested-loop = Unsupported. S14: -O2 differential row landed — every case (10 examples + 640 jobs × raw+rewritten) now at -O0 AND -O2; first run 1280 compile-and-runs, zero divergences; open item closed. S20 (marathon): by-ref array captures (#6 — ptr components in body-input products; matmul64_cap 2.08 s → 0.01 s), trap-aware fn attributes readonly/nounwind/willreturn + noalias nocapture readonly on clean by-ref params (#7), by-ref call args (#8/BL5 — loop-form matmul64 0.33 s → 0.01 s), last-use Update memcpy elision (#2 — llvm.memcpy 2→1 in the matmul4-class pin), proven-Index guard elision + #13 Div/Mod credit (S20b/c — matmul64_cap whole-file icmp count 3); 1280-run differential green at both opt levels under all of it; emit_iota/emit_fill (ADR-0029). S21: Widen = sext/sitofp/fpext; WP3b — first-class aggregate array moves eliminated (pointer-only staging fields + llvm.memcpy for owed copies; array Phi = select over ptrs): matmul256_cap clang -O2 OOM-kill → 0.08 s/57 MB, matmul512 same — the llvm N≥128 bench legs are live; capture_array_staging_never_loads_whole pin; 1,280-run differential green post-change. S24: parallel orchestrator v1 (Sapir-ratified plan)mapal_main consumes mapal_ir::path_plan: graph paths → @task{i} fns + one %Frame + a static rank/dep table on mapal-rt's new work-stealing pool (mapal_par_*; MAPAL_PAR env; GRAIN 4096); bulk sites slice across cores; speculate-and-order trap protocol (record + dummy-zero + host-fired exit-101 at threshold checkpoints) keeps output byte-equal to the oracle at any thread count; single-path fns and all non-entry fns byte-identical sequential text; two review finds fixed pre-merge (pinned pool race; effectful-loop seed hoist); R-PAR pinned live (trap prefix order, env matrix, run-twice, big-N) at -O0+-O2. S25: tile emission v1tile_plan-recognized sites emit the TILE_J=16 register micro-kernel (per-cell chain order exact ⇒ stdout byte-equal; tile-vs-untile + oracle differential -O0/-O2); EmitOpts{tiling,perf_timing} + MAPAL_PERF compute timer; measured 2.5–4.6× 1t tile-vs-untile local, NEON confirmed. S27: FMA contraction product face (EmitOpts::contract default-off; tile-kernel contract flags — conformance bit-exact untouched) + rung 3 packing (j-tile-major 64-aligned b panels; wrapper-task pack → nested slice dispatch; Seq/loop sites pack per iteration — loop-carried-b hole caught in review, regression pinned) + per-width TJ (f64→8) + k-unroll ×2 + prefetch; local @1024 1t fma 19.3 ms f32 (1.81× vs S26) / 38.6 f64 (1.73×). S27b: panel-residence jt-outer nest (per-thread b-traffic ÷4–16, zero spill, byte-exact) + the lifted-matmul4 tiled acceptance; 47 ✅ (23 differential + 24 golden). S28: the shapes ladder — FIR 1-D window rung (window1d_siteemit_tiled_map_blocked_1d, the rung-2 dual: TI×TJ lane blocks, ONE scalar w[k] shared per k, constant-TJ main, ×2 k-unroll; zero mapal-ir change) + conv2d k-split micro-kernel (conv_siteemit_tiled_map_conv: 9 taps at compile-time offsets, zero div/mod — cashes A1's ksplit record; non-conv ksplit keeps the untiled fallback; packing_site excludes ksplit). fir both tables WON local (fma-par 0.213 vs cpp-mt 0.239; fma-1t 0.216 vs cpp-1t 0.924 = 4.3×); conv kernel ≈0.04 = 3× over cpp-mt 0.133 (legs gen-dominated — the gen measurement boundary, recorded); matmul .ll byte-identical; 53 ✅ (28 differential + 25 golden). S31: the six hardcoded machine constants become one named TargetProfile table plus arithmetic (src/profile.rs; generic default / apple-m / zen3, EmitOpts::target, --target=) — tile_j_for/TILE_I/TILE_KC/tile_nc_for/HEAP_MIN_BYTES stop existing as literals, generic reproduces every one, and tile_i = vec_regs/(2×acc_vecs_per_row) reproduces S26's swept 4 including its recorded TI=8 spill. The KC gate closes by derivation (16 MB L2 ⟹ kc=4096 ≥ K), pinned in the strong form: apple-m WITH --kc is byte-equal to generic WITHOUT it — a threshold, not an off-switch (it reopens past K=4096, where the derivation disagrees with S30's measurement). TILE_I was two quantities — the FIR window rung's lane block is now WINDOW_SUBROWS, unbounded by any register budget. Value-invariance proven under zen3 (TJ 16→32, TI 4→2, oracle-equal at -O0/-O2); rule 1 checked as 66 byte-identical A/B emissions vs HEAD, because the 72 checked-in bench .ll are stale at HEAD (pre-existing, separate debt). status
runtime tested 21 ✅ NEW COMPONENT (S50). mapal-rt — the CPU backend's execution site: threads, queues, and the moment work crosses between them. Modelled for the first time; the 2026-08-01 review recorded that neither its Locs nor its Trms were in the architecture map. S50 rebuilt the handoff: per-lane wake (a dispatch wakes only the lanes it filled — empty dispatch 5.1 → 0.46 µs at 14 lanes); the run graph is sealed once and the completion path is lock-free (execute takes no lock; every slice but the last is one atomic; no per-completion Vec clone); the host spins before parking while workers do not (spinning workers REFUTED, 0.80–0.93× on real shapes); Placement collapsed to one variant (a dep-unlocked task is SEEDED, not put on its producer's lane — worth 1.49–1.58×); a one-slice task runs on the calling thread rather than crossing a boundary for nothing; one executor per lane, host included, so MAPAL_PAR now means TOTAL executors. Composition rule earned the hard way: retire unlocks dependents BEFORE publishing Done — the reverse let a waiter beat the unlock and trip mapal_par_run_pinned's assertion, survivable only while the host always parked. Two instruments added: examples/dispatch_cost.rs (a dispatch with no kernel under it) and examples/handoff_floor.rs (the wave round trip alone). status
backend-cuda tested 161 ✅ P6/M3 (S15): full textual-CUDA emitter emit(&CategoryIr) -> Result<String, EmitError> per the S14 review-hardened DESIGN — host/device split (zero whole-array D→H), sequential single-thread Fold (oracle order), Update full-copy kernel, cudaMalloc'd trap flag (kind+1) + check after every launch, exit-102 infra protocol, 64-bit indexing, -fmad=false, BC8 qualifiers (callerward Twin propagation), guard-first loop quartets (host + per-thread inline). M3 green on an RTX 4090: 10 examples + 320 testgen, raw+rewritten — 640 compile-and-runs, 0 divergences; traps 101; fmod parity; matmul Update. S15 review: 2 blockers + 4 majors fixed pre-GPU (trap-encoding collision; return-escape guard; Twin propagation; float Neg; +2 recorded cells). S20 (marathon): emitter-quality wave (#17 kernel shape dedup, #12 dead-host-twin elimination, #13 constant-divisor guard elision, #14 trap-param trimming) + smart arenas v1.0 (#18: per-fn zone with capacity deduced from the graph, abi_sizeof, 256 B offsets, 4 GiB EmitError guard; capture matmul 8→1 malloc; zone-release escape range-veto) + #19a emit_with_opts/--perf CUDA-event MAPAL_PERF timing (default byte-identical) + in-place Update + back-edge freeing (#2 — borrowed-init veto keeps matmul's full copy, .cu byte-identical) + bounds-proof guard elision & TrapCaps proofs (S20a/b/c — matmul64_cap map kernel: no trap param, 0 bounds guards); remote differential 151 green / 640 runs on the 4090 under all of it. S21 (ADR-0029 stage 2): Iota/Fill kernel realization — bulk-site family, arena members, NO trap plumbing, count as a launch arg (cross-count #17 dedup: iota(4)/iota(5) → one kernel), device-twin local loops; Widen C-cast arms (scalar, no kernel site); the 5th Unsupported cell discharged (4 remain). S23: WP-D invariant hoisting (assemble_body_arg split; d_fn4's loop body = 2 assigns + call) + the Fold force-Named fix (the box-differential catch) + emit_sweep.rs local emission sweep + hardware verification: remote differential 15/15 green over the S22+WP-D emitters; the S23 matrix measured (kernel-alone 1.60× from naive-CUDA at 512 f32); 163 ✅. status
backend-verilog not-started Feedforward + single-loop FSM Verilog (E1); not begun. status
cli not-started mapal build|run|dump-ir|test; not begun. status

Status vocabulary: not-started · design · building · tested · stable · blocked

Backend capability matrix

Feature interp llvm cuda verilog
pipelines / operator-shorthand planned
functions planned
guards → Phi (gated arms, S39) ✅¹ planned
loops / trace (guard-first, ADR-0016) planned
parallel fanout (pure) planned
seq + print (IO) planned
tuples / named types / fixed arrays planned
map / fold inline-block planned
zip / enumerate planned
iota / fill (ADR-0029, stage 2) planned
widen_i64/f32/f64 (ADR-0029 amendment) planned
array update (c[i] <- x) planned
time clock read (plan-time-builtin) planned

Legend: ✅ supported · ✋ rejected-with-error · planned

¹ CUDA: guards are gated. Map/fold bodies are emitted as __host__ __device__ fns by func.rs, so a guard inside a map emits a real if/else on the device — verified by emitting one. kernel.rs still holds a strict-select arm, but it is measured unreachable: probed with a panic across 106 emissions (every example, bench shape and matmul, raw and --rewrite) and the 163-test CUDA suite, it fired zero times. The S39 CUDA change has no hardware run behind it (emitted C++ was syntax-checked only), and at S39 close Sapir accepted that rather than owing it: "we are going to transition from cuda to nvptx anyways".

Standing Verilog restriction (HANDOFF §4.3): the Verilog backend supports only feedforward pipelines + single-loop FSMs (with the E1 done protocol). Everything else is rejected-with-error when implemented.

Blockers

The workspace gate is GREEN as of S44–S47 (2026-07-31): 1047 passed, 0 failed — re-verified at S48 start (cargo test --workspace --release, same tally).

RESOLVED (S40) — S39's defect: gating was not stable across LiftLoops. An arm that is not taken must not run (S39), but guard_plan v1 refused to gate a site whose arm work touched a loop SCC while LiftLoops made the SCC disappear — strict before the pass, gated after, eval ∘ rewrite = eval violated. Fixed in aaaa5dd ("guards gate the flow; arms own loops as units"): guard_plan gained LoopUnit (crates/mapal-ir/src/algo.rs) — an arm owns a whole loop atomically through its LoopEnter handle(s) and the driver fires the internals, with the invariant refused raw ⇒ refused rewritten, no stability hole (a non-canonical unit never joins an arm, and LiftLoops consumes the same loop_plan facts). The three proptest seeds in crates/mapal-rewrite/tests/property.proptest-regressions stay pinned and pass.

Previously (S39, superseded): the gate was RED (2026-07-28): 992 passed, 2 failed. Before that, GREEN as of S38 (2026-07-27): 981 passed, 0 failed. cargo test --workspace --release passes with both pinned proptest seeds retained, the LLVM differential ran (37/37, 403.93 s) rather than skipping, and cargo fmt --check is clean.

RESOLVED (S38) — Inline could change which trap fires. topo_order broke scheduling ties on object insertion order, a property of the compiler rather than of the program; the dataflow graph imposes no order between two independent trapping operations, and rewriting reshuffles insertion order, so Inline turned Trapped(IndexOob) into Trapped(DivZero) — a violation of eval ∘ rewrite = eval, the rewriter's one rule. Fixed by making the tie-break (loc.start, loc.end, insertion index): source position is the only key intrinsic to the program. Approach A (fix the one shared order) rather than B (fix the interpreter's selection key), because the oracle, sequential emission and the parallel runtime's CAS-min all derive from topo_order today and B would create a second key to keep in sync across every backend forever. The class was wider than the counterexample: the same root cause let a trap swallow output written before it in source order, which the 1,280-run differential cannot see because it discards stdout on trapping runs (expect_native(None, 101)). Both are now pinned — mapal-rewrite/tests/inline.rs::open_inline (seed retained) and backends/llvm/tests/differential.rs::differential_trap_preserves_preceding_output (verified negative control). Full account: components/ir/plans/plan-s38-trap-order-is-source-order.md.

Known consequence, accepted, not a blocker: emission order moved everywhere, so three of seven ladder cells shifted 2–5% (saxpy 1t +5.3% reproduced three times; conv2d 1t faster and par slower; mm1024 +2.6% on the conformance face only). Values are byte-identical throughout and vector instruction counts are unchanged, so this is scheduling rather than codegen. Mechanism is not isolated and is deliberately not being chased.

RESOLVED — the rewriter deleted a trap that must fire. Found by CI on its first run, pre-existing at 1daddaa. Root cause: map(id) → id, not fusion. is_identity_body (crates/mapal-rewrite/src/functor_laws.rs) asked only what a map body returns, so a body returning its parameter while also computing a dead trapping Div read as the identity and the whole Map was aliased away. The oracle evaluates a body's entire graph — that is exactly why DCE pins impure dead cones live (R4) — so such a body denotes id ∘ trap : A ⇀ A, a partial morphism, and the functor law List(id) = id does not apply to it. The guard now quantifies over the body's whole morphism set through graph_rewrites::is_pure, one predicate shared with DCE. All four failing entry points were this single bug. Two new pins (property.rs): identity_map_body_with_dead_trap_stays_trapped and its positive control pure_identity_map_is_still_eliminated, both negative-controlled. Known cost, recorded not papered over: a body containing Widen/Iota/Fill no longer forwards. Full account: components/rewrite/plans/plan-s34-identity-map-trap.md.

cargo test -q -p mapal-rewrite --release --test property   # 11 green, ~0.3 s

Remaining P0, not gate-visible: mapal_par_wait lets workers run ahead of the clock, so a kernel can finish before the clock meant to bracket it is read (3–4% of threaded runs; one live case read 0.0001 ms). Measurement-only, but in the direction that flatters us. A runtime-only fix was built and reverted — components/backend-llvm/plans/plan-s33b-clock-read-barrier.md §4 says why not to retry it.

Errata/ADR ledger

ID Title Status Applied to spec?
E1 Mapal-Cat cannot be both total and traced-cartesian (loops are partial / guarded trace + done protocol) accepted (ADR-0002) yes
E2 Parallel effects rule — no effects in parallel fanout; seq or KPN channels accepted (ADR-0003) yes
E3 Memory-model guarantee scoped to first-order non-cyclic core accepted (ADR-0004) yes
E4 Operator-precedence example fixed; a flow is a statement, not a value accepted (ADR-0005) yes
E5 Rename surface keyword categorytype accepted (ADR-0006) — veto window closed 2026-06-11, no veto; rename final yes
ADR-0001 Mapal-Core scope accepted n/a
ADR-0007 Tech stack accepted n/a
ADR-0008 Editor tooling & LSP plan accepted n/a
ADR-0009 Collection-operator syntax — postfix inline block; input tuple ↔ block params positionally ((init, array) -> fold { acc, item -> ... }) accepted yes (LC-2)
ADR-0010 Guard arrows are single lexemes — adjacency + statement-initial context gate; -7->x; is a guard arm, write -7 -> x; accepted — flagged to Sapir (next-session.md), revisable by superseding ADR n/a (spec silent; no text patched)
ADR-0011 Mapal-Core loop labels are loop only; statement-initial Ident { disambiguation accepted — amended by ADR-0012 (scan demoted to hint; Ident { always a struct literal) n/a
ADR-0012 Labeled blocks :label { … } / jumps -> :label; (prefix sigil both ends); enclosing-targets-only (E1/Verilog reducibility); loop keyword unchanged; labels stay Core+1 (P0110) accepted — decided with Sapir, Session 03 yes (LC-3: user-guide §3.5/§8.5 patched)
ADR-0013 IR realization: all dataflow is edges (per-slot Pair, constants-as-objects); Core op set (+Neg/Index/Map/Fold/Print/loop edges/Output, −Identity/Const/Trace); loops as inline SCC-visible cycles; IO as linear world token (signature synthesis, token sink, token-in⇒token-out) accepted — ratified by Sapir S13 (+ S13 amendment: div-zero trap is integer-only, float ÷0 is IEEE — closes IN6) yes (LC-4: category-ir §4.1/§5.3 marked, §3.3 pointer note)
ADR-0014 FRAMEWORK.md adopted as the Level-B categorical model layer for compiler-internal design (firewalled from Mapal-Cat / Level A); mandatory ## Categorical model (Dat + Trn) DESIGN lead-section + FRAMEWORK §8 reconcile-gate line; new docs/architecture/{categorical-model,INDEX}.md accepted — ratified by Sapir 2026-06-14 (Session 06) n/a (methodology; Level A untouched — HANDOFF §7.1.5/§7.2 patched, not spec)
ADR-0015 Split the print effect: print raw, println appends \n; one IR op Print { newline }; Core effect surface {print}{print, println} accepted — decided with Sapir (Session 07) n/a (scope/methodology; HANDOFF §4.1 patched; Level-A spec untouched)
ADR-0016 Loop branch evaluation is guard-first — the continue-branch (inr(U) arm) is NOT speculatively evaluated on the exit step (E1 refinement). Found while implementing the interp oracle: eager-both miscompiled fir (coeffs[4] OOB-trap at the exit state k=4 instead of 5.375). Fix lives in the oracle + this ADR; all backends inherit it (differential-tested). accepted — ratified by Sapir S13 n/a (operational refinement of E1; category-ir.md untouched — frozen Level A already says Elgot; no ERRATA entry)
ADR-0017 Category-architect docs tree adopted (extends ADR-0014): per-component IMPLEMENTATION.md (model → file:symbol functor) + suggestions.md + plans/reviews/general; top-level architecture-map.md/IMPLEMENTATION.md/suggestions.md; immutable docs/sessions/ logs; DESIGN.md doubles as the ARCHITECTURE doc accepted — user-directed (Session 09) n/a (methodology; HANDOFF §6/§7.2 patched)
ADR-0018 zip/enumerate join the realized Core op set as pure collection primitives (ADR-0013 delta +2): Zip : ([A;n],[B;n]) → [(A,B);n], Enumerate : [A;n] → [(i32,A);n] (index pinned i32, n ≤ i32::MAX); builtins resolved by name in lower like print (name-collision → L1009); effect-free (no token, fanout-legal). iota dropped (deduced map π₀ ∘ enumerate). accepted — decided with Sapir (Session 09) n/a (realized-set delta; HANDOFF §4.1 patched; Level-A spec untouched, as ADR-0013/0015/0016)
ADR-0019 seq { … } is a statement block (StageKind::SeqBlock(Block)), not a fanout kind — FanoutKind shrinks to Plain | Void, the Seq summand migrates to the new node. No IR footprint (ordering = ADR-0013 token thread, pin d; interp/ir untouched); the silent non-Chain fanout-block drop becomes diagnostic P0117; effectful-seq-in-branch is a composite rejected by E2, so CK5 upgrades pin→theorem and OQ-C1 is closed by construction. accepted — proposed by orchestrator, ratified by Sapir (Session 11) yes (LC-5: user-guide §5.2 prose + canonical example patched)
ADR-0020 Backend emission contract: one convention emit(&CategoryIr) -> Result<String, EmitError> (the String IS the TargetText; per-crate structured EmitError with Unsupported = the ✋ cell); one shared runtime crate mapal-rt (print via Rust Display = oracle render parity by construction; mapal_trap → exit 101); semantics parity = rewrite R1 (wrapping ints — no nsw; guarded div/mod/index traps); differential duty per backend (examples + testgen, raw + rewritten IR); toolchain absence ⇒ skip-with-reason; deterministic emission. accepted — ratified by Sapir S13 n/a (realization layer; Level-A spec untouched — same class as ADR-0013/0015/0016)
ADR-0021 Array element update: pure Update : (Array{T,n} × I × T) → Array{T,n} (fresh array, OOB ⇒ IndexOob trap, same class as Index; no token, fanout-legal) + surface rebind sugar c[i] <- x; (extends bind-stmt; desugars onto the existing mut-rebind path — Phi-arm/loop rules inherited); rewrite const-index laws; naive-copy semantics everywhere, in-place via last-use = recorded headroom; dynamic sizes stay out (E3 scope) accepted — decided with Sapir S13 (note Option A ratified + "I need the dynamic array access") n/a (realized-set delta, ADR-0018 class; Level-A spec untouched)
ADR-0022 Truth-in-docs + Level-B maintenance freeze: docs/spec/mapal-as-implemented.md designated the operative index of the language; "frozen" retired for an explicit authority order (spec v0.2 + E1–E5 + LC-1–5 + ADRs, newest wins, oracle behavior final arbiter); Level-B upkeep frozen except backend-seam invocations — amended at ratification: the exception is standing practice for new backends, paid in the backend's DESIGN; the re-evaluation trigger is scoped to defects outside backend seams (its original letter fired at the S14 keep-test and was judged to support the exception, not general upkeep); VISION related-work correction ratified as factual maintenance accepted — ratified by Sapir 2026-07-18 (S14), D2 amended at ratification n/a (governance; HANDOFF amended in 7 marked spots S14; ERRATA.md preamble patched at ratification)
ADR-0023 Dynamic-sized arrays: fixed-but-unknown-size heap arrays (length in the value, immutable), explicit promotion from literals, LenMismatch trap for zip, E3 restated with content, per-component impact (backend-llvm the only large cell). Sapir's allocator amendment: every backend provides a "size allocator" store component (ASIC/FPGA: fixed-capacity, possibly Mapal-authored); allocation is fallible and must be handled by calling code — checked errors as values, riding the coproducts ADR accepted — decided by Sapir 2026-07-18 (S14), post-M5; Q1/Q2/Q4 open; depends on the coproducts ADR n/a (patch list in-file §Spec impact applies on landing)
ADR-0024 Templates: T1 type-only monomorphization at a mapal-lower Pass 0 (zero IR change — Ty is monomorphic); T2 size-parametric [A; N] with call-site N-unification (joint design with ADR-0023); honest C++-lesson costs candidate — proposed 2026-07-18 · NOT decided · number provisional n/a
ADR-0025 TT backend (F_TT): post-M5 third target, the O5 proof-case; TT-Metalium host+kernel C++ emission over TTIR (stability); fold in-order host-staged (oracle parity); ttsim CI + on-metal validation on Blackhole (Sapir's chosen generation) accepted — decided by Sapir 2026-07-18 (S14): the P8 target, post-M5; Q3 (O5-demo framing) open n/a (realization layer; crates/backends/tt (package flow-backend-tt) + docs tree arrive at P8)
ADR-0027 Capture semantics: map/fold/enumerate bodies may read enclosing bindings (pure read captures) — realized as broadcast edges / hidden body-fn params (no new IR op); legality is a graph property (no path from fanout output back into a captured binding; no captured rebind targets; violations reported as the path); L1108 narrows to mutation/effect cases with a teaching diagnostic; E2 unchanged. The S16 GEMM-expressibility unblock accepted — ratified by Sapir 2026-07-21 (S17) as proposed (Q1–Q5 resolved in-file) pending (implementation S18; user-guide §5 + mapal-as-implemented + L1108 narrowed on landing)
ADR-0028 Tree-reduction of exact-op folds (wrapping Add/Mul, Min/Max, And/Or/Xor, int/Bool); f32/f64 folds stay sequential-pinned accepted — decided under the S20 delegated mandate n/a (realization layer; implementation queued — the S21+ tree-fold wave)
ADR-0029 Array-construction builtins iota(n)/fill(x,n) + explicit widen (amendment: Operation::Widen, four-edge lattice, widen_i64/widen_f32/widen_f64 builtin family) accepted — stage 1 shipped S20; stage 2 + amendment shipped S21 (cuda kernels, widen full-pipeline, procedural v2 benches) n/a (realized-set delta, ADR-0018 class; mapal-as-implemented patch on ledger close)
ADR-0030 External-backend protocol + SDK: serialized versioned CategoryIr bundle (+ deduced queries) over a subprocess contract; conformance = the oracle differential; in-tree folder move crates/backends/{b} (names unchanged) candidate — direction requested by Sapir S21; folder move EXECUTED S22 (2026-07-22); protocol NOT scheduled n/a
ADR-0031 iota/fill surface → pipeline form (n -> iota, (x, n) -> fill); the ADR-0029 stage-2a call-expression carve removed, P0108 uniform again; lower rides fill_from/iota builder entries; zero IR change (interp/rewrite/backends untouched) accepted — Sapir directive S22 (in-session) pending (mapal-as-implemented iota/fill surface row on ledger close)

Session log (newest first)

NN date focus outcome
28 2026-07-24 Shapes ladder (fir/conv2d) + S27 box run Sapir's S28 focus shipped: FIR 1-D window rung + conv2d k-split micro-kernel. A1: TileRead.ksplit records the fold's (k÷div,k%div) derived axes (the map-body move one level down; XOR-raw-k + rectangular-window rules; conv2d_16 recognized). B: emit_tiled_map_blocked_1d (rung-2 dual; zero mapal-ir change) — fir both tables WON local + box. A3: conv_siteemit_tiled_map_conv (9 unrolled taps, zero div/mod; non-conv ksplit = untiled fallback) — conv kernel 3× over cpp-mt; box par table won; M4 par leg OPEN on the gen measurement boundary (suggestion #14). Box: the "balance 0" misread corrected (credit+autobill); box #1 vanished mid-run → on-demand relaunch (45712913, destroyed ≈$0.45): disasm gates pass (S26 vfmadd finding closed), 2048/4096 rows clean, OpenBLAS frontier measured (threaded 9.7×/5.9×; 1t 2.7×), GRAIN quantization measured (#15). Gate 69/69 green; matmul .ll byte-identical. Log: sessions/2026-07-24-s28-shapes-ladder.md.
27c 2026-07-24 Local measurement campaign + fairness fix The S27 matchup exists — local, same-machine (M4 Pro), Sapir-directed: matmul/s27.md + results-s27-local.csv (runner gained macOS probes + MAX_N clamp). Fair compute-basis verdicts (Sapir's wall-vs-selftimed catch, now a standing rule): flow-fma ahead at EVERY size — par 2.8×@256 → 19.6×@1024 f32 over cpp-mt; 1t 42.7×@1024 over cpp. numpy = Accelerate-AMX (labeled, different silicon). Shapes cross-language baselines NEW (benches/shapes/shapes_{baseline.cpp,baseline.rs,numpy.py,ab.sh}, outputs byte-verified): fir mid-pack (no 1-D rung), conv2d 7.9× behind 1t-cpp — the derived-var demand gate FIRED → the S28 focus (Sapir: lift fir/conv2d to winning; tile_plan must record non-affine). Mapal 2048 local = macOS stack wall (heap lowering = enabler). Log: sessions/2026-07-24-s27c-local-measurement-campaign.md.
27b 2026-07-24 Loop→map/fold lifting + panel residence Sapir's unlock live: naive loop matmul lifts → inlines → tiles automatically (LiftLoops after Inline; R-LF/R-LM off loop_plan, K≥1 adjudicated, rejections pinned; matmul4 -275/3748 exact at O0/O2 × {default, MAPAL_PAR=1} — S26 non-tiling pin INVERTED; fir fold-lifts; testgen lift shapes feed the 1280-run sweep). Panel residence: jt-outer nest, per-thread b-traffic ÷4–16, zero spill (k-panel deferral corrected per Sapir; KC-split box-gated). Loop-form bench legs regenerated tiled. Codex ×2 (one correct STOP on the Iota(0) hole); orchestrator line-by-line. 72 suites green. Log: sessions/2026-07-24-s27b-loop-lift-panel-residence.md.
27 2026-07-24 FMA + rung 3 packing + fn-strip + loop→map design All three numpy-gap closers shipped + measured locally; fn-strip wired; loop→map designed. FMA: per-instruction contract flags (hypothesis verified — driver flag never retrofits textual IR), product face only, conformance bit-exact. Packing: j-tile-major 64-aligned b panels (wrapper-task pack → nested slice dispatch; Seq/loop sites pack per iteration — review caught codex's uninitialized-buffer hole, loop-carried regression pinned). Finishing: f64 TJ=8, k-unroll ×2, prefetch. Local @1024 1t: f32 32.8→fma 19.3 ms (1.81× vs S26), f64 64.0→38.6 (1.73×). fn-strip: Inline default+first, loop-guard, cap 256 — map-body helper Calls tile now. Harness to 4096 + fma legs; .cu byte-stable. Workspace green. Box BLOCKED: vast balance 0s27_box.sh prepped, S28 opener. Logs: sessions/2026-07-24-s27-fma-packing-fnstrip.md.
25 2026-07-23 Tile emission v1 (BLAS rung 1) + pool quota + timer Bit-exact SIMD shipped and measured: tile_plan (affine-triple recognition, 2-D + 1-D lane modes) + the backend-llvm TILE_J=16 micro-kernel — per-cell chain order exact, stdout byte-equal at any thread count/opt level; one review-caught R1 hole (skipped trap-capable Call/nested sites) fixed + pinned. Box (EPYC 7702P, 62-core quota — the S25 cgroup-aware pool live): flow 3–8.6× ahead of chapel-multicore at 512/1024 both widths; numpy gap 130× → 13.8× f64@1024. Shapes corpus: attn 2 chained tiled sites (4.6× local), fir large-K 2.2× (small-K clang self-SLPs — recorded boundary), conv2d/rowmajor refusals exact. Suggestion #9 closed as shipped-at-S20c; s24 reading 6 corrected. MAPAL_PERF compute timer ends wall-vs-floor estimates. 904 green; box ≈$0.55. Commits be4e827 + close-out.
24+24b 2026-07-23 Parallel orchestrator + fmad flip (see header + logs) path_plan task DAG → mapal-rt work-stealing pool → parallel mapal_main; N=1024 f32 flow ahead of chapel-multicore (184.0 vs 192.7 ms), 19.1× self-speedup; v3 speculate-and-order trap protocol; S24b: -fmad=true product default measured — f64 GEMM kernel at naive-CUDA-f64 parity (details: sessions/2026-07-23-s24-parallel-orchestrator.md, -s24b-fmad-flip-and-measure.md).
23 2026-07-22 WP-D + hardware verification + the S23 matrix The S22 mandate's done-bar met: the full measured matrix exists (docs/performance/matmul.md rewritten; 84 rows, one fresh 4090, all legs — flow-cuda loop/cap/kernel f64+f32, flow-llvm cap f64/f32, naive-CUDA, cuBLAS, numpy, rust, cpp, chapel; six-way output agreement at every shared N). WP-D loop-invariant hoisting shipped via codex (the delegation split restored — the S22 "network-dead" was a stdin wait; codex exec ... </dev/null is the standing fix); WP-E assessed and deferred with measurements (LLVM has no expression nesting; sroa already recovers the alloca shuffle — suggestions #10). The first hardware run of the S22 emitters caught a real bug (in-twin Fold scalar result Inline-classed → silently dropped; neg operand panic class, 27/640): fixed by mirroring the host force-Named rule, pinned, plus emit_sweep.rs — the deterministic emission sweep now runs locally without nvcc. Differential 15/15 at 16-core pinning (new gotcha: big-vCPU fan-out starves CUDA context init). Kernel-alone GEMM 1.60× from naive-CUDA (N=512 f32). Box #45539759 destroyed, ≈$0.42; workspace 853 green.
19 2026-07-21 Close-out: by-value diagnosis + perf docs + S19 agenda The flow-llvm capture slowness is diagnosed to the byte ({ [4096 x double], … } = 64 KB of by-value arrays per body call × 262k calls ≈ 17 GB of memcpy at N=64 = the 2,076 ms; the fix — by-reference captures — is backend-llvm suggestion #6). docs/performance/ created — tables-only benchmark home (Sapir directive), measurement-kind trap fixed (flow-cuda 277 ms beats the naive-C binary's 333 ms process wall at equal terms). S19 agenda recorded (Sapir-ordered): smart arena allocation (capacity computed from the execution graph, arenas by capacity not per-fn, system max guard, per-variable arena pointers — suggestions #18 elaborated), consolidation (kernel shape dedup first, #17), real kernel time (emitter timing blocks; a language-level time builtin deferred to Sapir, #19). Docs-only session; workspace 731 green; nothing live.
18 2026-07-21 ADR-0027 capture semantics + re-benchmark Captures shipped; the one-kernel GEMM is real; workspace 731 green (199+117+146+29+56+45+18+1+119). Sapir ratified ADR-0027; implementation by swarm (ir captures: u32 + map/fold_captured; lower free-var analysis + L1108 narrowed to mutation cases with a teaching diagnostic; interp read-at-position; rewrite identical-capture fusion; llvm/cuda capture args; testgen 5 shapes, 2048-draw stress zero failures). 4-lens review: 16 findings, all fixed — incl. the session's headline: a latent pre-existing loop_plan substrate bug (computed exit payloads never scheduled — the S15 unexplained interp panic; zero captures needed; fixed by widening the decide cone to route feeders' backward cone, merge-reachability bounded; verified t*2 -> ret → 6 and the exit-arm captured map on both engines). Also: poison-bit L1107 bypass, scope leaks, D3 rebind message, erased-capture builder rejection, stale-test conversion + oracle pins, replay arity, token→L1605. S18 re-benchmark (fresh 4090, ≈$0.30, destroyed): capture-form matmul N=64 277 ms (44×), N=128 278 ms (359×), N=256 278 ms — flat, startup-bound, faster than the naive-C (311–333 ms) and cuBLAS (469–496 ms) process walls on the same box; ONE __global__ with the inner fold as a per-thread loop; flow-llvm on the same source slower than its loop form (2,076 ms — by-value aggregate cost, recorded). Sapir's emitter-quality notes recorded (dedup #17, arena #18, +4 rows) with dedup/arena pulled forward pre-region.
17 2026-07-21 Design direction: parallel-first + region emission Two design artifacts, awaiting Sapir. Directives recorded: perf-per-step gate; parallel-first language; regions, not instructions — strip functions to the primitive graph. Written: components/backend-cuda/plans/plan-region-emission.md (v2 strategy: inline pass in mapal-rewrite → backend-independent region partition with per-backend cost models → one kernel per region; matmul64 acceptance 12.16 s → ~0.1 s → one kernel; traps re-grained per region; v1 stays as reference until v2's differential matches) and decisions/ADR-0027-capture-semantics-candidate.md (pure read captures as hidden body-fn params; L1108 narrows; Q1–Q5). Parallel-first package queued: canonical-tree reduce + par loops (additive, oracle-pinned). Honest reframe recorded: CT bought correctness-with-structure (the R1 license to optimize safely), not performance — the walls are physical, the fixes are mapping + expressibility. Design-only session; workspace 673 green; nothing live.
16 2026-07-21 Matmul benchmark (Mapal vs CUDA vs cuBLAS vs CPU) First measured perf numbers — the walls are named and priced. Mapal-CUDA matmul today: Θ(N³) kernel launches at ~24 µs/op (12.16 s @64, 99.8 s @128) — the DESIGN §3 correct-first price; same algorithm as one kernel = 3.8M× at N=64; same .mapal source on flow-llvm = 38× faster on a laptop CPU (backend strategy dominates). The real GEMM blocker is expressibility: L1108 (no captured map bodies) — capture/cartesian ADR queued for Sapir. Context: naive one-kernel CUDA 3.1 TF/s peak; cuBLAS 55.9 TF/s @4096 (22× algorithm gap); numpy 1.33 TF/s; flow-llvm hits the W3 N⁴ wall at N=64. Artifacts in-tree: examples/matmul4.mapal, benches/matmul/, docs/notes/bench-matmul.md, two emit cargo examples. Sapir's PTX question recorded (keep nvcc/C++; NVRTC = the interesting variant). Box ≈ $0.18, destroyed. Workspace 673 green.
15 2026-07-18..21 P6 backend-cuda (M3) P6 complete — M3 reached; workspace 673 green (199+106+139+29+44+27+13+1+115). Implementation per the S14 DESIGN by sequenced-TDD WPs (orchestrator + swarm; 6-agent parallel phase: loops ∥ differential harness ∥ 4 review lenses). 4-lens implementation review found 2 blockers the S14 design review missed: the trap-flag encoding collision (mapal-rt kind 0 = div_zero vs the memset-0 quiescent — a device div-zero read back as "no trap", an R1 class cross; fixed kind+1) and the return-escape class (lower's Dest::Ret bypasses Output; product/Phi returns freed escaping buffers — fixed via the epilogue pointer-value guard). Majors: Twin-ness not propagated through calls (illegal CUDA), float Neg of a negative constant (--1.5e0, ill-formed), products-with-arrays on device (→ recorded Unsupported cell), +2 minors→cells (erased-array graceful degradation; 16 KiB per-thread budget cell). Fixer TDD'd all 10 findings (115 green); orchestrator line-by-line review of the whole crate. GPU leg (fresh vast.ai 4090, ~$0.25, destroyed after): 10 examples + 320 testgen raw+rewritten — 640 nvcc compile-and-runs, zero divergences; traps 101; fmod/IEEE-÷0 parity; matmul Update; link tail verified. One flaky host aborted at 18 min loading (abort-and-record). Docs reconciled bottom-up; DESIGN §"As-built (S15)".
14 2026-07-18 Mitigation pass + ADR candidates + P6 DESIGN P6 DESIGN complete, review-hardened; implementation gate pending Sapir. Mitigation pass (7 workstreams) landed: getting-started seq fix + sweep; user-guide 18 empirically-verified badges + vector.mapal header; related-work + VISION O4 correction; mapal-as-implemented.md (operative index); ADR-0022 (pending ratification); backend-llvm -O2 row — 1280 compile-and-runs, 0 divergences; array-scale plan. ADR-0023/24/25 candidates recorded. backend-cuda DESIGN authored model-first (first real Loc/Trm pair); 4-lens adversarial review → 22/22 fixed (2 blockers: fn-walk driver-ownership skip, allocation-based buffer ownership); D2 keep-test = Level-B's first finding. Six non-backend crates green (orchestrator-verified); backend-llvm suite green incl. -O2. Category-architect adopted as the standing session contract.
13 2026-07-18 ADR-0021 array update + P5 LLVM backend (M2) P5 complete — M2 reached; workspace 558 green (199+106+139+29+44+27+13+1). Sapir ratified the S12 stack (ADR-0013+IN6 amendment, 0016, 0020, RW2, array-update note Option A) and asked for dynamic array access → ADR-0021: pure Update : (Array{T,n} × I × T) → Array{T,n} + c[i] <- x rebind sugar (P0013–15), landed pipeline-wide. Pre-code 38-agent design review: 15/15 findings CONFIRMED (srem MIN/-1 → 0 not MIN; u8 index zext — 3 lenses converged; rebind-inheritance false — 3 wiring points made explicit; rewrite L-b/L-c unimplementable → scoped to L-a + reoperand headroom; replay exhaustive-match break; open-mode testgen has no native observable; two-sequential-loops untested). Implementation: sequenced-TDD workflows (2 died on infra — internet outage, structured-output caps — orchestrator recovered both from cache/disk); reviewer catches: lower capture-check hole (L1108 evasion in map bodies), blind identity-seed matmul flagship test. P5: mapal-rt (render-parity + exit-101) + mapal_ir::loop_plan (BL7 — attribution predicate exported once, interp migrated) + full emitter (select-Phi, Div/Mod split guards, type-directed index guards, token erasure arity-0/1/≥2, Update memcpy, guard-first loop CFG) + differential harness (10 examples + 320-case testgen sweep, raw+rewritten, u8 ABI — caught the zeroext-before-type invalid-LLVM bug at first real clang contact, native loop-driven matmul 8\n136\n). Orchestrator line-by-line review of every diff: fixed exit-only-payload double-emission (walk now plan-owned; a duplicated exit-arm Print would have broken R1), overruled one fixer rejection (loop-carried-Update differential added), capped perf N honestly (array-literal module chokes clang -O2; ~80× native at N=4096).
12 2026-07-17/18 P4 mapal-rewrite + 2 interp P0s P4 complete — workspace 511 green (192+102+128+29+37+23). rewrite DESIGN model-first (plan+replay, R1 oracle-equality with ⊥-identified traps, P1–P3 plan laws) hardened by a 17-agent 4-lens adversarial review (8 confirmed findings — 3 lenses converged on Return-writer dropping; loop-state constify → LoopBackOutsideScc live-repro'd; fusion Diverged↔Trapped flip; dead pure loops lower-reachable — 5 refuted); implemented by a 10-agent sequenced-TDD workflow (WP3 died on API 529, WP5 self-healed the gap; fixer killed a token-forward TokenNotLinear miscompile); orchestrator line-by-line review then fixed per-merge exit attribution (two-loop fns rewritable). Two interp P0s found via Sapir's matmul exploration: multi-hop computed loop-invariants read-before-write-panicked the driver (fixed at root — ir topo_order LoopEnter deferral; invariants-before-header is now a theorem backends inherit) and two sequential loops per fn tripped the M1 assert (per-fn SCC-union attribution → per-merge; both pinned, matmul4 = permanent regression). P5 prep: ADR-0020 (emission contract + mapal-rt) + backend-llvm DESIGN written; verilator+icarus installed; vastai verified. Array-update design note (notes/array-update-design.md) for Sapir.
11 2026-07-17 ADR-0019 seq statement block seq is its own node; OQ-C1 closed; workspace 484 green (192+101+128+29+34). ADR-0019 (ratified by Sapir) + cross-component plan executed by a 17-agent workflow (3 Opus implementers TDD + 8 adversarial reviewers + 3 fixers), orchestrator line-by-line review. syntax: StageKind::SeqBlock(Block), FanoutKindPlain|Void, P0117 (direct-push — cooldown collapse was a review blocker; void-tail silent drop also fixed). lower: emit_seq_block (enclosing scope, HeadlessSeed, tail value, no IR footprint), L1611 + ChainCtx::RetValue; review caught 3 pre-existing miscompiles — Phi-arm scan, loop carried-set, capture check didn't descend Fanout (or the new SeqBlock): effectful-fanout-in-Phi-arm hoisted unconditionally, loop-mut-in-branch dropped, L1108→L1101 — all fixed + named regressions, seq-wrapped sum_to_n oracle-pinned 55. check: effects walk node-kind-keyed (Fanout unconditional, SeqBlock sticky), CK5 pin→theorem. Spec: LC-5 (§5.2 example had never parsed — patched to statement form); HANDOFF §4.1. New seq_demo.mapal golden through the full pipeline (36\n12\n, token thread only, pin d proven). As-built deltas (P0004/P0005/P0106+P0006 guard codes; P0117 void-only) reconciled into ADR/plan/ERRATA. vast.ai GPU access recorded for P6.
10 2026-07-16 P3 mapal-check (the owed checks) P3 complete — mapal-check designed + built; workspace 448 green (178+101+112+32+25). Plan from a 7-reader fan-out; DESIGN model-first, 4-lens adversarial design review pre-code (2 blockers: source param — Name is a span; lower §12 typing supersession · headline major: seq{} parses to the same StageKind::Fanout node as parallel — the walk keys on FanoutKind::Plain). Passes: T0101 Return exclusivity (strict CK3, hand-built-IR territory) + T0201 E2 (effectful? deduced from token signatures — no Pass-B recompute; sticky Plain context, inner seq does not legalize, CK5→OQ-C1 for Sapir). Typing discharged by construction (builder I2 + edge_type_ok; lower §12 amended). E3 vacuous-by-proof (no heap ops/ref types; zero code; reopen trigger = first heap-op ADR). Impl by Opus TDD + 3 adversarial reviewers + fixer (cross-pass-order fixture added, fails-on-reorder verified; 1 finding refuted); orchestrator line-by-line review. Discoveries: calc.mapal fully in-Core, first pipeline pass; local-shadow case is L1105 upstream. fmt+clippy clean.
09 2026-07-16 ADR-0018: zip/enumerate (same session, part 3) zip/enumerate are Core builtins, oracle-defined. Design discussion (sizes/generics/execution graphs → docs/notes/) produced ADR-0018; decided with Sapir, plan written model-first (ir/plans/plan-zip-enumerate.md), built by a sequenced workflow (ir → lower ∥ interp → examples; 4 Opus implementers + 3 adversarial reviewers + 1 fixer), every diff then orchestrator-reviewed line-by-line. ir: Zip/Enumerate ops, builder constructors (enumerate n ≤ i32::MAX bound, IrError/IrViolation twins), independent validate arms + 13 typing-table-golden rows, proptest generator coverage. lower: builtins routed like print (reserved L1009), 5 new L-codes L1606–L1610, D1 chain_seeded root-cause fix (fanout branches now typed with the scrutinee wire — pre-existing gap). interp (oracle-normative): elementwise-pair / (i as i32, x) arms + value contracts. Examples zip_demo/vector_add rewritten to the builtin form, golden through parse→lower→interp (c[0]=100, c[15]=115, sum=1720, e[k]=2k). Workspace 423 green (178+101+112+32), fmt+clippy clean. Capability matrix row added.
09 2026-07-16 Docs tree (ADR-0017) + CT-suggestion triage Category-architect tree live (ADR-0017): per-component IMPLEMENTATION.md (model → file:symbol functor, adversarially verified) + suggestions.md + plans/reviews/general; top-level architecture-map.md (§4.5 checklist: all six laws PASS), IMPLEMENTATION.md, suggestions.md, immutable sessions/ logs; HANDOFF §6/§7.2 patched. Suggestions triaged: 3 applied (lower resolve_tykind consolidation; ir §5.1 golden oracle test, +1 → 93; syntax LineIndex<'a> borrow — no source copy), 1 refuted on vet (interp width-seam: premise overstated), 4 parked with reasons. Every diff orchestrator-reviewed line-by-line after adversarial agent review; two extra §5.1 pins added (Eq-on-Str rejected, u8 index ok). Workspace 394 green (174+93+100+27), fmt+clippy clean. Found pre-existing uncommitted/unrecorded work (spec Mermaid lint diff, nvim, VISION.md, 3 out-of-Core generics examples, stray 2) — flagged for Sapir.
08 2026-06-15 P3 mapal-interp (the oracle) M1 oracle established — mapal-interp implemented; all six examples run correctly on CPU. Implementation found a blocker in the review-hardened interp/DESIGN §4 loop driver: the eager-both ("run the whole body, then test the guard") reading speculatively evaluated the continue-branch on the exit step, OOB-trapping fir (coeffs[4] at the exit state k=4) instead of 5.375. Resolved as a semantics question via ADR-0016 (guard-first loop evaluation) — grounded in E1/Elgot (U → B ⊕ U, the inr arm is not taken on exit), the "pure-branches-only → Phi" Core rule (a loop continue-branch can trap, so speculation is unsound), ADR-0013 traps, and the functor thesis (define meaning once → all backends inherit, differential-tested). DESIGN §4 rewritten guard-first (decide/exit cone → guard → advance). Built via a dynamic workflow (1 Opus implementer TDD-to-green + 4 adversarial reviewers incl. a deep Opus loop-driver reviewer): 0 blockers / 0 majors / 8 minors. Orchestrator fixed the one real finding — Le/Ge float NaN ordering now IEEE (was !Lt; +nan_ordering_is_ieee test) — and reconciled two code↔doc gaps (slotmap dep in §12; body_order degenerate-guard route-pack clause). Workspace 393 green (366 + 27 interp; fmt+clippy clean); interp_scale bench baseline recorded.
07 2026-06-14 interp DESIGN + println split mapal-interp DESIGN written + adversarially review-hardened (6-dimension review, 22 confirmed findings incl. 6 blockers — headline: the SCC loop-driver read the exit payload from an out-of-SCC route object → would have miscompiled every loop example; fixed via an incident-SCC body partition). Leads with its ADR-0014 categorical-model section; INDEX interp→modeled. ADR-0015 (print/println split) decided + implemented: Operation::Print{newline} + println builder (ir, +1 test); is_print_builtin helper routes the 9 effect/typing/emit sites that special-cased print (lower — one of them had regressed); examples use println for line output + print for pipeline's label; dump_ir example added (file → Category-IR Mermaid). Workspace 366 green (174 syntax + 92 ir + 100 lower; fmt+clippy clean); 13 snapshots regenerated + hand-verified Print→Println-only. Interpreter still unimplemented — next session.
06 2026-06-13 FRAMEWORK / categorical model layer (Level B) ADR-0014 accepted (ratified by Sapir 2026-06-14): FRAMEWORK.md adopted as the compiler-internal modeling method under a strict two-level firewall (Mapal-Cat = Level A, frozen, untouched; the compiler itself = Level B). New docs/architecture/categorical-model.md (incl. §7 reduction audit — 11/12 findings survived adversarial verification; one firm ADR candidate: the backend strategy-2-category / TargetText contract) + INDEX.md (10 component rows: syntax/ir/lower modeled, 7 planned). HANDOFF.md §7.1.5/§7.2 amended (mandatory ## Categorical model (Dat + Trn) DESIGN lead-section + FRAMEWORK §8 coherence line on the reconcile gate); syntax/ir/lower DESIGNs gained firewalled §0 model sections (point to the cross-cutting doc, no duplication). Built by a 6-phase dynamic workflow (73 agents: map → reduction-audit → synthesize → author → adversarial coherence review → critic); Fable-5 outage mid-run repaired by Opus. Methodology only — no spec/code/test change; workspace still 365 green.
05 2026-06-12 P2 mapal-lower mapal-lower complete — P2 done (workspace 365 tests green: 174 syntax + 91 ir + 100 lower). Binding lower/DESIGN.md written (passes A–E, literal-width unification, derives-from-merge tags, 46-code L1xxx catalogue, LD1–LD24 ledger) + 3-way Fable adversarial design review with per-finding Sonnet verification (38 confirmed findings applied; killed three would-be miscompiles: Phi-arm mut leaks, the 55→66 snapshot bug, Block-tail routing guards). Implementation by Opus agents + 2 impl reviews + Fable soundness attack (19 attack findings; all real ones fixed with named regressions — headline ATK-02: effectful-call loops now carry the token). mapal-ir empty-struct seal/validate hole (TY-1) fixed (+4 tests). All 8 golden snaps hand-read against §9 shape contracts (incl. orchestrator re-read after the head-naming fix). lower_scale bench recorded. Open: DESIGN §16 OQ1–OQ8 for Sapir.
04 2026-06-12 P2 mapal-ir mapal-ir complete (87 tests green; workspace 32 targets all ok): ADR-0013 + ERRATA LC-4 (spec's Pair-metadata/rhs_const dataflow conflict resolved → edges-only), DESIGN.md written + 3-way adversarial design review (26 findings applied — incl. the three token laws and exit-value pin, sum_to_n exits 55), implementation by Opus workflow agents + 2-way impl review + soundness attack (3 real breaches found: Str-param seal/validate gap, I5 route-vs-state SCC hole, u64→u32 arity truncation) + fix round with regressions. ir_scale bench recorded (100k morphisms: build+seal 65ms). lower/DESIGN §0.1 seeded with 5 pinned lowering obligations. Cross-builder id mixing pinned as UB — flagged to Sapir (nonce ADR if ever needed).
03 2026-06-12 P1 parser mapal-syntax parser complete — P1 done (174 tests green incl. the ADR-0012 amendment): two-tier grammar (ADR-0005), thin spanned parse tree, P0001–P0012 + P0101–P0116 diagnostics with recovery, ADR-0011 (loop labels / Ident { scan), golden parse trees for all 6 examples (zero diags, independently re-derived), full_surface precise rejection, criterion lex+parse bench (1–7.5 µs/example). Design 3-way + impl 2-way adversarially reviewed; totality stack-overflow defect found & fixed pre-merge. Post-review with Sapir: ADR-0012 labeled-block sigil (:label) decided + implemented (LC-3 spec patch); lower/DESIGN.md seeded with the parse-tree-obligations extract.
02 2026-06-11 P1 lexer mapal-syntax lexer complete (74 tests green): full-surface token set, ADR-0010 guard lexing, L0001–L0008 diagnostics, golden snapshots for all 6 examples + C8 fixture, proptest totality. Design + implementation each adversarially review-verified.
01 2026-06-11 M0 bootstrap M0 complete — skeleton green, E1–E5 applied + ERRATA, ADR-0001…0007, docs system, 6 examples.