Kernel Engineering Guide

This is the guide a new contributor reads to learn how Iron kernels are built, analysed, tested, measured, and improved — and where to start. It assumes you know GPUs and Rust; it does not re-teach either.

It is a synthesis, not a replacement, for the docs it points at. Where this guide and a more specific doc disagree, the specific doc wins — file an issue or a docs PR to reconcile them.

Table of contents

1. Why this guide / how to use it
2. Architecture analysis
3. Design principles for a kernel
4. The improvement loop
5. Correctness and testing
6. Measurement protocol
7. Gates and PR hygiene
8. Gotchas catalogue
9. Where to start (for new contributors)
10. Glossary
11. References

---

1. Why this guide / how to use it

Iron compiles one Rust-embedded kernel definition to four GPU backends — Metal (MSL), CUDA, HIP/ROCm, and Vulkan (SPIR-V) — through a shared IR and a backend-independent optimization pipeline. Writing a correct kernel is a few hours of work once you know the DSL. Writing a kernel that is correct, tested, portable across backends, and demonstrably faster than what it replaces is a different, longer exercise — and this repository has strong, sometimes counter-intuitive opinions about how to do it safely. This guide collects those opinions in one place.

Read it straight through once. After that, treat it as a reference: jump to §3 when designing a new kernel, §4 when you’re trying to make an existing one faster, §5 and §6 before you trust any result you produce, and §8 when something silently doesn’t work.

Two documents this guide leans on constantly and does not duplicate in full: docs/developing.md, whose “Kernel-authoring hazards” section is mandatory reading before you write or modify a kernel (one hazard is “a wrong dispatch geometry can freeze the machine”), and docs/STYLE_GUIDE.md, the authority on the mechanical shape of one kernel file. This guide is the map; those two are the law.

A note on scope: this guide documents the DSL as it stands, with the kernel-performance work landing from PR #299 into dev — the CUDA advanced tier in §3.3 and the levers, tests, and gotchas built on it throughout this guide are that work, verified against its own branch and recorded here as first-class content.

---

2. Architecture analysis

2.1 The pipeline

A kernel’s trip from Rust source to a running GPU thread has five stages:

flowchart LR
    A["#[kernel] fn\n(Rust DSL)"] -->|proc-macro parses body| B["Iron IR\n(Kernel / Op graph)"]
    B --> C["optimization pipeline\n14 passes, backend-neutral"]
    C --> D{"backend lowering"}
    D --> E["MSL\n(Metal)"]
    D --> F["CUDA\n(C++ / PTX)"]
    D --> G["HIP\n(ROCm)"]
    D --> H["SPIR-V\n(Vulkan)"]
    E --> I["wh-iron-runtime\ncompile + dispatch"]
    F --> I
    G --> I
    H --> I
    I --> J["GPU"]

1. Parse. #[kernel] is a proc-macro (wh-iron-macros) that reads the annotated Rust function’s body as a small imperative language — loads, stores, arithmetic, if/for, calls to a fixed set of DSL primitives — and lowers it to Iron’s IR at compile time. This is a real parser with real diagnostics, not a token-soup macro: unsupported syntax is a compile error, not a silent no-op (with one documented exception — see §8, “inner macro_rules! traps”).
2. IR. The IR (wh-iron-core::ir) is a typed, dtype-generic graph of Ops inside a Kernel: one entry block (kernel.body) plus a map of auxiliary blocks (kernel.blocks) for if/for bodies. It is backend- and precision-neutral — the same IR node lowers differently depending on target and dtype.
3. Optimize. wh-iron-codegen runs a fixed, ordered pipeline of passes over the IR (below) before any backend sees it.
4. Lower. A CodegenBackend implementation per target (msl::MslGenerator, cuda::CudaGenerator, and the HIP / SPIR-V equivalents) walks the optimized IR and emits target source: MSL text, CUDA C++/inline-PTX text, or a SPIR-V module.
5. Compile and dispatch. wh-iron-runtime turns emitted source into an executable pipeline per backend (a Metal MTLComputePipelineState from a newLibraryWithSource:-compiled .metallib, an NVRTC-compiled CUDA module, …), caches it, and dispatches it against real device buffers.

2.2 Crate map

Seven of the eight workspace crates form a strict dependency chain from shared data model up to the CLI; the eighth is a standalone tool.

Crate	Role	Depends on
wh-iron-core	The shared IR — Kernel/Op graph types, DType, Shape, error types. Pure data, no logic. Everything else depends on it.	—
wh-iron-macros	The #[kernel] / #[bench] / #[test_kernel] proc-macros: the DSL’s grammar, its compile-time diagnostics, and the body parser that turns a Rust fn into IR.	wh-iron-core
wh-iron-codegen	The optimizing compiler: the 14-pass pipeline plus the four backend lowerings (msl, cuda, hip, spirv) and the Target / TargetProfile types that describe what each backend/architecture supports.	wh-iron-core
wh-iron-runtime	The execution layer: per-backend Device implementations (metal_device.rs, device/cuda/, device/hip/, device/vulkan/), buffer pools, PSO/pipeline caching, and dispatch.	wh-iron-core, wh-iron-codegen
wh-iron	The facade — re-exports core, macros, codegen, runtime behind one prelude, so downstream code depends on a single crate. No logic of its own.	all of the above
wh-iron-std	The kernel standard library: every production #[kernel] definition, organized under kernels/<family>/ (gemm, sdpa, moe, norm, rope, convolution, ssm, quant, audio, vision, sampling, kv_cache, ops, …), plus each kernel’s #[test_kernel] correctness cases and #[bench] benchmarks, and hand-written GPU-correctness / stress tests under tests/. Where new kernels land.	wh-iron
wh-iron-cli	The iron binary — bench, build, test, inspect, device, snap, diff — the developer entry point. Nothing depends on it.	wh-iron, wh-iron-std
wh-iron-msl	A standalone MSL frontend: it parses externally-authored Metal Shading Language source into Iron’s IR, so kernels not written in the DSL can still be routed through Iron’s backends. Independent of the pipeline above — it produces IR, it doesn’t consume it — and out of scope for day-to-day kernel authoring.	wh-iron-core

wh-iron-cli is a thin protocol parser: the actual codegen, MSL/CUDA emission, and every bench/test dispatch runs inside __iron_runner, a separate [[bin]] of wh-iron-std that iron spawns as a sibling process. This matters in practice — see the “stale runner” entry in §8.

2.3 The optimization pipeline

The 14 passes run in this fixed order (PassRegistry::order() in crates/wh-iron-codegen/src/passes/mod.rs):

kernel_inline → type_check → const_fold → algebraic_simplify → copy_prop → cse → licm → if_conversion → value_sink → fusion → fma_fusion → unroll → schedule → vectorize → dead_store_elim

A few are worth knowing by name:

• fusion builds FusedElementwise chains out of independent elementwise ops so they emit as one expression instead of round-tripping through registers or memory between them.
• fma_fusion rewrites Add(Mul(a, b), c) into a single fused multiply-add. It runs after fusion and after type_check (it needs inferred types), and its own dead-code sweep cleans up the now-dead Mul.
• vectorize turns groups of adjacent scalar loads/stores into wide vector ops (half2/half4, bfloat2/bfloat4, …) when the addresses are provably consecutive. This is the pass most sensitive to how you write index arithmetic — see §8.
• cse (common-subexpression elimination) unifies identical value computations; it does not unify repeated address subtree computations — binding a repeated base + offset to a let yourself is what lets vectorize see it as one address (see §8).
• dead_store_elim is last for a reason: every pass before it can leave behind now-unused values, and running it once at the end is cheaper and simpler than threading dead-code cleanup through each pass individually.

iron inspect <kernel> --pass <name> (or --pass all) prints the IR after any stage; --stats prints the per-pass op-count delta. Both are the fastest way to find out which pass changed a kernel’s shape when a change regresses or a fusion doesn’t fire.

2.4 Backends

