# Native TileIR / TVMx vs PyTorch

This is an opt-in **correctness-checked, multi-shape** benchmark, not a CTest
performance threshold. It compares FP32 GEMM, add, row sum, softmax, RMSNorm,
LayerNorm, fused residual LayerNorm and per-row cross-entropy on CPU and actual
Metal / PyTorch MPS.
GEMM includes small/large squares, tall/wide matrices, and non-multiple tail
sizes; reductions vary both row count and width.

`--operations gemm_relu,gemm_gelu` also measures matrix-based expression graphs:
`max(0.125 * (A @ B) + 0.25, 0)` and tanh-GELU of that affine transform.
They share `--gemm-shapes` and matrix schedule controls with plain GEMM.
Both outputs are checked against the complete FP64 expression. Torch uses
preallocated intermediates and eager `mm.out`, scale/shift and activation;
this is not a compiled Torch or MPSGraph fusion baseline. Direct MPS/BLAS
comparison remains limited to plain GEMM, which has matching semantics.

With `--cooperative-matrix --matrix-realization mpp-views`, the optional
`--fuse-matrix-epilogues` enables the closed same-owner scalar-DAG candidate.
It requires the optional MPP element extension and otherwise retains the
established realization. The default is off: legality does not establish
profitability. [Fixed-schedule evidence](results/m1-max-20260907-fragment-epilogue/notes.md)
retains both improvements and large regressions. Operator names in this driver
select test graphs, not production matching rules.

The default GEMM cohort stops at 1024³. For explicit scale coverage, use
`run.py --gemm-shapes 2048x2048x2048,8192x8192x8192,256x11008x4096`;
`--row-shapes` independently selects non-GEMM M×N cases. Explicit lists keep
their order and override `--quick` only for that operator family. Dimensions
are bounded by the native driver's 16384 limit; the caller must still budget
FP64-oracle and device memory.

`compare_lowerings.py --metal-device-timing PATH_TO_HELPER` adds separate
no-counter GPU command-buffer controls for every comparison path. Host tables
remain E2E; GPU tables never substitute compute-pass probe times or host times.
Handwritten MPP keeps its direct command-buffer timer. Both GPU and host
single-call samples are retained in JSON, and missing controls fail the row.

Latest M1 Max evidence: [shape-held-out service policy, including small-case regressions](results/m1-max-20260905-service-policy-validation/notes.md),
[joint resource/width GPU and E2E acceptance](results/m1-max-20260905-access-demand-validation/notes.md),
[Metal subgroup sum/softmax/RMSNorm cohort](results/m1-max-20260905-metal-subgroup-reductions/notes.md),
[balanced same-binary RMSNorm lowering A/B](results/m1-max-20260905-metal-subgroup-rmsnorm-replay/notes.md),
[Metal subgroup LayerNorm/cross-entropy cohort](results/m1-max-20260905-metal-subgroup-row-extensions/notes.md),
[balanced LayerNorm/cross-entropy lowering A/B](results/m1-max-20260905-metal-subgroup-row-extensions-replay/notes.md),
[residual-LayerNorm materialization search](results/m1-max-20260905-residual-layernorm-materialization-search/notes.md),
[balanced shared-Tile materialization A/B](results/m1-max-20260905-residual-layernorm-materialization-replay/notes.md),
[bounded residual-LayerNorm thread search](results/m1-max-20260905-residual-layernorm-bounded-thread-search/notes.md),
[CPU residual-LayerNorm materialization search](results/m1-max-20260905-cpu-residual-layernorm-materialization-search/notes.md),
[MPP cost-model v1→v2 calibration-cohort study](results/m1-max-20260905-mpp-cost-v2-search/notes.md),
[frozen seven-path MPP-v2 replay](results/m1-max-20260905-mpp-cost-v2-replay/notes.md),
[proved CPU CBLAS plan](results/m1-max-20260905-cpu-cblas-v2-plan/notes.md),
[six-order CPU CBLAS/Torch/direct-BLAS replay](results/m1-max-20260905-cpu-cblas-v2-replay/notes.md),
[CPU reference add/sum/softmax plan](results/m1-max-20260905-cpu-reference-ops-v2/notes.md),
[CPU Accelerate add/sum/softmax plan](results/m1-max-20260905-cpu-accelerate-ops-v2/notes.md),
[six-round CPU array-math policy A/B](results/m1-max-20260905-cpu-accelerate-ops-replay/notes.md),
[CPU provider/code/docs validation](results/m1-max-20260905-cpu-provider-validation/notes.md),
[ragged CPU guard repair A/B](results/m1-max-20260905-cpu-guards-replay/notes.md),
[fresh guarded-view CPU/Torch/BLAS comparison](results/m1-max-20260905-cpu-guards-system/notes.md),
[seven-path Metal maintenance smoke check](results/m1-max-20260905-guards-lowerings-smoke/notes.md),
[other-operator guard smoke check](results/m1-max-20260905-cpu-guards-ops/notes.md),
[retained scalarization failure](results/m1-max-20260904-cpu-guard-plan/notes.md),
[CPU input-view A/B](results/m1-max-20260904-cpu-views-replay/notes.md),
[fresh input-view CPU/Torch/BLAS comparison](results/m1-max-20260904-cpu-views-system/notes.md),
[CPU Cartesian packing A/B](results/m1-max-20260904-cpu-cartesian-replay/notes.md),
[fresh Cartesian CPU/Torch/BLAS comparison](results/m1-max-20260904-cpu-cartesian-system/notes.md),
[seven-way explicit subgroup-sync candidate](results/m1-max-20260904-subgroup-sync-lowerings/notes.md),
[fixed-geometry subgroup-fence A/B replay](results/m1-max-20260904-subgroup-sync-replay/notes.md),
[seven-way replay with proved read-only TIRx views](results/m1-max-20260904-tirx-views/notes.md),
[independent TIRx MPP codegen and its retained regression](results/m1-max-20260904-tirx-mpp/notes.md),
[seven-way native/TIRx/MPP/MPS/Torch and Runtime controls](results/m1-max-20260904-runtime-controls/notes.md),
[rejected fragment-load batching, with saved patch and six-round replay](results/m1-max-20260904-fragment-batch/notes.md),
[five-path native/TIRx/MPP/MPS/Torch replay](results/m1-max-20260904-joint-lowerings/notes.md),
[larger tiles and controlled cooperative-copy batching](results/m1-max-20260903-copy-plan.md),
[six-round direct Accelerate/BLAS and MPS comparison](results/m1-max-20260903-system-baselines.md),
[execution planner, cost-model errors, and controlled comparisons](results/m1-max-20260903-planner.md),
[proved direct accumulator output and controlled measurements](results/m1-max-20260903-direct-store.md),
[dependence-aware synchronization and controlled measurements](results/m1-max-20260903-barrier-plan.md),
[matrix/copy layout sensitivity missing from the model](results/m1-max-20260903-layout-tie.md),
[rejected structural experiments and joint execution search](results/m1-max-20260904-tirx-structure.md),
[Staged/JIT and four-round comparisons](results/m1-max-20260903-jit.md),
and [actual PyTorch dispatch / Xcode profiling](results/m1-max-20260903-profile.md).
The reports include unsuccessful tuning choices and remaining library gaps.

## Top-K and sort across SIMD and Metal

`ranking.py` adds a separate multi-output ranking matrix. It exercises the
existing finite-FP32, stable-index `topk`/`sort` library composition, currently
quadratic reference code rather than tuned selection/sorting. The shared
`tile_rank_test_utils.h` fixture also drives Runtime regression tests.
Values are FP32 and indices are int64; both outputs, input immutability and
all three allocations' guards are checked. Correctness tests additionally
cover unique values, all-equal values and signed zero.

Complete the configured **full build first**; this driver never builds:

```bash
cmake --build BUILD --parallel 6
uv run --no-project --python 3.13 --with numpy --with torch==2.14.0 \
  python scripts/benchmark/tile_torch/ranking.py \
  --native xir-simd=BUILD/bin/benchmark_tile_xir \
  --native xir-metal4=BUILD/bin/benchmark_tile_xir \
  --native tirx-metal=BUILD/bin/benchmark_tile_native \
  --torch cpu --torch mps \
  --case topk:17,65,8 --case sort:17,65,65 --direction both \
  --samples 5 --sample-ms 20 --warmup-ms 100 \
  --metal4-device-timing \
  --metal-device-timing BUILD/bin/libluisa-benchmark-metal-timing.dylib \
  --output NEW_DIRECTORY
```

Without `--case`, ten R/N/K combinations cover K=1, small K, K near N,
power-of-two boundaries and widths through 1025, in both directions. Add
explicit larger widths with `--case topk:1,4097,32`; budget for the quadratic
reference implementation. Native timeouts and errors remain in the
predeclared matrix and do not suppress independent Torch baselines.
Default rounds are twice the route count, balancing route position and pair
precedence; a manually shortened run is a **screen**, not a balanced ranking.

The timed input has 31 distinct dyadic values, including negatives and ties.
This is not representative of every logits distribution. Tile and stable
Torch sort require original-index tie order. Torch top-k may select different
tied indices, but must return the exact sorted top-K value multiset, unique
in-range indices and correct value/index correspondence. NaNs are outside
this fixture's finite-value contract.

Native and Torch results are warm synchronized **E2E** timings with
preallocated outputs, not pure CPU kernel times. Optional Metal4 dispatch,
legacy Metal compute-pass and no-counter command-buffer intervals are
reported separately. Torch MPS CPU fallback is disabled. The `rank` mode of
`benchmark_tile_native` means **TIRx/Metal**, not native MPP. Each native route
runs in its own process, avoiding TVM/native LLVM coexistence.

For a direct XIR diagnostic invocation, `LUISA_TILE_BENCH_XIR_REGION_WORK`
overrides the shared planner's structured-map expansion budget (4096 by
default; zero preserves the old element-only decision). Record this separately
from algorithmic complexity or performance: it bounds duplicated map-body work,
not native instructions, kernel cycles or the quadratic reference algorithm.

The XIR `llm`/`rank` diagnostic also accepts
`LUISA_TILE_BENCH_XIR_MMA_OUTPUT_BLOCK=1|2|4` (default 1). Normal API callers
set `PlannerOptions::mma_output_block`; metadata distinguishes the request
from admitted `blocked_mmas`. This is an opt-in contiguous-output contraction
candidate, not an operator-name rule or a calibrated automatic policy. The
[attention interaction experiment](results/m1-max-20260913-attention-mma-block/notes.md)
retains both gains and a code-size/full-packet-specialization regression.

`LUISA_TILE_BENCH_XIR_MMA_UNROLL_TERMS` sets the additional MMA-only unroll
cap (`PlannerOptions::max_unrolled_mma_terms`, default 0 inherits the existing
Tile threshold). It can keep K as a runtime loop without changing the global
Tile traversal threshold or other reductions. Metadata reports the requested
cap and actual `rolled_mmas`; this remains an opt-in, uncalibrated candidate.
Smaller generated code is not sufficient evidence of faster execution.

## Common LLM captures across SIMD and Metal

`compare_llm.py` reuses the unit tests' Tile kernels, exports the exact native
inputs for Torch, checks every output against an independent FP64 formula,
and retains failures rather than omitting unsupported/slow shapes. It first
builds the complete selected configuration. Use `benchmark_tile_xir` for CPU
XIR/SIMD or the `llm` mode of `benchmark_tile_native` for **TIRx/Metal** through
Luisa Runtime; the latter is not the restricted native-MPP matrix emitter.

```bash
uv run --no-project --python 3.13 --with numpy --with torch \
  python scripts/benchmark/tile_torch/compare_llm.py \
  --native BUILD/bin/benchmark_tile_native --build-dir BUILD --backend metal \
  --case swiglu:1024,4096 --case rope:4096,4096 \
  --metal-device-timing BUILD/bin/libluisa-benchmark-metal-timing.dylib \
  --rounds 4 --samples 5 --sample-ms 30 --warmup-ms 100 \
  --output NEW_EMPTY_DIRECTORY
```

