Don't Write Kernels, Derive Them

1 Introduction

A modern model has a few dozen distinct operators. A modern framework has thousands of kernels. Every (operator, backend) pair needs one. So does every fusion that pays, such as attention with its softmax, a matmul with its bias and activation, or a norm with the projection after it. FlashAttention ¹ is the best-known case. Its large speedup came from rewriting one operator by hand so its score matrix never reaches memory. Its backward pass took a second hand-written kernel. Hopper took a third ², and Blackwell’s FP4 tensor cores are taking a fourth.

Each hand kernel is a separate copy of the same arithmetic, and each one freezes a schedule that was right for one shape on one device. We measured how often that schedule is wrong: the fastest schedule changes with the size, the chip and even between runs on the same machine (§7.1). The winning schedule is a property of the graph, the chip and the size, not of the operator.

Our thesis is that a kernel is a theorem, not a program. Every fast kernel we know of is a claim of the form “this comprehension, reassociated under this monoid and placed in this memory, equals the naive one over the reals.” Such a claim can be derived, checked and searched. It does not need to be written. It can also be bounded: the least traffic any box-tiled schedule of a comprehension can move, and the tile that moves it, are dual solutions of one linear program. A theorem fixes what is computed. How fast it runs is still a schedule, and §11 lists the schedules tang does not yet have. If that is right, the kernel library should be replaced by an algebra small enough to reason about, together with a search over its schedules. Figure 1 shows one graph under the schedules that the field knows by name.

The traffic column is not the time column. Checkpointing moves a third of the naive schedule’s traffic and is 2.2× slower. The score contraction accounts for more than the whole difference: 7.12 ms against 1.51 ms, with the same arithmetic and the same output. The naive schedule writes in its own small transpose kernel, so a warp’s 32 lanes then read 32 adjacent floats. Checkpointing folds the transpose into the reader’s view, so each lane reads its own row of . Nsight shows the same 67.6 M L1 requests carrying 5.4 GB in one case and 35.5 GB in the other. The output is exact to , so the lowering is correct but loses coalescing, and a traffic model that counts DRAM bytes once per buffer cannot see it. Flash, which also has the least traffic, wins.

The backward tells a less flattering story. tang’s fused attention backward kernel computes in bf16 tensor cores even when the graph is f32. On this graph its gradients miss the experiment’s tolerance of , measured as max error over max reference against the f64 interpreter (error in ), so the search rejects it. Without the fused kernel the searched placement passes at in 1.91 ms, against 0.74 ms for the failing fused version. The search also rejected one placement whose output was simply wrong, at error 7.0; that bug is caught by the check but not yet explained.

We propose writing the model once, as math, and letting the compiler derive every kernel. That works only if the IR is small enough to reason about and expressive enough that nothing important falls outside it. tang is a no_std Rust library that started as a geometry and physics math crate. Its tang-expr crate traces Rust closures into scalar expression graphs that it can differentiate, simplify and compile, and its tensor IR uses those graphs as its body language. The IR has two constructs.

• Views: affine index maps that express broadcast, transpose, slicing, striding, padding and per-head column selection without touching memory.
• Comprehensions: , where body is a scalar expression graph and is a monoid.

Contributions:

1. Kernels from rewrites. Fusion, flash attention, logsumexp and Welford variance are instances of a small set of rewrites (§2).
2. Derived backward passes. Reverse-mode differentiation maps comprehensions to comprehensions, including the online-softmax monoid. A proof procedure removes the stabilizing max’s gradient, and with nested contractions this yields an attention backward with no tensors (§3).
3. Precision as a schedule. An emulating interpreter checks a format choice against an error budget before it runs, with failure modes we report (§4).
4. One graph, six targets. Lowerings from shared emitters, measured against hand kernels, PyTorch and the bandwidth limit on four chips (§5, §7).
5. A production engine, derived. Top- routing as one monoid, integer tensor-core and narrow GEMV schedules, a segmented-sum gather backward, chunked scans, and Hopper attention schedules. With them, derived kernels replace hand kernels in an 80B MoE engine on an H200, faster in some regimes and slower in others, which we report (§8, §11). A whole derived decode step matches the hand engine on Qwen3-8B and trails it by 13–35% on the 80B MoE.
6. The speed of light. The traffic lower bound and the optimal tile are dual solutions of one linear program read off a kernel’s views, with weak duality in Lean. A time floor built on it never exceeds measured time on 24 GEMV-class and epilogue kernels on two GPUs; it is close on the large kernels of the RTX 3090 and loose on the H200 (§9).

