Apple GPU: Gumbel-max inference sampler (#18-safe, reproducible) by gstoner · Pull Request #31 · gstoner/tessera

gstoner · 2026-05-30T15:55:31Z

A GPU categorical sampler for inference token sampling — the #18-safe
counterpart to the deferred training RNG.

ids = argmax(logits/T + g),   g_i = -log(-log(u_i))    draws from softmax(logits/T)

#18-safe by design

The Gumbel noise g is generated on the host from the canonical Philox
stream (tessera.rng), so sampling is deterministic, reproducible (same
key + logits ⇒ same tokens), and needs no on-GPU RNG — it doesn't touch
Decision #18's bit-exactness concern at all. The per-row vocab argmax runs
on-GPU via MPSGraph reductionArgMaximum.

API

runtime._apple_gpu_gumbel_sample(logits, key=, temperature=, top_k=, top_p=, greedy=) — host-side top-k / top-p masking, Philox Gumbel noise, GPU argmax,
numpy fallback. greedy=True (or temperature==0) → plain argmax. Returns
int64 ids shaped like the leading dims of logits.

Tests (10)

greedy==argmax, T=0, 20k-sample distribution convergence to softmax (one
batched call), reproducibility, key-sensitivity, top-k / top-p candidate
restriction, batched [B,V], scalar single-row.

Honest benchmark finding

benchmark_gumbel_sampler.py shows the GPU path is upload-bound today — host
numpy argmax is faster (e.g. 8×128k: GPU ~12.8 ms vs host ~0.6 ms) because
uploading [B,V] logits + noise dominates the cheap argmax. This matches the
earlier analysis: the GPU win needs the logits to stay GPU-resident (a
fully-fused decode that doesn't read them back) or a Philox-MSL noise
generator to remove the noise upload. The kernel + API are correct and ready
for that; the value today is the reproducible, #18-safe sampler surface with
top-k/top-p (and a correct host fallback everywhere).

Verification (local, Apple Silicon)

gumbel suite: 10/10; apple_gpu + buffer-pool + ABI + reductions sweep: 513 passed
buffer-pool RAII gate + ABI audit: pass (dashboard 124 → 125 symbols)
mypy ratchet: clean (only the pre-existing environmental torch-import error)

CI on this repo is uniformly red on main (Python 3.8–3.11 matrix, missing
optional deps) — same state PRs #17–#30 merged through. The local signal above
is green.

🤖 Generated with Claude Code

Adds a GPU categorical sampler for inference token sampling, distinct from the deferred training-side RNG (Decision #18). ids = argmax(logits/T + g), g_i = -log(-log(u_i)) draws from softmax(logits/T) The Gumbel noise g is generated on the host from the canonical Philox stream (tessera.rng), so sampling is deterministic, reproducible (same key + logits -> same tokens), and #18-safe — no on-GPU RNG. The per-row vocab argmax runs on-GPU via MPSGraph reductionArgMaximum. - apple_gpu_runtime.mm: mpsg_run_gumbel_argmax + tessera_apple_gpu_gumbel_argmax_f32 (scale + Gumbel-add + per-row argmax, cached graph, RAII buffers, host reference fallback). Stub parity. - runtime.py: _apple_gpu_gumbel_sample(logits, key=, temperature=, top_k=, top_p=, greedy=) — host top-k/top-p masking, Philox Gumbel noise, GPU argmax, numpy fallback; _gumbel_noise_from_key + _apply_topk_topp_mask helpers. - tests/unit/test_apple_gpu_gumbel_sampler.py: 10 tests — greedy==argmax, T=0, 20k-sample distribution convergence to softmax (one batched call), repro, key-sensitivity, top-k / top-p candidate restriction, batched [B,V] shapes, scalar single-row. - benchmarks/apple_gpu/benchmark_gumbel_sampler.py: GPU vs host argmax sweep. HONEST FINDING: the GPU path is upload-bound today (host numpy argmax is faster) — the win needs the logits to stay GPU-resident (fused decode) or a Philox-MSL noise generator to remove the noise upload. The kernel + API are correct and ready for that; the value today is the reproducible #18-safe sampler surface with top-k/top-p. - docs: plan Tier-3 RNG row documents the sampler + the upload-bound finding. runtime_abi dashboard regenerated (124 -> 125 symbols). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>

chatgpt-codex-connector

💡 Codex Review

Here are some automated review suggestions for this pull request.

Reviewed commit: a085c19294

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

Open a pull request for review
Mark a draft as ready
Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".

chatgpt-codex-connector · 2026-05-30T15:58:23Z


+def _apple_gpu_gumbel_argmax_f32() -> Any:
+    runtime = _load_apple_gpu_runtime()
+    sym = getattr(runtime, "tessera_apple_gpu_gumbel_argmax_f32", None)


Require the new runtime symbol before accepting cached builds

When a developer already has build/src/compiler/codegen/Tessera_Apple_Backend/libTesseraAppleRuntime.* from the previous revision, _load_apple_gpu_runtime() can accept that cached library because its acceptance gate was not updated to require tessera_apple_gpu_gumbel_argmax_f32. In that environment this lookup returns None, so the advertised GPU sampler silently runs the host fallback and test_gumbel_symbol_exported fails until the user manually cleans/rebuilds; add the new symbol to the loader's required-symbol checks or otherwise force a rebuild when it is absent.

Useful? React with 👍 / 👎.

chatgpt-codex-connector · 2026-05-30T15:58:23Z

+        kth = np.partition(out, -top_k, axis=-1)[:, -top_k][:, None]
+        out = np.where(out < kth, neg_inf, out)


Mask top-k by indices so ties do not bypass k

For inputs with ties at the kth logit, this threshold mask keeps every value equal to kth, so top_k=1 on all-equal logits leaves the whole vocabulary eligible and Gumbel noise can sample any token. That violates the top-k restriction used by callers/tests (argsort(... )[:k] gives exactly k candidates) and is especially likely with quantized or deliberately uniform logits; build an explicit per-row top-k index mask instead of comparing only against the cutoff value.

Useful? React with 👍 / 👎.

… drops Three findings from PR #500 review, all correct. ── P1: wire the verifier into production lowering pipelines ── Standalone registration was the only reference, so ordinary compilation got no boundary checking and a future drop would stay silent unless a caller reproduced the fixture's CLI by hand. Both passes now bracket `TileIRLoweringPass` in the two named pipelines. That immediately proved the point: `flash_attn_full.mlir` and `nvidia_pipeline_alias.mlir` both reported a real, undeclared re-expression the moment the gate ran on them. Graph IR states `tessera.layout = "row_major"` and the pass restates it as `#tile.layout<shard = ...>`. The name survives, the value does not, and the verifier cannot tell a re-encoding from a replaced accumulator policy -- so the lowering now says which it is, via a new `re_expressed` reason stamped per function and only where a re-expression actually happened (a blanket declaration would itself be refused as STALE_DECLARATION). `re_expressed` is accepted ONLY while the name survives; if the attribute is gone entirely, nothing was re-expressed and it is a silent drop wearing the most permissive label. ── P1: track values, not just names ── The snapshot recorded the SET OF NAMES per function, which is a false-negative generator, exactly as review said. Two cases it waved through, both now fixtures and both now firing: * two ops carry `numeric_policy` and only one survives -- the name is still somewhere in the function, so the surviving occurrence covered for the lost one; * a policy REPLACED (`accum = "fp32"` -> `accum = "fp16"`) -- the name never moved, so nothing fired. This is the instruction-selection corruption the verifier exists to prevent, and it was invisible to it. The snapshot is now a multiset of (name -> printed value -> count), and a missing value raises METADATA_OBLIGATION_VALUE_DROP. ── P2: reject drop declarations absent from both inventories ── STALE_DECLARATION fired only when the attribute was still present. A declaration for an attribute the function NEVER had was silently accepted -- the more dangerous shape, since it looks harmless right up until the function acquires that attribute, at which point it licenses a real drop nobody reviewed. Both shapes are refused now, with the diagnostic naming which. ── One self-inflicted bug worth recording ── The first version of the `re_expressed` stamping looked for a plain `"layout"` attribute and found nothing, so it declared nothing and the two production fixtures kept failing. The attribute is spelled `tessera.layout`. That is the same normalization the verifier already does deliberately, re-implemented wrongly three functions away -- Decision #31's one-implementation rule at function scale. The helper now shares the rule. 294 lit, 14425 unit, mypy 0, ruff clean, docs in sync. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

…kend plans Three findings from PR #502 review, all correct. ── P1: record the shared type migration in the backend plans ── AGENTS.md requires a shared-IR change to assess all backends in the same PR and record the outcome. Added, under a common sync key TILE-FRAGMENT-TYPE-PARAM-2026-08-03: nvidia follow-up required -- 8 files; step 2b (FragmentZeroOp accumulator, 3 sites) blocks the K-loop from LOWERING even once it verifies rocm follow-up required -- 5 files, same 3-site blocker; also the reason `family` is in the type at all (ROCMFragmentLayout wave 32 vs 64) apple not applicable -- 0 references; the Apple lane lowers to func.call on runtime symbols and its simdgroup_matrix path is in the Python synthesizer, not this type. Flagged to revisit when that seam closes, so Apple acquires a `family` rather than a parallel fragment concept. x86 not applicable -- 0 references; x86 carries its own !tessera_x86.tile and has no cooperative-matrix fragment. ── P2: validate parameters before constructing ── `get()` was unchecked, so a nonpositive shape, an empty dtype, or an unknown role/layout/family was constructible. The sharpest case is the one review named: a PARTIALLY populated tuple is not `isUnknown()`, so step 2's type-based verification would have read it as a stated contract when the producer filled in half of it. `FragmentType::verify` now mirrors `TileMmaDescAttr::verify` value for value -- the two describe the same instruction contract from opposite sides, and drift between them would let a fragment state a family the descriptor rejects (Decision #31). The parser goes through `getChecked`, since it is the only entry point for textual IR. `role`'s legal set is {a, b, acc, scale_a, scale_b}, derived from MMAOp::verify and the two backend lowerings. Worth stating because the name is badly overloaded here: producer/consumer/manager are WARP roles and input/scratch are BUFFER roles, all spelled `role = ` on other ops -- a plausible value from a neighbouring vocabulary is the likely mistake, and it now has a fixture. ── P2: escape string parameters when printing ── `parseString` decodes escapes; the printer concatenated raw bytes between quotes. A value containing a quote or backslash therefore parsed once and emitted IR the second parse could not read. Each field now prints through a StringAttr. The round-trip fixture used only clean values and would not have caught this, so it gains a case with an embedded quote -- and that case is meaningful because `elem`/`acc` are open dtype names with no closed set, so printer correctness cannot lean on the new verifier. Five diagnostic codes registered (the registry gate caught them first). 296 lit, 14425 unit, mypy 0, ruff clean, docs in sync. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

…ogue Three findings from PR #503 review, all correct. ── P1: cross-check descriptor element types and layouts ── The descriptor cross-check compared only m/n/k/family/accType, so a retained descriptor could contradict the type on exactly the fields codegen reads: a bf16 A-fragment paired with `a = "f16"`, or a row_major fragment paired with `a_layout = "col_major"`. Both NVIDIALowering.cpp and TileToROCM.cpp select the instruction variant AND the physical register layout from the DESCRIPTOR, so that IR verified while codegen followed a different contract than the type stated -- silently wrong, in the one place the two sources of truth still overlap. The layout case is the sharper one: it decides the register layout, so the wrong answer is a transposed operand. `descriptorAgreesWithFragment()` now compares per role, since the descriptor has always carried A and B separately (aType/bType, a_layout/b_layout) -- the same asymmetry that is why the shared operand checks deliberately exclude `elem` and `layout`. One implementation, used by the consumer side, both producer sides, and the unpack (Decision #31), rather than the two partial copies it replaces. ── P2: let descriptorless typed results reach fragment_unpack ── "The descriptor is optional" was only half true. `FragmentUnpackOp::verify()` still demanded `mmaDescAttr(producer)`, so a descriptorless typed `tile.mma` verified right up until it fed the ordinary epilogue -- it worked only while the result was left packed. That verifier also producer-chased, so it carried the block-argument problem this whole item exists to remove: a K-loop accumulator unpacked AFTER the loop is an `scf.for` result, whose defining op is the loop and carries no descriptor. It now reads the input type when one is present. The K-loop fixture gains the full chain through `fragment_unpack`. ── P2: record the zero-role diagnostic's real origin ── All 11 codes were registered with `pass_origin = MMAOp::verify`, which is wrong for three of them: TILE_FRAGMENT_ZERO_ROLE comes from FragmentZeroOp::verify, TILE_FRAGMENT_ROLE_DISAGREES from the shared producer helper, and TILE_MMA_DESC_DISAGREES now from `descriptorAgreesWithFragment`. Corrected, and TILE_FRAGMENT_UNPACK_ROLE registered. 298 lit, 14425 unit, mypy 0, ruff clean, docs in sync. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

Two P1s, both correct. ── P1: non-ELF output bypassed the funnel ── Worse than reported: TWO sites, not one (the shared elementwise builder and the norm-backward lane). The cause was my transformation, not the branches. I converted raise sites by MATCHING MESSAGE TEXT ("was not an ELF hsaco"), which silently skipped "{pass_name}: gpu.binary not an ELF hsaco" (no "was") "compiled ROCm norm backward lane: ... was not ELF" (no "hsaco") Seven of nine ELF checks converted; two kept masking a malformed-compiler-output failure under strict dispatch -- exactly the class the change exists to protect. That is the third time this thread that deriving a set from one syntactic pattern produced a wrong set, so the fix is structural rather than another message list: every `[:4] != b"\x7fELF"` guard is a compiler-output-validity check BY CONSTRUCTION, whatever its wording. Both sites routed that way, and a new gate, `test_every_non_elf_check_routes_through_the_funnel`, enumerates by that structure. Verified it has teeth by reverting one site: it fails. A structural audit of the remaining `raise _RocmCompiledUnavailable` sites confirms the classification is now complete -- every one is guarded by an envelope condition (`hip is None`, hipInit, hipModuleLoadData, `lib is None`, `opt is None`, dtype/rank/arch), none by output validity. ── P1: record sibling-backend outcomes ── AGENTS.md covers runtime contracts and I did not add the entries on this PR. Added under ROCM-COMPILED-STRICT-DISPATCH-2026-08-04: rocm follow-up required -- owns the change; 18 sites + the structural gate apple parity validated -- already routes failure-class through the same funnel; its sites are among the 18 pre-existing strict failures, and it is the precedent this change follows rather than a parallel mechanism (Decision #31) nvidia not applicable -- no _RocmCompiledUnavailable sites. Recorded, not created: NVIDIA has no equivalent failure/envelope split, so the same masking may exist there; establishing that needs an sm_120 host, which this box is not x86 not applicable -- its raises are `lib is None` / missing-symbol only, envelope by construction since there is no compile step to malform 14439 unit, mypy 0, ruff clean, docs in sync. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

Implements finding C1 from DIFFERENTIABLE_PROGRAMMING_REVIEW.md. For a linear map ∂l(w)[v] = l(v) and the VJP is the adjoint l*; for a multilinear op the JVP is the product-rule sum of the forward with one arg replaced by its tangent (Blondel & Roulet §4.5.4). So a linear primitive's JVP is *derived* from its forward — the hand-maintained jvp.py/vjp.py duplication for linear ops is mechanical. * autodiff/linear.py: MULTILINEAR_PRIMITIVES registry, make_linear_jvp (JVP from linearity), register_derived_linear_jvps (additive gap-filler, never overrides a hand-written JVP — Decision #31 ordering). * custom.py: `linear=True`/`linear_args` on @custom_primitive. A linear primitive's declared transpose_rule is now *used* — registered as the VJP (the adjoint), with the JVP derived from linearity. Previously transpose_rule was stored and reported in metadata but never consumed (Decision #29 violation); it is now load-bearing. test_linear_transposition.py cross-checks the derived JVP against the hand-written one for matmul/gemm/transpose/reshape (proving the duplication is mechanical before anything is removed — Decision #31), checks the linearity identity vs finite differences, asserts the gap-filler is additive, and drives a linear custom primitive's grad end-to-end through its transpose_rule. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

compiler: TileRT-assessment dead-surface cleanups (Decisions #29/#31)

…ter claim Three repo-owner decisions applied to the CuTe IR assessment. 1. L0 decided: ONE C++ implementation, Python binds through the ctypes ABI. This reverses the recommendation in the first draft, which is recorded as withdrawn rather than quietly edited. The decision is the stronger one against our own governance: a single implementation satisfies Decision #31 by construction instead of by declared-oracle exemption, and it makes the MLIR consumers (L3's boundary verifier, L5's carrier) first-class rather than deferred, so the FORGE `⊑` query and the emitter index math end up calling the same code. The cost I raised is not wished away — it is folded into L1 acceptance. Coupling emit/ to a build artifact is the failure the Apple dylib already demonstrates (a stale dylib fails 32 tests rather than skipping), and layout algebra sits under every emitter. So: A1, the binding fails closed with one named diagnostic and ships NO fallback path, because a fallback is a second implementation in disguise (Decision #21a); A2, the build dependency is declared in the ordinary target set with a binding-loads test, so breakage surfaces at test time with a fix instruction instead of inside an emitter. Knock-on sizing: L1 grows ~1w -> ~2w (ABI surface, build wiring, A1/A2); L5 shrinks ~2-3w -> ~1-2w, since the algebra it would have introduced now already exists and there is no second implementation to differential-test. 2. LAYOUT-ALG-1 bound in INTEGRATED_COMPILER_PLAN.md §4, so the work has an owning ID rather than living only in a scoped assessment. Ordered L1..L5 with L1 gated on L2 being committed (otherwise L1 is a Decision #29 violation by construction), and L5 sequenced after W1.1 step 4 per the same #31 ordering caveat W1_1_TYPING_DESIGN.md makes for #tile.mma_desc. The assessment keeps mathematical and acceptance authority; the integrated plan keeps order and promotion. 3. The stale rasterization claim corrected in both documents that carry it -- but only the half that is actually stale. Verified by EXECUTING emit/nvidia_cuda.py's _raster_launch, not by reading imports: row_major, column_major and grouped_m emit materially different block-index code, all four emitters consume tile_rasterization.py, and the MLX swizzle_log heuristic was retired rather than promoted. So "no emitter consumes it" is closed. TileSight's substantive conclusion SURVIVES and is preserved as such: raster_order is carried, not swept -- row_major remains the production choice everywhere and automatic enumeration is deliberately withheld pending an architecture-owned correlation/retain verdict, because ROCM-CALIB-1 established that an unvalidated locality metric must not change a production raster choice. The lever is now expressible on every backend and still unpulled; the blocker moved from codegen to measured device evidence. Correcting this to "the finding is stale" would have been an over-correction that erased a live gap. That also rescopes L4 down (~2w -> ~1-2w): the emitter plumbing exists, so L4 is a consolidation onto shared algebra with bit-identical output as the gate -- a stronger acceptance test than the row-major-only check first proposed -- and the measured half stays out of scope entirely. Gates: 25/25 layout-algebra tests pass, generated-doc drift gate ok (25 docs in sync). Mac / Homebrew python3 3.14.6. Docs plus one hermetic test; no production code and no device lanes. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>

… drift - Route-consistency (codex P2): the eager op namespace and the Apple-lane numpy fallback still delegated tessera.maximum/minimum to np.maximum/np.minimum, whose ±0 tie sign is host-ISA-dependent (SSE second operand, NEON IEEE) — so the fleet tie contract depended on the execution route. Both now share one reference implementation (tessera/_ieee_minmax.py, Decision #31); host-independent tests pin the helper and both consuming routes (test_ieee_minmax_reference.py). - Sibling records (codex P1): x86 outcomes for IEEE-MINMAX-CONTRACT-2026-08-23 and JIT-MATH-AUDIT-FIXES-2026-08-23 (both validated on the AVX-512 host in this change), and assessed NVIDIA outcomes for both keys — emitters surveyed (NaN-propagating maximumf throughout; maxnum only in the Philox floor, input cannot be NaN; no eps-floor pattern), with the exact-device tie probe and the adafactor_vjp NaN follow-up recorded as NR2 Pro work, since no evidence transfers from gfx1151 or the AVX-512 host. - CI drift: regenerated docs/audit/generated/test_coverage.{csv,md} (the new adafactor NaN tests changed the adafactor row's counts). Gates on this branch: rocm/x86 binary + optimizer device suites and the new reference tests 140 passed; check_generated_docs.sh clean; check-tessera-rocm 63/63. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

…ODOs, bump freshness P1 — the row introduces a shared Schedule→Tile contract, so AGENTS.md L81-85 applies to a plan row exactly as it does to code. New sync key NUMPOL-CARRIER-1-2026-08-24 in all four backend queues, each with an architecture-specific outcome rather than a copied sentence. All four are "follow-up required" — the row is newly owned and nothing is implemented yet, so claiming parity anywhere would be false — but the obligations differ: ROCm owns the worked reference (the W1.1 fragment accumulator, which must be RE-EXPRESSED as an instance, not duplicated, #31); x86 has no carrier at all today and is both consumer and regression gate; NVIDIA must sequence behind its open W1.1 typed-fragment producers or the two collide at the same seam; Apple must not require a third policy representation across its Python-synthesizer / C++-pipeline seam. P2 — the resolved ownership TODOs are retired together with the row that resolved them: CORE_SUBSTRATE_VIEW's "four flagged inputs" paragraph now records S5 as closed (three remain open) so it cannot be re-proposed, and AUTODIFF_NEXTGEN_PLAN's two references (§2.3 key table and the AD-JET-IR-1 gate table) point at queue row 3b instead of "no owning row". The one remaining "no owning row" in that plan is the unrelated vmap/batching_rule item and is deliberately left. P2 — CORE_SUBSTRATE_VIEW frontmatter bumped 2026-08-15 → 2026-08-24 and the owning freshness view regenerated: exactly one line changes (that doc's row, 9 days stale → 0), no date churn elsewhere. test_audit_docs.py green; generated docs in sync. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

…eck identity Four findings, all real. The first two were the same defect seen from two sides: admission was too weak, and it was implemented twice. P1 — the gate admitted an unverified carrier. `carriesKeyedRngProduct` checked only the class string and that the digest was 64 hex characters, so the fixture's own placeholder digest passed. The recorded product is supposed to be *evidence*, and evidence nobody checks is decoration. It is now a hash chain: supported schema, lowercase 64-hex digest, `sha256(payload) == digest`, and a payload that names THIS op and THIS effect class — so a fabricated digest, a payload-less digest, and a valid product copied from another operation are each refused with their own reason. The positive fixture now carries a payload generated by the real E1 carrier rather than a placeholder, which is what made the weakness visible in the first place. P2 — admission inside regions. `RegionAdjointInterface::isReplayable` rejected every non-pure op, so the same keyed dropout that E2 admitted at the top level failed with AUTODIFF_REGION_ADJOINT the moment it appeared in an `scf.if` body. The newly admitted family was admissible only in straight-line code, which is not a useful family. Rather than write a second check, the verifier moved into `SemanticEffects.cpp` and all three call sites — the paired pass, the structured region walk, and the structurized-CFG body walk — call it (#31). A nested-region fixture proves both directions: a verified draw differentiates and saves the predicate rather than redrawing it; the same nesting with a product naming another op is still refused, which is the teeth for the sharing claim. P2 — mutation identity. `verify_recorded_state` compared only the content digest, answering "do these bytes match" rather than "is this the same state at the same version". Zero-initialised optimizer state is the everyday counterexample: every lineage's first moment is the same bytes, so a replay reattached to the wrong buffer verified clean. Lineage and version are now required keyword arguments — no default, because a defaulted identity is the permissive answer to a semantic question (#21a) — and the two failure directions get separate messages. P2 — the paired pass no longer emits AUTODIFF_STOCHASTIC_EFFECT, so its metadata stops listing it and names the two codes it does emit. A listed code the pass cannot emit is a declaration with no producer, and it reads as "this family is still refused wholesale" (#29). The in-place `tessera-autodiff` pass still emits the old code and keeps it. Evidence, this box (Strix Halo, gfx1151 + Zen 5 AVX-512): lit tests/tessera-ir/ 414/414 phase_f4 (incl. the new fixture) 62/62 test_recorded_product.py 38 passed test_w4_effects_physical_family.py 5 passed, native_gpu/native_cpu unit -m "not slow" failure profile unchanged from the recorded baseline (fp8 20, apple 17, scheduled_matmul 3, wmma 1 = 41) scripts/check_generated_docs.sh 26 in sync Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>

The design call is settled by project direction: the Tessera foundation is core MLIR/LLVM -> Tile IR -> codegen, and hand-written NVIDIA/CUDA kernels are not what should fall out of a compile. driver.py:526 already encoded it -- the scheduled route wins whenever tessera-opt exists and package_matmul is the fallback -- so nvidia_schedule steers only that fallback and must not select the route. Records the measured outcome (test_e2e_spine_native.py 304 passed / 0 failed on sm_120 with the full driver, was 4 failed at three different assertions), why the benchmark's call site is deliberately left alone, and the Decision #31 follow-on: two packagers now serve one boundary, so the fallback must become a declared oracle with a differential test or be retired -- with a coverage comparison first, per #31's own ordering caveat.

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 merged commit a991f84 into main May 30, 2026
8 of 26 checks passed

gstoner deleted the apple-gpu-gumbel-sampler branch May 30, 2026 15:58

chatgpt-codex-connector Bot reviewed May 30, 2026

View reviewed changes

gstoner mentioned this pull request Aug 4, 2026

W1.3: Decision #32 metadata lowering obligation + boundary verifier #500

Merged

gstoner mentioned this pull request Aug 4, 2026

W1.1 step 1: parameterize !tile.fragment #502

Merged

This was referenced Aug 4, 2026

W1.1 step 2: tile.mma reads its contract from the operand types #503

Merged

runtime: a compiled-ROCm failure stops masquerading as a native launch #507

Merged

gstoner mentioned this pull request Aug 5, 2026

ROCm: map the executing lane frontend-to-hsaco, and price W1.1's adoption cost #522

Merged

gstoner added a commit that referenced this pull request Aug 11, 2026

Merge pull request #544 from gstoner/integrate/tilert-cleanups

e81ee82

compiler: TileRT-assessment dead-surface cleanups (Decisions #29/#31)

This was referenced Aug 25, 2026

docs: NUMPOL-CARRIER-1 — the S5 numeric_policy carrier gets an owning queue row #624

Merged

compiler: SO-3 — the MegaMoE producer consumes inferred dependence edges #625

Merged

This was referenced Aug 25, 2026

docs: W4-EFFECTS-1 — scope the admissible-effects recorded-product ABI #628

Merged

compiler: W4-EFFECTS-1 slices E2–E5 — admissible effects, from the gate to real hardware #630

Merged

Navigation Menu

Apple GPU: Gumbel-max inference sampler (#18-safe, reproducible) - #31

Apple GPU: Gumbel-max inference sampler (#18-safe, reproducible)#31
gstoner merged 1 commit into
mainfrom
apple-gpu-gumbel-sampler

gstoner commented May 30, 2026

Uh oh!

Uh oh!

chatgpt-codex-connector Bot left a comment

Uh oh!

chatgpt-codex-connector Bot May 30, 2026

Uh oh!

chatgpt-codex-connector Bot May 30, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant

		kth = np.partition(out, -top_k, axis=-1)[:, -top_k][:, None]
		out = np.where(out < kth, neg_inf, out)

Conversation

gstoner commented May 30, 2026

#18-safe by design

API

Tests (10)

Honest benchmark finding

Verification (local, Apple Silicon)

Uh oh!

Uh oh!

chatgpt-codex-connector Bot left a comment

Choose a reason for hiding this comment

💡 Codex Review

Uh oh!

chatgpt-codex-connector Bot May 30, 2026

Choose a reason for hiding this comment

Uh oh!

chatgpt-codex-connector Bot May 30, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

1 participant