`--case attention:B,Hq,Hkv,Q,K,D,Dv --attention-block BQ BK` measures the same
online causal/GQA recurrence. Queries occupy the last Q KV positions; the
Torch SDPA reference therefore uses an explicit bottom-right-aligned mask.
SwiGLU and RoPE preallocate Torch intermediates/output. Norm, residual-GELU,
masked softmax and SDPA use functional eager APIs with their allocations
inside timing. They are not compiled/fused Torch baselines.

Use the `llm` entry of `benchmark_tile_tirx` for controlled Metal mapping
experiments: `--subgroup-reductions`, `--group-threads N` and
`--forward-input-views` are acknowledged separately in the recorded metadata.
The current subgroup family also attempts input-view forwarding, so its
on/off switch alone does **not** isolate collective emission. To isolate that
choice, request input views and the same exact group width on both controls;
also inspect the emitted source. An explicit view request enables only proved
forwarding, not arbitrary removal of snapshots or an altered numerical policy.

`--attention-qk reduce` and `--attention-pv reduce` are independent,
explicit **benchmark-only** decomposition probes. QK becomes
`reduce(query * key, d, add)`; PV becomes
`acc * alpha + reduce(probability * value, n, add)`. Both default to `mma`.
These test whether different contraction access directions need different lane
distributions. They are not planner improvements or new DSL primitives.
The native metadata must acknowledge each explicit probe; inherited probe
environment settings are cleared. The probes may change floating-point
evaluation order and must pass the same complete FP64 oracle, not a relaxed
tolerance or a claim of bitwise identity. Inspect actual phase mappings and
intermediate storage before attributing any timing difference to collectives.

The same controls are available through this driver with `--backend cpu|metal4`
and `--native BUILD/bin/benchmark_tile_xir` (the executable itself takes
`LUISA_TILE_BENCH_XIR_BACKEND=simd|metal4`). Non-attention cases in a mixed
matrix do not inherit attention-only controls. Torch runs in an isolated process
per visit: OS stdout/stderr, including native-library messages, are retained.
Known GPU failure diagnostics invalidate all cohort comparisons even at exit
code zero; remaining GPU visits are NotRun until queue health is re-established.
The diagnostic list is not exhaustive. Fresh-worker startup is outside warm
samples, but this changes the process-lifetime protocol of older reports.

For actual single-thread SIMD entry replay, capture with
`LUISA_TILE_BENCH_DUMP_SOURCE=OUTPUT.f32.source.txt` and
`LUISA_SIMD_DUMP_ASSEMBLY_DIR=NEW_OBJECT_DIRECTORY`. The LLM entry also exports
`OUTPUT.f32.expected.f64` and actual dispatch metadata. Prepare with
`native_tile.py prepare --capture-kind llm --prefix OUTPUT.f32 --log LOG
--objects NEW_OBJECT_DIRECTORY --output NEW_PREPARED_DIRECTORY --name VARIANT
--llvm LLVM_BIN_DIRECTORY`, then use its ordinary two-entry `replay` command.
This links the captured ORC object, not a recompiled LLVM approximation; all
three inputs and the complete output are checked. See the
[attention native-entry checkpoint](results/m1-max-20260913-attention-native/notes.md)
for the negative decomposition results and exact timing boundary.

For Apple arm64/macOS 15+ diagnostic references,
`native_attention_probe.py prepare --tile TILE/prepared.json --output NEW_PROBE
--llvm LLVM_BIN_DIRECTORY` freezes that actual Tile object beside handwritten
online NEON and dense Accelerate attention. Then use `replay --bundle
NEW_PROBE/probe.json --variant online_neon|dense_accelerate --output NEW_RESULTS`.
Both arms share the unchanged native timer; BLAS threading is set and queried
on the calling thread. `selftest --output NEW_VALIDATION` runs 24 numerical/guard
checks without timing loops. Run commands in externally supervised subprocesses
with complete stdout/stderr capture and process-group timeouts. These references
change reduction order/storage and are **not new Tile lowering paths**; see the
[six-shape comparison and boundaries](results/m1-max-20260913-attention-native-reference/notes.md).

For the experimental **compiler-generated** per-program SIMD MMA realization,
set `LUISA_SIMD_NATIVE_MMA_VECTOR_WIDTH=2|4|8` when capturing the Tile kernel;
`0` (the default) preserves the existing lowering. This is an inner vector
width, independent of the program packet width. Admission uses typed MMA
axes, snapshot strides and the operation's math permissions, not operator
names. The realization reports native call counts and actual snapshot costs;
it is not yet an automatically cost-selected policy. Compare actual captured
ORC objects with `native_tile.py`; storage/interleaving and authorized reduction
order can change, so validate each arm against FP64 and require each replay to
match its own capture instead of imposing bitwise equality between arms.

`--baseline FROZEN/bin/benchmark_tile_native --rounds 6` balances all six
orders of old/new/Torch. Freeze all adjacent Luisa libraries together; the
driver prepends that directory to the baseline process's loader path. Pass
the actual external compiler libraries with repeated `--compiler-artifact`.
Host batch/single-call timings and separate instrumented/uninstrumented GPU
phases retain their different meanings. Large tensor exports are temporary;
JSON records their SHA256 and complete validation, and retains generated
source and process logs. No timing selects a schedule. See
[the first LLM coverage and partitioned-output replay](results/m1-max-20260907-llm-coverage/notes.md).

## Actual native row entries without Runtime

`native_rows.py` compares actual emitted XIR/SIMD ORC entries and generated
TorchInductor C++ entries through one C++ callback timer. It handles static
FP32 input order, scratch reuse, const-correct pointer ABIs and partitioned
output aliases by inspecting the real generated wrapper; it does not
reimplement the operators. Unknown wrapper effects fail closed.

Capture, preparation and timing are separate phases. This checkpoint's
capture deliberately verifies the source and binary closure against
`results/m1-max-20260909-cohort-private/provenance.json`; it is not an
unversioned benchmark for arbitrary builds. Establish a new recorded
baseline when updating compiler sources. Each output directory must be new.

```bash
uv run --offline --no-project --python 3.13 --with numpy \
  python scripts/benchmark/tile_torch/native_rows.py capture \
  --binary VERIFIED_BUILD/bin/benchmark_tile_xir --output NEW_CAPTURE
uv run --offline --no-project --python 3.13 --with numpy --with torch==2.14.0 \
  python scripts/benchmark/tile_torch/native_rows.py prepare \
  --capture NEW_CAPTURE --output NEW_PREPARED
uv run --offline --no-project --python 3.13 --with numpy --with torch==2.14.0 \
  python scripts/benchmark/tile_torch/native_rows.py replay \
  --prepared NEW_PREPARED --output NEW_REPLAY
uv run --offline --no-project --python 3.13 --with numpy \
  python scripts/benchmark/tile_torch/audit_native_rows.py \
  --capture NEW_CAPTURE --prepared NEW_PREPARED --replay NEW_REPLAY \
  --output NEW_AUDIT.json
```

The timer excludes Runtime/Python/JIT and caller allocations. It includes
the common callback, actual entry, required block traversal/launch resets,
and compiler-emitted libc/internal allocations. No build, test or other
benchmark should run during replay. Six balanced orders, all samples,
complete FP64 checks, guards and repeated-output hashes are retained.
See the [24-case native report, including all losses](results/m1-max-20260909-native-rows/notes.md).
This is single-thread native evidence, not multi-thread E2E or GPU timing.

## Direct XIR/SIMD planner comparison

`compare_xir.py` is the independent CPU pilot for the direct
TileIR→XIR→SIMD Runtime path. It does not call or relabel the TIRx binary. It
balances automatic planning, a fixed `{root order [0,1], 64 workers/block}`
control and eager Torch over all six orders, retains full outputs and LLVM,
and rechecks an FP64 oracle before accepting a timing row:

```bash
uv run --no-project --python 3.13 --with numpy --with torch \
  python scripts/benchmark/tile_torch/compare_xir.py \
  --native BUILD/bin/benchmark_tile_xir \
  --compiler-artifact /path/to/the/actual/libLLVM.dylib \
  --output NEW_EMPTY_DIRECTORY \
  --rounds 6 --samples 5 --sample-ms 20 --warmup-ms 100 --threads 8
```

The C++ binary also accepts `planned`, `canonical`, or `reversed`, followed
optionally by an exact block-worker count. Planner policy, realized order,
cost decomposition, source identity, cold/JIT phases and both throughput and
single-call latency samples are emitted as JSON. The report hashes the binary,
all adjacent Luisa dynamic libraries and explicit compiler artifacts before
and after timing. Do not run builds, tests or profilers concurrently.

The first saved pilot is
[m1-max-20260905-xir-simd](results/m1-max-20260905-xir-simd/notes.md).
It is intentionally a negative-result report: a narrow mapping search does
not replace packed/register-blocked CPU matrix realization.

For a compiler A/B instead of a mapping A/B, add
`--baseline FROZEN_BUILD/bin/benchmark_tile_xir`. Copy its complete adjacent
backend/library set before rebuilding; copying only the executable still loads
the new backend. Both arms then request `planned`, and six orders balance the
old compiler, new compiler and Torch. Repeat `--shape M,N,K` to cover larger
or ragged matrices, and use `--block M,N,K` to fix the same Tile specialization
for both arms. Do not choose a new block during the frozen replay. The
[packet-index proof report](results/m1-max-20260906-xir-packet/notes.md)
also independently checks the plans, source/output hashes and both dispatch
metrics; its CPU timings are not kernel-only durations.

## Direct BLAS and MPS GEMM baselines

On macOS the build also creates `benchmark_tile_system`, an independent
Objective-C++ executable linked only to system frameworks, not TileIR or TVM.
It uses CPU Accelerate `cblas_sgemm` (classic LP64 API) and GPU
`MPSMatrixMultiplication` (`MPSKernelOptionsNone`, no reduced-precision option).
These are **comparison baselines, not replacement lowering paths**.
PyTorch's default MPS execution is listed separately; it is not assumed to use
the same internal kernel as the direct MPS matrix API.

Both library baselines use compact row-major FP32 `C=A*B`, alpha=1, beta=0,
no transpose, and the same deterministic values as the Tile/PyTorch driver.
MPS inputs and output use private device buffers, uploaded once before warm
timing; no padded strides or input prepacking are introduced. Every repeated
call writes the same preallocated output. One timed MPS batch encodes all
calls into one command buffer, commits and waits once. Single-call latency
uses a separate command buffer per call. API/encoding/submission costs remain
in warm timings; library-internal scratch management is not artificially
subtracted. Timing is host wall time, **not pure GPU time**.

After a complete build, validate both native paths explicitly:

```sh
ctest --test-dir cmake-build-tirx -R '^test_tile_system_' --output-on-failure
```

The CPU self-test is a host unit test; the Metal self-test is an explicit
hardware integration test. Both check every output against FP64 on four
shapes, including tails and repeated beta=0 calls starting with NaN output.
Add `--system-baseline cmake-build-tirx/bin/benchmark_tile_system` to `run.py`
to include the appropriate system baseline for each GEMM case; other
operations retain the existing two-implementation comparison.

For published comparisons, freeze valid CPU/Metal schedules in ordinary
`run.py` reports, then replay all three implementations in six balanced orders:

```sh
uv run --no-project --python 3.13 --with torch --with numpy python \
  scripts/benchmark/tile_torch/compare_system.py \
  --native cmake-build-tirx/bin/benchmark_tile_tirx \
  --system-baseline cmake-build-tirx/bin/benchmark_tile_system \
  --plan /path/to/cpu/results.json --plan /path/to/metal/results.json \
  --output /tmp/tile-system-comparison --rounds 6 \
  --samples 9 --sample-ms 40 --warmup-ms 200 --threads 8
```