2 The IR

2.1 Views

A view maps a point of an iteration space to a source tensor’s coordinates with one affine form per source dimension, . Broadcast is a zero form, transpose a permutation, slice an offset, stride a coefficient. A form with two terms expresses “head reads columns “. Views compose by substitution. They are the only way a comprehension reads its inputs, so every layout decision is a view and every view can be analyzed.

A view may be guarded, so that an out-of-range read returns the monoid’s identity (zero for sum, for max) rather than raising a shape error. Padded, dilated and transposed (polyphase) convolutions and overlap-add are guarded reads, with no padding operator. The adjoint of a guarded read is itself a guarded read, not a scatter. For , reads through the same guard, so out-of-range (padding) terms never reach .

Grouped-query attention’s is not affine. The IR expresses heads as two dimensions, , and the views do the rest. Dimensions can be fixed or symbolic (Dim::Sym("seq")). Construction, rewriting and differentiation work with symbols, and bind makes them concrete for scheduling and lowering.

A view may also gather. One source coordinate is then data, read from an index tensor that is a sibling argument of the same comprehension, . MoE routing (), paged KV caches () and embedding lookups become ordinary reads rather than special operators. The gradient of a gathered read is a scatter-add, written as a dense gated sum, which costs times the needed work unless the keys are declared sorted (§8). An affine map read through a gather still fuses into it, and only nested gathers materialize.

2.2 Comprehensions and monoids

The only compute operation is

A map is a comprehension with no reduce dimensions. Sum and max are monoids. Matmul is . Attention scores, convolution and gather are the same operation with different views. The monoid’s combine is itself a scalar graph, and a comprehension may carry tuple-valued state. Online softmax is the monoid on , the running max and the rescaled sum, and attention carries the weighted values with it as ; arg-max is the monoid on .

2.3 Rewrites

Fusion is one rule. A map is a pure function of its inputs, so substitute_maps splices its body into every consumer and composes the views. Map→map, map→reduce and transpose→anything are the same case. The rule is cost-aware, so a map is recomputed inside a reduction only when it is cheap. Reductions are never substituted. They are inlined by a schedule decision, which separates algebra from scheduling. A row-major reshape that only merges dimensions is a view, not a copy.

Epilogues run on the reducing thread. A map that reads a reduction point for point (a bias, an activation, the that finishes softmax) runs on the thread that holds the final accumulator.

Sibling reductions merge. Two reductions over the same space and monoid with the same readers become one reduction under the product monoid (merge_siblings). The gate and up projections of a SwiGLU MLP become one tuple-valued matmul, the activation becomes its epilogue, and the normalized input is read once.

Flash attention is a rewrite. stabilize_softmax finds and every over the same canonical , with masks and scales allowed, and replaces them with one reduction under the online-softmax monoid. The value column of becomes a point dimension (an output dimension, not reduced over) of the fused comprehension. Nothing is materialized. Logsumexp, Welford variance and top-k are the same rewrite with different monoids. Figure 2 shows the attention case.

The rewrite is not specific to plain attention. Anything that changes the score or the mask leaves the monoid structure alone, so it derives the same way. The lab below runs tang’s compiler, built to WebAssembly from the same crates as this paper, in your browser. Type a variant: it builds the naive graph, runs the real pass pipeline, then runs both programs on your GPU through WebGPU and checks them against tang’s f64 interpreter.

2.4 Rediscoveries

The rewrites above are general, so known algorithms fall out of them. Table 1 lists the ones tang derives today. Each row was a paper or a hand kernel. In tang each is a monoid, a view or a schedule field, and none adds an operator.

2.5 Schedules

