CLI

iron is the command-line driver for benchmarking, building, and inspecting kernels. Install it, or run it through cargo from a checkout.

cargo install --path crates/wh-iron-cli      # installs the `iron` binary
# or, from a checkout, without installing:
cargo run -p wh-iron-cli -- <command> …

make bench wraps iron bench; for the other subcommands run iron (or the cargo run form) directly.

iron bench — benchmark Iron kernels

Benchmarks the Iron kernels and reports wall-clock latency, throughput (GB/s), compute throughput (GFLOP/s), and roofline figures. By default it benches only the Iron kernels; pass --mlx to also run each kernel’s MLX reference for a side-by-side speed A/B plus an output-equivalence check.

iron bench [-f <substr>] [--mlx] [-v|-vv] [-o <file.json>] [--allow-dirty]
           [--diff] [--baseline-ref <git-ref>]

Flag	Effect
-f, --filter <substr>	only run kernels whose name contains <substr>
--mlx (alias --reference)	also run each kernel’s MLX reference: the Ref / Iron % columns and the output-equivalence check. Off by default (the wh-iron kernels have superseded the references; correctness lives in iron test); roughly doubles bench time
-v / -vv	-v adds the roofline (%BW / %FLOP / arithmetic intensity), occupancy/registers, and a bottleneck verdict (plus the reference latency when --mlx is set); -vv adds the GPU timing distribution (p95 / p99 / cv%)
-o, --json <file>	also write results as JSON
--allow-dirty	run on a dirty working tree (default: refuses, so numbers tie to a clean SHA)
--diff	opt into the post-bench diff against the target-branch baseline
--baseline-ref <ref>	git ref whose baselines/<chip>.json to diff against (default: first of origin/dev, upstream/dev, dev)

Metrics

The default table shows, per kernel/dtype: Iron(µs) (wall-clock latency, the min sample — the metric that makes “which precision is fastest” directly readable), Iron (GB/s bandwidth), GFLOP/s (compute throughput, blank for memory-bound kernels), and ok (correctness). With --mlx it also fills the Ref (MLX GB/s) and Iron % (Iron-vs-MLX ratio) columns; without it those stay blank.

-v adds the roofline view: %BW (achieved ÷ the device’s peak DRAM bandwidth), %FLOP (achieved ÷ peak compute — the M5 Neural-Accelerator FP16 ceiling where applicable, the SIMD pipe otherwise), AI (arithmetic intensity, FLOPs/byte), the estimated occ%/regs, and a combined bottleneck verdict (memory-bound / compute-bound / occupancy-limited / register-limited / latency-bound). Peak ceilings come from a per-device table (crates/wh-iron/src/runner/device_specs.rs); an unknown GPU leaves the roofline columns blank rather than failing.