The old reports supply **parameters only**, never timing scores. Each shape
gets all six permutations of Tile/PyTorch/system order; case order rotates.
All outputs use the same full FP64 check. Failures remain failures, not dropped
rounds, and binaries are hash-checked before/after the run. Raw timings,
individually synchronized latency, API/storage/stride metadata and correctness
errors are retained. Thread settings are requests, not measured library worker
counts. Compare device-specific baselines within each backend, not CPU BLAS
against Metal as if they used the same resources.

## Tile/PyTorch driver

For the LLVM path, `--cpu-stack-bytes 8192` enables the bounded
compiler-temporary stack realization. Zero (the default) retains workspace
allocation. The budget is cumulative with alignment padding, has a 65536-byte
ceiling, and does not override explicit Memory or permit pointer escapes.
It changes storage realization, not the tile geometry or arithmetic policy.
`repeat.py --candidate-cpu-stack-bytes 8192` changes only the candidate, so the
same frozen schedule and a pre-change binary can serve as the reference.
`compare_system.py --cpu-stack-bytes 8192` overrides CPU cases only and still
compares with eager PyTorch and direct Accelerate BLAS in six balanced orders.

`--auto-vectorize --cpu-vector-lanes 64` enables Cartesian CPU register packs.
The lane budget accepts 16/32/64/128; 16 preserves the single-row realization.
Larger packs keep separate contiguous row vectors inside a common serial
recurrence, exposing multi-row operand reuse without reassociating K. The
budget is not a hardware vector width. Unsupported regions keep the old
single-row packing, and CPU stack planning is independent.
`repeat.py --cpu-stack-bytes 8192 --candidate-cpu-vector-lanes 64` holds storage
fixed for both variants while changing only the candidate's pack budget.
`compare_system.py --cpu-vector-lanes 64` changes CPU cases only. Replay keeps
the exact reported policy; missing legacy lane metadata means 16, never an
inferred target-specific default.

`--cpu-input-views` independently enables proved immutable input expressions
on LLVM. Padded reads retain their original lazy guard and fill; explicit
Memory, mutable or aliased inputs, pointer escapes, and unproved consumer
indices keep snapshots. This is not MPP, and does not change its strict view
policy or the default reference lowering.
Use `repeat.py --cpu-stack-bytes 8192 --cpu-vector-lanes 64
--candidate-cpu-input-views` to hold both storage and register packing fixed
while changing only input forwarding. `compare_system.py --cpu-input-views`
applies to CPU cases only, preserving fresh Torch/BLAS comparisons. Keep
regressions: direct global reads can lose the benefits of compact packing.

`--cpu-matrix-backend cblas` is restricted to CPU GEMM with automatic root
binding. The bridge accepts it only for a proved whole compact rank-two FP32
`C=A*B` contract with noalias buffers and a registered
`tvm.contrib.cblas.matmul` provider. The native JSON must report
`cpu_matrix_backend="cblas"` and exactly one semantic external matrix call;
the runner rejects a missing or unexpected call. This is an actual Tile
lowering candidate, unlike `--system-baseline`, which is a separate executable
used only as a control. Keep both in reports when measuring provider overhead:

```sh
uv run --no-project --python 3.13 --with torch --with numpy python \
  scripts/benchmark/tile_torch/run.py \
  --native cmake-build-tirx/bin/benchmark_tile_tirx \
  --system-baseline cmake-build-tirx/bin/benchmark_tile_system \
  --output NEW_EMPTY_DIRECTORY --backends cpu --operations gemm \
  --cpu-model native --cpu-matrix-backend cblas \
  --samples 9 --sample-ms 40 --warmup-ms 200 --threads 8 --capture-sources
```

`--cpu-math-backend accelerate` is restricted to CPU but can be combined with
add/sum/softmax. It leaves add unchanged, realizes structurally proved
contiguous add/max/min reductions with vDSP, and realizes only a versioned
compiler-owned shared FP32 exp map with vForce. The static call-site diagnostic
is zero for add and nonzero for eligible sum/softmax. It can be larger than the
number of provider kinds when a small serial root is unrolled, so it is not a
dynamic call counter.

This policy permits provider reduction order and vForce denormal/floating-
exception behavior; reference is the default. Compare the two policies with
separate initial reports and `repeat.py`, which preserves each report's
`cpu_math_backend` field:

```sh
uv run --no-project --python 3.13 --with torch --with numpy python \
  scripts/benchmark/tile_torch/repeat.py \
  --reference /path/to/reference/results.json \
  --candidate /path/to/accelerate/results.json \
  --native cmake-build-tirx/bin/benchmark_tile_tirx \
  --output NEW_EMPTY_DIRECTORY --operations add,sum,softmax \
  --rounds 6 --samples 7 --sample-ms 40 --warmup-ms 200 \
  --threads 8 --capture-sources
```

Automatic CPU roots also report `cpu_parallel_task_threshold` (currently 64).
Below it, cheap automatic roots stay serial; explicit worker roots and bodies
with transcendental/opaque calls retain parallel mapping. This is a scheduling
prior, not a change to the source `parallel` semantics.

`--shared-tile-materialization preserve` is the structural TIRx default. It
retains every pure multi-consumer Tile SSA definition so a target pass can
still choose a compact resource realization. `expensive-only` is an explicit
diagnostic/tuning candidate matching the earlier policy: shared
transcendentals remain materialized while cheap arithmetic is cloned into its
consumers. Neither spelling is a source-level `Memory`, and neither changes
the TileIR computation.

Use `--tune-shared-tile-materializations preserve,expensive-only` to recapture,
JIT, fully validate and measure both candidates. The choice joins the same
finite Cartesian product as block/window/thread/copy options and counts against
`--max-tuning-candidates`; the winner is then recaptured and revalidated. The
current Metal residual-LayerNorm cohort selects `preserve`, while its CPU LLVM
counterpart selects `expensive-only`. This is why the runner records the
policy in global metadata, every candidate and every frozen replay.

`--capture-sources` archives both LLVM IR (`.ll`) and Metal (`.metal`) by SHA256.
The repeat/system runners fingerprint both executables and their adjacent
shared libraries; use repeatable `--compiler-artifact PATH` arguments to cover
external TVM compiler/runtime libraries. Every replay checks the reported CPU
budget exactly; older reports without the field mean the old zero-budget path.

The [M1 Max CPU storage A/B](results/m1-max-20260904-cpu-stack-replay/notes.md)
and [six-order Torch/BLAS comparison](results/m1-max-20260904-cpu-stack-system/notes.md)
record the first measured realization. Its median gains over the old lowering
do not establish library parity; 512³ regresses in two A/B pairs and the policy
remains default-off. Raw LLVM, complete output checks, and artifact fingerprints
are retained alongside the reports.

First configure TVMx support as described in the Tile design document, and
complete the full build and correctness tests. The driver never builds or
changes the build configuration:

```sh
cmake --build cmake-build-tirx -j 8
ctest --test-dir cmake-build-tirx -L unit_tile --output-on-failure
uv run --no-project --python 3.13 --with torch --with numpy python \
  scripts/benchmark/tile_torch/run.py \
  --native cmake-build-tirx/bin/benchmark_tile_tirx \
  --output /tmp/tile-torch-results \
  --threads 8
```

The output directory must not exist. Use `--quick` for a smoke run only;
published comparisons should include the full default matrix. `--gemm-block
8,8,16` changes ordinary host configuration and therefore captures/JITs another
native variant. Record that setting, and do not silently select the best of a
larger tuning search against an untuned baseline.

An explicit search can recapture/JIT ordinary host configurations per shape:

```sh
uv run --no-project --python 3.13 --with torch --with numpy python \
  scripts/benchmark/tile_torch/run.py \
  --native cmake-build-tirx/bin/benchmark_tile_tirx \
  --output /tmp/tile-metal-jit-search \
  --backends metal --operations gemm --execution-scope group \
  --cooperative-matrix \
  --tune-gemm-blocks '8,8,16;16,32,32;32,32,32' \
  --tune-pipeline-windows '1,2'
```

Metal group searches can also specify `--tune-group-threads '128,256'` and
`--tune-copy-batches '1,4,8'`. Any backend may add
`--tune-shared-tile-materializations preserve,expensive-only`. These form a
joint product with the block/window lists; they do not alter the frontend
kernel semantics. Zero in the thread list invokes the planner's automatic
choice, not a zero-thread launch.
Unspecified dimensions retain their ordinary command-line setting. Duplicate
configurations are removed; a product exceeding `--max-tuning-candidates`
(default 256 per shape) is rejected, never silently truncated. Include the
incumbent schedule in the lists when assessing an optimization.

No tuning happens by default. Every trial
uses the full FP64 correctness check; rejected candidates remain in JSON and
cannot win. Candidate order rotates across shapes. After selection, the driver
recaptures/JITs and measures the entire winning block/window/thread/copy
configuration again: the published table is that
fresh result, not the search minimum. Failed revalidation is never replaced by
an earlier favorable trial. Search cost (including validation/framework timing)
is separate from warm execution cost. This does not autotune PyTorch or
establish a globally optimal configuration.

On CPU, `--auto-vectorize` opts in to experimental independent-element SIMD
packing. It is off by default: the current heuristic has both wins and
substantial regressions. Packing uses semantic independence and preserves
inner serial/reduction order; it does not infer a new reduction permission or
narrow input precision. Explicit vector execution scopes still work without
this option. `--no-vectorize` disables TIRx vectorization altogether; it does
not disable LLVM's own optimizations. The two flags are mutually exclusive.
Reports record the executable hash and adjacent Tile/bridge shared-library
hashes separately, because rebuilding a dynamic bridge need not change the
executable itself. These hashes are not a complete runtime dependency trace.

Use `--backends metal --execution-scope group` for cooperative threadgroups,
or `--execution-scope worker` for one logical instance per worker. Both use the
same kernel source and block configuration; the requested mapping is recorded
and checked in every native result. `auto` retains the reference worker
mapping. CPU does not yet support `group` and rejects it explicitly. Run each
variant separately into a new output directory; do not change the block shape
when attributing a difference to execution mapping alone.

`--group-threads 128` supplies an exact Metal group worker-count constraint,
separate from `--threads 8` (the host/PyTorch CPU thread setting). Zero, the
default, lets the compiler planner choose. This supports controlled cost-model
ranking tests with the same TileIR and tile shape. The native report records
both the requested constraint and the realized `execution_plans`; the driver
checks their agreement. Repeated comparisons preserve the constraint. An
unsupported count fails rather than being silently clamped. The prior
reference worker launch width of 256 is not a hardware limit on group plans.

`--copy-batch 4` opts into up to four independent reads/computed values before
their shared-memory stores. The default, one, preserves the reference sequence.
This needs Metal group execution and is independent of the matrix capability.
Only supported compiler-owned shared destinations are transformed; external
writes and opaque effects retain their original lowering. Full worker chunks
are batched, while bounded accesses and the remainder remain guarded. There
is no async-copy, vector-alignment, or barrier-elision assumption. Native JSON
records `max_copy_batch` and `batched_copy_operations` per group; replay checks
and preserves the policy. The matrix score does not yet rank copy-batch sizes.

Use `--pipeline-window 1` and `--pipeline-window 2` for a same-binary GEMM
schedule comparison. Each ordinary host configuration captures/JITs a fresh
candidate: window 1 keeps ordered execution, while window 2 permits the safe
software-prefetch planner. The requested window is recorded in native output
and verified by the driver. Keep execution scope, tile shape, inputs, and
timing settings fixed, and run the two modes sequentially. This does not imply
hardware-asynchronous copies; the flag has no effect on non-pipelined operators.

Use `--backends metal --execution-scope group --cooperative-matrix` to opt in
to native FP32 SIMD-group matrices on a compatible device (Apple GPU family
7+, including this M1 Max). The flag is a device-capability assertion, not
automatic GPU detection; the helper supplies `thread_warp_size=32`. It is off
by default. Eligible M/N/K tile extents must be multiples of eight. Other
scopes, types, and shapes keep the reference loops; global ragged edges are
handled by the existing bounded Tile accesses. MMA allows reassociation at
the declared FP32 types, without TF32 or other input-precision reduction.

