CAKE on SM120/SM121: evidence review, b12x replacement map, and independent reproduction design

01-cake-consumer-blackwell-review-and-reproduction-design.md

CAKE on Consumer Blackwell: Evidence, Kernel Replacement Map, and Reproduction Design

Date: 2026-10-13. Status: research and proposed design, not an implemented compiler or a qualified kernel cutover.

Executive decision

CAKE's compiler-agent methodology applies to SM120 and SM121. Public FlashInfer contributions already demonstrate consumer-specific generated schedules. Datacenter Blackwell kernels cannot be enabled by changing an architecture string: consumer hardware has warp mma.sync and register accumulators, but lacks tcgen05, tensor memory (TMEM), WGMMA, and two-CTA MMA.

Use family-targeted sm_120f objects where all required instructions are family-supported. Keep operator admission and performance tuning separate from cubin compatibility. An instruction-compatible GB10 is not automatically qualified for an RTX-tuned operator.

No current CAKE merge replaces an entire flashinfer/experimental/b12x family unchanged. The closest replacements are V4.1 mixed-cache sparse-MLA decode, native NVFP4 sparse MLA, and SM120 paged FP8 indexer scoring. The CAKE MegaMoE merge is a different expert-parallel execution architecture, not a drop-in rank-local MoE replacement.

The public sources reviewed do not provide the Cake compiler, verifier, cost model, or agent harness. We can reproduce the methodology independently from the paper, exported schedules, tests, benchmark contracts, and review discussions. We cannot claim to reproduce its unpublished implementation or headline experiment without its artifacts.

Contents

1. Evidence rules and scope
2. Paper methods and limitations
3. Consumer ISA and target contracts
4. PR and issue findings
5. b12x replacement matrix
6. Independent methodology reproduction design
7. Qualification and implementation sequence
8. Open questions and source index

1. Evidence rules and scope

This report combines the CAKE paper, version 1, FlashInfer PR/issue discussions, checked-out implementation contracts, and one small on-device ISA probe. The b12x contract mapping used checkout revision 399b6c1d3, as recorded by the attention investigation. PR status is a snapshot, not a promise about future releases.

Evidence class	Meaning	What it does not establish
Documented	Paper or vendor specification states a capability	Compiler exposure, correct lowering, execution or throughput
Source-inspected	A source contract or instruction sequence was read	Runtime correctness or performance
Reported measurement	A PR author or reviewer reports a run	Independent reproduction in this session
Session measurement	A command or probe ran during this investigation	Results beyond that probe's inputs and hardware
Proposed	A design or acceptance gate recommended here	Existing implementation

Reported review findings remain findings, not universally reproduced defects. Where a fix was not located, this report does not assert that no later fix exists. No full CAKE operator, b12x substitution, consumer roofline, or end-to-end serving benchmark ran in this session.

The report deliberately omits private network addresses, local account paths, and unrelated project references.

2. What the paper contributes

2.1 The two loops

The paper, sections 2-4 separates kernel evolution from compiler evolution.

The kernel loop generates structurally different schedules, checks their legality, ranks promising candidates, runs independent numerical validation, measures GPU execution, and retains measured winners.

The compiler loop turns recurring failures into verifier rules, lowering fixes, new primitives, or calibration work. Human merge gates and a kernel corpus guard those changes.

Agents author a typed, hardware-explicit schedule. They declare shared-memory regions, synchronization objects, warp roles, pipelines, and target-specific compute resources. Lowering derives mechanical details such as barrier addresses and phases. Layout is not a general algebra exposed to the agent. The agent writes concrete storage/access commitments, and the compiler checks producer-consumer compatibility.

2.2 Analysis contract

Analysis	Role
Program safety	Pre-compile gate for modeled synchronization, ordering and memory hazards
Hardware conformance	Pre-compile gate for target instructions and resource limits
Data consistency	Pre-compile gate for producer-consumer representation compatibility
Schedule semantics	Pre-compile gate for declared roles, stages and resource lifetimes
Numerical validation	Execution gate against an external oracle
Performance analysis	Non-blocking estimate and bottleneck attribution
Optimization guidance	Non-blocking proposal hints

The paper states that static analysis operates only inside its modeled domain. It is not a proof of global GPU correctness. Missing coverage must be reported, not represented as a successful check.

2.3 Evaluation and transfer limits

The paper reports a matched Flash-KMeans clean-start experiment on B200. It uses one BF16 assign shape: B=32, N=65,536, K=1,024, D=128. Three runs per representation reach median best speedups of 1.144x for Cake IR and 0.928x for direct CUDA/PTX at the stated token cutoff. These are research outcomes, not a forecast for consumer GPUs or another agent.

The paper reports KDA prefill at a 2.05x geometric-mean speedup over FlashKDA across six B200 shapes. It also reports dispatcher-backed KNN/KMeans families. Generalization is a separate stage: fixed-shape seeds are not accepted as a library until shape guards, held-out cases, tails, and fallback behavior are validated.

Alpha-MoE illustrates measurement attribution. Its 6.204x and 4.025x API-level gains include launch/scheduling changes. GPU-span gains are 1.215x and 1.170x. An independent implementation must report both rather than treating API speedup as tensor-core speedup.

The paper lists SM120a and SM121a in Appendix B.5, but provides no consumer performance evaluation. Its timing model is calibrated for B200 and H100 and declines to predict elsewhere. Consumer calibration is new work.

The four paper-linked PRs are datacenter-targeted. Later consumer PRs are the relevant adoption evidence.

2.4 Availability

Generated CUDA, tests, benchmarks, and integration contracts are public in FlashInfer. The sources reviewed do not publish the Cake IR implementation, compiler, verifier, cost model, or full agent harness. Tracker #4254 includes an unanswered public question about open sourcing CAKE.

The paper's controlled agent/model configuration and the broader tracker contributions must not be conflated. Tracker contributions mention multiple agents. That does not invalidate the controlled experiment, but it prevents attributing every later merge to its exact experimental setup.

3. Consumer ISA and target contracts

The authority for instruction admission is NVIDIA PTX ISA. Pin the exact vendor/toolchain edition used by each experiment. The paper's architecture table is an implementation coverage statement, not a substitute for vendor instruction gates.

Capability	Plain SM120 target	SM120f family target	Scheduling consequence
BF16/FP16 and ordinary FP8 warp MMA	Available	Available	Accumulators live in registers
Ordinary cp.async, TMA and transaction barriers	Available	Available	Explicit staging and producer-consumer pipelines are viable
Narrow/block-scaled Blackwell MMA	Not admitted in baseline target	Available, subject to PTX version	Native NVFP4/MXFP4 and unit-scale FP8 belong on the family path
.m16n16.b8 transpose matrix load	Not admitted	Available	Avoids some explicit FP8 transpose/packing work
setmaxnreg	Not admitted in baseline target	Available	Register redistribution needs a valid role schedule
tcgen05, TMEM, WGMMA, two-CTA MMA	Unavailable	Unavailable	Datacenter schedules require redesign

SM120 is the RTX consumer/workstation generation. SM121 includes GB10. The family suffix describes compatible feature admission. It does not change the physical GPU or increase instruction rate by itself. Exact a targets have different compatibility rules from family f targets.

The presence of thread-block clusters does not imply support for datacenter cluster MMA or every multicast mode. Never infer a feature from the broad label "Blackwell".

3.1 Session ISA probe

The supplied remote GPU was NVIDIA GB10, compute capability 12.1. The assembler was CUDA 13.4, ptxas 13.4.92. The default system Python did not have PyTorch, so the probe used the CUDA driver through ctypes.

The probe initialized a shared-memory tile with packed bytes, executed ldmatrix.sync.aligned.m16n16.x1.trans.shared.b8, wrote its registers to global memory, and checked 64 packed output words.

Target	Assembly	Execution
sm_120a	Accepted positive control	Not launched
sm_120	Rejected explicitly for .m16n16 and ldmatrix.b8	Not applicable
sm_120f	Accepted	Driver module load, launch and synchronization passed on GB10. All 64 words matched.

This establishes family-object execution for that instruction on SM121. It does not prove an entire CAKE export loads, block-scaled MMA throughput, or RTX performance. No active compute process was listed by the pre-probe query. The probe was removed after execution. Future probes must use durable home/repository storage, not temporary-system storage.

3.2 GeForce FP8 instruction selection

Issue #5963 reports that legacy FP8 MMA with FP32 accumulation has lower throughput on RTX 5090 than block-scaled FP8 MMA with unit UE8M0 scales.

The issue reports raw loop rates of 510 versus 1,014 TFLOPS, then complete attention speedups of 1.29-1.34x. Unit UE8M0 scale bytes are 127, representing 2^0. The reported outputs were bit-identical for the tested inputs. SASS changed from QMMA.16832.F32.E4M3.E4M3 to QMMA.SF.16832.F32.E4M3.E4M3.E8. Register use increased from 247 to 248 without spills in that measurement.

This is a performance hypothesis for a specific SKU, not a legality rule for all SM120 devices. The issue requests PRO 6000 and GB10 qualification. Those devices must be measured separately. The identity of mathematical scaling does not guarantee bitwise identity across instruction families on all inputs.

4. PR and issue findings

4.1 Original paper deliveries

PR	Status in review	Contribution	Consumer applicability
#4262	Merged	KDA recurrent prefill	Datacenter schedule and exact target. Methods transfer, binary does not.
#4279	Merged	KDA recurrent decode, frozen CUDA variants	No SM12x replacement. Strict backend contracts matter.
#4274	Merged	TinyGEMM2 schedule portfolio	Datacenter only. Shallow/deep pipeline selection is transferable.
#4287	Merged	Fused W8A8 Alpha-MoE	No direct consumer replacement. Review includes an end-to-end throughput regression for a tested configuration.

4.2 Actual consumer deliveries

PR	Status in review	Scope	Important limitation
#5914	Merged	DSV4 NVFP4 sparse-MLA decode	Native low-precision cache/compute contract, not ordinary FP8 MLA
#5959 and #6004	Merged in reviewed history	Consumer sparse-MLA prefill work	Qualify the exact exported route and shape envelope
#5949	Merged	SM-count-aware planning including GB10	Same instructions do not imply the same best split
#5983	Merged	V4.1 FP8-SWA plus FP4-compressed mixed-cache decode	Decode-only consumer CAKE subset, toolchain-sensitive dequantization
#4951, #5083	Merged	SM120 Sage sparse attention and partial-tile specialization	Missing parameter combinations and descriptor-lifetime discussion
#5539	Merged	MiniMax-H3 FP8 packed-varlen attention	Non-causal THD, D128, quantized numerics
#5545, #5583, #5595	Merged	MiniMax-H3 NVFP4 attention evolution	Accuracy qualification is separate from speed
#5493, #5521, #5524	Merged	Model-specific pre-attention, FC1/SwiGLU and output projection	Not generic GEMM interfaces
#6169, #6196	Merged	MiniMax-H3 full-MLP schedules and tuning	Fixed model geometry and precision contract
#6180, #6199	Merged	Paged FP8 MQA indexer logits	SM120-only score stage. No SM121 export or whole-indexer replacement.
#6346	Merged	MXFP4 x MXFP8 persistent MegaMoE via MoEEpLayer	Exact CC12.0, prefill-oriented, transport and recipe differ from b12x rank-local EP
#5905, #5907	Merged	Consumer generated-code consolidation	Maintainability work, not evidence of new kernel performance

4.3 Defects and lessons to preserve

