|
| 1 | +# Layered GEMM-rewrite verification: z3-first, escalate-on-unknown |
| 2 | + |
| 3 | +§V1 — TLA+/TLC race-freedom + z3-concrete / egglog algebraic equivalence |
| 4 | + |
| 5 | +This documents the standalone escalation MACHINERY for the GEMM-rewrite prover: |
| 6 | +**z3 FIRST**, and on `z3.unknown` **classify the obligation and escalate** to the |
| 7 | +tool that is natural for its category. It plugs into the EXISTING z3 interface |
| 8 | +(`_async_barrier_plan._z3_proves_simdgroup_output_isolated`, which returns |
| 9 | +`(z3_used, z3_proved, reason)` and the literal `"z3 returned unknown ..."` string |
| 10 | +on `z3.unknown`) and never rewrites it — it wraps that same verdict and dispatches. |
| 11 | + |
| 12 | +## RULE #1 honesty contract (no silent overclaim) |
| 13 | + |
| 14 | +| verdict / layer | `proved_by` | `bounded` | claim strength | |
| 15 | +|---|---|---|---| |
| 16 | +| z3 `unsat` | `z3` | `False` | ∀ over the z3-encoded (quantified) index domain | |
| 17 | +| z3 `sat` (counter-witness) | `unproven` | — | rewrite is WRONG; surfaced loud, **NOT** escalated to mask it | |
| 18 | +| TLC clean | `tlc-bounded` | `True` | race-free **at the checked grid only**, NOT ∀N | |
| 19 | +| z3-concrete-tiles | `z3-concrete-tiles` | `True` | bijection **for these concrete tiles only**, NOT ∀tile | |
| 20 | +| egglog | `egglog` | `True` | multiset equality **for this index domain only**, NOT ∀N | |
| 21 | +| no layer resolves / backend missing | `unproven` | — | reported UNPROVEN (never silently passed); fetch cmd surfaced | |
| 22 | + |
| 23 | +TLC is **BOUNDED model checking**: it proves the property AT the checked sizes, |
| 24 | +not ∀N. Every TLC pass is labelled `"verified BOUNDED @ dim=X tile=Y"` and |
| 25 | +`bounded=True`. The z3-concrete / egglog layers likewise prove only the concrete |
| 26 | +index domain. The only `bounded=False` claim is a z3 `unsat`, which is quantified |
| 27 | +over the encoded domain. Numeric fp parity over ALL elements is the empirical |
| 28 | +ground-truth backstop and is **NOT** a substitute for these structural claims. |
| 29 | + |
| 30 | +## Components |
| 31 | + |
| 32 | +- `cppmega_mlx/nn/_tilelang/verify_escalation.py` — top-level dispatcher |
| 33 | + `verify_with_escalation(name, z3_probe, obligation, prefer_egglog)`. Consumes the |
| 34 | + existing prover's `(z3_used, z3_proved, reason)` verdict (ONE z3 path), classifies |
| 35 | + proved / refuted / unknown, and on unknown dispatches by `kind`. |
| 36 | +- `cppmega_mlx/nn/_tilelang/_gemm_rewrite_escalation.py` — TLA+ spec auto-generator |
| 37 | + + TLC runner (`prove_race_freedom_tlc`, `render_writer_tla`), the z3-concrete |
| 38 | + bijection prover, the egglog multiset prover, and the representative-grid shrink |
| 39 | + (`representative_race_shape`) that preserves the overlap relation. |
| 40 | +- `cppmega_mlx/nn/_tilelang/verify_algebraic.py` — Track-B algebraic arm |
| 41 | + (`prove_bijection_z3_concrete`, `prove_multiset_equiv_egglog`, scale/mask identity |
| 42 | + refuters, and the DESIGN-ONLY Lean `finProdFinEquiv` lemma sketch). |
| 43 | +- `tools/tla2tools.jar` — the TLC engine (single jar, 2.27 MB). |
| 44 | +- Tests: `tests/test_verify_escalation.py`, `tests/test_gemm_rewrite_escalation.py`, |
| 45 | + `tests/test_verify_algebraic.py` — **53 passing**, lock in non-vacuity. |
| 46 | + |
| 47 | +## The TLA+ writer spec (race-freedom) |
| 48 | + |
| 49 | +`render_writer_tla` emits a MODULE of the tiled-GEMM writer state machine: block |
| 50 | +`(bm,bn)` owns the rectangle `[bm*MStep, +TileM) x [bn*NStep, +TileN)`; `Next` |
| 51 | +lets any block write any owned cell in any interleaving; `wrote[c]` accumulates the |
| 52 | +SET of distinct block-ids that wrote cell `c`; the safety invariant |
| 53 | + |
| 54 | +``` |
| 55 | +NoDoubleWrite == \A c \in Cells : Cardinality(wrote[c]) <= 1 |
| 56 | +``` |
| 57 | + |
| 58 | +says no output cell is ever claimed by two distinct blocks. TLC explores EVERY |
| 59 | +interleaving exhaustively over the bounded grid. A correct disjoint tiling |
| 60 | +(`step == tile`) holds; an overlapping tiling (`tile > step`) yields a TLC trace |
| 61 | +witnessing the double write. |
| 62 | + |
| 63 | +**Representative-grid honesty:** the production grid (64×64, 4×4 blocks) is |
| 64 | +exponential for TLC. We shrink each axis's `(tile, step)` to the smallest pair with |
| 65 | +the SAME overlap relation (disjoint ⇔ `tile==step`, overlap ⇔ `tile>step`, gap ⇔ |
| 66 | +`tile<step`) and cap blocks at 2 (NoDoubleWrite quantifies over PAIRS of blocks, and |
| 67 | +all blocks are identical, so two adjacent blocks exhibit every overlap pattern). The |
| 68 | +result is labelled with BOTH the production shape AND the representative grid it |
| 69 | +stands in for — the bounded guarantee stays honest. |
| 70 | + |
| 71 | +## MEASURED RESULTS (this run, macOS arm64 CPU) |
| 72 | + |
| 73 | +### Obligation (a): F0/B2 tiled-writer RACE — parallel threads writing dC / dchunk_states / DYX cells |
| 74 | + |
| 75 | +z3 returns UNKNOWN (race-freedom phrased as a concurrency interleaving, not a static |
| 76 | +ownership predicate) → escalate to **TLA+/TLC**. |
| 77 | + |
| 78 | +- prod `dim=64×64 tile=16×16 step=16×16 blocks=4×4`, via representative `dim=4×4 tile=2×2 step=2×2 blocks=2×2` |
| 79 | +- **`proved=True`, `proved_by="tlc-bounded"`, `bounded=True`** |
| 80 | +- TLC: `524289 states generated, 65536 distinct states found, 0 left on queue` — exhaustive over all interleavings |
| 81 | +- Label: *"NoDoubleWrite holds — verified BOUNDED @ dim=4×4 tile=2×2 step=2×2 blocks=2×2; NOT proven ∀N"* |
| 82 | + |
| 83 | +### Obligation (b): tiling-BIJECTION algebraic equivalence — serial reduction (i,j,k) domain ≡ tiled-GEMM domain |
| 84 | + |
| 85 | +z3 returns UNKNOWN (the bijection `k = kt*tile_k + kk` is nonlinear when `tile_k` is |
| 86 | +symbolic) → escalate to **z3-concrete-tiles** (PRIMARY), **egglog** (SECONDARY). |
| 87 | + |
| 88 | +- `K=64 tile_k=16 k_steps=4` |
| 89 | +- PRIMARY: **`proved=True`, `proved_by="z3-concrete-tiles"`, `bounded=True`** — injective + covering (`ForAll`) + no-overshoot, all `unsat` of the negation, on concrete divisors (Presburger/decidable). *"proves ONLY these tiles, not ∀tile size"* |
| 90 | +- SECONDARY egglog cross-check: **`proved=True`, `proved_by="egglog"`, `bounded=True`** — tiled product multiset == serial product multiset under comm+assoc of `+` (exact in the abstract ring; fp rounding is the separate numeric gate). *"proves THIS index domain only"* |
| 91 | + |
| 92 | +## NON-VACUITY (a known-WRONG rewrite is REFUTED by the right layer) |
| 93 | + |
| 94 | +| injected fault | layer | verdict | witness | |
| 95 | +|---|---|---|---| |
| 96 | +| overlapping writers `tile=20 > step=16` (real race) | TLC | **REFUTED** `unproven` | TLC trace: *"NoDoubleWrite is violated, two distinct blocks own the same cell"* @ rep `dim=5×5 tile=3×3 step=2×2` | |
| 97 | +| dropped k-tile `k_steps=3` (GAP, tiled sum drops products) | z3-concrete / egglog | **REFUTED** `unproven` | *"k present in serial but not tiled: [48,49,50,51,52,53,54,55]"* | |
| 98 | +| transposed operand ⇒ flipped causal mask (`k≤32` vs `k≥32`) | algebraic mask refuter | **REFUTED** `unproven` | *"mask mismatch at k=0: serial=True gemm=False"* | |
| 99 | +| z3 `sat` counter-witness | dispatcher | **REFUTED** `unproven`, NOT escalated | counter-witness carried; escalation must never paper over a real bug | |
| 100 | + |
| 101 | +All four faults are caught: the layers are non-vacuous, not trivially-passing. |
| 102 | + |
| 103 | +## z3 base-case (what z3 settles directly, no escalation) |
| 104 | + |
| 105 | +- z3 `unsat` → `proved_by="z3"`, `bounded=False`, `escalated=False` — the only ∀ claim (quantified over the encoded index domain). |
| 106 | +- z3 `sat` → `proved_by="unproven"`, `z3_resolved=True`, `escalated=False` — the rewrite is WRONG; surfaced loud, **NOT** escalated. |
| 107 | +- z3 `unknown` → `escalated=True` → the fallback layer above (or `unproven` if no layer resolves). |
| 108 | + |
| 109 | +In THIS session both real obligations hit `z3.unknown` (race = concurrency |
| 110 | +interleaving; bijection = nonlinear symbolic tile) and were resolved ONLY by the |
| 111 | +escalation layers, **bounded**. No real obligation here was settled ∀N by z3 alone; |
| 112 | +the ∀N path is exercised by the unit tests' synthetic `unsat` probe. |
| 113 | + |
| 114 | +## Coverage summary (honest) |
| 115 | + |
| 116 | +| obligation | z3 ∀N? | escalated? | resolved by | bound | |
| 117 | +|---|---|---|---|---| |
| 118 | +| F0/B2 tiled-writer race | no (unknown) | yes | `tlc-bounded` | dim=4×4 rep of prod 64×64, all interleavings | |
| 119 | +| tiling bijection | no (unknown) | yes | `z3-concrete-tiles` (+ egglog cross-check) | K=64 tile_k=16 k_steps=4 | |
| 120 | + |
| 121 | +No obligation in this run is left `unproven` — but the labelling is BOUNDED, never ∀N. |
| 122 | +The ∀N tiling bijection is the DESIGN-ONLY Lean `finProdFinEquiv` lemma sketch |
| 123 | +(`lean_bijection_lemma_sketch()`, ~40 LOC, not built this round — no Lean install). |
| 124 | + |
| 125 | +## Deferred wiring |
| 126 | + |
| 127 | +Final integration into the running w8ctouyfx z3 GEMM-equivalence prover (track C) is |
| 128 | +DEFERRED until that lands. This workflow builds the escalation machinery standalone |
| 129 | +against the `_async_barrier_plan` / `_gemm_rewrite_proof` interface and tests it on |
| 130 | +the two real obligations. Wiring is a drop-in: at each site where the track-C prover |
| 131 | +gets a z3 `unknown`, construct the matching `RaceObligation` / `AlgebraicObligation` |
| 132 | +and call `verify_with_escalation(...)`; the `MetalReductionSyncPlan` receipt gains |
| 133 | +`proved_by ∈ {z3, tlc-bounded, z3-concrete-tiles, egglog, unproven}` via |
| 134 | +`as_feature_dict()`. |
0 commit comments