Apple GPU MLA: same-length B>1 batching in decode_batch (throughput) - #28
Merged
Merged
Conversation
MLABlockPagedCache.decode_batch now groups concurrent sequences by cached length (same length -> same RoPE positions) and dispatches each group in a single B = group_size kernel call instead of looping per sequence. - mla_paged.py: absorb_decode_batch(q_nope[G,H,dn], q_rope[G,H,dr], c_kv[G,S,Dl], k_rope[G,S,dr], ...) -> [G,H,dv]. Reuses the absorbed kernel's existing B>1 support (one MPSGraph dispatch over B=G); per-sequence numpy fallback. - mla_block_paged.py: decode_batch groups queries by length, gathers + stacks each group's windows, and runs one absorb_decode_batch per length group. Validates per-query q shapes; ragged batches just form multiple groups (singletons included). - tests/unit/test_mla_block_paged_cache.py (+3): same-length grouping matches per-sequence decode, mixed-length grouping (groups of 2/3/1), and absorb_decode_batch == stacked absorb_decode_one. - docs: plan item 5 (e) updated; native f16/bf16 MLA kernels noted as the last remaining follow-on. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
gstoner
added a commit
that referenced
this pull request
Jun 1, 2026
ests: 68/68 green across the new modules + the regression sweep (pipeline_gates + conformance_matrix + G5/G6 ABI + execution_matrix + envelope). A — Op×target conformance matrix (complete) python/tessera/compiler/conformance_matrix.py — 7-step proof ladder across 7 ops × 6 targets = 42 cells docs/audit/op_target_conformance.md — generated, drift-gated CLI: python -m tessera.cli.conformance_matrix --render | --check Signal: 5 complete / 13 partial / 24 missing — surfaced one real upstream gap (relu on apple_gpu has runtime envelope but no manifest entry) B.1 — Named pipeline capability gates (complete) python/tessera/compiler/pipeline_gates.py — pure aggregator over capabilities + backend_manifest + execution_matrix + toolchain probes 7 canonical gates: legality → codegen → toolchain → link → runtime_abi → hardware_smoke → numerical API: evaluate(target, op) returns full table; first_failing_gate(target, op) returns the named gate + detail string Every fail carries a non-empty reason (locked by test) Dashboard now reports the first failing gate per cell, with concrete diagnostics: nvidia × matmul: toolchain — nvcc not on PATH (CUDA Toolkit 13.2.1 not installed) rocm × matmul: toolchain — hipcc not on PATH (ROCm 7.2.3 not installed) metalium × matmul: link — artifact_only — IR emits but no linked-kernel path today This directly addresses the audit's framing: "the pipeline should say exactly which gate failed" is now true at the dashboard layer. B.2 (open follow-up) — runtime wiring The gate evaluator exists; runtime.launch() doesn't call it yet. When @jit(target="nvidia") is invoked on a CUDA-less host, the error should be "unsupported: first failing gate = toolchain (nvcc not on PATH)" rather than the current generic dispatcher miss. Created task #28 — small change, separate sequencing. Files (2 new, 4 modified): python/tessera/compiler/pipeline_gates.py — new tests/unit/test_pipeline_gates.py — new (13 tests) python/tessera/compiler/conformance_matrix.py — added first_failing_gate to ProofCell; new dashboard column tests/unit/test_op_target_conformance.py — extended import allowlist to include pipeline_gates docs/audit/op_target_conformance.md — regenerated with gate column docs/audit/compiler_layer_gap_remediation.md — §7 (A) + §8 (B.1) narrative
gstoner
added a commit
that referenced
this pull request
Jul 2, 2026
* Add compiler north-star plan pair + fix sm_120 capability drift Compiler direction (docs): - New paired plan + theory: COMPILER_THEORY_OF_OPERATION.md (three-tier kernel model, accuracy-budgeted measured arbiter, three-system fleet, W1-W8 scope register) and COMPILER_REFACTOR_PLAN.md (workstreams A-E spine + F-K world-class, coordination + §9 source-verified seam verdicts). - Reassess OPTIMIZING_COMPILER_PLAN.md: F0-F5 landed on Apple; rewrite F6 (the backend-build seam) since its "Mac can't run CUDA/ROCm" premise is dead (Strix Halo gfx1151 + NR2 Pro sm_120 now execute) and scope the anti-goal. - Dynamic-shapes decision pulled into the spine: symbolic-dim-aware KernelEmitter/TargetPlugin API + shape-bucket arbiter key, bucket-specialize first. - Wire the north star into MASTER_AUDIT.md, docs/audit/README.md, README.md, docs/README.md, and CLAUDE.md (new Decision #28 + reference row). sm_120 capability fix (code): - gpu_target.py: route all coarse capability properties (supports_wgmma / tcgen05 / tmem / cta_pairs / mbarrier / tma / block_scaled_mma) through the authoritative _CUDA_13_3_FEATURES matrix via cuda_feature_status instead of isa >= SM_x. Fixes consumer Blackwell sm_120 wrongly reporting Hopper wgmma + datacenter tcgen05/TMEM/CTA-pairs as supported (it is NOT a superset of sm_100; its matrix path is mma.sync). Also fixes sm_120 wrongly inheriting the SM_90 FA-4 attn default in jit.py. - test_gpu_target.py: bug-encoding test_sm120_runtime_arch becomes a test_sm120_consumer_blackwell_capabilities regression guard; drop stale "rubin_placeholder" naming. Doc drift cleanups: - CLAUDE.md: fix stale "Execution reality" (gfx1151 + sm_120 now execute). - docs/README.md: ROCm row artifact-only -> gfx1151 hardware-runtime. - CANONICAL_API.md: sm_120 "Rubin placeholder" -> Blackwell consumer; correct the WGMMA column + footnote to match the fixed properties. Gates: mypy clean, generated-doc drift in sync, doc lint passed, test_gpu_target + test_audit_docs green. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * Sync string-alias capability path with sm_120 feature matrix Address PR review: GPUTargetProfile(SM_120) was fixed, but the string-alias capability path (get_target_capability / backend_capabilities) still advertised invalid consumer-Blackwell features. - capabilities.py: drop wgmma / wgmma_sparse / tcgen05 / tcgen05_pair / tmem from nvidia_sm120.features (consumer Blackwell is NOT a superset of datacenter sm_100; FP4 goes through mma.sync.block_scale). Now mirrors the cuda_feature_set(SM_120) "ready" flags. - test_compiler_capabilities.py: add test_nvidia_features_match_cuda_matrix — a single-source-of-truth guard asserting no NVIDIA capability entry advertises a feature the CUDA matrix marks not_supported, plus a positive lock that sm_120 excludes the datacenter/Hopper flags. Prevents this drift from recurring. - Regenerate test_coverage dashboards (deterministic negative_refs count shift from the added guard test). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * Note AOCL-DLP as the x86 Tier-3 candidate for Zen Fold the AOCL-DLP reference (amd/aocl-dlp — AMD's BLIS-family DL primitives: low-precision GEMM/batch GEMM, pre/post-ops, INT4/FP16, symmetric quant, OpenMP) into the north-star plan: - Theory Tier-3 list: add it to the x86 line (CPU analog of cuBLAS/rocWMMA). - Refactor Plan C1: the x86 TargetPlugin registers AOCL-DLP as a Zen-family Tier-3 candidate — AVX512-based (fits the Zen 5 fleet box, no AMX), fills the x86 OpenMP + INT4/FP16 gaps, opt-in behind a build flag (BLAS-family lib like Accelerate, Decision #23-clean), arbiter-selected only where measured faster; license check before a shipped lane. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
gstoner
added a commit
that referenced
this pull request
Jul 7, 2026
) * C3 tail + C1b + D1-core: WMMA/AOCL-DLP Tier-3 candidates + arbiter Introduce the D1 candidate registry (emit/candidate.py) — the accuracy-budgeted arbiter seam from Decision #28's three-tier model — and wire the crown-jewel hand-tuned kernels through it as first-class candidates: D1 core (emit/candidate.py): - Candidate ABC (tier/target/op/available()/applies_to()/run()) + a registry keyed per (target, op); Tier enum {SYNTHESIZED, EMITTED, HAND_TUNED}. - arbitrate()/run_arbitrated(): filter by applicability+availability, F4-gate each candidate through the SAME universal oracle (a KernelRunner adapter reuses fusion_core.verify_synthesized_*), then select by tier priority (crown-jewel first — lead-safe) with a `measure` hook (D2 seam) and a `force` escape (E3). C3 tail (ROCm, live-proven gfx1151): - runtime._rocm_wmma_fused_2d: direct fused WMMA GEMM+bias+{relu,gelu,silu} via the generate-wmma-gemm-kernel Generate* pass (f16 storage / f32 accum). - emit/rocm_hip.py registers RocmWmmaGemmCandidate (Tier-3), RocmGenericHipCandidate (Tier-1), RocmFlashAttnCandidate (Tier-3). MFMA stays analytical (gfx1151 is RDNA3.5/WMMA; MFMA needs CDNA silicon). C1b (x86, opt-in): - emit/x86_aocl_dlp.py registers X86AoclDlpCandidate (Tier-3), availability-gated on $TESSERA_AOCL_DLP_LIB/$TESSERA_AOCL_DLP_SGEMM; the concrete post-op ABI is deliberately not guessed (declines until bound against real headers + license review). emit/x86_llvm.py registers X86GenericCCandidate (Tier-1), proven on Zen 5. Tests: test_candidate_arbiter.py (13 host-free), test_rocm_plugin.py §4 (live gfx1151), test_x86_plugin.py C1b block. Sweep 168 passed. mypy + ruff clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * Address PR #289 review: cache isolation + reference-decliner arbitration P1 — candidate F4 probes shared the backend target as their oracle cache key, so a Tier-3 probe's verdict could be reused for the real backend runner or a sibling candidate on the same target. The _as_runner adapter now namespaces its target as `candidate::<target>::<name>`, landing each probe on a candidate-private verdict-cache key that never collides. New regression test asserts a failing candidate probe on target T does not poison a later non-forced verification of a real runner on T. P2 — a candidate that declined to the numpy reference on its F4 probe still counted as verified, so it could win arbitration by tier and then hand back only the reference, starving a working lower tier. Fixed at two levels: - Arbiter (root cause): verify_candidate now records the probe's execution tag and returns False when the candidate declined (REFERENCE_EXECUTIONS) — a decliner is not a viable arbitration winner. New tests cover the drop + fall-to-lower-tier. - ROCm: RocmWmmaGemmCandidate.available() now probes the ACTUAL fused path (runtime._rocm_wmma_fused_available: tessera-opt + generated kernel), not just the shipped GEMM symbol. - x86 AOCL-DLP: gated behind a `_ABI_WIRED` flag (False until the concrete post-op ctypes ABI is bound against real headers + license review), so a resolvable symbol can no longer make the lane "available" while run() still declines. Sweep 171 passed (+3 regression tests). mypy + ruff clean. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> --------- Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
gstoner
added a commit
that referenced
this pull request
Jul 7, 2026
…293) Unblock the D1 arbiter for plain GEMM — the candidate registry only had fused_region/attention/gated/pointwise, no bare matmul, so the emitted mma.sync GEMM lane had nowhere to plug in. * fusion_core: MatmulRegion (D = A @ B, 16-bit storage / f32 accumulate) + verify_synthesized_matmul (arbiter-only F4 oracle, dtype-rounded reference) + a _round_to_storage helper; re-exported through the fusion facade. * candidate.py: OP_MATMUL op-kind + its (verify, run_matmul) map entry + a run_matmul method on the _as_runner adapter. * runtime.py: two 2D GEMM execution helpers the candidates call — _nvidia_mma_gemm_2d (shipped libtessera_nvidia_gemm, row-major B) and _nvidia_ptx_gemm_2d (compiler-emitted ptx_emit via the launch bridge, col-major B), keyed by 16-bit dtype; + a bridge loader mirroring the shipped-GEMM one. * emit/nvidia_cuda.py: NvidiaMmaGemmShippedCandidate (Tier-3 hand-tuned) and NvidiaMmaGemmEmittedCandidate (Tier-2 emitted, aligned-only) registered under (nvidia, matmul); both F4-gated, f16 accuracy budget. Tier-priority picks the shipped lane by default (lead-safe, Decision #28); the E3 force hatch selects the emitted lane. So NVIDIA gains its Tier-3 hand-tuned GEMM candidate (previously only reachable via the jit nvidia_mma executor) next to the Tier-2 emitted lane. D2's measured loop (lets Tier-2 win where faster) is the follow-on. Live-proven on sm_120 (RTX 5070 Ti): both lanes verify + execute + match the dtype-rounded reference across bf16/f16 x 16x8x16/32x16x32/64x64x64; arbiter selects shipped by default, force selects emitted (test_nvidia_plugin.py). Co-authored-by: angst <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
This was referenced Jul 7, 2026
gstoner
added a commit
that referenced
this pull request
Jul 8, 2026
…he D1 arbiter (#307) * spectral: real Stockham FFT kernels + dynamic-N legalize + LLVM-22 dialect port Make the ts-spectral-opt solver real and buildable end-to-end. Kernels (TargetHooks/{CPU,AMD,NVIDIA}/StockhamRadix4.*): replace the twiddle-less single-butterfly placeholders with a complete mixed-radix (radix-4 + radix-2 tail) Stockham autosort FFT, identical math across all three backends. CPU verified vs naive fp64 DFT; AMD proven on gfx1151 (benchmarks/spectral/hip/stockham_gfx_harness.hip, verdict=pass N=64..1M, ~96 GFLOP/s). Symbols now ts_stockham_r{4,2}_{cpu,nvidia,amd} + a runtime driver ts_fft_stockham_{...}; LowerToTargetIR emits them (dropped the wrong gfx94x/sm90 arch-suffixed names) and tags tessera.target_ir.arbiter_op. LegalizeSpectral: dynamic axes no longer fabricate a bogus radix-4 stage — they defer to the runtime driver via dynamic_shape/dynamic_axes; removed a stray std::reverse that ordered stages tail-first. Dialect ported to LLVM 22 so ts-spectral-opt builds standalone: fix the CMake tablegen (LLVM_TARGET_DEFINITIONS, binary include dir, MLIROptLib/ Func link), the .td (StrAttr, Pure, drop invalid trait/assemblyFormats, add a Plan TypeDef), and the dialect C++ glue (generated classes at global scope). Rewrote the lit fixtures to use a real plan op; 10/10 lit pass. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * tpp: make all 7 space-time passes real + correctness sentinel + LLVM-22 build Turn the tpp-space-time stub passes into real transforms and make the solver buildable/testable standalone. - HaloInfer: derive per-dim halo from the actual access pattern (grad order -> radius; stencil radius from the kernel operand's shape); structured tpp.halo array (was a hardcoded "1,1,0" string). - DistributeHalo: materialise real tpp.halo.exchange ops (new dialect op) in front of each halo consumer, carrying widths/mesh-axes/overlap; fusion-aware. - VectorizeTPP: compute real vector width + tile shape from field shape/dtype. - LowerTPPToTargetIR: annotate ops with hardware-free Target-IR call symbols (cpu/nvidia/amd) + tessera.target_ir.arbiter_op; keep lowered.bc.masked. - LegalizeSpaceTime: default+validate stencil scheme/order and time.step scheme -> stages/order/dt; unknown scheme is a hard error. - FuseStencilTime: group sibling stencils reading the same field (the %Hx/%Hy case) into one shared halo exchange (union halo). - time.step gets NoTerminator so it parses (its region needed an undefined tpp.yield); added a CPU stencil target hook (ts_stencil_grad_cpu). Numerical sentinel: benchmarks/correctness_microbench.cpp (tpp-correctness) implements the linearised shallow-water semantics and checks gradient 2nd-order convergence, periodic-BC correctness, and mass/energy conservation + a traveling-wave match (verdict=pass). Standalone tessera-tpp-opt driver added; 13/13 lit pass. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * arbiter: retarget spectral FFT + TPP stencil onto the D1 candidate arbiter Point both solvers' lower-*-to-target-ir seams at the Workstream-C D1 arbiter (Decision #28): the Target-IR call symbols the passes emit (ts_fft_stockham_cpu, ts_stencil_grad_cpu) become first-class F4-gated arbiter candidates. - candidate.py: additive register_op_kind() + _OP_KIND_VERIFY + a verify_by_reference() helper so solver op-kinds (with their own numpy reference, not a fusion_core region) plug into the same enumerate -> F4-gate -> select pipeline. Existing op-kinds unchanged. - emit/spectral_candidates.py: OP_SPECTRAL_FFT, SpectralFFTRegion (reference = numpy.fft). CPU + ROCm candidates run the real shipped Stockham kernels via ctypes (CPU host-portable; ROCm on gfx1151 via the new ts_fft_stockham_amd_hostptr wrapper). - emit/tpp_candidates.py: OP_TPP_STENCIL, StencilGradRegion (reference = periodic central-difference). CPU candidate runs the shipped stencil kernel. Proven: CPU FFT + stencil match numpy through the arbiter; the ROCm FFT lane runs the real gfx1151 kernel and F4-passes vs numpy.fft; a wrong higher-tier candidate is F4-rejected. 13 new tests + 16 existing arbiter tests pass; mypy clean on the changed files. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> --------- Co-authored-by: gstoner <angstroms01@gmail.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
This was referenced Jul 8, 2026
gstoner
pushed a commit
that referenced
this pull request
Jul 11, 2026
…o Phase 5 The Phase 4 exit (§4) — CompileResult truthfully exposes fwd/bwd states independently + the ledger's runtime_bound rung is matrix-sourced — is satisfied by A2 (ledger) + A3 (provenance/native_required). Record the A1 finding: the ROCm backward already dispatches natively via the autodiff.vjp rules, and there is only ONE executing backward candidate (the Tier-1 synthesized backward is IR-only until Phase 5), so routing backward through the emit arbiter is premature — a one-candidate lane isn't clean. A1 moves to Phase 5, taken when a second executing backward candidate exists (a real Decision-#28 choice, not mechanical wiring). test_coverage regenerated (new flash_attn test references; drift gate: 20 in sync). Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
gstoner
added a commit
that referenced
this pull request
Jul 27, 2026
Adds docs/audit/compiler/LSE_CHECKPOINT_CONTRACT.md recording what tessera_attn.lse.save / lse.load are supposed to mean, what Tessera actually implements, and who owns the open question. The FA-2 design: the per-row log-sum-exp is the coefficient that rescales partial outputs across KV blocks, and the backward needs it to form P = exp(S - LSE). FA-2 stores L to HBM in the forward and loads it in the backward, recomputing only S and P; in Triton and other MLIR kernel languages that is ordinary tl.store/tl.load pointer traffic, minimized to a single committed round-trip rather than per-block. Tessera's ops do not implement that. lse.save takes no destination, lse.load declares no arguments block at all, nothing links a load to a save, and the single emission site discards the result while typing it scalar f32 instead of the per-row [tile_q] vector. It is a declared-but-unimplemented contract, scaffolded in the v1.2/v1.3 design and never wired to memory — the same name-free global-state modeling that #tile.buffer_ref -> !tile.buffer and annotation-only #tile.pipeline_state -> threaded SSA already replaced elsewhere. No backend consumes it, and all three with an attention backward chose recompute, because each sells a zero-workspace determinism property a saved LSE would give back. That is defensible, but not obviously right at long context — which is a bandwidth question the lead performance targets own (Decision #28). NVIDIA-LSE-1 and ROCM-LSE-1 file that evaluation: measure where storing and reloading a [B*H*Sq] fp32 vector beats recomputing L with an extra pass over K, pricing in the workspace and determinism given up, then take one of three outcomes jointly. The third outcome is documented as the preferred landing and is available immediately, independent of the measurement: fix it at the source — stop TileIRLoweringPass emitting a destination-less lse.save and revert LseSaveOp to non-Pure, leaving the op honestly side-effecting and ready for a real implementation. That removes the defect rather than tolerating it, drops a dead op from NVIDIA and ROCm forwards too, takes the Pure trap out of shared ground, and keeps implement-or-retire open. Until then tests/unit/test_lse_checkpoint_contract.py is a tripwire: it fails the build if LseSaveOp acquires a destination-shaped operand, pointer/memref type, or memory-effect interface while still declaring Pure — the case where DCE would silently delete a real store and the backward would read uninitialized LSE. Verified by mutation: adding a memref destination makes it fail. The correct response is to drop Pure, not to relax the test. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
gstoner
added a commit
that referenced
this pull request
Jul 29, 2026
Routing Apple execution through the generic synth->compile->cache loop (kernel_cache.build) was supposed to be a call-site swap. It is not, and the reason is worth recording: AppleMSLEmitter.emit() only ever produced the canonical *scalar* matmul-epilogue body, while run_fused_region prefers the coopmat (simdgroup_matrix) kernel for any eligible region at f16/f32/bf16. Switching production to build() first would therefore have silently downgraded every matrix-unit kernel to the scalar one. Both bodies compute the same numbers, so the F4 oracle would not have caught it — it would have shown up only as a large, unexplained throughput regression. That is precisely the "shared infra must never cap the lead backend's ceiling" rule in Decision #28. So the emitter learns the variants first: * `select_fused_variant(region, dtype)` is now the single rule. `emit()` resolves AUTO through it and `run_fused_region` consults the same function instead of re-deriving the predicate inline, so the kernel the generic loop emits is the kernel the launch path would run. (Same one-predicate-two- callers discipline as the reduce placement fix.) * An explicit `variant=` pins the body. That is the hook a measured arbiter needs to emit every candidate and time them, rather than inheriting this preference order as if it were a decision. * Pinning a variant a region cannot express (coopmat with a reduction, residual, or prologue) raises EmitError rather than quietly returning the scalar body under a coopmat label (Decision #21). Verified: build() now yields `synth_matmul_epi_coopmat` with a simdgroup body for an eligible region, dims sharing a bucket reuse one cache entry, and a dtype variant does not alias. Launch-path numerics unchanged on Metal (f32 max err 0.0, f16 6e-5). Sweep unchanged at 47 pre-existing failures, +3 new passes. test_apple_emitter_wraps_matmul_epilogue_byte_identical encoded the old contract ("the emitter yields the canonical scalar form"). It now pins SCALAR to keep testing byte-identical passthrough, and the AUTO behaviour it used to assert is covered by its own test. Still open before production routes through build(): run_fused_region does synthesis and dispatch in one call, so using a pre-built KernelSource needs those split. That is the next slice, not a rename. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
gstoner
added a commit
that referenced
this pull request
Jul 29, 2026
Track 2a. `apple_gpu` is compile-on-launch — its compile_fn returns None and
Metal builds the kernel inside run_* via newLibraryWithSource:. Right default
for a JIT, wrong one for measuring: every distinct kernel pays a front-end
compile on first launch and nothing survives the process.
`apple_gpu_air` is a second registered target emitting the *same* MSL and
compiling it ahead of time:
MSL --`xcrun metal -c`--> .air --`xcrun metallib`--> .metallib
Verified on this host with the synthesizer's real coopmat output: 7296-byte
.air, 7390-byte .metallib, `deferred=False` with the artifact on disk, while
`apple_gpu` still defers and keys separately.
A separate target rather than a flag, so AOT-vs-JIT is a measured arbiter
candidate per (op, shape-bucket, dtype, target) per Decision #28 — both can be
built and timed — instead of a build-time switch nobody revisits.
The emitter delegates to AppleMSLEmitter rather than duplicating the synthesis
dispatch. A second copy would drift, and then "AOT vs JIT" would be comparing
two different kernels while claiming to compare compile strategies.
Without the Metal toolchain the lane raises MetalToolchainError with the
xcode-select / downloadComponent commands, and never falls back to the JIT
path — an AOT measurement that was quietly a JIT one is worse than no
measurement. Compile failures name the failing stage (Decision #21).
Artifacts are content-addressed on (source, entry), so identical source
compiles once per machine; TESSERA_APPLE_AIR_CACHE relocates the cache.
Grounding for the harder question (Decision #26a): the .air is LLVM bitcode —
magic dec0170b, `target triple = "air64_v28-apple-macosx26.0.0"`, and our
pinned LLVM 23 llvm-dis reads it. A test asserts the bitcode magic so a
toolchain change that stopped producing it is caught here rather than by
whoever later attempts direct AIR emission. Nothing in this lane depends on
that; it only bypasses the MSL front end.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
This was referenced Jul 29, 2026
gstoner
added a commit
that referenced
this pull request
Jul 29, 2026
Extends the two survey documents with six more ROCm projects. Documentation only. Compiler survey gains section 4.7 on rocisa, TensileLite's nanobind assembly generator — the same Python-driving-C++ shape we have. Three findings worth copying: IR nodes carry a mandatory clone() deep-copy contract; exporting a vector to Python is a copy, so elements are mutable through their shared_ptr but cannot be assigned or replaced; and import raises if any C++ source is newer than the built extension. That last one is added to the take list — we lost time this session to a tessera-opt binary that silently did not match its sources. Patterns doc gains four project briefs and a rocWMMA re-read: rocFFT has the best cache design in the ecosystem. The kernel name is the cache key, with every differentiating parameter encoded into it, so profiler output and cache identity are the same string and the cache needs no schema update when a new parameter appears. Three further key fields guard staleness — architecture, HIP version, and generator version. A read-only system cache ships with the library alongside a read-write user cache, the shipped one populated at build time by a helper that shares the generator but is not installed. AOT and JIT are one path with a policy knob rather than two lanes. Also records that hipRTC holds process-wide locks, so parallel compilation needs a helper process. rocPRIM turns tuning output into generated headers, and its fallback_config is a typed fallback ladder: an untuned type inherits the config of a representative matched on size range and floating-pointness. That is dtype bucketing, the same move Decision #28 makes for shapes. rocRAND is the one with a direct bearing on us. Under dynamic ordering it picks launch geometry per device, and AMD states plainly that the number of generators and the sequence of generated numbers can vary as a result. So reproducibility versus performance is a named opt-in mode, not an emergent property. Worth confirming the same holds for Decision #18: if a tuned launch configuration ever fed an RNG offset scheme, autotuning would silently change numerical output. rocALUTION is included as a contrast, not a pattern. It selects execution location at run time via RTTI and silently migrates an object back to the host when a routine is unavailable on the accelerator. That is the opposite of Decision #21, which requires a diagnostic naming the op and target. Both are defensible for their audience; the contrast is worth recording because silent host migration is how a performance cliff hides. rocWMMA re-read adds that collaborative fragments are a movement concept and are explicitly unsupported in MMA functions, that partial and oversized tiles became the library's problem in 2.0.0, and that the wavefront-centric contract is undefined behaviour rather than a hint. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
gstoner
added a commit
that referenced
this pull request
Jul 29, 2026
Extends the two survey documents with six more ROCm projects. Documentation only. Compiler survey gains section 4.7 on rocisa, TensileLite's nanobind assembly generator — the same Python-driving-C++ shape we have. Three findings worth copying: IR nodes carry a mandatory clone() deep-copy contract; exporting a vector to Python is a copy, so elements are mutable through their shared_ptr but cannot be assigned or replaced; and import raises if any C++ source is newer than the built extension. That last one is added to the take list — we lost time this session to a tessera-opt binary that silently did not match its sources. Patterns doc gains four project briefs and a rocWMMA re-read: rocFFT has the best cache design in the ecosystem. The kernel name is the cache key, with every differentiating parameter encoded into it, so profiler output and cache identity are the same string and the cache needs no schema update when a new parameter appears. Three further key fields guard staleness — architecture, HIP version, and generator version. A read-only system cache ships with the library alongside a read-write user cache, the shipped one populated at build time by a helper that shares the generator but is not installed. AOT and JIT are one path with a policy knob rather than two lanes. Also records that hipRTC holds process-wide locks, so parallel compilation needs a helper process. rocPRIM turns tuning output into generated headers, and its fallback_config is a typed fallback ladder: an untuned type inherits the config of a representative matched on size range and floating-pointness. That is dtype bucketing, the same move Decision #28 makes for shapes. rocRAND is the one with a direct bearing on us. Under dynamic ordering it picks launch geometry per device, and AMD states plainly that the number of generators and the sequence of generated numbers can vary as a result. So reproducibility versus performance is a named opt-in mode, not an emergent property. Worth confirming the same holds for Decision #18: if a tuned launch configuration ever fed an RNG offset scheme, autotuning would silently change numerical output. rocALUTION is included as a contrast, not a pattern. It selects execution location at run time via RTTI and silently migrates an object back to the host when a routine is unavailable on the accelerator. That is the opposite of Decision #21, which requires a diagnostic naming the op and target. Both are defensible for their audience; the contrast is worth recording because silent host migration is how a performance cliff hides. rocWMMA re-read adds that collaborative fragments are a movement concept and are explicitly unsupported in MMA functions, that partial and oversized tiles became the library's problem in 2.0.0, and that the wavefront-centric contract is undefined behaviour rather than a hint. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
gstoner
pushed a commit
that referenced
this pull request
Aug 4, 2026
Measured on starting step 3. `FragmentPackOp::verify` requires exactly one
`!tile.tile` input. What the five `tile.mma` construction sites actually pass:
TileIRLoweringPass x2 tile.async_copy results -- TENSORS
GenerateWMMAGemmKernel lane-level VECTORS (toFrag -> vector::BitCastOp)
GenerateWMMALinearAttnKernel same
GenerateWMMAFlashAttnKernel same
Zero producers pass a `!tile.tile` or a `tile.view` result, so no operand can be
wrapped in `fragment_pack`. The inventory warned to expect per-producer
surprises in step 3; the surprise is not per-producer, it is all of them.
This is a division-of-labour mismatch rather than a syntax gap. The typed form
assumes the COMPILER does the lane mapping (`materializeFragmentPack`); the
hand-written generators do it themselves and hand over finished vectors. Both
are coherent, and they are different models.
So step 3 as written is a rewrite of working, numerically-verified generators --
including the production ROCm GEMM lane -- and step 5 ("delete the permissive
branch") is unreachable, since deleting it breaks every existing producer.
Three options recorded, recommending (c): scope the typed form to synthesized
kernels (the Decision #28 lane) and treat the permissive branch as a DECLARED
compatibility boundary between two legitimate models. (a) rewrites proven
kernels for no measured benefit; (b) widening fragment_pack discards what the
typed contract buys.
Not choosing unilaterally -- it changes W1.1's endpoint. No code written for
steps 3/4, because their premise does not hold.
Docs only.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
This was referenced Aug 4, 2026
This was referenced Aug 18, 2026
Merged
gstoner
added a commit
that referenced
this pull request
Aug 30, 2026
The direction is settled (prune the Python bootstrap backend path, build the core through MLIR/LLVM), so the question that now blocks it is coverage: which families does the mainline compiler already serve, and which would lose their only lowering? This derives that from the live sources and drift-gates it, because the point is to watch a gap close and a hand table would be stale by the second landing. Measured: 49 package_* functions across four backends, of which 34 are bootstrap (they take a GraphIRModule and emit target code, bypassing Schedule and Tile) and 15 are compiled-route packagers that consume an already-lowered artifact and are NOT prune targets. Of 24 classified families, 6 resolve to a compiled admission predicate and 18 do not. The classification is by first-parameter type, not by name. A first version keyed on the name suffix and wrongly counted six compiled-route packagers -- package_scheduled_attention and siblings -- as bootstrap surface, overstating the prune. GraphIRModule vs Scheduled*Artifact is a data-flow fact; the name is a convention. Counts are AST-derived. The one thing that cannot be derived is which compiled predicate serves which family, since that lives in driver.py control flow -- so it is declared and VERIFIED: a renamed predicate raises rather than reporting a family as covered. That is the single error that would make a prune silently lossy, so it fails closed, and a test asserts the guard actually fires. The doc states the only two ways a family may leave the table -- absorbed by the mainline compiler, or re-expressed through a declared Target IR boundary so the Decision #28 arbiter can score it -- so it cannot be read as a delete-list.
gstoner
added a commit
that referenced
this pull request
Aug 30, 2026
Step 1 of the prune sequence. The contract landed in the previous commit verified but unread -- nothing scored a delegate, which was the stated gap. This closes it. `DelegatedCandidate` derives BOTH arbiter-relevant facts from the IR contract rather than accepting them at the registration site: tier comes from `provenance` and the F4 budget from `accuracy`/`tolerance`. A delegate therefore cannot claim in Python a budget it did not declare to the verifier. Everything else about arbitration is unchanged, which is the point -- a delegate is enumerated, F4-gated and selected by the same arbitrate() path as compiled candidates, so Decision #28 scores a hand-tuned kernel AGAINST compiled output rather than above it. Both provenances map to Tier.HAND_TUNED: a vendor library and an in-tree hand-tuned kernel differ in origin but not in what the arbiter must do with them. `provenance` is still carried because it is what distinguishes them in a dispatch log. `reference_exact` yields atol=None -- the oracle's default budget, the same standard compiled output is held to -- deliberately not 0.0, which would reject a correct candidate. An exact claim is "no worse than the reference", not "bit-identical in floating point". The Python validator is not a second contract. It rejects exactly what TesseraNVIDIADialect.cpp rejects, and test_delegate_contract.py asserts that case for case, including a count check so neither side can grow a case alone. Two enforcers of one contract is the shape that produced the Apple two-disconnected-compilers defect; the differential is what makes this a bridge rather than another instance of it. Fixture grew to 9 cases (added unknown_provenance, empty_arch) to match. Evidence: 19 Python tests, NVIDIA lit 60/60 on The-Super-Bear, ruff and mypy clean, 201 arbiter-adjacent tests pass.
gstoner
added a commit
that referenced
this pull request
Aug 30, 2026
A sweep of all 32 decisions (plus sub-decisions) against the direction: MLIR/LLVM core, prune the Python bootstrap backend path, contract- carrying Target IR, measured three-tier arbiter. Most support it. Six did not, and one was an unresolved conflict. #28 vs #31 -- the one that mattered. #31 says one production lowering per boundary and delete the second; #28 keeps three tiers of kernels deliberately competing for the same op. Read literally, every Tier-3 candidate is a #31 violation, and the delegation contract sits exactly on the ambiguity. The bootstrap prune therefore had no principled stopping point: #31 could be cited to delete the whole Tier-3 population, which is the ceiling #28 exists to protect. Reconciled in both places -- #31 governs lowering PATHS (how IR descends a level), #28 governs implementation SELECTION (which kernel runs for one op at one level). The test is not how many kernels exist but how many authorities decide what the next level looks like. #1 was actively harmful, not merely stale: it named AMX as the only execution path (retired) and gated GPU work behind isa >= SM_90, which reads as excluding sm_120 -- the live NVIDIA lane. Applied literally it gates off working hardware. #11 keys the autotune cache on {op, shape, dtype, arch, layout, numeric_policy, movement} with nothing versioned. Under #28 a cached entry is a measurement, and a measurement is only valid for the code that produced it; a toolkit upgrade silently invalidates every entry without invalidating the cache. Same failure as the Krylov ratchet, latent in a database instead of a JSON file. #12's schema cannot say which route produced a latency, so three competing tiers are not comparable. Practice was already ahead of the rule -- record_sm120_packet.py stamps `route` -- so this is a schema gap, and the added field is additive. #26a's "revisit on architectural grounds" trigger arrived. The architectural gap is real (Apple's Target IR declares dispatch containers and no machine primitives, while apple_msl.py already models simdgroup_matrix) but it is answered by up-levelling the dialect, NOT by emitting AIR -- NVVM and ROCDL sit above LLVM IR too. That strengthens the deferral rather than reversing it. #29 gains the sequencing corollary that keeps operator expansion honest: add each op only when its producer and consumer land with it. Gates: docs lint, 36 governance/audit tests, 28 generated docs in sync.
gstoner
added a commit
that referenced
this pull request
Aug 30, 2026
Closes the fusion-foreclosure bias, and the bias was worse than the plan described. arbitrate() picks by TIER by default and HAND_TUNED is the highest, so a delegate won outright before anything was measured; on the measured path it won because min(key=measure) scores the candidate in isolation and excludes the work it displaced. Both paths preferred delegates on exactly the graphs where fusion is the win. Fixed structurally, not with a penalty. A delegate declares `covers` (root_only | whole_region) and DelegatedCandidate.applies_to declines a region carrying epilogue / reduction / prologue / residual it does not implement. Why declining rather than penalising: a penalty is a guess at foregone DRAM traffic that then has to outweigh a tier bonus, and a wrong guess fails silently in the direction of the bias it was meant to remove. "This candidate does not serve this region" is a fact the delegate declared. A bare GEMM facing a matmul+epilogue region is not a cheaper way to do the same work -- it is a different plan (delegate + separate epilogue + DRAM round-trip), and comparing plans is not what arbitrate() does. Coverage is declared rather than inferred because an external kernel cannot be introspected, and guessing is how a partial candidate wins a whole-region comparison. Whole-region hand-tuned kernels still compete unchanged, so #28's governing rule -- shared infra never caps the leads -- is preserved; the test asserts that direction explicitly rather than only the decline. Evidence: 26 contract tests including the scenario itself, NVIDIA lit 60/60 on The-Super-Bear with the new required attribute, mypy and ruff clean.
gstoner
added a commit
that referenced
this pull request
Aug 31, 2026
…t time it The device timer I added for `nvidia_mma_gemm_emitted` returned rc=5 on sm_120 every time, and the cause is a real divergence rather than a missing table entry. The launch bridge's `benchmarkTileGemm16` launches `gx = ceil(N/tileN)`, `gy = ceil(M/tileM)` -- x maps to N. The NVIDIA Tile lowering agrees (`NVIDIALowering.cpp`: `mt = blockY*16`, `nt = blockX*8`), which is why both Tile candidates time correctly through it and why their existing latencies are sound. `ptx_emit` uses the opposite convention (`mt = ctaid.x*16`, `nt = ctaid.y*8`), as does the shipped AOT kernel. So the NVIDIA backend carries two block-index conventions. Registering the emitted kernel's geometry in `tileLaunchConfig` would launch it transposed: at 512x512 that covers rows to 1024 and columns only to 256, leaving half the output unwritten while still reporting a plausible latency -- a number that is worse than no number. Swapping the tile dims is not a fix either; it lines up only when M == N. Rather than ship that, the timer is removed, the divergence is recorded where the code is, and the device test names the one candidate without a timer so an unexplained `None` stays a regression while this one stays a tracked gap. Decision #28's displacement test is unaffected: the delegate is measured against the two compiled Tile lanes. A unit test now pins `ptx_emit`'s axis mapping, so "fixing" one side of the divergence fails loudly instead of silently producing a mis-shaped grid.
gstoner
added a commit
that referenced
this pull request
Aug 31, 2026
The device timer's first real output contradicts the arbiter's default.
On sm_120, f16, square, device-resident timing with 0.000-0.008 ms
spreads across repeats:
shape shipped(T3) tile_shared(T2) faster max|err|
512^3 0.043 ms 0.059 ms delegate, 37% both 2.48e-05
1024^3 0.320 ms 0.312 ms compiled, 2.3% both 6.10e-05
2048^3 2.448 ms 2.051 ms compiled, 16.2% both 1.54e-04
Errors are equal at every shape, so Decision #28's in-budget half is
satisfied outright and the displacement condition holds at 1024^3 and
above -- while `arbitrate()` still returns the delegate, because tier
priority is the default.
The test I first wrote asserted "the delegate is measurably the fastest"
and checked only 512^3, the one shape where that is true. It passed, and
it would have reported green for a default that is 16% wrong at 2048^3.
It now asserts the crossover in both directions, and checks that the
faster compiled kernel is no less accurate -- otherwise "faster" is not a
displacement argument.
This is not an argument for deleting the delegate: it wins by 37% at
512^3, and a flat "measurement beats tier" switch would regress that.
It is an argument for shape-bucketed measured selection, which is what
Decision #28's lead-safety is for.
Recorded in the NVIDIA queue with both follow-ups: wiring measured
selection into the OP_MATMUL path, and the two block-index conventions
that keep the emitted lane unmeasurable.
This was referenced Aug 31, 2026
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Compute-side optimization #1 for the MLA serving stack:
MLABlockPagedCache.decode_batchnow groups concurrent sequences by cached length and dispatches each group in
a single
B = group_sizekernel call instead of looping per sequence.Why it's correct to batch
Sequences sharing a cached length share their RoPE positions (
arange(S)), so alength-group can stack into one
[B,H,1,*]tensor and reuse the one cos/sintable set — exactly what the absorbed kernel's existing
B>1path consumes.The per-sequence latent
c_kv/k_ropediffer and stack alongB; the weightsare shared across the batch.
What
mla_paged.py:absorb_decode_batch(q_nope[G,H,dn], q_rope[G,H,dr], c_kv[G,S,Dl], k_rope[G,S,dr], …) -> [G,H,dv]— one MPSGraph dispatch overB=G, per-sequence numpy fallback.mla_block_paged.py:decode_batchgroups queries by length, gathers + stackseach group's windows, runs one
absorb_decode_batchper group. Ragged batchesjust form multiple groups (singletons included), so the public behavior is
unchanged — only the dispatch count drops.
Tests (+3, 11 total in the file)
test_decode_batch_same_length_grouped— 4 same-length sequences in one groupmatch per-sequence single decode.
test_decode_batch_mixed_lengths_grouping— groups of 2/3/1 each match theirsingle decode.
test_absorb_decode_batch_matches_per_seq— the batched helper equals stackingabsorb_decode_one.Verification (local, Apple Silicon)
torch-import error)B>1support)Remaining follow-on
Native f16/bf16 for the MLA decode kernels (f32 today) — the next PR.
🤖 Generated with Claude Code