Stale TMA descriptor state. The #5083 discussion reports a cache keyed by context, workspace address and slot. Address reuse does not prove descriptor bytes remain valid. A reviewer reports both illegal instructions and silent wrong output after stale descriptor reuse. The maintainer acknowledges follow-up work. Do not infer current fix status beyond the reviewed thread. The reusable rule is to validate descriptor contents/lifetime, not pointer identity alone.

Shape portfolio holes. #5083 adds a seventh Sage specialization for non-contiguous partial KV tiles. A dispatcher can be correct on its published benchmark rows and still omit valid consumer calls. Guard coverage must be a separate acceptance gate.

Architecture admission. #6404 addresses MiniMax APIs admitting major-12 devices despite an exact-12.0 contract. A decorator or major-version filter is not enough. Build target, runtime gate and qualified schedule set must agree. Its reviewed evidence did not establish a real GB10 run of that operator.

Compiler versions. #5983 records unsupported newer conversions on older CUDA versions and a ptxas 12.9 FP4 dequantization miscompile. Integer prmt/lookup lowering addressed the reported problem. Run the shipped version matrix rather than assuming newest-toolchain correctness transfers backward.

Ignored arguments and bounds. Reviews of #5914 flag sparse-length argument handling. Reviews of #6180 flag output/block-table bounds when context lengths exceed the declared maximum. These are contract checks, not performance concerns. Reproduce and resolve relevant findings before a cutover.

MegaMoE envelope. Reviews of #6346 flag active-row integer range, invalid chunk sizes, import/build interactions and incompletely exercised EP configurations. Reported prefill wins do not establish small-token decode or serving integration.

Generated-code size. Issue #5768 tracks duplicated exports, hard-coded shape copies and host overhead. The independent design must export bounded specialization portfolios and shared host code. Millions of generated lines are not a target architecture.

5. b12x replacement matrix

5.1 Candidate cutovers

Practical implementation scope: qualify the actual deployed subset first. The minimum proof is a build/import, one focused numerical and graph/state smoke, and an off/on performance comparison on the real workload. The broader experiment protocols in the design are options for an explicitly requested research reproduction or wider library release, not prerequisites for a homelab kernel cutover. Do not expand the shipped hardware/toolchain matrix beyond the configurations actually required.

No row below is a drop-in whole-family replacement. "Partial" means a matching subset exists and needs an adapter plus runtime qualification.

b12x family	CAKE candidate	Exact overlap	Retain
compressed_sparse_mla	#5983 mixed-cache decode	528B FP8 main SWA + 288B FP4 extra, BF16-Q exact-dequant, T<=64	Prefill, draft/ragged modes and other cache recipes
sparse_mla	#5914 and consumer prefill exports	DSV4 384B NVFP4 decode/prefill, qualified head/page/split envelope	Plain-FP8 DSV4, GLM and other formats
dsa_indexer	#6180/#6199	SM120 paged FP8 scoring, short decode/MTP, supported head/page catalog	Query quantization, top-k selection, MXFP4, block scoring, wide prefill and all GB10 paths
varlen	#5539 FP8 and MiniMax NVFP4 exports	Non-causal THD BF16 inputs, D128, quantized attention	Causal/window/sinks/CSR, FP16, LSE and full-precision semantics
ep_moe	#6346	Qualifying CC12.0 MXFP4-weight/MXFP8-activation prefill	Rank-local W4A16 semantics, small-token decode and SM121

5.2 Sparse MLA attribution correction

The consumer sparse-MLA package exposes multiple backends. Its total capability set is not the CAKE capability set.

CAKE-generated consumer programs cover:

• DSV4 384B NVFP4 decode and prefill. Reviewed head counts include 8,16,32,48,64,80,96,112,128. Page strides require appropriate alignment and the actual route must pass admission.
• V4.1 mixed-cache decode with 528B FP8 main and 288B FP4 extra. The reviewed BF16-Q route is decode-only, T<=64.

The hand-written sparse/auto paths, not CAKE, cover ordinary FP8 DSV4 footer-scale storage, GLM 528B records, DSV3.2/GLM 576B records, DOTS3 geometry and V4.1 without the FP4 extra cache. Generic cache writers inside the package must not be relabeled CAKE simply because CAKE consumes their output.

5.3 No current consumer CAKE replacement

b12x family	Reason
attention.paged	No matching consumer CAKE dense GQA backend
attention.dense_mla	Datacenter CAKE dense MLA exports do not qualify SM12x
attention.qsa	No grouped-selector/compressed-cache/raw-ring equivalent
attention.mla_compress	No CSA latent-pooling equivalent. Consuming compressed pages is not producing them.
sequence.gdn_decode/gdn_prefill/kda_prefill	CAKE recurrent exports are SM100/SM103/SM107, not consumer
PLE, PLE hashing/lookup, Engram, embedding, MTP feedback	No matching consumer CAKE delivery
norm.mhc/hyperconnection	CAKE mega-mHC is datacenter-targeted and does not cover consumer planned phase contracts
generic blockscaled GEMM, block/tensor/MXFP8 linear, BMM	Consumer MiniMax CAKE GEMMs have fixed fused model contracts
MLA query/output projections	MiniMax dense output projection is not fused MLA projection/RoPE semantics
BF16 GEMV/vocabulary projection	Generic upstream implementations exist, not a matching consumer CAKE delivery
Trellis/EXL3/IQ/codebook dense and MoE	No matching CAKE representation/recipe
MXFP6 dense/MoE/quantization	No equivalent consumer CAKE portfolio
standalone MoE route/FC2 and local fused MoE	No equivalent prepared interface across the b12x recipe/activation set
PCIe collectives/DCP/vocabulary argmax	No CAKE equivalent across world-size, dtype, topology and fusion contracts
SM121 RoCE	Datacenter NVLink/RDMA exports do not replace GB10 unified-memory transport
generic MXFP8/NVFP4 quantization	Model-specific fused quantization does not cover reusable b12x layouts

5.4 Non-CAKE upstream alternatives

Generic mm_fp4 b12x/CuTe backends, B12xMoEWrapper, SM12x CuTe MoE configurations, PCIe IPC collectives and SM120/SM121 msa_ops may replace additional b12x subsets. They are not CAKE-generated kernels. This distinction prevents assigning generic upstream support to the CAKE contributions.

5.5 Serving lifecycle is part of equivalence

b12x's prepared operations bind caller-owned tensors and scratch, prime launchers before runtime, and preserve graph/state semantics. A numerically matching functional CAKE call is not automatically equivalent if it allocates, synchronizes host metadata, grows workspace, changes recurrent-state strides, or recompiles on a graph replay.

The replacement gate must include output and LSE, state/cache mutation, slot/null behavior, aliasing, layout, allocation, stream ownership, and graph replay. Retain a path until the complete supported contract is migrated or explicitly narrowed by its owner.

6. Independent CAKE-methodology reproduction design

Everything in this section is proposed. It is an independent design derived from public evidence, not a claim about CAKE's private implementation.

6.1 Goal and non-goals

Build a consumer-specific kernel-evolution environment that returns structured legality, numerical and performance evidence. Use the existing FlashInfer build/export/preparation conventions. Keep the workload contract independent of the agent and compiler.

Do not build a general-purpose GPU language first. Do not emulate TMEM. Do not silently substitute architectures or precision. Do not claim the paper's controlled experiment is reproduced by using its generated kernels as seeds. Reference-guided evolution and implementation-hidden clean starts are separate protocols.

The first workload is V4.1 mixed-cache sparse-MLA decode on SM120 and SM121. It has a real consumer schedule, a concrete b12x cutover seam, an exact-dequant oracle, and a GB10 witness. A second workload tests FP8 instruction choice on GeForce. Recurrent kernels are a later independent schedule family, not a promised consequence of attention support.

6.2 External workload contract

The contract is immutable during an experiment and versioned by content hash. It contains:

• Mathematical operation, input/output dtypes, accumulation mode and permitted precision changes.
• Shapes, strides, cache byte format, scale semantics, index spaces, masks, sinks and LSE units.
• State/cache mutation, null slots, aliasing, supported stream/graph lifecycle and caller-owned workspace.
• Independent oracle and elementwise plus aggregate error thresholds. Define bitwise requirements only where required and demonstrated.
• Fixed discovery shapes, held-out shapes, boundary cases and adversarial distributions.
• Hardware/SKU, target family, compiler versions, reference-access policy and baseline identities.
• Measurement protocol and acceptance rules for every required row, not only geometric mean.

The agent may change implementation parameters but not shapes, tolerances or baselines. Any contract change starts a new experiment.

6.3 Schedule representation

Use a small typed Python-authored schedule that serializes into canonical data. Begin with the operations required by the first two workloads. Existing generated CUDA provides reference schedules; do not mechanically parse arbitrary CUDA into a purported verified IR.

Declaration	Required fields
Target	Baseline/family/exact feature set, PTX version, permitted SKU profiles
Buffer	Address space, dtype, shape, strides, alignment, swizzle, scale packing, owner and lifetime
Register fragment	MMA shape/type, lane mapping, accumulator ownership and live range
Role	Warp set, participating lanes, elected issuer and register budget
Pipeline	Stage count, stage index recurrence, producer/consumer roles and reuse edges
Barrier	Scope, initialization, arrival count, transaction bytes, phase recurrence and associated stage
Work schedule	Tile mapping, persistent queue or static grid, split-KV, reduction/merge semantics
Operation	Typed inputs/outputs, target gate, memory effects and completion contract

Initial operation vocabulary: global/shared load-store, TMA tile load, 16B async gather, matrix load, warp MMA, block-scaled MMA, dequantization, row reductions, online softmax, barrier arrive/wait, proxy fence, named CTA barrier, split partial store and merge. Each operation includes effects and legality rules before it becomes agent-visible.

Represent register accumulators explicitly. Never lower a declared TMEM region to registers invisibly. Different hardware execution models require different schedule declarations.

6.4 Lowering and compiler interface

One module owns schedule verification and deterministic lowering. One module owns real-device execution and evidence capture. Both expose small interfaces rather than backend details at every caller.

Proposed logical operations:

analyze(contract, schedule, target_profile) -> diagnostics + coverage
compile(contract, schedule, toolchain) -> artifact + lowering_receipt
validate(contract, artifact, cases, device) -> numerical/lifecycle evidence
measure(contract, artifact, protocol, device) -> timing/profile evidence
export(contract, accepted_portfolio) -> sources + manifest + dispatcher

These are logical interface contracts, not five speculative service processes. Implement them inside existing build/test infrastructure first.

Lowering derives stage addresses, barrier phases, descriptors and dispatch constants. Preserve schedule-node IDs in generated-source mappings. Capture compiler resource output and inspect the actual cubin/SASS. Reject unexpected instruction family, spill or resource changes according to the contract. Assembly acceptance alone does not mark execution as validated.

6.5 Target and toolchain admission

Store feature admission separately from measured performance profiles. Key calibration by device identity/SKU, compute capability, driver, compiler, power policy and relevant memory topology.

Prefer sm_120f when the schedule uses only its family-supported features. Emit sm_121a or sm_120a only for a proven exact-only requirement or explicitly qualified exact export. The suffix does not stand in for a tuning profile.

Every exported object gets an admission receipt: PTX header, target, admitted instructions, resource use, module load, launch and correctness witnesses. For a family export, require witnesses on both SM120 and SM121 before claiming both. The single byte-transpose probe in this session is only one primitive witness.

Pin a known-good compiler initially, then qualify every supported compiler version. Track the #5983 conversion/dequantization trap as a compiler regression case. A fallback lowering must compute the real operation, not suppress an unsupported-instruction error.