GFLOP/s, latency, and the roofline figures only appear for kernels that declared a FLOP count (#[bench(flops = …)] or BenchSetup::flops) — matmul, attention, and convolution; memory-bound elementwise/reduction kernels leave them blank. The JSON (-o) is additive: it keeps the ref/iron (GB/s) keys baseline diffing consumes and adds latency_us, gflops, pct_peak_bw, pct_peak_flops, and arith_intensity.

iron build — compile kernels to MSL

Compiles every kernel and reports errors; with --emit, writes artifacts.

iron build [-f <substr>] [--dtypes f32,f16,bf16] [-v]
           [--emit msl,metallib,swift,ir,all] [-o <dir>] [--sdk <sdk>] [-t]

Flag	Effect
-f, --filter <substr>	only build matching kernels
--dtypes <list>	comma-separated dtypes to build (f32,f16,bf16)
-v	print the generated MSL for each kernel
--emit <list>	emit artifacts — msl, metallib, swift, ir, or all
-o, --out <dir>	output directory (required when --emit is set)
--sdk <sdk>	xcrun SDK for the Metal toolchain (default: macosx)
-t, --time-passes	run the pass pipeline 25× per kernel, print per-pass median wall time instead of emitting

Codegen smoke check — emit everything and confirm xcrun metal accepts it: iron build --emit all -o /tmp/iron-smoke.

Metal AIR intermediates are content-addressed under $CARGO_TARGET_DIR/tile-build-air. Successful links prune inactive artifacts older than 30 days and abandoned temporary files older than 24 hours. Set IRON_AIR_CACHE_MAX_AGE_DAYS to change the grace period, IRON_DISABLE_AIR_CACHE_PRUNE=1 to retain every artifact, or IRON_DISABLE_AIR_CACHE=1 to force a clean compiler comparison.

The output layout matches a SwiftPM Sources/<Target>/ convention so --out can point directly at a target directory:

<out>/Resources/kernels/<name>.metal
<out>/Resources/kernels.metallib
<out>/Resources/manifest.json
<out>/Generated/IronKernels.swift

The generated Swift surface

IronKernels.swift is one public enum IronKernels whose static functions mirror the kernel ABI (buffer slots, constexpr bytes, the Elementwise _n_elems tail guard) exactly as the emitted MSL declares it. Each kernel gets, per dispatch shape, a pair of overloads that share one body:

Function	Encoder	Dispatch
<name>(…, gridSize:, threadgroupSize:, encoder:)	caller’s open MTLComputeCommandEncoder — sets the PSO, binds, dispatches; never creates or ends one	dispatchThreads (grid in threads)
<name>(…, gridSize:, threadgroupSize:, on:)	opens one on the MTLCommandBuffer, forwards to the encoder: overload, ends it	same
<name>_threadgroups(…, encoder:) / (…, on:)	same pair	dispatchThreadgroups (grid in threadgroups) — required for coop_tile / simdgroup-matrix kernels
<name>_indirect(…, encoder:) / (…, on:)	same pair; opt-in per kernel (Kernel::wants_indirect_variant)	grid from an MTLBuffer
<name>_record(…, into:) + <name>_params_size	records into an MTLIndirectComputeCommand (no encoder exists, so no encoder: form)	concurrentDispatchThreads

The encoder: overload is the only place the ABI is spelled out; the on: overload is a forwarder. To batch N dispatches of one kernel on a single encoder, open the encoder once and call the encoder: overload N times:

guard let enc = cmd.makeComputeCommandEncoder() else { return }
for (src, cache) in pairs {
    IronKernels.iron_kv_cache_update_bf16(
        src: src.buffer, srcOffset: src.offset,
        cache: cache.buffer, cacheOffset: cache.offset,
        head_dim: headDim, max_seq: maxSeq, position: pos,
        gridSize: grid, threadgroupSize: tg, encoder: enc)
}
enc.endEncoding()

Kernels emitted for several dtypes (<family>_f32 / _f16 / _bf16 …) also get a runtime selector so the variant is picked without string concatenation: IronKernels.DType (raw value = the name suffix), and per family <family>_name(dtype:) -> String plus <family>(dtype:, …, encoder:) / (…, on:) / <family>_threadgroups(dtype:, …) that forward to the selected variant — IronKernels.iron_swiglu(dtype: .bf16, …, encoder: enc). A dtype the family was not emitted for traps with the available list. A family whose variants take different Swift argument types (a constexpr typed on the kernel’s T) gets only the _name selector.

Families expanded from a #[kernel(variants(...))] axis set (iron_aura_flash_sdpa_kb{KB}_vb{VB}_d{DIM}, iron_aura_flash_pass2_d{DIM}, iron_sdpa_prefill_{NAME} …) also get one typed selector per axis set, so a caller never picks among the emitted rows by hand:

Function	Selects on
<family>_name(kb:, vb:, d:, dtype:) -> String	the emitted entry-point name
<family>(kb:, vb:, d:, dtype:, …, gridSize:, threadgroupSize:, encoder:) / (…, on:)	dispatchThreads — the on: form opens an encoder and forwards to the encoder: form
<family>_threadgroups(kb:, …, encoder:) / (…, on:)	dispatchThreadgroups

IronKernels.iron_aura_flash_sdpa(
    kb: 8, vb: 4, d: 128, dtype: .bf16,
    q_rot: q.buffer, /* … the kernel ABI … */
    gridSize: grid, threadgroupSize: tg, encoder: enc)
IronKernels.iron_aura_flash_pass2(d: 128, dtype: .bf16, …, on: cmd)
IronKernels.iron_sdpa_prefill_name(name: "d256_sink", dtype: .bf16)

The selector is a switch over the emitted rows; a combination that was not emitted traps (<family>_unavailable) with the list of emitted rows, exactly like an unknown dtype. A family is one #[kernel(variants(...))] — the base function name and its suffix template — and its selector is named by the base name plus the template’s literal-only segments, so the name and the arguments together spell the entry point: iron_moe_gather_qmm_mma emitted with b{BITS} and with int{BITS}_bm16_mpp gets iron_moe_gather_qmm_mma(b:, …) and iron_moe_gather_qmm_mma_bm16_mpp(int:, …); the block-scaled fn iron kernels with {FMT}_qmm_mma get iron_qmm_mma(fmt:, …). The rules for the parameters:

• Only the axes the suffix template references — a parameter is an argument when the template names it as a bare {NAME} placeholder. That is what separates the values a caller chooses (KB, VB, DIM) from the literal columns derived from them (KL, VL, DPL), which the template never mentions; an arithmetic placeholder ({DIM * 32}) selects nothing. Without a template the auto-suffix spells every parameter out, so every parameter is an argument.
• The label is the literal in front of the placeholder within its _-separated segment, lower-cased — kb{KB} → kb:, d{DIM} → d:, int{BITS} → int: — so a call reads like the entry point it selects (kb: 8, vb: 4, d: 128 ↔ kb8_vb4_d128). A bare {NAME} falls back to the lower-cased parameter name (name:), as do all labels when two would collide, and so does a label that would reuse one of the kernel’s own arguments (iron_ssm_step_record binds a buffer d, so its d{DH} axis is dh:; <name>_axis if that is taken too).
• Integer axes are Int; named axes are String ({NAME} → name: "d256_sink", {FMT} → fmt: "mxfp4"). An integer axis is already an open value checked by the switch at run time, so a named axis as a plain String keeps one rule — every axis is the value spelled in the kernel name, and an unemitted combination traps — where a per-family Swift enum would add a generated type per axis for no extra safety.
• dtype: is present only when every row carries a dtype suffix; a family whose rows disagree on their Swift signature (an only_when constexpr on some rows) keeps _name and the trap only, with a comment. Two templates that would yield the same selector name and labels are both omitted with a comment rather than emitted as one signature twice.
• The per-row dtype selectors (iron_aura_flash_sdpa_kb8_vb4_d128(dtype:)) stay; the plain per-kernel wrappers and their ABI are unchanged.

PSOs come from the hand-written PSOCache.shared in the consuming package (IronKernelsSwift in Butter); every overload resolves the PSO through it.

Capability sets

Alongside the dispatch wrappers, IronKernels exposes two static Set<String> properties, sorted and derived straight from the manifest v2 capability metadata (KernelManifest::requires / live_compile — see docs/specs/TENSOROPS_PRIMITIVE.md §4.9):

public static let liveCompile: Set<String> = [ /* kernel names, sorted */ ]
public static let requiresCooperativeTensors: Set<String> = [ /* … */ ]

liveCompile is the set of kernels whose MSL must be compiled at run time from Resources/kernels/<name>.metal rather than loaded from kernels.metallib — the MetalPerformancePrimitives header-skew hazard §4.9 describes: an offline-compiled metallib whose baked-in cooperative-tensor types disagree with the runtime device’s produces bit-deterministic wrong output. It defaults to every kernel Kernel::requires_cooperative_tensors() flags, and drops a kernel only when #[kernel(live_compile = false)] records an evidenced exception (a passed qualification probe for that SDK/runtime pair) — never as an opt-in. A name absent from liveCompile is not a claim of safety; it may simply be a kernel iron build --emit did not emit (§5.6’s not-yet-migrated hand-written MPP sources). requiresCooperativeTensors is the plain source-requirement set — every kernel whose body lowers mpp::tensor_ops::matmul2d — and stays populated even for a kernel that opts out of live-compiling, since the opt-out changes compilation policy, not what the kernel’s MSL actually contains.

iron inspect — IR and MSL for one kernel

iron inspect [<kernel>] [--filter <substr>] [--all] [--ir] [--stats]
             [--pass <name>] [--dtype <f32|f16|bf16|i32|u32>] [-o <dir>]

Flag	Effect
(no flag)	print the final generated MSL
--ir	print the raw IR before any passes
--pass <name>	print the IR after a specific pass (--pass all for every stage)
--stats	print the per-pass op-count reduction table
--dtype <d>	dtype override for monomorphisation
--filter <substr> / --all	inspect many kernels at once
-o, --dir <dir>	write output files instead of printing to stdout

Omit the kernel name to list every registered kernel. See Developing → debugging a kernel.

iron device — GPU info

Prints the Metal device name, Metal version, Apple GPU family, and the supported feature flags (native bfloat, simdgroup matrix, etc.). Add --json for machine-readable output.

iron snap — save a perf regression baseline

iron snap [-o <file>] [--from <file.json>] [--note <text>] [-f <substr>]

Flag	Effect
-o, --out <file>	write the snapshot here (default: .iron-snapshots/<sha>.json)
--from <file.json>	promote an existing bench JSON instead of re-running the bench
--note <text>	attach a note to the snapshot
-f, --filter <substr>	only include kernels whose name contains <substr>

iron diff — compare against a baseline

iron diff <baseline> [<current>] [-f <substr>] [--threshold <pct>]
          [--sort name|delta|pct] [--only-regressions] [--only-improvements]

<baseline> is a saved snapshot JSON; <current> is an optional bench JSON — omit it and diff runs the bench itself.

Flag	Effect
-f, --filter <substr>	only show kernels whose name contains <substr>
--threshold <pct>	highlight regressions larger than this percentage (default: 5)
--sort <key>	sort rows by name, delta, or pct (default: name)
--only-regressions	show only regressed kernels
--only-improvements	show only improved kernels