A comprehension says what to compute. Schedule fields say how, and they cannot change the math (precision changes rounding, §4):

• inline: recompute inside the consumer instead of writing a buffer.
• split: one workgroup per point, with threads striding the reduce space and tree-combining. This is valid for any monoid because combine is associative, including online softmax.
• tile, tile_2d: blocks of threads, each owning a register tile of points. The emitter tracks every address as an affine form over loop variables. Anything that does not depend on the lane axis is staged in workgroup memory once per reduce chunk. That covers buffer reads (K and V tiles, matmul operand tiles) and inlined sub-computations (the scores).
• precision: storage and accumulation formats (§4).

schedule_auto encodes rules of thumb distilled from measurement. autotune_kernels times each candidate mapping per materialized node on the device, using the graph’s real inputs, and caches the winner by device and kernel source. With substitution and the online-softmax rewrite, the staging rule produces the structure of flash attention: rows × lanes per workgroup, cooperative score tiles, and online softmax in registers.

3 Gradients are comprehensions

For , the gradient with respect to through a pick view (one whose forms are each a single coordinate) is found by renaming coordinates. The coordinates that picks become ‘s output dimensions, the rest become reduce dimensions, and every other view is pushed through the same renaming. The body is differentiated by tang-expr’s symbolic diff, which covers its nine arithmetic operations (add, multiply, negate, reciprocal, square root, sine, atan2, , ) plus a select. A dimension read through any other form (an offset, stride, diagonal or broadcast) stays an output dimension, and iota, a node that turns coordinates into values, gates the body by . Max gates its term to the arg-max. The backward graph is an ordinary graph, so the same passes optimize it. Rotary embedding and GQA (grouped-query attention) gradients come from the gated scatter with no special handling.

The online-softmax monoid has its own rule. With the forward’s and accumulators read back, , where is key ‘s value, so the flash-fused graph differentiates and its gradients equal the naive graph’s.

A stabilizing max has no gradient. That the max can be stopped is standard practice in logsumexp implementations. What is new here is a proof the compiler runs per graph. Softmax is invariant to shifting its scores, so the terms sum to zero, but only in exact arithmetic. While those terms exist, the max must be materialized and masked tiles cannot be skipped. backward drops them when it can prove the loss does not depend on . The proof tracks how each node downstream of scales with . In softmax, has degree 1 in , so does , and so does ; their ratio has degree 0, a function of the scores alone. In general a node is either or with free of , combined by exact identities.¹ If every path from ends at a degree-0 node, and the gradient is not built. The f64 evaluation of a training step is unchanged (largest absolute difference ).

The flash backward recomputes inside tiles (Figure 3). hoist_invariant_factors factors row-invariant terms out of sums (). schedule_nested_contractions then inlines a short contraction (one reduce dimension of at most 128) into every tile that reads it. The attention backward reads and only through and , which the dQ, dK and dV tiles recompute per staged block on tensor cores. On CUDA this removed S, , dP, the max over S and three correction kernels from a transformer training step, reducing 63 kernels to 49 and the summed best-per-kernel time from 1.80 to 1.22 ms at sequence length 512 (RTX 3090). WGSL and Metal now nest the same contractions with scalar dot products in workgroup memory, so no GPU backend we measured (CUDA, WGSL, Metal) writes anything . For now this costs speed. Forward plus backward runs at 0.68–0.93× the materialized schedule on Metal and 0.25–0.34× under Vulkan on a 3090, until the nested dots are register-blocked the way the mma path blocks them.

4 Precision is a schedule

f16 is a different way to run the same program, not a different program. Each node has a storage format, applied where a value is written to memory, and an accumulation format, applied to a reduction’s running state after each combine. The graph stays exact math.