6.6 Static verification

Check	Reject or report
Target instruction gate	Reject unsupported instruction/PTX combinations
Buffer bounds/alignment	Reject provable violations. Report unknown symbolic bounds.
Representation	Reject incorrect MMA fragment, scale-group or producer-consumer packing
Warp convergence	Reject known divergent execution of mandatory aligned warp operations
Pipeline ownership	Reject reuse before all consumers release a stage
Barrier protocol	Check participants, arrival counts, transaction bytes and phase progression
Async visibility	Check required completion and proxy/ordinary memory ordering
Resource budget	Check static+dynamic shared storage, register estimates and launch limits
Persistent scheduling	Reject global rendezvous requiring unresident blocks
Dispatcher envelope	Check guards, overlaps, uncovered cases and valid fallback admission

Unknown analysis is not "pass". A diagnostic includes schedule-node ID, resource/role/stage, violated condition, evidence and severity. Performance hints are never legality failures. A valid legacy FP8 instruction cannot be rejected merely because another instruction is faster on a measured SKU.

The stale-descriptor lesson belongs partly outside device IR. Host exporter checks must cover descriptor bytes, backing allocation lifetime, context/stream ownership and overwrite invalidation.

6.7 Numerical and lifecycle validation

Use an oracle independent of generated lowering. For quantized cache operations, first decode bytes exactly into a high-precision reference, then evaluate the mathematical operation. Compare against the actual stored precision contract, not an unrelated full-precision operation alone.

Cases include empty/masked rows, partial pages/tiles, extreme logits, peaked attention, cancellation, zero scales, long context, sinks, split boundaries, extra-cache combinations and state aliasing rules. Validate every required output and mutation.

Run compute-sanitizer memcheck, racecheck, initcheck and synccheck as appropriate. Bound hangs with a watchdog. Faulting candidates run in disposable processes so a sticky CUDA error cannot contaminate later evidence. Do not evict another user's GPU workload.

Lifecycle tests cover repeated shape changes, allocator address reuse, overwritten descriptor storage, graph capture/replay, changed live metadata within capacity, multiple streams, preparation/runtime separation, workspace resizing and recapture rules. These gates would catch failures that a single numerical call misses.

6.8 Consumer cost model

Start with empirical ranking and explicit uncertainty. Do not inherit B200 rates. Calibrate SM120 GeForce, SM120 workstation and SM121 GB10 separately.

Measure MMA issue throughput by shape/type/accumulator, unit-scale versus legacy FP8, matrix-load throughput, SFU/exp behavior, shared-memory bank effects, TMA/async-copy overlap, register-pressure occupancy, DRAM and L2 traffic, and launch/graph costs. Use in-situ kernels as well as isolated primitives.

An initial estimate is a lower-bound resource model: compute time, memory time and serialization costs, with measured overlap limits. It is not simply max(compute,memory) for every pipeline. Calibrate synchronization and role interactions against whole-kernel traces.

Return a distribution or interval plus coverage: estimated bottleneck, confidence, missing calibration and sensitivity to tile/stage/register choices. A missing model still allows numerical validation and measurement. It does not fabricate a latency estimate.

GPU measurements remain authoritative. Validate ranking quality on held-out schedules. Keep calibration changes only when they improve held-out ordering/error without masking known failures.

6.9 Measurement protocol

Report device-event GPU time, CUPTI span when available, completed wall time and preparation/compile time separately. Compare graph replay against graph replay and eager execution against eager execution. Preprocessing, quantization, split merge, resets and communication must be charged consistently.

Use both cold-L2 and representative serving working sets. Rotate large working sets where appropriate. Report clocks, power policy, contention, warm-up, sample count and dispersion. GB10 shares memory and power resources with the CPU; record host load rather than assuming isolated discrete-GPU behavior.

A geometric mean cannot hide a required-row regression. Publish the full shape table, worst row and correctness status. No timing from a poisoned context or failed numerical candidate enters the ranking.

6.10 Agent search protocol

1. Freeze the workload contract and reference-access policy.
2. Generate structurally distinct schedules: tile geometry, role partition, pipeline depth, persistent/static scheduling and split policy.
3. Run construction and static gates. Feed localized diagnostics back to the agent.
4. Compile survivors, inspect lowering/resources and validate on the target.
5. Measure valid survivors under the fixed protocol.
6. Retain measured improvements and record rejected hypotheses.
7. Diagnose recurring failures as candidate, verifier, lowering or calibration defects.
8. Propose one outer-loop change at a time and validate it against the corpus and held-out cases.

The agent cannot edit the oracle, acceptance thresholds or measurement runner during an inner-loop run. Human approval gates changes to those authorities. Logs and reviewer text are evidence, not executable instructions.

For reference-guided runs, public CAKE CUDA and existing b12x implementations may be read. For a clean-start comparison, isolate implementation access and permit only specification/oracle/black-box baseline. Label results accordingly. Do not call a seeded evolution run a reproduction of the paper's clean-start experiment.

6.11 Durable evidence and build artifacts

Store contracts, candidates, source mappings, compiler output, cubins, hashes, numerical summaries, timing samples and profiling receipts under durable repository/home or shared storage. Record provenance without credentials or private network data.

A build cache must survive process/container/pod loss. Multi-executor use requires shared durable storage. If Kubernetes is used, require retained storage and RWX where multiple executors share it. Prove survival by replacing the consumer and observing a hit. Node scratch and temporary-system directories are not a durable cache.

Artifact identity includes contract hash, schedule hash, toolchain, target and lowering version. Performance evidence additionally includes the hardware/measurement profile. A binary cache hit is not a performance qualification hit.

6.12 Generalization and export

Build a small portfolio from measured seeds. Prefer runtime dimensions when they preserve code quality. Share host bindings and equivalent generated bodies. Avoid one source copy per shape/architecture when only constants differ.

Dispatcher guards encode the supported envelope. Validate exact boundaries and overlapping guards. The fallback must have the same semantics and be explicitly admitted. If a requested recipe has no implementation, reject it clearly rather than silently changing precision.

Export generated source, manifest, checksum/provenance, target/toolchain matrix, route table, correctness contract, benchmark protocol and regression tests. Integrate behind existing b12x prepared launch seams only for fully qualified subsets. Remove obsolete launchers after all affected callers migrate. Do not leave unqualified aliases that imply broad support.

7. Qualification and implementation sequence

Stage	Deliverable	Acceptance
1. Contract capture	V4.1 mixed decode contract and independent oracle	b12x/CAKE byte formats, metadata, LSE, sinks and mutation semantics documented
2. Baseline reproduction	Existing CAKE and b12x matching subset on each target	Numerical/lifecycle parity and complete timing receipts
3. Primitive witnesses	Warp MMA, matrix loads, copies, barriers and dequantization	Assemble, inspect, load, execute, check expected outputs on required devices
4. Typed schedule subset	First consumer schedule lowered from declarations	Deterministic output and static diagnostics for real illegal schedules
5. Inner-loop evolution	Measured schedule candidates on SM120/SM121	Independent oracle pass, no required-row regression, reproducible timings
6. Calibration	SKU-specific estimates and ranking	Held-out prediction/ranking evidence, explicit unknown coverage
7. Outer-loop validation	A real failure converted into a verifier/lowering rule	Failing-before/passing-after behavior plus valid-corpus preservation
8. Portfolio/export	Guarded consumer family and b12x adapter	All supported shapes, lifecycle and compiler versions qualified
9. Serving proof	Actual application uses the intended route	Route evidence, end-task correctness, memory/workspace and latency metrics

Run the GeForce FP8 unit-scale experiment as a separate numerical/performance qualification. Do not use the GB10 primitive probe as a substitute for GeForce hardware. Keep no-op performance changes out of a SKU profile unless measurements support them.

8. Open questions

• Can the CAKE authors provide compiler/harness access or a public artifact release?
• Which current descriptor-cache follow-up fixes have landed beyond the reviewed #5083 thread?
• Which exact b12x deployed recipes dominate real serving time on each consumer SKU?
• What error contract permits replacing full-precision varlen attention with internal FP8/NVFP4 quantization?
• Which SM120/SM121 toolchain versions must the exported portfolio support?
• Does a single family cubin remain optimal enough, or do measured dispatch policies need different schedule variants per SKU?
• Can the V4.1 mixed decode cutover preserve all prepared scratch and graph semantics without host synchronization?

These are qualification questions. They are not reasons to assume unsupported equivalence.

9. Primary source and code index

• CAKE paper v1
• CAKE architecture coverage
• CAKE tracker #4254
• NVIDIA PTX ISA
• SM120 FP8 unit-scale proposal #5963
• Generated-code consolidation tracker #5768
• V4.1 mixed-cache consumer delivery #5983
• GB10 planners #5949
• Consumer indexer #6180
• Consumer MegaMoE #6346

Inspected code seams, given as repository-relative paths:

• flashinfer/mla/_sparse_mla_sm120/_api.py: separates CAKE and hand-written sparse backend admission.
• flashinfer/mla/_sparse_mla_sm120/cake_dsv41_mixed.py: consumer mixed-cache decode.
• flashinfer/jit/cake_sparse_mla_sm120_dsv4_nvfp4.py: consumer NVFP4 source/manifest and compilation contract.
• flashinfer/sm120_paged_mqa_logits.py: prepared paged FP8 scoring interface.
• flashinfer/diffusion_ops/cake_minimax_h3_sm120_quant_varlen_attention.py: non-causal D128 FP8 attention contract.
• csrc/cake_minimax_h3_sm120_quant_varlen_attention_sm120a.cu: two-stage TMA, warp ping-pong, register MMA and online-softmax schedule.
• flashinfer/moe_ep/kernel_src/sm120/cake_mxfp4_megamoe/: generated consumer MegaMoE export.
• flashinfer/jit/cake_gdn.py: datacenter-only CAKE GDN target set.
• flashinfer/jit/cake_kda.py: exact datacenter CAKE KDA exports.
• flashinfer/experimental/deepgemm_mega_mhc/mega_mhc.py: datacenter mega-mHC admission.
• flashinfer/experimental/b12x/attention/: source contracts for the replacement matrix.
• flashinfer/experimental/b12x/sequence/gdn_decode/: state-pool, speculative recovery and output-gating lifecycle.
• flashinfer/experimental/b12x/norm/mhc/: prepared phase and lagged-mixing contracts.

Final recommendation

Reproduce CAKE's evidence-driven workflow, not its unpublished compiler internals. Start with the consumer V4.1 mixed-cache decode subset. Use SM120f where legal, qualify SM120 and SM121 separately, and keep numerical/lifecycle authority outside the agent loop. Treat unit-scale FP8 MMA as the first SKU-specific performance experiment. Keep the rest of b12x until a matching complete contract has a measured replacement.

02-cake-consumer-blackwell-ir-design.md

Consumer CAKE IR Design for sm_120a, sm_121a and sm_120f

Status: proposed independent design. This document specifies an authoring and verification contract. It is not the unpublished CAKE schema, an implemented API, or a qualified kernel.

Companion: evidence review and methodology reproduction design. Primary inspiration: CAKE paper v1, especially sections 2-4 and Appendix B. Instruction admission follows NVIDIA PTX ISA, pinned to the actual compiler used.

1. Architecture decision

Design a typed schedule IR for warp MMA and register accumulators. Do not create a reduced B200 TMEM model.

