Add gfx1250/MI450 target reference and AMD kernel-compiler survey

docs/audit/backend/apple/APPLE_AUDIT.md

@@ -1,5 +1,5 @@
 ---
-last_updated: 2026-07-13
+last_updated: 2026-07-28
 audit_role: sub_audit
 ---
 
@@ -291,6 +291,55 @@ not current support counts.
    hardware-free bridge remains `python/tessera/compiler/microscaling.py`.
 7. **Cross-backend real-hardware proof** for NVIDIA / ROCm remains tracked in
    the per-platform audits, not here.
+8. **Be the validation site for the hardware-free cost model.** Added 2026-07-28
+   from the AMD kernel-compiler survey
+   ([`../../compiler/AMD_KERNEL_COMPILER_SURVEY.md`](../../compiler/AMD_KERNEL_COMPILER_SURVEY.md)),
+   which found two *static, device-free* quality metrics in production AMD code:
+   a step-distance locality histogram over a materialized access order (§3.7) and
+   a bank-conflict analyzer that computes N-way conflict from a descriptor alone
+   (§3.8). Both are computable on any target and neither needs silicon.
+
+   The action is not "add a metric" but **calibrate one**: compute the
+   locality/conflict score for kernel families that are already measured, and
+   check the score against recorded latency. A metric that does not rank measured
+   kernels correctly should not be trusted to rank unmeasured ones. This is the
+   concrete follow-on to the mock-cost-model finding in
+   [`TILESIGHT_ASSESSMENT.md`](../../compiler/TILESIGHT_ASSESSMENT.md) §2, and it
+   gates how much weight the arbiter's hardware-free tier can carry.
+
+   **This is a cross-backend calibration, not an Apple-only one** (corrected
+   2026-07-29 — an earlier draft of this item said Apple was "the only backend
+   that executes broadly enough" and referred to "ROCm or NVIDIA kernels we
+   cannot measure"; both claims were false, written from the pre-bring-up framing
+   CLAUDE.md itself retired). NVIDIA has a committed, consumed, device-keyed
+   `nvidia:sm_120` autotune corpus over 64/256/512/1024/2048 square buckets plus
+   fused GEMM and causal attention
+   ([`../nvidia/NVIDIA_AUDIT.md`](../nvidia/NVIDIA_AUDIT.md); generator at
+   `benchmarks/nvidia/record_autotune_corpus.py`). ROCm has measured gfx1151
+   evidence including the size-adaptive grouped-GEMM tile selector, the hot-path
+   ratchet, and a measured resident-GPU crossover for sparse attention
+   ([`../rocm/ROCM_AUDIT.md`](../rocm/ROCM_AUDIT.md);
+   `rocm_gfx1151_compiler_retune_2026_07_15.json`).
+
+   Excluding those corpora would not just be unfair to the sibling backends — it
+   would make the calibration *weaker*. A locality metric fitted against one
+   architecture is the exact failure mode the same assessment records for learned
+   predictors: TileSight §5.2 shows NeuSight leading on the A100 in its training
+   distribution and losing that lead on every newer part. A score validated on
+   Apple alone would carry the same defect by construction.
+
+   Apple's actual distinctive contribution is **op breadth**, not exclusivity:
+   the widest F4-verified family envelope in the fleet, so it can test whether a
+   score generalizes *across op kinds*. NVIDIA and ROCm contribute **shape depth**
+   within GEMM/attention. The calibration needs both axes, so it is owned across
+   the queues under sync key `COSTMODEL-CALIB-2026-07-29` — see the per-backend
+   items named there.
+
+   A second, smaller item from the same survey: the MSL synthesizer currently
+   *authors* access patterns, whereas CK *derives* vector width, access count and
+   traversal order from a distribution encoding (§3.9). That is a design question
+   for the synthesizer, not a task — record it when the codegen path is next
+   revisited, and do not treat it as blocking.
 
 ## Hardware capability reference (grounded 2026-06-17)
 

docs/audit/backend/apple/todo.md

@@ -3,11 +3,46 @@ audit_role: plan
 plan_state: landing
 owner: Apple backend
 target: apple_gpu
-last_updated: 2026-07-28
+last_updated: 2026-07-29
 ---
 
 # Apple compiler, exact-device, and performance plan
 