The interpreter gives an error budget, with known failure modes. eval is pure f64. eval_emulated rounds every stored value and every accumulator step to its scheduled format (round to nearest even, subnormals, overflow), with body arithmetic in f64. An ErrorBudget accepts or rejects a schedule before it runs. Emulation matches the GPU’s f16 kernels to 1 ulp when the reduction order is the same. Otherwise the order is part of the error: AdaIN’s variance over with an f16 accumulator has error under a 256-thread split and with one thread per row, while the sequential emulation predicts . Emulation models one sequential fold; a tree-combining split beats it, and the per-row GPU kernel, though also sequential, comes out somewhat worse, which we attribute to its different accumulation order and fused arithmetic rather than having verified it. The prediction can be far off when the reduction length or order differs: a conv1d sum of 3,584 terms measures against a predicted , a ConvNeXt LayerNorm statistic against , and a pointwise matmul against . Those predictions were made on smaller instances, whose shorter sums round less. At equal reduction length (matmul, AdaIN) prediction and measurement agree within about 1.7×. tune_precision lowers buffers to half precision one move at a time and keeps a move only if the program stays within budget and gets faster.

5 Lowering

Shared emitters produce the point, split and tiled schedules, register tiles, guarded views, epilogue fusion and tuple monoids once. A dialect switch then handles buffer access and surface syntax.

• WGSL (portable default, through wgpu on Metal, Vulkan and DX12). f16 storage is packed two per u32, because wgpu’s shader compiler (naga) rejected enable f16 on our targets.
• CUDA. Matmul-shaped tiles run on mma.sync (TF32, the tensor cores’ 19-bit float mode, m16n8k8, or f16/bf16 m16n8k16 when the operands are half precision) with cp.async staging rings. Block shapes are tuning candidates. cuda_flash_kernel recognizes the online-softmax comprehension and keeps , the row max and the partial sums in mma registers. The score accumulator is reused as the A operand of after permuting keys within each 8-wide step, which is legal because the reduction is commutative and associative.
• Metal (MSL). Native half and bfloat. Contractions with f16 operands run on simdgroup_half8x8 fragments. Reduce dimensions are flattened, so convolution becomes implicit GEMM, with dilated inputs staged once per chunk as a zero-padded, phase-de-interleaved window.
• OpenCL and GLSL for Adreno GPUs that have neither CUDA nor Metal. On one Adreno driver, WGSL is routed through GLSL because the direct SPIR-V path was miscompiled.
• C, for batch-1 control loops where a GPU round trip costs more than the arithmetic. A comprehension already is a loop nest, so to_c emits it directly with a vectorizable innermost point dimension, and tang_run() performs no allocation.

Program::tuned times each kernel’s native lowering against the portable one and keeps the faster. Kernel sources are canonicalized, so the same kernel at another layer is the same string and compiles once.

6 Verification

Speed claims depend on every kernel being correct. tang-verify runs each GPU kernel against its CPU oracle on the same input, starting from the device’s own state. Before each kernel runs, it fills the kernel’s outputs with NaN, so unwritten memory cannot pass as a plausible 0. When a check fails, it names the kernel at fault by provenance: which earlier kernel’s correct output the failing buffer contains. We built it after four structural bugs in a hand-written engine slipped past every unit test. One was a byte offset added where a float offset was meant, which wrote k’s values over q. Derived kernels are less exposed to that class of bug, because addresses come from affine forms rather than pointer arithmetic. On the IR path there is a second check, because the f64 interpreter runs the same graph as the GPU.

The rewrites themselves are mechanized in Lean 4 with Mathlib, with no sorry (proofs/, listed in docs/proofs.md). Over the reals, view composition, fusion, split and tiled reductions, online softmax as a commutative semigroup (the proof has no identity element, and the compiler supplies ), its equality with two-pass softmax, shift invariance, gradients of sum comprehensions, the zero gradient of a stabilizing max, and the flash backward identity are theorems. One combine exported from the compiler is proved equal to the verified online-softmax combine. Floating point is not modelled, and guarded and gathered views are outside the fusion theorem.

7 Evaluation

Setup. Apple M4 Max (Metal), NVIDIA RTX 3090 (Vulkan and CUDA 12), a Qualcomm QCS8550 in a Booster K1 humanoid robot (Adreno 740, Cortex-X3), and, for §8, an NVIDIA H200. PyTorch 2.10 for the shared-GPU runs and 2.8 for the idle-GPU runs, both cu128, timed with CUDA events. Unless a caption says otherwise (Table 2 includes submit and readback), tang times are kernel-only. Several runs shared the GPU with other tenants. We mark those, and report minima rather than means, because contention only adds time. Derived results are checked against an f64 interpreter, except where a caption says otherwise. The tolerance differs per experiment and is stated with it: to relative for the kernels of §7.1–7.2, – for bf16 attention at , and arg-max agreement at 257 positions for the Qwen3-8B step.