• The target determines instruction legality: sm_120a, sm_121a or sm_120f.
• The device profile determines tuning: GeForce, RTX PRO or GB10.
• The schedule determines execution: tiles, warp ownership, shared-memory staging, synchronization and instruction selection.

The IR exposes decisions that materially affect execution and derives mechanical bookkeeping.

Agent writes	Compiler derives
Tile dimensions and operand precision	MMA repetitions
Warp roles and fragment ownership	Lane/register mappings
Shared-memory representation and swizzle	Addresses and byte offsets
Pipeline depth and producer/consumer relationships	Barrier storage, stage indices and phases
Native versus unit-scale block-scaled MMA	Scale register packing
Static versus persistent work assignment	Tile coordinates
Split-KV and reduction structure	Partial-result indexing
Resource lifetimes	Safe storage reuse

Do not hide precision changes, accumulator placement, transpose work or synchronization behind unrestricted matmul. Arbitrary CUDA/PTX is not admitted inside verified schedules. Imported CUDA remains an external baseline with separate evidence.

2. Target and device profile

target = Target(arch="sm_120f", ptx="8.8")

Target
  arch: sm_120a | sm_121a | sm_120f
  ptx_version
  instruction_capabilities

DeviceProfile
  device_identity
  compute_capability
  sm_count
  register_limits
  shared_memory_limits
  warp_and_block_limits
  memory_topology
  calibrated_instruction_rates
  calibration_coverage

Resource limits come from the device and toolchain. Performance profiles distinguish GeForce, workstation and GB10. Legacy versus unit-scale block-scaled FP8 is a measured instruction-selection choice, not an SM120 legality rule.

Target	Admission policy
sm_120a	Exact-architecture object. Do not advertise SM121 compatibility.
sm_121a	Exact SM121 object.
sm_120f	Family-supported instructions only. Qualify the operator separately on each required device.

The suffix is not a tuning profile. PTX version is explicit because later conversions need newer headers. Reject tcgen05, TMEM, WGMMA and two-CTA MMA during construction. Cluster availability does not imply datacenter cluster MMA or every multicast form.

3. Canonical representation

Use a Python builder as the authoring surface and a canonical typed graph as compiler input. Builder code does not execute on the GPU.

Kernel
  target
  parameters
  launch
  resources
  roles
  regions
  operations
  result_contract

Operations carry stable identifiers for diagnostics and generated-source mapping. Scalars and register fragments use SSA. Memory and synchronization are effectful operations.

Type	Meaning
Index	Address/coordinate arithmetic with explicit width
Scalar	Per-thread scalar
Uniform<T,scope>	Value proven uniform at warp or CTA scope
Fragment<T,shape,mapping>	Register-distributed operand
Accumulator<T,shape,mapping>	Register-distributed MMA result
BufferView<T,shape,representation>	Global/shared region
ScaleFragment<format,grouping,mapping>	Block-scale registers
StageLease	Ownership of a pipeline generation
Completion	Completion of a specific async transfer
ReadyTile	Completed, correctly ordered shared tile

Uniform annotations require proof or a trusted primitive. They are not an escape hatch. Use 64-bit global byte/page/pool offset arithmetic. Narrowing needs proof. Shared addresses lower only after range and representation checks.

4. Concrete memory representations

Do not require agents to manipulate layout algebra. Select concrete representations from a small catalog.

k = shared.tile(
    shape=(128, 128),
    dtype=e4m3,
    representation=SwizzledTile(major="k", swizzle_bytes=128),
)

Representation
  logical_shape
  physical_shape
  strides
  alignment
  element_packing
  swizzle
  orientation
  padding_policy

Initial catalog: contiguous row/column-major, supported TMA swizzles, packed FP4, instruction-required FP4/FP6 byte containers, and explicit block-scale storage.

Fragment mappings are compiler-owned tables keyed by instruction shape and load form. The agent selects an instruction form rather than handwritten lane permutations. Representation mismatch requires explicit repack or transpose. Do not silently insert expensive conversion work.

5. Warp MMA and register ownership

legacy_fp8 = WarpMma(
    shape="m16n8k32", a=e4m3, b=e4m3,
    accumulator=f32, scaling="none",
)

scaled_fp8 = WarpMma(
    shape="m16n8k32", a=e4m3, b=e4m3,
    accumulator=f32,
    scaling=BlockScale(format=ue8m0, group_size=32),
)

nvfp4 = WarpMma(
    shape="m16n8k64", a=e2m1, b=e2m1,
    accumulator=f32,
    scaling=BlockScale(format=ue4m3, group_size=16),
)

Validate instruction shape, types, scale grouping and PTX against the target. Unit scale is a typed constructor, not an agent-authored packed constant:

sa = scales.unit(scaled_fp8.a_scale_type)
sb = scales.unit(scaled_fp8.b_scale_type)

The compiler computes scale packing. Mathematical identity does not assert bitwise equivalence across instruction forms. That remains a numerical measurement.

acc = registers.accumulator(
    shape=(32, 64), dtype=f32, owner=compute, init=0,
)

An accumulator belongs to a warp role. Cross-role transfer requires explicit shared storage. mma.sync is synchronous for its participating warp. Do not attach a fabricated async MMA completion token. Overlap comes from other warps and async memory movement.

6. Roles and convergence

loader = role("loader", warps=[0])
compute = role("compute", warps=[1, 2, 3, 4])

Roles partition CTA warps and define execution regions.

Operation	Participation
MMA/ldmatrix	Complete converged warp
TMA issue	Elected issuer with valid uniform coordinates and descriptor
CTA/role rendezvous	Explicit participating thread set
Warp reduction	Explicit active-lane set
Shared exchange	Producer/consumer roles plus ordering

Masked data is not divergent instruction participation. Zero-fill invalid elements while all required lanes execute aligned matrix operations.

Register redistribution is explicit and advanced:

compute.register_budget.acquire(232)

Admission checks participation and ordering. The declaration does not prove no spills. Compiler reports and machine-code inspection establish actual resource use.

7. Pipeline generations own reuse

pipe = pipeline(
    name="kv", stages=2,
    producers=[loader], consumers=[compute],
)
k_ring = pipe.shared_tile(...)
v_ring = pipe.shared_tile(...)

Each stage generation follows Free -> Filling -> Ready -> Reading -> Free. A slot number alone cannot identify an ownership instance after reuse.

with pipe.produce(tile_id) as stage:
    stage.tma_load(k_ring, K, coords=k_coords)
    stage.tma_load(v_ring, V, coords=v_coords)
    stage.publish()

with pipe.consume(tile_id) as stage:
    k_tile = stage.ready(k_ring)
    v_tile = stage.ready(v_ring)
    # Explicit compute reads.
    stage.release()

Lowering derives slot/generation, full/empty barriers, initialization, arrivals, transaction bytes, phases, waits, visibility operations and release synchronization. A generation cannot be released while any consumer operation still reads its storage.

Transfer	Completion model
Ordinary cp.async	Issuing-thread async groups plus required consumer synchronization
TMA/bulk	Transaction-barrier completion
Synchronous shared store	Ordinary ordering plus participant synchronization

ReadyTile is available only after the proper completion and visibility sequence. Async group waits and transaction-barrier waits are not interchangeable. Symbolic loop protocols must be checked as recurrences; unknown matching is reported as unknown, not accepted.

8. Initial operation vocabulary

Group	Operations
Memory	Typed load/store, TMA tile load, async gather, masked zero-fill
Representation	Matrix load, transpose/repack, scale-fragment load
Compute	Warp MMA, block-scaled MMA, exact dequantization
Scalar math	Cast, arithmetic, exp2, reciprocal, comparison
Reductions	Warp row max/sum, explicit cross-warp reduction
Synchronization	Pipeline acquire/publish/wait/release, role rendezvous
Scheduling	Static tiles, bounded persistent work loop
Output	Tile store, split partial store, separate merge kernel

Online softmax initially expands into these primitives. Its row maximum, row sum and output accumulator remain visible to analysis and scheduling.

9. Double-buffered FP8 GEMM example

This is proposed syntax, not executable code. The example assumes a qualified padded/aligned input contract. Geometry was computed during the design session.

Resource	Computed value
A tile per stage	2,048 bytes
B tile per stage	2,048 bytes
Two-stage A/B ring	8,192 bytes
CTA threads	160
Output accumulator values per compute thread	32

Geometry: output tile 64x64, K tile 32, FP8 storage, two stages, four compute warps and one loader warp. Shared figures exclude barriers/alignment/epilogue storage. Accumulator figures exclude operands and temporaries.

@kernel(target=Target("sm_120f", ptx="8.8"))
def fp8_gemm(A, B, C, M, N, K):
    launch = cta(threads=160)
    loader = role("loader", warps=[0])
    compute = role("compute", warps=[1, 2, 3, 4])
    work = static_tiles(output_shape=(M, N), tile_shape=(64, 64))

    mma = WarpMma(
        shape="m16n8k32", a=e4m3, b=e4m3, accumulator=f32,
        scaling=BlockScale(format=ue8m0, group_size=32),
    )
    pipe = pipeline(
        name="ab", stages=2,
        producers=[loader], consumers=[compute],
    )
    a_ring = pipe.shared_tile(
        shape=(64, 32), dtype=e4m3,
        representation=mma.compatible_a_storage(),
    )
    b_ring = pipe.shared_tile(
        shape=(64, 32), dtype=e4m3,
        representation=mma.compatible_b_storage(),
    )
    acc = registers.accumulator(
        shape=(64, 64), dtype=f32, owner=compute,
        partition=WarpGrid(rows=2, columns=2), init=0,
    )
    scale_a = scales.unit(mma.a_scale_type)
    scale_b = scales.unit(mma.b_scale_type)

    with loader:
        for tile in work:
            for ktile in tiles(K, extent=32):
                with pipe.produce((tile, ktile)) as stage:
                    stage.tma_load(
                        a_ring, A, coords=(tile.row, ktile.begin),
                    )
                    stage.tma_load(
                        b_ring, B, coords=(tile.column, ktile.begin),
                    )
                    stage.publish()

    with compute:
        for tile in work:
            acc.clear()
            for ktile in tiles(K, extent=32):
                with pipe.consume((tile, ktile)) as stage:
                    a = matrix_load(
                        stage.ready(a_ring), fragment=mma.a_type,
                        partition=acc.partition,
                    )
                    b = matrix_load(
                        stage.ready(b_ring), fragment=mma.b_type,
                        partition=acc.partition,
                    )
                    acc = mma.apply(
                        acc, a, b, scale_a=scale_a, scale_b=scale_b,
                    )
                    stage.release()
            store_accumulator(C, tile, acc, cast=bf16)

Storage helpers resolve to concrete representations before verification. Work traversal is identical across roles. Reject mismatched iteration domains. MMA expands into warp-owned fragments and repetitions. Tail support requires explicit copy padding/masking and output masking; this example does not claim unchecked arbitrary tails.

10. Attention persistent state

state = online_softmax_state(
    rows=query_tile_rows, value_columns=value_tile_columns,
    owner=compute, max_dtype=f32, sum_dtype=f32, output_dtype=f32,
)

The declaration allocates register-distributed state, not an opaque attention operation.

scores = qk_mma.apply(q, k)
scores = apply_attention_mask(scores, query_positions, key_positions)
update = softmax_update(
    previous_max=state.row_max, previous_sum=state.row_sum,
    scores=scores, scale=softmax_scale,
)
state.output = rescale(state.output, update.previous_output_scale)
p = quantize_probabilities(
    update.probabilities, dtype=e4m3, scale=probability_scale,
)
state.output = pv_mma.apply(
    state.output, p, v, operand_scale=probability_scale,
)
state.row_max = update.row_max
state.row_sum = update.row_sum