+## APPLE-CALIB-1: contribute op breadth to the hardware-free score calibration
+
+Cross-backend sync `COSTMODEL-CALIB-2026-07-29` — **follow-up required, owning
+host M1 Max (apple7).** Apple is the *breadth* axis, not the sole site.
+
+**What is being calibrated.** Two static, device-free quality metrics found in
+production AMD code and recorded in
+[`../../compiler/AMD_KERNEL_COMPILER_SURVEY.md`](../../compiler/AMD_KERNEL_COMPILER_SURVEY.md)
+§3.7–3.8: a step-distance locality histogram over a materialized access order,
+and an N-way bank-conflict analyzer computed from a descriptor alone. Both are
+computable on any target with no silicon. The question is whether either
+*predicts measured latency* — which decides how much weight the arbiter's
+hardware-free tier can carry, per
+[`TILESIGHT_ASSESSMENT.md`](../../compiler/TILESIGHT_ASSESSMENT.md) §2.
+
+**Apple's role.** The widest F4-verified op-family envelope in the fleet, so it
+answers *does the score generalize across op kinds* — norm chains, attention with
+online softmax, pointwise-reduce, gated matmul, coopmat `simdgroup_matrix`, not
+just GEMM. NVIDIA and ROCm supply shape depth within GEMM/attention
+(`NVIDIA-CALIB-1`, `ROCM-CALIB-1`). Both axes are required: a score fitted on one
+architecture reproduces the overfit that assessment §5.2 records for NeuSight,
+which led on the A100 inside its training distribution and lost that lead on
+every newer part.
+
+**Apple-specific caveat.** The bank-conflict half was derived for LDS with a
+known bank count and a 4-phase wave64 access pattern. Metal threadgroup memory is
+not LDS and its banking is not documented to the same level (Decision #27 — do
+not assert a Metal hardware detail without a real source), so the conflict metric
+may be **not applicable** on Apple even where the locality metric is not. Report
+that split rather than one blended verdict.
+
+**Missing exact-device evidence.** Rank correlation between each score and
+recorded M1 Max latency, per op family, over the families the Apple lane already
+measures. A score that does not rank measured Apple kernels correctly is not
+trustworthy for unmeasured kernels anywhere.
 ## APPLE-RASTER-1: reconcile the MLX-inherited swizzle with the shared contract
 
 Cross-backend sync `RASTER-CONTRACT-2026-07-28` — **follow-up required, owning

docs/audit/backend/nvidia/todo.md

@@ -3,11 +3,49 @@ audit_role: plan
 plan_state: landing
 owner: NVIDIA backend
 target: nvidia_sm120
-last_updated: 2026-07-28
+last_updated: 2026-07-29
 ---
 
 # NVIDIA compiler test-suite evaluation and rearchitecture
 
+## NVIDIA-CALIB-1: supply the sm_120 corpus to the hardware-free score calibration
+
+Cross-backend sync `COSTMODEL-CALIB-2026-07-29` — **follow-up required, owning
+host NR2 Pro (RTX 5070 Ti, sm_120).**
+
+**Correction that created this item.** `APPLE_AUDIT.md` originally scoped this
+calibration to Apple alone, on the stated grounds that ROCm and NVIDIA kernels
+"cannot be measured". That was false for NVIDIA: this backend already has a
+committed, **consumed**, device-keyed `nvidia:sm_120` autotune corpus covering
+64/256/512/1024/2048 square buckets plus fused GEMM and causal attention,
+generated by `benchmarks/nvidia/record_autotune_corpus.py`. Excluding it would
+have discarded the deepest per-shape latency evidence in the fleet.
+
+**What is being calibrated.** Two static, device-free scores from
+[`../../compiler/AMD_KERNEL_COMPILER_SURVEY.md`](../../compiler/AMD_KERNEL_COMPILER_SURVEY.md)
+§3.7–3.8 — a step-distance locality histogram and an N-way bank-conflict
+analyzer — against recorded latency, to decide how much weight the arbiter's
+hardware-free tier can carry
+([`TILESIGHT_ASSESSMENT.md`](../../compiler/TILESIGHT_ASSESSMENT.md) §2).
+
+**NVIDIA's role: shape depth.** The committed corpus already varies the shape
+axis within GEMM and attention at fixed op kind, which is exactly the axis Apple
+cannot supply and the one a locality score most needs to be tested against —
+locality changes with shape at constant op. Apple supplies op breadth
+(`APPLE-CALIB-1`); ROCm supplies a second, independent architecture
+(`ROCM-CALIB-1`). Fitting on any one of the three reproduces the single-arch
+overfit the assessment records for NeuSight.
+
+**Note on translating the metrics.** Both were derived for AMD LDS with a known
+bank count and a wave64 4-phase access pattern (survey §5.1). CUDA shared memory
+is 32-bank and warp-synchronous; the *method* transfers but every constant must
+be re-derived for sm_120 before a conflict number here means anything. Do not
+port AMD constants.
+
+**Missing exact-device evidence.** Rank correlation between each score and the
+recorded sm_120 latencies already in the corpus — this needs no new hardware run,
+only an analysis pass over committed data, which makes it the cheapest of the
+three contributions to land.
 ## NVIDIA-RASTER-1: consume the shared block-rasterization contract
 
 Cross-backend sync `RASTER-CONTRACT-2026-07-28` — **follow-up required, owning