hanzoai/papers
2
0
Fork 0
mirror of https://github.com/hanzoai/papers.git synced 2026-10-22 17:20:17 +00:00
Code Issues Actions Packages Projects Releases Wiki Activity

papers: split hanzo-kernel DSL into a pedagogical Paper 07 series

Browse Source 
Restructure the 12-page monolith into one paper per concept, each self-
contained and teachable (concrete failure -> concept -> repo code -> eval for
that concept only), with hanzo-kernel-dsl.tex slimmed to a ~5-page umbrella that
cites the six rather than inlining their depth. Each fact lives in exactly one
paper (DRY). Shared preamble in shared/kernel-preamble.tex.

  07   hanzo-kernel-dsl                  umbrella: thesis + guided tour
  07.1 hanzo-fusion-is-composition       functor law, fences, fused_interp, DAG
  07.2 hanzo-comptime-autotune           schedule-as-value, byte-identical variants
  07.3 hanzo-intrinsic-islands           oracle-anchored escape hatch vs forks
  07.4 hanzo-bitexact-oracle             CPU oracle law + nightly 3-backend CI
  07.5 hanzo-int8-mmq-negative-result    opaque WMMA fragment, sync-bound proof
  07.6 hanzo-collapse-methodology        ledger + verify-then-delete + migrations