Backend	Crate module	Status on dev	Notes
Metal (MSL)	wh-iron-codegen::msl	Default, most mature	Every kernel’s baseline; the macOS CI runner exercises it with a real GPU.
CUDA	wh-iron-codegen::cuda, wh-iron-runtime::device::cuda	Production	NVRTC (JIT) + the CUDA Driver API for launch. Cooperative-tile ops lower to inline PTX ldmatrix + mma.sync (MmaStrategy::MmaSync16x8x16, the profile default on sm_80+) with a wmma-fragment fallback (Wmma16x16x16) and a software fallback for ineligible tiles/dtypes/older architectures. `IRON_CUDA_MMA=software
HIP (ROCm)	wh-iron-codegen::hip, wh-iron-runtime::device::hip	Narrower validated set	Same IR and pipeline; a smaller kernel set is currently exercised end to end on this backend.
Vulkan (SPIR-V)	wh-iron-codegen::spirv, wh-iron-runtime::device::vulkan	Narrower validated set	Includes a real cooperative-matrix (coopmat) path for tiled matmuls, not just a scalar fallback.

All four share the same IR and the same 14-pass pipeline; only the final lowering step and the runtime’s compile/dispatch machinery differ per backend. A kernel written once against the DSL’s dtype-generic <T> parameter is emitted for f32/f16/bf16 on every backend that supports those dtypes, without four hand-written copies.

2.5 Dispatch modes and selectors

Every kernel declares a KernelMode, which controls which Metal-style built-ins the emitter binds (this vocabulary — thread vs. threadgroup position — is shared across all four backends, not Metal-specific):

Mode	Binds	Used for
Elementwise (default)	flat tid	One thread per output element.
Reduction	tgid/lsize/tid (3-axis, .x/.y aliases)	Row-reduction kernels — softmax, norm, GEMV — where program_id::<N>() means the threadgroup index and threads within a group cooperate via simd_*/threadgroup_*.
Grid3D	3-axis thread position	Kernels that need a genuine 3-D thread grid (e.g. RoPE).
Tile2D	2-D threadgroup-local + threadgroup position	Tiled 2-D kernels (GEMV, small matmuls).
SimdGroup2D	threadgroup position + simd lane/group, 3-axis	Simdgroup-matrix-tiled kernels (the steel GEMM family) and any kernel needing a third grid axis, e.g. a batched SDPA prefill split by K.

Getting the mode wrong doesn’t just misbehave — it’s rejected: a kernel body that reads tgid_x without declaring mode = Reduction (or an equivalent) fails a registry-level gate before it ever reaches a backend, precisely because program_id::<N>()’s meaning — thread index vs. threadgroup index — flips between Elementwise and Reduction, and a kernel dispatched under the wrong assumption produces silently wrong output rather than a crash.

Compile-time variants. #[kernel(variants(AXIS = [v1, v2, ...], suffix = "..."))] expands one generic body into several concrete kernels at compile time — tile shapes, block widths, feature flags — each getting its own entry-point name from the suffix template. The split-K GEMM family, for example, declares variants(BM = [64, 32], BN = [64, 32], WM = [2, 2], WN = [2, 2], suffix = "{BM}x{BN}x16_{WM}x{WN}") and gets two concrete kernels, iron_steel_gemm_splitk_64x64x16_2x2 and _32x32x16_2x2, from one function body.

Selectors. A production kernel family — several tile shapes or strategies for the same logical operation — is chosen at dispatch time by a plain Rust function, conventionally named <family>_for(...), that inspects shape, dtype, and device capability and returns which concrete kernel to call. sdpa_prefill_mma_for(dtype, family) is a representative example: it picks a bf16-specialized kernel body on pre-Apple9 GPUs and the general one elsewhere, and opts a narrowing f32→bf16 output cast in for the MMA path because that path’s error is already inside the correctness tolerance attention allows. Selector logic like this is itself unit-tested (a “selector test” is a plain #[test] that asserts which kernel a given shape/dtype/capability routes to) — see §5.

On the consumer side, the generated Swift surface (iron build --emit swift) mirrors this: a kernel family expanded from variants(...) gets a typed Swift selector (IronKernels.<family>(kb:, vb:, d:, dtype:, ...)) that switches over the emitted rows so a caller never assembles an entry-point name by hand, and a combination that was never emitted traps loudly instead of silently miscompiling.

Dumping and inspecting a kernel. iron inspect <kernel> prints final MSL; --ir prints raw IR before any pass runs; --dtype picks a monomorphization. iron build --emit all -o <dir> emits every kernel’s source (plus, for Metal, the compiled .metallib and the generated Swift bindings) for every backend the build was compiled with — the standard smoke check for “does everything I touched still compile.”

2.6 Runtime dispatch

wh-iron-runtime::Context owns one backend device and exposes dispatch_with_grid(kernel, buffers, constexprs, grid_xyz, tg_xyz). Buffers bind by name against the kernel’s declared parameters; #[constexpr] scalars go in as small uniform buffers. grid_xyz is always counted in threadgroups, not threads — total dispatched threads is grid.{x·y·z} · tg.{x·y·z} — which is the single most common source of a degenerate or catastrophic dispatch; see §8.

Compiled pipelines are cached per backend (a Metal PSO cache, an NVRTC-module cache, …) so a kernel compiles once per process. Multi-kernel sequences — a split-K GEMM’s partial-then-accumulate pair, a multi-pass SDPA prefill — are ordinary sequential dispatches through the same Context; the DSL has no special “pipeline” primitive for this, the handoff is just an ordinary device buffer one kernel writes and the next reads.

---

3. Design principles for a kernel

3.1 Start from the shape, not the algorithm

The right kernel for an operation depends on the production shape, not the operation in the abstract. The clearest illustration in the standard library is GEMV: iron_gemv puts one threadgroup (256 threads) on one output row, which is fine when K is large — plenty of work per thread. At a small-K, large-M shape (a [10240, 320] down-projection, K=320), that same layout leaves roughly one element per thread and spins up 10,240 near-idle threadgroups, measured at well under half the device’s streaming bandwidth floor. iron_gemv_warp_rows fixes it by putting one output row on one simdgroup instead of one threadgroup, so a single threadgroup covers 8 rows at once — an 8× reduction in threadgroup count and a real 4-wide unrolled work item per lane instead of a near-idle stride. Same operation, different production shape, different optimal tile — that’s the rule, not the exception. Before designing a kernel, know the exact shapes it will run at in the consumer.

3.2 Occupancy math

Occupancy — how many threadgroups/blocks can be co-resident on one compute unit — is bounded by whichever resource runs out first: registers, shared/threadgroup memory, or threads-per-block/threadgroup limits. The limits differ by backend and, on CUDA, by exact SKU; concrete numbers help more than the abstract rule.

Metal. Apple GPU families cap at 1024 threads/threadgroup and roughly 32 KB of threadgroup memory across the M1–M5 generation; where they differ is registers. Pre-M3 GPUs (Apple7/8) use a fixed ~128-register-per-thread guide; from M3 on, the on-chip Occupancy Management Unit allocates registers dynamically, so the 128-register figure becomes a soft heuristic rather than a hard ceiling, and the codegen models register pressure as a gradual occupancy degradation rather than a cliff. Iron’s own estimate (crates/wh-iron-codegen/src/passes/register_estimate.rs) is a conservative linear-scan liveness count over the IR — useful for comparing two tile-size candidates against each other, not a substitute for the real compiler’s allocation. It feeds an autotuner pass, not the pipeline’s correctness path.

CUDA — worked example on a GB10 (48 SMs, compute capability 12.1, ~24 MiB L2, up to ~99 KB of shared memory per block via an explicit cuFuncSetAttribute opt-in above the default 48 KB). For a split-K GEMM tile with (BM, BN) = (32, 64), four warps (128 threads) per block, and a double-buffered K-loop staging tile: shared memory is the two staging buffers times two (double buffer) times the element size, and the fit against the 99 KB ceiling directly determines whether the SM can host one block or two concurrently. Two blocks per SM at 128 threads each is 256 resident threads/SM — well under the architectural per-SM thread cap, so register pressure (not threads) is usually the binding constraint at that occupancy target; nvcc/ptxas -Xptxas -v (wired into the build in crates/wh-iron-runtime/build.rs) reports the actual register count and any spills for a compiled kernel, and ncu (NVIDIA Nsight Compute) reports achieved occupancy, warp-active percentage, and memory-throughput percentages against the roofline for a real launch. Concretely: a real round in this codebase’s kernel work took a split-K pass-1 tile from 96 registers with zero spills to a register budget that supported two resident blocks per SM simply by re-tiling (BM, BN) from 32×32 to 32×64 — no algorithmic change, purely a shared-memory/occupancy trade.

The general workflow, either backend: change one shape knob, recompile, re-read the register/shared-memory report, and treat “does occupancy actually improve” as an empirical question — the estimate is a compass, the compiler’s actual allocation and a profiler’s actual occupancy counter are the ground truth.

3.3 Memory paths available in the DSL

In increasing order of how much the DSL does for you:

• Plain scalar loads/stores — load(buf[idx]) / store(buf[idx], v). Always correct, never contended for register/shared-memory budget beyond the value itself. The default; reach for something else only when you have a stated reason.
• Vectorized loads/stores — the vectorize pass promotes a run of provably-adjacent scalar loads/stores into one wide op (half2/half4, bfloat2/bfloat4, …) automatically; you write scalar DSL and get wide memory transactions when the addresses cooperate. See §8 for the address-binding trick that makes this fire reliably.
• Cooperative/simdgroup-matrix tiles — simdgroup_alloc::<T, R, C>() / simdgroup_elem_store on Metal, and the backend-neutral coop_tile_* family (coop_tile_setup, coop_tile_load_a/_b, coop_tile_run, coop_tile_get) that lowers to Metal simdgroup_matrix, CUDA wmma/inline-PTX mma.sync, or Vulkan coopmat depending on backend and eligibility (dtype, tile shape, architecture). This is the DSL’s tensor- core-class primitive, and it’s genuinely backend-portable: the same coop_tile_* call sequence, written once, is what the split-K GEMM and the SDPA MMA-prefill families use across all three backends that support it.
• Shared/threadgroup memory (threadgroup_alloc) — a named, function-scope allocation every kernel that stages a tile through shared memory uses. Every allocation is hoisted to function scope regardless of which branch declares it, which has a naming-collision trap covered in §8.

The CUDA-only advanced tier. Four more primitives, landing via PR #299, round out the memory-path vocabulary on CUDA. Each is documented here in full because it is exactly the knowledge a contributor needs before touching the kernels that use it — check iron inspect <kernel> against your own dev checkout before depending on any one of them in new code.

• cp_async_wait(n) — the DSL surface over cp.async-staged copies: asynchronous global→shared bulk copies that overlap with compute instead of blocking a warp on a synchronous load. n is a macro-time literal (the CUDA lowering bakes it straight into the PTX cp.async.wait_group operand) giving the number of outstanding cp.async groups still allowed to be in flight — cp_async_wait(0) drains all of them. Used to double-buffer a GEMM K-loop’s staging tile so the next tile’s load overlaps the current tile’s compute.
• 2-D tensor-map (TMA) bulk copies — TensorMap<T, box_cols, box_rows> (fixed box shape) and its dynamic sibling TensorMapDyn<T, box_cols, box_rows> (box shape resolved at dispatch time; encodes exact M/N so a boundary box zero-fills rather than needing a separate tail kernel) are kernel parameter types backed by a hardware bulk-copy descriptor. tma_load_2d(dst_tg, dst_offset, "map_param", x, y, "swizzle_mode") issues the copy — dst_offset is in bytes, not elements, and swizzle_mode ("full", "kfull", "swizzle128", …) must match the box’s actual row byte width or the copy faults at synchronization (see §8). coop_tile_load_a/coop_tile_load_b accept the same swizzle-mode string as a trailing argument (coop_tile_load_b("nt0", "Ks", true, T, 64u32, 16u32, 0u32, "swizzle128")) to read a swizzled TMA box straight into cooperative-tile fragments via ldmatrix.
• mbarriers — the barrier primitive a TMA copy signals through: mbarrier_init("name", count) declares an 8-byte mbarrier object backed by a threadgroup_alloc("name", 1, "u64"); mbarrier_arrive("name") and mbarrier_expect_tx("name", bytes) are the split producer-side signal (the arrive count and the expected-bytes transaction count), with mbarrier_arrive_expect_tx("name", bytes) as the race-free fused replacement for that pair — prefer the fused call in any new pipeline, per its own doc comment in wh-iron-macros; mbarrier_wait("name", phase) is the consumer-side wait. Ordering matters beyond the call sequence: the wait must lower to an unbounded try_wait.parity loop (a bounded version can exhaust its attempt budget and fall through silently under contention — this is a landed fix, not a hazard to design around, but it’s why mbarrier_wait never takes an attempt-count argument), and a producer’s last generic-proxy read of a staged block needs a fence_proxy_async() call — the DSL surface over PTX fence.proxy.async.shared::cta — before the arrive that lets the TMA engine’s async proxy overwrite that memory; skipping it is a real, timing-dependent race, not a theoretical one (see §8 for how it was found).
• threadgroup_alias("name", "base", byte_offset [, dtype]) — a zero-cost typed view into an existing threadgroup_alloc’d buffer base, starting at a 16-byte-aligned byte_offset: no additional shared memory is allocated, so one physical allocation can be read back under a different element type without a second declaration. byte_offset must be a literal or const-evaluable expression; the codegen’s alias validator checks the offset stays in bounds of base’s actual allocated size.
• Native FP8 matrix-multiply-accumulate — coop_tile_setup("name", m, n, k, .., u8, .., "e4m3") tags a cooperative tile as packed-E4M3, so coop_tile_run lowers to mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e4m3.f32 (sm_89+) instead of the software or wmma-fragment fallback — a real hardware tensor-core path for FP8 GEMM, not a promote-to-f16 shim. CUDA-only; coop_tile_setup(.., "e4m3") is rejected on other backends.

3.4 Determinism

Whether a kernel’s output is bit-exact run-to-run for identical inputs depends entirely on whether its accumulation uses atomics.

The split-K GEMM kernels in wh-iron-std are deterministic by construction: pass 1 writes each K-split’s partial [M, N] product to a distinct slice of a [n_splits, M, N] fp32 partials buffer (no cross-split write conflict), and a separate pass 2 sums the partials in a fixed order. No atomics anywhere in the path.

CUDA’s large dense-GEMM path is different: it calls into the vendor tensor-core GEMM library the CUDA toolkit provides, whose fastest heuristic can select a split-K strategy that reduces partial sums via atomic accumulation — and atomic accumulation order depends on which thread block happens to arrive at the atomic first, which is not guaranteed run-to-run. Iron’s CUDA runtime pins the handle to forbid the atomic-accumulation (split-K) code path by default specifically so identical inputs produce bit-exact identical outputs across runs — the alternative, left as an explicit opt-in (IRON_GEMM_ATOMICS=1) for A/B performance comparisons only, has shipped measurable “logit jitter” (small per-element differences that occasionally flip an argmax) on some MoE-prefill shapes where the nondeterministic heuristic was allowed to run. A second env var, IRON_GEMM_ALGO=N, lets you pin a specific deterministic algorithm variant for controlled experimentation without touching the atomics policy at all.

The rule this generalizes to: any kernel or kernel pair whose correctness depends on floating-point accumulation order should state, in its doc comment, whether it is bit-exact deterministic — and if it isn’t, why (atomics, or an intentionally non-associative fast-math reduction), and what the fallback is for a caller that needs determinism.

3.5 Portability across backends

The DSL’s contract is: one generic <T> kernel body, written once, compiles correctly for every dtype (f32/f16/bf16, …) it’s monomorphized over and for every backend it’s built for. Backend divergence is handled at the lowering layer, not by hand-writing four kernel bodies:

• Capability-gated strategy selection. MmaStrategy is the clearest example: the same coop_tile_* DSL call becomes simdgroup_matrix on Metal, wmma fragments or inline-PTX ldmatrix/mma.sync on CUDA depending on architecture and eligibility (dtype, tile shape being multiples of the fragment size, an accumulator dtype the hardware supports), or a plain software accumulation loop when nothing better is eligible — never a compile failure, always a working (if slower) fallback.
• Alignment and layout assumptions must be re-verified per backend. A “legal stride” that lets a load go through zero-copy on one backend (e.g. a leading dimension that is a multiple of the fragment width) can force a scalar repack on another if the alignment assumption doesn’t transfer — treat “does this shape hit the fast path” as a per-backend question, not something you can infer once and reuse.
• The DSL primitives that don’t exist everywhere are the exception, not the rule. The CUDA-only advanced tier in §3.3 (cp.async, TMA/mbarriers, threadgroup_alias, FP8 MMA) is CUDA-specific by nature — TMA and native FP8 MMA are hardware features of recent NVIDIA architectures with no Metal or Vulkan analog. Everything else documented in this guide is available to write against on every backend it’s relevant to.

---

4. The improvement loop

This is the method, distilled from the standing rules in docs/rust-efficiency.md and from a large body of kernel-optimization work: an internal performance campaign ran on the order of 70 rounds against this method against GEMM, GEMV, RMSNorm, SwiGLU, RoPE, and SDPA prefill/decode kernels across Metal, CUDA, and Vulkan, and the lessons below are what survived that process, each restated as this codebase’s own before/after (never as a comparison to any other system).

4.1 The loop

1. Baseline. Start from the current shippable path — correct, reproducible, with a passing correctness test. You cannot measure an improvement without a trustworthy starting point.
2. Pick a shape family with real inputs. Not a synthetic round number — the exact dtype, dimensions, batch width, and backend the consumer actually dispatches. A large-matrix result does not predict a one-row projection’s behavior; a short-context attention result does not predict long-context.
3. Establish the oracle first. Before touching performance, know what “correct” means for this exact shape and dtype — a CPU reference computed in f32 (or f64 for tighter numerical work) with a calibrated tolerance. See §5.
4. Profile and classify. Read the emitted source (iron inspect) and classify what’s happening: scalar vs. vectorized loads, shared-memory staging, barrier count, register count and spills. Then profile for real — ncu on CUDA, Metal’s GPU counters/--stats/timeline capture on Metal — and let the numbers, not intuition, say what’s actually the bottleneck (memory-bound, compute-bound, occupancy-limited, register-limited, or latency-bound; iron bench -v reports this verdict directly for kernels with a declared FLOP count).
5. Choose exactly one lever. Do not combine a memory-transfer change, a reduction-order change, and a dtype change in one experiment — doing so destroys your ability to attribute the result to a cause. Candidate levers, in roughly the order the campaign found them cheap-to-expensive: re-tiling (BM/BN/WM/WN, or the CUDA analog), vectorizing loads/stores, restructuring the reduction tree, fusing a producer’s epilogue into a consumer, and (CUDA-specific, on hardware that supports it) moving from software or wmma-fragment cooperative tiles to inline- PTX ldmatrix/mma.sync.
6. A/B honestly. See §6 in full — this is the step most optimization work gets subtly wrong, and it’s why a “flat” curve or a huge one-shot win deserves suspicion before belief.
7. Keep only real wins; record everything else. A candidate that loses, or that wins only in a regime that never reaches production, is not wasted work if it’s written down as a documented negative result next to the kernel or in the relevant benchmark doc — the next person (or agent) who has the same idea should find your answer instead of re-deriving it.

4.2 The levers catalogue

What tends to win, and what tends to lose, in this codebase’s kernel work — each stated as an Iron-internal before/after, not a comparison to any external system:

Tends to win:

• Re-tiling to fit two blocks per SM instead of one, when the limiting resource is shared memory rather than registers — a GEMM split-K pass-1 tile moved from 32×32 to 32×64 this way and measured a consistent several-percent reduction in pass-1 latency across repeated, order- reversed A/B sessions.
• Vectorizing the memory path once the compute is already tight — widening scalar loads/stores to 2- or 4-wide, once index arithmetic makes the addresses provably adjacent (see §8), reduces transaction count without touching the arithmetic.
• Choosing a dispatch shape that matches the operand geometry — the GEMV small-K regime in §3.1 is the canonical example: matching threadgroup/simdgroup granularity to the actual row width, not defaulting to “biggest threadgroup,” produced the largest single win in that family.
• Removing a redundant intermediate buffer by fusing a producer’s store directly into a consumer’s load, when the two share a compatible reduction/threadgroup shape — a materialize-then-reload round trip is pure overhead when nothing else needs the intermediate.

Tends to lose, or is a documented dead end:

• Wider tiles that don’t fit shared memory at the target occupancy — a wider CUDA GEMM tile variant was tried and rejected specifically because the extra shared-memory footprint dropped achievable occupancy from two resident blocks per SM to one, more than erasing the tile’s compute advantage.
• Software-pipelining a K-loop without first confirming registers have headroom — a pipelined variant can look identical in the emitted instruction count while quietly raising live-register count enough to spill, which shows up as a regression only under profiling, not by reading the source.
• Threadgroup-memory staging for a value that’s already well-served by cache — adding a shared-memory round trip for reuse the hardware cache already captures adds barrier overhead without removing real traffic; measure against the best already-resident baseline before assuming shared memory helps.
• A “flat” latency curve across meaningfully different input sizes — this is almost never a real result. It means the harness, not the kernel, is being measured (see §6.3).

The CUDA advanced-tier levers, using the §3.3 primitives, land via PR #299. Each below is a real, measured win, restated as this codebase’s own before/after — the numbers are the exact A/B results recorded for the round, not re-run for this guide:

• TMA double-buffered K-loop, zero __syncthreads. The split-K GEMM champion’s K loop rewritten onto 2-D TMA boxes (A/Bᵀ tiles as tensor-map boxes at SWIZZLE_64B, two slots with full/empty mbarriers, one fused mbarrier_arrive_expect_tx plus two tma_load_2d calls per k-step, parity re-derived from the loop index) drops register count from 96 to 72 and shared-memory footprint from 42.0 KB to 36.9 KB at the same two-blocks/SM occupancy, with zero __syncthreads calls in the K loop — the mbarrier pair replaces the block-wide barrier entirely. Measured [96,2048]×[2048,512] 0.0126 → 0.0086 ms (−32%), [256,2048]×[2048,512] 0.0270 → 0.0188 ms (−30%) (both orders, 20/40 session, PASS receipts before timing).
• An occupancy-driven split count for TMA split-K GEMM. Rather than a fixed split factor, deriving it from clamp(ceil(2·SMs/(⌈M/32⌉·⌈N/128⌉)), 1, K/64) snapped to a divisor of K/64 lets one kernel cover a whole shape family instead of one shape: measured against this codebase’s own row-oriented GEMV kernel across M ∈ {1, 26, 32, 49, 96, 256} at N=96, K=5120, M=1 0.0085 → 0.0044–0.0065 ms (−24…−48%), …, M=256 0.1601 → 0.0106 ms (−93%), direction held across 3 sessions and both dispatch orders.
• Native FP8 (e4m3) tensor-core MMA for block-scaled quantized GEMM. coop_tile_setup(.., u8, .., "e4m3") driving a real mma.sync.aligned.m16n8k32...e4m3.e4m3.f32 path, measured against this codebase’s own FP16-decode reference kernel: M=1 1331 → 883 µs (1.51×), M=37 2581 → 1072 µs (2.41×), M=128 4357 → 3288 µs (1.33×). (An earlier pass at this same kernel reported far larger multiples; those numbers were retracted — see the ei/degenerate-fixture gotcha in §8 for why, and §6.1’s roofline-sanity rule for the general lesson.)

A fused single-dispatch attention prefill that keeps softmax state in cooperative-tile registers across the whole KV loop (instead of materializing an intermediate score matrix) is in progress on the same PR; its numbers are not final enough to restate here yet.

---

5. Correctness and testing

Correctness is checked in layers. No single layer is sufficient — each catches something the layer above cannot, and a change that only satisfies the cheap layers can still ship an all-zeros kernel.

Layer	Catches	Lives in
DSL / codegen unit tests	Pass correctness, body-parser arms, IR variants, emit paths	wh-iron-codegen, wh-iron-core, wh-iron-macros
Test-inventory witness	Silent registration shrinkage, dead-code stripping, renamed cases, dtype/tolerance drift	crates/wh-iron-std/tests/kernel_test_inventory.rs
Test-inventory drift check	A file with #[test]/#[test_kernel]/#[bench] and no audit row; a stale rendered inventory doc	make check-test-inventory
MSL snapshots	Codegen output drift — a reviewable text diff in the PR	insta fixtures in wh-iron-codegen/tests/msl_snapshots.rs
GPU correctness	Numeric disagreement vs. a naive CPU oracle, on a real device	<kernel>_gpu_correctness.rs or an in-file #[test_kernel]
Backend-specific GPU corpus	Target lowering, packed layout, launch geometry, concurrency/lifetime edge cases	wh-iron-runtime/tests/, wh-iron-std/tests/
Side-by-side bench	Throughput and output-equivalence vs. a fixed reference kernel	iron bench --reference (opt-in, local-only)

5.1 Writing a correctness test

Every non-trivial kernel ships a GPU correctness test in the same commit that adds or changes it. Two equivalent ways to write one:

• Declarative, next to the kernel (#[test_kernel], the newer, preferred style): a kernel_tests module in the same file as the kernel, each test a fn(dt: DType) -> TestSetup that builds inputs, computes an f32 (or f64) CPU reference, and calls .expect(...). Dtypes and per-dtype tolerances are declared right on the attribute: #[test_kernel(dtypes = [f32, f16, bf16], tol = [1e-3, 1e-1, 1.0])].
• Hand-written, in crates/wh-iron-std/tests/<kernel>_gpu_correctness.rs for cases that need more setup than the declarative form supports — dispatches the kernel via Context::dispatch_with_grid directly against a real Metal device and asserts against a naive CPU reference.

The naive CPU reference is the contract. If the kernel and the reference disagree, work out which one is wrong before merging — never loosen a tolerance to make a disagreement go away.

Calibrate tolerances from the dtype’s real precision, not a copy-pasted default: an f32 reduction over few elements can hold 1e-5; a transcendental (exp/log/sqrt/softplus) kernel usually needs on the order of 1e-3– 1e-4 in f32 because GPU fast-math intrinsics diverge from a CPU libm reference and the exact divergence varies by GPU family — a tolerance tuned tight enough to pass on one device can fail on another. f16 and bf16 tolerances are set by the dtype’s storage precision, not by wanting a smaller number: something in the range of 1e-1–0.5 relative error is often correct and expected for bf16’s ~3 decimal digits of precision, and tightening it “to be safe” just makes the test flaky.

5.2 Running tests

# The whole workspace: codegen, runtime, GPU correctness.
make test

# One kernel's declarative correctness cases, on one backend:
iron test -f <family> [--backend metal|cuda|hip|vulkan]

# One kernel's hand-written correctness test:
cargo test -p wh-iron-std --test <kernel>_gpu_correctness

# The same test's #[ignore]'d perf companion:
cargo test --release -p wh-iron-std --test <kernel>_gpu_correctness -- --ignored --nocapture

iron test defaults to the host’s native backend; --backend targets a different one when the hardware and feature-gated build are available. A hardware test that requires a device/feature/fixture that isn’t present must fail closed when explicitly selected — a silent early return is not skip accounting and cannot support a coverage claim (see rst-testing.md’s “campaign lessons”).

5.3 Ignored hardware tests, and stress under contention

A test that needs specific hardware or a slow full sweep is marked #[ignore] and run explicitly (-- --ignored --nocapture), not run by default. Every ordinary GPU correctness test above dispatches one kernel at a time on an otherwise-idle device — a condition under which pool- and chain-ordering defects stay invisible. Two tests exist specifically to close that gap by running a background-load Context on its own thread while asserting invariants under contention: metal_buffer_pool_stress_gpu.rs (a multi-pass fused dispatch chain’s untracked buffers stay correctly ordered across passes under concurrent GPU load, pooled buffers never hand back a stale payload, and a per-device buffer pool stays coherent across multiple Contexts and threads) and its CUDA twin cuda_buffer_pool_stress_gpu.rs (the raw allocator doesn’t hand a freed-but-still-in-flight pointer back to a new allocation, and pooled churn under concurrent GPU load stays bit-exact).

Repeated-launch bitwise stress, and stress under contention. Any kernel using the §3.3 CUDA advanced tier (mbarriers and cp.async in particular) needs a stronger check than the single-launch tests above: dispatch the same kernel several hundred times back to back and assert every launch’s output is byte-identical. A single launch essentially never surfaces a low-probability mbarrier/cp.async race — one recorded case landed as a 1-in-10 silent corruption rate under a bounded mbarrier_wait before the fix that made the wait unbounded. CudaDevice::run_kernel_stress(kernel, buffers, grid, block, out_name, launches) makes this practical to run routinely: it prepares the kernel once and loops launches with a full device sync between them (matching run_kernel’s per-launch semantics), so 300 launches take on the order of seconds rather than the tens of minutes a naive prepare-per-launch loop costs.

Uncontended 0/N is necessary but not sufficient. A bitwise stress that passes with the GPU idle but shows mismatches once another process shares it is the signature of a synchronization race, not measurement noise — GPU contention changes timing, never the arithmetic of a race-free kernel. The standing requirement for any mbarrier/TMA kernel before it lands is 0 mismatches under deliberate contention (a concurrent synthetic GPU load running alongside the stress loop), not just 0/N on a quiet device. This is exactly how a real ordering bug was root-caused rather than dismissed as noise: an SDPA prefill kernel’s TMA path showed 1/999 launches contended even after tightening mbarrier_wait to .acquire.cta, tracing to a missing fence_proxy_async() between a warp’s last read of a staged block and the arrive that lets the TMA engine overwrite it (see §3.3 and §8); adding it took the same kernel to 0 mismatches across roughly 3,000 launches at both a contended and an uncontended sequence length.

5.4 Selector tests, MSL snapshots, and the test inventory

Selector tests are ordinary #[test]s that assert a dispatch-selector function (<family>_for(...), see §2.5) routes a given shape/dtype/capability combination to the kernel you expect — cheap, host-only, and the thing that actually protects a multi-kernel family from a routing regression when a new tile shape is added.

MSL snapshots (insta fixtures in wh-iron-codegen/tests/msl_snapshots.rs) pin the exact emitted MSL for a hand-built kernel through MslGenerator. They exist to make codegen drift a reviewable text diff, not to be exhaustive — add one when a new DSL primitive or fusion pattern lands that no existing snapshot exercises. Refresh intentional changes with cargo insta review or cargo insta test --accept. A snapshot pins output; it proves nothing about whether that output is correct — only a GPU correctness test does that.

The test inventory and its witness. Every file that declares a #[test], #[test_kernel], or #[bench] needs a row in docs/test-inventory-status.tsv; docs/test-inventory.md is a generated, reviewable render of that file (scripts/render-test-inventory.sh). make check-test-inventory fails if the rendered doc is stale or a declaring file has no row — this is a required gate, including for a docs-only change, because it’s cheap to run and catches the inventory silently drifting out of sync with the source tree. Separately, crates/wh-iron-std/tests/kernel_test_inventory.rs is a witness test: it hashes the sorted, canonicalized set of every registered #[test_kernel] case (file, name, dtype, tolerance bits) against a pinned expected count and FNV-1a digest. It exists to catch silent registration shrinkage — a case that link-time dead-stripping removed, a case someone deleted without noticing, a tolerance that got quietly loosened — that a normal “tests still pass” run would never surface. An intentional corpus change updates the pinned count and digest as part of that PR’s review, not as a drive-by fix.

---

6. Measurement protocol

A performance number is only as good as the discipline that produced it. The rules below are standing project policy (docs/rust-efficiency.md’s “Benchmark discipline” and “Optimization loop”), reinforced by dev-branch CUDA benchmark reports (docs/benchmarks/) that apply them concretely.

6.1 Before you time anything

• Realistic inputs. Random bytes are not safe input for a floating-point kernel — they alias to inf/nan often enough to poison a comparison. Benchmarks seed inputs from a bounded, deterministic, nan-free domain.
• A PASS receipt before every timed run. Never time a candidate whose correctness you haven’t already established at that exact shape and dtype in that same session — a “fast” wrong answer is worthless and a documented failure mode (a small quant-format oracle bug in this codebase passed unnoticed until a correctness test caught it; nothing about its benchmark numbers would have flagged it).
• Resident buffers for inputs constant across iterations. Otherwise you measure host→GPU upload bandwidth, not the kernel — this exact mistake produced a “flat ~215 µs regardless of context length” result that turned out to be upload overhead; switching to resident buffers dropped the floor and revealed the real, shape-dependent curve.
• Warm the clock before the first sample. A cold GPU runs at a lower DVFS clock state; a dummy dispatch before the timed loop avoids charging that penalty to whichever shape happens to run first.
• A win far beyond the roofline is a bug, not a result. Before publishing a GEMM/GEMV number, compute achieved bytes/s (or FLOPs/s) and compare it to the device’s memory (or compute) roofline — a number that implies bandwidth above what the part can physically deliver is not a discovery, it’s a measurement or correctness defect. A block-scaled FP8 GEMM kernel once reported a per-element time on a GB10 that implied roughly 360 GB/s moving an 85 MiB matrix, on a part whose ceiling is roughly 250 GB/s; the number was impossible, and it was — the honest, fixed kernel measured roughly 100 GB/s. See the ei/degenerate-fixture gotcha in §8 for the bug that produced it.

6.2 While timing

• Warmup, then enough samples to see the distribution — the standing convention in this codebase’s benches is on the order of 20 warmup iterations followed by enough measured samples (dozens to low hundreds, depending on variance) to report a median, not a single sample. iron bench -vv additionally reports the distribution shape (p95/p99, coefficient of variation).
• Both dispatch orders. Run candidate-then-baseline and baseline-then-candidate, not just one order — a directional bias (thermal drift, cache warm-up asymmetry) that only shows up in one order is a real confound, and requiring both orders to agree is cheap insurance against it.
• Multiple independent sessions, not one long run — process-level effects (a stray background load, a driver JIT cache miss) can bias a single session in a way repeated fresh processes wash out.
• A GPU lock so runs don’t overlap. On a shared machine, an unguarded concurrent run silently corrupts every number it touches; the standing convention here is a simple file lock (flock) around the timed section so two invocations never contend for the device at once.
• Report the min and the median, not just one — the min is closer to the kernel’s true achievable latency (freest of transient noise), the median is closer to what a real deployment experiences.
• A keep bar, not a vibe. A candidate that wins by less than roughly 3% is inside typical measurement noise for a single kernel launch and is not a documented win — the project’s convention is a ≥3% bar per shape, usually checked as both the raw median improvement and a bootstrap or paired-repetition confidence interval clearing that bar, not a single best-of-N sample.

6.3 Flat curves are a harness bug until proven otherwise

A latency that does not scale with input size when input size changes meaningfully is the single most common tell that you are measuring the harness, not the kernel — see the resident-buffers example above. Before publishing any number, sanity-check that the curve has the shape physics predicts (roughly linear in bytes moved for a memory-bound kernel, roughly linear in FLOPs for a compute-bound one at large enough sizes) and treat any of the following as a harness warning, not a result: latency doesn’t scale when work scales, a result changes under reordering, the first measured cell is consistently the slowest, an explicitly unsupported code path reports as a pass, or a microbenchmark improves while the shape that actually matters to the consumer does not move.

6.4 Kernel-level wins don’t always transfer

A kernel-level A/B win is necessary but not sufficient — it has to survive contact with the actual consuming pipeline. A win at the kernel’s own launch boundary can be swallowed entirely by a phase that dominates end-to-end time more than the kernel does, or by a fixed per-launch host-side dispatch cost that a tiny kernel’s compute time doesn’t clear. The standing rule (rust-efficiency.md’s optimization loop, step 8) is: bound the plausible end-to-end gain before investing in a kernel change (multiply the candidate phase’s share of total time by the plausible local improvement), and re-measure at the consumer boundary once the kernel change lands, not just at the kernel’s own boundary — a kernel win with no consumer-level signal is either genuinely too small to matter yet, or evidence the wrong phase was optimized.

---

7. Gates and PR hygiene

7.1 Local gates

make fmt-check              # cargo fmt --all -- --check
make clippy                 # cargo clippy --all-targets --all-features -- -D warnings
make typos                  # typos --config .github/configs/typos-cli.toml
make test                   # cargo test --workspace
make check-test-inventory   # docs/test-inventory.md matches the rendered inventory

make hooks installs three git hooks that mirror the cheapest of these locally: pre-commit runs fmt-check + typos; commit-msg scans for AI-attribution trailers and lookalike trailer shapes; pre-push runs clippy plus the kernel-registry consistency check. They catch the fast, common failures before a slower CI round-trip; the full GPU correctness suite and the full bench harness still only run in CI (or locally against real hardware).

7.2 CI

Job	Runs
Lint	typos, buffer-binding check, make check-test-inventory, cargo fmt --check, clippy — Linux, no GPU
Build / test / bench	iron build, iron test (GPU correctness vs. CPU oracle), iron bench — macOS GPU runner
Coverage	cargo llvm-cov --workspace on pushes touching crates/, Cargo.*, or the toolchain pin
Spark (opt-in, PR label)	A CUDA correctness suite (binary add/multiply, GEMV, RMSNorm — all vs. a CPU oracle) on real CUDA hardware, plus an optional short performance sample; requested per-PR via a spark-test / spark-test-long label, not run by default

A docs-only change (no crates/, Cargo.*, or toolchain-file touch) skips the heavy build/test/bench and coverage jobs but still runs Lint — meaning typos and make check-test-inventory are real, required gates for a documentation PR, this one included.

7.3 Commits and PRs

• Conventional-commit format (feat:, fix:, perf:, docs:, test:, chore:, …) — CI validates the PR title against it, and it drives release-note categorization.
• No Co-Authored-By: trailers, no AI-attribution footers. The commit-msg hook scans for exactly this and for trailer-shaped lookalikes. Disclosing AI use in prose (the PR body, a code comment) is fine and, per CONTRIBUTING.md, expected — it’s trailers and footers in the commit message that are rejected.
• One logical change per PR, tests and docs landing with the code that needs them — not as a follow-up.
• Keep losers as documented negatives. A rejected optimization attempt, written down next to the kernel or in a benchmark doc with the numbers and the reason it lost, is worth as much to the next contributor as a shipped win — it stops the same idea from being re-tried blind.
• If bench numbers changed, paste the rows that moved in the PR body, not a description of them.

---

8. Gotchas catalogue

Each entry: the symptom, then the fix. Pulled from docs/developing.md’s “Kernel-authoring hazards”, from the PR #299 record, and verified against the current source in each case.

• A wrong dispatch geometry can freeze the machine. Metal dispatch is non-preemptive — an infinite loop inside a kernel never yields, the compositor starves of GPU time, and a hard power-cycle is the only recovery. The concrete trap: a reduction kernel’s simdgroup count is lsize / 32 (integer division); a loop strided by that count becomes an infinite GPU loop when it’s dispatched with fewer than 32 threads per threadgroup. Fix: threads-per-threadgroup for any kernel using simd_*/threadgroup_* must be a multiple of 32 and at least 32, derived from the kernel’s own invariants — never from an unrelated “number of elements” count.
• Grid is counted in threadgroups, not threads. grid_xyz · tg_xyz must equal exactly the thread count the kernel expects; grid=[N,1,1] tg=[N,1,1] for a one-thread-per-element kernel dispatches N² threads, most with garbage indices. Fix: derive grid from the kernel’s actual per-threadgroup output count, not from the input element count directly.
• An inner macro_rules! call inside a #[kernel] body silently empties it. The proc-macro doesn’t expand inner declarative macros — it sees the call as opaque tokens, drops them, and emits a kernel with no body. xcrun metal and MSL snapshots both accept it; only a GPU correctness test sees the resulting all-zeros output. Fix: wrap the entire #[kernel] fn declaration in an outer macro_rules! instead — the compiler expands that before the proc-macro runs, so the body parser sees concrete tokens. This shape is now a compile error rather than a silent failure; crates/wh-iron-std/src/kernels/sdpa/sdpa_prefill_sink.rs is the canonical reference for the outer-macro pattern.
• threadgroup_alloc names collide across branches. Every call is hoisted to one function-scope declaration regardless of which if branch it’s written in, so two branches each naming a buffer "tg_max" collide into a duplicate declaration — a loud Metal compile error, but the fix isn’t obvious from the error alone. Fix: give every threadgroup buffer a name unique across the whole kernel, not just its branch.
• Pass ordering can leave an empty loop body. A pass that eliminates a loop’s body but not its header, or that consumes a Const before the pass that produces it has run, can leave for (...) { } — accepted by xcrun metal, invisible to MSL snapshots, ships all-zeros. Fix, if you’re writing a pass: always walk both kernel.body and every entry in kernel.blocks; a pass that removes a loop body must also remove the loop header; order dependencies between passes (e.g. fma_fusion after type_check) are load-bearing, not incidental.
• cargo run -p wh-iron-cli can silently run stale codegen. The real work — codegen, MSL emission, every bench/test dispatch — runs inside __iron_runner, a sibling subprocess iron spawns; cargo run -p wh-iron-cli only rebuilds iron itself, never the runner, and iron doesn’t check the runner it finds is fresh. Fix: go through make (every wrapper depends on make runner, a cheap fingerprint-checked rebuild); if you invoke the CLI directly, rebuild the runner first with the same profile.
• select does not guard a load. select(cond, load(buf[i]), other) evaluates both operands — the load happens regardless of cond — so an “optional buffer, never read when the flag is off” contract silently breaks the moment the placeholder buffer is smaller than what an unconditional load would index. Fix: put the load inside real if control flow, not a select branch; the if-conversion pass leaves an assignment-in-an-arm shape alone, so the emitted code stays a real branch.
• CSE does not unify repeated address subtrees. Common-subexpression elimination unifies identical value computations, not repeated base + offset address arithmetic written out twice — which also means the vectorizer, which needs to see adjacent addresses as one computed base, may not fire. Fix: bind a repeated address expression to a let yourself before issuing the loads/stores that use it.
• A repeated reduction can read a shared scratch slot mid-overwrite. A barrier before a shared-memory read does not protect that read from a write in the next loop iteration — one SIMD group can finish reading a broadcast scratch slot while another group is still writing the next iteration’s value into it. Fix: every SIMD group must reduce shared partials into a register before the next barrier that could let another group overwrite the scratch; a custom broadcast-through-shared-memory reduction needs an extra barrier after every group has consumed the slot.
• A “too flat to be physical” latency is a harness artifact, not a discovery — see §6.3.
• A TMA swizzle atom must match the box’s actual row byte width, or the copy faults at synchronization time rather than at issue time, which makes the fault hard to trace back to the call that caused it. Fix: pick the tma_load_2d/coop_tile_load_a/coop_tile_load_b swizzle-mode string ("full", "kfull", "swizzle128", …) to match the box’s real row width in bytes, not its width in elements, and re-verify it whenever the box shape changes.
• A TMA dst_offset is in bytes, not elements. Passing an element-count offset where a byte offset is expected silently overlaps transactions in shared memory instead of failing loudly. Fix: multiply by the element size explicitly at every tma_load_2d call site; don’t reuse an index variable that’s counted in elements elsewhere in the same kernel.
• A swizzle applied before folding in a call’s own offset corrupts the swizzle for every sub-tile read at a non-zero offset inside a box. A real round of this codebase’s kernel work shipped coop_tile_load_b(.., "swizzle128") applying the XOR to (row, column) before folding in the call’s ptr_offset, so the first sub-tile in a box read correctly and every subsequent one didn’t. Fix: fold any per-call offset into (row, column) first, then apply the swizzle transform — order matters here, and a test that only exercises offset zero cannot catch getting it backwards.
• An mbarrier wait implemented as a bounded retry loop can fall through silently instead of actually blocking on the barrier — a version bounded at roughly 2M attempts fell through under contention and produced a 1-in-10 silent-corruption rate at one sequence length, not a hang (which would at least have been noticed). Fix: mbarrier_wait must lower to an unbounded try_wait.parity loop or trap explicitly — never “proceed anyway” after N attempts.
• A missing async-proxy fence before an mbarrier arrive is a real, timing-dependent race, not a theoretical one. Two separate PTX memory-model gaps can each let a TMA engine overwrite a staged block while a warp is still reading it through the generic proxy: .acquire missing from mbarrier.try_wait.parity (a stale cache line can satisfy a post-wait read), and no fence.proxy.async.shared::cta separating a warp’s last generic-proxy read from the arrive that signals the buffer is free. Both were needed together — fixing only one still left roughly 1-in-1000 launches contended. Fix: .acquire.cta on every mbarrier wait, and a fence_proxy_async() call before every producer-role arrive/arrive_expect_tx that follows a generic-proxy read of the same memory; verify with the repeated-launch stress protocol under deliberate contention (see §5.3), not just a single correctness run.
• A non-literal stride argument to a cooperative-tile load can silently become zero instead of failing to compile. coop_tile_load_a/_b’s row-stride argument was originally parsed as a macro-time literal only; passing a runtime value silently lowered to 0, zeroing that term of every load’s address so every row read the same bytes. The bug shipped undetected because the correctness fixture that exercised it was degenerate — its per-row hash happened to be byte-identical (or sign-negated) across every legal stride, so a wrong-but-structured answer still looked plausible enough not to trip the oracle’s tolerance. The benchmark numbers taken against the broken kernel looked like a huge win (see the roofline-sanity rule in §6.1) — the honest, fixed kernel’s numbers are a fraction of what was originally reported. Fix, generalized: a macro-time-literal-only argument that silently accepts a runtime expression is a miscompile waiting to happen — make the parser reject the non-literal case loudly instead of defaulting it; and calibrate correctness fixtures for full avalanche (every legal input should produce a value that changes when any input byte changes), not just numeric plausibility — a fixture that happens to be order-invariant or sign-symmetric can pass a broken kernel for a long time.
• A kernel-level A/B win does not always survive the real dispatch path. A CSE gap — the DSL repeating column*k+base as a fresh address computation at each of several loads, which ConstFold+CSE never unifies into one subtree — meant the vectorizer saw no contiguous run even though the underlying loads were genuinely adjacent; binding the repeated address to one let let it fire and won a few percent in an isolated kernel harness. Through the actual consumer dispatch path, though, the same change measured no improvement at one input size and a real regression at another — a roughly constant per-dispatch floor (encode/graph-sync overhead common to every code path at that shape) was large enough to swamp the tens of microseconds the wider load bought. Fix: treat a kernel-harness win as a hypothesis, not a result, until it’s re-measured at the actual consumer dispatch boundary (see §6.4); when it doesn’t transfer, the next lever to try is usually one that changes the dispatch shape itself (fewer, larger launches), not a further micro-optimization of the kernel body.
• A per-tile-name codegen helper can have a side effect that’s easy to miss. A cooperative-tile capacity/coordinate helper keyed by tile name silently zeroed two of three M-blocks in one internal experiment because the side effect was per-name rather than per-call-site — documented in the relevant codegen module as a reminder to check side effects against call site, not just tile identity, when writing a new cooperative-tile helper.
• A harness’s own per-launch host cost can dominate a small kernel’s measured time. For a kernel small enough that its GPU compute time is comparable to command-buffer/stream launch overhead, that fixed host cost can be the entire measured signal — always sanity-check a small kernel’s measured latency against a known launch-floor figure before attributing it to the kernel’s own work.

---

9. Where to start (for new contributors)

9.1 First task — add a small kernel

1. Read docs/developing.md’s kernel-authoring hazards section and docs/STYLE_GUIDE.md end to end first — both are short, and both prevent real mistakes.
2. Pick a small, self-contained elementwise or reduction operation that doesn’t exist yet, or a variant of one that does (a new dtype, a new fixed shape specialization).
3. Write the kernel file following the file skeleton in STYLE_GUIDE.md: the kernel function, a kernel_tests module with at least one #[test_kernel] covering every dtype you support with calibrated per-dtype tolerances, and a kernel_benches module with at least one #[bench] at a realistic shape.
4. Put the file in crates/wh-iron-std/src/kernels/<family>/, matched to the operation, not to whether a reference implementation exists for comparison.
5. Add its row to docs/test-inventory-status.tsv and regenerate docs/test-inventory.md (scripts/render-test-inventory.sh > docs/test-inventory.md, or just run make check-test-inventory and follow its diff).
6. Run the full local gate list from §7.1 before opening a PR. iron test -f <your kernel> to confirm it passes on the backend(s) you can reach locally.

9.2 Second task — vectorize an existing kernel and A/B it

1. Pick an existing kernel with a #[bench] already registered.
2. iron inspect <kernel> --stats to see its current pass-by-pass op counts, then iron inspect <kernel> to read the emitted source and confirm whether its loads/stores are already vectorized.
3. If not, and the addresses are adjacent, try binding the shared address expression to a let (see the CSE/vectorize gotcha in §8) so the vectorizer can see it, and re-check iron inspect for the wider load/store ops appearing.
4. A/B against the pre-change kernel following the full protocol in §6 — both orders, multiple sessions, a correctness PASS before every timed run.
5. If it’s a real win (≥3%, held across orders and sessions), keep it and note the before/after numbers in the PR body. If it isn’t, that’s a valid, useful outcome — say so and why.

9.3 Picking a shape family, and where the numbers live

Pick shapes from what the consumer actually dispatches, not round numbers — if you don’t know the production shape for an operation, that’s worth finding out before optimizing it. docs/benchmarks/<topic>-<date>/ holds past benchmark reports in the format described in §6; reading two or three of the CUDA ones is the fastest way to see the measurement protocol applied for real, with actual numbers, actual sample counts, and actual correctness receipts.

9.4 Whom to ask, how to record findings

For anything beyond a trivial fix, open an issue first to align on scope — a short exchange there saves rework on the PR (CONTRIBUTING.md). Record a negative result the same place you’d record a positive one: next to the kernel in a doc comment, or in a new or existing docs/benchmarks/ report, with enough detail (shape, both orders, session count, the actual numbers) that someone else can trust it without re-running it themselves.

---

10. Glossary

Term	Meaning
DSL	The Rust-embedded language inside a #[kernel] function body.
IR	Iron’s typed kernel graph (Kernel/Op), produced by the proc-macro, consumed by the codegen pipeline.
MSL	Metal Shading Language — Apple’s GPU shader language and the Metal backend’s emitted output.
PSO	Pipeline state object — a compiled Metal compute pipeline; runtime-cached so a kernel compiles once.
Pass	One transformation stage in the codegen pipeline (e.g. ConstFold, Vectorize).
Threadgroup / block	A group of GPU threads that share local memory and can barrier-synchronize (Metal: threadgroup; CUDA: block).
Simdgroup / warp	A hardware-scheduled lockstep group of threads within a threadgroup/block (Metal: simdgroup, 32 lanes; CUDA: warp, 32 threads).
TPG	Threads per threadgroup — the threadgroup dimension of a dispatch.
#[kernel]	The proc-macro that parses a Rust function body into Iron IR.
#[test_kernel]	The declarative correctness-test attribute; runs the kernel on a real device against a CPU oracle.
#[bench]	The declarative benchmark attribute; run via iron bench.
variants(...)	A #[kernel] attribute argument that expands one generic body into several concrete kernels at compile time over an axis (tile shape, block width, …).
Selector	A plain Rust function, <family>_for(...), that picks which concrete kernel a shape/dtype/capability combination should dispatch to.
Cooperative tile (coop_tile_*)	The DSL’s backend-neutral tensor-core-class primitive; lowers to simdgroup_matrix (Metal), wmma/mma.sync (CUDA), or coopmat (Vulkan).
Split-K	Partitioning a GEMM’s reduction (K) dimension across threadgroups/blocks so a large-K, skinny-M/N matmul still saturates the device; requires either a two-pass partials-buffer accumulate (deterministic) or atomic accumulation (not deterministic).
Occupancy	How many threadgroups/blocks can be simultaneously resident on one compute unit, bounded by whichever of registers, shared memory, or thread-count limits runs out first.
Witness test	A test that hashes a registered corpus’s identity (names, dtypes, tolerances) against a pinned count and digest, to catch silent registration shrinkage that a normal test run wouldn’t flag.
cp.async	CUDA’s asynchronous global→shared copy instruction family; the DSL surfaces its drain point as cp_async_wait(n).
TMA (tensor map)	Tensor Memory Accelerator — a hardware bulk-copy descriptor (TensorMap/TensorMapDyn DSL types) that moves a 2-D box between global and shared memory via tma_load_2d, signaled through an mbarrier.
mbarrier	The barrier primitive a TMA (or cp.async) copy signals completion through; DSL surface: mbarrier_init/_arrive/_expect_tx/_arrive_expect_tx/_wait.
threadgroup_alias	A zero-cost typed view into an existing threadgroup_alloc’d buffer at a given byte offset, with no additional shared memory allocated.
FP8 (e4m3) MMA	Native 8-bit-float matrix-multiply-accumulate on supporting CUDA hardware, reached via coop_tile_setup(.., u8, .., "e4m3"); lowers to mma.sync...e4m3.e4m3.f32 rather than a promote-then-multiply fallback.

---

11. References

Everything below lives under docs/ in this repository unless noted:

• developing.md — repo layout, dev loop, branching, commits, and the kernel-authoring hazards.
• STYLE_GUIDE.md — the authority on how to write one kernel file: naming, variants(...), the CPU oracle, the bench.
• testing.md — the test layers, what runs where, and the documented gaps in the test infrastructure.
• rst-testing.md — the project’s Rapid Software Testing method: risk framing, oracles, coverage models, and reporting discipline for a change or investigation session.
• rust-efficiency.md — the project’s standing rules for compiler/runtime efficiency, kernel fusion, occupancy, and benchmark discipline; this guide’s §4 and §6 are built directly on it.
• spark-ci.md — the opt-in CUDA hardware CI lane.
• cli.md — the full iron CLI reference.
• specs/CUDA_BACKEND_SPEC.md, specs/AMD_BACKEND_SPEC.md, specs/VULKAN_BACKEND_SPEC.md — the backend-seam design and each backend’s delta on it.
• specs/ARCHITECTURE.md — the source-to-shader path and evidence lifecycle in more depth than §2 covers.
• docs/benchmarks/ — dated benchmark reports; the closest thing to a worked example of §6 applied for real.
• CONTRIBUTING.md — PR process, the agentic- contribution disclosure requirement, and dependency policy.