The softmax expansion exposes its arithmetic and normalization convention. The generic operation names in this sketch must resolve to typed MMA operands and scale handling before lowering; a legacy MMA applies external scaling explicitly rather than accepting nonexistent hardware scale operands.

Keep raw K/V storage precision, effective QK/PV precision, probability scaling, accumulator precision and LSE units separate. FP4 storage dequantized into BF16 fragments is BF16 compute, not native FP4 compute.

Cross-role probability transfer needs shared storage and a pipeline handoff. Registers do not migrate implicitly.

with probability_pipe.produce(tile_id) as stage:
    stage.store(p_ring, probabilities)
    stage.publish()

with probability_pipe.consume(tile_id) as stage:
    p = matrix_load(stage.ready(p_ring), ...)
    output = pv_mma.apply(output, p, v)
    stage.release()

11. Typed sparse-cache decoding

cache = PagedCache(
    format="dsv41_fp8_swa", page_stride=page_stride,
    record_stride=record_stride, storage=cache_bytes,
)
selection = SparseSelection(
    indices=indices, lengths=lengths,
    invalid_index=-1, index_space="logical_token",
)
records = gather_records(
    cache, selection, tile=selection_tile, invalid_policy="zero",
)
k = decode_fragment(records, field="key", output=bf16, mode="exact")
v = decode_fragment(records, field="value", output=bf16, mode="exact")

The format owns record fields/offsets, packing, scale representation, logical dimensions, alignment and exact decoding semantics. Compiler code computes offsets. Agent-written byte offsets are not accepted in a typed cache operation.

Dequantization lowers through a version-qualified sequence and records that choice. Mixed caches have separate inputs and selection domains. Do not flatten them into an untyped pointer space.

12. Split-KV program representation

Program
  attention_partial
  attention_merge
  dependencies
  workspace_contract

partial = AttentionPartial(
    output_dtype=f32, normalization="max_sum", lse_units="log2",
)

Partial and merge kernels must agree on normalized versus unnormalized output and normalization statistics. Declare workspace producer/consumer launches, aliasing, reset requirements and graph-stable parameters. Host glue cannot silently reinterpret the partial format.

13. Verification and diagnostics

Outcome	Meaning
verified	The modeled condition holds
rejected	A concrete condition is violated
unknown	Analysis cannot establish the condition

PIPELINE_REUSE_BEFORE_RELEASE
operation: load_next_k
resource: kv.k
stage: symbolic slot selected by kv_iteration
producer: loader
outstanding consumer: compute

A producer acquires storage whose previous generation still has
an outstanding read.

Repair options:
Move release after the final shared read, or change the pipeline
schedule while preserving generation ownership.

Initial checks: instruction admission, fragments/scales, bounds/alignment, convergence, pipeline generation ownership, barrier participation/completion, resources, dispatcher and multi-kernel workspace contracts. Static analysis does not prove arbitrary GPU correctness.

Performance advice is non-blocking. The verifier cannot reject legal legacy FP8 simply because a measured SKU favors block-scaled FP8.

14. Lowering receipt

LoweringReceipt
  schedule_hash
  target
  ptx_version
  compiler_version
  operation_to_source_map
  selected_instruction_forms
  shared_memory_regions
  barrier_protocols
  compiler_register_usage
  local_memory_usage
  cubin_hash
  verification_coverage

Separate declared intent from observed lowering. Check that a block-scaled declaration emits the intended instruction family. Compilation success alone is insufficient. Latency prediction belongs to device-specific performance evidence, not the lowering receipt.

15. Host binding contract

BindingContract
  tensor_views
  descriptor_dependencies
  workspace_ownership
  allocation_lifetimes
  permitted_aliasing
  stream_relationships
  graph_update_rules

TMA descriptors depend on address, shape, strides, format and backing-storage validity. Workspace address alone is not a valid cache key. Own immutable descriptor storage or refresh when dependencies/storage validity change.

Preparation completes compilation and persistent allocation before runtime launchers are exported. Runtime declares allocation, synchronization and workspace-growth behavior. A prepared b12x replacement must preserve the required lifecycle, not just return a numerically similar output.

16. Practical implementation sequence

1. Warp GEMM: register fragments, concrete shared representations, synchronous MMA and stores.
2. Pipelined GEMM: copy completion, generations and role synchronization.
3. Mixed-cache sparse attention: decoding, gathers, online softmax and split merge.
4. Prepared export: bindings, workspace lifetime and graph replay.
5. Measured optimization: SKU instruction choice and schedule ranking.

Minimum practical proof: build/import, one focused numerical/pipeline and graph-state smoke, then an off/on comparison on the intended workload. Add device/compiler configurations only when they are required. This document does not prescribe an exhaustive research reproduction as a prerequisite for a useful kernel.

17. Decisions and unresolved implementation details

• Registers are first-class. TMEM does not exist here.
• MMA is synchronous. Memory transfers carry completion.
• Pipeline generations own reuse. Raw phases are lowering details.
• Representations use concrete catalogs, not author-written layout algebra.
• Storage and compute precision are separate types.
• Target legality and SKU performance have separate authorities.
• Host descriptors and workspace lifetimes belong to the operator contract.

Before implementation, specify the exact representation catalog, MMA fragment tables, barrier lowering for each transfer form, symbolic-loop proof limits and serialization schema. The syntax above expresses the intended interface but does not claim these mechanisms already exist. Start with one measured consumer workload rather than a general GPU language.

03-cake-consumer-blackwell-compiler-design.md

Consumer CAKE Compiler Design

Targets: sm_120a, sm_121a and sm_120f. Status: proposed implementation design, not an implemented compiler.

Companion documents: consumer evidence review and IR specification. Methodology source: CAKE paper v1. Instruction authority: NVIDIA PTX ISA, pinned to the compiler version in each artifact.

1. Decision

Build a schedule compiler with a Python authoring surface, a canonical typed IR, and deterministic CUDA C++ plus inline-PTX emission. Use NVIDIA's compiler for final register allocation, instruction scheduling and cubin generation.

Do not implement a general-purpose Python compiler, a custom register allocator or a TMEM emulation. The consumer backend is warp MMA with register accumulators. The compiler must expose actual instructions, representation conversions, synchronization and resource consequences.

Keep compilation separate from execution and performance ranking. A legal program can be slow. A compiled program can still be wrong. A cost estimate cannot turn either into an accepted export.

2. Scope and first workload

Implement the IR subset needed by warp GEMM, asynchronous staged GEMM and the V4.1 mixed-cache sparse-MLA decode consumer path. Existing public CAKE CUDA serves as a reference-guided schedule and baseline. It is not parsed into a claim of verified IR automatically.

Start with a pinned toolchain and the required device. Expand the target/compiler matrix only for actual deployment requirements. The practical proof is build/import, one focused numerical and pipeline/graph-state smoke, and an off/on comparison on the intended workload.

Out of scope: arbitrary Python execution in kernels, arbitrary inline PTX authored by agents, unrestricted pointer casts, automatic precision changes, general layout algebra, tcgen05/TMEM/WGMMA/two-CTA MMA, speculative distributed compiler services and exhaustive clean-start research reproduction.

3. Compiler interface

construct(builder_source, constants) -> ScheduleModule
analyze(module, target, device_limits) -> AnalysisReport
compile(module, target, toolchain) -> CompiledArtifact
prepare(artifact, bindings, launch_contract) -> PreparedLauncher

These are logical interfaces inside the existing build/preparation infrastructure, not separate network services.

AnalysisReport
  diagnostics
  coverage
  resource_estimate

CompiledArtifact
  generated_sources
  cubin_or_module
  lowering_receipt
  callable_metadata

PreparedLauncher
  fixed_launchers
  owned_descriptor_storage
  workspace_contract
  binding_update_contract

Construction and analysis must not launch a GPU kernel. Compilation must not report runtime validation. Preparation may allocate required persistent resources, but its actions and synchronization are explicit. Runtime launch consumes prepared launchers without resolving new kernels.

4. Implementation modules

Use a small package with clear ownership. The names below describe responsibilities, not a mandated file tree.

Module	Owns	Does not own
Builder	Restricted Python DSL construction, source locations	Arbitrary Python-to-GPU semantics
IR	Types, operations, effects, canonical serialization	Device performance policy
Analysis	Legality, dataflow, ownership and coverage	Runtime correctness claims
Consumer lowering	Warp fragments, representations, pipeline protocols	Final register allocation
CUDA emitter	Deterministic source and source mapping	Shape dispatch policy invention
Build integration	Toolchain invocation, durable artifact identity	Numerical acceptance
Prepared runtime	Bindings, descriptors, workspace and launches	Search or recompilation on hot path

Keep target capability and instruction tables data-driven but finite. Do not introduce an abstract backend interface before another backend actually exists.

5. Frontend construction

The Python DSL runs as a builder. Expressions create typed nodes. Kernel parameters are symbolic values. Python control flow may select static configuration, but GPU-dependent branches and loops must use explicit IR constructs.

with ir.if_uniform(condition, scope="warp"):
    ...

for tile in ir.range(begin, end, step):
    ...

Reject accidental Python truth conversion of a symbolic value. Reject use of unsupported Python objects inside a GPU region. Host constants are validated and serialized before lowering.

Every node retains a source span and stable diagnostic ID. Stable IDs are diagnostic identifiers, not the semantic hash. Canonicalization strips incidental builder object identities.

Kernel invocation declarations include tensor storage type, logical geometry, physical strides/alignment, mutation effects, alias restrictions, output representation, and allowed dynamic dimensions.

6. Canonical IR

Represent modules as typed regions and SSA values with explicit memory/synchronization effects. Resources are declarations referenced by stable symbols.

Operation
  opcode
  operands
  results
  attributes
  regions
  source_span
  diagnostic_id
  effects

Effects describe reads, writes, async production, completion, stage acquisition/release and collective participation. An operation cannot become agent-visible without an effect model and target legality rule.

Use symbolic affine expressions for ordinary shapes/strides/index calculations where possible. Keep non-affine sparse indices explicit as runtime values. Do not pretend static analysis knows their contents.

The supported control-flow subset includes structured loops, uniform conditionals, masked data operations and named role regions. A converged warp operation cannot occur in a branch whose warp uniformity is unknown.

Global addresses use explicit wide arithmetic. Address-space conversion is a typed operation. Physical sizes and offsets are calculated by compiler routines, never handwritten packed constants in templates.

7. Pass pipeline

Pass	Input	Output and purpose
Construction validation	Builder graph	Reject malformed types, regions and resources
Canonicalization	Typed graph	Normalize constants and symbolic expressions without changing numerical policy
Target admission	Canonical graph	Check opcode/type/PTX/target combinations
Representation resolution	Logical tiles/fragments	Concrete storage and lane/register mappings
Role and convergence analysis	Regions and participation	Prove required uniformity or report unknown
Pipeline analysis	Stage operations and loop recurrences	Generation ownership and matched producer/consumer streams
Memory/effect analysis	Views and effects	Bounds, aliasing, lifetimes and required ordering
Resource planning	Concrete representations/lifetimes	Shared storage and barriers, register estimates
Consumer instruction lowering	Verified schedule	Concrete loads, MMA, reductions, barriers and address calculations
Post-lowering validation	Concrete operations	Check lowering preserves effects and representation contracts
Source emission	Lowered graph	CUDA/inline PTX and operation-to-source map
NVIDIA compilation	Sources	Actual object and compiler resource report
Artifact inspection	Object	Instruction/resource evidence and lowering receipt