All 7 build via the Makefile (auto-discovered) with zero undefined refs.
This commit is contained in:
zeekay
2026-07-08 08:55:25 -07:00
parent 37c9e79f83
commit 54135d00ba
15 changed files with 1375 additions and 596 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hanzo-bitexact-oracle.pdf
BIN
View File
Binary file not shown.
hanzo-bitexact-oracle.tex
+182
View File
@@ -0,0 +1,182 @@
\documentclass[11pt]{article}
\input{shared/kernel-preamble.tex}
\title{\textbf{One Oracle, Every Backend:\\ The CPU Bit-Exactness Law and How It Is Enforced}}
\author{Hanzo AI --- ML Systems}
\date{hanzo-kernel Paper 07.4 --- part of the \emph{One Source, Every Backend} series}
\begin{document}
\maketitle
\begin{abstract}
The most expensive bugs in a multi-backend inference engine are not crashes; they are silent numeric
divergences, where the same model produces coherent text on one GPU and degenerate repetition on another
from byte-identical weights. This paper isolates the cause --- ``the same operation, a different kernel
per backend, therefore a different numeric behavior'' --- and states the law that makes the class
structurally impossible: \textbf{the CPU runtime is the single bit-exact reference for every backend.} It
is not an accelerator target; it is the \emph{oracle}. Every operation ships with a plain-Rust reference
whose scalar semantics the device code mirrors line-for-line; because the CPU executes a pointwise
operation in program order, its result is the \emph{definition}, and a GPU backend may differ only by
floating-point re-association where a reduction reorders a sum --- a bounded, explained difference, not a
bug to hunt. We show how the law is anchored (a mirror between host and device code), why it holds
(index-locality plus in-order execution), and how it is enforced continuously: a nightly three-backend
GPU CI on self-hosted NVIDIA, AMD, and Apple machines that runs the same source against the oracle and
fails on any divergence. The cross-backend bit-exactness results --- including a single attention
implementation that is the structural cure for the repetition collapse --- are the evidence that the law
holds in practice.
\end{abstract}
\section{The pain: coherent on Metal, collapsing on CUDA}
\label{sec:pain}
A concrete failure names the whole problem. Qwen3-8B-Q4\_K\_M, a single quantized model file, generates
coherent text on Apple Metal --- ``The capital of France is \textbf{Paris}'' --- and, from the
byte-identical file (SHA-256 verified) with the same sampler, \emph{collapses into repeated-token
degeneration} on CUDA and Vulkan during long generation. Nothing about the model changed. What changed is
the kernel: flash attention's accumulated numeric behavior on the CUDA/Vulkan tensor-core path differs
from the Metal path, and only the former drifts into the repetition attractor. The same operation
(attention) has three implementations and therefore three numeric behaviors, and one of them is wrong in
a way no unit test on that backend alone would flag as ``wrong'' --- it produces finite, plausible
numbers that steer generation into a ditch.
This is the defining bug class of a multi-backend engine, and it has three properties that make it
brutal: it is \emph{silent} (no crash, no NaN), it is \emph{backend-specific} (passes on the box you
tested), and it is \emph{emergent} (visible only in long generation, not in a single forward). You cannot
test four copies of a kernel \emph{into} agreement; agreement has to be a structural property of having
one definition.
\section{The idea: the CPU is the definition, not a backend}
\label{sec:idea}
The law is one sentence: \textbf{the CPU runtime is the single bit-exact reference every backend is
measured against.} The CPU is not ``another target that also happens to work.'' It is the \emph{oracle} ---
the place where the operation's meaning is defined --- and every GPU backend is correct exactly insofar as
it reproduces the oracle's bits.
Why can the CPU play this role, and why bit-\emph{exact} rather than merely close? Because a pointwise
(Map) operation is \emph{index-local} --- output element $i$ depends only on inputs at index $i$ --- and
the CPU runs it in program order. There is no reordering, no re-association, no fused-multiply-add
surprise: the CPU result is the arithmetic the source literally specifies. A GPU backend running the same
source computes the same index-local arithmetic and must agree \emph{to the bit}. The only sanctioned
source of difference is a \textbf{reduction}: a \code{sum} over a row, which a GPU computes with a
tree-reduction across threads and therefore in a different \emph{order} than the CPU's left fold. Float
addition is not associative, so the reordered sum can differ in the last bits. That difference is
\emph{bounded} (it is f32 re-association noise, and a small-$k$ control confirms it does not grow with the
reduction length --- the signature of per-element rounding, not a decode defect) and \emph{explained}
(it lives only at reductions, nowhere else). Everything that is index-local is bit-identical; everything
that reduces is identical up to a characterized rounding. Nothing is unexplained.
\section{The mechanism: a mirror, held by test}
\label{sec:mech}
The law is anchored by a discipline: for every op, the device code is a \emph{line-for-line mirror} of a
plain-Rust reference. The fusion engine's device morphism table mirrors its host table; every island's
\code{default} arm \emph{is} the CPU semantics (\cite{island}); every autotune variant is checked against
the oracle (\cite{tune}). Listing~\ref{lst:mirror} shows the shape --- the device \code{apply\_un\_dev}
and the host \code{apply\_un} are the same expressions --- and that identity is \emph{why} a fused kernel
equals its reference bit-exactly rather than approximately.
\begin{lstlisting}[language=Rust,caption={The mirror. The device morphism table (left, a \code{\#[device]}
fn) and the host reference (right, plain Rust) are the same arithmetic written twice. Bit-exactness is a
consequence of the mirror, not a tolerance that happens to pass.},label={lst:mirror}]
// device (#[device]) // host reference (plain Rust)
UnOp::Silu => x * dsigmoid::<F>(x), UnOp::Silu => x / (1.0 + (-x).exp()),
UnOp::Gelu => { let c = F::new(0.797_884_56); UnOp::Gelu => { const C: f32 = 0.797_884_56;
let inner = c*(x + F::new(0.044715)*x*x*x); 0.5*x*(1.0 + (C*(x + 0.044715*x*x*x)).tanh()) }
F::new(0.5)*x*(F::new(1.0)+inner.tanh()) }
\end{lstlisting}
The oracle is enforced across the whole op library --- \code{rms\_norm}, \code{layer\_norm}, \code{rope}
(both the NeoX half-split and GPT-J interleaved conventions), \code{sdpa} (online/flash softmax), the GDN
\code{conv1d}/\code{gating}/\code{scan}, the quantized matvecs, and the fusion chains --- across 36 CPU
tests. Each asserts the device result equals the plain-Rust reference within the characterized reduction
tolerance (and, for index-local ops, to the bit).
\section{Enforcement: a nightly three-backend GPU CI}
\label{sec:ci}
A law nobody checks is a hope. The oracle is enforced continuously by a GitHub Actions workflow that runs
a three-backend matrix on \emph{self-hosted} boxes --- one runner label per backend
(Listing~\ref{lst:ci}): \code{cuda} $\to$ spark (NVIDIA GB10, sm\_121), \code{rocm} $\to$ evo (AMD Strix
Halo, gfx1151, ROCm 7.x), \code{metal} $\to$ dbc (Apple M4 Max). Each runner builds the engine for its
backend and runs the same bit-exactness tests against the same CPU oracle.
\begin{lstlisting}[language=Rust,caption={The nightly matrix (\code{.github/workflows/gpu-tests.yml}). One
self-hosted runner per backend; the same tests, the same oracle. Serial execution keeps a single GPU
consumer at a time --- a necessity, not fastidiousness (\S\ref{sec:ci}).},label={lst:ci}]
on: { workflow_dispatch: {}, schedule: [ { cron: '0 8 * * *' } ] } // nightly + on demand
jobs:
test:
strategy: { matrix: { backend: [cuda, rocm, metal] } } // cuda->spark rocm->evo metal->dbc
runs-on: [self-hosted, "${{ matrix.backend }}"]
steps:
- uses: ./.github/actions/gpu-guard // one GPU consumer at a time
- run: cargo test -p hanzo-ml --features ${{ matrix.backend }} -- --test-threads 1
\end{lstlisting}
Two operational details are load-bearing. The workflow runs \emph{nightly and on demand}, not per push,
because these boxes double as development and benchmark machines. And it runs \code{--test-threads 1}
behind a GPU-guard step because concurrent GPU initializations on the ROCm box race in the kernel-fusion
driver (\code{kfd}) and wedge it until reboot --- so ``one GPU consumer at a time'' is a hard
serialization requirement, enforced by both the guard and the concurrency group. The practical result is
that ``the same op has the same numbers on every backend'' is a \emph{checked invariant}: a divergence
introduced on any backend fails the next nightly run against the oracle.
\section{Evaluation: the same source, every runtime, bit-exact}
\label{sec:eval}
Table~\ref{tab:bitexact} is the evidence that the law holds where it counts --- not merely that the DSL
\emph{compiles} for each backend, but that it \emph{dispatches} through each backend's real runtime and
matches the oracle. The ROCm figure is the falsifiable end-to-end proof that the toolchain seam to ROCm
7.13 works: the same \code{rms\_norm} source, lowered through cubecl-hip and dispatched on gfx1151 (device
query, module compile, launch, readback), matches the CPU oracle to $\text{max\_rel}=2.43\times10^{-7}$ ---
well under the $2\times10^{-3}$ reduction-tolerance gate, i.e. pure f32-reorder noise.
\begin{table}[h]
\centering\small
\setlength{\tabcolsep}{5pt}
\begin{tabular}{@{}lll@{}}
\toprule
\textbf{Path} & \textbf{Op(s)} & \textbf{Result} \\
\midrule
CPU oracle & rms/layer\_norm, rope$\times$2, sdpa, gdn, quant, fuse & 10/10 bit-exact (36 tests) \\
DSL $\to$ Vulkan (gfx1151) & rms\_norm, softmax, matvec, sdpa & bit-exact, real dispatch \\
DSL $\to$ Metal (M4 Max) & rms\_norm, layer\_norm, rope, sdpa & bit-exact ($2.35\mathrm{e}{-}7$) \\
DSL $\to$ ROCm (gfx1151) & rms\_norm & $\text{max\_rel}=2.43\mathrm{e}{-}7$ \\
\bottomrule
\end{tabular}
\caption{The same source, dispatched through each backend's real runtime, matches the CPU oracle. Values
are reduction-order f32 noise, well under the $2\mathrm{e}{-}3$ gate.}
\label{tab:bitexact}
\end{table}
The last row of the story returns to \S\ref{sec:pain}. Because attention is written \emph{once} in the
DSL --- one \code{sdpa} kernel with an online (flash-style) softmax, lowered to every backend --- the
``flash vs.\ eager vs.\ Metal, three numeric behaviors'' fork \emph{cannot occur}. There is nothing to
diverge. The repetition collapse is not a bug to fix on CUDA and Vulkan; under one implementation it is a
state that is unreachable. That is the difference between testing four kernels toward agreement and
having one definition: the first is endless, the second is structural.
\section{The takeaway}
A multi-backend engine must hold one \emph{numerically vetted} implementation per operation, because ``the
same op, a different kernel per backend'' is a correctness hazard first and a performance question second.
The CPU oracle is what makes ``numerically vetted'' concrete: the reference is the CPU's in-order
arithmetic, the device code is its mirror, the only sanctioned difference is characterized
reduction-order rounding, and a nightly three-backend CI turns the whole thing from a principle into a
gate. Fusion (\cite{fuse}), autotuning (\cite{tune}), and islands (\cite{island}) all rest on this one
foundation --- and the collapse method (\cite{method}) uses it as the bit-exact half of every
verify-then-delete decision.
\begin{thebibliography}{9}
\small
\bibitem{umbrella} Hanzo AI. \emph{One Source, Every Backend: The hanzo-kernel DSL.} Paper 07 (umbrella).
\bibitem{fuse} Hanzo AI. \emph{Fusion Is Composition: Kernel Fusion as the Functor Law.} Paper 07.1.
\bibitem{tune} Hanzo AI. \emph{The Schedule Is a Value: Comptime-Specialized Autotuning.} Paper 07.2.
\bibitem{island} Hanzo AI. \emph{Intrinsic Islands: An Oracle-Anchored Escape Hatch.} Paper 07.3.
\bibitem{oracle} Hanzo AI. \emph{One Oracle, Every Backend: The CPU Bit-Exactness Law.} Paper 07.4 (this paper).
\bibitem{mmq} Hanzo AI. \emph{When the Product Won't Collapse: An int8 Tensor-Core MMQ Negative Result.} Paper 07.5.
\bibitem{method} Hanzo AI. \emph{Verify, Then Delete: Collapsing a Kernel Cartesian Product.} Paper 07.6.
\end{thebibliography}
\end{document}
hanzo-collapse-methodology.pdf
BIN
View File
Binary file not shown.
hanzo-collapse-methodology.tex
+185
View File
@@ -0,0 +1,185 @@
\documentclass[11pt]{article}
\input{shared/kernel-preamble.tex}
\title{\textbf{Verify, Then Delete:\\ A Method for Collapsing a Kernel Cartesian Product}}
\author{Hanzo AI --- ML Systems}
\date{hanzo-kernel Paper 07.6 --- part of the \emph{One Source, Every Backend} series}
\begin{document}
\maketitle
\begin{abstract}
The other papers in this series build the machinery to write a GPU kernel once and lower it to every
backend. This paper is about the migration: how you drain $52{,}345$ lines of hand-written kernel
arithmetic across five backend trees down toward a single source \emph{without} a big-bang rewrite and
without ever shipping a slower engine. The method is a discipline we call \textbf{verify-then-delete}:
never retire a hand-tuned kernel on faith, and never all at once. A hand-written kernel is deleted only
when its single-source twin is (a) bit-exact against the shared CPU oracle \emph{and} (b) at least as fast
on a kernel-only benchmark --- in practice $\ge 97\%$ of the incumbent's throughput. A
per-\code{(op, type, backend)} registry flips one cell to the DSL the moment its twin passes both gates,
so the product drains to a sum one measured cell at a time, with a working, faster-or-equal engine at
every intermediate state. We give the ledger, the protocol, and the measured outcomes --- including a
migration that was \emph{7.1$\times$ faster} than the shader it replaced and bit-exact, and, just as
importantly, the migrations the gate correctly \emph{declined}.
\end{abstract}
\section{The ledger: what there is to collapse}
\label{sec:ledger}
Table~\ref{tab:ledger} is the denominator. It is the hand-written kernel arithmetic in the Hanzo ML and
engine repositories: the same $\sim$20 operations, specialized over $\sim$22 GGUF quantization formats,
re-implemented once per backend across five trees and five shading languages.
\begin{table}[h]
\centering
\begin{tabular}{llr}
\toprule
\textbf{Tree} & \textbf{Language} & \textbf{Lines} \\
\midrule
CUDA (ML) & CUDA C++ & 13{,}123 \\
ROCm & HIP C++ & 6{,}725 \\
Metal & Metal Shading Language & 3{,}530 \\
Vulkan & GLSL $\rightarrow$ SPIR-V & 7{,}105 \\
CUDA (engine) & CUDA C++ & 21{,}862 \\
\midrule
\textbf{Total} & 5 trees & \textbf{52{,}345} \\
\bottomrule
\end{tabular}
\caption{The collapse ledger: hand-written, non-generated, non-vendored kernel arithmetic.\protect\footnotemark{}
Roughly 85\% is a deletable target --- the duplicated arithmetic --- as opposed to the genuinely peak,
genuinely per-backend kernels (block-quant tensor-core MMQ, \cite{mmq}) that stay behind islands.}
\label{tab:ledger}
\end{table}
\footnotetext{These are the \emph{deletable} lines. A na\"ive \code{wc -l} over all
\code{*.cu/*.cuh/*.hip/*.hpp/*.metal/*.comp} is roughly $2\x$ larger ($\sim$112k) because it also counts
C++ headers, a \code{target/package/} build-artifact copy, machine-generated codebook tables (the i-quant
grids), and vendored upstream kernels (FlashInfer, Marlin, GPTQ) --- none of which are hand-written
duplication or a DSL collapse target.}
The product is not merely a maintenance cost; it is the breeding ground for the divergence bugs the oracle
paper catalogs (\cite{oracle}). Draining it is therefore a correctness project as much as a code-size one.
But it must be drained safely: a rewrite that regresses performance, or that lands as one enormous
untestable change, is worse than the duplication it removes.
\section{The protocol: verify, then delete}
\label{sec:protocol}
The method is three rules, applied per cell of the $\text{types}\times\text{backends}$ matrix.
\begin{enumerate}[leftmargin=1.6em,itemsep=2pt]
\item \textbf{Coverage first.} The DSL fills the entire matrix from one source. Many cells were previously
\emph{missing} or buggy; those gain correct coverage immediately, before any deletion --- a strict
improvement with nothing to retire.
\item \textbf{The two-part gate.} A hand-tuned kernel is retired for its single-source twin \emph{only}
when the twin is both:
\begin{itemize}[leftmargin=1.4em,itemsep=0pt]
\item \textbf{bit-exact} against the shared CPU oracle (the correctness half; \cite{oracle}), and
\item \textbf{at least as fast} on a kernel-only benchmark --- buffers created once, $N$ dispatches
enqueued back-to-back, one final sync, so the number reflects the kernel and not the host round-trip. In
practice the bar is $\ge 97\%$ of the incumbent's throughput.
\end{itemize}
The number decides, per quant type, per backend. Not a human's confidence --- the benchmark.
\item \textbf{Delete on parity.} When a twin passes both gates, its hand-written copies are deleted and a
per-\code{(op, type, backend)} registry flips that cell to the DSL, keyed by the gate result. Until it
passes, the cell keeps its hand-tuned kernel. The registry is the seam that lets the product drain to a
sum one measured cell at a time, with a working, faster-or-equal engine at every intermediate state.
\end{enumerate}
The protocol is reversible and never a downgrade by construction. It is a scoreboard with a gate, not a
rewrite.
\section{Measured migrations: the wins}
\label{sec:wins}
The method is not hypothetical; cells have flipped on \code{main}. All figures below are kernel-only
benchmarks on the named hardware.\footnote{Throughput and latency are session benchmarks on evo (AMD Strix
Halo, gfx1151; RADV/Vulkan and ROCm 7.13). The bit-exactness each accompanies is enforced by the committed
tests of the oracle paper (\cite{oracle}).}
\paragraph{Portability costs nothing (\dpa{} matvec).} The central fear of any portable DSL is a
performance downgrade. It does not materialize on the bandwidth-bound matvec that dominates decode. A tuned
single-source \dpa{} Q8 matvec --- \code{Vector<i8,4>}$\to$\code{OpSDot}, one thread block per output row
with threads striding the groups so adjacent threads read adjacent packed weights (coalesced 256-byte
transactions), partials tree-reduced in shared memory --- measures \textbf{196~GB/s} on evo ($8192^2$,
cache-busting), which is \textbf{118\% of the $\sim$166~GB/s hand-tuned kernel} it replaces and $\sim$77\%
of the DRAM roofline, bit-exact. The na\"ive one-thread-per-row version that started at $\sim$17~GB/s was
memory-\emph{uncoalesced}, not DSL-limited; the fix was two standard steps --- coalescing and reduction ---
both expressed in the DSL. The DSL reaches the same hardware integer-dot instruction and, with standard
coalescing, \emph{beats} the hand-written code.
\paragraph{A migration that improved performance (\code{add\_rmsnorm}).} Two hand-written Vulkan kernels
were drained on \code{main}: the DSL block-per-row \code{rms\_norm} (270~GB/s) and \code{softmax}
(268~GB/s) replaced na\"ive \code{.comp} shaders, each now the production kernel with its \code{.comp}
deleted. The sharpest case is the fused residual-add\,+\,RMSNorm: the DSL \code{add\_rmsnorm} twin ran
\textbf{7.1$\times$ faster} than the incumbent shader ($166.5\to23.4$~ms at $4096^2$ on RADV) \emph{and}
was bit-exact ($s=x{+}\text{res}$ byte-identical, $y$ within $4.8\times10^{-7}$). The ``uniformity''
change was also a large speedup, because the incumbent was na\"ive and the DSL block kernel is coalesced.
Table~\ref{tab:migr} collects the flipped cells.
\begin{table}[h]
\centering\small
\setlength{\tabcolsep}{6pt}
\begin{tabular}{@{}llll@{}}
\toprule
\textbf{Kernel} & \textbf{Backend} & \textbf{Outcome} & \textbf{Correctness} \\
\midrule
\dpa{} Q8 matvec & Vulkan (gfx1151) & 196~GB/s, 118\% of hand-tuned & bit-exact \\
\code{rms\_norm} & Vulkan (gfx1151) & 270~GB/s, \code{.comp} deleted & bit-exact \\
\code{softmax} & Vulkan (gfx1151) & 268~GB/s, \code{.comp} deleted & bit-exact \\
\code{add\_rmsnorm} & Vulkan (RADV) & 7.1$\times$ ($166.5\to23.4$~ms) & bit-exact \\
\bottomrule
\end{tabular}
\caption{Cells flipped to the DSL on \code{main}. Every deletion cleared both gates: bit-exact against the
oracle and at least as fast (here, faster). The engine is faster-or-equal at every step.}
\label{tab:migr}
\end{table}
\section{The discipline cuts both ways: the declines}
\label{sec:declines}
A gate that never says no is not a gate. Where a hand-tuned kernel is \emph{already optimal} --- the
Vulkan quant matvec, Metal's \code{simd\_sum} reductions, the fused \code{rope\_norm} epilogue --- a
bit-exact DSL twin was built, measured, and \emph{correctly not swapped}. A uniformity-only change with
zero speedup and real regression risk is churn, not migration. Such a twin is kept as the portable
\emph{fallback} for backends that lack the tuned version, and the tuned peak stays the specialized fast
path. ``One source'' is the goal, but the gate serves it only by also refusing the swaps that would trade
speed for tidiness. We only ever add coverage and correctness.
\section{An engineering note: opening a backend so it can flip}
\label{sec:rocm}
A cell cannot flip on a backend that does not build. Shipping the DSL on AMD required fixing two
independent upstream defects \emph{without} advancing the lowering engine past its newest published
release (no version escape exists for either). First, the HIP FFI bindings are selected from
\code{hipconfig --version}; ROCm 7.13 reports a patch past the newest shipped binding, so the upstream
build emitted a feature with no bindings module and every symbol went undefined. The fix is a vendored
\code{[patch.crates-io]} fork whose \code{build.rs} clamps an unknown-newer patch to the latest shipped
binding --- ABI-safe because AMD versions every changed symbol (R0600) and 7.13 still maps
\code{hipGetDeviceProperties} to that revision. Second, the HIP server uses a \code{MultiStream} gated
behind a \code{std}-only configuration the \code{hip}/\code{rocm} features do not pull, so the ROCm feature
is declared to include \code{cubecl/std}. Both are recorded because a portable DSL lives or dies on the
least-loved backend building cleanly --- and only a backend that builds can have its cells verified and
its hand-written kernels deleted.
\section{The takeaway}
Collapsing a Cartesian product of kernels is a migration problem, and migrations fail when they are
faith-based or big-bang. Verify-then-delete makes both failure modes structurally impossible: the CPU
oracle supplies an objective correctness gate (\cite{oracle}), the kernel-only benchmark supplies an
objective performance gate, and the per-cell registry supplies incrementality. The result is an engine
that is faster-or-equal at every intermediate state while its five hand-written kernel trees drain toward
one source --- and a boundary, honestly drawn, where one op (\cite{mmq}) stays hand-written behind an
island. The method is repeatable: point it at the next cell, run both gates, let the number decide.
\begin{thebibliography}{9}
\small
\bibitem{umbrella} Hanzo AI. \emph{One Source, Every Backend: The hanzo-kernel DSL.} Paper 07 (umbrella).
\bibitem{fuse} Hanzo AI. \emph{Fusion Is Composition: Kernel Fusion as the Functor Law.} Paper 07.1.
\bibitem{tune} Hanzo AI. \emph{The Schedule Is a Value: Comptime-Specialized Autotuning.} Paper 07.2.
\bibitem{island} Hanzo AI. \emph{Intrinsic Islands: An Oracle-Anchored Escape Hatch.} Paper 07.3.
\bibitem{oracle} Hanzo AI. \emph{One Oracle, Every Backend: The CPU Bit-Exactness Law.} Paper 07.4.
\bibitem{mmq} Hanzo AI. \emph{When the Product Won't Collapse: An int8 Tensor-Core MMQ Negative Result.} Paper 07.5.
\bibitem{method} Hanzo AI. \emph{Verify, Then Delete: Collapsing a Kernel Cartesian Product.} Paper 07.6 (this paper).
\end{thebibliography}
\end{document}
hanzo-comptime-autotune.pdf
BIN
View File
Binary file not shown.
hanzo-comptime-autotune.tex
+204
View File
@@ -0,0 +1,204 @@
\documentclass[11pt]{article}
\input{shared/kernel-preamble.tex}
\title{\textbf{The Schedule Is a Value:\\ Comptime-Specialized Autotuning Without Forks}}
\author{Hanzo AI --- ML Systems}
\date{hanzo-kernel Paper 07.2 --- part of the \emph{One Source, Every Backend} series}
\begin{document}
\maketitle
\begin{abstract}
A single-source kernel DSL promises to kill the ``one operation, $N$ backends, $N$ numeric behaviors''
bug class by writing each op once. But there is a subtler place the same bug hides: the \emph{schedule}
--- the tile size, vector width, unroll factor, and thread-block dimension a kernel is tuned with, which
differs per device. The usual fix is to fork a hand-tuned kernel per device, and a fork is exactly what
reintroduces the bug: two schedules, drifting independently, one of them one day reordering an
accumulation and changing the numbers. This paper shows how to autotune without forking. The schedule is
exposed as a set of compile-time (\code{\#[comptime]}) knobs on the \emph{one} source; each knob value is
a distinct compiled kernel the lowering engine monomorphizes; a small set of these variants is timed once
per \code{(device, op, shape)} and the winner cached to disk. Because the knobs never touch the
arithmetic, \emph{every variant is byte-identical} --- the tuner buys speed and can never spend
correctness. We give the three-value algebra (\code{Variant}, \code{Tuned}, \code{Tuner}), the
per-device on-disk cache, the two shipped variant tables, and the offline tests that prove
select/cache/reload works without a GPU.
\end{abstract}
\section{The pain: a fork that drifts}
\label{sec:pain}
Suppose you have written a matvec kernel once, in a portable DSL, and it lowers correctly to every
backend. You are not done, because ``correct everywhere'' is not ``fast everywhere.'' The fastest
schedule for an AMD iGPU is not the fastest for an NVIDIA datacenter part: they differ in warp/subgroup
width, in how much shared memory pays off, in the ideal number of output elements a thread should carry
(the unroll factor), in the block size that fills the machine. The reflex is to fork --- keep a
\code{matvec\_amd} and a \code{matvec\_nvidia}, each hand-tuned.
That reflex is the original sin the DSL exists to prevent, wearing a schedule's clothes. Two forks are
two kernels, maintained by two hands, drifting apart. The day someone reorders a loop in one fork ``to
help the unroll,'' its accumulation order changes, its rounding changes, and now the same operation
produces different bits on two backends --- the precise failure a single-source DSL was supposed to make
impossible (\cite{oracle}). Worse, the drift is invisible until a model degenerates on one device and not
another.
We want the performance a per-device schedule buys \emph{without} the second kernel that buys the bug.
\section{The idea: a schedule is a compile-time value}
\label{sec:idea}
The DSL already gives one source lowered to every backend. The remaining fork --- the schedule --- can be
folded into that one source if we treat the schedule not as \emph{code} but as a \emph{value}. In \hk{},
a kernel parameter marked \code{\#[comptime]} is resolved at compile time, and the lowering engine
\emph{monomorphizes} the kernel on it: each distinct comptime value produces a distinct compiled kernel,
specialized as if you had hand-written it. So we lift the schedule knobs to \code{\#[comptime]}
parameters of the one kernel (Listing~\ref{lst:tunedkernel}). A ``variant'' is then not a forked file but
a \emph{choice of comptime tuple}. There is one source; only the chosen tuple differs per device.
\begin{lstlisting}[language=Rust,caption={One autotuned source. The unroll factor \code{rpt}
(rows-per-thread) is a comptime knob; the block size is the host's knob. Every \code{(rpt, block)} is a
distinct compiled kernel yet bit-identical to the reference: the per-row fold is the same order, so only
the thread$\to$row mapping and occupancy differ.},label={lst:tunedkernel}]
#[kernel(targets(cuda, metal, vulkan, webgpu, cpu), unchecked)]
pub fn rms_norm_tuned<F: Float>(x: &Array<F>, w: &Array<F>, out: &mut Array<F>,
eps: &Array<F>, #[comptime] n: usize, #[comptime] rpt: usize) {
let nrows = out.len() / n;
let base = ABSOLUTE_POS * rpt;
#[unroll]
for r in 0..rpt {
let row = base + r;
if row < nrows {
let b = row * n;
let mut ss = F::new(0.0);
for i in 0..n { let v = x[b + i]; ss += v * v; } // same fold, same order, every variant
let denom = (ss / F::cast_from(n as u32) + eps[0]).sqrt();
for i in 0..n { out[b + i] = x[b + i] / denom * w[i]; }
}
}
}
\end{lstlisting}
The critical invariant is stated in the comment and enforced by test: the knobs change the
thread-to-work mapping and the occupancy, \emph{not} the arithmetic. The per-row sum-of-squares is the
same left fold in the same order for every \code{(rpt, block)}. Therefore all variants are
\textbf{byte-identical} to each other and to the reference. The autotuner is choosing among numerically
indistinguishable kernels; it cannot pick a ``faster but slightly wrong'' one, because there is no wrong
one to pick.
\section{The mechanism: three orthogonal values}
\label{sec:mech}
Autotuning decomposes into three plain values, each testable in isolation (Listing~\ref{lst:tune}).
\begin{itemize}[leftmargin=1.4em,itemsep=2pt]
\item A \textbf{\code{Variant}} is a named schedule: a closure \code{f(iters) -> (output, ms)} that
launches the schedule \code{iters} times after a fixed warmup and returns its output and mean
milliseconds per dispatch. The same closure both \emph{benchmarks} (many iters) and \emph{executes} (one
iter) --- this is the crate's established \code{*\_bench} shape, so no separate ``run'' and ``time'' paths
can disagree.
\item A \textbf{\code{Tuned}} is the fluent builder: collect the variants for one
\code{(op, shape-key)}, then \code{run}. It reads like the fusion builder --- fluent, no macro, no
ceremony.
\item A \textbf{\code{Tuner}} is the cache: an in-memory \code{(device, op, key) -> winner} map backed by
a per-device on-disk file. It is a \emph{plain value} constructed over an explicit cache root, so it is
fully unit-testable offline --- no GPU, no environment, no global state. The process-wide default roots
at the XDG cache directory.
\end{itemize}
\begin{lstlisting}[language=Rust,caption={The production surface. One source per op, autotuned per device,
no forks. On a cache hit only the winner launches (once, to produce its output); the losers are never
run.},label={lst:tune}]
let y = Tuned::new("rms_norm", format!("rows={rows},n={n}"))
.variant("b64_r1", |it| rms_norm_tuned_bench(&c, x, w, rows, n, eps, 1, 64, it))
.variant("b128_r1", |it| rms_norm_tuned_bench(&c, x, w, rows, n, eps, 1, 128, it))
.variant("b256_r1", |it| rms_norm_tuned_bench(&c, x, w, rows, n, eps, 1, 256, it))
.variant("b64_r2", |it| rms_norm_tuned_bench(&c, x, w, rows, n, eps, 2, 64, it))
.run(&c); // times each once per (device,op,shape), caches the winner, thereafter runs it
\end{lstlisting}
The selection logic is deliberately small. On a cache \emph{miss}, every variant is timed
\code{TUNE\_ITERS} times (25), the fastest is kept, its winning name is recorded, and its
already-computed output is returned --- no redundant re-run. On a cache \emph{hit}, only the winner is
launched, once, to produce its output; the losers are never launched. If a cached winner name is no
longer among the current variants (the set changed), that is treated as a miss and re-tuned rather than
trusting a stale name.
\paragraph{The cache is a value, not a database.} The on-disk format is one TSV file per device ---
\code{\$XDG\_CACHE\_HOME/hanzo-kernel/autotune/<device>.tsv} --- each line \code{op\textbackslash tkey\textbackslash twinner}.
No serde, no new dependency, no schema migration. Cross-process writes append and are idempotent (the
loader keeps the last line for a key), which is exactly how an advisory autotune cache should behave: a
hint that is safe to lose and safe to race on.
\section{Evaluation: the variant tables and the offline proof}
\label{sec:eval}
Two ops ship autotuned (Table~\ref{tab:variants}). Both variant sets are proven on the CPU runtime, which
needs no GPU to establish the two things under test: that the variants are numerically indistinguishable,
and that the select/cache machinery works.
\begin{table}[h]
\centering\small
\setlength{\tabcolsep}{5pt}
\begin{tabular}{@{}lll@{}}
\toprule
\textbf{Op} & \textbf{Knobs} & \textbf{Variants (all byte-identical)} \\
\midrule
\code{rms\_norm} & cube dim $\{64,128,256\}\times$ unroll $\{1,2\}$ & \code{b64\_r1 b128\_r1 b256\_r1 b64\_r2} \\
\code{matvec\_q8\_dp4a} & cube dim $\{64,128,256\}\times$ vec.\ width $\{1,2,4,8\}$ & \code{b64\_v1 b64\_v4 b128\_v2 b256\_v8} \\
\bottomrule
\end{tabular}
\caption{The autotune variant sets (\hk{} 0.2.17). A knob value is a distinct compiled kernel; the tuner
caches the winner per \code{(device, op, shape)}. For \dpa{} matvec the vector-width knob is the ILP
factor (groups per unrolled step); the per-group int8 dot is accumulated in the same order for every
width, so all four are byte-identical.}
\label{tab:variants}
\end{table}
The committed tests assert, for both ops:
\begin{enumerate}[leftmargin=1.5em,itemsep=1pt]
\item \textbf{Byte-identity.} Every schedule's output has identical bits to the reference (RMSNorm:
exactly equal; \dpa{}: byte-identical across variants and within f32-reorder tolerance
$2\times10^{-3}$ of the per-element oracle). The knobs do not touch the numerics.
\item \textbf{Select and cache.} A first call is a miss that benchmarks all four variants
(\code{benched == 4}) and records the winner. A second call is a hit that benchmarks \emph{zero}
(\code{benched == 0}) and returns the winner's output. A fresh \code{Tuner} over the same directory
reloads the winner \emph{from disk} and skips timing. A distinct shape-key is an independent decision. A
shrunk variant set (winner absent) re-tunes rather than trusting a stale name.
\end{enumerate}
Timing on the CPU proves the select/cache mechanism; a GPU run additionally makes the \emph{winner}
meaningful, since on the CPU the schedules are near-equivalent and the chosen winner is
arbitrary-but-stable. The mechanism, not the specific winner, is what the offline tests pin down.
\section{Why first-party, not the engine's autotuner}
The lowering engine (CubeCL) ships an autotuner; it was studied and declined for three concrete reasons,
not out of not-invented-here. First, it is re-exported only from \code{cubecl\_runtime}, so adopting it
would force a \emph{second} engine crate into a facade architected so that only one line of one
\code{Cargo.toml} names the engine at all --- a deliberate ``values, not places'' boundary the whole DSL
is built to preserve. Second, its persistent cache lives under \code{target/} (via an MD5 checksum and
serde), not the XDG path this layer is required to use. Third, its generic-associated-type input
machinery is a large surface for what is, in the end, ``pick the fastest of $N$ named closures.'' The
concepts above --- variant, tuned set, cache keyed by device and shape --- are \emph{deliberately} the
same as the engine's; they are sized to this crate and pointed at XDG. A thin analog, not a new
philosophy.
\section{The takeaway}
The schedule was the last place a backend forked a kernel, and the fork was the last hiding place for the
divergence bug. Making the schedule a compile-time value folds that fork into the one source: the tuner
enumerates numerically-identical variants, times them per device, and caches the winner, so the engine is
autotuned and single-source at once. Speed is chosen; correctness is never spent. Combined with fusion as
composition (\cite{fuse}) and the oracle law (\cite{oracle}), the schedule stops being a place where two
implementations can disagree, because there is only ever one.
\begin{thebibliography}{9}
\small
\bibitem{umbrella} Hanzo AI. \emph{One Source, Every Backend: The hanzo-kernel DSL.} Paper 07 (umbrella).
\bibitem{fuse} Hanzo AI. \emph{Fusion Is Composition: Kernel Fusion as the Functor Law.} Paper 07.1.
\bibitem{tune} Hanzo AI. \emph{The Schedule Is a Value: Comptime-Specialized Autotuning.} Paper 07.2 (this paper).
\bibitem{island} Hanzo AI. \emph{Intrinsic Islands: An Oracle-Anchored Escape Hatch.} Paper 07.3.
\bibitem{oracle} Hanzo AI. \emph{One Oracle, Every Backend: The CPU Bit-Exactness Law.} Paper 07.4.
\bibitem{mmq} Hanzo AI. \emph{When the Product Won't Collapse: An int8 Tensor-Core MMQ Negative Result.} Paper 07.5.
\bibitem{method} Hanzo AI. \emph{Verify, Then Delete: Collapsing a Kernel Cartesian Product.} Paper 07.6.
\end{thebibliography}
\end{document}
hanzo-fusion-is-composition.pdf
BIN
View File
Binary file not shown.
hanzo-fusion-is-composition.tex
+283
View File
@@ -0,0 +1,283 @@
\documentclass[11pt]{article}
\input{shared/kernel-preamble.tex}
\title{\textbf{Fusion Is Composition:\\ Kernel Fusion as the Functor Law}}
\author{Hanzo AI --- ML Systems}
\date{hanzo-kernel Paper 07.1 --- concept core of the \emph{One Source, Every Backend} series}
\begin{document}
\maketitle
\begin{abstract}
Kernel fusion is usually presented as a compiler optimization: a pass that inspects a graph and, when
a cost model approves, merges adjacent operations to avoid writing an intermediate array to memory. This
paper takes the opposite stance --- fusion is not a pass, it is an \emph{algebraic identity}. An
elementwise GPU kernel is the lift of a scalar function to every array index, and such lifts obey the
functor law $\mathrm{map}(g)\circ\mathrm{map}(f)=\mathrm{map}(g\circ f)$. Read right-to-left, that law
\emph{is} fusion: one kernel computing the composed function, no intermediate. Reduce and Scan are the
points where the law stops applying --- they are algebraic \emph{fences}, decidable from the operation's
class, not from a pattern matcher. We show how this single idea becomes a working, one-launch fusion
engine in \hk{}: a tiny op algebra whose fusibility rule is a total function of op class; one generic
kernel that interprets any composed chain as straight-line arithmetic; and the same algebra lifted from a
line to an arbitrary DAG. Every claim is gated bit-exactly on a CPU reference. The concrete payoff is
measured on a real model pattern --- the per-layer Gated-DeltaNet chain fuses from 9 kernel launches to
5, byte-for-byte identical to the unfused path.
\end{abstract}
\section{The pain: an intermediate nobody wanted}
\label{sec:pain}
Consider the tail of a Transformer feed-forward block, after its two matmuls have produced two arrays
\code{gate} and \code{up}, each of length $n$. The SwiGLU activation is
\[
y \;=\; \mathrm{silu}(\code{gate}) \cdot \code{up}, \qquad \mathrm{silu}(x)=x\cdot\sigma(x).
\]
A \emph{kernel} is a function the GPU runs once per array element in parallel; a kernel \emph{launch} is
one dispatch of that function over the whole array. The obvious way to compute $y$ is two kernels: one
computes $t=\mathrm{silu}(\code{gate})$, writing $n$ floats to device memory (DRAM); a second reads $t$
and \code{up} back from DRAM and writes $y$. That middle array $t$ is \emph{materialized} --- written to
DRAM and immediately read back --- for no reason but that we wrote two kernels. On a decode step, where
the whole workload is memory-bound (you are limited by how fast you can move bytes, not by arithmetic),
that round-trip is pure waste.
Now scale the pattern. Qwen3.5 and Qwen3.6 are hybrid models whose bulk layers are Gated-DeltaNet
(GDN) linear attention, and each GDN layer emits a short elementwise chain ---
$\mathrm{rms\_norm}(x)\cdot w\cdot\mathrm{silu}(g)$ and its siblings --- as a \emph{storm} of tiny kernel
launches: roughly \textbf{9 launches per layer} on the ops-composed CUDA path. Each launch reads its
inputs from DRAM, does a trivial amount of arithmetic, and writes its output back to DRAM for the next
launch to read. Multiply by dozens of layers and every decoded token pays for a parade of intermediates.
This is not a corner case; it is a measurable part of why an ops-composed engine trails a hand-fused one.
The question this paper answers is: \emph{when may two adjacent kernels be merged into one, and how do
you do it once, correctly, for every backend?} The usual answer is ``run a fusion pass with a pattern
matcher and a cost model.'' Ours is simpler and total.
\section{The idea: an elementwise kernel is a lifted function}
\label{sec:idea}
Strip a pointwise (elementwise) kernel to its essence and it is a scalar function $f:F\to F$ applied
independently at every index $i$:
\[
\mathrm{map}(f)\,:\;\text{Array}\to\text{Array},\qquad \mathrm{map}(f)(A)[i]=f(A[i]).
\]
The word \emph{map} is exact: it is the same map a Rust programmer writes as \code{v.iter().map(f)}. Two
such kernels in sequence apply $f$, store the result, then apply $g$. But applying $g$ after $f$ at every
index is, by definition, applying $g\circ f$ at every index. That is the \textbf{functor law}:
\[
\boxed{\;\mathrm{map}(g)\circ\mathrm{map}(f)\;=\;\mathrm{map}(g\circ f)\;}
\]
Read left-to-right, the equation is two kernels with a materialized array between them. Read
right-to-left, it is \emph{one} kernel computing the composed function $g\circ f$ with \emph{no}
intermediate. \textbf{Kernel fusion is nothing but that law, read right-to-left.}
If this feels familiar it is the transducer insight in disguise: \code{map(f)} and \code{map(g)} are
transformations that do not mention the collection they run over, so you may compose the
\emph{transformations} and apply the composite to the collection \emph{once}, never allocating the
intermediate collection. Here the collection is a GPU array and ``once'' is a single kernel launch.
Categorically, each Map is a morphism $\text{Array}\to\text{Array}$ and \emph{composition of morphisms is
fusion}. There is no separate ``fusion pass'' to invent; there is only the associativity of $\circ$.
\section{The class decides fusibility}
\label{sec:class}
Not every op is a Map, and the ones that are not are exactly the ones that cannot be fused this way. The
whole surface is two classes:
\begin{itemize}[leftmargin=1.4em,itemsep=2pt]
\item \textbf{Map} (a pointwise morphism): output element $i$ depends only on input(s) at index $i$
(plus compile-time constants). It is shape-preserving and \emph{index-local}. Maps compose freely: any
run of adjacent Maps is one Map, by the functor law. This is the fusible family ---
\code{Add Sub Mul Div Neg Recip Abs Exp Rsqrt Sigmoid Tanh Silu Gelu}. (\code{Silu} and \code{Gelu} are
not special cases: they are \emph{defined} as compositions of the primitives, so they are already one
composed morphism.)
\item \textbf{Reduce / Scan} (a contraction): output element $i$ depends on a \emph{whole row or axis}
--- \code{sum}, \code{max}, \code{mean}, and the operations built on them (RMSNorm, softmax, matmul).
This is \emph{not} index-local. It cannot be absorbed into an adjacent Map by the same lift, because it
needs cross-lane communication: threads must exchange partial results (through shared memory or a
subgroup reduction) to compute one output. A Reduce is therefore a \textbf{fence}.
\end{itemize}
The decision ``can these two adjacent ops fuse?'' is then a total function of their class --- no pattern
matcher, no cost model:
\[
\mathrm{fusible}(a,b)\;:=\;\mathrm{is\_map}(a)\wedge\mathrm{is\_map}(b).
\]
At a fence there is still fusion, just attached at the fence's \emph{ports}: a Map run \emph{feeding} the
reduction fuses into its input read (a \emph{prologue}), and a Map run \emph{consuming} the reduction's
result fuses into its output write (an \emph{epilogue}). The reduction kernel itself is the seam where
the pattern changes. RMSNorm is the canonical shape: \code{square}(Map, prologue) $\to$ \code{sum}(fence)
$\to$ \code{scale-by-rsqrt}(Map, epilogue) --- three regions, the two Map runs each collapsed to one.
\section{The mechanism: reify, fold, interpret}
\label{sec:mech}
Three small steps turn the algebra into a running engine. All of it lives in one crate module and is
proven on a CPU reference (\S\ref{sec:eval}).
\paragraph{Reify.} A pointwise expression is recorded as a tree of scalar morphisms rather than computed
(Listing~\ref{lst:expr}). \code{Cur} is the value flowing down the pipeline at the current element;
\code{In(id)} is a side-input array read at the current index; \code{Const} is a compile-time scalar.
Operator overloading makes building one read like the math it denotes. Building an \code{Expr} computes
nothing --- it records the trace.
\begin{lstlisting}[language=Rust,caption={The op algebra. The whole fusible surface is a tree of unary
and binary pointwise morphisms over the current value, side inputs, and constants.},label={lst:expr}]
enum Expr { Cur, In(usize), Const(f32), Un(UnOp, Box<Expr>), Bin(BinOp, Box<Expr>, Box<Expr>) }
// Morphism composition: substitute `inner` for every `Cur` leaf. This is the ONLY operation
// the fusion pass needs -- map(g) . map(f) becomes map(g[Cur := f]).
fn compose_cur(&self, inner: &Expr) -> Expr {
match self {
Expr::Cur => inner.clone(),
Expr::Un(op, a) => Expr::Un(*op, Box::new(a.compose_cur(inner))),
Expr::Bin(op,a,b)=> Expr::Bin(*op, Box::new(a.compose_cur(inner)),
Box::new(b.compose_cur(inner))),
leaf => leaf.clone(),
}
}
\end{lstlisting}
\paragraph{Fold.} The fusion pass is one fold over the chain of steps: accumulate adjacent Maps by
composing their expressions (\code{compose\_cur} \emph{is} $g[\code{Cur}:=f]$), and cut a new region at
every Reduce. That is the entire pass (Listing~\ref{lst:fold}). Its output is a list of regions; the
number of regions is the number of kernels, and it is strictly fewer than the number of steps whenever
any two Maps were adjacent --- that gap \emph{is} the fusion.
\begin{lstlisting}[language=Rust,caption={The fusion pass: fold adjacent Maps by composing their
expressions; cut at every Reduce. No matcher, no cost model --- just the class.},label={lst:fold}]
fn fuse(&self) -> Vec<Region> {
let (mut regions, mut acc) = (Vec::new(), None);
for step in &self.steps {
match step {
Step::Map(e) => acc = Some(match acc.take() { // g after f ==> g[Cur := f]
None => e.clone(), Some(prev) => e.compose_cur(&prev),
}),
Step::Reduce(r) => { if let Some(e) = acc.take() { regions.push(Region::Map(e)); }
regions.push(Region::Reduce(*r)); }
}
}
if let Some(e) = acc.take() { regions.push(Region::Map(e)); }
regions
}
\end{lstlisting}
\paragraph{Interpret.} A fused Map region is lowered to a small linear ``op program'' carried as a
compile-time value, and \emph{one} generic kernel walks that program in a compile-time-\emph{unrolled}
loop (Listing~\ref{lst:interp}). Because the loop and the branch are resolved at compile time (the DSL's
lowering engine monomorphizes on the program value), each instruction becomes straight-line arithmetic
in one kernel body: the input is read once, threaded through every op, and written once. There are no
intermediate arrays and exactly one launch, whatever the chain length. This is not a hand-written fixed
kernel --- any linear elementwise chain runs through this one \code{fused\_interp}.
\begin{lstlisting}[language=Rust,caption={The one fused kernel. \code{prog} rides as a compile-time
value, so \code{\#[unroll]} plus a comptime \code{match} lower the whole chain to straight-line code:
one launch, zero intermediates.},label={lst:interp}]
#[kernel(targets(cuda, metal, vulkan, webgpu, cpu), unchecked)]
pub fn fused_interp<F: Float>(x: &Array<F>, sides: &Sequence<Array<F>>,
out: &mut Array<F>, #[comptime] prog: Program) {
let i = ABSOLUTE_POS;
if i < out.len() {
let mut acc = x[i];
#[unroll]
for step in 0..prog.instrs.len() {
match comptime!(prog.instrs[step]) {
Instr::Un(op) => acc = apply_un_dev::<F>(op, acc),
Instr::BinIn(op, slot) => acc = apply_bin_dev::<F>(op, acc, sides.index(slot)[i]),
Instr::BinConst(op, b) => { let c = comptime!(f32::from_bits(b));
acc = apply_bin_dev::<F>(op, acc, F::new(c)) }
}
}
out[i] = acc;
}
}
\end{lstlisting}
The ergonomic surface is a builder that reads like the composition it denotes:
\code{Fuse::new(a).\allowbreak{}mul(w).\allowbreak{}add(b).\allowbreak{}silu().\allowbreak{}run(c)}
appends one instruction per call, compiles them to one program, and dispatches exactly one
\code{fused\_interp}. The constant leaf is stored as \code{f32::to\_bits} so the program is
\code{Eq + Hash} --- the lowering engine keys one compiled kernel per distinct fused chain.
\section{From a line to a DAG}
\label{sec:dag}
A real forward pass is not a straight line. A value \emph{fans out} (is consumed by several later ops)
and a binary op \emph{fans in} (reads two independently computed values). The expression
$\mathrm{silu}(a)\cdot b + a$ reuses $a$ twice; a linear chain cannot express it in one region, but a
DAG can, and it still fuses to \emph{one} kernel. The same algebra lifts from a path to a graph with no
new ideas:
\begin{enumerate}[leftmargin=1.5em,itemsep=1pt]
\item \textbf{Trace.} A \code{NodeId}-taping newtype records an arbitrary pointwise DAG as a
\code{Vec<Node>} (operator overloads push nodes instead of computing). The node algebra is the \emph{same}
\code{UnOp}/\code{BinOp}/\code{Red} as the line case --- one IR, reused.
\item \textbf{Partition.} Flood-fill maximal Map-only connected regions; cut every edge into a fence
(a Reduce, or any non-elementwise op --- matmul, gather, conv). This is the identical \code{Class}
predicate from \S\ref{sec:class}, applied to graph edges instead of list-adjacency.
\item \textbf{Lower.} Each Map region compiles to a small DAG program over a per-element \emph{slot}
array, and one generic \code{dag\_interp} kernel interprets it. Fan-out is a slot read by several later
instructions; fan-in is a binary reading two slots. This generalizes \code{fused\_interp}'s single
accumulator (a line) to $n$ accumulators (a graph): one launch, zero intermediates, any DAG shape.
\end{enumerate}
\section{Evaluation: the bit-exact gate}
\label{sec:eval}
Fusion must not change the numbers. Every fused chain is checked three ways on the CPU runtime, which
executes a Map in program order and is therefore the definitional reference (the companion oracle paper,
\cite{oracle}, makes this law and its cross-backend enforcement precise):
\begin{enumerate}[leftmargin=1.5em,itemsep=1pt]
\item the fused \emph{one} kernel,
\item the na\"ive \emph{$N$}-kernel pipeline that materializes $N{-}1$ device intermediates,
\item a plain-Rust reference.
\end{enumerate}
The device morphism table (\code{apply\_un\_dev}) is a line-for-line mirror of the host table
(\code{apply\_un}) --- which is \emph{why} fused equals reference bit-exactly, not merely within
tolerance. The committed tests exercise the two patterns the llama.cpp fusion audit named plus a stress
case (Table~\ref{tab:fuse}). The SwiGLU tail and the 3-op chain are \emph{byte-identical} fused-versus-
na\"ive; a 7-op mixed chain runs in one launch and matches its host oracle. And on the real engine
pattern that motivated \S\ref{sec:pain}, the DAG path fuses the per-GDN-layer chain
$\mathrm{rms\_norm}(x)\cdot w\cdot\mathrm{silu}(g)$ from \textbf{9 launches to 5}, bit-exact
($\text{max}=2.98\times10^{-8}$) against the unfused per-op path --- draining the launch storm directly.
\begin{table}[h]
\centering\small
\setlength{\tabcolsep}{6pt}
\begin{tabular}{@{}lll@{}}
\toprule
\textbf{Chain} & \textbf{Launches (na\"ive $\to$ fused)} & \textbf{Result vs na\"ive / reference} \\
\midrule
$\mathrm{silu}(a)\cdot b$ (SwiGLU tail) & $2 \to 1$ & byte-identical / $\le 4.8\mathrm{e}{-}7$ \\
$\mathrm{silu}(a\cdot w + b)$ & $3 \to 1$ & byte-identical / $\le 4.8\mathrm{e}{-}7$ \\
7-op mixed chain (all families) & $7 \to 1$ & matches host oracle \\
$\mathrm{rms\_norm}(x)\cdot w\cdot\mathrm{silu}(g)$ (GDN, DAG) & $9 \to 5$ & bit-exact ($2.98\mathrm{e}{-}8$) \\
\bottomrule
\end{tabular}
\caption{The fusion gate. Fused one-kernel $=$ na\"ive $N$-kernel $=$ plain-Rust reference on the CPU
runtime. The launch-count drop is the fusion; the byte-identity is the guarantee that it changed nothing
but the launch count.}
\label{tab:fuse}
\end{table}
\section{Why this is the right base}
Fusibility here is a property of the operation's \emph{type}, computed once, not a fact a matcher
rediscovers for each new op combination. Adding an op adds one algebra variant and its device mirror; it
adds \emph{zero} fusion rules. And the fusion composes \emph{with} the peak kernels rather than competing
with them: the hand-tuned reduction and matmul kernels stay, and the Map regions attach to them as
read-prologue and write-epilogue at the fence. That is the whole point of making Reduce a first-class
fence rather than something a fusion pass must be taught to stop at. The elementwise fabric of an entire
model --- the SwiGLU tails, the norm epilogues, the GDN gating --- collapses to a minimal set of kernels
by one identity, and every step is proven to change only the launch count, never the result.
\begin{thebibliography}{9}
\small
\bibitem{umbrella} Hanzo AI. \emph{One Source, Every Backend: The hanzo-kernel DSL.} Paper 07 (umbrella).
\bibitem{fuse} Hanzo AI. \emph{Fusion Is Composition: Kernel Fusion as the Functor Law.} Paper 07.1 (this paper).
\bibitem{tune} Hanzo AI. \emph{The Schedule Is a Value: Comptime-Specialized Autotuning.} Paper 07.2.
\bibitem{island} Hanzo AI. \emph{Intrinsic Islands: An Oracle-Anchored Escape Hatch.} Paper 07.3.
\bibitem{oracle} Hanzo AI. \emph{One Oracle, Every Backend: The CPU Bit-Exactness Law.} Paper 07.4.
\bibitem{mmq} Hanzo AI. \emph{When the Product Won't Collapse: An int8 Tensor-Core MMQ Negative Result.} Paper 07.5.
\bibitem{method} Hanzo AI. \emph{Verify, Then Delete: Collapsing a Kernel Cartesian Product.} Paper 07.6.
\end{thebibliography}
\end{document}
hanzo-int8-mmq-negative-result.pdf
BIN
View File
Binary file not shown.
hanzo-int8-mmq-negative-result.tex
+139
View File
@@ -0,0 +1,139 @@
\documentclass[11pt]{article}
\input{shared/kernel-preamble.tex}
\title{\textbf{When the Product Won't Collapse:\\ A Negative Result on int8 Tensor-Core MMQ}}
\author{Hanzo AI --- ML Systems}
\date{hanzo-kernel Paper 07.5 --- part of the \emph{One Source, Every Backend} series}
\begin{document}
\maketitle
\begin{abstract}
A single-source kernel DSL collapses the backends$\times$ops$\times$types product for the decode matvecs,
norms, RoPE, attention, and the whole elementwise fabric of a model. This paper reports, in full, the one
operation that resists: block-quantized integer matrix--matrix multiply on the tensor cores (MMQ), the
kernel that governs prefill throughput. The result is negative and structural, and we report it as such
because it is more useful than a success story. Through the portable high-level WMMA (cooperative-matrix)
API, MMQ is capped at roughly one-tenth of the hand-tuned kernel --- not for want of tuning, but because
the API's accumulator fragment is \emph{opaque}, forcing the per-32-block quantization scale to be applied
via a store$\to$synchronize$\to$reload round-trip through shared memory once per block. The proof that
this serialization is the wall, not occupancy, is that a \emph{better-tiled} version is \emph{slower}. The
higher-ceiling route --- manual \code{mma.sync} with an exposed fragment-to-lane map, so the scale is
applied in register --- is viable but per-backend by nature (the DSL's int8 matrix-core intrinsic is
CUDA-only), so it does not collapse the product; it writes one more backend-specific kernel, in Rust.
The honest conclusion: MMQ stays hand-written behind benchmark gates, held by an intrinsic island.
\end{abstract}
\section{The target: the prefill gap is an MMQ gap}
\label{sec:target}
Decode (generating one token at a time) is memory-bound and, on this stack, at parity with the hand-tuned
baseline on every backend. Prefill (processing a prompt) is compute-bound, and there the DSL-composed
engine trails llama.cpp by a consistent margin (dense $0.6$--$0.9\x$). That gap has a single name:
tensor-core MMQ. llama.cpp quantizes the activation to \code{q8\_1} (int8) and contracts block-quantized
weights (Q8\_0, Q4\_K) against it on the hardware \emph{matrix cores} --- the units that do a small
matrix multiply as one instruction. Closing prefill means expressing that kernel well. The question this
paper answers is whether it can be expressed \emph{once}, in the DSL, and lowered to the matrix cores of
every backend --- the same collapse the rest of the series achieves for other ops.
\section{The obstruction: an opaque accumulator}
\label{sec:obstruction}
The DSL exposes two routes to the int8 matrix cores, both registered by the lowering engine on NVIDIA:
\begin{itemize}[leftmargin=1.4em,itemsep=2pt]
\item \textbf{High-level WMMA} (\code{cmma::Matrix}, i8$\to$i32 at $16{\times}16{\times}16$): fragments are
loaded and stored automatically, one tile at a time. Ergonomic and portable --- but the fragment layout
is \emph{opaque} to the programmer.
\item \textbf{Manual MMA} (\code{cmma::MmaDefinition}, \code{mma.sync} at $16{\times}8{\times}32$): the
fragment element-to-lane map \emph{is} exposed (\code{position\_of\_nth}), so values can be manipulated in
register.
\end{itemize}
Now the arithmetic MMQ actually needs. For a Q8\_0 weight against a q8\_1 activation, the contraction is
\[
\code{out}[m,n]\;=\;\sum_{\text{32-blocks }kb}\big(\code{xs}[m,kb]\cdot \code{wd}[n,kb]\big)\cdot
\sum_{k\in kb}\code{xq}[m,k]\cdot\code{wq}[n,k],
\]
where the inner \code{int8} dot is the tensor-core op and the per-block \code{f32} scale
$\code{xs}\cdot\code{wd}$ is applied \emph{once per 32-element block}. This is not a kernel choice: int8
accumulation is only valid \emph{within} one constant-scale block, so the scale must be applied at every
block boundary, inherent to the format.
Here is the collision. With the high-level WMMA API the accumulator fragment is opaque, so the only place
to apply the per-block scale is to \code{cmma::store} the integer accumulator to shared memory,
synchronize, reload it as f32, and scale --- \textbf{once per 32-block}. That store$\to$sync$\to$reload is
a full block-wide serialization inserted into the inner loop. It is the wall.
\section{The proof: better tiling is slower}
\label{sec:proof}
A cap could be an artifact of a na\"ive kernel structure --- too few warps, poor operand reuse --- rather
than the API. To separate ``API limit'' from ``na\"ive structure,'' we built the \emph{fair} ceiling: a
$32{\times}64$ output tile per block, 8 warps, with A and B staged to shared memory once per 32-block and
reused across the tile (A across four N-subtiles, B across two M-subtiles). Occupancy and operand reuse
are now fixed; the \emph{only} remaining WMMA-API tax is exactly the per-block scale round-trip through
the opaque fragment (Listing~\ref{lst:tiled}).
\begin{lstlisting}[language=Rust,caption={The fair high-level-WMMA ceiling (\code{mmq\_q8\_wmma\_blk}). A
$32{\times}64$ tile, 8 warps, operands staged once per block and reused. Everything that would help a
na\"ive kernel is already done; the residual cost is the opaque-fragment scale round-trip. Note the
target gate: the DSL's int8 matrix-core intrinsic is CUDA-only.},label={lst:tiled}]
#[kernel(targets(cuda), unchecked)] // <-- CUDA-ONLY: int8 MMA is not portable
pub fn mmq_q8_wmma_blk(xq: &Array<i8>, xs: &Array<f32>, wq: &Array<i8>, wd: &Array<f32>,
out: &mut Array<f32>, #[comptime] m: usize, #[comptime] n: usize,
#[comptime] k: usize) {
let (warp, lane) = (UNIT_POS as usize / 32, UNIT_POS as usize % 32); // 8 warps, 2x4 subtiles
// ... stage A,B to shared once per 32-block, reuse across the tile ...
// ... per block: cmma::execute (i8->i32), then store->sync->reload->scale (the wall) ...
}
\end{lstlisting}
The measurement is the result. On GB10/Blackwell, the high-level WMMA MMQ reaches only
$\sim$\textbf{7.6 TFLOP/s}, roughly \textbf{one-tenth} of the hand-tuned kernel --- and the better-tiled
version above is \emph{slower} than the na\"ive one.\footnote{The structural argument (opaque fragment
$\Rightarrow$ per-block shared-memory scale round-trip; manual MMA with an exposed lane map as the
higher-ceiling path) is documented in the committed \code{mmq.rs} probe kernels. The TFLOP/s figure and
the tiled-is-slower comparison are session measurements on spark (GB10); the probe commit records CUDA
verify/bench as pending, so we report them as measured, not as a committed test.} That is the diagnostic
signature of a \emph{sync-bound} kernel: adding tiling adds output per block, which adds
scale-round-trips, which adds synchronization --- so ``more efficient tiling'' buys more of the exact cost
that dominates. If the kernel were occupancy- or compute-bound, better tiling would help; it hurts, which
proves the bottleneck is the serialization the opaque fragment forces.
\section{The viable route is per-backend anyway}
\label{sec:viable}
The higher-ceiling path is manual MMA (\code{mma.sync}), where the fragment element-to-lane map is
exposed, so the block scale can be applied \emph{in register} with no shared-memory round-trip --- exactly
the serialization removed. But that path is per-backend by nature: the DSL's int8 matrix-core intrinsic
is CUDA-only (the probe kernels are literally \code{\#[kernel(targets(cuda))]}, as in
Listing~\ref{lst:tiled}). Writing it therefore does \emph{not} collapse the product; it produces one more
backend-specific kernel, just authored in Rust instead of CUDA C++. The portable-API collapse and the
high-ceiling implementation are, for this one operation, mutually exclusive.
\section{The honest conclusion, and the seam that holds it}
\label{sec:conclusion}
\textbf{Block-quantized MMQ stays hand-written behind benchmark gates.} This is a boundary of the
``one source, every backend'' program, not a refutation of it. The decode matvecs, the norms, RoPE,
attention, GDN, and the entire elementwise fabric --- the bulk of the hand-written kernel ledger and the
whole of the divergence-bug surface (\cite{method}, \cite{oracle}) --- do collapse. The one op that
resists is exactly the one the series already built a seam for: an \emph{intrinsic island} (\cite{island})
holds the hand-written tensor-core arm behind a normative, CPU-checkable \code{default}, so even the
un-collapsed kernel keeps one signature, one oracle, one test. A negative result that names its boundary
precisely --- \emph{which} op, \emph{why} (the opaque fragment), \emph{how you know} (better tiling is
slower), and \emph{what contains it} (the island) --- is worth more than a portability claim that quietly
excludes the hardest kernel. We report it in full so the boundary is known and the seam is used.
\begin{thebibliography}{9}
\small
\bibitem{umbrella} Hanzo AI. \emph{One Source, Every Backend: The hanzo-kernel DSL.} Paper 07 (umbrella).
\bibitem{fuse} Hanzo AI. \emph{Fusion Is Composition: Kernel Fusion as the Functor Law.} Paper 07.1.
\bibitem{tune} Hanzo AI. \emph{The Schedule Is a Value: Comptime-Specialized Autotuning.} Paper 07.2.
\bibitem{island} Hanzo AI. \emph{Intrinsic Islands: An Oracle-Anchored Escape Hatch.} Paper 07.3.
\bibitem{oracle} Hanzo AI. \emph{One Oracle, Every Backend: The CPU Bit-Exactness Law.} Paper 07.4.
\bibitem{mmq} Hanzo AI. \emph{When the Product Won't Collapse: An int8 Tensor-Core MMQ Negative Result.} Paper 07.5 (this paper).
\bibitem{method} Hanzo AI. \emph{Verify, Then Delete: Collapsing a Kernel Cartesian Product.} Paper 07.6.
\end{thebibliography}
\end{document}
hanzo-intrinsic-islands.pdf
BIN
View File
Binary file not shown.
hanzo-intrinsic-islands.tex
+191
View File
@@ -0,0 +1,191 @@
\documentclass[11pt]{article}
\input{shared/kernel-preamble.tex}
\title{\textbf{Intrinsic Islands:\\ An Oracle-Anchored Escape Hatch for a Portable Kernel}}
\author{Hanzo AI --- ML Systems}
\date{hanzo-kernel Paper 07.3 --- part of the \emph{One Source, Every Backend} series}
\begin{document}
\maketitle
\begin{abstract}
Every portable kernel DSL eventually meets a loop that wants a target-specific instruction: the hardware
4-wide int8 dot on one backend, a subgroup shuffle on another. The tempting response is to fork a whole
per-backend kernel for that one line --- which rebuilds, in miniature, the very Cartesian product the DSL
exists to collapse. This paper presents a disciplined alternative: the \emph{intrinsic island}, a scoped,
target-gated block written \emph{inside} the one kernel. The kernel keeps one signature, one launch
contract, one reference, and one test; only the marked inner expression differs per target. The design
rests on two rules. First, selection is a \emph{compile-time} \code{match} on a target tag the launch
layer fills in and the caller never sees, so exactly one arm is lowered per backend and nothing else
compiles. Second, the \code{default} arm is \emph{normative}: it \emph{is} the reference semantics, the
CPU always takes it, and every accelerated arm must equal it byte-for-byte on the CPU oracle before it can
ship. An island is thus an escape hatch that cannot escape correctness --- the accelerated instruction is
never trusted, only proven equivalent. We show the mechanism, the one-kernel-one-oracle test, and why an
island beats a fork on exactly the axis that matters.
\end{abstract}
\section{The pain: one line wants a special instruction}
\label{sec:pain}
A quantized matvec spends its inner loop on a 4-wide integer dot product: four \code{int8} weights times
four \code{int8} activations, summed to an \code{i32}. On NVIDIA, AMD, Apple, and Vulkan hardware this is
a \emph{single instruction} (\dpa{} / \code{OpSDotAccSat}) --- an order of magnitude cheaper than four
scalar multiply-adds. A portable kernel that emits four scalar MACs leaves that instruction on the table;
one that emits the hardware dot cannot run on a target without it (the CPU reference has no such op).
The reflex is to fork: a hand-written \code{matvec\_cuda} that uses the intrinsic, and a portable
\code{matvec\_fallback} that does not. But a fork for one inner line drags the entire kernel with it ---
its signature, its launch bounds, its indexing, its reference, its test --- all now duplicated and free to
drift. That is the Cartesian product (backends $\times$ ops) reappearing at the granularity of a single
expression. We want the hardware instruction where it exists and the portable arithmetic where it does
not, \emph{without} a second copy of the kernel.
\section{The idea: gate the expression, not the kernel}
\label{sec:idea}
An island moves the fork inward, from the whole kernel to the one expression that actually differs
(Listing~\ref{lst:island}). The kernel is written once. At the one spot that needs a target-specific
idiom, an \code{island!} block lists accelerated arms by backend and a mandatory \code{default} arm. The
rest of the kernel --- the loop, the indexing, the accumulation, the bounds --- is shared verbatim.
\begin{lstlisting}[language=Rust,caption={An island in a \dpa{} matvec. The accelerated arm issues the
hardware 4-wide int8 dot; the \code{default} arm --- the normative oracle, taken by the CPU and any
unlisted target --- computes the identical contraction with portable scalar lane MACs. Same math, two
idioms.},label={lst:island}]
let dp = island! {
// Accelerated: the hardware 4-wide int8 dot (dp4a / OpSDot), one instruction.
cuda | rocm | metal | vulkan => { w.dot(x) }
// NORMATIVE oracle: the identical integer contraction, portable scalar lane MACs.
default => { w[0]*x[0] + w[1]*x[1] + w[2]*x[2] + w[3]*x[3] }
};
\end{lstlisting}
Two properties make this safe rather than merely convenient, and both are structural.
\paragraph{The branch site names backends, not a target.} The author writes \code{cuda | rocm | metal |
vulkan => ...}; they never write ``if I am on CUDA.'' The \emph{tag} that selects the arm is filled in
elsewhere (\S\ref{sec:mech}), so a kernel author cannot accidentally couple the math to a runtime string.
\paragraph{\code{default} is the definition.} The accelerated arms are optimizations \emph{of} the
\code{default} arm. \code{default} is not a fallback in the sense of ``lower quality if unlucky'' --- it
is the \emph{normative} semantics that the accelerated arms must reproduce. This is the inversion that
makes an island trustworthy: the portable arithmetic is the truth, and the fast instruction has to match
it.
\section{The mechanism: target selection at compile time}
\label{sec:mech}
The subtlety is that the lowering engine compiles the same target-agnostic IR for every backend; at
kernel-expansion time there is no ``which backend am I'' string available inside the kernel, only a
capability set. But the backend brand \emph{is} known on the \emph{host}, from the runtime's name
(\code{"cpu"}, \code{"cuda"}, \code{"hip"}, \code{"wgpu<spirv>"}, \code{"wgpu<msl>"}). So an island is
selected by a compile-time tag threaded in from the host (Listing~\ref{lst:target}):
\begin{lstlisting}[language=Rust,caption={The target tag is derived on the host from the runtime name and
threaded into the kernel as a \code{\#[comptime]} parameter. The \code{\#[kernel]} macro rewrites
\code{island!\{...\}} into a comptime \code{match target}. Unknown names fall to \code{Cpu} --- fail
safe, never a wrong accelerated arm.},label={lst:target}]
pub enum Target { Cpu, Cuda, Rocm, Metal, Vulkan, WebGpu } // #[comptime] tag: Eq + Hash + Copy
pub fn from_runtime_name(name: &str) -> Target {
match name {
"cuda" => Target::Cuda,
"hip" => Target::Rocm,
n if n.contains("spirv") => Target::Vulkan,
n if n.contains("msl") => Target::Metal,
n if n.starts_with("wgpu") => Target::WebGpu,
_ => Target::Cpu, // fail safe -> the oracle arm
}
}
\end{lstlisting}
Because \code{target} is a \code{\#[comptime]} value, the \code{match} it drives is a comptime match over
a constant scrutinee: \textbf{only the selected arm is lowered.} Every backend compiles exactly its own
arm and nothing else --- there is no runtime branch, no dead code for other targets in the shipped kernel.
The tag is kept out of the authoring surface twice over: the branch site lists backends (not a target),
and the host launch wrapper fills the tag (\code{Target::of::<R>}), so callers pass only data. One
signature, one launch, every backend.
\section{The oracle contract, tested}
\label{sec:oracle}
Here is the rule that makes an island an escape hatch that cannot escape correctness. The CPU runtime is
\emph{always} tagged \code{Target::Cpu} --- a variant no accelerated arm ever claims --- so on the CPU
every island resolves to \code{default}. The CPU runtime is therefore the single bit-exact oracle for
every island (the general form of this law, and its cross-backend enforcement, is the companion oracle
paper, \cite{oracle}). The test states the contract directly and mechanically: run the kernel on the CPU
with each accelerated tag \emph{forced}, and it must equal the \code{default} result byte-for-byte
(Listing~\ref{lst:oracletest}).
\begin{lstlisting}[language=Rust,caption={The oracle-equivalence test. Every accelerated arm, executed on
the CPU runtime with its tag forced, must be byte-identical to the \code{default} arm and to the
plain-Rust reference. An accelerated instruction is never trusted --- it is proven equal to the normative
arm.},label={lst:oracletest}]
// The default arm (CPU is tagged Cpu -> default) must equal the plain-Rust oracle.
let default = matvec_island_run_with::<CpuRuntime>(&c, &wq, &xq, &wd, rows, k, Target::Cpu);
assert_eq!(bits(&default), bits(&matvec_island_ref(&wq, &xq, &wd, rows, k)));
// Every accelerated arm, forced onto the CPU runtime, must equal the oracle byte-for-byte.
for accel in [Target::Cuda, Target::Rocm, Target::Metal, Target::Vulkan, Target::WebGpu] {
let out = matvec_island_run_with::<CpuRuntime>(&c, &wq, &xq, &wd, rows, k, accel);
assert_eq!(bits(&out), bits(&default), "island arm {accel:?} diverged from default");
}
\end{lstlisting}
A separate test asserts the \emph{production} launch --- the one that derives the tag from the runtime and
hides it --- resolves to the oracle arm on the CPU without the caller naming a target. So the property is
proven both under forced tags (every arm equals \code{default}) and under the real launch path (the
caller-facing API does the right thing by default).
\section{Why an island beats a fork}
\label{sec:vsfork}
The comparison is not aesthetic; it is about what can drift and what can be verified
(Table~\ref{tab:vsfork}). A fork gives two kernels: two signatures to keep in sync, two references to
trust, two tests to keep honest, and no mechanical guarantee that the fast kernel computes what the
portable one does --- the equivalence lives only in a human's intention. An island gives one kernel: one
signature, one reference (the \code{default} arm), one test that \emph{proves} every accelerated arm
equals it. The accelerated instruction is contained, not trusted.
\begin{table}[h]
\centering\small
\setlength{\tabcolsep}{7pt}
\begin{tabular}{@{}lll@{}}
\toprule
& \textbf{Fork (per-backend kernel)} & \textbf{Island (scoped block)} \\
\midrule
Signatures & one per backend, drift-prone & one, shared \\
Reference & one per fork, trust-based & one \code{default} arm, normative \\
Equivalence & human intention only & proven byte-for-byte on the CPU oracle \\
What differs & the whole kernel & the one gated expression \\
New backend & another full kernel & another arm, or falls to \code{default} \\
\bottomrule
\end{tabular}
\caption{An island localizes the difference to the expression that actually differs and anchors it to a
CPU-checkable normative arm. A fork localizes nothing and anchors nothing.}
\label{tab:vsfork}
\end{table}
\section{The designed seam}
An island is not a licence to sprinkle intrinsics; it is the \emph{one} sanctioned place a per-backend
idiom may live, precisely because it comes with the oracle contract attached. That matters most for the
operation this series shows does \emph{not} collapse into portable form: block-quantized tensor-core MMQ,
whose viable implementation is per-backend by nature (\cite{mmq}). Even there, an island holds the
hand-written tensor-core arm behind a normative, CPU-checkable \code{default}, so the one un-collapsed
kernel still keeps one signature, one oracle, one test. The escape hatch exists so that the rare
irreducibly-per-backend kernel does not force a return to forks --- it is the boundary condition of the
whole ``one source, every backend'' program, made safe.
\begin{thebibliography}{9}
\small
\bibitem{umbrella} Hanzo AI. \emph{One Source, Every Backend: The hanzo-kernel DSL.} Paper 07 (umbrella).
\bibitem{fuse} Hanzo AI. \emph{Fusion Is Composition: Kernel Fusion as the Functor Law.} Paper 07.1.
\bibitem{tune} Hanzo AI. \emph{The Schedule Is a Value: Comptime-Specialized Autotuning.} Paper 07.2.
\bibitem{island} Hanzo AI. \emph{Intrinsic Islands: An Oracle-Anchored Escape Hatch.} Paper 07.3 (this paper).
\bibitem{oracle} Hanzo AI. \emph{One Oracle, Every Backend: The CPU Bit-Exactness Law.} Paper 07.4.
\bibitem{mmq} Hanzo AI. \emph{When the Product Won't Collapse: An int8 Tensor-Core MMQ Negative Result.} Paper 07.5.
\bibitem{method} Hanzo AI. \emph{Verify, Then Delete: Collapsing a Kernel Cartesian Product.} Paper 07.6.
\end{thebibliography}
\end{document}
hanzo-kernel-dsl.pdf
BIN
View File
Binary file not shown.
hanzo-kernel-dsl.tex
+150 -596
View File
@@ -1,170 +1,38 @@
\documentclass[11pt]{article}
\usepackage[margin=1in]{geometry}
\usepackage{booktabs}
\usepackage{amsmath}
\usepackage{amssymb}
\usepackage{graphicx}
\usepackage{xcolor}
\usepackage{hyperref}
\hypersetup{colorlinks=true, linkcolor=blue, urlcolor=blue, citecolor=blue}
\usepackage{enumitem}
\usepackage{listings}
\usepackage{array}
\input{shared/kernel-preamble.tex}
\definecolor{kw}{rgb}{0.13,0.29,0.53}
\definecolor{cmt}{rgb}{0.40,0.40,0.40}
\definecolor{str}{rgb}{0.16,0.50,0.16}
\lstset{
basicstyle=\ttfamily\footnotesize,
keywordstyle=\color{kw}\bfseries,
commentstyle=\color{cmt}\itshape,
stringstyle=\color{str},
showstringspaces=false,
breaklines=true,
columns=fullflexible,
frame=single,
framesep=4pt,
}
\lstdefinelanguage{Rust}{
morekeywords={fn,let,mut,pub,struct,impl,trait,enum,match,if,else,for,while,loop,return,use,mod,crate,self,super,async,await,where,type,const,static,unsafe,ref,move,dyn,as,true,false},
morecomment=[l]{//}, morecomment=[s]{/*}{*/}, morestring=[b]", sensitive=true,
}
\newcommand{\x}{\ensuremath{\times}}
\newcommand{\dpa}{\texttt{dp4a}}
\newcommand{\hk}{\texttt{hanzo-kernel}}
\newcommand{\code}[1]{\texttt{#1}}
% Absorb minor prose overfulls without rivers; keep line-breaking sane.
\tolerance=800
\emergencystretch=3em
\title{\textbf{Fusion Is Composition:\\
A Category-Theoretic Kernel DSL that Collapses the\\
Backends\,\x\,Ops\,\x\,Types Product}}
\title{\textbf{One Source, Every Backend:\\ The \hk{} DSL --- Thesis and Guided Tour}}
\author{Hanzo AI --- ML Systems}
\date{Paper 07 --- companion to Paper 06 (\emph{When Does a Kernel Win Surface?})}
\date{Paper 07 (umbrella) --- companion to Paper 06 (\emph{When Does a Kernel Win Surface?})}
\begin{document}
\maketitle
\begin{abstract}
A multi-backend inference engine hand-writes the same $\sim$20 operations, specialized over $\sim$22
GGUF quantization formats, once per GPU backend --- a Cartesian product of five kernel trees and
$\sim$52{,}000 lines that is not merely a maintenance cost but the breeding ground for
divergence bugs, because one operation with five implementations has five numeric behaviors. We
present \hk{}, a first-party Rust GPU-kernel DSL that collapses this product to a sum, built on one
organizing idea: \textbf{fusion is composition} --- the functor law $\mathrm{map}(g)\circ\mathrm{map}(f)
= \mathrm{map}(g\circ f)$ made mechanical, with Reduce/Scan as algebraic fences, so ``can these two
ops fuse?'' is a total function of the op's \emph{class}, not a pattern matcher or a cost model.
On this base sit three mechanisms that each attack one axis of the product: \emph{comptime-specialized
autotune} (one kernel source, its schedule exposed as compile-time knobs the tuner picks per device,
every variant byte-identical); \emph{intrinsic islands} (scoped target-gated blocks inside one kernel
whose \code{default} arm is the normative oracle); and \emph{a CPU runtime as the bit-exact oracle}
for every backend. We report the evidence: 10/10 op families bit-exact against the CPU oracle;
DSL$\to$Vulkan and DSL$\to$ROCm dispatch bit-exact on real hardware (rms\_norm on gfx1151,
$\text{max\_rel}=2.43\times10^{-7}$); a single-source \dpa{} matvec at \textbf{118\% of the hand-tuned
kernel} it replaces (196~GB/s, gfx1151); a shipped Vulkan migration where the DSL twin was
\textbf{7.1\x{} faster} than the hand-written shader \emph{and} bit-exact; and a nightly three-backend
GPU CI (self-hosted CUDA/ROCm/Metal) that enforces the bit-exactness law. We also report a
scientifically load-bearing \emph{negative} result: int8 tensor-core MMQ via the high-level WMMA API is
structurally sync-bound (the opaque accumulator fragment forces a store$\to$sync$\to$reload$\to$scale per
32-element block; a better-tiled version is \emph{slower}), so block-quantized GEMM stays hand-written
behind benchmark gates --- which is exactly what islands are designed to hold. None of Triton, XLA,
torch.compile, TVM, Halide, Burn, or Mojo models fusion as morphism composition; we state precisely
what is and is not novel.
The kernel layer of a multi-backend LLM engine is a Cartesian product: the same $\sim$20 operations,
specialized over $\sim$22 quantization formats, hand-written once per GPU backend --- five kernel trees
and $\sim$52{,}000 lines, and the breeding ground for the divergence bugs that cost the most to chase,
because one operation with five implementations has five numeric behaviors. This umbrella paper states the
thesis that collapses the product to a sum and gives a guided tour of the six companion papers that make
it concrete. The thesis rests on one idea --- \textbf{fusion is composition}, the functor law made
mechanical --- extended by three mechanisms that each attack one axis of the product: comptime-specialized
autotune (the schedule as a value, not a fork), intrinsic islands (a per-backend idiom held behind a
normative oracle arm), and a CPU runtime that is the bit-exact reference for every backend. We describe the
facade that ties them together (a first-party namespace over the CubeCL lowering engine, plugged in as a
code generator rather than a second runtime), summarize the evidence --- portability that costs nothing, a
shipped migration that improved performance, an honest structural negative result --- and place the work
against Triton, XLA, torch.compile, TVM, Halide, Burn, and Mojo. Each fact in this series lives in exactly
one paper; here we connect them.
\end{abstract}
\section{Introduction}
\section{The thesis: a product that should be a sum}
\label{sec:thesis}
The kernel layer of a multi-backend LLM engine is a product. Table~\ref{tab:dup} is the current
inventory in the Hanzo ML and engine repositories: five hand-written kernel trees --- CUDA C++, HIP
C++, Metal Shading Language, GLSL$\to$SPIR-V, and the engine's own CUDA --- carrying on the order of
$52{,}000$ lines of hand-written kernel arithmetic. The same $\sim$20 operations (matvec, matmul,
flash attention, RMSNorm, RoPE, softmax, the Gated-DeltaNet linear-attention ops) are re-implemented in
each, and within each they are specialized over $\sim$22 GGUF quantization formats. Three axes ---
\emph{backends}, \emph{ops}, \emph{types} --- are multiplied by hand.
That product is where the correctness bugs breed, and Paper~06 catalogs the class concretely: a
Qwen3-8B model that generates coherently on Metal collapses into token repetition under CUDA
flash-attention, byte-identical weights and sampler, because \emph{the same operation has a different
kernel per backend and therefore a different numeric behavior}. Uniform numeric behavior across
backends is not something four or five copies can be tested \emph{into}; it is a structural property of
having exactly one implementation. The engineering goal, then, is not ``write the kernels more
carefully'' but ``write each operation once and lower it everywhere,'' turning the product into a sum:
\[
\underbrace{\text{backends} \x \text{types} \x \text{ops}}_{\text{hand-written today}}
\;\longrightarrow\;
\underbrace{\text{backends} + \text{types} + \text{ops}}_{\text{one source, lowered}}.
\]
\begin{table}[h]
\centering
\begin{tabular}{llrr}
\toprule
\textbf{Tree} & \textbf{Language} & \textbf{Files} & \textbf{Lines} \\
\midrule
CUDA (ML) & CUDA C++ & -- & 13{,}123 \\
ROCm & HIP C++ & 19 & 6{,}725 \\
Metal & Metal Shading Language & -- & 3{,}530 \\
Vulkan & GLSL $\rightarrow$ SPIR-V & 121 & 7{,}105 \\
CUDA (engine) & CUDA C++ & -- & 21{,}862 \\
\midrule
\textbf{Total} & 5 trees & & \textbf{52{,}345} \\
\bottomrule
\end{tabular}
\caption{The collapse ledger: hand-written kernel arithmetic across the five backend trees.\protect\footnotemark{}
The DSL's job is to drain this to op-specs $+$ type-specs $+$ backend lowerings.}
\label{tab:dup}
\end{table}
\footnotetext{These are hand-written, non-generated, non-vendored kernel lines --- the \emph{deletable}
duplication. A na\"ive \code{wc -l} over all \code{*.cu/*.cuh/*.hip/*.hpp/*.metal/*.comp} is roughly
$2\x$ larger ($\sim$112k) because it also counts C++ headers, a \code{target/package/} build-artifact
copy, machine-generated codebook tables (e.g. the i-quant grids), and vendored upstream kernels
(FlashInfer, Marlin, GPTQ) that are neither hand-written duplication nor a DSL collapse target.}
\paragraph{Thesis and contributions.}
This paper argues that the collapse is not an engineering slog of retyping kernels into a portable
syntax but follows from one algebraic idea, applied three ways. The idea is that \textbf{kernel fusion
is morphism composition}: an elementwise kernel is the lift of a scalar function to every array index,
and the functor law $\mathrm{map}(g)\circ\mathrm{map}(f)=\mathrm{map}(g\circ f)$ read right-to-left
\emph{is} fusion --- one kernel computing the composite, no intermediate array. Reduce/Scan are the
\emph{fences} where the pattern changes, decidable from the op's class, not from a pattern matcher.
Concretely, we contribute:
\begin{enumerate}[leftmargin=1.6em,itemsep=1pt]
\item \textbf{Fusion as composition} (\S\ref{sec:fuse}): an op algebra in which legal fusion is a total
function of op \emph{class}, one fused-interpreter kernel that lowers any composed chain to a single
launch with zero intermediates, and its generalization from a line to an arbitrary DAG --- all
bit-exact-gated on the CPU (fused one-kernel $=$ na\"ive $N$-kernel $=$ plain-Rust reference).
\item \textbf{Comptime-specialized autotune} (\S\ref{sec:tune}): the schedule of a single-source kernel
exposed as \code{\#[comptime]} knobs the engine monomorphizes over; a small variant set timed once per
$(\text{device},\text{op},\text{shape})$ and cached to XDG; every variant byte-identical, so the tuner
picks speed without touching numerics --- the ``one op, $N$ device forks'' bug class killed.
\item \textbf{Intrinsic islands} (\S\ref{sec:island}): a target-gated block \emph{inside} one kernel
whose \code{default} arm is normative and whose accelerated arms are proven byte-for-byte equal to it
on the CPU oracle --- the disciplined seam for the rare instruction that truly needs a per-backend idiom.
\item \textbf{The CPU oracle law and its CI} (\S\ref{sec:oracle}): the CPU runtime is the single
bit-exact reference for every backend, enforced nightly on self-hosted CUDA/ROCm/Metal runners.
\item \textbf{Evidence, positive and negative} (\S\ref{sec:eval}--\ref{sec:negative}): cross-backend
bit-exactness, a portable kernel that \emph{beats} the hand-tuned one, a shipped migration that improved
performance, and an honest structural negative result on int8 tensor-core MMQ.
\end{enumerate}
We adopt CubeCL as the lowering engine and wrap it in a first-party facade so that kernel source names
only \code{hanzo\_kernel::*}; the string \code{cubecl} appears in exactly one line of one
\code{Cargo.toml}. That is ``values, not places'': CubeCL is the implementation \emph{value};
\hk{} is the stable \emph{namespace} we build against, upgradable or forkable without editing a kernel.
\section{Design: the facade and the three axes}
\label{sec:design}
\subsection{Three orthogonal axes, one source}
The kernel code braids three independent concerns. Pulled apart, each is written once and they compose
(Table~\ref{tab:axes}). The transformation is not new in spirit --- ROCm already decomplected two of
the three axes, writing its kernels as \code{qmatvec\_core<WTYPE>}: one contraction, the quantization
type injected as a compile-time template parameter. The DSL is exactly that pattern with the
\emph{backend} as the other compile-time axis. We are generalizing a decomplect already trusted in
production, not inventing an untested one.
Table~\ref{tab:axes} is the disease. A multi-backend engine braids three independent concerns into one
hand-multiplied product: the operations (the math), the quantization types (the bit layout), and the
backends (the lowering). Written by hand, that is $\sim$20 $\times$ 22 $\times$ 5 kernels; in the Hanzo
repositories it is $52{,}345$ lines across five trees and five shading languages (the ledger and its
audit live in the collapse-method paper, \cite{method}).
\begin{table}[h]
\centering
@@ -177,29 +45,39 @@ Quant types & $\sim$22 & the bit layout: Q4\_K, Q6\_K, Q8\_0, the IQ codebooks \
Backends & 5 & the lowering: threads, memory, tensor cores, per shading language \\
\bottomrule
\end{tabular}
\caption{The three axes the hand-written product multiplies. The DSL makes each a summand.}
\caption{Three orthogonal axes, multiplied by hand today. The DSL's job is to make each a summand.}
\label{tab:axes}
\end{table}
\subsection{A code generator, not a second runtime}
\hk{} lowers a \code{\#[kernel]} to SPIR-V for Vulkan, to CUDA/PTX for CUDA, to MSL for Metal, and
through cubecl-hip for ROCm. Crucially it is a \emph{code generator}, not a competing GPU runtime: the
engine's Vulkan path is already \code{include\_bytes!(<name>.spv)} $\to$ \code{create\_shader\_module}
$\to$ dispatch, and CubeCL emits exactly those bytes. A DSL RMSNorm dumped and inspected is a
\code{GLCompute} entry point with buffers on descriptor bindings $0,1,2$ --- byte-for-byte the shape
the engine's hand-written \code{.comp}$\to$\code{.spv} shaders already take. So the DSL plugs into each
backend's \emph{existing} dispatch machinery; there is no dual runtime and no cross-runtime copy. The
one bounded integration step is a three-line SPIR-V post-pass (rename the entry point to \code{main};
strip CubeCL's unused info-uniform at binding~3 via \code{spirv-opt --remove-unused-interface-variables};
dispatch), automated by a reusable emitter.
That product is not merely a maintenance cost. It is where correctness bugs breed: a single quantized
model that generates coherent text on one backend and collapses into token repetition on another, from
byte-identical weights, because the same operation has a different kernel --- and therefore a different
numeric behavior --- per backend (the bug and its cure are the subject of the oracle paper,
\cite{oracle}). Uniform numeric behavior across backends cannot be tested \emph{into} five copies; it is a
structural property of having one implementation. So the goal is to turn the product into a sum:
\[
\underbrace{\text{backends} \x \text{types} \x \text{ops}}_{\text{hand-written today}}
\;\longrightarrow\;
\underbrace{\text{backends} + \text{types} + \text{ops}}_{\text{one source, lowered}}.
\]
This is a decomplect we already trust in miniature: the ROCm tree writes \code{qmatvec\_core<WTYPE>} --- one
contraction, the quant type injected as a compile-time parameter. The DSL generalizes exactly that pattern
with the \emph{backend} as the other compile-time axis.
\subsection{The kernel-authoring surface}
Listing~\ref{lst:names} shows the facade's first-principles vocabulary: \code{\#[kernel]} says what it
is (a GPU kernel), not \code{cube} (a brand); \code{Grid} and \code{Block} name the launch shape
directly. A kernel's only import is \code{use hanzo\_kernel::prelude::*;}.
\section{The facade: a namespace, and a code generator}
\label{sec:facade}
\begin{lstlisting}[language=Rust,caption={The Q8\_0 matvec: one source, no backend in it. The decode
(\code{wd*wq}) and the contraction are the arithmetic each backend hand-writes today.},label={lst:names}]
\hk{} is a first-party Rust DSL: a kernel is written once, in Rust, and lowered to CUDA, ROCm, Vulkan, and
Metal. It is a \emph{facade} over the CubeCL lowering engine --- kernel source names only
\code{hanzo\_kernel::*}, and the string \code{cubecl} appears in exactly one line of one \code{Cargo.toml}.
That is ``values, not places'': CubeCL is the implementation \emph{value}; \hk{} is the stable
\emph{namespace} we build against, upgradable or forkable without editing a kernel. The authoring surface
uses first-principles names --- \code{\#[kernel]} says what a thing is (a GPU kernel), not \code{cube} (a
brand); \code{Grid} and \code{Block} name the launch shape directly (Listing~\ref{lst:matvec}).
\begin{lstlisting}[language=Rust,caption={A quantized matvec: one source, no backend in it. The decode
(\code{wd*wq}) and the contraction are the arithmetic each backend hand-writes today. This one function
lowers to CUDA, ROCm, Vulkan, Metal, and CPU.},label={lst:matvec}]
use hanzo_kernel::prelude::*;
#[kernel(targets(cuda, metal, vulkan, webgpu, cpu), unchecked)]
@@ -215,446 +93,122 @@ pub fn matvec_q8<F: Float>(wd: &Array<F>, wq: &Array<i32>, x: &Array<F>,
}
\end{lstlisting}
The perf primitives the hand-tuned kernels rely on are all reachable and each lowers to the native
instruction: \code{dot} on a vectorized \code{Line<i8>} $\to$ \dpa{}\,/\,\code{OpSDotAccSat};
\code{cmma} $\to$ WMMA\,/\,cooperative-matrix\,/\,simdgroup; \code{SharedMemory}, the \code{plane\_*}
subgroup reductions, \code{Atomic}, and barriers. The DSL is a superset of what we express by hand, so
nothing about portability forces a performance downgrade --- a claim we make good on in \S\ref{sec:eval}
and qualify honestly in \S\ref{sec:negative}.
Crucially the DSL is a \emph{code generator}, not a competing GPU runtime. The engine's Vulkan path is
already ``load SPIR-V bytes, create a shader module, dispatch,'' and the DSL \emph{emits} exactly those
bytes: a compiled kernel is a \code{GLCompute} entry point with buffers on descriptor bindings $0,1,2$ ---
byte-for-byte the shape the hand-written \code{.comp}$\to$\code{.spv} shaders already take. So the DSL plugs
into each backend's \emph{existing} dispatch machinery; there is no dual runtime and no cross-runtime copy.
The one bounded integration step is a three-line SPIR-V post-pass (rename the entry point to \code{main};
strip the engine's unused info-uniform; dispatch). The perf primitives the hand-tuned kernels rely on are
all reachable and each lowers to the native instruction --- \code{dot} on a vectorized \code{Line<i8>}
$\to$ \dpa{}, \code{cmma} $\to$ tensor cores, \code{SharedMemory} and the \code{plane\_*} subgroup
reductions --- so portability does not, by construction, force a performance downgrade.
\section{Mechanism I: fusion is composition}
\label{sec:fuse}
\section{A guided tour of the mechanisms}
\label{sec:tour}
\subsection{The functor law, read right-to-left}
An elementwise (pointwise) kernel is the lift of a scalar morphism $f:F\to F$ applied at every index.
Pointwise lifts obey the functor law
\[
\mathrm{map}(g)\circ\mathrm{map}(f) \;=\; \mathrm{map}(g\circ f).
\]
Read left-to-right this is two kernels with a materialized intermediate array between them; read
right-to-left it is \emph{one} kernel computing the composed scalar function with \emph{no}
intermediate. Kernel fusion is nothing but that law read right-to-left. It is exactly the transducer
story: \code{map(f)} and \code{map(g)} are transformations independent of the collection, so you compose
the \emph{transformations} and apply the composite to the collection once. Here the collection is a GPU
array and ``once'' is one launch. Categorically the Maps are morphisms $\mathrm{Array}\to\mathrm{Array}$
and composition of morphisms \emph{is} fusion --- there is no separate ``fusion pass'' concept, only the
associativity of $\circ$.
The collapse is not a syntactic retyping of kernels; it follows from one algebraic idea applied three
ways. Each is developed in full in its own companion paper; here is the map.
\subsection{The class decides fusibility}
The whole surface is a tiny op algebra, and legal fusion falls out of the op's \emph{class}, not from a
pattern matcher over concrete ops:
\begin{itemize}[leftmargin=1.4em,itemsep=1pt]
\item \textbf{Map} (pointwise morphism): output element $i$ depends only on input(s) at index $i$.
Shape-preserving, index-local. Maps compose freely: any run of Maps is one Map (the functor law). This
is the fusible family --- \code{Add Sub Mul Div Neg Recip Abs Exp Rsqrt Sigmoid Tanh Silu Gelu}.
\item \textbf{Reduce\,/\,Scan} (contraction): output $i$ depends on a whole row/axis
(\code{sum}, \code{max}, \code{mean} $\to$ RMSNorm, softmax, matmul). Not index-local, so it cannot be
absorbed into an adjacent Map by the same lift; it needs cross-lane communication. A Reduce is therefore
a \emph{fence}: a preceding Map run fuses into its input \emph{read} (prologue) and a following Map run
fuses into its output \emph{write} (epilogue), but the reduction is the seam where the kernel pattern
changes.
\end{itemize}
The decision ``can these two adjacent ops fuse?'' reduces to
$\mathrm{fusible}(a,b) := \mathrm{is\_map}(a)\wedge\mathrm{is\_map}(b)$, plus prologue/epilogue
attachment at a Reduce's ports. Nothing else. The fusion pass is one fold: walk the chain, compose
adjacent Maps by substituting the previous expression for the current-value leaf (\code{compose\_cur},
which \emph{is} $g[\,\code{Cur}:=f\,]$), and cut a new region at every Reduce.
\paragraph{Fusion is composition (\cite{fuse}).} An elementwise kernel is the lift of a scalar function to
every index, and such lifts obey the functor law $\mathrm{map}(g)\circ\mathrm{map}(f)=\mathrm{map}(g\circ
f)$. Read right-to-left, that law \emph{is} fusion: one kernel computing the composite, no intermediate
array. Reduce and Scan are algebraic \emph{fences} where the pattern changes, decidable from the op's
class rather than a pattern matcher. This is the theory core: the fusibility of two ops is a total function
of their type, one fused-interpreter kernel lowers any composed chain to a single launch, and the same
algebra lifts from a line to an arbitrary DAG --- draining, for instance, a per-layer Gated-DeltaNet chain
from 9 kernel launches to 5, bit-exact.
\subsection{One kernel interprets the whole chain}
The fused Map region is lowered to a small linear ``op program'' carried as a \code{\#[comptime]} value
and walked in a comptime-\emph{unrolled} loop, so each op becomes straight-line arithmetic in one kernel
body --- no runtime interpreter, no intermediate arrays, no extra launches
(Listing~\ref{lst:fused}). It is not a hand-written fixed kernel: any linear elementwise chain the
builder produces runs through this one \code{fused\_interp}, whatever its length.
\paragraph{The schedule is a value (\cite{tune}).} The last place a backend forks a kernel is the schedule
--- tile size, vector width, unroll, block dimension. A fork is exactly what reintroduces the divergence
bug. Autotune folds the fork into the one source by exposing the schedule as compile-time knobs the engine
monomorphizes; a small variant set is timed per device and the winner cached to disk. Because the knobs
never touch the arithmetic, every variant is byte-identical --- the tuner buys speed and can never spend
correctness.
\begin{lstlisting}[language=Rust,caption={The one fused kernel. It reads the primary input once,
threads it through the comptime program (unrolled to straight-line code), writes once. Every
\code{Instr} is inlined; one launch, zero intermediates.},label={lst:fused}]
#[kernel(targets(cuda, metal, vulkan, webgpu, cpu), unchecked)]
pub fn fused_interp<F: Float>(x: &Array<F>, sides: &Sequence<Array<F>>,
out: &mut Array<F>, #[comptime] prog: Program) {
let i = ABSOLUTE_POS;
if i < out.len() {
let mut acc = x[i];
#[unroll]
for step in 0..prog.instrs.len() {
match comptime!(prog.instrs[step]) {
Instr::Un(op) => acc = apply_un_dev::<F>(op, acc),
Instr::BinIn(op, slot) => acc = apply_bin_dev::<F>(op, acc, sides.index(slot)[i]),
Instr::BinConst(op, b) => { let c = comptime!(f32::from_bits(b));
acc = apply_bin_dev::<F>(op, acc, F::new(c)) }
}
}
out[i] = acc;
}
}
\end{lstlisting}
\paragraph{Intrinsic islands (\cite{island}).} When one inner line wants a target-specific instruction, you
do not fork the kernel; you write a scoped, target-gated block \emph{inside} it. Selection is a
compile-time \code{match} on a tag the launch layer fills in, so exactly one arm is lowered per backend.
The \code{default} arm is \emph{normative} --- the CPU always takes it, and every accelerated arm is proven
byte-for-byte equal to it on the CPU oracle before it ships. An escape hatch that cannot escape correctness.
The ergonomic surface is a builder that reads like the composition it denotes:
\code{Fuse::new(a).\allowbreak{}mul(w).\allowbreak{}add(b).\allowbreak{}silu().\allowbreak{}run(c)} appends one \code{Instr} per call, compiles
them to one \code{Program}, and dispatches exactly one \code{fused\_interp}. Because the \code{Program}
rides as a comptime value, the constant leaf is stored as \code{f32::to\_bits} so the program stays
\code{Eq + Hash} --- the lowering engine keys the compiled kernel on it, one compiled kernel per distinct
fused chain.
\paragraph{One oracle, every backend (\cite{oracle}).} The foundation under all three: the CPU runtime is
the single bit-exact reference for every backend. Every op's device code mirrors a plain-Rust reference
line-for-line; because the CPU runs pointwise ops in program order, its result is the definition, and a GPU
differs only by bounded reduction-order rounding. The law is enforced nightly by a three-backend GPU CI on
self-hosted NVIDIA, AMD, and Apple machines --- so ``the same op has the same numbers everywhere'' is a
checked invariant, and the repetition collapse becomes an unreachable state rather than a bug to fix.
\subsection{From a line to a DAG}
A real forward pass is not a line: a value fans out (used by several ops) and a binary fans in (reads two
independently computed values). \code{silu(a)*b + a} reuses $a$ twice and a linear chain cannot express
it in one region; a DAG can, and it still fuses to one kernel. The same algebra lifts from a path to a
graph: a \code{NodeId}-taping newtype records an arbitrary pointwise DAG as a \code{Vec<Node>} (operator
overloads push nodes instead of computing); a flood-fill partitions maximal Map-only connected regions,
cutting every edge into a fence --- the identical \code{Class} predicate applied to graph edges instead
of list-adjacency; and one generic \code{dag\_interp} kernel interprets a per-region DAG program over a
per-element slot array (fan-out $=$ a slot read by several later instructions; fan-in $=$ a binary
reading two slots), generalizing \code{fused\_interp}'s single accumulator to $n$ accumulators. The
concrete payoff is measured on the real engine pattern: the per-GDN-layer chain
\code{rms\_norm(x)*w*silu(g)} fuses from \textbf{9 launches to 5}, bit-exact
($\text{max}=2.98\times10^{-8}$) against the na\"ive per-op path --- directly draining the prefill
launch storm these hybrid models suffer.
\paragraph{Where the product won't collapse (\cite{mmq}).} Honesty is part of the thesis. Block-quantized
tensor-core MMQ --- the kernel that governs prefill --- does \emph{not} collapse through the portable API:
its accumulator fragment is opaque, forcing a per-block scale round-trip that caps it near one-tenth of
hand-tuned (proven sync-bound because better tiling is \emph{slower}). The viable route is per-backend
anyway. So MMQ stays hand-written behind exactly the island seam built for it --- a named boundary, not a
quiet exclusion.
\subsection{The bit-exact gate}
Every fused chain is checked three ways on the CPU runtime: (i) the fused one kernel, (ii) the na\"ive
$N$-kernel pipeline that materializes $N{-}1$ device intermediates, and (iii) a plain-Rust reference ---
all byte-for-byte identical. The tests exercise the two patterns the llama.cpp fusion audit named plus a
stress case: the SwiGLU tail \code{silu(a)*b} (2 kernels $\to$ 1) and \code{silu(a*w+b)} (3 $\to$ 1) are
\emph{byte-identical} fused-versus-na\"ive with reference error $\le 4.8\times10^{-7}$, and a 7-op mixed
chain (\code{gelu((|2a{-}1|+b)/sigmoid(w)).tanh()}) runs in one launch, matching its host oracle. The
device morphism table \code{apply\_un\_dev} is a line-for-line mirror of the host \code{apply\_un} ---
which is \emph{why} fused equals oracle bit-exactly, not merely within tolerance.
\paragraph{Verify, then delete (\cite{method}).} The migration method ties it together: a hand-written
kernel is retired only when its single-source twin is both bit-exact against the oracle and at least as
fast on a kernel-only benchmark, cell by cell through a per-\code{(op,type,backend)} registry, so the
engine is faster-or-equal at every step. Cells have already flipped on \code{main} --- including a Vulkan
\code{add\_rmsnorm} twin that was \emph{7.1$\times$ faster} than the shader it replaced and bit-exact.
\paragraph{Why this is the right base.} Fusibility is a property of the op's \emph{type}, computed once,
not rediscovered by a matcher for each new op combination. Adding an op adds one algebra variant and its
device mirror; it does not add fusion rules. The reduction and matmul kernels stay hand-tuned and the
Map regions attach to them as prologue/epilogue --- the fusion composes \emph{with} the peak kernels
rather than competing with them.
\section{Evidence, in one breath}
\label{sec:evidence}
\section{Mechanism II: comptime-specialized autotune}
\label{sec:tune}
\subsection{The collapse, closed on the schedule}
The DSL already gives one source lowered to every backend. Autotune closes the loop on the last place a
backend used to fork a kernel: the \emph{schedule}. Hand-tuning a schedule per device --- tile size,
vector width, unroll factor, cube dimension --- is the ``same op, $N$ forks, $N$ drifting numeric
behaviors'' bug class in a subtler dress. The fix is to expose the schedule as \code{\#[comptime]} knobs
on the \emph{one} source. Each knob value is a distinct compiled kernel (CubeCL monomorphizes on comptime
state), so a small variant set is enumerated, timed once per $(\text{device},\text{op},\text{shape})$,
and the winner cached. There are no forks: the source is one; only the chosen comptime tuple differs per
device. Crucially the knobs never change the numerics --- the per-row reduction is the same fold in the
same order for every schedule --- so \emph{every variant is byte-identical}, and the tuner buys speed
without ever risking a numeric change.
\subsection{The algebra: Variant, Tuned, Tuner}
Three orthogonal values (Listing~\ref{lst:tune}): a \code{Variant} is a named schedule --- a closure
\code{f(iters) -> (output, ms)} that both benchmarks (many iters) and executes (one iter); \code{Tuned}
is the fluent builder collecting variants for one $(\text{op},\text{shape-key})$; \code{Tuner} is the
cache, an in-memory $(\text{device},\text{op},\text{key})\to\text{winner}$ map backed by a per-device
on-disk file. The \code{Tuner} is a plain value over an explicit cache root, so it is fully unit-testable
offline --- no GPU, no env, no globals; the process-wide default roots at the XDG cache dir and the disk
format is one TSV line per key (no serde, no new dependency).
\begin{lstlisting}[language=Rust,caption={The production surface. One source per op, autotuned per
device, no forks. On a cache hit only the winner launches (once, to produce its output); the losers are
never run.},label={lst:tune}]
let y = Tuned::new("rms_norm", format!("rows={rows},n={n}"))
.variant("b64_r1", |it| rms_norm_tuned_bench(&c, x, w, rows, n, eps, 1, 64, it))
.variant("b128_r1", |it| rms_norm_tuned_bench(&c, x, w, rows, n, eps, 1, 128, it))
.variant("b256_r1", |it| rms_norm_tuned_bench(&c, x, w, rows, n, eps, 1, 256, it))
.variant("b64_r2", |it| rms_norm_tuned_bench(&c, x, w, rows, n, eps, 2, 64, it))
.run(&c); // times each once per (device,op,shape), caches the winner, thereafter runs it
\end{lstlisting}
\subsection{The variant tables and their proof}
Two ops ship autotuned (Table~\ref{tab:variants}). Both variant sets are proven on the CPU runtime:
every schedule is bit-identical to the oracle (rms\_norm: byte-for-byte equal, the fold order is
invariant; \dpa{}: byte-identical across variants, within $2\times10^{-3}$ f32-reorder of the
per-element oracle), and the cache mechanism is exercised directly --- a miss benchmarks all four and
records the winner, a second call is a hit that benchmarks zero, a fresh \code{Tuner} over the same
directory reloads the winner from disk and skips timing, and a changed variant set is treated as a miss
and re-tuned rather than trusting a stale name. Timing on the CPU proves the select/cache machinery; a
GPU run additionally makes the \emph{winner} meaningful.
\begin{table}[h]
\centering\small
\setlength{\tabcolsep}{5pt}
\begin{tabular}{@{}lll@{}}
\toprule
\textbf{Op} & \textbf{Knobs} & \textbf{Variants (all byte-identical)} \\
\midrule
\code{rms\_norm} & cube dim $\{64,128,256\}\times$ unroll $\{1,2\}$ & \code{b64\_r1 b128\_r1 b256\_r1 b64\_r2} \\
\code{matvec\_q8\_dp4a} & cube dim $\{64,128,256\}\times$ vec.\ width $\{1,2,4,8\}$ & \code{b64\_v1 b64\_v4 b128\_v2 b256\_v8} \\
\bottomrule
\end{tabular}
\caption{The autotune variant sets (\hk{} 0.2.17). A knob value is a distinct compiled kernel;
the tuner caches the winner per $(\text{device},\text{op},\text{shape})$ to
\code{\$XDG\_CACHE\_HOME/hanzo-kernel/autotune/<device>.tsv}.}
\label{tab:variants}
\end{table}
\subsection{Why first-party}
CubeCL ships an autotuner; it was studied and declined for three concrete reasons, not NIH. It is
re-exported only from \code{cubecl\_runtime}, so adopting it would force a \emph{second} engine crate
into a facade architected so only \code{Cargo.toml} names the engine; its persistent cache lives under
\code{target/} rather than the XDG path this layer requires; and its generic-associated-type input
machinery is a large surface for ``pick the fastest of $N$ named closures.'' The concepts above are
deliberately the same as CubeCL's, sized to this crate and pointed at XDG --- a thin analog, not a new
philosophy.
\section{Mechanism III: intrinsic islands}
\label{sec:island}
\subsection{Killing the whole-kernel escape hatch}
Sometimes one inner loop needs a target-specific instruction or idiom. The wrong response is to fork a
parallel per-backend kernel --- that is the product, back again. The right response is a scoped,
target-gated \code{island!} block \emph{inside} the one DSL kernel (Listing~\ref{lst:island}). The kernel
keeps one signature, one launch contract, one CPU oracle, one test. The \code{default} arm is the
normative semantics; accelerated arms are optimizations that must equal it.
\begin{lstlisting}[language=Rust,caption={An island in a \dpa{} matvec. The accelerated arm issues the
hardware 4-wide int8 dot; \code{default} --- the normative oracle, taken by CPU and any unlisted target
--- computes the identical contraction with portable scalar lane MACs.},label={lst:island}]
let dp = island! {
cuda | rocm | metal | vulkan => { w.dot(x) } // hardware dp4a / OpSDot
default => { w[0]*x[0] + w[1]*x[1] + w[2]*x[2] + w[3]*x[3] } // NORMATIVE oracle
};
\end{lstlisting}
\subsection{The mechanism: target selection at comptime}
CubeCL compiles the same target-agnostic IR for every backend; at \code{\#[cube]} expand time there is no
``which backend am I'' string, only a capability set. But the backend brand \emph{is} known on the host,
via the runtime name. So an island is selected by a \code{\#[comptime] target: Target} tag the launch
layer derives from the runtime (\code{Target::of::<R>}) and threads in; the \code{\#[kernel]} macro
rewrites \code{island!\{...\}} into a comptime \code{match target}. A comptime match over a const
scrutinee lowers \emph{only} the selected arm --- every backend compiles exactly its own arm, nothing
else. The tag is out of the authoring surface twice over: the branch site names backends, not a target,
and the launch wrapper fills the tag, so callers pass only data.
\subsection{The oracle contract, tested}
The CPU runtime is always tagged \code{Target::Cpu} --- a variant no accelerated arm claims --- so on CPU
every island resolves to \code{default}. This makes the CPU runtime the single bit-exact oracle for every
island, and the test proves it directly: run the kernel on CPU with each accelerated tag \emph{forced},
and it must equal the \code{default} result byte-for-byte. The shipped test drives all five accelerated
tags (\code{Cuda, Rocm, Metal, Vulkan, WebGpu}) through the CPU runtime and asserts each is byte-identical
to \code{default} and to the plain-Rust reference, and separately asserts the production launch (which
derives the tag from the runtime) resolves to the oracle on CPU without the caller naming a target. An
accelerated arm is thus never trusted; it is proven equivalent to the normative arm before it can ship.
This is the designed seam for the one case the rest of the paper cannot make portable ---
block-quantized tensor-core GEMM (\S\ref{sec:negative}).
\section{The CPU runtime as bit-exact oracle, and its CI}
\label{sec:oracle}
The three mechanisms share one foundation: \textbf{the CPU runtime is the single bit-exact reference for
every backend.} It is not an accelerator target --- it is the oracle. Every op ships with a plain-Rust
reference whose scalar semantics the device code mirrors line-for-line; the fusion engine's device
morphism table mirrors its host table; every island's \code{default} arm \emph{is} the CPU semantics; and
every autotune variant is checked against the oracle. Because a Map's arithmetic is index-local and the
CPU runs it in program order, the CPU result is the definition, and a GPU backend differs only by f32
re-association where a reduction reorders a sum --- a difference bounded and explained, not a bug to
hunt. The op library --- \code{rms\_norm}, \code{layer\_norm}, \code{rope} (both NeoX and GPT-J
conventions), \code{sdpa} (online/flash softmax), the GDN \code{conv1d}/\code{gating}/\code{scan}, the
quant matvecs, and the fusion chains --- is gated this way across 36 CPU tests.
This law is enforced continuously. A nightly GitHub Actions workflow
(\code{.github/\allowbreak{}workflows/\allowbreak{}gpu-tests.yml}) runs a three-backend matrix on self-hosted boxes --- one runner
label per backend: \code{cuda} $\to$ spark (NVIDIA GB10, sm\_121), \code{rocm} $\to$ evo (AMD Strix Halo
gfx1151, ROCm~7.x), \code{metal} $\to$ dbc (Apple M4~Max) --- each running
\code{cargo test -p hanzo-ml} \code{-{}-features <backend>} \code{-{}-test-threads 1} behind a
GPU-guard step. Serial
execution and the guard keep a single GPU consumer at a time per backend, which is not fastidiousness but
necessity: concurrent ROCm inits race in the kernel-fusion driver and wedge it until reboot. The
practical result is that ``the same op has the same numbers on every backend'' is a checked invariant,
not an aspiration.
\section{Evaluation}
\label{sec:eval}
We separate what is a committed, reproducible test (bit-exactness) from what is a measured benchmark on
named hardware.\footnote{Bit-exactness figures are enforced by the committed CPU/GPU tests cited.
Throughput and latency figures are session benchmarks on the named machines: evo (AMD Strix Halo,
gfx1151, RADV/Vulkan and ROCm~7.13), spark (NVIDIA GB10, sm\_121), dbc (Apple M4~Max, Metal).
Benchmarks are kernel-only (buffers created once, $N$ dispatches enqueued back-to-back, one final
sync).}
\subsection{Cross-backend bit-exactness}
The op library is bit-exact against the CPU oracle for every family it targets, and --- the part that
matters --- the DSL \emph{dispatches} bit-exact through real backend runtimes, not merely compiles
(Table~\ref{tab:bitexact}). The ROCm figure is the falsifiable end-to-end proof that the
cubecl-hip$\to$ROCm~7.13 seam works: the same \code{rms\_norm} source lowered through cubecl-hip and
dispatched on gfx1151 (device query via the R0600 bindings $+$ HIP-RTC module compile $+$ launch $+$
readback) matches the CPU oracle to $\text{max\_rel}=2.43\times10^{-7}$, well under the $2\times10^{-3}$
gate --- pure f32-reorder noise, not a decode defect.
\begin{table}[h]
\centering\small
\setlength{\tabcolsep}{5pt}
\begin{tabular}{@{}lll@{}}
\toprule
\textbf{Path} & \textbf{Op(s)} & \textbf{Result} \\
\midrule
CPU oracle & rms/layer\_norm, rope$\times$2, sdpa, gdn, quant, fuse & 10/10 bit-exact (36 tests) \\
DSL $\to$ Vulkan (gfx1151) & rms\_norm, softmax, matvec, sdpa & bit-exact, real dispatch \\
DSL $\to$ Metal (M4 Max) & rms\_norm, layer\_norm, rope, sdpa & bit-exact ($2.35\mathrm{e}{-}7$) \\
DSL $\to$ ROCm (gfx1151) & rms\_norm & $\text{max\_rel}=2.43\mathrm{e}{-}7$ \\
\bottomrule
\end{tabular}
\caption{The same source, dispatched through each backend's real runtime, matches the CPU oracle. The
\code{sdpa} online-softmax being a \emph{single} implementation is the structural cure for the 8B
repetition-collapse: with one attention kernel there is no flash-vs-eager-vs-Metal fork to diverge.}
\label{tab:bitexact}
\end{table}
\subsection{Portability costs nothing: the DSL matches and beats hand-tuned}
The central risk of any portable DSL is a performance downgrade. It does not materialize on the
bandwidth-bound kernels that dominate decode. A tuned single-source \dpa{} Q8 matvec ---
\code{Vector<i8,4>}$\to$\code{OpSDot}, one thread block per output row with its threads striding the
$k/4$ groups so adjacent threads read adjacent 4-byte packed weights (coalesced 256-byte transactions),
partials tree-reduced in shared memory --- measures \textbf{196~GB/s} on evo (gfx1151, Vulkan, $8192^2$,
cache-busting), which is \textbf{118\% of the $\sim$166~GB/s hand-tuned \dpa{} kernel} it would replace
and $\sim$77\% of the $\sim$256~GB/s DRAM roofline, bit-exact against the CPU oracle. The na\"ive
one-thread-per-row kernel that started at $\sim$17~GB/s was memory-\emph{uncoalesced}, not DSL-limited;
the fix was two standard decomplected steps --- coalescing and reduction --- both expressed in the DSL.
The DSL reaches the same hardware integer-dot instruction and, with standard coalescing, \emph{beats} the
hand-written code.
\subsection{A migration that improved performance}
Retirement is real, not hypothetical, and it can be a strict win. Two hand-written Vulkan kernels were
drained on \code{main}: the DSL block-per-row \code{rms\_norm} (270~GB/s) and \code{softmax} (268~GB/s)
replaced na\"ive hand-written \code{.comp} shaders, each now the production kernel with its \code{.comp}
deleted, dispatching bit-exact through the engine's real Vulkan pipeline. The sharpest example is the
fused residual-add\,+\,RMSNorm: the DSL \code{add\_rmsnorm} twin ran \textbf{7.1\x{} faster} than the
incumbent hand-written shader ($166.5\to23.4$~ms at $4096^2$ on RADV) \emph{and} was bit-exact ---
$s=x{+}\text{res}$ byte-identical, $y$ within $4.8\times10^{-7}$. A ``uniformity'' change turned out to
also be a large speedup, because the incumbent was na\"ive and the DSL block kernel is coalesced. The
migration is faster-or-equal at every step by construction: a kernel is retired only when its DSL twin is
bit-exact \emph{and} at least as fast.
\subsection{The discipline cuts both ways}
The gate also \emph{declines} migrations. Where a hand-tuned kernel is already optimal --- the Vulkan
quant matvec, Metal's \code{simd\_sum} reductions, the fused \code{rope\_norm} epilogue --- agents built
bit-exact DSL twins, measured, and correctly did \emph{not} swap: a uniformity-only change with zero
speedup and real regression risk is churn, not migration. ``One source'' is the goal; a bit-exact,
no-faster twin is kept as the portable fallback for backends that lack the tuned version, and the tuned
peak stays the specialized fast path. We only ever add coverage and correctness.
\subsection{The ROCm-7 seam, as an engineering note}
Shipping the DSL on AMD required fixing two independent upstream defects without bumping CubeCL (0.10 is
the newest published; no version escapes either). First, cubecl-hip-sys selects its FFI bindings from
\code{hipconfig --version}; ROCm~7.13 reports patch 99004, past the newest shipped binding, so upstream
emitted a feature with no bindings module and every symbol went undefined. The fix is a vendored
\code{[patch.crates-io]} fork whose \code{build.rs} clamps an unknown-newer patch to the latest shipped
binding --- ABI-safe because AMD versions every changed symbol (R0600) and 7.13 still maps
\code{hipGetDeviceProperties} to that revision. Second, cubecl-hip's server uses a \code{MultiStream}
gated behind a \code{std}-only cfg the \code{hip}/\code{rocm} features do not pull, so the ROCm feature is
declared \code{rocm = ["cubecl/hip", "cubecl/rocm", "cubecl/std"]}. Both are recorded because a portable
DSL lives or dies on the least-loved backend building cleanly.
\section{A negative result: int8 tensor-core MMQ}
\label{sec:negative}
Not everything collapses. The universal prefill gap Paper~06 documents (dense $0.6$--$0.9\x$ of
llama.cpp on every backend) is a tensor-core-MMQ gap: llama quantizes the activation to q8\_1 (int8) and
contracts block-quantized weights against it on the matrix cores. We asked whether that kernel, too, can
be written once in \hk{} and lowered to the hardware int8 matrix cores instead of hand-written four times
(CUDA \code{fast\_mmq}, Vulkan \code{mul\_mm\_q8\_mmq}, ROCm \code{qmmq}, Metal simdgroup). The answer,
honestly, is \emph{not through the portable API}, and the reason is structural rather than a missing
optimization.
\subsection{The structural cap}
CubeCL exposes two routes to the int8 matrix cores. The high-level WMMA API (\code{cmma::Matrix},
$16{\times}16{\times}16$, i8$\to$i32) loads and stores accumulator fragments automatically, one tile at a
time --- but \emph{the fragment layout is opaque}. Block-quantized GEMM needs a per-32-block f32 scale
applied to the integer accumulator, and int8 accumulation is only valid \emph{within} one constant-scale
32-block (this is inherent to the format, not a kernel choice). With an opaque fragment the only place to
apply the scale is a store$\to$sync$\to$reload round-trip through shared memory, \emph{once per block}.
That serialization is the wall. Measured on GB10/Blackwell, the high-level WMMA MMQ reaches only
$\sim$7.6~TFLOP/s, roughly one-tenth of the hand-tuned kernel, and --- the diagnostic that proves it is
sync-bound, not occupancy-bound --- a \emph{better-tiled} version (a $32{\times}64$ output tile, 8 warps,
operands staged to shared memory once per block and reused across the tile, which fixes occupancy and
operand reuse) is \emph{slower}, because the only remaining tax is exactly the per-block scale round-trip
through the opaque fragment, and better tiling adds more of them.\footnote{The structural argument (opaque
fragment $\Rightarrow$ per-block shared-memory scale round-trip; manual MMA with exposed lane map as the
higher-ceiling path) is documented in the committed \code{mmq.rs} probe. The TFLOP/s figures and the
tiled-is-slower comparison are session measurements on spark (GB10); the probe commit records CUDA
verify/bench as pending, so we report them as measured, not as a committed test.}
\subsection{The viable route is per-backend anyway}
The higher-ceiling path is manual MMA (\code{cmma::MmaDefinition}, \code{mma.sync} at
$16{\times}8{\times}32$), where the fragment element-to-lane map \emph{is} exposed, so the block scale can
be applied in register with no shared-memory round-trip. But that path is per-backend by nature ---
CubeCL 0.10's int8 MMA is CUDA-only (the probe kernels are literally \code{\#[kernel(targets(cuda))]}) ---
so writing it does not collapse the product; it writes one more backend-specific kernel, just in Rust.
The honest conclusion is therefore: \textbf{block-quantized MMQ stays hand-written behind benchmark
gates.} This is not a defeat of the thesis but a boundary the thesis anticipates. The decode matvecs,
norms, RoPE, attention, GDN, and the entire elementwise fabric --- the bulk of the ledger and the whole
of the divergence-bug surface --- collapse. The one op that resists is exactly the one for which
\S\ref{sec:island} built the seam: an island holds the hand-written tensor-core arm behind a normative,
CPU-checkable \code{default}, so even the un-collapsed kernel keeps one signature, one oracle, one test.
The negative result names the boundary and the mechanism that contains it.
\section{The collapse ledger}
\label{sec:ledger}
Table~\ref{tab:dup} is the denominator; this section is the protocol that drains it. Roughly 85\% of the
$52{,}345$ lines are a deletable target --- the hand-written arithmetic duplicated across backends, as
opposed to the genuinely peak, genuinely per-backend kernels (the MMQ GEMMs of \S\ref{sec:negative}) that
stay behind islands. Draining is never all-or-nothing and never a downgrade, because it is governed by a
\textbf{verify-then-delete protocol}:
\begin{enumerate}[leftmargin=1.6em,itemsep=1pt]
\item \textbf{Coverage first.} The DSL fills the entire $22\,\text{types}\x5\,\text{backends}$ matrix from
one source. Cells that were missing or buggy become uniform immediately.
\item \textbf{The gate.} A hand-tuned kernel is retired for its DSL twin \emph{only} when the twin is
(a)~bit-exact against the shared CPU oracle and (b)~within performance noise on a kernel-only benchmark
--- in practice $\ge 97\%$ of the incumbent's throughput. The number decides, per type, per backend.
\item \textbf{Delete on parity.} As each twin passes, its hand-written copies are deleted and a
per-$(\text{op},\text{type},\text{backend})$ registry flips that cell to the DSL, keyed by the gate
result. The product drains to a sum one measured cell at a time, with a working, faster-or-equal engine
at every intermediate state.
\end{enumerate}
Executed instances already on \code{main}: the Vulkan \code{rms\_norm}, \code{softmax}, and
\code{add\_rmsnorm} shaders deleted in favor of bit-exact DSL twins (\S\ref{sec:eval}, the last a 7.1\x{}
win); the ROCm-7 seam opened so gfx1151 dispatches from the same source; and the DAG fusion cutting the
GDN launch storm 9$\to$5. Instances correctly \emph{not} executed: the Vulkan quant matvec and Metal
\code{simd\_sum}, where the DSL twin was bit-exact but not faster. The ledger is a scoreboard with a
gate, not a rewrite.
The claims the tour makes are backed by measurement, detailed in the companion papers: the op library is
bit-exact against the CPU oracle and dispatches bit-exact through real Vulkan, Metal, and ROCm runtimes
(rms\_norm on gfx1151 to $2.43\times10^{-7}$; \cite{oracle}); a single-source \dpa{} matvec runs at 118\%
of the hand-tuned kernel it replaces (196~GB/s on gfx1151; \cite{method}); a shipped migration was
7.1$\times$ faster \emph{and} bit-exact (\cite{method}); the autotune variants are byte-identical with a
per-device cached winner (\cite{tune}); and the one op that resists is reported in full as a structural
negative result (\cite{mmq}). Portability costs nothing where it counts, the collapse is real and
reversible, and the boundary is drawn honestly.
\section{Related work}
\label{sec:related}
Portable kernel authoring and kernel fusion are both well-trodden; the specific claim we make is narrow
and we state it precisely. \textbf{Triton} gives a tile-level Python DSL that lowers to GPU code, but
fusion is the programmer's job (you write a fused kernel) or the compiler's inside a larger stack; it has
no notion of fusion as an algebraic law over an op class. \textbf{XLA} fuses HLO graphs via
producer-consumer heuristics and a cost model; \textbf{torch.compile}/TorchInductor captures an FX graph
and fuses by a scheduler's heuristics. Both fuse by \emph{analysis over a graph}, not by a fusibility
Portable kernel authoring and kernel fusion are both well-trodden; our specific claim is narrow, and we
state it precisely. \textbf{Triton} gives a tile-level Python DSL, but fusion is the programmer's job or a
larger stack's; it has no notion of fusion as an algebraic law over an op class. \textbf{XLA} fuses HLO
graphs by producer-consumer heuristics and a cost model; \textbf{torch.compile}/TorchInductor captures a
graph and fuses by a scheduler's heuristics --- both fuse by \emph{analysis over a graph}, not by a
predicate that is a function of the op's type. \textbf{Halide} and \textbf{TVM} pioneered separating the
\emph{compute} from the \emph{schedule} and autotuning the latter --- this is the nearest neighbor of our
autotune mechanism (\S\ref{sec:tune}), and we claim no novelty for schedule/compute separation itself;
what differs is that our schedule knobs are \code{\#[comptime]} monomorphizations of one kernel proven
byte-identical, so autotuning is guaranteed not to perturb numerics. \textbf{Burn}'s \code{burn-fusion}
is the closest relative to our fusion engine: it records an op stream and JITs fused elementwise kernels
over the \emph{same} CubeCL substrate. The difference is where fusibility comes from --- \code{burn-fusion}
traces Burn tensor ops and applies its own cut/epilogue policy, coupling to Burn's tensor abstraction;
ours is the functor law over an op \emph{class}, one predicate, on our own tensor type, composing with our
hand-tuned peak kernels. \textbf{Mojo} is a language for kernels, not a fusion-as-composition model.
Across all of these, we are unaware of a system that presents kernel fusion \emph{as} morphism composition
--- the functor law $\mathrm{map}(g)\circ\mathrm{map}(f)=\mathrm{map}(g\circ f)$ as the mechanical rule
and Reduce/Scan as the algebraic fences --- with a bit-exact CPU oracle as the definitional semantics.
That framing, and the three mechanisms it organizes, is the contribution; the substrate (CubeCL) and the
individual ideas (portable IR, autotuning, oracle testing) are prior art we build on and cite.
\emph{compute} from the \emph{schedule} and autotuning the latter; this is the nearest neighbor of our
autotune mechanism (\cite{tune}), and we claim no novelty for schedule/compute separation itself --- what
differs is that our schedule knobs are compile-time monomorphizations of one kernel proven byte-identical,
so autotuning cannot perturb numerics. \textbf{Burn}'s \code{burn-fusion} is the closest relative to our
fusion engine: it records an op stream and JITs fused elementwise kernels over the \emph{same} CubeCL
substrate, but it traces Burn tensor ops with its own cut/epilogue policy, coupling to Burn's tensor
abstraction; ours is the functor law over an op \emph{class}, one predicate, on our own tensor type,
composing with our hand-tuned peak kernels. \textbf{Mojo} is a language for kernels, not a
fusion-as-composition model. Across all of these we are unaware of a system that presents kernel fusion
\emph{as} morphism composition --- the functor law as the mechanical rule, Reduce/Scan as the algebraic
fences --- anchored by a bit-exact CPU oracle as the definitional semantics. That framing, and the three
mechanisms it organizes, is the contribution; the substrate (CubeCL) and the individual ideas (portable
IR, autotuning, oracle testing) are prior art we build on and cite.
\section{Conclusion}
The kernel layer is a Cartesian product of backends, quantization types, and operations, hand-written in
five shading languages, and it is the breeding ground for the divergence bugs that cost the most to
chase. \hk{} collapses that product to a sum on one idea --- fusion is composition --- extended by three
mechanisms that each attack an axis: autotune folds the schedule fork into comptime knobs proven
byte-identical; islands hold the rare per-backend idiom behind a normative, CPU-checkable arm; and the CPU
runtime is the single bit-exact oracle every backend is measured against, nightly. The evidence is that
portability costs nothing where it counts (a single-source \dpa{} matvec at 118\% of hand-tuned, a shipped
migration 7.1\x{} faster and bit-exact), that the collapse is real and reversible (a verify-then-delete
ledger with kernels already deleted on \code{main}), and --- stated as plainly as the wins --- that one
op resists: block-quantized tensor-core MMQ is structurally sync-bound through the portable WMMA API and
stays hand-written behind exactly the island seam we built for it. What remains is a disciplined,
measured migration whose engine is faster-or-equal at every step.
five shading languages, and it is the breeding ground for the divergence bugs that cost the most to chase.
\hk{} collapses that product to a sum on one idea --- fusion is composition (\cite{fuse}) --- extended by
three mechanisms that each attack an axis: autotune folds the schedule fork into byte-identical comptime
knobs (\cite{tune}); islands hold the rare per-backend idiom behind a normative, CPU-checkable arm
(\cite{island}); and the CPU runtime is the single bit-exact oracle every backend is measured against,
nightly (\cite{oracle}). Portability costs nothing where it counts, the collapse is real and reversible via
a verify-then-delete method (\cite{method}), and --- stated as plainly as the wins --- one op resists and
stays hand-written behind an island (\cite{mmq}). Read the six companion papers for the depth; this one is
the map.
\begin{thebibliography}{9}
\small
\bibitem{umbrella} Hanzo AI. \emph{One Source, Every Backend: The hanzo-kernel DSL.} Paper 07 (this umbrella).
\bibitem{fuse} Hanzo AI. \emph{Fusion Is Composition: Kernel Fusion as the Functor Law.} Paper 07.1.
\bibitem{tune} Hanzo AI. \emph{The Schedule Is a Value: Comptime-Specialized Autotuning.} Paper 07.2.
\bibitem{island} Hanzo AI. \emph{Intrinsic Islands: An Oracle-Anchored Escape Hatch.} Paper 07.3.
\bibitem{oracle} Hanzo AI. \emph{One Oracle, Every Backend: The CPU Bit-Exactness Law.} Paper 07.4.
\bibitem{mmq} Hanzo AI. \emph{When the Product Won't Collapse: An int8 Tensor-Core MMQ Negative Result.} Paper 07.5.
\bibitem{method} Hanzo AI. \emph{Verify, Then Delete: Collapsing a Kernel Cartesian Product.} Paper 07.6.
\end{thebibliography}
\end{document}
shared/kernel-preamble.tex
+41
View File
@@ -0,0 +1,41 @@
% Shared preamble for the hanzo-kernel Paper 07 series.
% One style, one vocabulary, inputted by every paper so the series reads as one voice.
\usepackage[margin=1in]{geometry}
\usepackage{booktabs}
\usepackage{amsmath}
\usepackage{amssymb}
\usepackage{graphicx}
\usepackage{xcolor}
\usepackage{hyperref}
\hypersetup{colorlinks=true, linkcolor=blue, urlcolor=blue, citecolor=blue}
\usepackage{enumitem}
\usepackage{listings}
\usepackage{array}
\definecolor{kw}{rgb}{0.13,0.29,0.53}
\definecolor{cmt}{rgb}{0.40,0.40,0.40}
\definecolor{str}{rgb}{0.16,0.50,0.16}
\lstset{
basicstyle=\ttfamily\footnotesize,
keywordstyle=\color{kw}\bfseries,
commentstyle=\color{cmt}\itshape,
stringstyle=\color{str},
showstringspaces=false,
breaklines=true,
columns=fullflexible,
frame=single,
framesep=4pt,
}
\lstdefinelanguage{Rust}{
morekeywords={fn,let,mut,pub,struct,impl,trait,enum,match,if,else,for,while,loop,return,use,mod,crate,self,super,async,await,where,type,const,static,unsafe,ref,move,dyn,as,true,false},
morecomment=[l]{//}, morecomment=[s]{/*}{*/}, morestring=[b]", sensitive=true,
}
\newcommand{\x}{\ensuremath{\times}}
\newcommand{\dpa}{\texttt{dp4a}}
\newcommand{\hk}{\texttt{hanzo-kernel}}
\newcommand{\code}[1]{\texttt{#1}}
% Absorb minor prose overfulls without rivers; keep line-breaking sane.
\tolerance=800
\emergencystretch=3em