Compare runs with and without this flag using the same binary, block shape,
pipeline window, execution scope, and timing settings. Native output records
both the requested capability and the actual number of static
`simdgroup_multiply_accumulate` call sites in generated Metal source. The
driver checks that eligible cases select those instructions and fallback
cases do not. A call-site count is not a dynamic instruction counter.

### Metal SIMD-group reductions

`--metal-subgroup-reductions` opts automatic Metal sum, softmax, RMSNorm,
LayerNorm, residual LayerNorm and cross-entropy into the proved FP32
add/max/min collective
realization. It requires the ordinary `simdgroup` TIRx path, automatic root
execution, no cooperative matrix option and only those operations. The option is also explicit
permission to replace the reference FP32 left fold with a tree order; it is
never inferred merely because the target supports SIMD collectives.

The planner enumerates every whole-SIMD-group width per logical row, up to
the target's thread limit and the two-level collective's 32-subgroup bound.
The TVM test/benchmark adapter queries the physical Metal thread limit and
sets both target thread-capacity attributes instead of inheriting TVM's
conservative 256-thread default. This is a device bound; compilation and
dispatch must still respect each generated pipeline's resource constraints.
One-group programs may be packed independently into a wider threadgroup;
multi-group programs use uniform SIMD collectives, small shared partial arrays
and a group barrier. Eligible reused compiler-owned Tiles are compacted to
worker-private stripes only after an affine access/ownership proof. Native JSON
records groups per program, whole-group threads, shared bytes, private stripe
size, reduction counts, barrier sites, queried device limit, physical group
count, scalar rounds, useful lane-work fraction and the separately versioned
`metal_subgroup_reduction_v1` score.

Run the automatic planner with:

```sh
uv run --no-project --python 3.13 --with torch --with numpy python \
  scripts/benchmark/tile_torch/run.py \
  --native cmake-build-tirx/bin/benchmark_tile_tirx \
  --output NEW_EMPTY_DIRECTORY --backends metal \
  --operations sum,softmax,rmsnorm,layernorm,residual_layernorm,cross_entropy \
  --metal-subgroup-reductions \
  --samples 11 --sample-ms 100 --warmup-ms 100 --capture-sources
```

Without explicit packing, `--group-threads 128` is an exact cooperating-worker
constraint for this realization. A measured staged/JIT sweep can instead use
`--tune-group-threads '32,64,128,256'`. Each width is separately captured,
compiled and fully validated; invalid trials remain in JSON. The measured
winner is then captured/JIT-compiled and measured again, so the reported row
is not a search minimum. Preserved logical Tiles count against the default
`max_reduction_striped_scalars_per_worker=64` compiler-state bound after
ownership mapping; over-budget widths fail before code generation. This is
not a claimed hardware-register count. Materialization and thread candidates
may be searched jointly. GEMM block/pipeline and copy-batch tuning are rejected
for this reduction mode.

`--reduction-programs-per-group P` independently fixes program packing.
`P>1` can use one or several SIMD groups per program. An explicit thread count
is the **total** group size `32*S*P`; for example, `--group-threads 512
--reduction-programs-per-group 2` assigns 256 workers to each program, not 512.
With `--group-threads 0`, the solver searches fitting cooperating widths for
that P. Automatic packing (`P=0`) retains its existing single-subgroup packing
family. Cooperating packed tails require a uniform-fence/read-write proof;
unsupported cases and over-budget resources fail before timing. The
[fixed packing experiment](results/m1-max-20260905-cooperating-packing/notes.md)
retains substantial regressions, so this is an explicit candidate, not a
default promotion. `--reduction-unroll U` partially unrolls ordered worker
stripes without changing their accumulation order. Its default remains one;
the legal range is 1..16. Both policies and their realized values are recorded
and checked before accepting an output or replaying a plan.

`--reduction-lane-elements V` selects a blocked-cyclic ownership layout:
`i = (chunk * workers + worker) * V + element`, with `0 <= element < V`.
Legal widths are 1, 2, 4 and 8; the default remains 1. Unlike stripe unrolling,
changing V changes each worker's FP32 recurrence and requires the existing
SIMD-group reduction-tree permission. The same map distributes reductions,
elementwise consumers and compiler-local Tile storage; manual memory is not
silently reinterpreted. Full packs have no tail predicate; only the last
partial pack is guarded. Consecutive elements expose contiguous accesses to
codegen, but do **not** promise hardware vector instructions. Resource bounds
use the maximum live private slots under this layout, not the old V=1 count.

The JIT product can include `--tune-reduction-packing '0,4'` and
`--tune-reduction-unroll '1,4'`, `--tune-reduction-lane-elements '1,2,4'` and
`--tune-reduction-input-caches 'reload,cache'` together with
`--tune-group-threads '0,128'`. Each input-reuse choice is a separate capture
and JIT, not a mutation of a captured function. Zero includes the automatic planner as a
candidate; do not accidentally exclude its packed short-row realization by
searching only exact cooperating widths. Illegal/resource-heavy combinations
remain rejected trials, and the full Cartesian product is budgeted. An
independent `repeat.py` replay—not the search minimum—is the performance
claim. `--row-shapes '1x4096,64x4096,1024x4096,17x257,1024x257'` replaces the
non-GEMM shape set and separates program-count scaling from reduction width.

JIT selection remains host-wall throughput by default. Use the explicit
`--tuning-metric gpu-control` with `--metal-device-timing` to minimize the
no-counter GPU command-buffer throughput instead. It does not fall back to
host time or instrumented compute-pass time if the control is absent. Raw
trials record the objective and score; model regret uses that same objective.
The selected layout is recaptured and measured afresh, then can be independently
replayed with GPU and E2E measurements retained side by side.

Reduction plans also export an availability flag and per-program/per-worker
global/private payload read/write demand. These are conservative IR access
features, not measured DRAM traffic or physical register/spill counts. The
backend policy's optional access coefficients default to zero until calibrated.
See the [access-demand checkpoint](results/m1-max-20260905-access-demand-validation/notes.md)
for the precise scope and tests; the legacy score does not prune joint JIT
resource candidates.

Automatic GPU elementwise programs now also admit a fused program/element
grid. `--element-grid reference` disables this family for same-binary A/B;
`--element-grid auto` permits it, subject to proof and explicit-scope rules.
Plans report `elementwise_elements_per_program`; an enabled option alone is
not evidence that a particular function used the map. Older frozen binaries
receive no new CLI slots unless a new option is requested.

The saved [12-case base report](results/m1-max-20260905-metal-subgroup-reductions/notes.md),
[eight-case LayerNorm/cross-entropy extension](results/m1-max-20260905-metal-subgroup-row-extensions/notes.md),
[balanced RMSNorm A/B](results/m1-max-20260905-metal-subgroup-rmsnorm-replay/notes.md)
and [balanced extension A/B](results/m1-max-20260905-metal-subgroup-row-extensions-replay/notes.md),
plus the [residual-LayerNorm materialization search](results/m1-max-20260905-residual-layernorm-materialization-search/notes.md),
[balanced materialization A/B](results/m1-max-20260905-residual-layernorm-materialization-replay/notes.md),
[bounded exact-width search](results/m1-max-20260905-residual-layernorm-bounded-thread-search/notes.md)
and [CPU target counterexample](results/m1-max-20260905-cpu-residual-layernorm-materialization-search/notes.md)
document exact plans, commands, hashes, complete errors and interpretation
limits. They are an M1 Max FP32 cohort, not all-device or production-LLM parity.

Measurement contract:

- Identical deterministic contiguous FP32 inputs; full outputs checked against
  a CPU FP64 reference with the same tolerances for both implementations.
- CPU thread environment is set before importing either framework. MPS or
  native Metal unavailability is an error, never a CPU fallback.
- Inputs and native outputs are allocated before warm measurements. PyTorch
  uses eager `out=` operations under inference mode where the operator exposes
  one. The functional RMSNorm, LayerNorm, residual LayerNorm and cross-entropy
  calls used here return new outputs, so their output allocation remains inside
  warm timing.
  Every row records either `output_policy=framework_return_value` or
  `preallocated_out`.
- Capture, native compilation, allocation/upload, first invocation, and
  download are reported separately. First-call timings are not a claim of an
  empty OS/driver cache, and PyTorch's process is reused across cases.
- Warmup, calibrated batch duration and sample count are explicit command-line
  inputs and recorded in JSON. The original saved 20-case subgroup cohorts use
  100 ms warmup and 11 × 100 ms samples. The four residual-LayerNorm search
  rows use 100 ms warmup and 9 × 60 ms samples. Four-round causal replays use
  80 ms warmup and 7 or 9 × 60 ms samples. Native/PyTorch order alternates.
- Warm measurements use a host clock around dispatch plus synchronization.
  They exclude transfers but include C++/Python binding and launch overhead;
  they are **not pure GPU-event kernel timings**. Do not run alongside builds.
- The default native schedule is a reference realization. Explicit Metal
  `group` partitions independent elements across workers and shares
  group-owned compiler temporaries. The opt-in matrix selector replaces only
  proven compatible MMA bodies; hardware-asynchronous transfers and parallel
  reduction trees remain separate work. A correctness pass must not be
  described as competitive performance.

`results.json` contains raw samples, numerical errors, setup phases, compiler
and hardware information, thread settings, the binary hash, and source
revision. `results.md` is the readable comparison. Failed cases are retained
and cause a nonzero exit code; no speed ratio is published for an invalid case.

## Five-path Tile lowering comparison

### Separate GPU and dispatch measurements

`run.py --metal-device-timing <libluisa-benchmark-metal-timing.dylib>` adds
device-side compute-pass timestamps to both TIRx and eager PyTorch, and to the
optional direct MPS GEMM baseline. `repeat.py` accepts the same option for
counterbalanced replay. Complete the full build and run
`ctest --test-dir cmake-build-tirx -R '^test_tile_metal_timing$'` first.
For example:

```sh
uv run --no-project --python 3.13 --with torch --with numpy python \
  scripts/benchmark/tile_torch/run.py \
  --native cmake-build-tirx/bin/benchmark_tile_tirx \
  --backends metal --operations add,gelu_add --input-views \
  --metal-device-timing cmake-build-tirx/bin/libluisa-benchmark-metal-timing.dylib \
  --output /tmp/tile-dual-timing
```

The ordinary `throughput_us` and `latency_us` remain **uninstrumented host
wall time**: allocation/transfer/JIT excluded, dispatch/binding and final
synchronization included. Optional `device_timing` is a **separate phase**
using the actual framework dispatches and resources, not generated-source
replay. It records sum of GPU compute-pass intervals, their first-to-last
span, command-buffer GPU duration, pass/buffer counts, raw clock calibration,
and its own repetition count (at most 64). The JSON preserves every sample;
the report shows both GPU execution and end-to-end dispatch.

The current helper ABI is **2**. Every device phase also measures a matched
**command-buffer control with no encoder hooks or GPU counter attachments**.
It reads completed `GPUStartTime`/`GPUEndTime` through a commit-only observer.
Probe/control order alternates at the same batch size. The primary GPU table
uses this control; the compute-pass table is diagnostic. The control includes
GPU work and gaps inside each command buffer, including any blits, so it is
not an individual-kernel timestamp. Both scopes exclude host API time itself,
but still reflect GPU scheduling/contention. Neither guarantees an idle device.

This distinction is necessary: in the M1 Max integration probe, eager Torch
softmax's instrumented/control GPU command-buffer ratio was 3.08–5.85×. The
probe path is therefore **not transparent for performance**, even though it
preserves numerical output and encoder boundaries. Raw paired ratios diagnose
observer effects; do not use them to subtract or normalize counter time into
an imagined uninstrumented kernel time. The current diagnostic cohort also
has background-load variability and is not a stable speed ranking.