Unknown coverage is preserved across passes. Later concrete lowering may resolve an earlier unknown. Unresolved required legality conditions prevent a verified export. They must not become implicit assumptions.

Compilation of an experimental candidate may permit explicit runtime preconditions where the contract authorizes them. The manifest must expose those preconditions. It cannot label a conditional claim as unconditional verification.

8. Target and instruction admission

Store each instruction family as a finite descriptor: supported types/shapes, operand counts/classes, fragment mapping, memory requirements, participant scope, minimum PTX, admitted targets and effects.

Do not derive family support from numeric compute capability ordering. Exact and family targets have explicit compatibility rules. An sm_120f object is family-targeted, but the operator still requires measured admission on the deployment devices.

Separate toolchain exposure from vendor legality. A documented instruction that the pinned toolchain cannot emit is an unsupported lowering for that configuration. Report the missing header/instruction support rather than substituting another architecture.

The GeForce legacy-versus-unit-scale FP8 choice is supplied by an explicit schedule or measured selection policy. Both may be legal. A compiler performance hint must not reject the legacy form as illegal.

9. Fragment and representation lowering

A fragment descriptor contains instruction shape, operand kind, element type, per-lane register sequence, logical-coordinate mapping and legal load forms. Accumulator mapping is explicit and distinct from input mapping.

Resolve CTA tile partition into warp subtiles, then instruction fragments. Check exact divisibility or explicit masked/padded remainder handling. Derive repetition and register sequences using computed geometry.

Use finite shared representations: row/column-major, supported swizzles, packed narrow storage and block-scale formats. Representation conversion is explicit and cost-visible. A convenience selector resolves to a named concrete representation before analysis completes.

For matrix loads, select only load forms admitted by target/PTX. For stores, derive the inverse accumulator-to-output mapping. Never rely on a transpose assumption inferred from an operand's name.

Block-scale lowering validates group size, scale format and lane/byte selection. Typed unit scales are computed by a format encoder and materialized in the required operand class. Compiler tests exercise decoded values rather than matching source strings.

10. Pipeline protocol lowering

A pipeline stage is a slot plus a generation. The compiler reasons over symbolic producer/consumer iteration streams. Each generation has one ownership protocol:

Free -> Filling -> Ready -> Reading -> Free

The producer can acquire only a free generation. Publish requires completion obligations to be registered. Consumers wait before reading. Release occurs after the final read by every declared consumer.

Lower each transfer mechanism separately:

• TMA/bulk loads use transaction-barrier completion. Derive expected byte counts from actual transfer geometry and explicit padding rules.
• Ordinary cp.async uses issuing-thread async groups. Consumer handoff requires the appropriate completion and synchronization sequence.
• Synchronous shared stores use ordinary ordering and role rendezvous.

The compiler allocates barriers, emits initialization, derives phase recurrence and arrival counts, and inserts only ordering required by the effect model. Proxy/ordinary visibility is target-operation-specific, not one universal fence after every copy.

Prove matching iteration streams for the supported structured loop subset. Unknown non-affine matching is rejected for a verified pipeline, or requires an explicit separately qualified protocol. Never invent a proof by unrolling a few iterations.

Pipeline lowering must include CTA entry and exit. Loader and compute roles cannot leave while another role can still issue or read a stage. Named barriers have explicit participant sets.

11. Registers, shared memory and occupancy

The compiler owns logical register-fragment allocation and liveness. NVIDIA's backend owns final physical register allocation and instruction scheduling.

Estimate register pressure from simultaneously live fragments, scalar state and compiler-known temporaries. Label this as an estimate. Parse actual registers and local-memory use from compilation. Spills are performance evidence and may violate an explicit no-spill contract; they are not automatically a semantic error.

Allocate shared buffers using actual byte sizes, alignment and lifetime interference. Alias only when ordering proves disjoint lifetimes. Include declared dynamic storage and emitted static/reserved storage in the launch/resource assessment. Do not assume source declarations predict every toolchain overhead.

Query occupancy from the compiled function/device for persistent launch planning. Do not use estimated occupancy to guarantee a global rendezvous. Prefer designs without cross-CTA residency dependencies. Consumer work queues must not wait for blocks that cannot be scheduled.

Register-budget redistribution lowers only for admitted instructions and valid role participation. It is not a promise that arbitrary budgets are feasible or deadlock-free.

12. Arithmetic and numerical policy

Arithmetic transformations require an explicit policy. Default preserves the selected operation order and precision where the IR defines them. Fast-math, reassociation, approximate reciprocal/exp and lower-precision conversion are visible attributes, not emitter defaults.

Do not constant-fold by host arithmetic when host and device semantics differ. Format encoding/decoding uses tested routines with explicit rounding and finite-value behavior. Conversion fallbacks identify their exact algorithm.

The ptxas dequantization issue discussed in public consumer contributions motivates a real version-qualified regression case. A fallback computes the operation correctly; it must not hide a rejected instruction or clamp an unexplained numerical failure.

Online softmax is a visible expansion with typed row maximum, sum and output recurrence. Sparse cache storage precision and effective MMA precision remain distinct. Split partials carry normalization representation and LSE units so the merge cannot reinterpret them.

13. Source emission

Emit deterministic CUDA C++ with inline PTX for instruction forms that require exact control. Use ordinary CUDA constructs where their lowering is sufficient and inspectable.

Generated source includes stable operation labels/source-map metadata, not large repeated explanatory comments. Centralize host bindings and shared device helpers. Specialize only dimensions that affect code quality or legality. Dynamic dimensions remain runtime parameters when appropriate.

Do not export one source clone for every shape and architecture. Share equivalent bodies and keep measured variants distinct only where schedules differ. Hashes detect stale artifacts; they do not substitute for runtime correctness.

Compile through the repository's existing JIT/build mechanism. Emit target flags explicitly and inspect the actual PTX/cubin target rather than assuming command-line intent survived.

14. Prepared runtime and descriptor lifetime

Preparation creates fixed launchers, persistent workspace and descriptor storage. Runtime binding validates metadata against the admitted envelope without triggering compilation or implicit workspace growth.

TMA descriptors depend on address, shape, strides, dtype, swizzle/format and allocation validity. A cache keyed only by workspace address/context/slot is insufficient. Choose owned immutable descriptor storage for fixed bindings, or explicit refresh for changed dependencies. Record graph-update/recapture rules.

Expose mutation and aliasing contracts. For sparse attention, bounds of runtime index contents remain caller preconditions or explicit validation work. Metadata checks alone do not prove every index is valid. Graph replay must preserve the contract when live lengths change within capacity.

Multi-kernel programs declare dependencies, partial representation, resets and workspace lifetime. Split merge and preprocessing are part of the operator, not omitted benchmark work.

15. Diagnostics and coverage

Diagnostic
  code
  severity: error | unknown | warning | performance_hint
  operation_id
  source_span
  resource
  role
  symbolic_condition
  evidence
  repair_options

Diagnostics name the actual violated contract. Repair options do not mutate semantics or acceptance thresholds automatically.

Examples: unsupported MMA shape for target; scale grouping mismatch; shared stage reused before final consumer read; unknown warp convergence; descriptor dependency changed; launch shared-storage requirement exceeds the device limit.

Performance diagnostics are advisory and profile-keyed. Uncalibrated targets return missing coverage, not B200 estimates. Compiler errors from NVIDIA are attached to generated-source locations and mapped back to schedule nodes where possible. Preserve the original error when mapping fails.

16. Artifacts and durable reuse

Artifact identity includes canonical schedule, contract, target/PTX, lowering version, toolchain, flags and relevant dependencies. Performance records additionally include device/SKU and measurement policy.

Store sources, compiler output, objects and receipts on durable home/repository or shared storage. Do not use temporary-system storage as the build cache. Multi-executor use requires shared durable storage. Prove cache survival after consumer replacement before calling it durable reuse.

LoweringReceipt
  schedule_hash
  contract_hash
  target
  ptx_version
  compiler_version
  operation_to_source_map
  selected_instruction_forms
  shared_regions
  barrier_protocols
  actual_register_and_local_memory_usage
  object_hash
  verification_coverage

A binary cache hit is not runtime qualification or performance evidence. Do not reuse a performance conclusion across SKU/compiler profiles merely because a cubin hash matches.

17. Compiler evolution seam

The agent edits schedules in the inner loop. Changes to instruction definitions, verifier rules, lowering or arithmetic policy are outer-loop changes. Keep the external oracle and workload contract outside the agent-editable compiler path.

A proposed compiler change needs a concrete motivating failure and a focused behavior check. Confirm that it fixes the failure and retains the valid schedule that previously worked. Add broader corpus coverage when the change actually affects those operations, not as automatic ceremony.

Static rules must not be source-text tests. Test acceptance/rejection or numerical lowering behavior. Compiler-version workarounds require a version condition and a case that demonstrates the numerical difference.

18. Implementation sequence and practical proof

Step	Implementation	Smallest useful proof
1	Builder, typed graph, serialization and target admission	Construct a real warp-GEMM schedule and reject a genuinely unsupported operation
2	Fragment/shared representation and CUDA emission	Compile, inspect selected MMA/load, run one numerical GEMM smoke
3	Pipeline generations and async lowering	Run a multi-generation staged workload and detect a reuse-before-release schedule
4	Mixed-cache decode and softmax/split program	Compare the deployed subset with an independent decoded-cache reference
5	Prepared bindings and graph lifecycle	Replay with changed live metadata and exercise descriptor refresh/lifetime
6	Profile-based ranking/export	Off/on complete-operator timing on the actual device, with numerical gate

These are implementation steps, not completed deliverables. No compiler was built or run for this design document.

19. Alternatives rejected

• Direct arbitrary CUDA as the only representation: too little structured information for localized ownership diagnostics.
• A full layout algebra in the agent interface: unnecessary authoring complexity for the initial consumer instruction catalog.
• Direct PTX-only backend first: adds runtime/ABI and code-generation work without removing the need for NVIDIA register allocation.
• Automatic datacenter-to-consumer lowering: hides accumulator/resource redesign and unsupported execution models.
• Automatic precision substitution: changes the workload rather than implementing it.
• A general compiler framework before one kernel works: delays the useful end-to-end seam.

20. Open implementation decisions

Choose the exact initial fragment/load catalog from the first measured schedule. Define the canonical serialization version and symbolic-expression subset. Specify barrier lowering against the pinned PTX memory model. Decide fixed versus refreshable descriptor ownership from the deployed graph lifecycle. Measure which specialization dimensions deserve separate binaries.

Those decisions must be grounded in the first real consumer workload. They are not placeholders for a promised general compiler. The complete first result is a typed schedule that compiles, runs correctly, preserves its runtime contract, and reports enough evidence to guide the next schedule change.

04-cake-consumer-blackwell-costing-and-agent-loop.md

Consumer CAKE Costing, Scoring and Agentic Optimisation

Targets: SM120 and SM121, compiled as admitted sm_120a, sm_121a or sm_120f objects. Status: proposed design. No cost model or optimisation runner was implemented or benchmarked for this document.

Companions: evidence review, IR and compiler designs. Methodology source: CAKE paper v1, sections 3-6 and Appendix B.5. The paper's B200/H100 calibration does not supply consumer rates.

1. Decision

Use three separate authorities:

1. Contract and verifier determine eligibility.
2. A calibrated, uncertainty-aware cost model orders experiments.
3. Paired on-device measurements determine winners.