7.1 Derived vs hand-written kernels

tang shipped hand-written WGSL kernels before the IR existed, which makes a controlled comparison possible: same runtime, same chip, same inputs. Here “as written” is the unfused graph, one kernel per comprehension; “derived” is any schedule the compiler produces, and the autotuner keeps the as-written graph’s naive schedule as one candidate, which is why it can win. On the M4 Max the fastest schedule for causal attention flipped between two runs of the same code: in one run was fastest naive ( threads) and tiled, and Table 2, a later run, has the reverse. The naive graph wins when its scores fit in cache and parallelism, not bandwidth, is the bottleneck, which depends on the size and the machine’s power state. A shared-memory matmul tile wins at and loses at , and rules of thumb tuned on the M4 picked the wrong matmul schedule on an RTX 3090. Table 2 is one consistent run on the M4 Max, including about 1.8 ms of submit and readback per measurement.

Leaving aside the training step, the derived kernels win or lose by small margins on the small ops. With the 1.8 ms overhead subtracted, matmul is 1.32× faster than hand and rmsnorm 1.25× slower. The schedules that win differ across operators and sizes. On the 3090 under Vulkan, the derived transformer block first lost to the hand path. Three chip-agnostic fixes followed: cost-aware substitution, keeping reductions schedulable under their epilogues, and a bank-conflict-free tile layout. After them, the autotuned block won there too (2.5 vs 3.1 ms at seq 128, 6.2 vs 10.4 ms at seq 512, both including about 1.25 ms of submit and readback).

7.2 Against PyTorch

Table 3 tracks a 9-head GQA transformer block (, , rotary embeddings, SwiGLU) on the RTX 3090 as the backends matured. Through WGSL, which has no access to tensor cores, derived kernels stayed 2.6–10× behind PyTorch. With the CUDA lowering, flash attention on mma.sync and the backward passes of §3, the seq-512 forward block passes PyTorch (Figure 4).

On the CUDA path, the derived inference attention kernel alone is 0.032 ms at seq 512, down from 0.178 ms before the mma.sync rewrite. Table 3 compares against eager PyTorch, which launches about ten cuBLAS and flash kernels per block forward and is bound by launches at these sizes. Against a fair baseline the picture is narrower (Table 4). On an idle RTX 3090, PyTorch captured in a CUDA graph with TF32 beats tang’s forward at seq 128 (0.162 against 0.196 ms) and ties forward plus backward at seq 128 and 512. tang still wins the seq-512 forward against every PyTorch mode, 0.331 ms against 0.396 ms for the CUDA graph and 0.407 ms for torch.compile. A run on a contended GPU had tang’s tuned training step at 5.05 ms against PyTorch’s 4.21 ms, so the training-step ranking is not established.

We do not claim a general win over cuBLAS: a lone matmul on the derived CUDA TF32 tile takes 0.103 ms at the default block shape and 0.088 ms at its best one, against cuBLAS TF32′s 0.079 ms, 1.1× slower. The gain comes from fusion across operator boundaries that a library of separately tuned operators cannot do.

7.3 Attention variants against FlexAttention and FA2

We ran the six attention variants from §2.3 on the RTX 3090 against FlexAttention ⁷ (torch.compile, fp32 with TF32 and bf16) and FlashAttention-2 in bf16 where it supports the mask. FA2 here means the kernels bundled with PyTorch, run through its scaled-dot-product-attention (SDPA) FLASH_ATTENTION backend, because the flash_attn package could not be built on the test machine. Flashlight and EqiForge publish no code. Every row was checked against an f64 reference; tang is pure fp32 or TF32.

