15 — Compilers and kernel DSLs: what actually generates the kernel

Every engine in 03-engine-selection.md is, underneath, a choice of who writes the kernels: a hand-written CUDA/HIP library, a template library, a DSL, or a graph compiler. That choice decides three things this repo keeps rediscovering as gotchas — which hardware is reachable at all, where the compile cost lands (build time vs. first request), and whether an artifact is portable. The 22 compiler records answer those three questions; this doc reads them together.

0. The decision, short form

1. The four families

Kernel DSLs — write the kernel, the compiler schedules it

Template libraries — hand-tuned building blocks, no compiler

Graph compilers — take a model, emit fused kernels

Infrastructure and runtimes

2. Where the compile cost lands

The strategy column is the cost model. A JIT stack pays at first request: Triton compiles each kernel on first call and autotuning compiles every candidate config (Triton); NVRTC moves compilation onto process start (CUDA NVRTC (runtime compilation library)); torch.compile recompiles on new shapes (torch.compile / TorchInductor); Ascend's GraphEngine generates code at serve time, so cold start is spent there rather than at build (CANN / Ascend GraphEngine (GE)). An AOT stack pays at build time and ships an artifact: TensorRT (TensorRT), XLA (XLA), IREE (IREE), TVM — whose AOT compilation with autotuning the record puts at minutes or more (Apache TVM). Template instantiation pays at the library user's compile time (CUTLASS / CuTe, AMD Composable Kernel (CK)).

The cold-start consequences are in 21-deployment-topology.md; the measured warm-up traps are in 05-known-traps.md.

3. Portability is the exception, not the default

• TensorRT engines are not portable across GPU architectures without hardware-compatibility mode (TensorRT).
• Cerebras CSL binaries must match the WSE generation they run on (Cerebras CSL compiler (cslc) and the CSL kernel language).
• nvFuser wheels are bound to a specific PyTorch build (nvFuser (MosaicML fusion compiler, now under NVIDIA)).
• Modular MAX/Mojo's cross-vendor portability claim is vendor-published only; no third-party cross-vendor benchmark backs it in this corpus (MAX and Mojo (Modular)).
• Only four records list both cuda and rocm as targets for kernel generation: Triton, TileLang (Tile Language), ThunderKittens (and the HipKittens AMD port) (via its HipKittens port) and torch.compile / TorchInductor — plus the graph-level Apache TVM, IREE, MLIR and ONNX Runtime.

3a. Where a hand-written kernel still wins on Tenstorrent

TT-Metal's own programming model removes the usual reason to fuse. A tensor kept in per-core L1 is readable by the next op without a round trip to DRAM, so a chain of ops is already fused as long as every input and output lives in L1 (FlashAttention on Tenstorrent's Wormhole Architecture). The fused kernel is only worth writing when the algorithm itself changes the work. Attention is that case: the scores matrix must not materialize, and causal attention must not compute the upper triangle.

The Wormhole FlashAttention report measures that fused op at 20x the stack's own unfused softmax attention on average, and between 9x and 44x, at a fixed 16k tokens (TT-Metal Flash-Decode: split the KV cache across Tensix cores, reduce over the NoC). The same report shows the op is not purely bandwidth-bound: BFP8, at about half the bytes of BF16, is only moderately faster, not twice as fast. Decode is the bandwidth-bound half, and the kernel that ships for it splits the KV sequence across cores and reduces the partial softmax over the on-chip network — but only for multi-query attention on a single token (TT-Metal's decode attention kernel only implements multi-query attention). A grouped-query model, which is most models shipped since Llama-2, does not run on it. The report names the missing piece: multicast one K and V head to the cores holding the query heads of that group.

4. What the records cannot answer

• No compile-time benchmark records. compile_time_seconds is prose-level in most records; there is no measured, methodology-complete cold-start comparison across stacks.
• Two records are draft: MAX and Mojo (Modular) and Qualcomm AI Engine Direct (QNN) SDK (QNN docs sit behind a sign-in wall, so its mechanics are unverified from a primary source).
• Schema gap. target_backends has no slot for Ascend, Tenstorrent, Cerebras or Hexagon, so those records carry an empty list rather than a mislabelled one (TT-Forge (Tenstorrent's MLIR-based compiler; the TT-Bridge / graph-frontend layer)).
• No kernel-level performance join. Which compiler produced the kernel behind a given benchmark row is not recorded on the benchmark.

Related documents

• 03-engine-selection.md — the engines these compilers sit under.
• 02-flop-map.md — the operation classes the kernels implement.
• 13-model-compatibility.md — which architecture needs which kernel.
• 21-deployment-topology.md — what a cold start costs.