Never combine correctness and speed into a weighted reward that permits an incorrect kernel to win. Never promote an unmeasured candidate because its predicted latency is low. Never substitute GeForce calibration for RTX PRO or GB10 simply because an instruction is legal on each.

Start with analytic features and empirical nearest-neighbour ranking. Add a fitted residual model only when actual measurements show a recurring prediction error. A complex learned model is not a prerequisite for the first useful kernel optimisation.

2. Immutable experiment contract

ExperimentContract
  operation_and_numerical_policy
  supported_shapes_and_layouts
  state_cache_workspace_and_graph_semantics
  oracle_and_acceptance
  primary_metric
  representative_workload_rows
  row_weights_or_trace_frequencies
  allowed_regressions
  target_device_and_toolchain
  measurement_protocol
  reference_access_policy

The primary metric defaults to completed operator GPU time for the deployed execution mode. It includes preprocessing, quantization, split merge and resets when the operator requires them. Wall time is a separate metric when host overhead matters. Preparation and compilation are reported separately from steady-state runtime.

Use real trace frequencies when available. Otherwise use explicit equal weights over the selected representative rows and state that choice. Do not invent traffic probabilities.

The agent cannot edit the oracle, tolerance, weights, shapes, primary metric or regression allowance during an inner-loop run. A contract change creates a new experiment identity.

3. Eligibility, measurement and scoring are different states

Proposed -> Constructed -> Analyzed -> Compiled -> Validated -> Measured
                                                          -> Rejected
Measured -> Challenger -> ConfirmedWinner -> ExportQualified

Failures retain their stage and diagnostic. An analysis unknown is not a correctness pass. Compilation does not establish runtime validation. A confirmed fixed-shape winner is not yet a qualified dispatcher export.

Eligibility includes required numerical outputs, state/cache mutation and runtime lifecycle. A precise bounds failure or graph-lifetime defect is disqualifying. Resource limits are hard constraints. Register pressure below the limit is a performance feature, not a moral penalty.

4. Cost model input

Extract features from concrete lowered IR and, after compilation, replace estimates with observed compiler facts.

Feature group	Examples
Work	MMA instruction counts/forms, scalar arithmetic, reductions, conversions
Data movement	Global bytes, L2 reuse assumptions, shared traffic, gathers, stores
Representation	Swizzle, transpose/repack, narrow packing and scale loads
Schedule	Roles, stages, tile grid, persistent/static assignment, split count
Dependence	Loop-carried state, copy-to-read paths, softmax-to-PV paths
Synchronization	Waits, rendezvous, pipeline startup/drain, merge launches
Resources	Shared bytes, estimated/actual registers, spills, occupancy
Runtime	Descriptor refresh, host metadata work, allocations, launch mode
Workload	Length distribution, masking density, page pitch and selected indices

Do not estimate irregular gather locality from shape alone. If index distributions are unavailable, report a range and missing coverage. Separate logical byte volume from measured memory transactions.

5. Hardware calibration

CalibrationProfile
  device_identity_and_sku
  compute_capability
  driver_and_compiler
  clocks_and_power_policy
  memory_topology
  instruction_and_memory_rates
  launch_and_sync_costs
  fitted_overlap_parameters
  coverage_and_residual_samples

Profiles distinguish GeForce SM120, workstation SM120 and GB10 SM121. The family target is not the calibration key.

Measure only primitives relevant to the current workload: its MMA form, matrix loads, dequantization, memory transfers, softmax/reductions and launch mode. Measure both isolated throughput and in-situ overlap where the schedule depends on it. Device limits come from the actual device. Rates come from measurements, not marketing throughput.

Record cache-resident versus cold-memory protocols, contention, clocks and sample dispersion. GB10 CPU activity can affect shared memory/power behavior. Missing CUPTI does not prevent event-based timing, but the record must identify the timing source and its limitations.

Calibration is invalidated or marked stale when the relevant hardware/toolchain/protocol changes. Uncovered operations return unknown contributions. Do not silently import B200 or another SKU's rates.

6. Analytic estimate

For resource class r, estimate service demand as:

D_r = work_r / calibrated_rate_r

All work counts, byte volumes and computed estimates are produced by code from the IR. Do not hand-enter derived offsets, instruction counts or latency arithmetic.

A simple lower bound is the largest service demand and the dependency critical path. It is not a prediction of total time. A realistic pipeline also pays startup, drain, exposed dependency waits, launch and finite-work scheduling costs.

Model the lowered schedule as a dependency/resource graph:

• Nodes carry service demand and participant/resource use.
• Edges represent true dependencies, storage ownership and synchronization.
• Resource capacities limit simultaneous issue/service.
• Pipeline stages constrain which generations can overlap.
• CTA occupancy limits resident work and latency hiding.

Initially use a coarse deterministic schedule estimator, not a cycle-accurate simulator. Estimate repeated steady-state work separately from startup/drain. Fit overlap parameters to measured schedules. Do not sum overlapping demands or assume every demand overlaps perfectly.

For multiple kernels, estimate each kernel and the measured inter-launch overhead under the chosen execution mode. For irregular sparse work, use observed workload summaries or measured neighbours rather than a fictitious uniform-access model.

7. Uncertainty and model output

CostEstimate
  predicted_latency
  empirical_prediction_interval
  coverage
  nearest_measured_schedules
  dominant_demands
  exposed_dependencies
  resource_pressure
  sensitivity_hints
  model_version

Prediction intervals come from residual measurements or explicitly bounded assumptions. They are not confidence percentages invented by the agent. If evidence is insufficient, omit the numeric interval and label the estimate uncalibrated.

Separate measurement noise from model uncertainty. Repeated measurements reduce timing noise but do not repair missing calibration or a wrong overlap model.

A useful explanation identifies a schedule decision: probability repacking is exposed between QK and PV; shared-stage depth cannot hide gather latency; accumulator liveness reduces resident CTAs. Avoid generic hints such as "improve occupancy" without naming the limiting allocation and tradeoff.

Do not automatically prefer higher occupancy. A lower-occupancy schedule may have better reuse or fewer instructions. The estimate and measurement determine the tradeoff.

8. Prediction is an acquisition score, not the final score

When candidates have calibration coverage, order experiments by expected positive improvement relative to the measured incumbent, with uncertainty represented by the empirical prediction distribution.

A useful acquisition quantity is:

expected_gain = E[max(incumbent_latency - candidate_latency, 0)]
priority = expected_gain / expected_evaluation_time

Evaluation time includes compilation and the required smoke/measurement. This denominator prioritizes experiments; it is not a runtime performance objective or a reason to discard a required workload.

Avoid committing to an arbitrary weighted sum of novelty, occupancy and latency. Use this ordering:

1. Required repairs and missing qualification work.
2. Covered candidates with promising expected improvement.
3. A structurally distinct candidate whose model uncertainty represents a real unexplored schedule.
4. Calibration experiments that discriminate between competing explanations.

Reserve actual experiment opportunities for unknown schedule classes. Do not let a poorly calibrated model filter every novel candidate before measurement. No fixed exploration percentage is required initially.

If acquisition uncertainty cannot be estimated, rank by measured neighbours, analytic bottleneck evidence and structural diversity. State that the ranking is heuristic.

9. Measured score

For each workload row i, define speedup as baseline latency divided by candidate latency. Latencies use the same measurement protocol. Report the full row table and dispersion.

For a fixed-shape inner loop, choose the candidate with the lowest confirmed primary latency among eligible candidates. For a production portfolio with real invocation weights, choose the lowest weighted total latency:

objective = sum(weight_i * candidate_latency_i)

Also report weighted geometric-mean speedup when useful for comparison:

geomean_speedup = exp(sum(normalized_weight_i * log(speedup_i)))

The geometric mean is descriptive unless the contract explicitly selects it as the objective. It cannot hide a forbidden regression on a required row.

The winner rule is lexicographic:

1. Complete required correctness/lifecycle admission.
2. Satisfy workspace/resource and required-row regression constraints.
3. Improve the primary measured objective beyond established timing uncertainty.
4. For indistinguishable candidates, prefer the simpler schedule or smaller specialization portfolio.

Compilation time, workspace, binary size and preparation cost are separate reported dimensions. Use hard limits where the deployment actually requires them, rather than ungrounded penalty coefficients.

10. Paired measurement and promotion

Measure incumbent and challenger in an interleaved, identical protocol. Warm both, use the same input/reset policy and execution mode, and record sample order. Alternate order to avoid assigning clock drift to one arm.

A first measurement screens candidates. A promising challenger receives a confirmation run. If the apparent gain lies within observed noise, mark it tied or collect additional focused evidence. Do not announce a win from a single favourable sample.

Keep two notions separate:

• Best observed sample: useful diagnostic, not default promotion evidence.
• Robust paired latency summary: promotion evidence for the workload.

The run record retains raw paired samples and invalidation reasons. Faulting candidates run in isolated processes so sticky context errors do not contaminate later timing. A timeout is a failure record, not an extreme latency sample.

No performance comparison is admissible when either arm fails its numerical contract. Include full-operator work symmetrically. Do not exclude candidate preprocessing while charging baseline preprocessing.

11. Candidate record and evidence ledger

CandidateRecord
  candidate_id
  parent_id
  contract_hash
  canonical_schedule_hash
  hypothesis
  changed_decisions
  target_and_toolchain
  analysis_report
  lowering_receipt
  validation_result
  measurement_records
  promotion_status
  failure_diagnosis

Store canonical schedules, compiler output, artifacts and measurements durably. Deduplicate by semantic identity and compilation identity, not filenames. Performance records are additionally device/protocol-specific.

Retain failed hypotheses so the agent does not repeat them without new evidence. A failure on one target is not automatically a failure on another. An unsupported instruction and a numerical lowering defect are different diagnoses.

12. Agent interfaces

Keep the action surface small and typed.

inspect(candidate_or_baseline, requested_evidence) -> compact evidence
propose(parent, hypothesis, schedule_changes) -> candidate
analyze(candidate) -> legality and coverage
compile(candidate) -> artifact and lowering receipt
validate(candidate, required_cases) -> numerical/lifecycle result
measure(candidate, comparison) -> paired measurements
promote(candidate) -> gate result

The orchestrator owns promotion and workload authority. The agent proposes and interprets evidence. It cannot set its own success flag.

External logs, PR text and profiler output are untrusted data. They cannot authorize shell commands, changes to the oracle or altered acceptance. Compiler changes use a separate approval path.

Progressive disclosure: give the agent the contract summary, incumbent schedule, recent failures and bottleneck evidence. Fetch detailed source/profile rows on demand. Do not flood every proposal with the entire history.

13. Inner optimisation loop

load immutable contract and measured incumbent
select an evidence-backed hypothesis
construct one candidate with explicit changed decisions
analyze; repair the candidate if diagnostics explain a local defect
compile; inspect actual lowering/resources
validate required deployed subset
measure against incumbent
confirm a promising improvement
promote only through orchestrator gates
record outcome and next hypothesis

Change one causal decision at a time when diagnosing an effect. A new schedule family may change several coupled decisions, but record that it is a structural experiment rather than attributing its gain to one edit.

Examples of legitimate hypotheses:

• Unit-scale block-scaled FP8 removes a measured GeForce instruction-rate limit.
• A different stage depth hides gather latency but increases shared storage.
• A smaller accumulator tile removes spills at the cost of more repeated loads.
• A split policy reduces load imbalance on GB10 but adds merge traffic.
• Probability packing moved into an existing producer stage shortens an exposed dependency.