Where a variant changes the score (soft cap, ALiBi), the derived forward is up to 3.4× faster than FlexAttention in fp32 up to and beats it in bf16 up to (the low end of that range is within the noise of a shared GPU), because the score modifier becomes part of the fused body. tang loses in three places, each a missing schedule rather than missing algebra. FA2 runs bf16 tensor cores where tang’s tile is TF32. A sliding window is 25× slower because tang computes masked tiles that FlexAttention’s block mask skips. And at no tensor-core tile is chosen for the forward, a 150× gap at in Table 5. Across the full sweep the forward wins 0 of 8 against FA2 and the forward 0 of 24 against FlexAttention bf16. From the backward is 2.5–7× behind FlexAttention bf16, and 50× with a sliding window.

7.4 A whole model as one graph

Qwen3-8B with MLX 4-bit weights runs as one IR graph per decode step, lowered by to_cuda_decode into 584 kernels and captured as a CUDA graph. Nibble unpacking and the dequantization are a map the GEMV reads. That map is inlined and evaluated in registers as weights stream in. Figure 5 shows the graph.

The KV-cache update is a select over a runtime position, which the lowering turns into loop bounds, so a step touches only live cache rows. Table 6 compares it with llama.cpp on the same card (engineering notes in the appendix, Decode engineering notes), run back to back on the same weights. The model’s greedy output matches an exact f32 reference token for token.

Every change between the first build and the current one is a schedule or a lowering rule, not a kernel; the appendix lists them.

The same graph builder lowers to Metal. There, prefill and decode programs share one copy of the weights through buffer binding, prefill can start from a nonzero offset (300 + 200 tokens is bit-identical to 500), and the next token is chosen on the device by an arg-max monoid, so decoding reads back 8 bytes per token.

7.5 Data-dependent addressing

A top-k MoE FFN builds as one graph: router GEMM, top-k, softmax weights, SwiGLU experts reading through gather views, then the weighted combine. Top-k is exact, and outputs match the f64 interpreter to on Metal and on CUDA (the router runs as a TF32 tile). Table 7 has the timings.

7.6 Edge hardware

On the K1′s Adreno 740, a derived matmul reaches 322 GFLOP/s at , which is 21% of the measured fp32 peak, and a full transformer forward pass plus all seven weight gradients at seq 128 runs in 19.1 ms. Rules tuned for desktop GPUs under-fill the Adreno’s six compute units. The best schedule there (an register tile with 8-deep chunks) is 1.45× better than schedule_auto’s choice, which again argues for search over remembered rules. Through to_c on one Cortex-X3 core, a batch-1 GRU policy step (6 contractions plus gate math, fused into one kernel) takes 22.5 µs median and a locomotion MLP takes 6.7 µs, both within of f64.

8 Deriving a production engine

Every kernel in the sections above is derived. To show where derivation has to go next, we report the other half of tang: the hand-written serving engine in tang-llm, which serves Qwen3-Coder-Next (80B parameters, 3B active, 4.5-bit GGUF in Q4_K and Q5_K) on one NVIDIA H200 with kernels written by hand. Table 8′s numbers are not results of the IR; they are the target it has to reach.

The hand engine also shows why the contract matters. Ported to GLM-5.3, the same MoE kernel is correct but up to 2.5× slower than the engine’s generic path below 64 tokens per expert. A kernel tuned for one model’s expert shapes is a point in the schedule space, and on the next model it is the wrong point. It is one data point, but from a production engine, and it is §1′s claim.

Deriving this engine needed constructs the IR lacked: a gather view for routing tokens to experts, integer tensor-core fragments with fused dequantization, a top- reduction, scans, and Hopper’s asynchronous copies and warpgroup MMA as schedule fields. Each now exists, and each can replace hand kernels in the engine behind a switch, measured against them on the same bytes.

8.1 The MoE layer

One full MoE layer of the engine builds as a single graph: router, top-, expert-sorted gather, the gate|up and down expert GEMMs on GGUF Q4_K/Q5_K weights, SwiGLU and the weighted combine. The int8 path uses schedule fields only: GGUF blocks as an input storage format, int8 activations quantized by their writer, and an IntMma field that stages activations in fragment order for m16n8k32 tensor cores. There is no new operator.