This M1 Max supports stage-boundary counters, not per-dispatch counters.
A single-kernel pass gives that kernel's device interval. A multi-dispatch
pass gives batch time divided by repetitions; a multi-kernel eager operator
is still an operator sequence. Neither is just arithmetic-instruction time.
GPU pass intervals include dispatch/barrier work inside the pass; intervals
between passes are visible in span/command-buffer diagnostics, not hidden in
compute time. Do not subtract independently sampled GPU and host medians to
estimate CPU overhead. The tuning winner remains the explicitly documented
host-wall metric; collecting counters does not silently change that objective.

The benchmark-only library temporarily intercepts public Metal command-buffer
encoder factories and commit for counter samples, and **only commit for the
control**, on the concrete class discovered from the
system-default device. It preserves dispatch type, existing counter
attachments and encoder boundaries, then restores the methods on success or
failure. No Runtime, private TVM/PyTorch ABI, Python source export, or TVMx
version change is involved. Only the isolated single-device benchmark
process is supported. Unsupported counters, incomplete work, empty capture
and sample-capacity overflow fail closed. Do not combine it with another
profiler or concurrent workloads. Metal 4 command-buffer/counter-heap timing
is a different API and is not covered by this Metal command-buffer helper.

Native `benchmark_tile_tirx` (TVM or Luisa Runtime), `benchmark_tile_native`
and `benchmark_tile_system` can also load the helper through
`LUISA_TILE_BENCH_METAL_TIMING=/absolute/path/to/library.dylib`.
CPU-only runs cannot request this device timer; their synchronous execution
still includes runtime dispatch and is not relabeled as a pure kernel.

`gelu_add` is tanh-approximate `GELU(A+B)`. Native retains one shared Tile SSA
producer; eager Torch uses **two operations with preallocated output and add
scratch**. This compares a fused graph against eager, not Torch compilation.
`--input-views` enables proved immutable forwarding on CPU or Metal,
independently of `--element-grid auto|reference`. The older
`--cpu-input-views` option remains CPU-only. Plans expose the actual scalar
producer count as `elementwise_scalar_temporaries`.

