metakernel

An agent iterating on CUDA candidates is good at writing code and bad at trusting its own measurements. It needs numbers it cannot reliably improvise: timing that rejects contaminated samples, comparison that separates real bugs from reduction-order noise, and device facts measured rather than copied from a spec sheet. metakernel packages those as a loadable kernel, so they run anywhere get_kernel runs, including headless CI where Nsight and CUPTI are unavailable.

Usage

from kernels import get_kernel
mk = get_kernel("phanerozoic/metakernel", version=1, trust_remote_code=True)

verdict = mk.judge(cand_op, ref_op, inputs=(x,))     # compare + fuzz + bench
dossier = mk.probe_device()                          # measured facts, disk-cached
report  = mk.bench(lambda: cand.op(x), iters=200,
                   bytes=B, flops=F, dossier=dossier)
verdict = mk.compare(ref_op, cand_op, gen=make_inputs)
table   = mk.sweep(factory, {"BLOCK": [64, 128, 256]})

Every op returns plain data with a fixed schema. Reports carry a one-sentence verdict, failures carry hints, and mk.schema(op) returns field layouts at runtime.

API

Symbol	Purpose
`judge(cand, ref, inputs \| gen)`	compare, a quick fuzz, and a roofline bench in one pass/fail verdict
`probe_device(quick, cached, tenancy_check)`	measured dossier: bandwidths, latencies, bank-conflict and ILP sweeps, MMA rates with lane calibration, fp8 accumulate precision, fed-from-memory MMA, atomics, raw and torch launch overhead, occupancy, clock under load; disk-cached per device and driver; raises `device_busy` on a contended device
`probe_all_devices()`	per-device dossiers plus a P2P bandwidth matrix
`bench(fn, iters, bytes, flops, dossier, graph)`	cudaEvent medians with CI; rejection by SM clock for compute-bound work, by rate consistency for memory-bound; roofline placement; graph capture; allocation-growth and hidden-sync detection
`compare(ref, cand, gen \| inputs, order_shuffle)`	ULP evidence with flat and multi-dim indices, pointwise reduction-order bands, a one-sided-bias gate, corrective argument errors
`order_shuffles.last_dim()` / `.matmul_k()`	shuffles for the common contractual-noise axes
`fuzz_shapes(op, ref, make=None, isolate)`	edge shapes, layout mutations, 2^31 expansion, per-case repro recipes; pooled subprocess isolation
`sweep(factory, grid, check, budget_s, race)`	config search with a correctness gate; pruned points visible
`probe_capturable(fn)`	graph capturability tested in a child process
`alloc_stamps` / `read_stamps` / `stamped_demo`	clock64 phase stamps; `MK_STAMP_DRAIN` includes the store drain
`dossier_drift` / `pinned_dossiers`	fresh measurements against the release pin per architecture
`mma_tile_check(trials)`	fragment-path tiles against the reference GEMM, compared exactly
`measure_clock()` / `ulp_diff(a, b)` / `schema(op)`	primitives
`ops.mk_*`	the raw probe kernels

Method

Each dossier entry comes from a dedicated probe kernel: streaming triads and gather kernels for bandwidth, Sattolo-cycle pointer chases (minimum of three runs) for latency, register-resident FMA and mma_sync loops for compute, a shared-memory stride sweep for bank conflicts, a 1/2/4-chain ladder for latency-boundness, an atomic hammer, launch trains on the torch and raw C++ paths, and a cooperative kernel for grid-barrier cost. MMA flops derive from a sub-saturation accumulator read-back with a linearity check, never the declared loop shape; the fp8 accumulate width is read from the point where accumulator growth stalls. Every number carries the SM clock it ran under. Comparison treats the reference as a distribution: an order shuffle bounds the legitimate band pointwise, and the bias gate fails errors that sit above their band in one direction.

Measured

Release-pinned dossier for RTX 6000 Ada (sm89, 142 SMs, 96 MB L2, ECC on):

entry	value
clock under sustained FMA load	2.71 GHz max, 0.3% sag over 24 windows
device-memory bandwidth	triad 655-682 GB/s; read-only 767; write-only 767; random 16 B gather 218
L2 / shared-memory bandwidth	9,591 GB/s / 22,244 GB/s aggregate
shared-memory conflict penalty	1.0x at stride 1 to 14.1x at stride 32
host transfers	H2D 24.0 pinned; D2H 26.3; D2D 324 GB/s
dependent-load latency	smem 30.0 cy; L1 45.8; L2 292; DRAM 657
FMA pipes	f32 86.6 TFLOP/s; f64 1.26; ILP ladder 11.8/23.4/43.0 at 1/2/4 chains
MMA sustained	fp16 365, bf16 370, tf32 178, int8 750 TOP/s; 2:4-sparse 356 executed / 711 dense-equivalent; fp8 e4m3 772, e5m2 771
fp8 accumulate precision	8 effective mantissa bits (growth stalls at increment x 2^9)
fed-from-memory MMA	85 TFLOP/s through cp.async-staged smem tiles
atomics fp32 curve	0.71 Gop/s at 1 address to 414 at one per thread
launch overhead	raw 6.8 us queued / 13.9 round trip; torch 7.1 / 18.5
grid barrier	0.55 us per grid.sync

A second pinned dossier covers GeForce RTX 3070 Ti Laptop (sm86).

Certification

The suite asserts invariants that must hold on any correctly functioning device: bandwidth and latency hierarchies, tensor-core rates above the FMA pipe at the silicon's structural ratios with full verified K-lanes, a cuBLAS GEMM within [0.30, 1.05] of the dossier fp16 rate, the conflict cliff and ILP scaling, comparison passing reordered reductions while failing corrupted and uniformly biased candidates, and fuzzing that reports every case including the 2^31 expansion.

Limits

No hardware counters; entries are inferred from timing and the occupancy API and carry confidence fields.
One tenant, one device per dossier. A contended device is refused, not corrected.
The banded comparison cannot reject errors under band_margin * (band + 2) ULP at order-stable elements.
Stamps record instruction issue. Hidden-sync detection sees nothing under ~50 us per call.
MMA rates are register-resident ceilings; mma_fed is one staging pattern, not a frontier.
In-process fuzzing dies with a killed context; isolation costs a worker respawn per crash.
Published variants are Linux x86_64, torch 2.12+. Windows JIT-builds from source via load_local.py.