Routing is one reduction. Monoid::topk(k) keeps the best (index, value) pairs, sorted, with ties to the smaller index. Its combine is a merge of two sorted lists ( compare-selects), which is associative, so split and subgroup trees give the same answer. Top-10 over 512 experts is one 0.013 ms kernel in place of arg-max rounds, and a whole layer at batch 1 drops from 32 kernels and 0.199 ms to 13 kernels and 0.103 ms.

Batch 1 needs a different schedule. Stage::Gemv computes one dot per (position, row) on dp4a, with 16-byte loads of half super-blocks and partial sums combined by shuffles. Its producers fuse: the and nodes are marked inline, and the GEMV computes and quantizes its operand row into shared memory in its prologue, rounding exactly as the writer did.

With equal work the derived layer loses at batch 1 and 4, by 18% and 7%, ties at 8 and wins by 5% at 64. The derived expert half alone is faster than the hand layer at batch 1, but only because the hand layer’s time includes the residual, norm and quantization that ours leaves out. Fragment-order staging alone is worth about 20% on gate|up at batch 64. The hand MoE kernels took a dedicated tuning effort; the derived ones are a schedule search, and on equal terms they are not yet faster at small batch.

The backward of a gathered contraction was dense: every expert’s weight gradient summed over every pair under a gate, times the needed work. Declaring the key index sorted (set_sorted) tags that reduction with a segment, and the point loop binary-searches each key’s run and sums only over it. The gate stays in the body, so a lowering that ignores the tag is still correct. The expert-weight gradient (E = 512) drops from 8.27 to 3.89 ms on the H200, 16.3 to 2.67 ms on the 3090 and 39.7 to 2.65 ms on the M4 Max, with identical error.

8.2 Attention and scans in prefill

The engine’s model alternates gated-delta (GDN) layers with full attention at head dimension 256, and prefill is where both cost. At the old schedule left the online-softmax reduction untiled, and the score stayed an buffer at 2.8 TFLOP/s. Allowing the tile up to 256 lanes lets the score inline, and the Hopper flash shape takes over with no change to the graph or the emitters (Table 10).

SDPA is not the strongest baseline at this head dimension. Against FlashAttention-3 built for head dimension 256, on an idle H200 in a later run, tang’s forward runs at 0.43–0.54 of FA3′s speed and its backward at 0.22–0.40 (Table 11). In that run tang’s forward beats SDPA from up and loses at causal, and its backward loses to SDPA everywhere (SDPA’s backward is not in the table).

Gated delta rules, Mamba-2′s SSD and linear attention are associative scans over an affine monoid on a matrix state. A delta_scan comprehension lowers by a chunked plan: a WY transform per chunk, a sequential state scan, and a parallel readout. Its backward is the derived reverse scan at chunk granularity. On the 3090 it matches fla’s Triton kernels within 6% forward and beats mamba_ssm’s SSD forward 1.3–1.8×. Against the engine’s hand recurrence (a warp per state column, token by token), the chunked scan is 4.0–4.3× faster on the H200 and 2.1× on the 3090.

Both went into the engine behind switches, gated on the engine’s own accuracy checks: KL divergence against the reference at 3k and 8k tokens with int8 and bf16 KV caches, equal greedy output, bitwise-identical batched decode, and a cached-prefix continuation equal to a cold one.

These gains are over the engine’s own prior code, and that baseline is far from the state of the art. vLLM 0.33 on the official FP8 checkpoint, on an idle H200 with DeepGEMM off, reaches first token in 0.17 s at 8k tokens and 1.08 s at 44k. The engine, on its 4.5-bit GGUF, took 2.1 s and 13.8 s in the same session, and 12.9 s with derived attention: vLLM is about 12× faster. The weight formats differ, so the comparison is not matched for precision, but a gap of an order of magnitude does not come from the format. The engine’s 13.8 s in that run is also slower than the 8.33 s in Table 12, which came from a warm server; we have not reconciled the two measurement paths, so Table 12 should be read as relative.

8.3 Attention backward on Hopper