`sigmoid_pair` and `gelu_pair` measure multi-output graphs: each returns the
activation and the derivative of its forward formula in separate native
buffers. GELU uses the tanh approximation. Both native outputs are joined
**after timing** and checked completely against FP64, in value/derivative
order. Torch preallocates both outputs: sigmoid uses `sigmoid.out`, `sub.out`,
`mul.out`; GELU uses forward and backward out calls with a preallocated unit
gradient. These compare fused graphs with eager sequences, not compiled
Torch, and do not include a full training/backpropagation step.
The existing fixed element block is 1×256; `repeat.py` can replay these graphs
without a parameter search. `--element-grid reference` preserves their serial
program-per-worker control, while the automatic mapper admits same-domain
multi-output graphs only after the ownership and effect checks documented in
the [lowering reference](../../../docs/source/internals/tile/lowering.md#automatic-gpu-pointwise-graphs).

### Runtime routes

`benchmark_tile_native` is the actual TileIR→Metal-backend→MPP route, launched
through Luisa Runtime. It is distinct from the handwritten `benchmark_tile_mpp`
probe. `benchmark_tile_tirx` remains the native C++ TVM bridge comparison;
neither benchmark calls MPS as a fallback. Build with both Metal and the TIRx
bridge enabled, complete the **full build**, and run correctness tests before
timing. Do not build or run other GPU work during measurement.

```sh
uv run --no-project --python 3.13 --with numpy --with torch python \
  scripts/benchmark/tile_torch/compare_lowerings.py \
  --native cmake-build-tirx/bin/benchmark_tile_native \
  --tirx cmake-build-tirx/bin/benchmark_tile_tirx \
  --mpp cmake-build-tirx/bin/benchmark_tile_mpp \
  --mps cmake-build-tirx/bin/benchmark_tile_system \
  --mpp-plan scripts/benchmark/tile_torch/results/m1-max-20260904-mpp-search/results.json \
  --tirx-plan scripts/benchmark/tile_torch/results/m1-max-20260904-joint-search/results.json \
  --rounds 10 --samples 7 --sample-ms 30 --warmup-ms 200 \
  --output /tmp/tile-five-path-replay
```

The plans supply configurations only; all binaries recapture/recompile and
measure afresh. Without plan arguments, explicit fixed defaults are recorded,
not described as autotuned winners. The default eight shapes cover small and
large squares, rectangles and tails; `--shape MxNxK` is repeatable. Full FP64
validation uses identical deterministic inputs and preallocated outputs.
Ten rounds balance positions and pairwise precedence. Fewer rounds are allowed
only as clearly labeled unbalanced smoke/exploration runs. All five paths run
even when another fails, and failures cause a nonzero exit code.

Tables compare **host-wall** batched throughput for all five paths, including
their different Runtime/API overheads; they do not claim pure GPU speedups.
MPS/handwritten-MPP GPU intervals remain in raw JSON only. Native/hand-MPP
ratios use matched descriptors/cohorts but include different host runtimes.
The current native subset is FP32, dynamic K, inline tensors and cooperative
output; an incompatible MPP plan is rejected rather than silently changed.

FP32 here specifies tensor types, not identical compiler math policies or
bitwise reproducibility. Native/handwritten MPP explicitly disable fast math
and relaxed precision; MPS uses `MPSKernelOptionsNone`, and Torch uses its
recorded default MPS route. This driver does not independently override or
validate the external TVMx Metal runtime's MSL fast-math flag. The deterministic
benchmark inputs are dyadic and cannot alone certify multiplication precision;
the separate kernel tests use non-dyadic inputs, transposes, tails and changed
buffers. Keep those tests and policy checks alongside performance validation.

### Same-source TIRx Runtime controls

The Metal backend now accepts `tile::Lowering::TIRX` through the same Tile
factory and ordinary Runtime shader/stream path as native MPP. The standalone
`benchmark_tile_tirx` accepts a final `tvm` (default), `luisa`, or `luisa-fast`
argument. The latter two compile unchanged TIRx-generated Metal source with
fast math explicitly off/on. They share the same `capture()` function and
planner options as the TVM runtime benchmark. No LLVM host wrapper is JITed
for the Luisa path. Only a single static launch with FP32 buffer arguments is
currently supported; unsupported host work is rejected rather than dropped.

Add `--tirx-runtime-controls` to `compare_lowerings.py` to measure all seven
paths. Its default becomes **14 rounds**, balancing positions and pairwise
precedence across seven implementations. Every TIRx path dumps its generated
source outside timing; the runner stores content-addressed sources and requires
identical SHA-256 plus matching threadgroup widths before accepting a Runtime
comparison. Whole-output FP64 validation still applies independently to every
path. Compiler language/resource options and submission APIs may differ, so
an equal source hash does not imply identical binaries or isolate launch cost.

```sh
uv run --no-project --python 3.13 --with torch --with numpy python \
  scripts/benchmark/tile_torch/compare_lowerings.py \
  --native cmake-build-tirx/bin/benchmark_tile_native \
  --tirx cmake-build-tirx/bin/benchmark_tile_tirx \
  --mpp cmake-build-tirx/bin/benchmark_tile_mpp \
  --mps cmake-build-tirx/bin/benchmark_tile_system \
  --mpp-plan scripts/benchmark/tile_torch/results/m1-max-20260904-mpp-search/results.json \
  --tirx-plan scripts/benchmark/tile_torch/results/m1-max-20260904-joint-search/results.json \
  --tirx-runtime-controls --rounds 14 \
  --output /tmp/tile-runtime-comparison
```

### Independent TIRx MPP codegen

`--tirx-mpp` adds a sixth path through TVM's own Metal MPP code generator,
beside native MPP, default TIRx SIMD-group, handwritten MPP, MPS and Torch.
It does not call the native emitter. Build the optional
[pinned C++ TVM extension](../../../src/tile/bridge/tirx/patches/README.md)
in isolation; normal TIRx builds need no patch. Complete the full build and
correctness checks before comparing. The standalone matrix-mode argument
`mpp` explicitly selects this realization; an unavailable extension is an
error, never a fallback.

Use the command above with the isolated binaries, `--tirx-mpp --rounds 12`,
and repeat `--compiler-artifact /path/to/library` for `libtvm_compiler.dylib`,
`libtvm_runtime.dylib`, `libtvm_runtime_metal.dylib` and `libtvm_ffi.dylib`.
Those externally linked libraries are hashed before and after the run.
The two TIRx paths use the same frozen geometry. The additional path must
report actual generated MPP call sites, zero SIMD-group MMA call sites, and
`cost_basis=metal_mpp_memory_v2`. That separately versioned relative-work score
uses MPP memory/shape features, subgroup critical-path work and an outer
program-wave prior. It is not an internal instruction count, measured register
or occupancy value, or calibrated time prediction. Do not merge its timings
into the SIMD-group cost model. With both optional comparisons enabled, eight
paths require sixteen rounds instead.

The [2026-09-04 six-path replay](results/m1-max-20260904-tirx-mpp/notes.md)
records 576 valid complete outputs, library/source fingerprints, the retained
1024³ regression versus original TIRx, and the remaining native/MPS gap.

### Read-only snapshot forwarding before resource planning

The optional TIRx MPP extension also supports a separately measured read-only
view-forwarding family. Use `run.py --matrix-realization mpp-views` together with
`--cooperative-matrix --execution-scope group --backends metal --operations gemm`
and explicit tuning lists. Every candidate is captured/JIT-compiled again and
validated; rejected capacities or padded full-K candidates remain in the report.
Forwarding requires proved immutable, noalias, bounded input snapshots. Manual
memory remains explicit. The `metal_mpp_memory_v2` basis remains an analytic
prior, not a calibrated time model; Staged/JIT measurements choose the frozen
schedule.

For the independent comparison, add `--tirx-mpp --tirx-view-plan SEARCH/results.json`
to `compare_lowerings.py`, with the ordinary `--tirx-plan` and external
`--compiler-artifact` fingerprints. This retains original TIRx, staged TIRx MPP,
native MPP, handwritten MPP, MPS and Torch and adds forwarding MPP as a seventh
path. Fourteen rounds balance order. Both requested forwarding policy and actual
MPP calls are checked. The view schedule is recorded separately; a geometry
change is not presented as a same-geometry lowering-only speedup.

The [seven-path replay](results/m1-max-20260904-tirx-views/notes.md) records
784 valid complete outputs. On this M1 Max, 512³ measured 43.523 µs versus
Torch's 48.794 µs; 1024³ measured 291.736 µs versus Torch's 291.133 µs and
MPS's 278.687 µs. These are synchronized host-wall batched times, not GPU
kernel durations. The large-GEMM library gap is reduced, not closed.

The newer [MPP cost-model v2 search](results/m1-max-20260905-mpp-cost-v2-search/notes.md)
separates subgroup critical-path work from whole-device waves. On the same
finite calibration cohort it reduces mean/median/max model regret from
74.18/43.05/239.58% to 8.82/2.59/34.37%; this is explicitly not a held-out
result. Its [independent 14-round replay](results/m1-max-20260905-mpp-cost-v2-replay/notes.md)
validates all 784 outputs and beats Torch and MPS on all eight tested FP32
GEMMs. At 1024³, TIRx MPP views measure 270.675 µs versus MPS at 272.572 µs
and Torch at 284.654 µs; the paired ratios are 0.9938× and 0.9513×. This does
not establish other-device, low-precision or non-GEMM parity.

Subgroup-fence elision is a separate, default-off tuning choice, not an
automatic consequence of forwarding. `run.py --elide-independent-subgroup-barriers`
selects it for `mpp-views`; `repeat.py --candidate-subgroup-fences elide` and
`compare_lowerings.py --tirx-view-subgroup-fences elide` select it explicitly
for their candidate/view path. Frozen reports preserve the policy. The native
benchmark records both the requested policy and each group's independently
proved `independent_subgroups` fact. Unsafe groups retain their fences even
when elision is requested. The first four-round A/B found a **512³ regression**,
so fewer barriers must not be assumed profitable or enabled by a shape table.

For compiler implementation A/B tests, `repeat.py --capture-sources` requires
both binaries to dump generated Metal, archives content-addressed sources and
records each command. The runner fingerprints both executables, every adjacent
runtime library and any explicit `--compiler-artifact` before and after timing;
changing an artifact fails the run. No JIT/compiler boundary uses Python source
generation: these Python programs only orchestrate benchmarks and oracles.

## Handwritten MPP versus MPS

`benchmark_tile_mpp` is a standalone hand-written Metal Performance Primitives
GEMM probe, not a TileIR lowering and not an MPS library fallback. It needs a
macOS 26 SDK/runtime and supported Apple GPU. Its default `tensor_inline`
representation is an ordinary buffer plus extents/strides; it uses the classic
tracked command queue, like `benchmark_tile_system`. An optional tensor-handle
mode uses Metal 4 with explicit dispatch barriers and commit-feedback checks.

The configuration is
`tile_m,tile_n,operation_simdgroups,cooperative_output,static_k,inline_tensors[,group_simdgroups,cohort_rows]`.
One operation can use a single subgroup or the whole group. Independent
single-subgroup operations can form a spatial cohort; the group need not have
the same size as the operation. Static M/N slices are used only inside the
matrix. Dynamic slices retain bounds for ragged output and K tails. Static K
requires a multiple of 16 for this FP32 family.

For a standalone physical-traversal experiment, the **binary CLI** additionally
accepts `walk-rows [walk-columns]` after `group-simdgroups cohort-rows`.
Omitting them preserves the original 2D launch. `1 1` uses a linear row-major
launch; positive row/column spans enumerate row-major inside bounded program
rectangles. These spans count **programs**, not tensor elements. Partial
rectangles use actual dimensions, preserving launch count and ownership;
uint32 launch/product overflow is rejected before tensor allocation.
`LUISA_TILE_BENCH_DUMP_SOURCE` optionally archives the exact MSL passed to JIT.

This option is not part of `compare_mpp.py`'s selected-config format or a
production planner default. Its metadata validator rejects non-default walks
unless explicitly requested by a diagnostic caller. The
[predeclared walk screen](results/m1-max-20260906-mpp-grid-walk/notes.md) and
[K-partition screen](results/m1-max-20260906-mpp-k-partition/notes.md) retain
negative results and unstable controls; do not promote their minima.

After a full build, screen a small family, then freeze and replay it:

```sh
uv run --no-project --python 3.13 --with numpy python \
  scripts/benchmark/tile_torch/compare_mpp.py \
  --mpp cmake-build-tirx/bin/benchmark_tile_mpp \
  --mps cmake-build-tirx/bin/benchmark_tile_system \
  --config 64,64,4,1,0,1 \
  --config 32,32,1,1,0,1,4,4 \
  --config 16,16,1,1,0,1 \
  --output /tmp/mpp-search

uv run --no-project --python 3.13 --with numpy python \
  scripts/benchmark/tile_torch/compare_mpp.py \
  --mpp cmake-build-tirx/bin/benchmark_tile_mpp \
  --mps cmake-build-tirx/bin/benchmark_tile_system \
  --plan /tmp/mpp-search/results.json \
  --rounds 6 --samples 9 --sample-ms 40 --warmup-ms 200 \
  --output /tmp/mpp-replay
```

This runner permits only FP32, no fast math or relaxed precision, and checks
every full output against the same FP64 oracle as the Tile comparisons. Search
selects by GPU batch time; its minimum is not a validated speedup. Replay
executes all six orders of MPS/default-MPP/selected-MPP for every shape, rotates
shape order, verifies binary stability and records all failures. Keep binaries
unchanged and do not build or run other GPU work during either campaign.

Both probes also report `gpu_throughput_us` and `gpu_latency_us`, in addition
to host wall times. These are Metal command-buffer GPU intervals; they include
dispatch/barrier work and are not individual arithmetic-instruction timings.
The existing TileIR/PyTorch timers remain host-wall timers.

The [M1 Max experiment](results/m1-max-20260904-mpp.md) records the candidate
family, numerical checks, frozen replay and remaining gap. These samples can
calibrate MPP-specific ranking only after a matching emitter exists; they must
not be pasted into the native 8x8 atom model as interchangeable issue costs.

## Repeat and profile

Freeze schedules from two reports and repeat them without selecting new winners:

```sh
uv run --no-project --python 3.13 --with torch --with numpy python \
  scripts/benchmark/tile_torch/repeat.py \
  --native cmake-build-tirx/bin/benchmark_tile_tirx \
  --reference /path/to/reference/results.json \
  --candidate /path/to/selected/results.json \
  --output /tmp/tile-schedule-repeats --rounds 4
```

Each case keeps the recorded block, pipeline, mapping, and matrix/vectorization
policies. The even number of rounds balances schedule order and framework order
separately, rotates shape order, and validates every fresh capture/JIT result.
Reports missing the explicit `auto_vectorize` policy cannot be replayed because
the old implicit behavior is ambiguous. `--candidate-vector-mode auto-vectorize`
can isolate packing against the same report passed as both inputs. Use
`--operations gemm,add,sum,softmax` when both reports contain those cases.
The paired min–max ratios are observed ranges, **not confidence intervals**.
Use `--candidate-native` for a second, separately prebuilt executable when
comparing implementation changes at frozen schedules. The report fingerprints
both executables and their adjacent Tile libraries; verify their dynamic-loader
paths before relying on a copied baseline. A build must never run during replay.

`profile_torch.py` provides a long, checked, preallocated eager GEMM workload
for an external profiler. `--backend cpu` uses the actual installed CPU build;
`--backend metal --signposts` enables MPS signposts without per-dispatch waiting.
`--mps-path metal` explicitly selects PyTorch's alternative Metal implementation;
it is not the default-MPS baseline. `--capture-dir /tmp/new-capture-directory`
captures one warmed invocation for launch/resource inspection. PyTorch chooses
the filename inside that new directory; the script verifies it exists and
records the actual path. Profiler timings are not
mixed into uninstrumented benchmark results. Record the exact PyTorch version
and build commit before studying its dispatch heuristics.

For the native executable, set `LUISA_TILE_BENCH_DUMP_SOURCE` to a **new file**
to inspect generated LLVM IR (CPU) or MSL (Metal). Dumping is outside all timed
phases, and never substitutes an LLVM host wrapper for Metal device code. The
normal positional arguments are unchanged. External profilers can launch the
same executable with longer `sample-ms` / `warmup-ms` arguments. Complete the
full build first, and do not profile while compiling or running other tests.
Filter exported trace rows to the target process: system GPU tracks can include
other applications even for a process-targeted recording. Do not publish those
unfiltered traces with benchmark reports. Encoder intervals may contain many
dispatches; they are not automatically per-kernel GPU durations.

## Native matrix measurements

The [M1 Max matrix comparison, 2026-09-03](results/m1-max-20260903-matrix.md)
records 24 validated pairs at `37e1337fd`: eight Metal cases with/without
matrix selection at identical 8×8×16 tiles and a two-window pipeline, plus
eight fresh CPU cases. Metal 1024³ improved from 4.618 to 3.244 ms (1.42×),
but remained 9.41× slower than PyTorch. CPU measured 12.425 versus 1.103 ms.
Raw generated-instruction counts distinguish actual matrix selection from
permission to select it. No autotuning was performed and the default is unchanged.

## Pipeline-window measurements

The [M1 Max pipeline comparison, 2026-09-03](results/m1-max-20260903-pipeline.md)
records all 48 GEMM pairs at `aeeee8eaf`: the same binary and 8×8×16 block
shape, two pipeline windows, and CPU worker / Metal worker / Metal group.
All numerical checks passed. CPU 1024³ improved from 37.870 to 12.577 ms
(3.01×), but Metal group moved from 4.334 to 4.439 ms and remained about
12.8× slower than PyTorch. The report includes every shape, raw samples,
compilation costs, and measurement limits; legality alone does not establish
pipeline profitability.

## Cooperative execution measurements

The [M1 Max execution comparison, 2026-09-03](results/m1-max-20260903-execution.md)
records all final worker/group cases at `2f8f2a044`, plus the earlier
experimental results and a CPU GEMM repeat. At identical 8×8×16 tile shape,
1024³ Metal GEMM improved from 14.725 ms worker to 4.275 ms group (3.44×), but
remained about 12× slower than PyTorch. The CPU loop-flattening regression was
substantially reduced by retaining axes; the largest CPU case still measured
48.226/38.508 ms in two runs, above the historical 31.571 ms observation.
This remaining performance issue is not considered closed. All numerical
checks passed, and no default mapping was changed based on these timings.

## Recorded reference baseline

The [M1 Max baseline, 2026-09-03](results/m1-max-20260903-reference/results.md)
contains all 40 CPU/Metal cases, with [raw samples and metadata](results/m1-max-20260903-reference/results.json).
The implementation is commit `8591c599d`; all 40 pairs passed the shared
numerical checks. This is one machine/run and one untuned native configuration,
not a hardware limit or a best-of-tuning result.

For FP32 1024³ GEMM, native/PyTorch warm batched times were **31.571/1.192 ms
on CPU** and **14.876/0.353 ms on Metal**: native was about 26.5× and 42.1×
slower, respectively. These measurements include host dispatch and predate the
explicit cooperative-group mapping. They are retained as a historical
reference, not a measurement of the current cooperative implementation.
## Experimental reduction service policy

`--reduction-cost-profile service-v1,C,D,R,K,G,W,P` injects a typed,
backend-owned whole-launch objective into the Metal reduction solver. `C`
is a positive subgroup saturation prior; the six finite nonnegative
coefficients price dispatch, scalar rounds, collectives, per-program global
bytes, per-worker global bytes and per-worker private bytes. These are fitted
score units, not queried occupancy or measured DRAM/register traffic.
`analytic` remains the default, and an explicit profile with unavailable
payload features fails instead of falling back to different score units.

`--tuning-metric model` selects from separately captured/JIT-compiled
configurations by summed whole-kernel score only. Trials still execute for
correctness and diagnostic timing, and the winner is freshly captured again;
this does not claim compile-only search or use the trial times for selection.
Combine it with `--tune-reduction-input-caches reload,cache` to choose resource
reuse while the C++ solver independently chooses each legal execution width.
The exact profile string and generated objective are checked and preserved by
`repeat.py` for frozen replay. See the
[service-policy experiment](results/m1-max-20260905-service-policy-validation/notes.md)
for the calibration data, limitations and predeclared held-out protocol.

## Windows CUDA Tile and torch.compile matrix

`cuda_matrix.py` runs the actual `benchmark_tile_workloads` executable and
`cuda_torch_baseline.py` serially. Build the executable first, then supply the
completed-build JSON marker and the Python executable containing CUDA PyTorch
and Triton. The runner does not build or install dependencies. Choose an
affinity mask for the machine; the runner verifies its CPU topology and applies
the mask to each owned child before its primary thread starts. Timeouts close a
Windows job that owns that child and its descendants.

```powershell
python scripts/benchmark/tile_torch/cuda_matrix.py --suite smoke --list-cases
python scripts/benchmark/tile_torch/cuda_matrix.py `
  --build-dir build-msvc-llvm --build-marker <completed-build.json> `
  --torch-python <torch-environment/python.exe> --output <new-output-directory> `
  --affinity-mask <machine-specific-mask> --threads 4 --routes native,tirx `
  --path-prefix <CUDA-bin-directory> --path-prefix <TVM-DLL-directory>
```

The smoke suite has 17 FP32 cases covering row operators, reductions, inclusive
tree scan, ranking, GEMM/GEMV and attention, including tails and adversarial
inputs. `--suite broad` adds decode/prefill widths 512–8192 and bounded ranking
and attention cases. `--cases id1,id2` selects exact built-in IDs; alternatively,
pass a JSON file with `schema: 1` and a `cases` array in the format printed by
`--list-cases`. Repeat `--case-filter` to select ID globs. `--routes native` is
useful for an initial CUDA Tile IR comparison; `tirx` and optional `simd` remain
separate routes. Each run requires a fresh output directory.

`--cases scripts/benchmark/tile_torch/cuda_llm_primitives.json` selects the
78-case CUDA tuning matrix: GEMM/GEMV, decode/prefill attention, normalization,
activations, RoPE, reductions, scan, sort and Top-K, with FP32/FP16/BF16 and
small, wider and non-aligned dimensions. It includes explicit schedule,
fast-math and ranking-algorithm variants; 78 cases are not 78 distinct shapes.
Use `--graph-batch 100 --samples 7 --sample-ms 100 --warmup-ms 500` for the
adaptive graph protocol below. Availability and performance are measured by
the runner, not guaranteed by inclusion in this inventory.

`--precisions fp32,fp16,bf16` expands selected cases into separate precision
variants; a case JSON can also set `precision` directly. Inputs and outputs use
their actual declared storage dtype, including two-byte FP16/BF16 packets, and
each precision has its own exported FP64 oracle and bound. Ranking uses a
power-of-two padded input tile while storing only the requested logical output.
The optional JSON operation `scan_ordered` retains the sequential fold-left
reference; `scan` permits tree reassociation. The manifest records this policy,
so results with different contribution-order contracts should not be pooled.
Historical output directories are never rewritten when these algorithms change.

Defaults are seven 100 ms stream samples after 500 ms warmup, graph batches of
32, full-graph Inductor `max-autotune` and one compilation worker. Native compile
time and Torch's compile-plus-first-call time are separate cold measurements.
Synchronized host wall, native CUDA event stream spans and graph event spans
are recorded separately; event spans can include scheduling or submission gaps.
Native graphs reuse a preallocated output with write dependencies. Torch uses
normal functional returns and its capture memory pool. These are observable
allocation differences, so the runner does not infer isolated kernel speedups.

Starting with the 2026-10-01 phase3 measurements, both native and Torch graphs
receive a separate synchronized replay warmup for the configured `--warmup-ms`
before graph samples. Torch records `warmup_target_ms`, `warmup_actual_ms` and
`warmup_replays` inside each compiled/eager graph result; this excludes capture,
the first replay and measured samples. Earlier phase2 results retain their
original protocol: native warmed graph replay for that duration while Torch
only replayed once after its stream warmup. Historical result files remain
unchanged and should not be treated as having the new symmetric graph warmup.

Graph timing protocol `adaptive_replay_span_v2` first records both CUDA events
around the existing untimed first replay, then performs the configured graph
warmup. Calibration starts at one whole graph replay and permits four measured
attempts. If its event span is below 80% of `--sample-ms`, the next replay count
is the ceiling of the measured count scaled to the target, bounded by 65,536
replays and 10,000,000 logical operations per sample. The last measured count R
is fixed for every formal sample, even when a cap or four attempts prevented
reaching the target. Seven samples are retained without discarding outliers.

Each formal sample now executes R complete graphs of N logical operations and
normalizes both event and synchronized host spans by N*R. Results retain R,
calibration attempts, target/caps, actual full spans, and the existing per-op
fields. Native multi-stage workloads additionally count stage dispatches;
Torch's N counts logical calls, not generated kernels. This targets 100 ms per
sample by default and still allows submission gaps. Older phase3/candidate44
records used one replay per sample and remain unchanged; compare protocols
explicitly rather than treating their sample durations as equivalent.

Every native route exports the same seeded inputs and FP64 oracle with
per-element error bounds. Their file hashes must match before Torch reads the
packet. Unsupported native cases still run Torch when a complete packet exists.
The default ranking comparison uses ordinary Torch sort/topk, allowing tied
indices to differ while checking values, source correspondence and uniqueness;
`--ranking-contract stable` instead requires the native stable-index contract.
`--eager` adds a clearly separate eager baseline. No result is implied by the
presence of an operation in the suite.

Results checkpoint after every child and case, retaining stdout/stderr, raw
samples, input hashes, executable/source/library hashes, build marker, affinity
readback and `nvidia-smi` telemetry. Exit 0 means all requested cases passed;
exit 3 means completed with explicit unsupported routes; exit 1 means failure.

`test_tile_primitives host` checks capture and reduction-policy preservation
without a device. `test_tile_primitives simd native` and
`test_tile_primitives cuda tirx` exercise those runtime routes; for
`test_tile_primitives cuda native`, set `LUISA_CUDA_TILE_IR=1`. The six test
groups cover stable sorting/prefix selection, signed-zero bit preservation,
interior axes, guarded buffer views, sequential scan bitwise order and a
separate tree-scan FP64 error bound. Direct native top-k with a seven-element
Tile result may report its explicit power-of-two limitation; the padded
benchmark stores a logical seven-element prefix from a full sorted Tile.

`attention` retains FP32 probabilities and PV accumulation. The separate
`attention_tensorcore` operation accepts FP16/BF16 and explicitly rounds the
unnormalized probabilities to that storage type before PV MMA, while retaining
the FP32 denominator and accumulator. Its packet exports both the original
strict bound and a predeclared probability-rounding allowance weighted by
absolute V values; both implementations use the same primary bound. Strict
failures remain visible in a separate result field. Automatic PyTorch SDPA is
recorded through generated-call evidence; passing this common numerical
envelope does not prove identical internal arithmetic. Failed first-call
validation retains the actual typed outputs and compiler evidence.

Native CUDA integer sum/scan uses unsigned arithmetic for signed inputs to
preserve fixed-width wrapping, with signed comparisons retained for min/max.
The native conformance suite checks 32/64-bit signed and unsigned values,
including overflow and values beyond exact FP32/FP64 integer representation.
These conformance tests are separate from the floating-point workload matrix;
they do not imply integer PyTorch performance measurements.

Native CUDA MMA with `allow_reassociation=true` uses the CUDA Tile typed MMA
operation for FP32 as well as FP16/BF16 inputs and FP32 accumulation. FP32
inputs remain FP32; no TF32 conversion or reduced-precision permission is
introduced. This delegates layout and contraction scheduling to the Tile
compiler, without promising tensor-core instructions for FP32. The ordered
policy retains ascending-K FP32 FMA with preserved subnormals. Its constant
expansion is bounded by contraction length, output size and total expanded
work; larger contractions retain the dynamic loop. Conformance covers typed
precision boundaries, ordered results, and both sides of the expansion budget.


### Optional repeated-extrema Top-K schedule

`benchmark_tile_workloads` keeps full bitonic sorting plus prefix output as the
ranking default. For `topk` only, append `--ranking-algorithm repeated_extrema`
to explicitly select a serial sequence of FP32 maximum and INT32 minimum trees,
one row per program. Two singleton loop carries hold the previous winner's
FP32 value and INT32 original index. Each iteration considers valid elements
strictly after that key in descending-value/ascending-index order; the initial
key (+infinity, -1) includes real positive infinity. A point load retrieves the
winner's original storage bits without a dynamic extraction from the whole
register Tile. The input row remains loop-invariant. Input/output views must
be disjoint, as allocated by the workload runner and unit tests; this helper
does not provide in-place sorting. The same logical Top-K,
stable original-index tie ordering, and FP16/BF16/FP32 storage contracts remain.
NaNs are excluded and rejected before upload. The conformance test also checks
signed-zero value bits, infinities, all-negative-infinity rows, masked tails
and complete selection of all 33 values from a 64-wide padded Tile.
This optional schedule is not an automatic dtype/shape tuner or a speed claim.

A matrix JSON ranking case can set `"ranking_algorithm": "repeated_extrema"`;
`"full_sort_prefix"` is the default. Use distinct IDs to compare both algorithms
with identical dimensions, dtype, tile, seed and pattern, for example logical
`[4,65,7]` with tile `[1,128,1]` or `[4,4096,16]` with tile `[1,4096,1]`.
K7 needs no seven-wide native Tile: the repeated algorithm writes one value/index
per iteration. The manifest and result record the requested algorithm, and the
manifest labels the realization `stable_repeated_extrema`. Matrix validation
checks that label and includes the helper source in build/provenance hashes.

Torch still executes the same full logical `topk` expression over the entire
input row: ordinary `torch.topk` for the standard tie contract, or stable sort
plus prefix for `--ranking-contract stable`. Native's internal K iterations are
one shader invocation, so a graph batch of N still means N complete Top-K
operations, not N individual maxima. All stores and reductions are included in
native graph and host timing. Keep cold compile costs, symmetric graph warmup,
full output/index checks and the functional Torch versus preallocated native
output distinction. Compare same-window baselines; O(K*N) work and dependent
reductions do not guarantee an improvement over bitonic sorting.


Native CUDA Tile accepts the existing public fourth `ShaderOption` argument:
`tile::compile(device, kernel, {}, {.enable_fast_math = true})`. Its default
remains strict. The workload executable exposes `--fast-math 0|1`; a matrix
case can carry `"fast_math": true` (use distinct case IDs for strict/fast).
The flag is recorded and checked in both manifest and native results. The
input/oracle files, numerical acceptance bounds and Torch program are unchanged.
Fast results must pass the same declared oracle; the flag does not enlarge it.

The native policy `elementwise-fp32-approx-ftz-rsqrt-v2` relaxes FP32 elementwise
add/sub/mul (RNE plus FTZ), division/sqrt (approximate plus FTZ), and exp
(approximate, with no FTZ guarantee). Tanh retains full precision because the
SDK's approximate mode failed the unchanged elementwise error budget.
MMA precision/order, reduction/scan
ordering and native tree intrinsics, casts, comparisons and min/max keep their
prior behavior. Elementwise operations inside ordered or custom reduction
bodies follow the same explicit fast-math policy. It does not
enable blanket compiler fast math, TF32, general division-to-reciprocal rewrites or a fallback.
Other routes receive their existing `ShaderOption` flag and may implement a
different fast-math policy; compare realization metadata, not flag equality.

The explicit v2 fast policy additionally maps a FP32 `x / sqrt(y)` to
`ct::mul(x, ct::rsqrt(y, FTZ), RNE, FTZ)` only when the DIV operand is directly
the SQRT result. Both original named-axis alignments are retained, including
row broadcasts and scalar values within `map`. Casts, explicit broadcast
operations and other producers are not traversed. A shared SQRT remains
available to its other users. Default/strict generated source is unchanged.
The native route currently has no shader cache; the realization tag identifies
this policy separately from historical v1 results.

The SDK reciprocal square root has an explicit subnormal mode, but no separate
approximate rounding parameter. FP32 FTZ applies to its input/result and to
both multiplication inputs/result. This changes intermediate rounding in fast
mode; it is not a bit-exact rewrite of two separately rounded operations.
For finite positive normal y, reciprocal-sqrt is itself normal, avoiding the
intermediate reciprocal-underflow problem of rewriting arbitrary `x / y`.
Special values and signed zeros remain subject to the runtime conformance
tests. The existing elementwise and workload error bounds are not widened.
See the [CUDA 13.4.1 math API](https://docs.nvidia.com/cuda/archive/13.4.1/cuda-tile-cpp-api-reference/math_operations.html#cuda-tiles-rsqrt)
and [subnormal policies](https://docs.nvidia.com/cuda/cuda-tile-cpp-api-reference/general_principles.html#subnormals-rounding-mode).


### Optional complete chunked-sort pipelines

`ranking_algorithm: "chunked_bitonic_c256"` and `"chunked_bitonic_c512"` are explicit
`sort`-only compositions. The input/output ABI and exact stable index/value oracle
are unchanged. The row Tile remains `(1,bit_ceil(N),1)`, and its width must be at
least the named chunk size. The first kernel sorts alternating runs; each later
kernel merges a doubled span, ending with typed values and INT64 indices. No
public Tile primitive or lower-precision ranking key is introduced.

Every timed operation includes all stages. Scratch comprises two disjoint
FP32-value/INT32-index planes, each with 64 guards at both ends; allocation/upload
and compilation are outside warm execution timing. `pipeline` and
`pipeline_stages` describe the widths, buffers, stage sources, realizations and
individual compile times. The ordinary `compile_ms` is the sum of those stage
compile times. All sources are retained and hashed by the matrix runner.

With `--graph-batch 100`, the graph contains 100 complete pipelines. Protocol
`adaptive_replay_span_v2` calibrates R complete replays per sample; timing is
divided by 100*R, never by its number of kernels. Both actual total spans and
per-operation arrays remain in the receipt, along with stage dispatch counts. `CudaGraphExt` creates a buffer
hazard DAG, so independent stages in adjacent complete operations may overlap.
This is complete-operation throughput, not isolated single-sort latency; use
the existing synchronized single-call measurement for that separate quantity.
There is no inserted dummy dependency or omitted merge. Every operation starts
from the same original input rather than the previous sorted output. Native
stable tie order and the selected Torch ranking contract remain distinct when
the standard Torch unstable sort is requested. No performance advantage is
asserted by selecting this algorithm.

### Extended operator and size inventories

`cuda_llm_extended.json` adds BMM, larger GEMM/GEMV, and paired strict/fast
normalization configurations. Use the standard Torch ranking contract.
`cuda_sort_chunked.json` compares full sorting with both chunk sizes on identical
inputs; run it with `--ranking-contract stable` to require matching index ties.
Both inventories use the ordinary `cuda_matrix.py` command and graph timing
protocol above. A case list specifies coverage, not a claim of performance parity.

BMM dimensions are `[B,M,N,K]`, with contiguous `[B,M,K]` and `[B,K,N]`
inputs, FP32 accumulation, and typed `[B,M,N]` output. Each program handles a
singleton batch slice; batch broadcasting and dynamically grouped matrices
are outside this fixture. GEMM/GEMV/BMM permit at most `2^34` product terms
and `2^24` elements in each input or output. The complete FP64 oracle remains
enabled, so larger cases also cost more host time. Matrix compile failures and
timeouts remain failures; they do not become unsupported capability results.

The [core78 checkpoint](results/2026-10-01-core78-v2/README.md) records the
earlier implementation, its full samples and remaining performance gaps.

The [extended48 checkpoint](results/2026-10-01-extended48-v2/README.md) records
BMM, larger matrices, strict/fast normalization and complete chunked sorting.
It retains the sorting regressions and large-GEMM gaps as well as improvements.

`cuda_selection_embedding.json` adds stable first-index argmax and embedding
lookup with FP32/FP16/BF16 storage. Argmax dimensions are `[rows,width]`; each
call returns the original value and an INT64 index, including exact tie order
and signed-zero bits. Embedding dimensions are `[vocabulary,width,tokens]`,
with a real INT64 token-ID buffer and a typed output. Each program loads one
uniform row ID and one feature block. Invalid IDs are rejected before dispatch;
this fixture does not implement arbitrary lane-dependent Tile gathers. Both
operations use a complete exact selected-source oracle.

`cuda_row_blocking.json` compares one, four and eight independent rows per
program for inclusive tree scan, sum and max. `tile[0]` selects this explicit
schedule; the default remains one row. Column and row tails stay masked, and
the original arithmetic and numerical bounds are unchanged. Repeated schedules
are comparisons for the same workload, not additional shape coverage.

`cuda_gemm_schedules.json` compares four explicit tile shapes for FP16/BF16
1024-cubed GEMM. A larger tile is not assumed to be faster. These three new
inventories use the standard Torch contract and the same matrix command;
argmax always requires first-index ties even under that contract.

### Explicit scheduling and alignment experiments

For `swiglu`, `gelu_residual` and `rope`, `tile[1]` may also split the feature
dimension across programs. RoPE uses half the physical row width as its feature
extent, retaining a separate pair axis so tail stores cannot overwrite the
other half. These cases report `feature_tiled_pointwise_fp32_compute`. Whole-row
schedules keep their original capture, FP32 arithmetic and numerical bounds.
Other row operations still require a tile covering the whole feature extent.
`cuda_pointwise_schedules.json` includes whole-row controls, feature-block
candidates and tail cases.

Embedding additionally permits explicit `tile[0]` values of four or eight.
These schedules use an ordered SERIAL region over rows, with one uniform INT64
ID load per iteration; they report `serial_grouped_uniform_int64_row_gather`.
The default one-row schedule and its `uniform_int64_row_gather` identity remain
unchanged. Feature width and row grouping are separate tuning dimensions.
`cuda_embedding_schedules.json` covers these two dimensions separately.

`cuda_packed_sort.json` compares full stable sorting against
`ranking_algorithm: "packed_fp32"` on identical inputs. The candidate uses one
reversible UINT64-key network, including original indices and signed-zero bits,
and decodes values without another input gather. Its realization is
`stable_packed_fp32_full_sort_prefix`. Narrow storage is promoted exactly to
FP32 as in the existing ranking fixture. Run these comparisons with
`--ranking-contract stable`; both sides must satisfy the same tie contract.
The existing default and chunked algorithms remain selectable.

`cuda_matrix.py --native-aligned16` explicitly enables the native-only
`LUISA_CUDA_TILE_IR_ALIGNED16=1` experiment. Compilation retains the original
entry and adds a second entry only for FP16/BF16 buffer roots whose every access
has proved bounds, row stride and feature origin alignment. The runtime checks
the final device pointers, including BufferView offsets, before selecting that
entry. Direct launches and CUDA Graph construction share this selector; whole
graph updates rebuild the selection from the new bindings. No aliasing promise
is introduced. The flag is off by default, and is removed from inherited
environments for every matrix child unless explicitly requested for native CUDA.

Each native stage records an eligible-buffer mask, actual argument-address
residues modulo 16, and `expected_selected_entry`. The last field is a host-side
prediction, not an execution trace. Requested but ineligible cases still use
the original entry; distinguish them from specialized cases when reporting
performance. Controls and candidates require separate output directories and
otherwise identical case and graph-v2 settings. Compilation includes both
entries and its cost is retained separately from warm execution timing.
`cuda_alignment.json` provides matching off/on cases for GEMM, BMM, scan,
normalization, softmax and reductions, including ineligible tail controls.

`cuda_matrix.py --native-worker-warps 4` (or `8`) requests a CUDA Tile worker
warp hint for calibration on SM89. Zero, the default, emits the original source
without a hint. CUDA 13.4 supports only these two nonzero values; a hint is not
an observation of physical threads, occupancy or the compiler's final choice.
The Tile launch ABI remains unchanged. The runner clears inherited worker hints
for every route and passes the requested value only to native CUDA. Each stage
must report the matching hint, and its emitted source is hashed when the child
finishes. Keep each candidate and the unhinted control in separate output paths.

`--routes native --native-only` runs calibration without a Torch child. These
records explicitly mark Torch as not requested and contain no Torch comparison.
It can be used to compare several schedules against one separately recorded
native-plus-Torch baseline, followed by another unhinted native run to check
drift. Such comparisons must match the full input/oracle receipts, math policy,
measurement configuration and generated source apart from the scheduling hint.
Retain every sample and failed case; the fastest sample is not a cost model.

Admitted native collectives report `collective-work-v1` facts derived from Tile
IR: logical program count, collective width and independent elements, elementwise
work, nominal view bytes, and explicit Tile SSA liveness. These are inputs for
candidate selection and calibration, not estimates of physical registers, spills,
actual memory traffic or elapsed time. Ordered/custom reducers and unrecognized
constraints remain outside this analysis; numerical legality is established
before any scheduling decision.

`--native-streaming-scan 1024` (or `2048`) requests the separate
`LUISA_CUDA_TILE_STREAMING_SCAN` experiment. A closed, unordered FP32 prefix
with FP32/FP16/BF16 storage can load, scan and store consecutive chunks while
retaining an FP32 carry. The original entry remains available. The runtime
selects streaming only when the proved static input/output intervals, using
final BufferView pointers, are disjoint; overlapping views retain the original
snapshot behavior. Direct launches and graph construction share this selector.
Optional compilation or module-loading failure also restores the original
source. Benchmark calibration requires an available candidate and records the
expected host selection from actual pointers/range facts; it does not label a
fallback timing as streaming. This is a host prediction, not a device trace.
Both chunk sizes regressed on all eight configurations in the October 3
calibration. They remain explicit diagnostic candidates and are rejected for
automatic selection; reducing a logical intermediate does not establish a
physical register reduction or a runtime benefit.

`--native-program-rows 1` (or `2` / `4`) requests
`LUISA_CUDA_TILE_PROGRAM_ROWS`, a separate independent-program partition
experiment. The shared Tile IR planner proves independent ownership, reduction
closure, identity padding, complete coverage and exact root-view intervals.
The initial CUDA realization admits contiguous FP32/FP16/BF16 SUM/MAXIMUM with
original tiles of 4 or 8 rows and a strictly smaller divisor. It appends a
separate entry with its own grid; overlapping final input/output intervals and
unavailable optional compilation retain the original entry and grid together.
Direct dispatch and CUDA Graph construction/update share the same launch
selector. Calibration requires actual host bindings to select the candidate,
records both grids, and never labels an unavailable or alias fallback as a
partition timing. The receipts are host selection evidence; the GPU regression
test additionally inspects real graph CUfunction/grid and verifies alias
snapshots and graph updates. This experiment is off by default and mutually
exclusive with other schedule experiments. Its logical work facts are not
physical register or occupancy measurements.

`--native-collective-cost` enables `LUISA_CUDA_TILE_COLLECTIVE_COST=1`, a
default-disabled experimental worker-hint profile. A shared Tile IR feature
extractor and bounded decision-tree evaluator use logical collective work;
the fitted parameters and device/version guard belong to CUDA. The initial
`sm89-24-cuda134-v3` profile targets strict math, SM89 with 24 processors,
32-lane warps, 1536 maximum resident threads per processor, and CUDA toolkit /
Driver API version 13.4. This version-family guard does not check compiler
binary hashes or guarantee performance on every matching driver build.

The profile selects either the original schedule or the 8-worker compiler hint
from an uncalibrated relative log-time prediction. Prefix scans, unsupported
algebra/IR, fast math, and unmatched targets retain the original schedule.
Per-stage metadata records the profile, decision, reason, optional score and
actual requested hint. This experiment is mutually exclusive with explicit
worker hints, alignment specialization and the structural/streaming experiments.
Controls and other routes clear inherited cost settings. The profile contains
no benchmark-name or tensor-shape lookup; all negative calibration outcomes
remain part of the model's validation record.

The [independent collective validation](results/2026-10-03-collective-validation/README.md)
records the frozen profile's twelve unseen configurations, matched default
rechecks, fresh Torch baselines, and the rejected streaming candidates. The
profile improved the geometric mean of per-case time ratios by approximately
2% in both visits, with repeatable gains around 9% for the tested small RMSNorm
and 11% for the tested small Softmax. Pure reductions changed little, and the
tested Softmax and long half-precision scans still trail Torch. This is not a
claim that the profile is reliable across other devices, compilers or shapes.

`cuda_attention_mha.json` adds batch-two MHA decode and prefill with a ragged
key length, complementing the existing causal GQA/MQA attention cases.

The [coverage156 checkpoint](results/2026-10-01-coverage156-v2/README.md)
records the preceding `15d411fc` implementation: 51 operation/shape combinations
and 94 operation/shape/dtype combinations, including 6 long-scan compiler
failures. It preserves every schedule and raw sample, with separate strict and
common-envelope numerical acceptance. It is a baseline for these experiments,
not a result for the new packed, feature-blocked, SERIAL or aligned entries.