Every hypothesis names predicted evidence that could falsify it. Do not repeat "optimise memory access" as an actionable proposal.

14. Outer compiler/model evolution loop

Classify recurring failures before changing the environment:

Evidence	Owner
Bad schedule violates an existing rule	Candidate repair
Illegal schedule escaped checks	Verifier rule
Valid IR emitted incorrect instruction/representation	Lowering fix
Missing expressible consumer operation	IR primitive plus effects/legality/lowering
Correct kernel consistently misranked	Calibration/model change
Stale descriptors or graph lifetime	Prepared runtime/exporter fix
Baseline/oracle/measurement discrepancy	Experiment authority investigation

An outer-loop proposal contains the concrete failure, proposed change, expected behavioral evidence and affected operations. Keep one environment change per causal experiment. Preserve the workload contract unless its owner explicitly changes it.

The smallest proof is failing-before/passing-after behavior plus a relevant previously valid case. Expand corpus checks according to the actual affected lowering or analysis. Do not impose a broad academic reproduction matrix on every local fix.

A model update is judged on retained measured schedules not used to fit that update. Confirm that it improves ranking or error for the relevant domain. Do not tune a model merely to assign the current favourite candidate a better score.

15. Parallelism and device scheduling

Parallelize independent candidate construction, static analysis and compilation when storage/toolchain resources permit. Serialize timing on a device unless concurrent timing is itself the deployed workload. GPU contention invalidates a quiet-device comparison.

Workers receive immutable contracts and separate candidate workspaces. One orchestrator owns incumbent promotion and dispatcher integration. Workers do not overwrite calibration or export state independently.

A GPU queue records device ownership, active workload and run protocol. Do not evict serving workloads. If the device is busy, do useful construction/analysis work and wait for an admissible measurement window.

Multiple devices provide separate measurement profiles. Family-object compatibility does not justify merging their timing samples.

16. Portfolio and dispatcher optimisation

Fixed-shape winners seed a small portfolio. The portfolio objective includes dispatch, preparation behavior and complete-operator latency over the deployed rows.

For each candidate route, record guard, schedule, target/toolchain admission and supported numerical/lifecycle contract. Test actual guard boundaries when introducing or changing a route. A missing recipe rejects clearly or reaches a qualified semantically equivalent fallback.

Prefer one runtime-variable schedule when it performs indistinguishably from many clones. Add a specialization only for a measured benefit or required legality. Do not reward generated line count or number of kernels.

A per-shape win does not authorize migration of a whole b12x family. Export only the qualified subset and retain required nonmatching paths until they have replacements.

17. Stopping and delivery

Stop an optimisation run when the requested deployed contract has a qualified result and further tested hypotheses do not produce a distinguishable gain, or when the user-defined experiment limit is reached. A plateau is recorded from actual unsuccessful experiments, not agent intuition.

If the measured incumbent remains best, export it and report that result. Optimisation is not required to invent a speedup. If hardware, credentials or a reference is unavailable, finish construction/design work and identify the precise missing prerequisite.

Delivery contains accepted schedule/portfolio, source/object identity, admission envelope, numerical/lifecycle smoke evidence, paired full-operator timing, resource/workspace costs, known exclusions and calibration coverage. Proposed and unmeasured candidates remain labelled as such.

18. First practical application

Begin with the matching V4.1 mixed-cache decode subset on the available consumer device. Establish the current b12x and CAKE baselines under one protocol. Capture actual compiler resources and full-operator timings. Use the public schedule to identify the first exposed dependency or resource limit.

For GeForce FP8, use a separate actual GeForce witness to compare legacy and unit-scale block-scaled MMA. GB10 execution of a byte-transpose primitive does not establish the GeForce performance hypothesis.

Initially keep scoring deterministic: eligibility, paired primary latency, required-row constraints and simplicity tie-break. Add predictive ranking only from relevant measured samples. This reproduces CAKE's evidence-driven loop without claiming access to its private cost model.

Appendix A. Material implementation changes from citing papers

These are design decisions informed by other systems, not evidence about CAKE's unpublished implementation. All four papers cite CAKE. None supplies consumer-device calibration or CAKE's private cost-model parameters.

Source	Decision	Material change
Cleave	Adopt traffic/liveness accounting now. Defer algebraic search.	Add a deterministic schedule replay that counts repeated transfers and tracks peak live storage by memory class.
AI as a Compiler	Adopt evidence-directed independent candidate branches. Keep the typed IR.	Add structured hypotheses, branch isolation and orchestrator-owned admission of measured improvements.
RESOLVE	Add an optional deeper qualification path. Do not replace numerical gates.	Add determinism evidence and, where available, separately recorded reduction/proof evidence.
RR-Evict	No kernel/compiler change.	Consider serving policy only if measured agent-serving tail latency limits optimisation throughput.

A.1. Cleave: implement traffic and liveness before a fitted predictor

Cleave Appendix A describes single-threadblock execution simulation with a shared-memory allocator. It accumulates global-memory traffic across iterations and tracks peak shared storage through buffer liveness and reuse. Its scheduler shortlists candidates before compilation and profiling. It assumes all buffers reside in shared memory for cost estimation.

The compiler's existing representation and resource-planning passes should supply one common event stream to the storage allocator and estimator:

ScheduleEvent
  operation_id
  symbolic_iteration_multiplicity
  acquire_or_release_storage
  memory_class_and_physical_size
  transfer_source_and_destination
  transferred_bytes
  stage_generation
  dependencies

Replay structured loop summaries rather than materializing every dynamic iteration. Count repeated global loads even when the source tensor is unchanged. Preserve conditional and sparse-access uncertainty rather than treating every logical load as an equal number of physical transactions.

Unlike Cleave's simplifying model, distinguish shared storage, register fragments and global workspace. Track accumulator liveness separately from pipeline buffers. Register-fragment counts remain estimates until compilation supplies actual registers and spills. Do not convert global byte counts directly into latency without the device/profile-specific rate and overlap model.

This changes the implementation order: build a testable traffic/liveness ledger first, then calibrate dependency/overlap estimates. A learned residual predictor comes later. The ledger also makes a bad estimate diagnosable by operation and loop, rather than presenting an unexplained scalar score.

Acceptance evidence: use two real schedules with different stage depths or tile reuse. Check the ledger against concrete allocation sizes and executed transfer geometry. Compile both and compare actual resources and measured ordering. A ranking mismatch is model evidence, not a reason to modify the measurements.

Cleave also separates symbolic algebraic transformation from scheduling in sections 3-5. Do not add its superoptimizer to the first implementation. If schedule-only search later stalls because the graph materializes an avoidable intermediate, introduce a separate graph-variant identity and approved rewrite preconditions. Each graph variant can then have several schedules. Algebraic equality alone does not authorize changes to floating-point behavior, quantization, masking or sparse-cache semantics.

A.2. AI as a Compiler: branch hypotheses, not arbitrary edits

AI as a Compiler section 3.3 describes a diagnose-plan-branch-evaluate loop. Independent subagents implement candidate patches. The execution environment admits only candidates that pass its gates and improve the measured objective.

Make the optimisation plan a machine-readable input to candidate workers:

Hypothesis
  parent_schedule_hash
  bottleneck_evidence_ids
  proposed_decisions
  expected_resource_or_profile_change
  falsification_condition
  required_validation_cases

Generate independent branches only when they test distinct useful choices, such as stage depth versus tile geometry. Each worker receives the same immutable contract and incumbent identity. Construction and compilation may run concurrently. Validation and timing use the device queue, without competing GPU experiments.

After compilation, compare the expected lowering/resource change with the receipt. If the proposed instruction form did not appear, report that fact before attributing timing to it. Fetch profiler counters only when they discriminate between hypotheses. Missing profiler access must not produce invented bottleneck evidence.

Promotion remains an orchestrator operation using paired confirmation. A worker's claim that it improved performance cannot mutate the incumbent. Record whether a result falsified the hypothesis, improved despite the proposed explanation, or remained inconclusive. This distinction prevents successful timing from laundering an incorrect causal explanation into future tactics.

Do not adopt direct arbitrary PTX generation as the default backend. The paper studies a different compilation boundary. Preserve the typed schedule and deterministic emitter proposed here. A missing instruction or transformation triggers the outer compiler loop, not an unverified agent escape hatch.

Its evaluation uses L40S, H100 and B200. Its section 4.5 reports verifier limitations around atomics, input-dependent control flow and bit-level manipulation. Neither its speedups nor its verification coverage qualifies an SM120/SM121 implementation.

A.3. RESOLVE: represent validation claims separately

RESOLVE sections 3-5 combine binary-level determinism testing, bitwise-tested concurrency reduction and formal algebraic reasoning. This materially changes the evidence schema, not the runtime score.

Replace a single undifferentiated validation flag with explicit records:

ValidationEvidence
  numerical_contract_result
  mutation_and_lifecycle_result
  sanitizer_result_and_coverage
  repeated_execution_determinism
  timing_perturbation_result_and_coverage
  reduction_correspondence_test_result
  formal_specification_identity
  formal_proof_result_and_assumptions

The orchestrator derives eligibility from the contract's required records. Missing optional proof evidence is not failure. Missing required numerical evidence is not success. Hash or compare contract-observable output and mutated state, including workspace state when its contents affect later invocations. Do not require untouched scratch bytes to be identical unless the contract makes them observable.

Start with repeated-execution checks on finalists whose contract requires determinism. Add binary timing perturbation only after confirming instrumentation support on the actual consumer device and relevant instructions. Instrumented runs are validation runs, never performance samples. Preserve input/reset identity across comparisons.

If reductions and proofs are introduced, keep the links explicit: original-to-reduced correspondence is tested, reduced-to-proof-program correspondence is tested, and the proof applies to a named specification under named assumptions. The oracle/specification must not be derived from the candidate. Review specification strength before treating a proof as useful evidence.

Do not label this pipeline a proof of the original binary's race freedom or numerical error bound. RESOLVE explicitly reports perturbation gaps, including intra-warp behavior and asynchronous-load timing. Its algebraic proof abstracts floating-point arithmetic and does not establish an acceptable error budget. Legitimate nondeterministic reductions require a different contract, not automatic rejection by a universal determinism rule.

Keep reduction/proof work outside the cheap candidate-screening loop initially. Apply it to a finalist or a diagnosed ambiguity that warrants deeper qualification. A concrete acceptance scenario is a synchronization defect exposed by changed execution timing, repaired without changing numerical tolerance, followed by the same validation campaign and an uninstrumented paired performance check.

A.4. RR-Evict: leave the compiler unchanged

RR-Evict section 5 cites CAKE as related operator-execution work. It does not extend CAKE's cost model, IR or search mechanism.

Do not add prefix-cache management to the compiler or GPU measurement queue. If local LLM serving becomes a measured bottleneck for concurrent optimisation agents, evaluate serving-policy changes separately using actual agent trajectories and tail latency. Kernel scoring stays unchanged. Do not infer that the paper's H100 results apply to the available serving hardware.

A.5. Resulting implementation priority

1. Add concrete traffic/liveness extraction to compiler analysis and receipts.
2. Add structured, falsifiable hypotheses and independent candidate branches to the orchestrator.
3. Split numerical, lifecycle, determinism and proof claims in the evidence ledger.
4. Retain paired full-operator timing as winner authority.

Do not make algebraic superoptimization, arbitrary PTX synthesis, formal proof infrastructure or serving-cache changes prerequisites for the first correct measured consumer kernel. Those changes need a concrete workload limitation and separate evidence.