Training needs the backward. The derived attention backward was nine kernels that round-tripped P and dS through memory, at 13–25 TFLOP/s. A lowering recognizes the pattern the gradient pass produces, a tuple reduction over queries whose operands are pointwise maps of two short contractions, and emits one kernel per key block. It is a shape-keyed emitter. dK and dV stay in registers, and dQ is added with Hopper’s bulk asynchronous reduce. Variants on mma.sync, on wgmma with two consumer warpgroups, and with a producer warpgroup are all derived from the same match, and the autotuner picks per shape. The fused kernel alone reaches 0.81–0.96× of SDPA’s backward kernel at , and the whole program is 4–8× faster than the nine kernels. The whole backward program is slower than SDPA’s: 0.43–0.74× at in the run behind Table 11, which does not tabulate SDPA’s backward. The fused kernel also computes in bf16 tensor cores even for an f32 graph, which is fine for bf16 training but fails an f32 tolerance (§1).

8.4 The derived MoE in the model

A layer that wins in isolation still has to win inside the model. Behind a switch, the derived MoE replaces the engine’s expert kernels in every layer of Qwen3-Coder-Next. It passes the engine’s accuracy checks at 3k and 8k tokens and is bitwise-equal batched at 8 streams; the 48-stream agreement gate has a 99% bar that the derived path misses at 98.93%. The hand engine fails it too in the same run, at exactly 99.00% against a strict inequality. Getting there exposed one numerical choice: a block-split sum of squares in RMSNorm, split 256 ways, pushed the KL divergence at 3k tokens to , over the bar. Split 1024 ways it is , lower than the serial sum’s . Precision is a schedule here too.

At 1–8 streams the derived path was slower, so the engine takes it only for windows of at least 16 rows and keeps the hand kernels below that. Above that cutoff it is faster. Over five interleaved rounds on one idle H200, with hand and derived alternating in fresh processes, a decode step takes 13.99 ms against 15.84 ms at 16 streams, 21.36 against 23.44 at 32 and 28.43 against 30.00 at 48 (medians). The derived step takes 5–12% less time, and the largest spread across rounds is 0.47 ms, so every round shows it. Earlier single runs disagreed because they compared different pods; hand at 48 streams has read 29.67 to 30.26 ms across four H200s. With equal work the layer alone is slower at batch 1 and 4, which matches where the cutoff sits.

The whole decode step can be derived as well. On an idle RTX PRO 6000, single stream, Qwen3-8B’s derived step takes 4.72 ms at context 128 and 5.12 ms at 4k, against 4.82 and 5.19 ms for the hand engine, and produces the same 64 greedy tokens. Qwen3-Coder-Next’s derived step, 1,924 kernels, also matches the hand engine’s 64 greedy tokens but is slower: 5.89 ms against 5.21 at context 128 (13%) and 6.93 against 5.15 at 4k (35%). The gap grows with context; we have not yet broken it down by kernel.

9 The speed of light

“Derived kernels beat hand kernels” is a claim about two programs. It says nothing about whether either is fast. The question that matters is how far a kernel is from the fastest any implementation could be. tang can answer it from the same comprehension it compiles.

9.1 The floor and the tile are dual

For a comprehension over affine views, the least memory traffic of any tiled schedule comes from a linear program. Each view projects the iteration space onto one operand, and the Brascamp–Lieb inequality bounds how many iterations a tile can run for the data it holds ⁸. The program’s optimum is the traffic lower bound. Its dual is a tile shape that attains it. So the bound and the schedule are two solutions of one problem, and the duality gap is how far a schedule is from optimal. tang builds this program from a kernel’s views (TensorGraph::kernel_lp) and solves both sides.

It reproduces the classical bounds. Matmul gets exponent with a cube tile. Attention gets Saha and Ye’s ⁹, and its dual tile is FlashAttention’s blocks. We compared the dual with the tiles tang picks (Table 13). Two different things pick tiles: schedule_auto is a fixed rule that times nothing (in this comparison it gives 64 × 64 for matmul and 16 × 16 for attention; its tile depends on the kernel’s shape and the backend), and the autotuner is a timed search. The comparison does not show that the model predicts the autotuner.