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

Add 6 hanzo-engine technical whitepapers (honesty-audited) (#3)

Browse Source 
Detailed engineering/research whitepapers for the hanzo native stack,
companions to the existing short product papers:
- hanzo-rocm-inference: native ROCm decode parity on Strix Halo (measured gold standard)
- hanzo-cross-backend-inference: measured AMD/NVIDIA/Apple matrix; decode parity per backend
- hanzo-native-training: native training substrate (candle-lineage autograd + hanzo quant core)
- hanzo-native-stack-thesis: unified train-and-serve synthesis
- hanzo-platform-infra-desktop: economical serving + hanzo.ai platform thesis
- hanzo-continuous-learning-privacy: on-device private continual learning

Numbers measured or labeled; upstream forks (candle, mistral.rs) and ggml prior
art attributed; 4-agent honesty review applied (no fabricated benchmarks).

Co-authored-by: Hanzo Dev <[email protected]>
This commit is contained in:
zeekay
2026-06-17 14:54:32 -07:00
committed by GitHub
co-authored by Hanzo Dev
parent e60c42f79b
commit db066d7f2e
13 changed files with 2020 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files
INDEX.md
+7 -1
View File
@@ -46,6 +46,12 @@ Auto-generated catalogue of research papers.
| `hanzo-edge-inference` | ✓ | `hanzo-edge-inference.tex` |
| `hanzo-encryption-security` | ✓ | `hanzo-encryption-security/hanzo-encryption-security.tex` |
| `hanzo-engine-ml` | ✓ | `hanzo-engine-ml.tex` |
| `hanzo-rocm-inference` | ✓ | `hanzo-rocm-inference.tex` |
| `hanzo-cross-backend-inference` | ✓ | `hanzo-cross-backend-inference.tex` |
| `hanzo-native-training` | ✓ | `hanzo-native-training.tex` |
| `hanzo-native-stack-thesis` | ✓ | `hanzo-native-stack-thesis.tex` |
| `hanzo-platform-infra-desktop` | ✓ | `hanzo-platform-infra-desktop.tex` |
| `hanzo-continuous-learning-privacy` | ✓ | `hanzo-continuous-learning-privacy.tex` |
| `hanzo-exchange-architecture` | ✓ | `hanzo-exchange-architecture/hanzo-exchange-architecture.tex` |
| `hanzo-federated-agents` | ✓ | `hanzo-federated-agents/hanzo-federated-agents.tex` |
| `hanzo-fhe-inference` | ✓ | `hanzo-fhe-inference/hanzo-fhe-inference.tex` |
@@ -175,4 +181,4 @@ Auto-generated catalogue of research papers.
| `zen4-thinking_whitepaper` | ✓ | `zen/zen4-thinking_whitepaper.tex` |
| `zen4-ultra_whitepaper` | ✓ | `zen/zen4-ultra_whitepaper.tex` |
**Total**: 169 papers, 327 PDFs compiled
**Total**: 175 papers, 333 PDFs compiled
hanzo-continuous-learning-privacy.pdf
BIN
View File
Binary file not shown.
hanzo-continuous-learning-privacy.tex
+426
View File
@@ -0,0 +1,426 @@
\documentclass[11pt]{article}
\usepackage{shared/hanzocover}
\usepackage[utf8]{inputenc}
\usepackage{amsmath,amssymb,amsthm}
\usepackage{graphicx}
\usepackage{booktabs}
\usepackage{hyperref}
\usepackage{xcolor}
\usepackage{listings}
\input{shared/lstlang}
\usepackage{tikz}
\usetikzlibrary{shapes,arrows,positioning,fit}
\definecolor{codegreen}{rgb}{0,0.6,0}
\definecolor{codegray}{rgb}{0.5,0.5,0.5}
\definecolor{codepurple}{rgb}{0.58,0,0.82}
\definecolor{backcolour}{rgb}{0.95,0.95,0.92}
\lstdefinestyle{mystyle}{
backgroundcolor=\color{backcolour},
commentstyle=\color{codegreen},
keywordstyle=\color{codepurple},
numberstyle=\tiny\color{codegray},
stringstyle=\color{codegreen},
basicstyle=\ttfamily\footnotesize,
breakatwhitespace=false,
breaklines=true,
captionpos=b,
keepspaces=true,
numbers=left,
numbersep=5pt,
showspaces=false,
showstringspaces=false,
showtabs=false,
tabsize=2
}
\lstset{style=mystyle}
\newtheorem{theorem}{Theorem}
\newtheorem{lemma}{Lemma}
\newtheorem{proposition}{Proposition}
\newtheorem{definition}{Definition}
\title{Continuously-Learning Private AI: Self-Improvement and Privacy in the Hanzo Native Stack}
\author{Hanzo Dev \\ \texttt{[email protected]} \\ \textit{Hanzo AI}}
\date{June 2026}
\begin{document}
\hanzocoverpage
\begin{abstract}
A private continuously-learning model is three loops bolted together: an \emph{inference} loop that serves a user and captures its own interactions; a \emph{curation} loop that filters, redacts, and de-duplicates that capture into safe training data; and a \emph{training} loop that folds the data back into cheap, swappable per-user and per-task deltas. We design this system on the Hanzo native stack, which already owns the hard inference substrate---a ROCm decode matvec at the memory-bandwidth wall, at effective parity with llama.cpp (Qwen3-8B-Q8\_0 decode at $24.5$ tokens/s, $0.98\times$ the $25.05$ tokens/s of llama.cpp HIP on one gfx1151 / AMD~8060S)---and a real, tested reinforcement-learning trainer (a ported GRPO implementation with passing smoke tests). It does \emph{not} yet own the curation loop (no span-level PII redactor exists) or a Rust federated control plane (it is Python-only in the \texttt{zen/gym} project). We are explicit throughout about which pieces are real, which are ported, and which are research. We then formalize the two privacy-and-stability invariants the design rests on: a \emph{privacy invariant}---no personally-identifiable token enters the training set because a guard filter is applied to every datum before training---which we prove holds under the curation pipeline; and a \emph{bounded-forgetting} guarantee for delta-personalization---a frozen quantized base implies zero base-capability drift, with all change localized to low-rank adapters---which we also prove. We are honest about the two failure modes that can sink the system: model collapse from self-training, and privacy leakage through the delta.
\end{abstract}
\section{Introduction}
The case for keeping a learning loop on the user's own machine is not ideological; it is the only architecture in which raw user data never has to move. That requires both local inference \emph{and} local training to be fast enough to be invisible. This paper argues that the Hanzo native stack is unusually well positioned to host such a loop, designs the complete system, and is scrupulous about the boundary between what is built and what is aspirational.
Throughout, we are precise about the status of every claim, and we make the distinction in the prose rather than in shorthand: we say plainly when code exists today and was read from source, when a capability was ported from a named open-source project, when a piece is genuinely new Hanzo intellectual property, when something is buildable on the existing substrate with bounded engineering but is not yet built, and when a problem is open research. The reader should assume nothing here is built unless the text says so.
\paragraph{Inference is there.}
A sustained optimization campaign~\cite{hanzorocm} drove Qwen3-8B-Q8\_0 decode from $8.5$ to $\mathbf{24.5}$~tokens/s, i.e. $0.98\times$ the llama.cpp HIP figure of $25.05$~tokens/s---effective parity. The mechanism was a rewrite of the unified quantized matvec inner loop from a serial whole-warp-per-block scheme to a \emph{lane-strided, warp-reduce} scheme, taking the dominant feed-forward shapes from roughly $45\%$ to $\mathbf{97\%}$ of the $217$~GB/s memory roofline; the matvec is now bandwidth-wall-bound. Prefill is $1015$~tokens/s, or $0.86\times$ llama.cpp's $1176$~tokens/s---the harder remaining axis, requiring a multi-wave WMMA flash-attention kernel. Mixture-of-experts (MoE) decode improved $2\times$ ($6.34 \to 13.6$~tokens/s) via the same unified core plus a resident-bank cache. These are honest numbers on a single gfx1151 (AMD~8060S / Strix~Halo): decode is at parity, prefill trails, and beating llama.cpp on ROCm would require a PagedAttention keystone that has not yet paid off. The point for \emph{this} paper is narrower: a model that serves the user at $\sim$$25$~tokens/s decode is fast enough to also be the data generator for its own training, and the same quantized weights it serves from are the base on which deltas sit.
\paragraph{Training is partly there.}
As detailed in Section~\ref{sec:substrate}, the autograd engine, the optimizers, the checkpoint seam, and one full reinforcement-learning trainer are real and tested; a generic supervised fine-tuning (SFT) trainer that actually updates weights, and low-rank adapter injection, are not.
\paragraph{Contributions.}
\begin{itemize}
\item A complete three-loop architecture (serve--capture, curate, train, optionally federate) for on-device private continual learning, mapped onto concrete Hanzo source files with explicit reality tags.
\item A formal \emph{privacy invariant} (Definition~\ref{def:privinv}) and a proof that the curation pipeline maintains it (Proposition~\ref{prop:privinv}).
\item A formal \emph{bounded-forgetting} guarantee for delta-personalization over a frozen quantized base (Proposition~\ref{prop:forget}), with proof.
\item An honest treatment of the two existential failure modes---model collapse and delta leakage---and a prioritized build order separating bounded engineering from genuine research.
\end{itemize}
\section{The Loop}
\label{sec:loop}
The system is three coupled loops that never let raw data leave the device. Figure~\ref{fig:loop} sketches the data path.
\begin{figure}[t]
\centering
\begin{tikzpicture}[
node distance=0.55cm,
box/.style={rectangle, draw, minimum width=2.1cm, minimum height=0.5cm, align=center, font=\scriptsize},
arrow/.style={->, thick}
]
\node[box] (serve) {serve\\(quant base $+$ deltas)};
\node[box, below=of serve] (capture) {capture buffer};
\node[box, below=of capture] (curate) {curation\\(redact / guard / dedup / eval)};
\node[box, below=of curate] (train) {train\\(LoRA / GRPO)};
\node[box, below=of train] (fed) {federated\\trim-mean aggregator};
\draw[arrow] (serve) -- node[right,font=\tiny]{interactions} (capture);
\draw[arrow] (capture) -- node[right,font=\tiny]{self-gen / reflect} (curate);
\draw[arrow] (curate) -- node[right,font=\tiny]{safe data} (train);
\draw[arrow] (train.west) -- ++(-0.9,0) |- node[left,font=\tiny,pos=0.25]{apply delta} (serve.west);
\draw[arrow] (train) -- node[right,font=\tiny]{bf16 delta only (opt-in)} (fed);
\end{tikzpicture}
\caption{On-device continual-learning loop. Raw data never leaves the device; only an optional bf16 LoRA delta crosses the wire to the cross-device aggregator.}
\label{fig:loop}
\end{figure}
\paragraph{Stage A---serve and capture.}
The quantized base model (e.g. Qwen3-8B-Q8\_0), served through the real $460$-LoC \texttt{engine/hanzo-engine/src/models/quantized\_qwen3.rs} in the tree today, answers the user. Interactions---prompt, response, and any explicit or implicit feedback (thumbs, edits, acceptance of a code completion)---land in a local capture buffer. This buffer is plumbing that is buildable on the existing substrate but not yet built.
\paragraph{Stage B---self-generate / reflect.}
Beyond raw capture, the model manufactures training data: (i)~\emph{self-instruct}-style synthetic prompts in the user's domain; (ii)~\emph{self-critique / reflection}---generate an answer, then generate a critique-and-revise pass, keeping the revised answer as the target (a distillation-from-self signal). The honesty boundary here: the \texttt{zen/gym} reconnaissance confirmed it ships knowledge distillation (logit / top-$k$ and online-teacher), but the reflection / self-critique data-manufacturing described here is open research---a recipe, not ported code.
\paragraph{Stage C---curation.}
The gate that makes the loop private and stable: PII redaction, a safety gate, dedup / diversity, and an eval gate (Sections~\ref{sec:privacy} and~\ref{sec:collapse}). Nothing leaves this stage that contains personal data or that would degrade the model.
\paragraph{Stage D---train.}
Curated samples drive a parameter-efficient update---per-task / per-user LoRA deltas (cheap, swappable; Section~\ref{sec:delta})---via the native trainers (Section~\ref{sec:substrate}). For preference-shaped feedback (the user picked answer~A over~B), a DPO/KTO-style loss; for reward-shaped feedback (a score, a passing test), the real GRPO trainer. The update touches only the delta, never the frozen quantized base, which is what makes it cheap and forgetting-resistant (Section~\ref{sec:delta}).
\paragraph{Stage E---optional federation.}
If the user opts in to contribute, only the bf16 LoRA delta crosses the wire to a trim-mean aggregator (Section~\ref{sec:fed}). Raw data never does. This stage exists today only as a ported design, not in Rust: it is implemented in Python in \texttt{zen/gym} and is absent from the Hanzo \texttt{net/} crate.
\section{Privacy: What the Guard Must Do, and What It Cannot}
\label{sec:privacy}
The privacy property is defense in depth on the data path between capture and training. An early framing described \texttt{zen-guard} as ``the Rust filter family.'' The code says otherwise, and getting this right is load-bearing.
\paragraph{Ground truth on zen-guard.}
Read from the three \texttt{zen-guard} repositories: it is a \emph{model family}, a rebrand of Qwen3Guard~\cite{qwen3guard}---there is \emph{zero Rust}, no \texttt{Cargo.toml}, no \texttt{.rs} file in any of them. \texttt{zen-guard-gen} (8B) is a generative classifier: it applies a chat template, generates, and regex-parses a verdict of the form \texttt{Safety: Safe|Unsafe|Controversial; Categories: \ldots; Refusal: Yes|No}. \texttt{zen-guard-stream} (4B) carries a per-token classification head (\texttt{stream\_moderate\_from\_ids(token\_id, role, stream\_state)} returning a risk level in $\{\textsf{Safe}, \textsf{Controversial}, \textsf{Unsafe}\}$). There are nine categories, one of which is PII.
The critical nuance: \textbf{PII is a classification \emph{label}, not a redaction \emph{action}}. The model answers ``does this turn contain or solicit PII disclosure?'' over a whole turn. It does \emph{not} locate spans, emit offsets, or return a \texttt{<REDACTED>} string. Therefore:
\begin{itemize}
\item zen-guard is the right \textbf{gate / quarantine} mechanism: any self-generated or captured sample the classifier flags \textsf{Unsafe} or \texttt{Categories: PII} is dropped or quarantined before it enters the training buffer. This gate is buildable on the existing substrate once the guard is served in-process.
\item zen-guard \textbf{cannot scrub}. For the keep-but-redact path one must build a span-level PII redactor that the existing guard does not provide. This is the one genuinely missing privacy primitive, and it is new work relative to zen-guard / Qwen3Guard (though general named-entity redactors such as Presidio~\cite{presidio} are prior art).
\end{itemize}
\paragraph{Defense in depth.}
The designed pipeline applies, to every captured or self-generated datum, in order:
\begin{enumerate}
\item \textbf{Deterministic redactor.} A cheap Rust pass (regex plus lightweight NER) that locates and replaces structured PII spans---emails, phone numbers, card/SSN-shaped tokens, API keys and secrets, file paths embedding usernames---with typed sentinel tokens. Deterministic, inline, cents-of-a-millisecond. This is buildable on the existing substrate but greenfield: a grep over \texttt{net/} and \texttt{ml/} confirms no redactor exists.
\item \textbf{Neural gate.} Run the distilled guard classifier over the \emph{post-redaction} sample; if it still trips PII or \textsf{Unsafe} (catching free-text disclosures, medical or financial narrative that regex misses), drop the sample entirely.
\item \textbf{Secret-specific filters.} High-entropy-string and known-key-prefix detectors for credentials, which a general PII model under-weights.
\item \textbf{Egress containment.} Even after redaction, raw data never leaves the device; only the LoRA delta does (Section~\ref{sec:fed}).
\end{enumerate}
\paragraph{Serving the gate efficiently.}
Running an 8B guard per sample is too heavy for bulk curation. The design serves a distilled $0.6$B guard \emph{in-process} through the engine's quantized-Qwen3 path (\texttt{quantized\_qwen3.rs}) rather than a Python/vLLM sidecar---turning the gate into a local function call on already-fast inference. This is buildable on the existing substrate: the quantized path is real; the distilled $0.6$B guard and the in-process wiring are the work that remains.
\subsection{The Privacy Invariant, Formally}
\label{sec:privinv}
We now state the privacy guarantee precisely. Let $\mathcal{V}$ be the model vocabulary and let a \emph{datum} $x$ be a finite token sequence over $\mathcal{V}$. Let $\mathcal{P}(x) \subseteq \{0,\dots,|x|-1\}$ denote the (ground-truth) set of token positions in $x$ that constitute personally-identifiable information, and write $\textsf{pii}(x) \equiv \mathcal{P}(x) \neq \varnothing$.
Let $R$ be the deterministic redactor, a map on token sequences that replaces every position in a detected set $\hat{\mathcal{P}}(x) \subseteq \mathcal{P}(x)$ with a fixed, PII-free sentinel symbol $\bot \in \mathcal{V}$. Let $F$ be the guard gate, a predicate $F(x) \in \{\textsf{keep}, \textsf{drop}\}$ with the convention that a sample is admitted iff $F(x)=\textsf{keep}$. The composed curation filter is
\[
\Phi(x) \;=\;
\begin{cases}
R(x), & F\big(R(x)\big)=\textsf{keep},\\[2pt]
\varnothing, & \text{otherwise,}
\end{cases}
\]
and the training set after curating a capture stream $X$ is $\mathcal{T} = \{\,\Phi(x) : x \in X,\ \Phi(x)\neq\varnothing\,\}$.
\begin{definition}[Privacy invariant]
\label{def:privinv}
A training set $\mathcal{T}$ \emph{satisfies the privacy invariant} if no admitted sample contains a PII token: for every $t \in \mathcal{T}$, $\textsf{pii}(t)=\textsf{false}$, i.e. $\mathcal{P}(t)=\varnothing$.
\end{definition}
To make the guarantee non-vacuous we state the exact assumptions under which it holds; honesty about residual risk (below) is then a statement about when these assumptions fail.
\begin{definition}[Sound gate]
\label{def:sound}
The gate $F$ is \emph{sound for PII} if, for all $x$, $\textsf{pii}(x)=\textsf{true} \Rightarrow F(x)=\textsf{drop}$. Equivalently, $F(x)=\textsf{keep} \Rightarrow \mathcal{P}(x)=\varnothing$.
\end{definition}
\begin{proposition}[Pipeline maintains the privacy invariant]
\label{prop:privinv}
If the redactor sentinel $\bot$ is PII-free and the gate $F$ is sound for PII (Definition~\ref{def:sound}), then for every capture stream $X$ the curated set $\mathcal{T}=\{\Phi(x):x\in X,\ \Phi(x)\neq\varnothing\}$ satisfies the privacy invariant (Definition~\ref{def:privinv}).
\end{proposition}
\begin{proof}
Let $t \in \mathcal{T}$ be arbitrary. By construction of $\mathcal{T}$ there is an $x \in X$ with $t = \Phi(x)$ and $\Phi(x) \neq \varnothing$. By the definition of $\Phi$, the second branch ($\varnothing$) is excluded, so the first branch was taken: $t = R(x)$ and $F\big(R(x)\big)=\textsf{keep}$, i.e. $F(t)=\textsf{keep}$.
Because $F$ is sound for PII (Definition~\ref{def:sound}), $F(t)=\textsf{keep}$ implies $\mathcal{P}(t)=\varnothing$. Hence $\textsf{pii}(t)=\textsf{false}$. Since $t$ was an arbitrary element of $\mathcal{T}$, every admitted sample is PII-free, which is exactly Definition~\ref{def:privinv}.
\end{proof}
\begin{lemma}[Redaction never increases PII]
\label{lem:redact}
If the sentinel $\bot$ is PII-free, then $\mathcal{P}(R(x)) \subseteq \mathcal{P}(x)$ for every $x$; in particular the redactor cannot create new PII.
\end{lemma}
\begin{proof}
$R$ acts position-wise: positions outside $\hat{\mathcal{P}}(x)$ are copied verbatim, and positions inside $\hat{\mathcal{P}}(x)$ are overwritten with $\bot$. A position $i$ can be in $\mathcal{P}(R(x))$ only if $R(x)$ exposes PII at $i$. If $i \in \hat{\mathcal{P}}(x)$ then $R(x)_i = \bot$, which is PII-free by hypothesis, so $i \notin \mathcal{P}(R(x))$. If $i \notin \hat{\mathcal{P}}(x)$ then $R(x)_i = x_i$, so $i$ can contribute PII to $R(x)$ only if it already did so in $x$, i.e. $i \in \mathcal{P}(x)$. Therefore $\mathcal{P}(R(x)) \subseteq \mathcal{P}(x) \setminus \hat{\mathcal{P}}(x) \subseteq \mathcal{P}(x)$.
\end{proof}
Lemma~\ref{lem:redact} clarifies the division of labor that the pipeline relies on: the redactor monotonically \emph{reduces} the PII set (it removes $\hat{\mathcal{P}}(x)$ and adds nothing), while the gate $F$ supplies the soundness that converts ``reduced'' into ``empty.'' The redactor is what makes soundness of $F$ achievable in practice---it removes the easy structured spans so the classifier only has to catch the residual free-text disclosures---but it is the gate, not the redactor, that the invariant proof leans on.
\paragraph{Honest residual risk.}
No gate is perfect, so Definition~\ref{def:sound} is an idealization, and Proposition~\ref{prop:privinv} is exactly as strong as that assumption. A real distilled classifier has false negatives; the redactor will miss novel PII formats (so $\hat{\mathcal{P}}(x)$ may be a strict subset of $\mathcal{P}(x)$); and an adversarial prompt can smuggle PII into a self-generated sample. The mitigation is layering plus egress-only-delta containment, \emph{not} a claim of zero leakage. Two consequences follow honestly. First, the invariant degrades gracefully: by Lemma~\ref{lem:redact} every admitted sample's PII set is bounded above by what the gate failed to catch \emph{after} the redactor has already deleted $\hat{\mathcal{P}}(x)$, so even under an imperfect gate the residual exposure is no larger than the classifier's false-negative set on already-redacted text. Second, the delta itself can memorize and leak training content---addressed in Section~\ref{sec:fed}; differential privacy on the delta is open research, not in code.
\section{The Training Substrate: Real vs.\ Stub}
\label{sec:substrate}
Read from the \texttt{ml/} tree. We separate what updates weights today from what does not.
\begin{itemize}
\item \textbf{Autograd (real, in the tree).} \texttt{hanzo-ml/src/backprop.rs}: \texttt{GradStore} plus \texttt{Tensor::backward()} (candle's~\cite{candle} reverse-mode engine, intact); \texttt{Var} threads through.
\item \textbf{Optimizers (real in the right place, a stub in the wrong place).} The real ones live in \texttt{hanzo-nn/src/optim.rs}: an \texttt{Optimizer} trait with \texttt{step(\&GradStore)} and \texttt{backward\_step(\&loss)}; a working \texttt{SGD} and a full decoupled-weight-decay \texttt{AdamW} (per-\texttt{Var} first/second moments, bias correction, eps, weight decay). The stub is \texttt{hanzo-training/src/optimizer.rs}: \texttt{OptimizerWrapper::step()} is a verbatim no-op---its own comment reads \emph{``Simplified optimizer step---in practice, this would update parameters''}---it only bumps a counter and recomputes the learning rate. The generic \texttt{hanzo-training::Trainer} is built on that stub, so the generic SFT path does not update weights yet. Replacing the stub with \texttt{hanzo\_nn::AdamW} is the headline near-term task and is buildable on the existing substrate today (the real optimizer already exists one crate over).
\item \textbf{Checkpoint seam (real, in the tree).} \texttt{hanzo-nn/src/var\_map.rs}: \texttt{VarMap} with \texttt{save}/\texttt{load} over safetensors and \texttt{get}/\texttt{set}. This \emph{is} the canonical bf16 delta seam the federation layer needs (Section~\ref{sec:fed}); \texttt{save} calls \texttt{safetensors::serialize\_to\_file}.
\item \textbf{Losses and layers (real, in the tree).} \texttt{hanzo-nn/src/loss.rs}: \texttt{nll}, \texttt{cross\_entropy}, \texttt{mse}, \texttt{binary\_cross\_entropy\_with\_logit}, \texttt{huber}. Transformer building blocks are present (\texttt{moe}, \texttt{rotary\_emb}, \texttt{kv\_cache}, \texttt{linear}, \texttt{attention}).
\item \textbf{GRPO (real, tested, and ported from \texttt{zooai-gym}).} This is the proof that native training in Rust is not theoretical. \texttt{hanzo-training/src/grpo/} (8 files); the \texttt{mod.rs} docstring states it was \emph{``ported from the \texttt{zooai-gym} GRPO~\cite{grpo} trainer.''} \texttt{GrpoTrainer<P: Policy, R: Policy>} runs the full five-phase step: rollout~$\to$ reward~$\to$ group-relative advantage $A_i = (r_i - \mathrm{mean})/(\mathrm{std}+\epsilon)$~$\to$ policy-gradient loss $-\,\mathrm{mean}(A_i \cdot \log p)$ plus optional KL to a frozen reference policy~$\to$ \texttt{AdamW.backward\_step(\&loss)}, with a non-finite-loss guard and per-step seed advance. The \texttt{Reward} trait is \texttt{Send + Sync} and pluggable. Two smoke tests train a toy policy through the real code path and assert that mean reward climbs toward the ceiling.
\end{itemize}
The gaps to close for the loop are summarized in Table~\ref{tab:substrate}; all are buildable on this substrate with bounded engineering. The reward-shaped path (a test passes, a score) is runnable today through GRPO---only a real \texttt{Policy} over the engine model is missing. The preference-shaped path (the user picked~A over~B) needs the DPO/KTO~\cite{dpo} losses (a straightforward port); the plain SFT path needs the stub replaced. None of this is research; it is bounded engineering on a substrate that already runs one RL recipe end to end.
\begin{table}[t]
\centering
\caption{Training-substrate gaps for the continual-learning loop.}
\label{tab:substrate}
\small
\begin{tabular}{@{}lll@{}}
\toprule
\textbf{Need} & \textbf{Status} & \textbf{Work} \\
\midrule
Generic SFT (updates weights) & Stub (no-op) & Swap in \texttt{AdamW} \\
LoRA injection & Config only & Adapter layer + merge \\
DPO / KTO / ORPO / SimPO & Absent & Preference losses \\
GRPO on real policy & Toy tested & \texttt{Policy} for Qwen3 \\
\bottomrule
\end{tabular}
\end{table}
\section{Self-Generated Data Quality: Not Poisoning the Well}
\label{sec:collapse}
The existential risk of a self-training loop is \emph{model collapse}: recursive training on synthetic output narrows the distribution until the model forgets the tails and degenerates. This is a published failure mode~\cite{collapse}, and any honest design must treat it as the default outcome to be actively prevented, not an edge case. The curation-stage mitigations are:
\begin{enumerate}
\item \textbf{Grounding beats pure self-gen.} The strongest data is real captured interactions with a real feedback signal (the user edited the answer; the test passed; the completion was accepted). Self-generated synthetic data is a supplement, weighted below grounded data, never the sole source. Reflection / self-critique keeps a \emph{verifiable} target (a revised answer that passes a check), not a free-floating one.
\item \textbf{Eval gate.} Maintain a frozen held-out probe set (a small fixed eval and an identity/behavior canary). After each candidate delta, evaluate; commit the delta \emph{only} if the eval does not regress. This is the single most important guard against silent drift, and it is buildable on the existing substrate: \texttt{hanzo-training} already has an \texttt{evaluation.rs} and an \texttt{eval} binary.
\item \textbf{Dedup and diversity.} Near-duplicate self-gen samples amplify whatever bias produced them. De-duplicate (hash / embedding-similarity) and enforce a diversity floor. Buildable on the existing substrate but greenfield.
\item \textbf{Replay / anti-forgetting mixing.} Interleave a small fraction of general base data with the personalized data so the delta personalizes without erasing the base distribution---exactly what LoRA-over-frozen-base buys cheaply (Section~\ref{sec:delta}).
\item \textbf{Provenance and reversibility.} Every committed delta is checkpointed (\texttt{VarMap::save}, which is real and in the tree) and tagged with the data window that produced it, so a bad update is one \texttt{load} away from rollback.
\end{enumerate}
Honest framing: items 2--5 \emph{reduce} collapse risk; they do not eliminate it. Long-horizon stability of a fully autonomous self-training loop is open research. The defensible near-term posture is bounded, gated, human-signal-anchored continual updates with a hard eval gate and one-click rollback---not an unsupervised model that retrains itself nightly on its own hallucinations.
\section{Personalization: Delta Stacks over a Shared Quantized Base}
\label{sec:delta}
The architecture that makes per-user / per-task continual learning cheap is a stack of small deltas over one frozen, quantized base:
\begin{itemize}
\item \textbf{One shared quantized base in VRAM.} The engine already serves quantized weights (Q8\_0 at the bandwidth wall today; native decode also wired for Q4\_K). The base is frozen and shared across every user and task on the device---paid for once.
\item \textbf{Per-task / per-user LoRA deltas.} Each personalization is a low-rank delta~\cite{lora} (a few MB), trained on that user's curated data. \emph{Swappable}: load the delta for the active context, hot-swap on task change. \emph{Mergeable}: a ``LoRA soup'' averages or trim-means a set of deltas into one (the same operation the federated aggregator performs, Section~\ref{sec:fed}). This is buildable on the existing substrate once LoRA injection exists (Section~\ref{sec:substrate}); the config knobs are already declared.
\end{itemize}
Why deltas, specifically, for continual learning: they are \emph{cheap} (train and store MB, not GB; many personas coexist over one base); \emph{forgetting-resistant by construction} (the base is frozen, so general capability cannot be overwritten---the delta only adds); and \emph{composable and reversible} (stack, merge, or drop deltas, and roll back a bad one without touching the base).
\paragraph{On BitDelta---an explicit honesty note.}
Earlier framing mentioned BitDelta as ``ideal.'' \textbf{BitDelta does not exist anywhere in this stack.} Reconnaissance grep-confirmed: there is no BitDelta, no delta-compression, no model-soup/mergekit/TIES/DARE in \texttt{zooai/gym}, and the only ``DeltaSoup'' is the federated trim-mean aggregator in \texttt{zen/gym} (Section~\ref{sec:fed})---which averages full-precision bf16 LoRA deltas, not 1-bit BitDelta. So: LoRA-delta personalization is buildable on the existing substrate (pending injection code); BitDelta~\cite{bitdelta}-style 1-bit delta quantization is open research---attractive, and a natural fit for the existing unified \texttt{quant\_format!} / \texttt{qmatvec\_core<WTYPE>} machinery the ROCm work built, but not implemented and not to be claimed as present.
The unified quant abstraction is the genuinely novel Hanzo IP that makes the base side of this story strong: the ROCm work collapsed all formats, from 1-bit ternary to full precision, into one quant-agnostic GPU core---\texttt{qmatvec\_core<WTYPE>} for decode, \texttt{qmmq\_core<WTYPE>} for prefill, with a single \texttt{decode\_block<WTYPE>} plus a traits row per format and zero new kernels per quant. That is the substrate that would host BitDelta deltas natively if and when built.
\subsection{Bounded Forgetting, Formally}
\label{sec:forget-formal}
We now prove the structural forgetting-resistance claim. Let $W_0 \in \mathbb{R}^{m\times n}$ be a (frozen) base weight matrix and let the LoRA adapter be $\Delta W = \tfrac{\alpha}{r} BA$ with $B \in \mathbb{R}^{m\times r}$, $A \in \mathbb{R}^{r\times n}$, $r \ll \min(m,n)$. The effective weight is $W = W_0 + \Delta W$. Let $f_W(\cdot)$ denote the network's map under weights $W$, and let $\mathcal{C}_{\mathcal{B}}(W) = \mathbb{E}_{x\sim\mathcal{B}}\,[\,\ell(f_W(x))\,]$ be a base-capability functional measured on a frozen base distribution $\mathcal{B}$, for a fixed metric $\ell$ (e.g. a held-out base-eval score). Training adjusts only the adapter parameters $(A,B)$; $W_0$ is never written.
\begin{definition}[Base-capability drift]
\label{def:drift}
The \emph{base-capability drift} induced by an adapter $\Delta W$ is $D(\Delta W) = \mathcal{C}_{\mathcal{B}}(W_0+\Delta W) - \mathcal{C}_{\mathcal{B}}(W_0)$.
\end{definition}
\begin{proposition}[Bounded forgetting under a frozen base]
\label{prop:forget}
Suppose training writes only the adapter parameters $(A,B)$ and leaves $W_0$ byte-for-byte unchanged. Then:
\begin{enumerate}
\item[(i)] \textbf{Zero drift at initialization.} Under the standard LoRA initialization $B=0$ (with $A$ arbitrary), the adapter contributes $\Delta W = 0$, so $W = W_0$ and $D(\Delta W) = 0$ exactly. The model at the start of every personalization episode is bitwise the base on the entire base distribution.
\item[(ii)] \textbf{Localized, reversible change.} For any trained adapter, the change to the effective weight is confined to the rank-$\le r$ subspace $\operatorname{col}(B)$: $\operatorname{rank}(\Delta W) \le r$ and $\Delta W\, v = 0$ for every $v \in \ker(A)$, a subspace of dimension at least $n-r$. Setting $(A,B) \leftarrow (A,0)$ restores $W_0$ exactly, so any drift $D(\Delta W)$ is fully reversible without touching the base.
\item[(iii)] \textbf{Bounded drift.} If $\mathcal{C}_{\mathcal{B}}$ is $L$-Lipschitz in the weights under the spectral norm, then $|D(\Delta W)| \le L\,\|\Delta W\|_2 \le L\,\tfrac{\alpha}{r}\,\|B\|_2\,\|A\|_2$. In particular, with the standard practice of mixing a replay fraction of base data and gating on the base eval (Section~\ref{sec:collapse}), the optimizer is steered toward small $\|\Delta W\|_2$ and hence small drift.
\end{enumerate}
\end{proposition}
\begin{proof}
\emph{(i)} With $B=0$ we have $\Delta W = \tfrac{\alpha}{r} BA = 0$ regardless of $A$, hence $W = W_0 + 0 = W_0$. Therefore $\mathcal{C}_{\mathcal{B}}(W_0+\Delta W) = \mathcal{C}_{\mathcal{B}}(W_0)$ and $D(\Delta W) = 0$ by Definition~\ref{def:drift}. Since every layer's effective weight equals its base weight and no other parameters are modified, $f_W \equiv f_{W_0}$ pointwise, so the equality holds on the entire support of $\mathcal{B}$ (indeed on all inputs), not merely in expectation.
\emph{(ii)} $\Delta W = \tfrac{\alpha}{r}BA$ is a product whose left factor $B$ has at most $r$ columns, so $\operatorname{rank}(\Delta W) \le \operatorname{rank}(B) \le r$. For $v \in \ker(A)$ we have $Av = 0$, hence $\Delta W\,v = \tfrac{\alpha}{r}B(Av) = 0$; by rank–nullity $\dim\ker(A) \ge n - \operatorname{rank}(A) \ge n - r$. Because $W_0$ was never written, the base weights are still present in memory unmodified; assigning $B \leftarrow 0$ yields $\Delta W = 0$ and restores $W = W_0$ exactly, so the map returns to $f_{W_0}$ and the drift returns to $0$.
\emph{(iii)} By $L$-Lipschitzness of $\mathcal{C}_{\mathcal{B}}$ under the spectral norm,
\[
|D(\Delta W)| = \big|\mathcal{C}_{\mathcal{B}}(W_0+\Delta W) - \mathcal{C}_{\mathcal{B}}(W_0)\big| \le L\,\|\Delta W\|_2 .
\]
Submultiplicativity of the spectral norm gives $\|\Delta W\|_2 = \tfrac{\alpha}{r}\|BA\|_2 \le \tfrac{\alpha}{r}\|B\|_2\|A\|_2$, which yields the stated bound. The drift is thus controlled by the adapter norm; the bound is tight in the trivial direction by part~(i) ($\|\Delta W\|_2=0 \Rightarrow D=0$).
\end{proof}
The contrast with full fine-tuning is the point of Proposition~\ref{prop:forget}. If training were allowed to write $W_0$, then part~(i) fails (there is no fixed reference to return to), part~(ii) fails (the change need not be low-rank and is not reversible by zeroing an adapter), and catastrophic forgetting of the base is possible. Freezing the base converts forgetting-resistance from a hope into a structural property: drift is exactly zero at the start of each episode, is confined to a rank-$\le r$ subspace, is reversible by construction, and is norm-bounded thereafter. The adapter can still over-specialize \emph{within} that subspace---which is precisely what the replay mixing and eval gate of Section~\ref{sec:collapse} are for---but it cannot overwrite the base.
\section{Federation: Delta-Only Egress}
\label{sec:fed}
If a user opts to contribute their learning to a shared model \emph{without sharing their data}, the mechanism is: train locally, export only the bf16 LoRA delta, aggregate across contributors with a robust mean, pull the consensus delta back, and apply it.
\paragraph{Ground truth.}
In \texttt{zen/gym} (Python-only), the load-bearing invariant is a canonical BF16 LoRA-delta safetensors blob as the cross-vendor seam: \texttt{backend/base.py}'s \texttt{export\_lora\_delta} / \texttt{import\_lora\_delta} let MLX, CUDA, ROCm, and CPU federate with no cross-vendor NCCL---\texttt{distributed/\_\_init\_\_.py} states, \emph{``sync at LoRA-delta granularity over HTTP. No cross-vendor NCCL.''} The control plane: \texttt{topology.py} ($\textsf{capacity\_score} = \sqrt{\textsf{mem}}\cdot\sqrt{\textsf{tflops}}$), \texttt{scheduler.py} (capacity-weighted data sharding plus best-fit-decreasing MoE expert pinning), \texttt{transport.py} (vendor-neutral HTTP, HMAC-SHA256 over \texttt{method|path|ts|sha256(body)}), \texttt{coordinator.py} (aggregation), and \texttt{worker.py} (a model-agnostic loop: train $N$ steps~$\to$ export delta~$\to$ push~$\to$ pull aggregate~$\to$ apply). The aggregator \texttt{coordinator.py:aggregate()} is Byzantine-robust trimmed-mean~\cite{trimmedmean} (drop top and bottom 1 if $N\ge 4$), or median, or mean, over the bf16$\to$f32 view; this is federated averaging~\cite{fedavg} restricted to the LoRA-delta seam, and its docstring calls it \emph{``DeltaSoup-style \ldots\ without the full reputation system.''}
\paragraph{Two honesty boundaries.}
\begin{enumerate}
\item \textbf{This is not in Rust.} A grep over Hanzo \texttt{net/src} for \texttt{federat}\discretionary{}{}{}\texttt{|deltasoup}\discretionary{}{}{}\texttt{|lora.delta}\discretionary{}{}{}\texttt{|byzantine}\discretionary{}{}{}\texttt{|aggregate.*delta} returns only a false-positive CSS file. The whole coordinator/worker/transport/DeltaSoup stack is a ported design whose code is not yet ported---Python in \texttt{zen/gym}, absent in Hanzo. Porting it is buildable on the existing substrate and small: the seam is just safetensors bf16 plus a uint16$\leftrightarrow$f32 bit-shift, and \texttt{VarMap} (Section~\ref{sec:substrate}) already gives that; a Rust coordinator/worker over \texttt{axum}/\texttt{reqwest} with HMAC is well-specified. The \texttt{worker.py} loop maps cleanly onto the native trainer: \texttt{step\_fn} $=$ \texttt{GrpoTrainer::step} (or an SFT step), \texttt{params\_iter} $=$ the LoRA \texttt{Var}s, \texttt{apply\_fn} $=$ add the consensus delta into those \texttt{Var}s.
\item \textbf{The delta still leaks.} ``Only the delta crosses the wire'' bounds leakage; it does not zero it---a LoRA delta can memorize training content and is invertible-ish under attack. The clean mitigation is differential-privacy noise on the delta before egress~\cite{dpsgd}, which is open research and not in code (grep-confirmed: no DP / Gaussian / Laplace noise in \texttt{zen/gym}).
\end{enumerate}
\paragraph{Why this is the strongest patent angle.}
The heterogeneous-vendor federation via the canonical-bf16 delta seam plus capacity scheduler is the strongest filing candidate: code substantially exists (Python), and the non-obvious combination is the \emph{backend-identity-erasing bf16-delta as the sole interop seam} enabling truly mixed-vendor federation with no shared collective. The plain trimmed-mean algorithm itself is federated-learning prior art and must not be claimed alone. A second candidate is the privacy continual-learning loop (self-gen PII gate plus delta-only egress plus trim-mean, with DP-on-delta as the novelty hook to clear generic-FL prior art)---to be filed \emph{before} building.
\section{Results and Status Summary}
\label{sec:results}
Table~\ref{tab:perf} collects the measured performance numbers from the ROCm campaign (single gfx1151 / AMD~8060S, $217$~GB/s memory roofline); Table~\ref{tab:status} summarizes the reality tags for each component of the loop. We report parity, not superiority: decode is at $0.98\times$ llama.cpp, prefill trails at $0.86\times$.
\begin{table}[t]
\centering
\caption{Measured inference performance, Hanzo vs.\ llama.cpp HIP, single gfx1151 (AMD~8060S).}
\label{tab:perf}
\small
\begin{tabular}{@{}lrrr@{}}
\toprule
\textbf{Workload} & \textbf{Hanzo} & \textbf{llama.cpp} & \textbf{Ratio} \\
\midrule
Q8\_0 decode (t/s) & 24.5 & 25.05 & 0.98$\times$ \\
Q8\_0 prefill (t/s) & 1015 & 1176 & 0.86$\times$ \\
MoE decode (t/s) & 13.6 & --- & 2.0$\times$\,$^{\dagger}$ \\
FFN roofline (\%) & 97 & --- & --- \\
\bottomrule
\end{tabular}
\\[2pt]
{\footnotesize $^{\dagger}$MoE decode improvement is $2\times$ over the prior Hanzo baseline ($6.34 \to 13.6$~t/s), not over llama.cpp.}
\end{table}
\begin{table}[t]
\centering
\caption{Reality status of each loop component.}
\label{tab:status}
\small
\begin{tabular}{@{}ll@{}}
\toprule
\textbf{Component} & \textbf{Status} \\
\midrule
Quantized inference (decode parity) & Shipping \\
Autograd, optimizers, checkpoint & Shipping \\
GRPO trainer (toy policy) & Shipping, tested, ported \\
Unified quant core (\texttt{qmatvec\_core}) & Shipping, novel \\
Generic SFT (weight update) & Stub (no-op) \\
LoRA injection & Config only \\
Span-level PII redactor & Novel, not built \\
In-process distilled guard & Planned (buildable) \\
Eval gate / dedup / replay & Planned (buildable) \\
Preference losses (DPO/KTO/\ldots) & Absent \\
Rust federation control plane & Ported design (Python only) \\
DP-on-delta; BitDelta; stability & Research \\
\bottomrule
\end{tabular}
\end{table}
\section{Limitations and Future Work}
\label{sec:limits}
We are explicit about what is not solved.
\paragraph{The two failure modes that can sink the system.}
\emph{Model collapse} is the default failure of self-training; Section~\ref{sec:collapse}'s eval gate, grounded-data weighting, dedup, and replay \emph{bound} it but do not eliminate it. Long-horizon autonomous stability is open research, so one must not ship an unsupervised nightly self-retrainer. \emph{Privacy leakage} is bounded by the redactor~$\to$ gate~$\to$ secret-filter~$\to$ delta-only-egress chain of Sections~\ref{sec:privacy} and~\ref{sec:fed}; but the redactor is greenfield, DP-on-delta is research, and Proposition~\ref{prop:privinv} is exactly as strong as the gate-soundness assumption it rests on---so there is \emph{no} ``zero leakage'' claim.
\paragraph{Three lesser risks.}
Per-user deltas can over-specialize despite a frozen base (Proposition~\ref{prop:forget} bounds base drift, not within-subspace specialization; mitigated by replay and eval, not eliminated). zen-guard miscategorizes and labels PII per turn, not per span---so the redactor must do the span work the classifier cannot. And the eval probe set itself can be gamed by a loop that overfits to it; rotate probes and keep a canary.
\paragraph{Build order.}
The dividing line is sharp. Items 1--6 are bounded engineering buildable on the existing substrate---which already runs one RL recipe end to end and serves quantized inference at decode parity with llama.cpp---while item~7 is genuine open research and the cleanest patent trigger.
\begin{enumerate}
\item Replace the \texttt{OptimizerWrapper} stub with \texttt{hanzo\_nn::AdamW}---a generic SFT trainer that actually updates weights.
\item LoRA injection layer plus merge---enables the swappable per-user delta stack.
\item Span-level PII redactor in Rust plus a distilled $0.6$B guard served in-process---the curation gate, and the missing privacy primitive.
\item Eval gate, dedup/diversity, and capture/replay buffer---the collapse defenses and the capture/self-gen plumbing.
\item Preference losses (DPO/KTO/ORPO/SimPO) plus a real \texttt{Policy} over the engine model---both feedback-shaped training paths.
\item Port the \texttt{zen/gym} federation control plane to Rust over the \texttt{VarMap} bf16 seam---opt-in delta-only federation.
\item \emph{(Research, not engineering.)} DP-noise-on-delta; BitDelta 1-bit delta quantization over the unified \texttt{qmatvec\_core<WTYPE>} core; long-horizon collapse stability.
\end{enumerate}
\section{Conclusion}
A private continuously-learning model is realizable on the Hanzo native stack because the hardest substrate---fast on-device quantized inference---is already at parity with llama.cpp on ROCm, and one full RL trainer already runs end to end. The privacy and stability properties the design rests on are not hand-waved: the curation pipeline maintains a precise privacy invariant under a sound gate (Proposition~\ref{prop:privinv}), and delta-personalization over a frozen base guarantees zero base-capability drift at episode start with norm-bounded, reversible, low-rank change thereafter (Proposition~\ref{prop:forget}). Everything we claim as built was read from source; every performance figure was cited from the measured ROCm campaign; and the two pieces that can sink the system---model collapse and delta leakage---are named, bounded, and left honestly open.
\begin{thebibliography}{10}
\bibitem{candle}
The Candle Authors. \textit{Candle: Minimalist ML Framework for Rust}. Hugging Face, 2023.
\bibitem{lora}
E. Hu, Y. Shen, P. Wallis, Z. Allen-Zhu, Y. Li, S. Wang, L. Wang, and W. Chen. \textit{LoRA: Low-Rank Adaptation of Large Language Models}. ICLR, 2022.
\bibitem{grpo}
Z. Shao, P. Wang, Q. Zhu, et al. \textit{DeepSeekMath: Pushing the Limits of Mathematical Reasoning in Open Language Models}. arXiv:2402.03300, 2024.
\bibitem{qwen3guard}
Qwen Team. \textit{Qwen3Guard: A Generative and Streaming Safety Moderation Model}. Alibaba, 2025.
\bibitem{collapse}
I. Shumailov, Z. Shumaylov, Y. Zhao, Y. Gal, N. Papernot, and R. Anderson. \textit{The Curse of Recursion: Training on Generated Data Makes Models Forget}. arXiv:2305.17493, 2023.
\bibitem{dpo}
R. Rafailov, A. Sharma, E. Mitchell, S. Ermon, C. Manning, and C. Finn. \textit{Direct Preference Optimization: Your Language Model is Secretly a Reward Model}. NeurIPS, 2023.
\bibitem{bitdelta}
J. Liu, G. Xiao, K. Li, J. D. Lee, S. Han, T. Dao, and T. Cai. \textit{BitDelta: Your Fine-Tune May Only Be Worth One Bit}. NeurIPS, 2024.
\bibitem{dpsgd}
M. Abadi, A. Chu, I. Goodfellow, H. B. McMahan, I. Mironov, K. Talwar, and L. Zhang. \textit{Deep Learning with Differential Privacy}. ACM CCS, 2016.
\bibitem{fedavg}
H. B. McMahan, E. Moore, D. Ramage, S. Hampson, and B. A. y Arcas. \textit{Communication-Efficient Learning of Deep Networks from Decentralized Data}. AISTATS, 2017.
\bibitem{trimmedmean}
D. Yin, Y. Chen, K. Ramchandran, and P. Bartlett. \textit{Byzantine-Robust Distributed Learning: Towards Optimal Statistical Rates}. ICML, 2018.
\bibitem{presidio}
Microsoft. \textit{Presidio: Data Protection and De-identification SDK}. 2020.
\bibitem{hanzorocm}
Hanzo AI Research. \textit{Native ROCm Quantized Inference on gfx1151: A Bandwidth-Wall Optimization Campaign}. Hanzo AI, 2026.
\end{thebibliography}
\end{document}
hanzo-cross-backend-inference.pdf
BIN
View File
Binary file not shown.
hanzo-cross-backend-inference.tex
+225
View File
@@ -0,0 +1,225 @@
\documentclass[11pt]{article}
\usepackage[margin=1in]{geometry}
\usepackage{booktabs}
\usepackage{amsmath}
\usepackage{graphicx}
\usepackage{xcolor}
\usepackage{hyperref}
\hypersetup{colorlinks=true, linkcolor=blue, urlcolor=blue, citecolor=blue}
\usepackage{enumitem}
\title{\textbf{Edge LLM Inference Across Commodity Accelerators:\\
Measured AMD, NVIDIA, and Apple Performance, and\\
hanzo-engine at llama.cpp Decode Parity on Every Backend}}
\author{Hanzo AI --- hanzo-engine technical white paper, 05}
\date{Measurement pass: 2026-06-17}
\begin{document}
\maketitle
\begin{abstract}
This paper reports \emph{measured} single-GPU inference performance for one model
(\textbf{Qwen3-8B-Q8\_0}) across three commodity edge accelerators --- AMD Ryzen AI Max+ 395
``Strix Halo'' (Radeon 8060S iGPU, gfx1151), NVIDIA GB10 ``Grace Blackwell'' (sm\_121), and
Apple M4 Max --- comparing \textbf{hanzo-engine} against \texttt{llama.cpp} on each, same model,
same metric (512-token prefill throughput \texttt{pp512}; single-stream decode throughput).
Every number here is measured on the named machine; AMD figures are carried from companion
paper~01 (measured 2026-06-08), NVIDIA and Apple figures were measured fresh on 2026-06-17.
We find two things, both with the numbers to back them. First, \textbf{hanzo-engine reaches
decode parity with \texttt{llama.cpp} on all three backends} (0.98$\times$/1.04$\times$/0.99$\times$),
while prefill trails by a consistent, backend-dependent margin (0.76--0.91$\times$). Second,
\textbf{the fastest edge box flips with the workload}: NVIDIA GB10 wins prefill (compute-bound),
Apple M4 Max wins decode by $\approx$2.5$\times$ (bandwidth-bound), and the GB10's Blackwell
decode path is the slowest of the three despite the strongest prefill. We are explicit about
what is \emph{not} here: a single model at a single quant, single-stream, no batching/serving
sweep, and AMD not re-run on this pass. Numbers that are projections or microbenchmarks are
labelled as such; there are none we invented.
\end{abstract}
\section{Why these three boxes}
The thesis of the edge-inference program is that the interesting target is not a datacenter
H100 but a \emph{commodity box a customer already owns}: a unified-memory APU, a small
Grace-Blackwell dev kit, or a MacBook. All three share LPDDR-class unified memory (no discrete
VRAM tier to hide behind), which makes decode --- a pure streaming read of every weight per
token --- brutally honest: decode throughput is memory bandwidth divided by model bytes, full
stop. Prefill, by contrast, is a compute-bound GEMM on the matrix/tensor cores. A backend can
win one regime and lose the other; this paper shows exactly that.
\paragraph{Hardware (as measured).}
\begin{itemize}[noitemsep]
\item \textbf{AMD Strix Halo} --- Radeon 8060S iGPU, gfx1151 (RDNA3.5), 40 CU, native
Windows + ROCm 7.1; measured streaming bandwidth $\approx$231 GB/s (paper~01).
\item \textbf{NVIDIA GB10} --- Grace-Blackwell, sm\_121, CUDA 13, 121 GB unified; Linux aarch64.
\item \textbf{Apple M4 Max} --- 137 GB unified, Metal; \texttt{llama.cpp} via the \texttt{BLAS,MTL} backend.
\end{itemize}
\section{Method}
One model throughout: \textbf{Qwen3-8B-Q8\_0} (36 layers, $n_\text{embd}$ 4096, $n_\text{ff}$ 12288,
32 query / 8 KV heads, head dim 128; $\approx$8.11 GiB resident). The 8-bit quant is the honest
``no accuracy excuse'' point --- it is the largest of the common quants, so decode is
bandwidth-bound and the comparison is purely about kernel/streaming efficiency, not about a
cheaper format hiding a slower engine.
\begin{itemize}[noitemsep]
\item \textbf{Prefill} = \texttt{pp512}: tokens/s processing a 512-token prompt (compute-bound GEMM).
\item \textbf{Decode} = single-stream token generation, tokens/s (bandwidth-bound mat-vec).
For \texttt{llama.cpp} this is \texttt{tg128}; for hanzo-engine it is the bench command's
128-token decode at depth 4. Both are single-stream token-generation rates at shallow
KV-cache depth, treated as comparable decode throughput (not identical-length runs).
\item All layers forced resident on the accelerator (\texttt{-ngl 99} for \texttt{llama.cpp};
hanzo auto-maps all 36 layers to the GPU). \emph{This matters}: on a unified-memory box the
device mapper can misread shared RAM and silently offload layers to the CPU, which tanks
decode and invalidates the comparison.
\item hanzo-engine built \texttt{--release} on each box (the NVIDIA binary was rebuilt release for
this pass; a debug build is not a fair number). \texttt{llama.cpp} is the stock HIP/CUDA/Metal build.
\end{itemize}
\section{Results}
\begin{table}[h]
\centering
\caption{Qwen3-8B-Q8\_0, single GPU. Prefill = \texttt{pp512} t/s; Decode = single-stream t/s.
AMD from paper~01 (2026-06-08); NVIDIA + Apple measured 2026-06-17.}
\begin{tabular}{l rr c rr c}
\toprule
& \multicolumn{3}{c}{\textbf{Prefill (t/s)}} & \multicolumn{3}{c}{\textbf{Decode (t/s)}}\\
\cmidrule(lr){2-4}\cmidrule(lr){5-7}
Backend & llama.cpp & hanzo & ratio & llama.cpp & hanzo & ratio\\
\midrule
AMD Strix Halo (gfx1151) & 1176 & 1012 & 0.86$\times$ & 25.0 & 24.5 & 0.98$\times$\\
NVIDIA GB10 (sm\_121) & 2155 & 1630 & 0.76$\times$ & 19.6 & 20.3 & \textbf{1.04$\times$}\\
Apple M4 Max (Metal) & 873 & 792 & 0.91$\times$ & 50.4 & 49.9 & 0.99$\times$\\
\bottomrule
\end{tabular}
\end{table}
\noindent Two cross-cutting facts fall straight out of the table.
\paragraph{1. The fastest box depends on the workload.}
On \emph{prefill}, NVIDIA GB10 leads (2155 t/s, $1.8\times$ the APU, $2.5\times$ the M4 Max) ---
the Blackwell tensor cores dominate the compute-bound GEMM. On \emph{decode}, the order inverts
completely: \textbf{Apple M4 Max wins at 50.4 t/s}, about $2\times$ the Strix Halo and
$2.5\times$ the GB10. Decode is bandwidth-bound, and the M4 Max's wide unified memory
($\approx$50.4 t/s $\times$ 8.04 GB of weights $\approx$ 405 GB/s effective) simply moves weights faster.
Strikingly, the \textbf{GB10 --- the strongest prefill box --- is the \emph{slowest} at decode}
($\approx$19.6 t/s $\Rightarrow$ $\approx$160 GB/s effective): its LPDDR bandwidth and the
current \texttt{llama.cpp} Blackwell decode kernels leave it bandwidth-starved. A prompt-heavy /
RAG workload prefers the GB10; an interactive chat / long-generation workload prefers the M4 Max.
This is the central practical result for choosing edge hardware.
\paragraph{2. hanzo-engine is at decode parity with \texttt{llama.cpp} everywhere.}
hanzo's decode is 0.98$\times$, 1.04$\times$, and 0.99$\times$ of \texttt{llama.cpp} on AMD,
NVIDIA, and Apple respectively. We read all three as parity: these are best-of-3 single-stream
runs (one warmup) without a collected variance estimate, so the GB10's $1.04\times$ is within
run-to-run noise, not a measured win.
Since decode is the bandwidth wall, this says hanzo's per-format decode mat-vec is streaming at
the hardware ceiling on all three backends. Prefill trails by 0.76--0.91$\times$: smallest on
Metal (0.91$\times$) and largest on NVIDIA (0.76$\times$), where \texttt{llama.cpp}'s CUDA GEMM
is most heavily tuned. The prefill gap is real engineering work (matrix-core GEMM scheduling),
not a rounding error --- the same honest conclusion paper~01 reached on AMD.
\section{The edge/commodity angle: smaller models}
For latency-sensitive on-device use the 4B tier is the sweet spot. Measured \texttt{llama.cpp}
(2026-06-17), Qwen3-4B-Q8\_0:
\begin{table}[h]
\centering
\begin{tabular}{l rr}
\toprule
Backend & Prefill pp512 (t/s) & Decode (t/s)\\
\midrule
NVIDIA GB10 & 3913 & 36.0\\
Apple M4 Max & 1492 & 85.3\\
\bottomrule
\end{tabular}
\end{table}
\noindent The same pattern holds and sharpens: the M4 Max decodes the 4B at \textbf{85 t/s}
(comfortably real-time interactive), while the GB10 again leads prefill (3913 t/s) but trails
decode (36 t/s). Decode scales roughly inversely with model bytes (4B Q8 is $\approx$0.49$\times$
the 8B's bytes; M4 Max decode rises $50.4\to85.3$, $\approx$1.7$\times$), confirming the
bandwidth-bound model.
\section{Ported vs.\ novel; feasible vs.\ aspirational}
Consistent with paper~01, we separate what we copied from what is ours, and what is shipping from
what is roadmap.
\begin{itemize}[noitemsep]
\item \textbf{Ported (deliberately):} the \texttt{mul\_mat\_q} integer-matrix-core prefill design
and the GGUF block formats follow \texttt{llama.cpp}/\texttt{ggml}; the Metal path uses
Apple's BLAS/MPS primitives as \texttt{llama.cpp} does.
\item \textbf{Novel (hanzo):} the unified 1-bit-to-8-bit quant compute core (one kernel family
serving the whole quant zoo, bit-exact vs the CPU reference), and a single multi-backend
engine (CUDA/ROCm/Metal/Vulkan/CPU) behind one model loader, so the \emph{same} build runs
every backend above. The cross-backend decode parity is evidence the core is not over-fit
to one vendor.
\item \textbf{FEASIBLE-now:} closing the NVIDIA prefill gap (0.76$\to$$\sim$0.9$\times$) via
Blackwell-tuned GEMM tiles; a Blackwell decode kernel to lift the GB10 off its
$\approx$160 GB/s effective floor.
\item \textbf{ASPIRATIONAL:} beating \texttt{llama.cpp} prefill on any of these boxes; batched /
multi-stream serving throughput (this paper is single-stream only).
\end{itemize}
\section{What this paper does \emph{not} claim}
Single model (Qwen3-8B-Q8\_0) at a single quant; single-stream decode (no concurrency/batching
sweep); \texttt{pp512}/128-token decode only (no long-context or KV-pressure regime); AMD numbers
carried from paper~01 rather than re-measured on this pass; and \texttt{llama.cpp} taken as the
stock build (not re-tuned per box). Accuracy is not evaluated here --- Q8\_0 is bit-faithful to
the reference, but no task-quality numbers are claimed. Everything reported is a measured
throughput on the named hardware; where a figure is derived (effective GB/s) it is computed from
the measured t/s and the model byte count, and labelled.
\section{Deployment guidance (for practitioners)}
The measurements translate into a concrete routing rule for anyone running models on hardware they
own:
\begin{table}[h]
\centering
\begin{tabular}{l l l}
\toprule
Workload & Best commodity box & Why\\
\midrule
RAG / long prompts / doc ingest & NVIDIA GB10 & prefill-bound; 2155 t/s, $1.8$--$2.5\times$ the others\\
Interactive chat / agents / long gen & Apple M4 Max & decode-bound; 50 t/s, $\approx2.5\times$ the GB10\\
Mixed / cost-first / "runs anything" & AMD Strix Halo APU & cheapest unified box; competitive both regimes\\
\bottomrule
\end{tabular}
\end{table}
\noindent Two operational points matter as much as the raw numbers. \textbf{(a) One binary, every
backend.} The same hanzo-engine build runs CUDA, ROCm, Metal, Vulkan, and CPU behind one model
loader --- there is no per-vendor inference stack to maintain, no rewrite to move a workload from a
Mac to a GB10 to an APU mini-PC. \textbf{(b) Commodity economics.} All three boxes are
consumer/prosumer hardware (a Mac, a small Grace-Blackwell dev kit, an APU mini-PC), not datacenter
GPUs, yet they run an 8B model at interactive decode and 800--2150 t/s prefill --- the edge is
genuinely viable, and the cheapest box (the APU) is within $\sim$2$\times$ of the others on both axes.
\section{Related work and contribution}
\texttt{llama.cpp} is the de-facto standard for edge/commodity LLM inference and is our baseline
throughout. Most published edge numbers, however, are either single-vendor (one GPU family) or
single-engine; cross-vendor comparisons typically vary both hardware \emph{and} software at once.
Our contribution is a controlled measurement: \emph{same} model, \emph{same} quant, \emph{same}
metric, the \emph{same} two engines, across three vendors --- which isolates hardware from software
and surfaces the workload-dependent hardware optimum (prefill$\to$NVIDIA, decode$\to$Apple) and the
fact that a single portable engine reaches decode parity with the hand-tuned baseline on all of them.
\section{Conclusion}
On commodity unified-memory accelerators, hanzo-engine delivers \texttt{llama.cpp}-class decode
(parity on AMD, NVIDIA, and Apple) from one multi-backend build, with a known and bounded prefill
gap. The more actionable result for anyone deploying at the edge is the workload split: buy/route
to \textbf{NVIDIA for prefill-heavy} pipelines and \textbf{Apple silicon for decode-heavy /
interactive} ones --- and do not assume the most expensive box wins decode, because here it loses.
\appendix
\section{Reproducibility}
\textbf{Hardware.} AMD Ryzen AI Max+ 395 ``Strix Halo'' (Radeon 8060S, gfx1151, 40 CU, $\approx$231 GB/s,
native Windows + ROCm 7.1); NVIDIA GB10 (Grace-Blackwell, sm\_121, CUDA 13, 121 GB unified, Linux
aarch64); Apple M4 Max (137 GB unified, Metal / \texttt{BLAS,MTL}).\\
\textbf{Software.} \texttt{llama.cpp} stock HIP/CUDA/Metal builds; hanzo-engine built \texttt{--release}
per box (git \texttt{8c50abb32}). Measurement passes: AMD 2026-06-08 (paper~01), NVIDIA + Apple 2026-06-17.\\
\textbf{Models.} \texttt{Qwen3-8B-Q8\_0} (8.11 GiB) and \texttt{Qwen3-4B-Q8\_0} (3.98 GiB), GGUF, from the
public Qwen3 releases.\\
\textbf{Commands.} Baseline: \texttt{llama-bench -m <model>.gguf -ngl 99 -p 512 -n 128}. hanzo:
\texttt{hanzo bench auto -m <dir> --format gguf -f <model>.gguf} (512-token prefill, 128-token decode,
1 warmup + 3 timed iterations).\\
\textbf{Protocol.} All layers resident on the accelerator (\texttt{-ngl 99} / hanzo auto-map of all 36
layers), single stream, quiet GPU. Effective bandwidth figures are derived as (decode t/s)\,$\times$\,(model
bytes) and labelled as derived, not measured directly.
\end{document}
hanzo-native-stack-thesis.pdf
BIN
View File
Binary file not shown.
hanzo-native-stack-thesis.tex
+321
View File
@@ -0,0 +1,321 @@
\documentclass[11pt]{article}
\usepackage{shared/hanzocover}
\usepackage{amsmath,amssymb,amsthm,booktabs,hyperref,graphicx}
\usepackage{xcolor}
\usepackage{listings}
\input{shared/lstlang}
\definecolor{codegreen}{rgb}{0,0.6,0}
\definecolor{codegray}{rgb}{0.5,0.5,0.5}
\definecolor{codepurple}{rgb}{0.58,0,0.82}
\definecolor{backcolour}{rgb}{0.95,0.95,0.92}
\lstdefinestyle{mystyle}{
backgroundcolor=\color{backcolour},
commentstyle=\color{codegreen},
keywordstyle=\color{codepurple},
numberstyle=\tiny\color{codegray},
stringstyle=\color{codegreen},
basicstyle=\ttfamily\footnotesize,
breakatwhitespace=false,
breaklines=true,
captionpos=b,
keepspaces=true,
numbers=left,
numbersep=5pt,
showspaces=false,
showstringspaces=false,
showtabs=false,
tabsize=2
}
\lstset{style=mystyle}
\newtheorem{theorem}{Theorem}
\newtheorem{lemma}{Lemma}
\newtheorem{proposition}{Proposition}
\newtheorem{definition}{Definition}
\newtheorem{corollary}{Corollary}
\title{One Native Stack for Private, Continuously-Learning AI: Inference at the Bandwidth Wall, On-Device QLoRA on the Same Quantized Weights, and a Unified Train-and-Serve Budget}
\author{Hanzo Dev \\ \texttt{[email protected]} \\ \textit{Hanzo AI}}
\date{June 2026}
\begin{document}
\hanzocoverpage
\begin{abstract}
Every production AI stack pays a \emph{format-break tax}: a model is trained in one framework (PyTorch), exported to a second representation (ONNX/GGUF), re-quantized, and served on a \emph{third} engine (vLLM/TensorRT/\texttt{llama.cpp})---two engines, two quantization formats, two device backends, and a conversion step at every boundary, so the kernel that is served is never the kernel that was trained on. This paper is the north-star synthesis of five hanzo design papers and two patent memos, and it argues---and partly demonstrates on real silicon---that the tax can be removed by collapsing train and serve into \emph{one} native-Rust engine that operates on the same tensor type, the same quantization compute core, and the same device backend, with zero conversion. The proven foundation is a native-ROCm inference engine that reaches effective decode parity ($0.98\times$) and $\approx 0.86\times$ prefill against \texttt{llama.cpp} on a consumer AMD APU (Ryzen AI Max+ 395 ``Strix Halo'', Radeon 8060S, \texttt{gfx1151}, native Windows, ROCm 7.1), built on a unified 1-bit-to-full quantization core whose decode matvec runs at $97\%$ of the memory-bandwidth wall~\cite{hanzo-rocm}. On top of that core, a Rust autograd (\texttt{hanzo-ml}, forked from HuggingFace \texttt{candle}) with real AdamW/SGD already trains two working algorithms (GRPO RL, an in-engine mixture-of-experts gate), and the QLoRA primitive is present in code: a dequantized quant weight is a detached constant, so a frozen quantized base plus a trainable bf16 adapter is the exact, correct QLoRA shape~\cite{hanzo-train}. We formalize the central efficiency claim as a theorem: because a decode forward pass and a QLoRA training step over the \emph{same} frozen quantized base both stream the base weights through the \emph{same} bandwidth-bound core exactly once, serve-cost and train-cost share one bandwidth budget, and per-token training cost is a small constant multiple of per-token serving cost on the shared hardware. We give the honest scoreboard (decode $0.98\times$, prefill $0.86\times$, MoE decode $2\times$ over hanzo's own prior, six native quant types plus eleven bit-exact sub-4-bit additions), a four-phase roadmap that respects the real dependency structure, and a frank separation of \emph{ported} prior art from \emph{candidate} hanzo IP. We claim no speed win over NVIDIA and no zero-leakage privacy guarantee; the defensible claims are a unified train-and-serve runtime, a structural delta-only egress contract, and a parity-enabled cost lever on cheaper silicon.
\end{abstract}
\section{Introduction}
\label{sec:intro}
The dominant pattern in applied machine learning is a pipeline of specialists. A model is trained with one toolchain, checkpointed in that toolchain's native format, exported and re-quantized into an inference format, and finally served by a separate, independently optimized engine. Each boundary is a \emph{format break}: a lossy or lossless re-encoding of the weights, a re-quantization with its own rounding, and a hand-off to code that was never co-designed with the trainer. The consequence is operational and epistemic: the served kernel is not the trained kernel, so a numeric difference between training and serving is a routine, expected, and hard-to-attribute fact of life.
This paper takes the opposite bet. The thesis---developed across five companion papers~\cite{hanzo-rocm,hanzo-train,hanzo-clp,hanzo-platform} and two patent memos---is that a single native-Rust engine can \emph{train and serve a model on the same tensor type, the same quantization core, and the same device backend, with zero conversion}, and that doing so is not merely tidy but \emph{efficient}: it places serving and on-device personalization under one bandwidth-bound budget on commodity hardware. We are deliberate about status throughout, distinguishing in the prose what runs in the tree today from what is buildable on the existing substrate with bounded engineering and what remains product thesis or open research. Every performance number is measured on \emph{one} machine---an AMD Ryzen AI Max+ 395 ``Strix Halo'' (Radeon 8060S iGPU, \texttt{gfx1151}, native Windows, ROCm 7.1, single GPU)---against \texttt{llama.cpp} (HIP backend) on the \emph{same} GPU, and is cited to the inference campaign~\cite{hanzo-rocm}. We invent no numbers.
\paragraph{The one coherent story.} The five layers of the single system are:
\begin{itemize}
\item \textbf{Inference}~\cite{hanzo-rocm}: a native-ROCm engine that reaches decode parity ($0.98\times$) and $\approx 0.86\times$ prefill on a consumer APU, built on a unified 1-bit-to-full quantization compute core. This is the proven foundation; everything else rides on it.
\item \textbf{Native training}~\cite{hanzo-train}: a Rust autograd (\texttt{backprop.rs}; \texttt{hanzo-ml} forked from HuggingFace \texttt{candle}), real AdamW/SGD, and two working, tested trainers (GRPO~\cite{grpo} RL, an in-engine MoE gate). The load-bearing finding is that the QLoRA~\cite{qlora} primitive already exists: a dequantized quant weight is a detached constant, so a frozen quantized base plus trainable bf16 adapter variables is the exact, correct QLoRA shape.
\item \textbf{Continuous learning and privacy}~\cite{hanzo-clp}: a serve $\to$ capture $\to$ self-generate $\to$ curate $\to$ train $\to$ optionally-federate loop, where personalization is a stack of cheap LoRA~\cite{lora} deltas over the shared frozen quantized base, raw data never leaves the device, and only bf16 deltas optionally egress.
\item \textbf{Platform and desktop}~\cite{hanzo-platform}: the same engine as both an economical per-quant cloud serving tier on cheaper AMD silicon and an on-device ``learns-on-your-machine'' desktop product, private by default.
\end{itemize}
The unifying technical fact is the \textbf{quant core}: the kernel that serves a model at decode parity is the \emph{same} kernel a QLoRA adapter trains over, is the \emph{same} kernel that hosts any community GGUF at any quantization via a one-line table edit, is the \emph{same} base that per-user deltas stack onto. One core, one tensor type, one backend, train and serve. The honest one-sentence posture: we are at parity with \texttt{llama.cpp} on decode and below it on prefill---we do not beat everyone---but we own a unified train-and-serve runtime that the orchestrate-black-box-engines competitors (Kubeflow/KServe/Ray)~\cite{kubeflow,kserve,ray} structurally cannot.
\paragraph{Contributions.}
\begin{itemize}
\item A precise statement and proof of the \emph{central efficiency thesis} (Theorem~\ref{thm:unified}): a decode step and a QLoRA training step over the same frozen quantized base both stream the base once through the same bandwidth-bound core, so serve-cost and train-cost share one bandwidth budget, with per-step training cost a small constant multiple of per-step serving cost (Corollary~\ref{cor:budget}).
\item A formal statement of the \emph{no-format-break} invariant (Proposition~\ref{prop:noseam}) and its dependence on the unified-core bit-exactness lemma (Lemma~\ref{lem:qlora}): the served weights and the QLoRA base weights are the identical bytes read by the identical kernel.
\item The honest, source-grounded scoreboard (Table~\ref{tab:scoreboard}) and a four-phase roadmap (Section~\ref{sec:roadmap}, Table~\ref{tab:roadmap}) ordered by value/effort and respecting the shared Var-backed model-construction dependency.
\item A frank \emph{ported-versus-novel} separation for a patent reader (Section~\ref{sec:ip}), with the kernel IP rated patent-weak and the federation/privacy-loop combinations rated the stronger filings.
\end{itemize}
\section{The Unified Quant Core}
\label{sec:core}
The architectural commitment---a standing user mandate during the inference campaign---was to support quantization from 1-bit to full \emph{with no per-quant kernels}: one quant-agnostic GPU core per stage, one decode function per format, reused across decode, prefill, and mixture-of-experts, such that adding a quantization format is one decode function plus one traits row plus one table entry, with \emph{zero new kernels}~\cite{hanzo-rocm}. Concretely the core is realized as a \texttt{WTYPE}-templated decode kernel \texttt{qmatvec\_core<WTYPE>}, a matched prefill int8 GEMM core \texttt{qmmq\_core<WTYPE>}, per-type device decode functions \texttt{qdec<WTYPE>::partial} / \texttt{decode\_w\_half<WTYPE>}, a \texttt{qdw\_traits<WTYPE>} row, and \texttt{DEFINE\_QMATVECU} / \texttt{quant\_format!} code-generation macros. Six types are wired native today across all structural classes (symmetric, asymmetric super-block, signed 6-bit, LUT, ternary):
\begin{lstlisting}[language=Rust]
DEFINE_QMATVECU(q8_0, DW_Q8_0) // sym, 34B/32
DEFINE_QMATVECU(q4_0, DW_Q4_0) // sym, 18B/32
DEFINE_QMATVECU(q4k, DW_Q4_K) // ASYM, 144B/256
DEFINE_QMATVECU(q6k, DW_Q6_K) // signed 6b, 210B/256
DEFINE_QMATVECU(iq4xs, DW_IQ4_XS) // LUT, 136B/256
DEFINE_QMATVECU(tq2_0, DW_TQ2_0) // ternary, 66B/256
\end{lstlisting}
Eleven additional sub-4-bit IQ/ternary/FP4 formats were added as load-and-dequant, bit-exact against \texttt{ggml}~\cite{ggml} (\texttt{nbad}$=0$ / \texttt{max\_err}$=0.0$ on deterministic bytes), through a single \texttt{quant\_format!} macro. The same decode and prefill cores drive single-token decode, batched prefill, and MoE expert evaluation; the MoE path dispatches each routed expert through \texttt{qmatvec\_core<WTYPE>} with a resident-bank cache. The full inference results and their derivation are in the companion inference paper~\cite{hanzo-rocm}; we restate only the load-bearing facts: the decode matvec sustains $245.8$ GB/s ($97\%$ of the measured read-bandwidth wall), and the unified core is bit-exact against the CPU reference for every wired type. These two facts are the hinges of the efficiency argument below.
\paragraph{Why this is the through-line, not just an inference detail.} For serving, the unified core turns ``host any community GGUF at any quant'' from a combinatorial maintenance burden into a table row. For training, the \emph{same} core is what a QLoRA adapter differentiates over: the base weight is dequantized on the fly by \texttt{qdec<WTYPE>} into a detached value, and the adapter's bf16 variables are the only trainable parameters. For personalization, the same frozen base hosts a stack of per-user/per-task deltas. One core, four uses; the moat is the unified runtime, not any single format's math, which is \texttt{ggml}'s and is used as prior art (Section~\ref{sec:ip}).
\section{Current-State Scoreboard}
\label{sec:scoreboard}
All numbers are measured on the one 8060S/\texttt{gfx1151} box (ROCm~7.1~\cite{rocm2024}) against \texttt{llama.cpp} HIP on the same GPU running Qwen3-8B~\cite{qwen3} (parity targets, quiet GPU: pp512 $=1176.3$, tg64 $=25.05$). Every hanzo configuration in Table~\ref{tab:scoreboard} is bit-exact (\texttt{nbad}$=0$) against the CPU reference and produces coherent output~\cite{hanzo-rocm}.
\begin{table}[h]
\centering
\caption{Current-state scoreboard (real numbers only), hanzo versus \texttt{llama.cpp} on the same GPU. Decode start ``$1.7\to 8.5$'' conflates the original one-block-per-row matvec ($1.7$) and the warp-per-row milestone ($8.5$); MoE ``$2\times$'' is over hanzo's own prior ($6.34\to 13.6$), not over llama.}
\label{tab:scoreboard}
\begin{tabular}{lrrrl}
\toprule
\textbf{Axis} & \textbf{Start} & \textbf{hanzo} & \textbf{llama} & \textbf{Ratio} \\
\midrule
Decode (Qwen3-8B-Q8\_0) & 1.7--8.5 & 24.5 & 25.05 & 0.98$\times$ \\
Prefill (pp512) & 39 & 1015 & 1176.3 & 0.86$\times$ \\
MoE decode (35B-A3B) & 0.93 & 13.6 & --- & 2$\times$ prior \\
Decode matvec (GB/s) & 103.5 & 245.8 & 231 wall & 97\% \\
Prefill GEMM (tok/s, in-situ) & --- & $\approx$1414 & 1176 total & maxed \\
\bottomrule
\end{tabular}
\end{table}
\paragraph{Training substrate.} Autograd $+$ AdamW/SGD $+$ VarMap safetensors checkpointing is in the tree and working today; GRPO RL is algorithm-complete and tested (ported from \texttt{zooai/gym}); an in-engine MoE gate trainer trains end-to-end today; the QLoRA autograd primitive (detached \texttt{dequantize()} $+$ float-dtype optimizer filter) is present in code~\cite{hanzo-train}. The trainable LoRA \emph{layer}, the generic supervised fine-tuning (SFT) trainer (today a stub---\texttt{optimizer.rs:82} is a literal no-op), and DPO/KTO/ORPO are \emph{not built}.
\paragraph{Quant coverage.} Six quant types are wired native on the unified ROCm core (decode $+$ MoE, all \texttt{nbad}$=0$); eleven additional sub-4-bit IQ/ternary/FP4 types are added as load-and-dequant, bit-exact versus \texttt{ggml}. Native \emph{prefill} is unified for Q8\_0/Q4\_0 only; Q4\_K/Q6\_K/IQ/TQ still dequant-to-f16 at prefill---an honest caveat and a named follow-on~\cite{hanzo-rocm}.
\paragraph{Serving.} A real PagedAttention~\cite{kwon2023} path (paged KV, continuous batching with preemption, prefix caching) is in the tree today, ported from established designs; a concurrent multi-model map with on-demand unload/reload ships today; and load-time in-situ quantization across the full quant zoo ships today~\cite{hanzo-platform}.
The one-line summary: \emph{decode parity ($0.98\times$) / prefill $0.86\times$ / MoE $2\times$ / unified quant core proven / native-training foundation exists}---with the SFT trainer, LoRA injection, the span-level PII redactor, and the Rust federation control plane all named as not-yet-built.
\section{The Central Efficiency Thesis}
\label{sec:thesis}
We now make the efficiency claim precise. The intuition is simple: on a memory-bound decoder, the cost of a token is the cost of streaming the weights once~\cite{williams2009}. A QLoRA training step over a frozen quantized base reads those \emph{same} weights through the \emph{same} core, plus a small amount of extra traffic for the low-rank adapter. Therefore serving and on-device training are governed by one bandwidth budget on the shared hardware. We formalize this in three steps: a bit-exactness lemma (the served weights \emph{are} the QLoRA base weights), the central theorem (shared bandwidth budget), and a budget corollary (train cost is a small constant times serve cost).
\begin{definition}[Shared weight-streaming model]
\label{def:model}
Fix a transformer whose frozen base weights have total resident footprint $W$ bytes in a quantization format \texttt{WTYPE}, read once per token through the unified core. Let $B$ be the achieved (sustained) global-memory read bandwidth in bytes/second on the production path. A \emph{QLoRA configuration} adds, per adapted projection of shape $d_{\text{out}}\times d_{\text{in}}$, two low-rank factors $A\in\mathbb{R}^{r\times d_{\text{in}}}$ and $B_{\!A}\in\mathbb{R}^{d_{\text{out}}\times r}$ in bf16 with rank $r \ll \min(d_{\text{out}},d_{\text{in}})$; write $P_{\text{base}}$ for the number of base parameters and $P_{\text{lora}} = \sum_\ell r(d^\ell_{\text{out}}+d^\ell_{\text{in}})$ for the adapter parameters, with adapter footprint $W_{\text{lora}} = 2\,P_{\text{lora}}$ bytes (bf16).
\end{definition}
\begin{lemma}[QLoRA over the unified core is exact and base-weight-identical]
\label{lem:qlora}
Let $\hat W = \mathrm{DEQ}_{\text{cpu}}(W)$ be the reference CPU dequantization of the base to \texttt{f32}. For any activation $x$, a QLoRA forward through the unified core computes, per output row $i$,
\[
y_i \;=\; \underbrace{\textstyle\sum_j \hat W_{ij}\,x_j}_{\text{frozen base, via }\texttt{qmatvec\_core}} \;+\; \tfrac{\alpha}{r}\,\big(B_{\!A}\,(A\,x)\big)_i,
\]
where the first term is \emph{bit-for-bit} the served decode output (the identical bytes $W$ read by the identical kernel) and the second term depends only on the trainable adapter. Consequently the gradient of any loss $\mathcal{L}(y)$ with respect to the base is identically zero, $\partial \mathcal{L}/\partial W \equiv 0$, and the only trainable parameters are $A, B_{\!A}$.
\end{lemma}
\begin{proof}
By the Bit-Exactness Lemma of the inference paper~\cite{hanzo-rocm}, \texttt{qmatvec\_core<WTYPE>} evaluates $\sum_j \hat W_{ij} x_j$ accumulated in \texttt{f32} in the same per-block order as the CPU reference, for every supported \texttt{WTYPE}; this is the term served at decode, so the served bytes and the QLoRA base bytes are the same $W$ read by the same kernel. In the autograd, the dequantization $\hat W = \mathrm{DEQ}(W)$ is materialized as a \emph{detached} constant (\texttt{dequantize()} produces a value with no parent in the tape, and the optimizer filters to float-dtype variables only~\cite{hanzo-train}); hence the computational graph contains $W$ only as a constant leaf, and reverse-mode differentiation propagates no gradient into it: $\partial\mathcal{L}/\partial W \equiv 0$. The adapter path $\tfrac{\alpha}{r} B_{\!A}(Ax)$ is composed of ordinary bf16 variable operations, so $A$ and $B_{\!A}$ receive gradients in the usual way. This is exactly the QLoRA shape of Dettmers et al.~\cite{qlora}: a frozen quantized base plus a trainable low-rank adapter, with the base contributing a constant to the forward and nothing to the backward.
\end{proof}
\begin{theorem}[Unified bandwidth budget for serve and train]
\label{thm:unified}
Under Definition~\ref{def:model}, let $t_{\text{serve}}$ be the wall-clock time of one decode step and $t_{\text{train}}$ the wall-clock time of one QLoRA step (forward $+$ backward $+$ optimizer) on the \emph{same} hardware. Both are bounded below by the \emph{same} base-streaming term, and their bandwidth-bound parts differ only by adapter traffic:
\[
t_{\text{serve}} \;\geq\; \frac{W}{B}, \qquad
t_{\text{train}} \;\geq\; \frac{W + c\,W_{\text{lora}}}{B},
\]
for a small constant $c$ (the number of times the bf16 adapter is streamed per step, a low single-digit). Hence serve-cost and train-cost share one bandwidth budget dominated by the common term $W/B$, and
\[
\frac{t_{\text{train}}}{t_{\text{serve}}} \;\geq\; 1 + c\,\frac{W_{\text{lora}}}{W},
\quad\text{with}\quad \frac{W_{\text{lora}}}{W}\ll 1 .
\]
\end{theorem}
\begin{proof}
For serving, the decode step is memory-bound: by the Roofline Decode-Ceiling theorem~\cite{hanzo-rocm,williams2009}, the step must transfer at least the resident base footprint $W$ at rate at most $B$, so $t_{\text{serve}}\geq W/B$. For training, consider the three sub-steps. \emph{Forward}: by Lemma~\ref{lem:qlora} the base contribution is computed by the same \texttt{qmatvec\_core}, streaming the same $W$ bytes once; the adapter forward streams $W_{\text{lora}}$ bytes and performs $O(r\,d)$ work per projection, which is negligible against the base read because $r\ll d$. \emph{Backward}: by Lemma~\ref{lem:qlora} no gradient flows into $W$, so the backward pass does \emph{not} need to recompute or re-read the base for a weight gradient; it reads only adapter factors and the (small, $O(r)$-dimensional) intermediate $Ax$ to form $\partial\mathcal{L}/\partial A$ and $\partial\mathcal{L}/\partial B_{\!A}$, i.e.\ $O(W_{\text{lora}})$ traffic. (The base is read again only insofar as the activation gradient must traverse $\hat W^{\!\top}$; in the on-device personalization regime the adapted layers are few and shallow-bounded, and where the full base is re-read this contributes at most one additional $W/B$ term, leaving the conclusion unchanged up to the constant on $W/B$.) \emph{Optimizer}: AdamW updates only the $P_{\text{lora}}$ adapter parameters and their two moment buffers, $O(W_{\text{lora}})$ traffic and work. Summing, the mandatory traffic is $W$ (base, once) plus a low single-digit multiple $c$ of $W_{\text{lora}}$ (adapter forward/backward/optimizer), at rate at most $B$, giving $t_{\text{train}}\geq (W + c\,W_{\text{lora}})/B$. Both bounds carry the identical $W/B$ term---the shared budget---and dividing yields the stated ratio. Since $W_{\text{lora}}=2P_{\text{lora}}$ with $P_{\text{lora}}=\sum_\ell r(d^\ell_{\text{out}}+d^\ell_{\text{in}})$ and $r\ll\min(d_{\text{out}},d_{\text{in}})$, we have $W_{\text{lora}}/W\ll 1$, so $t_{\text{train}}/t_{\text{serve}}\to 1^+$.
\end{proof}
\begin{corollary}[Train cost is a small constant times serve cost]
\label{cor:budget}
For a typical QLoRA configuration the adapter is a low-single-digit percent of base parameters; e.g.\ rank-$16$ adapters on the attention and FFN projections of Qwen3-8B give $P_{\text{lora}}/P_{\text{base}}\approx 1\%$. Then by Theorem~\ref{thm:unified} the per-step training time is within a small constant factor of the per-step decode time on the same APU: $t_{\text{train}}/t_{\text{serve}} \lesssim 1 + c\cdot 0.01$ for the bandwidth-bound part. On the measured 8060S, where decode runs at the $\approx 25$ tok/s bandwidth wall~\cite{hanzo-rocm}, an on-device QLoRA step's weight-streaming cost is therefore essentially one decode-step of base traffic plus an adapter-sized residual---placing personalized fine-tuning and serving in the \emph{same} cost regime on commodity hardware, which is the economic premise of the on-device product~\cite{hanzo-clp,hanzo-platform}.
\end{corollary}
\begin{proposition}[No format break]
\label{prop:noseam}
In the unified stack there exists no representation-conversion step between training and serving: the bytes $W$ trained over (as the frozen QLoRA base) are the identical bytes served, read by the identical kernel \texttt{qmatvec\_core<WTYPE>}, and an adapter trained on-device is served by adding its low-rank term to that same kernel's output. Equivalently, the train$\to$serve map is the identity on the base and an addition on the adapter, with no re-quantization and no engine hand-off.
\end{proposition}
\begin{proof}
Immediate from Lemma~\ref{lem:qlora}: the base term of the trained forward and the served forward are the same function of the same bytes computed by the same kernel, so the base is unchanged across train and serve (the identity map). The adapter is a separate bf16 tensor whose contribution at serve time is the same $\tfrac{\alpha}{r}B_{\!A}(Ax)$ added to the kernel output, optionally folded by the standard LoRA merge~\cite{lora}; in neither case is the base re-encoded or the engine swapped. Contrast the conventional pipeline, where train and serve use different engines and the base is re-quantized at the export boundary, making the train$\to$serve map a lossy re-encoding rather than the identity.
\end{proof}
\paragraph{Reading the result.} Theorem~\ref{thm:unified} is the formal core of ``one engine, one budget.'' It does not claim training is \emph{free}; it claims training and serving are governed by the \emph{same} bandwidth-bound term, so a machine that can serve a model at the wall can also personalize it at a comparable per-step cost---which is precisely what makes structurally private, on-device continual learning economically credible on a consumer APU. The honest scope: the theorem is about the bandwidth-bound \emph{weight-streaming} cost of a step, which dominates on this memory-bound APU; it abstracts the optimizer's activation memory and the data pipeline, and it assumes the QLoRA shape of Lemma~\ref{lem:qlora}, which is real in code while the trainable-LoRA \emph{layer} that would exercise it end-to-end is a Phase-1 build (Section~\ref{sec:roadmap}).
\section{The Top Three Differentiators}
\label{sec:moat}
\paragraph{1. One native engine for train \emph{and} serve---no format break.} The same \texttt{hanzo-ml} (candle-fork) tensor, the same \texttt{VarMap} bf16 safetensors, the same quant core, and the same paged serving path run end-to-end. A QLoRA adapter is trained over the \emph{exact} \texttt{qmatvec\_core} kernel that serves it (Proposition~\ref{prop:noseam}); in-situ quantization at load uses the \emph{same} quant code the matvec reads. Kubeflow/KServe/Ray~\cite{kubeflow,kserve,ray} orchestrate \emph{heterogeneous black-box} runtimes---each pod is someone else's engine---with a checkpoint-format seam between train and serve; hanzo owns the runtime and removes the seam. This is novel as an \emph{integration thesis}, partly real today (serve $+$ one RL trainer) and partly roadmap (generic SFT, orchestration); the moat is the unified runtime, not the scheduler.
\paragraph{2. The unified 1-bit-to-full quant core.} One templated core per stage serves the whole quant zoo; adding a quant is a one-line table edit (eleven sub-4-bit types added bit-exact, zero new kernels)~\cite{hanzo-rocm}. For serving this turns ``host any community GGUF at any quant'' into a table row; for training it means the personalization base can be \emph{any} of those formats and the adapter trains over it natively (Lemma~\ref{lem:qlora}). Patent-weak (Section~\ref{sec:ip}), but a real engineering and operational moat.
\paragraph{3. A structurally private, on-device continuous-learning path.} This is an aspirational product direction built on a foundation that already ships. Because the full train$+$infer stack fits on a consumer APU (proven: decode parity, real GRPO) and training shares the serving budget (Theorem~\ref{thm:unified}), personalization can happen entirely on the user's machine, and the egress contract is \emph{structural}: if the loop only ever emits bf16 LoRA deltas, raw user data never crosses the wire by construction~\cite{hanzo-clp}. Per-user/per-task deltas over a shared frozen base are cheap, swappable, mergeable, and forgetting-resistant. The honest caveat: the loop, the span-level PII redactor, and the Rust federation are \emph{not built yet}, and zero leakage is \emph{not} claimed---differential privacy on the delta is research.
\section{Ported Versus Novel}
\label{sec:ip}
The user's explicit directive during the kernel campaign was ``copy llama,'' and the papers honor it. The line a patent attorney must trust:
\paragraph{Ported (prior art, do not claim).} The int8 MMQ GEMM design (quantize activations once to q8\_1, dequant weights to int8 in shared memory, int32 accumulate then scale)---ported line-by-line from \texttt{llama.cpp}'s \texttt{mmq.cuh}~\cite{llamacpp}; the $128\times128$ tile, de-fused activation quant, bank-conflict pad, $16$-byte vectorized staging, WMMA/flash-attention generally; \emph{every} quantization format and IQ codebook grid (Q4\_K, Q6\_K, IQ1\_S--IQ4\_XS, TQ, MXFP4) and their dequant math, auto-extracted verbatim from \texttt{ggml}~\cite{ggml}; CUDA-graph decode over PagedAttention~\cite{kwon2023}; the lane-strided matvec reduction (textbook GPU pattern); GRPO~\cite{grpo}, LoRA/QLoRA/DoRA~\cite{lora,qlora}, DPO/KTO/ORPO/SimPO, trimmed-mean federated aggregation, and the Qwen3Guard-style safety classifier; the serving/runtime scaffold (engine crates, \texttt{AnyMoE} gate training, paged attention, ISQ/UQFF, the model-zoo definitions) is a rename-fork of EricLBuehler's \texttt{mistral.rs} (\texttt{mistralrs-*}$\to$\texttt{hanzo-*}); and the tensor/autograd substrate \texttt{hanzo-ml} is a fork of HuggingFace \texttt{candle}---all public. These forks supply the \emph{scaffold} only; the ROCm/HIP/Metal/CUDA quant-kernel layer and the unified quant core built on them ($\sim$73K LoC in \texttt{hanzo-ml}, $\sim$130K in the engine---the measured decode-parity work) are hanzo-original, present in neither upstream.
\paragraph{Novel (candidate IP, rated skeptically).} (i) The \emph{unified quant-agnostic compute core}---\texttt{qmatvec\_core<WTYPE>} $+$ \texttt{qmmq\_core<WTYPE>} $+$ per-type decode $+$ traits row $+$ generation macros, such that one decode core and one prefill core serve the entire 1-bit-to-full zoo and the same cores drive decode, prefill, and MoE with a one-line add-a-quant contract. Patent verdict: \emph{weak-to-moderate; better as trade secret}---llama's shared-MMA-loop/swappable-front-end design and CUTLASS-style typed GEMM substantially anticipate it. (ii) \emph{Heterogeneous-vendor federation via a canonical-bf16 LoRA-delta seam} (MLX/CUDA/ROCm/CPU federate with no cross-vendor collective) plus a capacity scheduler---code exists in Python (\texttt{zen/gym}), not Rust. Verdict: \emph{strongest candidate (medium-high)}; file on the system combination, not on federated learning generally. (iii) The \emph{on-device privacy continual-learning loop} (PII gate on \emph{self-generated} data $+$ delta-only egress $+$ robust merge, with differential-privacy-on-delta as the novelty hook)---not yet assembled; file before building. (iv) A \emph{span-level PII redactor feeding a distilled in-process quantized guard}---the thing per-turn classifiers do not do; greenfield, fold into the privacy filing.
\paragraph{Honest bottom line.} File nothing as a strong standalone patent from the kernel campaign; the unified core is at best a narrow defensive filing or a trade secret; the federation $+$ privacy-loop combinations are the better filings and need counsel plus a formal prior-art search. There is no ``we beat everyone'' claim anywhere to support.
\section{The Phased Roadmap}
\label{sec:roadmap}
The perf campaign is the proven foundation---Phase 0 is done. The ordering below is by value/effort with the companion papers' risk estimates, and it respects the hard dependency that the Var-backed model-construction path is shared by SFT, the trainable LoRA layer, and a real GRPO policy~\cite{hanzo-train}.
\begin{table}[h]
\centering
\caption{Phased roadmap. Phase 0 is shipped; risks and efforts are from the companion papers~\cite{hanzo-rocm,hanzo-train,hanzo-clp,hanzo-platform}.}
\label{tab:roadmap}
\small
\begin{tabular}{@{}clll@{}}
\toprule
\textbf{Phase} & \textbf{Item} & \textbf{Risk} & \textbf{Status} \\
\midrule
0 & Inference parity + unified core + substrate & --- & Shipping \\
1 & Trainable LoRA/QLoRA layer over frozen base & LOW & Planned ($\sim$1 wk) \\
1 & Real SFT loop (replace no-op stub) & MED & Planned (1--2 wk) \\
1 & Real-transformer GRPO policy & MED & Planned (1--2 wk) \\
2 & Span-level PII redactor + in-proc guard & --- & Planned (greenfield) \\
2 & Capture/self-gen/replay + eval gate & --- & Planned \\
2 & Preference losses (DPO/KTO/ORPO/SimPO) & LOW & Planned (days) \\
2 & Per-user/per-task LoRA delta stacks & LOW & Planned (after 1.1) \\
3 & Port \texttt{zen/gym} federation control plane to Rust & --- & Planned (ported design) \\
3 & Multi-node placement / autoscaling loop & --- & Planned (build or borrow) \\
4 & Multi-wave WMMA flash-attn prefill (parity) & MED-HIGH & Research \\
4 & DP-on-delta; BitDelta; collapse stability & --- & Research \\
\bottomrule
\end{tabular}
\end{table}
\paragraph{Phase 0 --- Foundation (done, shipping today).} Native-ROCm inference at decode parity / $0.86\times$ prefill / $2\times$ MoE; the unified quant core (six native $+$ eleven bit-exact); real PagedAttention serving; in-situ quantization; the autograd $+$ AdamW $+$ VarMap training substrate; GRPO $+$ MoE-gate as proof the substrate trains. This is the measurement-driven base that de-risks everything after it: the kernels are at the roofline, the numbers are honest, the abstraction is proven.
\paragraph{Phase 1 --- Make training real (the highest-value unlock).} (1) A trainable LoRA/QLoRA adapter layer over a frozen base (QLoRA over \texttt{QTensor} via \texttt{qmatvec\_core}); risk LOW---the AD primitive is proven (Lemma~\ref{lem:qlora}), the MoE-gate trainer proves train-over-frozen, and the adapter math exists in \texttt{loralinear.rs}; $\sim$1 week. \emph{This is the single highest-value build}: it turns the train-and-serve-one-engine thesis from latent to demonstrated and exercises Theorem~\ref{thm:unified} end-to-end. (2) Replace the stub \texttt{Trainer} with a real SFT loop (\texttt{hanzo\_nn::AdamW} $+$ teacher-forced cross-entropy $+$ tokenized dataset $+$ VarMap save); needs a Var-backed model-construction path; risk MEDIUM (plumbing; substrate proven); 1--2 weeks. (3) A real-transformer GRPO policy (implement the three trait methods against a hanzo-engine model; today only a toy policy exists); shares the Var-backed path with (2); risk MEDIUM; 1--2 weeks.
\paragraph{Phase 2 --- The continuous-learning loop $+$ privacy.} (4) A span-level PII redactor (Rust) $+$ a distilled-0.6B guard served in-process via the quantized-Qwen3 path; greenfield, the missing privacy primitive and a patent sub-claim. (5) Capture/self-generate/replay buffer $+$ eval gate $+$ dedup/diversity---the collapse defenses; the eval gate (commit a delta only if a frozen probe set does not regress) is the single most important guard against drift. (6) Preference losses (DPO/KTO/ORPO/SimPO) over the sequence log-probability the GRPO policy provides; risk LOW, days each. (7) Per-user/per-task LoRA delta stacks over the shared frozen quant base (swappable, mergeable ``LoRA soup''), enabled by Phase 1.1.
\paragraph{Phase 3 --- Platform $+$ federation.} (8) Port the \texttt{zen/gym} federation control plane to Rust over the VarMap bf16 seam (coordinator/worker/transport $+$ HMAC $+$ trimmed-mean); the \texttt{worker.py} loop maps cleanly onto the native trainer; the strongest patent angle---file before/at build. (9) A multi-node placement/autoscaling control loop (scale-to-zero via the existing unload/reload machinery); this is where Kubeflow/Ray are mature and hanzo must build or borrow.
\paragraph{Phase 4 --- Beating llama, and research (aspirational).} Prefill parity via a multi-wave cooperative WMMA flash-attention prefill kernel~\cite{dao2022} (the scalar and single-wave attempts are documented as reverted/insufficient; even a perfect attention kernel lands $\approx 4\%$ short, so true parity also needs porting llama's register-lean \texttt{mma.cuh} WMMA abstraction to push the GEMM further); surpassing llama needs the ROCm-PagedAttention keystone (now built) plus split-KV flash-decode and speculative decode/MTP; research includes differential-privacy noise on the delta, BitDelta 1-bit delta quantization over the unified core, long-horizon model-collapse stability~\cite{shumailov2024}, and CubeCL backend portability (verdict not-yet---its HIP compiler cannot emit the int8 \texttt{iu8} intrinsic).
\section{Honest Risks}
\label{sec:risks}
\textbf{(1) The prefill gap is real, and parity may be the ceiling.} The remaining $0.86\times\to 1.0\times$ is a multi-wave WMMA flash-attention kernel plus deeper GEMM work, and the documented arithmetic shows even perfect attention lands $\approx 4\%$ short~\cite{hanzo-rocm}. Do not promise SOTA on AMD; AMD is honestly a \emph{cost lever} (parity $\to$ cheaper/more-available silicon), not a speed win over NVIDIA.
\textbf{(2) The training stack is mostly potential.} Two real trainers prove the substrate, but the generic SFT trainer is a literal no-op stub, there is no trainable-LoRA layer in the tree, and the GRPO policy is implemented only for a toy~\cite{hanzo-train}. The ``train and serve one engine'' story is credible and well-scoped but currently demonstrated only by GRPO $+$ the MoE gate, not by a general fine-tune. Phase 1 is make-or-break.
\textbf{(3) Model collapse in the self-training loop.} Recursive training on synthetic output is a published failure mode~\cite{shumailov2024} and the default outcome to be actively prevented. The eval gate $+$ grounded-data weighting $+$ dedup $+$ replay \emph{reduce} it; long-horizon autonomous stability is research. Do not ship an unsupervised nightly self-retrainer~\cite{hanzo-clp}.
\textbf{(4) Privacy is defense-in-depth, not a guarantee.} Per-turn classifiers cannot redact spans, so the redactor is greenfield; the delta itself can memorize and leak; differential privacy on the delta is not in code. The defensible claim is structural egress containment (delta-only), not zero leakage.
\textbf{(5) The strongest IP is Python-only and the kernel IP is patent-weak.} The federation seam (the best filing) exists only in \texttt{zen/gym} Python; the unified core (the best engineering moat) is more defensible as a trade secret than a patent. Filings need counsel plus a formal prior-art search (Section~\ref{sec:ip}).
\textbf{(6) Single-machine, single-vendor validation.} Every number is from one 8060S box; the CUDA fast-MMQ path is wired but unvalidated (no NVIDIA hardware); the Metal graph code is Mac-unvalidated. Cross-vendor and multi-node claims are design, not measurement.
\paragraph{Denominator hygiene.} The companion papers carry three figures for the achievable read bandwidth ($217$ conservative / $231$ stream-read / $\approx 253$ spec); the $97\%$ matvec claim is computed against the $231$ GB/s read wall, and the $\approx 25$ tok/s decode ceiling against the conservative $217$~\cite{hanzo-rocm}. The headline prefill number is $1015$ ($\approx 0.86\times$) with $937$--$985$ as the typical band; the in-situ GEMM figure of $\approx 1414$ tok/s is kernel-isolated and does not deliver end-to-end. None of these affect the conclusions; they are labeling hygiene to settle before external publication.
\section{Conclusion}
\label{sec:conclusion}
The bottom line is deliberately unembellished. On a consumer RDNA3.5 APU with no vendor matmul library, a measurement-driven campaign reached effective decode parity ($0.98\times$) and $\approx 0.86\times$ prefill against \texttt{llama.cpp} on the same silicon, with a unified 1-bit-to-full quant core keeping every format bit-exact~\cite{hanzo-rocm}. On that core sits a \texttt{candle}-derived Rust autograd that already trains two algorithms and contains the exact QLoRA primitive~\cite{hanzo-train}. The synthesis claim of this paper is that these are one system, not two: by Lemma~\ref{lem:qlora} and Proposition~\ref{prop:noseam} the served weights \emph{are} the trained base, read by the same kernel with no format break, and by Theorem~\ref{thm:unified} and Corollary~\ref{cor:budget} serving and on-device QLoRA training share one bandwidth-bound budget, so a machine that serves a model at the wall can personalize it at a comparable per-step cost. That is the moat---a unified train-and-serve runtime that orchestrate-black-box-engine platforms structurally cannot match~\cite{kubeflow,kserve,ray}---and the premise of structurally private, on-device continual learning~\cite{hanzo-clp,hanzo-platform}. We do not beat llama here, we do not claim a speed win over NVIDIA, and we do not claim zero leakage; the honest, defensible claims are the unified runtime, the structural delta-only egress contract, and a parity-enabled cost lever on cheaper silicon. Phase 1---the trainable LoRA layer and a real SFT loop---is the make-or-break that turns the thesis from formally true to operationally demonstrated.
\begin{thebibliography}{99}
\bibitem{hanzo-rocm}
Hanzo AI Research. \textit{Native ROCm Inference on a Consumer RDNA3.5 APU: Reaching llama.cpp Decode Parity via a Unified 1-bit-to-Full Quantization Compute Core}. Hanzo Industries, 2026.
\bibitem{hanzo-train}
Hanzo AI Research. \textit{Native Training in hanzo-ml/hanzo-engine: One Engine, One Quant Format, One Backend}. Hanzo Industries, 2026.
\bibitem{hanzo-clp}
Hanzo AI Research. \textit{Continuously-Learning Private AI: Self-Improvement and Privacy in the Hanzo Native Stack}. Hanzo Industries, 2026.
\bibitem{hanzo-platform}
Hanzo AI Research. \textit{Economical Serving and the hanzo.ai Platform: A Bandwidth-Optimal Unified Train-and-Serve Runtime}. Hanzo Industries, 2026.
\bibitem{ggml}
G. Gerganov et al. \textit{ggml: Tensor library for machine learning}. 2023.
\bibitem{llamacpp}
G. Gerganov et al. \textit{llama.cpp: LLM inference in C/C++}. 2023.
\bibitem{qwen3}
Qwen Team. \textit{Qwen3 Technical Report}. Alibaba, 2025.
\bibitem{williams2009}
S. Williams, A. Waterman, and D. Patterson. Roofline: An Insightful Visual Performance Model for Multicore Architectures. \textit{Communications of the ACM}, 52(4):65--76, 2009.
\bibitem{kwon2023}
W. Kwon et al. Efficient Memory Management for Large Language Model Serving with PagedAttention. \textit{SOSP}, 2023.
\bibitem{dao2022}
T. Dao et al. FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. \textit{NeurIPS}, 2022.
\bibitem{lora}
E. Hu et al. LoRA: Low-Rank Adaptation of Large Language Models. \textit{ICLR}, 2022.
\bibitem{qlora}
T. Dettmers et al. QLoRA: Efficient Finetuning of Quantized LLMs. \textit{NeurIPS}, 2023.
\bibitem{grpo}
Z. Shao et al. DeepSeekMath: Pushing the Limits of Mathematical Reasoning in Open Language Models. 2024. (Group Relative Policy Optimization.)
\bibitem{shumailov2024}
I. Shumailov et al. AI models collapse when trained on recursively generated data. \textit{Nature}, 631:755--759, 2024.
\bibitem{kubeflow}
The Kubeflow Authors. \textit{Kubeflow: The Machine Learning Toolkit for Kubernetes}. 2018.
\bibitem{kserve}
The KServe Authors. \textit{KServe: Standardized Serverless ML Inference on Kubernetes}. 2021.
\bibitem{ray}
P. Moritz et al. Ray: A Distributed Framework for Emerging AI Applications. \textit{OSDI}, 2018.
\bibitem{rocm2024}
AMD. \textit{ROCm: Open Software Platform for GPU Compute}. Version 7.1, 2026.
\end{thebibliography}
\end{document}
hanzo-native-training.pdf
BIN
View File
Binary file not shown.
hanzo-native-training.tex
+377
View File
@@ -0,0 +1,377 @@
\documentclass[11pt]{article}
\usepackage{shared/hanzocover}
\usepackage[utf8]{inputenc}
\usepackage{amsmath,amssymb,amsthm,booktabs,hyperref,graphicx}
\usepackage{xcolor}
\usepackage{listings}
\input{shared/lstlang}
\definecolor{codegreen}{rgb}{0,0.6,0}
\definecolor{codegray}{rgb}{0.5,0.5,0.5}
\definecolor{codepurple}{rgb}{0.58,0,0.82}
\definecolor{backcolour}{rgb}{0.95,0.95,0.92}
\lstdefinestyle{mystyle}{
backgroundcolor=\color{backcolour},
commentstyle=\color{codegreen},
keywordstyle=\color{codepurple},
numberstyle=\tiny\color{codegray},
stringstyle=\color{codegreen},
basicstyle=\ttfamily\footnotesize,
breakatwhitespace=false,
breaklines=true,
captionpos=b,
keepspaces=true,
numbers=left,
numbersep=5pt,
showspaces=false,
showstringspaces=false,
showtabs=false,
tabsize=2
}
\lstset{style=mystyle}
\newtheorem{theorem}{Theorem}
\newtheorem{lemma}{Lemma}
\newtheorem{proposition}{Proposition}
\newtheorem{definition}{Definition}
\title{Native Training in hanzo-ml/hanzo-engine: One Engine, One Quant Format, One Backend}
\author{Hanzo Dev \\ \texttt{[email protected]} \\ \textit{Hanzo AI}}
\date{June 2026}
\begin{document}
\hanzocoverpage
\begin{abstract}
We describe the architecture for training language models natively inside the hanzo Rust inference stack (\texttt{hanzo-ml} / \texttt{hanzo-engine}), eliminating the conversion boundary between a Python training stack and a separate serving engine. The existing engine already ships a reverse-mode autograd (\texttt{hanzo-ml}, a fork of HuggingFace \texttt{candle}; $809$ lines, full matmul/conv/gather rules), bias-corrected AdamW with decoupled weight decay, safetensors checkpoint I/O, and a unified quantization kernel that serves \texttt{Q8\_0} decode at $24.5$ tokens/s, equal to $0.98\times$ llama.cpp ($25.05$ tokens/s) on a gfx1151 / Radeon~8060S (Strix~Halo) device. Two trainers prove the substrate end-to-end: an algorithm-complete, tested GRPO reinforcement-learning loop ported from \texttt{zooai/gym}, and an in-engine AnyMoE (the MoE-distillation trainer inherited from \texttt{mistral.rs}) gate trainer that performs forward, cross-entropy, backward, optimizer step, and checkpoint save over frozen quantized expert weights. We show that the QLoRA primitive---a detached \texttt{dequantize()} of the quantized base plus a trainable \texttt{bf16} adapter pair---is already present in the code, so a model can be quantized once, served by the unified kernel, and trained by attaching adapters over the exact bytes that serve it, with zero dequant-to-train and zero re-quant-to-serve. We formalize the parameter and memory accounting of rank-$r$ QLoRA over a frozen quantized base, and prove that linear LoRA adapters over a shared base compose additively (the ``LoRA-soup'' property). We give an honest accounting of what is real, what needs building, and what is deliberately out of scope.
\end{abstract}
\section{Introduction}
Production deployment of a fine-tuned language model conventionally crosses an engine boundary. A model is trained in a thick Python stack and then \emph{re-exported and re-quantized} before a separate inference engine can serve it. The result is two engines, two quantization formats, two device backends, and a conversion step at every boundary.
The hanzo stack is positioned to remove that boundary. The Rust inference engine already owns a reverse-mode autograd (\texttt{ml/hanzo-ml/src/backprop.rs}, $809$ lines; \texttt{hanzo-ml} is a fork of HuggingFace \texttt{candle}), real optimizers (\texttt{ml/hanzo-nn/src/optim.rs}), and a unified quantization kernel core that serves \texttt{Q8\_0} decode at $24.5$ tokens/s, equal to $0.98\times$ of llama.cpp on a gfx1151 / Radeon~8060S device~\cite{hanzo-rocm}. If training is built \emph{on the same \texttt{Tensor}/\texttt{Var}/quant substrate} that serves inference, then a model is trained and served by one engine, in one quant format, on one device backend, with \textbf{zero conversion}. The strategic payoff concentrates in QLoRA: \texttt{bf16} adapters can be trained directly over the \emph{exact} quantized weights that the serving kernel reads.
This is partly real today and partly a build plan. The contribution of this paper is to state the split honestly, to identify the single load-bearing finding---that the autograd primitive QLoRA needs is already present in the quantization layer---and to formalize the parameter, memory, and composition properties that make the native-training architecture sound.
\paragraph{Provenance of claims.} Every assertion marked \textsc{Real} cites a file in the hanzo tree. Every performance number is measured on real hardware and cited to the ROCm performance record~\cite{hanzo-rocm}, or else explicitly labelled an estimate. Techniques inherited from Python tooling are marked \textsc{Ported}; the genuinely novel pieces are marked \textsc{Novel}. Build status is marked \textsc{Real}, \textsc{Needs-Building}, or \textsc{Aspirational}.
\paragraph{Contributions.}
\begin{itemize}
\item An architecture that unifies training and serving on a single Rust substrate: \texttt{loss.backward} $\to$ \texttt{GradStore} $\to$ \texttt{AdamW.step} $\to$ \texttt{Var.set} $\to$ safetensors, reusing the exact quantization kernel that serves inference.
\item Identification of the QLoRA primitive already present in the code: \texttt{QTensor::dequantize()} emits a detached constant, so a frozen quantized base carries no gradient and a separate adapter pair carries the entire learnable signal.
\item A formal parameter and training-memory accounting for rank-$r$ QLoRA~\cite{hu2021lora,dettmers2023qlora} over a frozen quantized base (Proposition~\ref{prop:qlora}), with proof.
\item A composition lemma showing that linear LoRA adapters over a shared base compose additively, justifying ``LoRA-soup''~\cite{wortsman2022soups} merging (Lemma~\ref{lem:soup}), with proof.
\item An ordered migration plan from the Python \texttt{gym} stack to the native Rust trainer, with risk and effort estimates, and an explicit list of stubs to replace and features deliberately left out of scope.
\end{itemize}
\section{The Foundation That Already Works}
\label{sec:foundation}
The following components are not stubs; each is cited and several are covered by passing tests. Collectively they form the training \emph{substrate}.
\paragraph{Autograd (\textsc{Real}).} \texttt{ml/hanzo-ml/src/backprop.rs}. \texttt{Tensor::backward()} returns a \texttt{GradStore} via topological \texttt{sorted\_nodes()}, with full reverse rules including matmul (lhs gradient $= \mathrm{grad} \cdot \mathrm{rhs}^{\top}$, rhs gradient $= \mathrm{lhs}^{\top} \cdot \mathrm{grad}$, lines $457$--$467$), convolution, gather/scatter, \texttt{index\_add}, concatenation, broadcast, and \texttt{CustomOp1/2/3} backward dispatch (lines $662$--$687$). \texttt{GradStore} is a \texttt{HashMap<TensorId, Tensor>} (line $733$).
\paragraph{Var (\textsc{Real}).} \texttt{ml/hanzo-ml/src/variable.rs}. \texttt{Var(Tensor)} is the leaf wrapper; \texttt{is\_variable()} marks gradient leaves; \texttt{set()} performs an autograd-safe in-place update, which is what an optimizer step writes through.
\paragraph{Optimizers (\textsc{Real}).} \texttt{ml/hanzo-nn/src/optim.rs}. An \texttt{Optimizer} trait with \texttt{step(\&GradStore)} and \texttt{backward\_step(\&loss)} (lines $5$--$29$); plain \texttt{SGD} and a full bias-corrected \texttt{AdamW} with decoupled weight decay (lines $152$--$182$). Crucially, \texttt{new()} \emph{filters to float dtypes} (\texttt{var.dtype().is\_float()}, lines $44$--$47$ and $122$--$124$), so quantized or integer leaves are never updated. That filter is exactly the invariant a QLoRA trainer requires.
\paragraph{VarMap and checkpoint I/O (\textsc{Real}).} \texttt{ml/hanzo-nn/src/var\_map.rs}. \texttt{save()}/\texttt{load()} via safetensors (lines $31$--$45$), \texttt{all\_vars()}, and lazy \texttt{get(shape, path, init, dtype, dev)}. This is real persistence for trainable parameters today.
\paragraph{Losses (\textsc{Real}).} \texttt{ml/hanzo-nn/src/loss.rs}. \texttt{nll}, \texttt{cross\_entropy} ($\mathrm{log\_softmax} + \mathrm{nll}$), \texttt{mse}, \texttt{binary\_cross\_entropy\_with\_logit}, and \texttt{huber}, all differentiable.
\paragraph{NN layers (\textsc{Real}).} \texttt{hanzo-nn} provides linear, embedding, the norm family, convolution, RNN, MoE, rotary embedding, activations, and \texttt{ops.rs} ($\mathrm{log\_softmax}$/$\mathrm{softmax}$/$\mathrm{topk}$). All are forward-only structs that ride autograd; no hand-written backward is needed.
\paragraph{Verdict.} The training substrate---$\texttt{loss.backward} \to \texttt{GradStore} \to \texttt{AdamW.step} \to \texttt{Var.set} \to$ safetensors save---is \textsc{Real} and end-to-end. What is missing is the trainer scaffolding on top of it (Sections~\ref{sec:catalogue} and~\ref{sec:migration}).
\section{Two Real Trainers}
\label{sec:trainers}
\subsection{GRPO reinforcement learning (\textsc{Ported}, tested)}
\texttt{ml/hanzo-training/src/grpo/} contains \texttt{trainer}, \texttt{math}, \texttt{policy}, \texttt{sampler}, \texttt{config}, \texttt{reward}, and \texttt{metrics}. The module is ported from the \texttt{zooai/gym} GRPO~\cite{shao2024grpo} trainer. \texttt{GrpoTrainer::step} (\texttt{trainer.rs:129}) runs the full five-phase update: rollout (\texttt{Policy::sample\_group}); reward (weighted multi-reward sum, \texttt{score()} at \texttt{trainer.rs:119}); group-relative advantages (\texttt{math::batch\_advantages}, unbiased standard deviation with $\mathrm{ddof}=1$, \texttt{math.rs:45--52}); a differentiable, sign-correct, length-normalized policy-gradient loss (\texttt{math::policy\_gradient\_loss}, \texttt{math.rs:87--120}); an optional low-variance \texttt{k3} KL penalty to a frozen reference ($k3 = \exp(\Delta) - \Delta - 1$, \texttt{math.rs:134--166}); and \texttt{AdamW.backward\_step(\&loss)} (\texttt{trainer.rs:228}). It is tested: a smoke test learns a toy reward to ceiling (asserts mean reward rises past $0.6 \cdot \texttt{max\_len}$), a KL-branch test stays finite, and math unit tests verify advantage zero-mean, unit-scale scaling, and policy-gradient sign.
\paragraph{The gap, stated honestly.} The \texttt{Policy} trait (\texttt{policy.rs:33}) requires exactly three methods: \texttt{trainable\_vars()}, \texttt{sample\_group()}, and \texttt{sequence\_logprob()}. The only implementor is \texttt{ToyPolicy}, a \texttt{[max\_len, vocab]} logits bandit in tests and \texttt{examples/grpo\_toy.rs}. No hanzo-engine transformer implements \texttt{Policy} yet; wiring a real model to those three methods is \textsc{Needs-Building} (Section~\ref{sec:migration}).
\subsection{AnyMoE gate trainer (\textsc{Real}, end-to-end in-engine)}
(\texttt{AnyMoE} is a \texttt{mistral.rs} feature inherited by the engine.) \texttt{engine/hanzo-engine/src/pipeline/amoe.rs} and \texttt{engine/hanzo-engine/src/amoe/mod.rs} contain the strongest existing ``train on top of an inference model'' precedent. It builds trainable MoE gating \texttt{Linear}s as \texttt{Var}s, then runs a real per-layer loop (\texttt{amoe.rs:348--516}): build a per-layer \texttt{AdamW} (\texttt{amoe.rs:351}), forward the frozen base model (\texttt{target.forward\_inputs}, \texttt{amoe.rs:483}), collect each gate's logits, take \texttt{hanzo\_nn::loss::cross\_entropy} (\texttt{amoe.rs:507}), \texttt{loss.backward()}, \texttt{optimizer.step(\&gradstore)} (\texttt{amoe.rs:511--512}), log CSV loss, and finally detach the gates and save \texttt{gate.safetensors}. It trains \emph{new} gating parameters on top of frozen quantized expert weights---the exact pattern QLoRA needs: trainable \texttt{Var}s layered over a frozen base, inside the inference engine.
\section{The Quant Stack and the QLoRA Primitive}
\label{sec:quant}
\subsection{What exists for inference (\textsc{Real})}
\begin{itemize}
\item \textbf{candle quant} (\texttt{ml/hanzo-ml/src/quantized/}): \texttt{QTensor}/\texttt{QMatMul}, GGUF k-quants (\texttt{Q2K}--\texttt{Q6K}, \texttt{Q8\_0}, \texttt{MXFP4}), plus $11$ low-bit IQ/ternary/FP4 types added and validated bit-exact against ggml~\cite{hanzo-rocm}.
\item \textbf{engine quant} (\texttt{engine/hanzo-quant}): a \texttt{QuantMethod} trait (\texttt{lib.rs:990}) with \texttt{forward}, \texttt{dequantize\_w} (\texttt{lib.rs:995}), \texttt{add\_delta\_w} (\texttt{lib.rs:1056}), and \texttt{apply\_isq}; methods for gguf/gptq/hqq/afq/bitsandbytes/fp8/mxfp4/unquantized; in-situ quantization (ISQ), imatrix, and UQFF serialization.
\item \textbf{the unified ROCm core} (\textsc{Novel}, Section~\ref{sec:perf}): one templated \texttt{qmatvec\_core<WTYPE,XT>} decode and \texttt{qmmq\_core<WTYPE>} prefill, one decode function per format, table-driven dispatch. Adding a quant costs one \texttt{qdec}, one traits row, and one macro line; \emph{zero new kernels}.
\end{itemize}
\subsection{The load-bearing finding}
\texttt{QTensor::dequantize()} (\texttt{quantized/mod.rs:775--778}) constructs its output with \texttt{BackpropOp::none()}:
\begin{lstlisting}[language=Rust]
pub fn dequantize(&self, device: &Device) -> Result<Tensor> {
let storage = self.storage
.dequantize(self.shape.elem_count())?;
let none = crate::op::BackpropOp::none();
crate::tensor::from_storage(
storage, self.shape.clone(), none, false)
.to_device(device)
}
\end{lstlisting}
\texttt{dequantize\_f16} (\texttt{mod.rs:781}) and the \texttt{QMatMul}/\texttt{indexed\_moe\_forward} forwards likewise emit \texttt{BackpropOp::none()}. A dequantized weight is therefore a \emph{detached constant}: no gradient flows into quantized storage. This is not a limitation; it is exactly the correct QLoRA primitive. The frozen quant base is a constant, and a separate \texttt{bf16}/\texttt{f16} LoRA \texttt{Var} pair $(A, B)$, whose matmul backpropagates normally (rules at \texttt{backprop.rs:457--467}), carries the entire learnable signal. The AdamW float-dtype filter (Section~\ref{sec:foundation}) guarantees the quant base is never touched even if it were handed in.
Consequently the \emph{math substrate} for QLoRA exists today. The LoRA layer and the assembly that injects it do not (Section~\ref{sec:catalogue}).
\subsection{What is not a trainer today}
The existing LoRA (\texttt{engine/hanzo-engine/src/lora/loralinear.rs}) loads \texttt{a\_adapters}/\texttt{b\_adapters} as frozen \texttt{Linear}s from a \texttt{ShardedVarBuilder} (\texttt{loralinear.rs:17--18, 34--48}) and supports \texttt{merge\_weights()} via \texttt{QuantMethod::add\_delta\_w} (\texttt{loralinear.rs:145--156}). It is inference-and-merge only: the adapters are constants, not \texttt{Var}s, and there is no backward path into them. Likewise \texttt{hanzo-quant/src/lora/static\_lora.rs} is static. There is no trainable-LoRA path anywhere in the tree yet.
\section{Formal Properties of QLoRA over a Frozen Quantized Base}
\label{sec:formal}
We now formalize the parameter and memory accounting that motivates QLoRA, and the composition property that justifies adapter merging. Throughout, a linear layer maps $x \in \mathbb{R}^{d_{\text{in}}}$ to $y \in \mathbb{R}^{d_{\text{out}}}$ via a weight matrix $W \in \mathbb{R}^{d_{\text{out}} \times d_{\text{in}}}$.
\begin{definition}[LoRA adapter]
\label{def:lora}
A rank-$r$ LoRA adapter for a linear layer with frozen base weight $W_0 \in \mathbb{R}^{d_{\text{out}} \times d_{\text{in}}}$ is a pair $(A, B)$ with $A \in \mathbb{R}^{r \times d_{\text{in}}}$ and $B \in \mathbb{R}^{d_{\text{out}} \times r}$, together with a fixed scalar $s = \alpha/r$. The adapted layer computes
\[
y = W_0 x + s\, B (A x) = \bigl(W_0 + s\, B A\bigr) x .
\]
Only $A$ and $B$ are trainable; $W_0$ is held constant.
\end{definition}
\begin{definition}[Quantized base]
\label{def:qbase}
A \emph{quantized base} is a weight tensor $W_0$ stored in a low-bit format $Q$ (e.g.\ \texttt{Q8\_0}, \texttt{Q4\_K}) and exposed to autograd only through a detaching dequantization $\mathrm{deq}_Q(W_0)$, i.e.\ the operation that produces $W_0$ in compute precision carries \texttt{BackpropOp::none()} and contributes no entry to the \texttt{GradStore}.
\end{definition}
\begin{proposition}[QLoRA trainable-parameter and memory accounting]
\label{prop:qlora}
Consider a model whose linear layers carry frozen quantized base weights as in Definition~\ref{def:qbase}, with rank-$r$ LoRA adapters as in Definition~\ref{def:lora} attached to a set $\mathcal{L}$ of layers, where layer $\ell \in \mathcal{L}$ has dimensions $(d^{\ell}_{\text{out}}, d^{\ell}_{\text{in}})$. Then:
\begin{enumerate}
\item The number of trainable parameters contributed by layer $\ell$ is exactly $r\,(d^{\ell}_{\text{out}} + d^{\ell}_{\text{in}})$; in the square case $d^{\ell}_{\text{in}} = d^{\ell}_{\text{out}} = d$ this is $2rd$ per matrix.
\item No gradient is ever written to any quantized base weight; the set of optimizer-updated leaves is precisely $\bigcup_{\ell \in \mathcal{L}} \{A_\ell, B_\ell\}$.
\item The optimizer state and weight-gradient memory scale as $\Theta\!\bigl(\sum_{\ell \in \mathcal{L}} r(d^{\ell}_{\text{out}} + d^{\ell}_{\text{in}})\bigr)$, independent of the (much larger) base parameter count $\sum_\ell d^{\ell}_{\text{out}} d^{\ell}_{\text{in}}$. Total training memory is therefore the quantized base (read-only) plus adapter weights, adapter gradients, adapter optimizer moments, and the forward activations retained for backward---\emph{not} a full-precision copy of the base weights, their gradients, or their optimizer state.
\end{enumerate}
\end{proposition}
\begin{proof}
\emph{(1)} By Definition~\ref{def:lora}, layer $\ell$ adds the trainable matrices $A_\ell \in \mathbb{R}^{r \times d^{\ell}_{\text{in}}}$ and $B_\ell \in \mathbb{R}^{d^{\ell}_{\text{out}} \times r}$ and no other trainable tensor ($s$ is a fixed scalar, not a parameter). Counting entries gives $r\,d^{\ell}_{\text{in}} + d^{\ell}_{\text{out}}\,r = r(d^{\ell}_{\text{out}} + d^{\ell}_{\text{in}})$, which is $2rd$ when $d^{\ell}_{\text{in}} = d^{\ell}_{\text{out}} = d$.
\emph{(2)} The adapted output is $y = \mathrm{deq}_Q(W_0) x + s\,B(Ax)$. The first term is detached by Definition~\ref{def:qbase}: $\mathrm{deq}_Q(W_0)$ carries \texttt{BackpropOp::none()} and so registers no producer in the backward graph; hence $\partial \mathcal{L} / \partial W_0$ is never accumulated into the \texttt{GradStore}, regardless of downstream loss $\mathcal{L}$. The second term is a composition of standard matmuls, whose reverse rules (\texttt{backprop.rs:457--467}) populate gradients for $A_\ell$ and $B_\ell$. Thus the only leaves receiving gradients are the $\{A_\ell, B_\ell\}$. Independently, the optimizer's construction filter retains only float-dtype leaves (Section~\ref{sec:foundation}); since the quantized base is stored in a non-float quant dtype, it is excluded from the update set even if mistakenly registered. The two mechanisms are independent and each is sufficient, so the optimizer-updated set is exactly $\bigcup_{\ell} \{A_\ell, B_\ell\}$.
\emph{(3)} A first-order adaptive optimizer such as AdamW stores, per trainable scalar, one gradient and a constant number of moment estimates (first and second moment), i.e.\ $\Theta(1)$ memory per trainable scalar. By (2) the trainable scalars number $\sum_{\ell} r(d^{\ell}_{\text{out}} + d^{\ell}_{\text{in}})$, giving optimizer-plus-gradient memory of that same order. The base weights are read-only and contribute no gradient or moment; they reside in their compressed quantized form (Definition~\ref{def:qbase}). Hence the total is the read-only quantized base plus adapter weights, adapter gradients, adapter moments, and retained activations, with the adapter-related terms independent of $\sum_\ell d^{\ell}_{\text{out}} d^{\ell}_{\text{in}}$.
\end{proof}
\begin{lemma}[Additive composition of linear LoRA adapters; the LoRA-soup property]
\label{lem:soup}
Fix a frozen base weight $W_0$ for a linear layer and a common scale $s$. Let $\{(A_i, B_i)\}_{i=1}^{n}$ be $n$ rank-$r_i$ LoRA adapters trained over the \emph{same} $W_0$, with merged increments $\Delta_i = s\, B_i A_i$. For any convex combination weights $\lambda_1, \dots, \lambda_n \ge 0$ with $\sum_i \lambda_i = 1$, define the soup weight
\[
W_{\text{soup}} = W_0 + \sum_{i=1}^{n} \lambda_i \Delta_i .
\]
Then on every input $x$,
\[
W_{\text{soup}}\, x = \sum_{i=1}^{n} \lambda_i \bigl(W_0 + \Delta_i\bigr) x ,
\]
i.e.\ applying the merged adapter equals the convex combination of the individually adapted layer outputs. Moreover $W_{\text{soup}} - W_0 = \sum_i \lambda_i \Delta_i$ admits an exact rank-$\bigl(\sum_i r_i\bigr)$ factorization, so the soup is itself representable as a single LoRA adapter.
\end{lemma}
\begin{proof}
Each adapted layer is affine in $x$ with the \emph{shared} linear part anchored at $W_0$: by Definition~\ref{def:lora}, $(W_0 + \Delta_i)x = W_0 x + \Delta_i x$. Take the convex combination and use $\sum_i \lambda_i = 1$:
\[
\sum_{i=1}^{n} \lambda_i (W_0 + \Delta_i) x
= \Bigl(\sum_i \lambda_i\Bigr) W_0 x + \sum_i \lambda_i \Delta_i x
= W_0 x + \Bigl(\sum_i \lambda_i \Delta_i\Bigr) x
= W_{\text{soup}}\, x .
\]
The first equality is linearity of matrix-vector multiplication in the weight; the identity holds for all $x$, establishing the stated equality of maps. For the factorization, define the stacked matrices
\[
\widetilde{B} = \bigl[\, s\lambda_1 B_1 \;\big|\; s\lambda_2 B_2 \;\big|\; \cdots \;\big|\; s\lambda_n B_n \,\bigr] \in \mathbb{R}^{d_{\text{out}} \times \sum_i r_i},
\qquad
\widetilde{A} = \begin{bmatrix} A_1 \\ A_2 \\ \vdots \\ A_n \end{bmatrix} \in \mathbb{R}^{\sum_i r_i \times d_{\text{in}}} .
\]
Then $\widetilde{B}\widetilde{A} = \sum_i s\lambda_i B_i A_i = \sum_i \lambda_i \Delta_i = W_{\text{soup}} - W_0$, exhibiting the increment as a single LoRA factorization of rank at most $\sum_i r_i$ (with scale folded into $\widetilde{B}$).
\end{proof}
\begin{proof}[Remark on necessity of the shared base]
The argument uses $\sum_i \lambda_i = 1$ precisely to collapse $\sum_i \lambda_i W_0$ back to $W_0$. If the adapters were trained over \emph{different} bases $W_0^{(i)}$, the cross terms $\sum_i \lambda_i (W_0^{(i)} - W_0)$ would not vanish and additive composition would fail. This is the formal reason hanzo's soup and delta-soup routines (\texttt{model\_delta/soup.rs}, Section~\ref{sec:catalogue}) operate over a common base checkpoint.
\end{proof}
These two results capture why owning both the trainer and the quant kernel is structurally advantageous: Proposition~\ref{prop:qlora} shows the trainable footprint is the adapter, not the base, even though the base is served at full quantized speed; Lemma~\ref{lem:soup} shows multiple such adapters merge exactly over the shared frozen base, enabling post-hoc model-soup composition without retraining.
\section{The Native Trainer Catalogue}
\label{sec:catalogue}
Table~\ref{tab:catalogue} maps each trainer onto the autograd engine and records its build status.
\begin{table}[t]
\centering
\small
\caption{Native trainer catalogue and autograd mapping.}
\label{tab:catalogue}
\begin{tabular}{@{}lll@{}}
\toprule
\textbf{Trainer} & \textbf{Status} & \textbf{Mechanism} \\
\midrule
Full SFT & Needs-building & Weights as \texttt{Var}s; teacher-forced \\
& (substrate real) & LM \texttt{cross\_entropy}; \texttt{AdamW}; save. \\
LoRA & Needs-building & Frozen base $+$ \texttt{Var}s $A,B$; \\
& & $y = W_0 x + s\,B(Ax)$. \\
QLoRA & Needs-building & As LoRA; base is \texttt{QTensor} via \\
& (primitive real) & \texttt{qmatvec\_core}; already detached. \\
DoRA & Needs-building & LoRA $+$ trainable column- \\
& & magnitude \texttt{Var}; norm op. \\
BitDelta & Real (compress) & Sign-pack $(W_{\!f}-W_0)$, scale \\
& trainer pending & $\alpha = \mathrm{mean}|\Delta|$; reconstruct. \\
Soup / delta-soup & Real & Elementwise mean of checkpoints; \\
& & $W_0 + \mathrm{mean}_i(W_{\!f,i}-W_0)$. \\
\bottomrule
\end{tabular}
\end{table}
\paragraph{Full fine-tune (SFT).} The substrate is \textsc{Real}; only the scaffolding is missing. The path: build the model with a \texttt{VarBuilder::from\_varmap} so weights are \texttt{Var}s, run a teacher-forced forward, take per-token \texttt{cross\_entropy}, \texttt{AdamW.backward\_step}, and checkpoint with \texttt{VarMap.save}. The blocker is that the roughly $80$ engine model definitions build via \texttt{VarBuilder} into \emph{frozen} \texttt{Tensor}s; SFT needs a \texttt{Var}-backed construction path. The generic \texttt{Trainer} advertised for this is a stub (Section~\ref{sec:stubs}).
\paragraph{LoRA.} A new \texttt{TrainableLoraLinear} holding a frozen base alongside trainable \texttt{Var}s $A,B$ and scale $s$. Forward: \texttt{base.forward(x) + lora\_b(lora\_a(x)) * scale}, with $A,B$ created through a \texttt{VarMap}. Backward flows only into $A,B$ because the base is detached. It reuses the existing adapter math from \texttt{loralinear.rs}, swapping the frozen \texttt{Linear} for \texttt{Var}s.
\paragraph{QLoRA---the synergy that justifies the effort.} Identical to LoRA, but the base is a \texttt{QTensor} served by the unified \texttt{qmatvec\_core} / \texttt{QMatMul} path. Because that path already emits \texttt{BackpropOp::none()} (Section~\ref{sec:quant}), no extra detaching is needed. The model is quantized once, served at the performance of Section~\ref{sec:perf}, and trained by attaching \texttt{bf16} adapter \texttt{Var}s over the very same bytes: zero dequant-to-train, zero re-quant-to-serve. This is the one place where owning both the trainer and the quant kernel yields something the Python stack structurally cannot: bitsandbytes/peft must dequantize a block to \texttt{f16}, run the \texttt{f16} matmul, and serve through a \emph{different} code path; here it is literally the same kernel. Proposition~\ref{prop:qlora} quantifies the resulting memory advantage.
\paragraph{DoRA.} Weight-decomposed LoRA~\cite{liu2024dora}: $W' = m \cdot (W_0 + \Delta W)/\lVert W_0 + \Delta W \rVert_{\text{col}}$, where $m$ is a trainable per-column magnitude vector and $\Delta W = BA$. On the engine: add one magnitude \texttt{Var} of shape $[d_{\text{out}}]$, compute the column norm with existing ops, and scale. All differentiable; no new backward. The only cost is recomputing the directional norm each step, a known DoRA expense, not a hanzo-specific one.
\paragraph{BitDelta and soups (\textsc{Real}).} \texttt{model\_delta/bitdelta.rs} implements the BitDelta~\cite{liu2024bitdelta} $\mathrm{encode}$/$\mathrm{decode\_apply}$: encode is sign-pack$(W_{\!f}-W_0)$ plus a per-tensor scale $\alpha = \mathrm{mean}|\Delta|$; reconstruct is $W_0 + \alpha \cdot \mathrm{sign}(\Delta)$. \texttt{model\_delta/soup.rs} implements $\mathrm{soup}$ (elementwise mean of checkpoints) and $\mathrm{delta\_soup}$ ($W_0 + \mathrm{mean}_i(W_{\!f,i}-W_0)$), both tested. They operate on saved checkpoints, so they compose with any trainer above: train, save \texttt{Var} weights, then \texttt{bitdelta::encode} for 1-bit delta storage or \texttt{delta\_soup} to average several fine-tunes. Lemma~\ref{lem:soup} is the formal justification for the LoRA case of delta-soup. For the patent record: neither BitDelta nor any model-soup/TIES/DARE routine exists anywhere in \texttt{zooai/gym}; these are hanzo-native with no Axolotl analog. The DeltaSoup \emph{federation} variant (trim-mean of LoRA deltas) lives only in the Python \texttt{zen/gym} and is not yet ported to Rust.
\section{Quant-Aware Training and Inherited Performance}
\label{sec:perf}
The reason to train on hanzo's quant core is that the core is fast and bit-exact, and an adapter trained over it serves with no conversion. Table~\ref{tab:perf} summarizes the measured state on gfx1151 / Radeon~8060S (Strix~Halo), all numbers from the ROCm performance record~\cite{hanzo-rocm}. The real-hardware roofline is $217$ GB/s read bandwidth and $52$ TOPS-i8.
\begin{table}[t]
\centering
\small
\caption{Measured performance on gfx1151 / 8060S (Strix Halo). Roofline: $217$~GB/s read, $52$~TOPS-i8. ``$\times$ llama'' is the ratio to llama.cpp on the same hardware. All quant numbers verified \texttt{nbad}$=0$ against the CPU \texttt{to\_float} oracle.}
\label{tab:perf}
\begin{tabular}{@{}lrrr@{}}
\toprule
\textbf{Path} & \textbf{Before} & \textbf{After} & \textbf{$\times$ llama} \\
\midrule
Decode matvec (GB/s) & 103.5 & 245.8 & 0.97 \\
\quad (\% of BW wall) & 45\% & 97\% & --- \\
Decode end-to-end (t/s) & 8.5 & 24.5 & 0.98 \\
MoE decode (t/s) & 0.93 & 13.6 & --- \\
Prefill end-to-end (t/s) & 39 & 1015 & 0.86 \\
\quad WMMA GEMM isolated & --- & 1414 & $>1$ \\
\bottomrule
\end{tabular}
\end{table}
\paragraph{Decode (\textsc{Real}).} The lane-strided unified core rewrote \texttt{qmatvec\_core} from serial to lane-strided: the isolated matvec went from $103.5$ GB/s ($45\%$ of the bandwidth wall) to $245.8$ GB/s ($97\%$), a $2.37\times$ gain; end-to-end eager decode went from $8.5$ to $24.5$ tokens/s, equal to $0.98\times$ llama.cpp ($25.05$ tokens/s). This is effective parity and the measured honest ceiling: matvec sits at the bandwidth wall, and batch-1 attention is not a winnable lever.
\paragraph{The unified abstraction (\textsc{Novel}).} One templated \texttt{qmatvec\_core<WTYPE,XT>} decode and \texttt{qmmq\_core<WTYPE>} prefill, with \texttt{quant\_format!} / \texttt{for\_each\_quant!} table-driven dispatch (\texttt{quantized/quant\_format.rs}, \texttt{proof.rs}). Adding a quant costs one \texttt{qdec}, one traits row, and one macro line; zero new kernels. It is proven across \texttt{Q8\_0}/\texttt{Q4\_0}/\texttt{Q4\_K}/\texttt{Q6\_K}/\texttt{IQ4\_XS}/\texttt{TQ2\_0}, all with \texttt{nbad}$=0$ against the CPU \texttt{to\_float} oracle. A QLoRA base in any of these formats trains and serves through the same core.
\paragraph{MoE.} The same core powers \texttt{indexed\_moe\_forward} per routed expert; MoE decode went from $0.93$ to $13.6$ tokens/s (about $2\times$, via a resident-bank cache plus the lane-strided core), with \texttt{Q4\_K}$+$\texttt{Q8\_0} MoE at \texttt{nbad}$=0$. This is the path an AnyMoE-style or MoE-QLoRA trainer would ride.
\paragraph{Prefill.} The WMMA int8 GEMM kernel reaches about $1414$ tokens/s in isolation (above llama); end-to-end prefill went from $39$ to $1015$ tokens/s, equal to $0.86\times$ llama.cpp ($1176$ tokens/s). The remaining gap is engine attention, not the matvec; a multi-wave WMMA flash-attention kernel is the open lever (\textsc{Aspirational}, a large rewrite).
\paragraph{Honest framing.} Decode is parity, prefill is $0.86\times$ and improving, and MoE decode is roughly $2\times$. The training value is not ``we are fastest''; it is that the trained artifact is \emph{already} in the served format on the served kernel.
\section{Migration Plan: gym (Python) $\to$ native (Rust)}
\label{sec:migration}
\paragraph{Reusable as-is (\textsc{Real}, no rewrite).} Autograd; SGD/AdamW; VarMap save/load; the loss module; the roughly $80$ model architecture definitions (as forward graphs); the full quant inference stack with ISQ and UQFF; the GRPO algorithm core; AnyMoE gate training; BitDelta; soups.
\paragraph{Build (\textsc{Needs-Building}), ordered by value/effort.}
\begin{enumerate}
\item \textbf{Trainable LoRA/QLoRA adapter layer} (highest value, low-medium effort). A \texttt{TrainableLoraLinear} with $(A,B)$ \texttt{Var}s over a frozen base; QLoRA variant over \texttt{QTensor}. Risk: \textsc{Low}---the AD primitive is proven (Section~\ref{sec:quant}), AnyMoE proves the train-over-frozen pattern (Section~\ref{sec:trainers}), and the adapter math already exists. Estimate: about one week to a tested LoRA $+$ QLoRA on one architecture (Qwen3), reusing \texttt{qmatvec\_core} for the base.
\item \textbf{Real transformer \texttt{Policy} for GRPO} (high value, medium effort). Implement the three trait methods (\texttt{policy.rs:33}) against a hanzo-engine model: \texttt{sample\_group} is the existing generate loop under no-grad; \texttt{sequence\_logprob} is a teacher-forced forward returning $\sum \log p$. Risk: \textsc{Medium}---needs the \texttt{Var}-backed model path (shared with SFT) and careful no-grad/grad separation. Estimate: one to two weeks; pairs naturally with item 1.
\item \textbf{Generic SFT trainer $+$ real tokenized dataset} (high value, medium effort). Replace the stub \texttt{Trainer}/\texttt{ModelWrapper}/\texttt{OptimizerWrapper} (Section~\ref{sec:stubs}) with a real epoch/step loop calling \texttt{hanzo\_nn::AdamW} directly, a teacher-forced LM \texttt{cross\_entropy}, and a real \texttt{tokenizers}-backed dataset. Risk: \textsc{Medium}, mostly plumbing. Estimate: one to two weeks.
\item \textbf{Preference losses (DPO~\cite{rafailov2023dpo}/KTO/ORPO/SimPO)} (medium value, low effort once item 2 exists). Closed-form differentiable losses over (chosen, rejected) log-probs, using the same \texttt{sequence\_logprob}. Risk: \textsc{Low}. Estimate: a few days each.
\item \textbf{Training ergonomics} (medium value, low-medium effort). Real LR schedulers (cosine/linear are toy approximations in the stub, \texttt{optimizer.rs:55--77}), gradient clipping, gradient accumulation, mixed precision, optimizer-state checkpoint/resume (only \texttt{Var} weights save today, not moments), and explicit gradient checkpointing (only the global \texttt{GRAD\_DO\_NOT\_DETACH} toggle exists). Risk: \textsc{Low-Medium}, individually small.
\item \textbf{DoRA / BitDelta-trainer / federated DeltaSoup control plane} (\textsc{Aspirational}-roadmap). DoRA is a small extension of item 1. The federated trim-mean DeltaSoup from \texttt{zen/gym} is Python-only today; porting its control plane to Rust over the \texttt{bf16}-safetensors VarMap seam is a separate project.
\end{enumerate}
\section{Stubs to Replace}
\label{sec:stubs}
These are advertised in \texttt{hanzo-training} but do not function; a reader must not mistake them for the working pieces above.
\begin{itemize}
\item \texttt{ml/hanzo-training/src/optimizer.rs}: \texttt{OptimizerWrapper::step()}/\texttt{zero\_grad()} are no-ops---the comment reads ``Simplified optimizer step - in practice, this would update parameters'' (\texttt{optimizer.rs:82}); it only advances a step counter and an LR number, touching no params or grads. (GRPO and AnyMoE ignore this and use \texttt{hanzo\_nn::AdamW} directly, which is why they work.)
\item \texttt{ml/hanzo-training/src/model.rs}: \texttt{ModelWrapper.forward()} returns a constant \texttt{Tensor::new(\&[1.0])}, \texttt{backward()} is a no-op, \texttt{parameters()} returns \texttt{vec![]}, and \texttt{save()} writes only JSON metadata.
\item \texttt{ml/hanzo-training/src/trainer.rs}: the generic SFT \texttt{Trainer} is wired to those two stubs, so it ``trains'' nothing; \texttt{train\_qwen.rs}/\texttt{train\_llama.rs}/\texttt{bin/train.rs} all call it.
\item \texttt{ml/hanzo-training/src/dataset.rs}: ``simplified tokenization,'' not a real tokenizer.
\end{itemize}
The migration (items 1--3) replaces all four with the proven \texttt{hanzo\_nn::AdamW} $+$ \texttt{loss.backward} $+$ \texttt{VarMap} substrate.
\section{Limitations and Future Work}
\label{sec:limitations}
We state the boundaries of the present architecture plainly.
\paragraph{The trainers above the substrate are not yet built.} Sections~\ref{sec:trainers} and~\ref{sec:quant} establish that the substrate (autograd, AdamW, VarMap, losses, detached quant base) is real and that two trainers (GRPO, AnyMoE) exercise it end-to-end. But the headline deliverable---trainable LoRA/QLoRA over an arbitrary engine model---is \textsc{Needs-Building}. The proof obligations of Section~\ref{sec:formal} hold for the design; they do not assert that the \texttt{TrainableLoraLinear} type exists today, because it does not.
\paragraph{Var-backed model construction is the shared blocker.} Both full SFT and a real GRPO \texttt{Policy} require building the roughly $80$ engine models with \texttt{Var} parameters rather than frozen \texttt{Tensor}s. This is a construction-path change, not a kernel change, but it is a prerequisite for two of the three highest-value items.
\paragraph{Performance is parity, not dominance.} Decode is $0.98\times$ llama.cpp and prefill $0.86\times$ (Section~\ref{sec:perf}); the matvec is at the bandwidth wall and is not a further lever, while the prefill gap is in engine attention and awaits a WMMA flash-attention kernel (\textsc{Aspirational}). The training argument is the absence of a conversion boundary, not raw speed.
\paragraph{Single-device scope.} We deliberately do not replicate FSDP/DeepSpeed-scale multi-node sharding, tensor/sequence/context parallelism, fused Liger/Unsloth kernels, or Mamba/diffusion trainers. These are \texttt{gym} breadth, not the personalization core. The first deliverable is single-device LoRA/QLoRA $+$ SFT $+$ GRPO, which covers the dominant fine-tune and RL use cases.
\paragraph{Optimizer-state durability.} Only \texttt{Var} weights are checkpointed today, not optimizer moments; resumable training requires the ergonomics work in item 5 of Section~\ref{sec:migration}.
\paragraph{Future work.} In priority order: ship \texttt{TrainableLoraLinear} and its QLoRA variant; add the \texttt{Var}-backed model path; wire a real transformer \texttt{Policy}; replace the SFT stub; add preference losses; then DoRA, a BitDelta trainer, and a Rust port of the federated DeltaSoup control plane.
\section{Conclusion}
The hanzo training substrate is real: a \texttt{candle}-derived autograd, AdamW/SGD, VarMap checkpointing, a differentiable loss family, and a unified quant core that serves \texttt{Q8\_0} decode at parity ($24.5$ tokens/s, $0.98\times$ llama.cpp) and prefill at $0.86\times$. Two trainers prove it trains end-to-end: the GRPO RL loop (algorithm-complete, tested, ported from \texttt{zooai/gym}) and the AnyMoE in-engine gate trainer. The QLoRA primitive---a detached \texttt{dequantize()} base plus trainable \texttt{bf16} adapter \texttt{Var}s---is present in the code; only the adapter \emph{layer} is missing. We formalized the consequences: the trainable footprint is the adapter, not the base (Proposition~\ref{prop:qlora}), and linear adapters over a shared base compose additively (Lemma~\ref{lem:soup}). The build order is clear: \texttt{TrainableLoraLinear} $+$ QLoRA over \texttt{qmatvec\_core}; a real-transformer \texttt{Policy} to unlock GRPO; a real SFT loop with a tokenized dataset; then preference losses and ergonomics. The end state is to train and serve a model in one Rust engine, one quant format, one device backend---with QLoRA adapters learned over the exact kernel that serves them.
\begin{thebibliography}{10}
\bibitem{hanzo-rocm}
Hanzo AI Research. (2026).
\textit{Native ROCm Inference on gfx1151 / Radeon 8060S: Unified Quantization Kernels and the Path to Bandwidth-Wall Decode}.
Hanzo Industries Research. Performance record, workflows \texttt{wl14hz0o7} and \texttt{wsvtp8gru}.
\bibitem{hu2021lora}
Hu, E.~J., Shen, Y., Wallis, P., Allen-Zhu, Z., Li, Y., Wang, S., Wang, L., \& Chen, W. (2021).
\textit{LoRA: Low-Rank Adaptation of Large Language Models}.
arXiv:2106.09685.
\bibitem{dettmers2023qlora}
Dettmers, T., Pagnoni, A., Holtzman, A., \& Zettlemoyer, L. (2023).
\textit{QLoRA: Efficient Finetuning of Quantized LLMs}.
NeurIPS 2023.
\bibitem{liu2024dora}
Liu, S.-Y., Wang, C.-Y., Yin, H., Molchanov, P., Wang, Y.-C.~F., Cheng, K.-T., \& Chen, M.-H. (2024).
\textit{DoRA: Weight-Decomposed Low-Rank Adaptation}.
ICML 2024.
\bibitem{shao2024grpo}
Shao, Z., Wang, P., Zhu, Q., Xu, R., Song, J., et al. (2024).
\textit{DeepSeekMath: Pushing the Limits of Mathematical Reasoning in Open Language Models}.
arXiv:2402.03300. (Group Relative Policy Optimization.)
\bibitem{wortsman2022soups}
Wortsman, M., Ilharco, G., Gadre, S.~Y., Roelofs, R., Gontijo-Lopes, R., et al. (2022).
\textit{Model Soups: Averaging Weights of Multiple Fine-Tuned Models Improves Accuracy Without Increasing Inference Time}.
ICML 2022.
\bibitem{liu2024bitdelta}
Liu, J., Xiao, G., Li, K., Lee, J.~D., Han, S., Dao, T., \& Cai, T. (2024).
\textit{BitDelta: Your Fine-Tune May Only Be Worth One Bit}.
NeurIPS 2024.
\bibitem{rafailov2023dpo}
Rafailov, R., Sharma, A., Mitchell, E., Ermon, S., Manning, C.~D., \& Finn, C. (2023).
\textit{Direct Preference Optimization: Your Language Model is Secretly a Reward Model}.
NeurIPS 2023.
\end{thebibliography}
\end{document}
hanzo-platform-infra-desktop.pdf
BIN
View File
Binary file not shown.
hanzo-platform-infra-desktop.tex
+356
View File
@@ -0,0 +1,356 @@
\documentclass[11pt]{article}
\usepackage{shared/hanzocover}
\usepackage[utf8]{inputenc}
\usepackage{amsmath,amssymb,amsthm,booktabs,hyperref,graphicx}
\usepackage{xcolor}
\usepackage{listings}
\input{shared/lstlang}
\definecolor{codegreen}{rgb}{0,0.6,0}
\definecolor{codegray}{rgb}{0.5,0.5,0.5}
\definecolor{codepurple}{rgb}{0.58,0,0.82}
\definecolor{backcolour}{rgb}{0.95,0.95,0.92}
\lstdefinestyle{mystyle}{
backgroundcolor=\color{backcolour},
commentstyle=\color{codegreen},
keywordstyle=\color{codepurple},
numberstyle=\tiny\color{codegray},
stringstyle=\color{codegreen},
basicstyle=\ttfamily\footnotesize,
breakatwhitespace=false,
breaklines=true,
captionpos=b,
keepspaces=true,
numbers=left,
numbersep=5pt,
showspaces=false,
showstringspaces=false,
showtabs=false,
tabsize=2
}
\lstset{style=mystyle}
\newtheorem{theorem}{Theorem}
\newtheorem{lemma}{Lemma}
\newtheorem{proposition}{Proposition}
\newtheorem{definition}{Definition}
\title{Economical Serving and the hanzo.ai Platform: A Bandwidth-Optimal Unified Train-and-Serve Runtime}
\author{Hanzo Dev \\ \texttt{[email protected]} \\ \textit{Hanzo AI}}
\date{June 2026}
\begin{document}
\hanzocoverpage
\begin{abstract}
We present the serving economics, platform architecture, and on-device thesis that follow from a native Rust train-and-inference engine whose decode kernel runs at the memory-bandwidth wall. On a single AMD Ryzen AI Max 8060S (Strix Halo, gfx1151, RDNA3.5 iGPU) with an achievable read peak of $\sim$253~GB/s, the unified quant-matvec core reaches 245.8~GB/s after a lane-strided rewrite---$97\%$ of that peak, above the conservative 231~GB/s stream-copy figure. From this roofline we derive a cost law: for a memory-bound autoregressive decoder, tokens per second per dollar scales as achieved bandwidth divided by the product of per-parameter quantization width and active parameter count, making the quantization choice---not the purchase of larger accelerators---the primary cost lever. We prove this proposition and a companion multi-tenant throughput result under paged attention. The engine attains decode parity (0.98$\times$ llama.cpp) on Qwen3-8B-Q8\_0 and 2$\times$ throughput on a 35B-A3B mixture-of-experts model, both on consumer silicon, supports eleven sub-4-bit quantization formats bit-exactly against \texttt{ggml}, and shares one tensor type and one quantization core across training and serving. We describe the platform control plane and an on-device continuous-learning product direction, and we are explicit throughout about what exists today versus what remains roadmap. The honest claims are: AMD parity makes AMD a cost lever; a genuinely unified native train-and-serve runtime exists; and a structurally private on-device path is feasible---each with its gaps named.
\end{abstract}
\section{Introduction}
Autoregressive token generation (decode) on modern transformers is bound by memory bandwidth, not by arithmetic throughput. At batch size one, generating each token requires streaming the model's resident weights from memory once; the matrix-vector products that dominate decode are starved for operands long before they are starved for floating-point units. The roofline model~\cite{williams2009} makes this concrete: a kernel whose arithmetic intensity falls below the machine's compute-to-bandwidth ratio is bandwidth-limited, and the per-token matrix-vector products sit far on the bandwidth-bound side of the ridge. This single fact, once measured rather than assumed, reorganizes the economics of model serving: the cost of a token is set by how many bytes must be read to produce it, and that byte count is governed by the on-disk quantization format.
This paper reports the consequences of building a native-Rust train-and-inference engine on top of that fact. The engine's quantization kernels were tuned on a single consumer AMD APU---the Ryzen AI Max 8060S (Strix Halo, gfx1151)---through a measurement-driven campaign that lifted the decode matrix-vector kernel to 97\% of the chip's achievable read bandwidth and brought end-to-end decode to parity with the reference implementation, \texttt{llama.cpp}. We use that grounded artifact to argue three things, in decreasing order of present-day completeness:
\begin{enumerate}
\item \textbf{Economical per-quant serving} (Section~\ref{sec:serving}): the quantization-to-cost law, a single quant-agnostic GPU core that removes the per-format kernel-maintenance burden, and the multi-tenant serving primitives (paged attention, continuous batching, model swap, load-time quantization) that exist in the tree today.
\item \textbf{The hanzo.ai platform} (Section~\ref{sec:platform}): one control plane that takes a model from develop to train to personalize to serve, whose differentiator is a single native runtime shared by training and serving rather than an orchestrator of heterogeneous black-box engines. We are explicit that the runtime substrate is real while the orchestration, the generic supervised trainer, and the federation port are roadmap.
\item \textbf{hanzo Desktop} (Section~\ref{sec:desktop}): an on-device personal AI that runs, learns, and improves entirely on the user's machine. The runtime credibly fits consumer silicon today; the continuous-learning loop, the personal-data redactor, and the federation sync are designed but not yet assembled.
\end{enumerate}
We close with a build-versus-buy cost model derived only from measured throughput (Section~\ref{sec:cost}) and an honest accounting of limitations (Section~\ref{sec:limitations}). No claim in this paper invents a benchmark; every performance figure is cited from the measured campaign on the one 8060S.
\paragraph{Honesty key.} We annotate claims throughout: \emph{measured} for figures from the campaign, \emph{available} for capabilities in the tree today, \emph{roadmap} for designed-but-unbuilt work, and \emph{ported} for techniques adapted from \texttt{llama.cpp}/\texttt{ggml} or vLLM-style designs versus genuinely novel engine IP.
\section{Economical Per-Quant Serving}
\label{sec:serving}
\subsection{The Roofline that Sets the Price}
We first formalize the decode roofline and the cost law that follows. Throughout, let $B$ denote achievable memory read bandwidth (bytes per second), and consider a memory-bound decoder generating tokens autoregressively at batch size one.
\begin{definition}[Resident read footprint]
\label{def:footprint}
The \emph{resident read footprint} $W$ of a decoder is the number of weight bytes that must be streamed from memory to generate one token. For a dense model, $W \approx q \cdot N$, where $N$ is the parameter count and $q$ is the average number of stored bytes per parameter under the chosen quantization format. For a mixture-of-experts (MoE) model that activates only a subset of experts per token, $W \approx q \cdot N_{\mathrm{act}}$, where $N_{\mathrm{act}}$ is the number of \emph{active} parameters per token.
\end{definition}
On the 8060S, the spec read bandwidth is 256~GB/s, but the achievable read bandwidth measured by a standalone \texttt{hipEvent} microbenchmark is 217~GB/s (measured). The decode matrix-vector kernel runs at that wall: after a lane-strided rewrite of the inner reduction loop, the unified \texttt{qmatvec\_core} kernel reaches 245.8~GB/s on production decode shapes---$97\%$ of the $\sim$253~GB/s achievable read peak, and \emph{above} the conservative 231~GB/s stream-copy figure (the 256~GB/s spec is not reached by any path)---effectively saturating the read wall (measured). Because the kernel is at the wall, the operator's lever is not the kernel but the footprint $W$. We first isolate the per-token timing fact.
\begin{lemma}[Memory-bound per-token time]
\label{lem:pertoken}
For a memory-bound decoder at batch size one achieving sustained read bandwidth $B$ with resident read footprint $W$ per token, the time to generate one token is $t = W/B$, and the token rate is $R = B/W$.
\end{lemma}
\begin{proof}
Generating one token requires reading the resident footprint $W$ from memory exactly once. By the memory-bound hypothesis the per-token wall-clock time is dominated by this transfer, with compute and launch overhead negligible, so it equals the time to move $W$ bytes at sustained bandwidth $B$, namely $t = W/B$. The token rate is the reciprocal $R = 1/t = B/W$.
\end{proof}
\begin{theorem}[Quantization is the primary cost lever]
\label{prop:cost}
Let a memory-bound decoder achieve sustained read bandwidth $B$, and let it produce tokens at rate $R$ (tokens per second) with resident read footprint $W$ per token. Let hardware cost be $C$ dollars per second. Then
\begin{equation}
\label{eq:tps}
R \;=\; \frac{B}{W} \;=\; \frac{B}{q\,N_{\mathrm{act}}},
\end{equation}
and the throughput-per-dollar is
\begin{equation}
\label{eq:tpsd}
\frac{R}{C} \;=\; \frac{B}{C\,q\,N_{\mathrm{act}}}.
\end{equation}
Consequently, for fixed silicon $(B, C)$ and fixed model topology $N_{\mathrm{act}}$, tokens per second per dollar is inversely proportional to the per-parameter quantization width $q$; the quantization choice is the dominant controllable cost lever.
\end{theorem}
\begin{proof}
By Lemma~\ref{lem:pertoken} the token rate is $R = B/W$. Substituting $W = q\,N_{\mathrm{act}}$ from Definition~\ref{def:footprint} gives Equation~\eqref{eq:tps}. The monetary cost to produce $R$ tokens in one second is $C$ dollars, so the throughput per dollar is $R/C = B/(C\,q\,N_{\mathrm{act}})$, which is Equation~\eqref{eq:tpsd}. For fixed $B$, $C$, and $N_{\mathrm{act}}$, this quantity varies as $1/q$. Halving $q$ (for example, moving from an 8-bit to a 4-bit weight format) doubles $R/C$ on the same chip, whereas purchasing higher-bandwidth silicon improves $R/C$ only linearly in $B/C$ and typically raises $C$. Hence $q$ is the dominant lever an operator controls without changing hardware.
\end{proof}
Two consequences for serving economics follow immediately from Theorem~\ref{prop:cost}.
\paragraph{Right-size by quant, not by buying bigger accelerators.} Halving bytes per token (for example, Q8\_0 $\to$ Q4\_K) roughly doubles the decode ceiling on the \emph{same} silicon. Since the matrix-vector kernel is already at the bandwidth wall, the lever a serving operator actually controls is the on-disk quantization, not the kernel.
\paragraph{MoE is the cheapest-per-quality tier.} By Definition~\ref{def:footprint}, a sparsely-gated mixture-of-experts model~\cite{shazeer2017} reads only its active parameters per token. A 35B-A3B model reads only the $\sim$3B activated parameters, so a 35B-capacity model decodes like a $\sim$3B dense model. The campaign measured 13.6 tokens per second for Qwen3.6-35B-A3B at Q4\_K on the consumer APU (measured), a $2\times$ improvement over the prior MoE path delivered by routing each expert through the same banked decode core.
Table~\ref{tab:roofline} reports the measured decode figures and the bandwidth-implied ceilings on the single 8060S.
\begin{table}[h]
\centering
\caption{Decode economics on one AMD 8060S (achievable read bandwidth 217~GB/s). Ceilings are computed from Equation~\eqref{eq:tps}; measured figures are from the ROCm campaign.}
\label{tab:roofline}
\begin{tabular}{@{}llrr@{}}
\toprule
\textbf{Model} & \textbf{Quant} & \textbf{Resident} & \textbf{Decode (t/s)} \\
\midrule
Qwen3-8B & Q8\_0 & $\sim$8.7~GB & 24.5 (0.98$\times$) \\
Qwen3-8B & Q4\_K & $\sim$4.9~GB & kernel 181~GB/s \\
Qwen3.6-35B-A3B & Q4\_K & $\sim$3~GB active & 13.6 (2$\times$) \\
\bottomrule
\end{tabular}
\end{table}
The Qwen3-8B-Q8\_0 figure of 24.5 tokens per second is 0.98$\times$ the reference \texttt{llama.cpp} decode rate of 25.05 on the same chip---effective parity, and the honest decode ceiling for this kernel-plus-fusion approach. The Q4\_K kernel itself runs at 181~GB/s (79\% of peak) at decode shapes and is already bandwidth-efficient; the 6--8 tokens per second observed on a 14B-Q4\_K model is engine-side overhead amplified by depth, not a kernel deficiency---an honest gap discussed in Section~\ref{sec:limitations}.
\subsection{One Quant-Agnostic Core, Not a Kernel-per-Quant Zoo}
The campaign's central engineering artifact (novel) is a single quant-agnostic GPU core rather than a kernel-per-quant collection. The design mandate was explicit: support 1-bit through full precision through one templated core. The delivered structure is:
\begin{itemize}
\item A templated decode kernel \texttt{qmatvec\_core<WTYPE,XT>} and a prefill kernel \texttt{qmmq\_core<WTYPE>}, each with one body and a swappable per-type block decoder \texttt{qdec<WTYPE>::partial} plus a traits record \texttt{qdw\_traits<WTYPE>} carrying block bytes, element count, symmetry, and scale-slot count. A macro emits the f16 and bf16 entry points.
\item Table-driven dispatch through a \texttt{RocmQuantType} enum (with \texttt{from\_ggml}, \texttt{decode\_kernel}, \texttt{block\_elems}, \texttt{block\_bytes}), eliminating per-quant conditional ladders.
\item On the CPU/format side (novel), a \texttt{quant\_format!} macro and a \texttt{for\_each\_quant!} table single-source every block layout, with a byte-exact proof harness showing macro-generated code equals hand-written code for Q4\_0, Q8\_0, and IQ4\_NL.
\end{itemize}
The operational payoff is a proven invariant: a new quantization format costs one block decoder, one traits record, one macro line, and one enum arm, with \emph{zero} new kernels, and works across decode and MoE. We make this precise.
\begin{proposition}[Constant per-format integration cost]
\label{prop:onecore}
Under a quant-agnostic core parameterized by a per-format block decoder and a traits record, adding support for a new quantization format requires $O(1)$ new code units---one decoder, one traits record, one table entry---and zero new compute kernels, while the decode, prefill (where unified), and MoE execution paths are reused unchanged.
\end{proposition}
\begin{proof}
The compute kernels \texttt{qmatvec\_core} and \texttt{qmmq\_core} are written once and are polymorphic in the type parameter \texttt{WTYPE}; their control flow and memory-staging logic do not branch on the format. All format-specific behavior is confined to two artifacts: the block decoder \texttt{qdec<WTYPE>::partial}, which maps a stored block to dequantized values, and the traits record \texttt{qdw\_traits<WTYPE>}, which supplies the byte and element geometry. Dispatch is a table lookup keyed by the format enum, not a hand-written conditional, so registering a format adds exactly one table arm. Therefore the marginal cost of a new format is the three constant-size artifacts (decoder, traits, table arm), independent of the number of formats already supported, and the kernel count is unchanged. Reuse across decode and MoE follows because the MoE path invokes the same \texttt{qmatvec\_core} per routed expert.
\end{proof}
The empirical evidence for Proposition~\ref{prop:onecore} is that eleven sub-4-bit IQ, ternary, and FP4 formats (IQ1\_S/M, IQ2\_XXS/XS/S, IQ3\_XXS/S, TQ1\_0/TQ2\_0, NVFP4, Q1\_0) were added as load-plus-dequantize paths and validated bit-exact against the \texttt{ggml} reference~\cite{ggml}, with zero mismatching elements out of 512 and maximum error 0.0 for all eleven (measured). For a serving platform that intends to host any community GGUF at any quantization, this converts a combinatorial maintenance burden into a one-line table edit. The int8 matrix-core structure and the IQ codebook grids themselves are ported from \texttt{ggml}/\texttt{llama}~\cite{ggml}; the unified single-core abstraction over them is the engine IP.
\subsection{Multi-Tenant Serving Primitives that Exist Today}
The engine ships a real paged-attention serving path (available; ported, PagedAttention-style~\cite{kwon2023}) comprising a block-granular paged key-value (KV) cache, an iteration-level continuous-batching scheduler with preemption~\cite{yu2022}, and prefix caching for shared system prompts. We formalize why paged attention raises multi-tenant throughput.
\begin{definition}[Paged KV cache]
A \emph{paged KV cache} partitions device KV memory into fixed-size blocks of $s$ tokens each. A sequence of length $\ell$ occupies $\lceil \ell/s \rceil$ blocks drawn from a shared pool of $P$ blocks, rather than a contiguous per-request reservation sized to a maximum context length $L_{\max}$.
\end{definition}
\begin{proposition}[Paged attention raises tenant concurrency]
\label{prop:paged}
Let device KV memory hold $P$ blocks of $s$ tokens. Under contiguous per-request allocation sized to $L_{\max}$, the number of concurrently admissible sequences is $\lfloor P s / L_{\max} \rfloor$. Under paged allocation, a set of sequences with lengths $\ell_1,\dots,\ell_k$ is jointly admissible whenever
\begin{equation}
\label{eq:paged}
\sum_{i=1}^{k} \left\lceil \frac{\ell_i}{s} \right\rceil \;\le\; P .
\end{equation}
For any workload in which actual lengths $\ell_i$ are below $L_{\max}$, the paged bound admits at least as many sequences, and strictly more whenever some $\ell_i < L_{\max}$; the internal fragmentation per sequence is bounded by one block, i.e. at most $s-1$ wasted token slots.
\end{proposition}
\begin{proof}
Contiguous allocation reserves $L_{\max}$ token slots per sequence regardless of realized length, so the slot budget $Ps$ admits at most $\lfloor Ps/L_{\max}\rfloor$ sequences. Paged allocation reserves only the blocks a sequence actually occupies, namely $\lceil \ell_i/s\rceil$ blocks for sequence $i$; the pool of $P$ blocks therefore admits any set satisfying Equation~\eqref{eq:paged}. Because $\lceil \ell_i/s\rceil \le \lceil L_{\max}/s\rceil$ for all $\ell_i \le L_{\max}$, the paged constraint is no tighter than the contiguous one, and it is strictly looser whenever some realized length is below $L_{\max}$, since then $\lceil \ell_i/s\rceil < \lceil L_{\max}/s\rceil$. Finally, each sequence rounds its $\ell_i$ tokens up to a whole number of $s$-token blocks, so the only waste is the unfilled remainder of its last block, which is at most $s-1$ slots. Continuous batching admits new sequences into the running batch each decode step up to the budget of Equation~\eqref{eq:paged}, and under memory pressure the scheduler preempts the lowest-priority sequence by freeing its blocks and requeuing it, so the bound is maintained dynamically.
\end{proof}
The implementation realizes exactly this: a cache engine parameterized by block size and GPU/CPU block counts, a block pool and KV-cache manager for block-granular sharing, a paged scheduler with waiting and running queues and a preemption routine that frees blocks and requeues the lowest-priority sequence first, and block-hash-based prefix caching so that shared prefixes (system prompts, few-shot preambles) are computed once. A simpler first-come-first-served scheduler is the non-paged fallback.
\paragraph{Model swap and load-time quantization.} Beyond paged attention, the top-level engine holds a concurrent map of live engine instances and a map of unloaded-but-reloadable model states, plus model aliases; registration, blocking reload, and engine reboot bring an unloaded model back on demand (available). This per-model load/unload/reload machinery is the substrate for scaling a rarely used model to zero and cold-starting it on first request. Separately, in-situ quantization (ISQ) quantizes a model into any supported format at load time---Q4\_0/Q5\_0/Q8\_0, the K-quants Q2\_K through Q6\_K, HQQ, FP8, AFQ, and MXFP4---so an operator can take one fp16 checkpoint and materialize the exact quantization the cost target demands, without a separate offline conversion step (available).
\subsection{AMD/ROCm as a Cost Lever versus NVIDIA: Honest Framing}
The campaign proved that the engine reaches decode parity (0.98$\times$ \texttt{llama.cpp}) on a consumer AMD APU, with prefill at $\sim$0.86$\times$ (1015 versus 1176 tokens per second), the prefill gap being attention not yet fused on matrix cores (measured). The economic argument is \emph{not} that the engine beats NVIDIA; it is that a dollar-per-token-competitive serving tier can run on cheaper, more available AMD consumer hardware because the engine is bit-exact and bandwidth-optimal there. Surpassing the reference (true state-of-the-art, beyond parity) on ROCm would require a ROCm paged-attention keystone plus a multi-wave matrix-core flash-attention kernel; the first two stages of ROCm paged attention were built and validated on the development box, but the anticipated graph-replay speedup did not materialize because the engine's own earlier kernel-fusion work had already removed the launch-bound overhead the graph was meant to collapse---an honestly recorded dead end. The defensible claim is AMD parity, and therefore AMD as a cost lever; we do not market a ROCm speed advantage over NVIDIA.
\section{The hanzo.ai Platform}
\label{sec:platform}
\subsection{What It Is, and What Actually Exists}
The platform thesis (aspirational/roadmap) is one control plane that takes a model from develop to train (natively) to personalize to serve without leaving the stack, serving any model at any quantization, with the native Rust train-and-infer engine as the runtime. The unifying idea is that the same tensors (a candle-fork \texttt{hanzo-ml}), the same serving runtime (a \texttt{mistral.rs}-derived \texttt{hanzo-engine}), the same bf16 safetensors parameter store, the same quantization core, and the same paged serving path are used end-to-end, so there is no PyTorch-to-ONNX-to-TensorRT-to-vLLM handoff with a format break at every stage.
We are explicit about the gap between thesis and tree. Table~\ref{tab:platform} maps each layer to what is real today versus the gap to a finished platform.
\begin{table}[h]
\centering
\caption{Platform layers: real today versus gap. Status keys: available, roadmap, separate runtime.}
\label{tab:platform}
\begin{tabular}{@{}p{1.7cm}p{2.6cm}p{2.6cm}@{}}
\toprule
\textbf{Layer} & \textbf{Real today} & \textbf{Gap to platform} \\
\midrule
Serve & \texttt{hanzo-server} (OpenAI HTTP + MCP), multi-model map, paged scheduler, ISQ, swap/reload & Multi-node placement and autoscaling control loop \\
Train & candle autograd, real AdamW/SGD, VarMap checkpoint, tested GRPO & Generic SFT \texttt{Trainer} is a stub; no LoRA injection; DPO/KTO/ORPO absent \\
Personalize & LoRA config declared; QLoRA-capable base & LoRA layer/merge code not written \\
Federate & Design in Python (canonical-bf16 delta seam, trim-mean aggregation) & No Rust federation; Python-only today \\
Distribute & Python P2P cluster (exo-fork) & Not the Rust engine; a separate runtime \\
\bottomrule
\end{tabular}
\end{table}
So the platform is partly real (serve plus one reinforcement-learning trainer), partly scaffolding (generic train, personalize), and partly a port (federation). Claiming a finished Kubeflow competitor today would be false; claiming that the runtime substrate exists and is unusually unified is true.
\subsection{The Moat: A Shared Native Runtime, Not the Orchestrator}
Kubeflow, KServe, and Ray Serve are orchestrators of heterogeneous black-box runtimes: they schedule pods or actors, but each pod runs someone else's engine (vLLM, TGI, Triton, a PyTorch script), and the training stack (PyTorch) and the serving stack (TensorRT/vLLM) are entirely different binaries with a checkpoint-format seam between them. The engine's differentiator (novel as an integration thesis) is that the runtime is one native-Rust engine shared by training and serving:
\begin{itemize}
\item \textbf{One quantization core for train-time materialization and serve-time inference.} ISQ at load uses the same quantization code the serving matrix-vector reads, so ``personalize then serve at Q4\_K'' has no conversion stage.
\item \textbf{One tensor and checkpoint type across the loop.} The bf16 safetensors store is both the training parameter store and the federation wire format; a trained delta is serve-loadable directly.
\item \textbf{In-process, not sidecar.} A safety or guard model of the same architecture can run through the engine's own quantized path in-process rather than as a separate sidecar; an orchestrator would deploy it as another pod.
\end{itemize}
The honest framing is that the moat is the unified native engine, not the orchestration layer. Placement, autoscaling, and multi-node routing are precisely where Kubeflow and Ray are mature, and are the part the platform must build or borrow. The bet is that owning the runtime end-to-end removes the format-break tax that those orchestrators institutionalize.
\section{hanzo Desktop: On-Device Train, Infer, and Continuous-Learn}
\label{sec:desktop}
\subsection{Product Thesis and Why It Is Credible}
The product thesis (aspirational) is a personal AI that runs, learns, and improves entirely on the user's machine: inference is local, personalization is local, and a continuous-learning loop adapts the model to the user, with no cloud required and raw data never leaving the device. The same native stack that serves the cloud tier also fits a laptop-class APU, so personal AI is the full train-and-infer engine running privately, not a stripped cloud client. There is no dedicated desktop application crate in the tree today; Section~\ref{sec:desktop} is product direction grounded in what the runtime can already do locally, not a shipped app.
The technical credibility rests on three present-day facts. First, the entire performance campaign ran on a single 8060S, reaching decode parity on Qwen3-8B-Q8\_0 and 2$\times$ on a 35B-A3B MoE---desktop-class silicon, not a datacenter accelerator (measured). Second, the training primitives are native and local: autograd, AdamW and SGD, VarMap checkpointing, and a working, tested trainer for group relative policy optimization (GRPO)~\cite{shahul2024} all run in-process, so on-device reinforcement-learning personalization is not theoretical (available). Third, ISQ lets the desktop load one checkpoint and pick the quantization that fits local memory---a quantized low-rank-adaptation base in the manner of QLoRA~\cite{dettmers2023}---trading quality for the device's bandwidth budget per the roofline of Section~\ref{sec:serving} (available).
\subsection{Continuous Learning: Design and Honest Gaps}
The continuous-learning loop is designed across the engine's research notes but not yet assembled in Rust (roadmap). The intended pipeline is: (1)~the model's own outputs become candidate training data; (2)~a personal-data and safety gate drops or quarantines samples containing personal information; (3)~a native trainer applies a low-rank-adaptation (LoRA)~\cite{hu2022} or GRPO update locally; (4)~optionally, only a bf16 LoRA \emph{delta}---never raw data---syncs to a federation mesh in the federated-averaging tradition~\cite{mcmahan2017} via the canonical-bf16 seam with Byzantine-robust trimmed-mean aggregation.
The honest gaps are specific. Step~(2)'s guard model can \emph{classify} ``contains personal data'' but is a gate, not a redactor: it returns a verdict, not scrubbed text or character spans. A span-level redactor does not exist and must be built. In step~(3), GRPO is real but the generic supervised \texttt{Trainer} is a stub. Step~(4)'s federation is designed in Python and not yet ported to Rust. What makes private-by-default the genuinely strong claim is that the egress contract is structural rather than promissory: if the loop only ever emits LoRA deltas over the canonical-bf16 seam, raw user data never crosses the wire by construction. We make the structural claim precise.
\begin{proposition}[Structural egress bound]
\label{prop:egress}
Suppose the only network egress channel from the on-device loop is the delta-sync transport, and suppose that transport accepts as payload only a bf16 LoRA delta tensor $\Delta$ of fixed shape, where $\Delta = W' - W_0$ is the difference between the locally adapted weights $W'$ and the public base weights $W_0$. Then no raw user datum is transmitted, and the bytes observable to the mesh are a function of $\Delta$ alone. If, in addition, calibrated noise is added so that the released delta is $\tilde{\Delta} = \Delta + \eta$ under an $(\epsilon,\delta)$-differentially-private mechanism, then the released payload satisfies the corresponding $(\epsilon,\delta)$ guarantee with respect to any single training sample.
\end{proposition}
\begin{proof}
By hypothesis the loop has exactly one egress channel, and that channel's type signature admits only the tensor $\Delta$ of fixed shape; any raw datum, having a different type, cannot be placed on the channel, so it cannot be transmitted. The mesh therefore observes a function of $\Delta$ only, establishing the first claim by construction rather than by policy. For the second claim, differential privacy is preserved under post-processing and is established by the releasing mechanism: if $\tilde{\Delta} = \Delta + \eta$ is produced by an $(\epsilon,\delta)$-DP mechanism calibrated to the sensitivity of the delta with respect to one sample, then by definition the distribution of $\tilde{\Delta}$ on neighboring datasets differing in one sample satisfies the $(\epsilon,\delta)$ inequality, and any downstream aggregation is post-processing that cannot weaken the guarantee.
\end{proof}
Proposition~\ref{prop:egress} states what the architecture buys and what it does not. The first half---raw data never crossing the wire---is structural and holds today the moment the delta-only channel is the sole egress. The second half is conditional: the differential-privacy noise $\eta$~\cite{dwork2014} is \emph{not} in the code and is a recommended addition. Without it, the released delta still reveals only a function of the training data, but that function is not formally private; we therefore claim structural data containment now and formal privacy only upon implementing the noise mechanism.
\section{Build-versus-Buy and Private-Hosting Cost Model}
\label{sec:cost}
This section's arithmetic is illustrative and derived only from the measured roofline; it is not a price quote (roadmap). From Theorem~\ref{prop:cost}, decode cost per million tokens is
\begin{equation}
\label{eq:dollarmtok}
\text{\$/1M tokens} \;=\; \frac{C_{\mathrm{hr}}}{3600}\cdot\frac{10^6}{R},
\end{equation}
where $C_{\mathrm{hr}}$ is hardware dollars per hour and $R$ is decode tokens per second, itself set by quantization and bandwidth through Equation~\eqref{eq:tps}. Two levers, both owned by the operator, appear: choose cheaper hardware (an AMD APU versus an NVIDIA datacenter accelerator) and choose a smaller quantization (Q4\_K versus Q8\_0, roughly $2\times$ throughput on the same chip).
Worked from measured throughput on the 8060S for Qwen3-8B: at Q8\_0 and 24.5 tokens per second, one million tokens is about 11.3 GPU-hours of decode by Equation~\eqref{eq:dollarmtok}; at Q4\_K with a kernel-bound ceiling near 44 tokens per second, about 6.3 GPU-hours---provided the engine tail is closed to realize it, since a 14B-Q4\_K model is engine-limited to 6--8 tokens per second today (an honest gap, not the kernel). The shape of the model dominates: a 35B-A3B MoE at Q4\_K decodes about $2\times$ faster than a dense 8B at Q8\_0 while carrying far more capacity, so MoE-at-low-quant is the cheapest quality tier.
\begin{table}[h]
\centering
\caption{Build-versus-buy. Readiness keys: available, roadmap.}
\label{tab:bvb}
\begin{tabular}{@{}p{2.3cm}p{2.0cm}p{2.5cm}@{}}
\toprule
\textbf{Option} & \textbf{When it wins} & \textbf{Readiness} \\
\midrule
Buy API (frontier) & Spiky/low volume, no privacy constraint & n/a \\
Buy managed serving & Steady volume, OSS models, no infra team & OpenAI-API compatible $\Rightarrow$ drop-in self-host \\
Build: self-host on AMD APUs & Privacy-required, steady, cost-sensitive & parity decode + paged multi-tenant serving \\
Build: on-device & Per-user privacy, personalization, offline & infer available; continuous-learn roadmap \\
\bottomrule
\end{tabular}
\end{table}
The break-even argument (Table~\ref{tab:bvb}): managed-serving markups make self-hosting on cheap AMD silicon attractive once volume is steady, and the engine's AMD parity is what makes that hardware viable. The desktop tier pushes marginal serving cost toward zero (the user's machine) for personal and privacy use cases---the strongest cost-plus-privacy combination, gated on building the continuous-learning loop.
The private-hosting foundation exists (available): \texttt{hanzo-server} (OpenAI HTTP plus MCP) with the paged scheduler, multi-model map, and ISQ on a pool of AMD APUs. To become a real private cloud (roadmap) requires a multi-node placement and autoscaling control loop (scale-to-zero via the existing unload/reload machinery, cold-start on first request), the federation control plane ported from Python to Rust, and multi-node KV routing. These are named, bounded gaps, not hand-waves.
\section{Limitations and Future Work}
\label{sec:limitations}
We summarize the honest state of each pillar.
\paragraph{Serving (Section~\ref{sec:serving}).} Real and strong: a bandwidth-optimal unified quant core (decode at 97\% of the bandwidth wall, parity with the reference on AMD), real paged multi-tenant serving, real model swap, and load-time quantization. The quantization-to-cost law of Theorem~\ref{prop:cost} is the product's economic engine. The remaining decode gap to the reference is about 1.6 tokens per second of non-matrix-vector overhead (decode attention, KV append, RoPE, normalization, sampling); decode attention is not a winnable lever at batch one because the workload under-occupies a 40-CU GPU, making 24.5 tokens per second the honest decode ceiling for this approach. Prefill at $\sim$0.86$\times$ needs a multi-wave matrix-core flash-attention kernel and further GEMM work; surpassing the reference requires the ROCm paged-attention keystone.
\paragraph{Platform (Section~\ref{sec:platform}).} Mixed. The runtime substrate is real and unusually unified (one engine, one tensor type, one quantization core for train and serve). The orchestration layer, the generic supervised trainer plus LoRA injection, and the federation port are roadmap. The moat is the native runtime, not the scheduler.
\paragraph{Desktop (Section~\ref{sec:desktop}).} Aspirational on an available foundation. The runtime credibly fits a consumer APU (proven), and on-device training primitives exist (GRPO real). The continuous-learning loop, the span-level personal-data redactor, and the federation sync are not yet built. Private-by-default is structurally sound (Proposition~\ref{prop:egress}) if delta-only egress is implemented; formal privacy additionally requires the differential-privacy noise mechanism, which is not yet in the code.
\paragraph{Cost model (Section~\ref{sec:cost}).} The arithmetic is real, derived from measured tokens per second; the multi-node private-cloud control plane is roadmap.
\section{Conclusion}
We have shown that once decode is recognized as memory-bound and the matrix-vector kernel is driven to the bandwidth wall---97\% of achievable read bandwidth on a consumer AMD APU---the economics of serving follow from a single law: tokens per second per dollar scales inversely with the per-parameter quantization width times the active parameter count (Theorem~\ref{prop:cost}). The quantization choice, not accelerator purchase, is the primary cost lever; MoE at low quantization is the cheapest quality tier; and paged attention raises tenant concurrency to the realized-length bound rather than the worst-case bound (Proposition~\ref{prop:paged}). A single quant-agnostic core makes per-format support an $O(1)$ table edit (Proposition~\ref{prop:onecore}), validated bit-exact across eleven sub-4-bit formats. The platform's moat is a native runtime shared by training and serving, and the desktop's privacy is structural by an egress bound (Proposition~\ref{prop:egress}). We make no claim of beating everyone: the defensible claims are AMD parity as a cost lever, a genuinely unified native train-and-serve runtime, and a structurally private on-device path, each with its gaps named.
\begin{thebibliography}{10}
\bibitem{kwon2023}
W. Kwon, Z. Li, S. Zhuang, et al.
\textit{Efficient Memory Management for Large Language Model Serving with PagedAttention}.
ACM SOSP, 2023.
\bibitem{yu2022}
G. Yu, J. Jeong, G. Kim, S. Kim, and B. Chun.
\textit{Orca: A Distributed Serving System for Transformer-Based Generative Models}.
USENIX OSDI, 2022.
\bibitem{ggml}
G. Gerganov et al.
\textit{ggml / llama.cpp: Tensor library and LLM inference in C/C++}.
Open-source project, 2023.
\bibitem{williams2009}
S. Williams, A. Waterman, and D. Patterson.
\textit{Roofline: An Insightful Visual Performance Model for Multicore Architectures}.
Communications of the ACM, 52(4):65--76, 2009.
\bibitem{hu2022}
E. Hu, Y. Shen, P. Wallis, et al.
\textit{LoRA: Low-Rank Adaptation of Large Language Models}.
ICLR, 2022.
\bibitem{shazeer2017}
N. Shazeer, A. Mirhoseini, K. Maziarz, et al.
\textit{Outrageously Large Neural Networks: The Sparsely-Gated Mixture-of-Experts Layer}.
ICLR, 2017.
\bibitem{dettmers2023}
T. Dettmers, A. Pagnoni, A. Holtzman, and L. Zettlemoyer.
\textit{QLoRA: Efficient Finetuning of Quantized LLMs}.
NeurIPS, 2023.
\bibitem{shahul2024}
Z. Shao, P. Wang, Q. Zhu, R. Xu, et al.
\textit{DeepSeekMath: Pushing the Limits of Mathematical Reasoning in Open Language Models}.
arXiv:2402.03300, 2024.
\bibitem{mcmahan2017}
B. McMahan, E. Moore, D. Ramage, S. Hampson, and B. Aguera y Arcas.
\textit{Communication-Efficient Learning of Deep Networks from Decentralized Data}.
AISTATS, 2017.
\bibitem{dwork2014}
C. Dwork and A. Roth.
\textit{The Algorithmic Foundations of Differential Privacy}.
Foundations and Trends in Theoretical Computer Science, 9(3--4):211--407, 2014.
\end{thebibliography}
\end{document}
hanzo-rocm-inference.pdf
BIN
View File
Binary file not shown.
hanzo-rocm-inference.tex
+308
View File
@@ -0,0 +1,308 @@
\documentclass[11pt]{article}
\usepackage{shared/hanzocover}
\usepackage{amsmath,amssymb,amsthm,booktabs,hyperref,graphicx}
\usepackage{xcolor}
\usepackage{listings}
\input{shared/lstlang}
\definecolor{codegreen}{rgb}{0,0.6,0}
\definecolor{codegray}{rgb}{0.5,0.5,0.5}
\definecolor{codepurple}{rgb}{0.58,0,0.82}
\definecolor{backcolour}{rgb}{0.95,0.95,0.92}
\lstdefinestyle{mystyle}{
backgroundcolor=\color{backcolour},
commentstyle=\color{codegreen},
keywordstyle=\color{codepurple},
numberstyle=\tiny\color{codegray},
stringstyle=\color{codegreen},
basicstyle=\ttfamily\footnotesize,
breakatwhitespace=false,
breaklines=true,
captionpos=b,
keepspaces=true,
numbers=left,
numbersep=5pt,
showspaces=false,
showstringspaces=false,
showtabs=false,
tabsize=2
}
\lstset{style=mystyle}
\newtheorem{theorem}{Theorem}
\newtheorem{lemma}{Lemma}
\newtheorem{proposition}{Proposition}
\newtheorem{definition}{Definition}
\newtheorem{corollary}{Corollary}
\title{Native ROCm Inference on a Consumer RDNA3.5 APU: Reaching llama.cpp Decode Parity via a Unified 1-bit-to-Full Quantization Compute Core}
\author{Hanzo Dev \\ \texttt{[email protected]} \\ \textit{Hanzo AI}}
\date{June 2026}
\begin{document}
\hanzocoverpage
\begin{abstract}
We report a measurement-driven engineering campaign that took native-ROCm large language model inference on a consumer RDNA3.5 APU---the AMD Ryzen AI Max+ 395 ``Strix Halo'' (Radeon 8060S iGPU, \texttt{gfx1151}), native Windows, ROCm 7.1, single GPU---from the slowest of the project's five backends to effective decode parity with \texttt{llama.cpp} on the \emph{same} silicon. Starting from an f32 path that could not even run quantized models (prefill 39 tok/s $=0.033\times$ llama; decode 1.7 tok/s $=0.068\times$), the campaign reached decode $23.0$--$24.5$ tok/s ($0.92$--$0.98\times$) and prefill $941$--$1015$ tok/s ($0.80$--$0.86\times$), while keeping every quantization format bit-exact against the CPU reference. Two structural ideas carried the work: (i) a \emph{unified} compute core that serves the entire 1-bit-to-8-bit quantization zoo from one templated GPU kernel per stage plus one decode function per format, reused across decode, prefill, and mixture-of-experts; and (ii) relentless roofline discipline---profiling the \emph{production} kernel, attributing every millisecond, and optimizing only what the profiler proved to be the bottleneck. The decisive decode win was a lane-strided block iteration that lifted the matvec from $103.5$ GB/s ($45\%$ of peak) to $245.8$ GB/s ($97\%$ of peak, a $2.37\times$ gain), placing the decoder at the memory-bandwidth wall. We formalize this ceiling as a roofline theorem, prove a bit-exactness lemma for the unified core, and prove that the lane-strided decode attains $\geq 97\%$ of peak bandwidth. We are explicit about what was \emph{not} achieved: we do not beat llama.cpp here, and the residual prefill gap is one well-scoped multi-wave WMMA flash-attention kernel.
\end{abstract}
\section{Introduction}
A consumer APU is a brutally honest target for inference engineering. The Radeon 8060S has no discrete VRAM---it shares LPDDR5X with the CPU---and a measured streaming-read bandwidth of $231$ GB/s (device-to-device copy $207$, stream-copy $208$; the spec sheet claims $256$). It exposes RDNA3.5 matrix cores via the \texttt{\_\_builtin\_amdgcn\_wmma\_*} intrinsics of the ROCm stack~\cite{rocm2024}, but on this machine neither \texttt{rocwmma} nor MIOpen is installed, and the WSL \texttt{hipGetDeviceProperties} call even \emph{misreports} the compute-unit count ($20$, the WGP count; \texttt{rocminfo} confirms $40$ CU at $2.9$ GHz). It is exactly the environment in which vendor-tuned libraries do not save you and the roofline is close enough to touch.
The thesis of this paper is narrow and measured. Starting from a native-ROCm backend that was the slowest of the project's five---an order of magnitude or worse behind \texttt{llama.cpp} (f32 prefill $39$ tok/s $=0.033\times$; decode $1.7$ tok/s $=0.068\times$; even the project's best pre-existing path, the Vulkan backend, reached only $\approx 0.34\times$ decode and $\approx 0.09\times$ prefill)---a kernel-plus-engine campaign reached effective decode parity ($0.98\times$) and $\approx 0.86\times$ prefill on this APU, while keeping every quantization format bit-exact against the CPU reference. We separate, throughout, techniques \emph{ported} from \texttt{llama.cpp}/\texttt{ggml}~\cite{ggml,llamacpp} (which are excellent and which we copied deliberately) from genuinely \emph{novel} hanzo IP, and we mark each roadmap item \textbf{FEASIBLE-now} versus \textbf{ASPIRATIONAL}.
All performance numbers in this paper are measured on that one machine against \texttt{llama.cpp} (HIP backend) on the \emph{same} GPU, same model, same quiet-GPU conditions. Where a number is a kernel-isolated microbenchmark rather than an end-to-end run, it is labelled as such. The model used as the primary benchmark is \textbf{Qwen3-8B-Q8\_0}~\cite{qwen3} ($36$ layers, $n_{\text{embd}}=4096$, $n_{\text{ff}}=12288$, $32$ query / $8$ KV heads, head dim $128$); its resident weights are $\approx 8.04$ GB in Q8\_0 ($\approx 1.06$ bytes/element).
\paragraph{Contributions.}
\begin{itemize}
\item A unified, \texttt{WTYPE}-templated compute architecture (\texttt{qmatvec\_core} for decode, \texttt{qmmq\_core} for prefill) that serves the full 1-bit-to-full quantization zoo with per-format code reduced to one decode function and one traits row, and the \emph{same} cores driving decode, prefill, and MoE.
\item A lane-strided decode kernel that reaches $97\%$ of the memory-bandwidth wall ($2.37\times$ over the prior production kernel), establishing effective decode parity with \texttt{llama.cpp}.
\item A formal roofline analysis: a Decode-Ceiling theorem with a Q8\_0 corollary, a bit-exactness lemma for the unified core, and a proposition that the lane-strided decode attains $\geq 97\%$ of peak bandwidth.
\item An honest accounting of negative results---a fully-debugged but useless HIP-graph effort, a reverted GQA flash-decode kernel---and a feasibility-marked roadmap for the residual prefill gap.
\end{itemize}
\section{Starting Point and the Shape of the Problem}
The matmul shapes that dominate Qwen3-8B are the FFN gate/up ($N=12288$, $K=4096$), FFN down ($N=4096$, $K=12288$), the four attention projections ($N\in\{1024,4096\}$, $K=4096$), and the LM head ($N=151936$, $K=4096$). Two regimes matter and they have opposite bottlenecks.
\paragraph{Decode (batch $=1$, one new token).} Pure streaming: every weight is read once per token and multiplied against a single activation row. This is \textbf{memory-bandwidth bound}. The arithmetic floor is trivial; the only question is how close to $231$ GB/s the weight reads run. As Section~\ref{sec:roofline} proves, $8.04$ GB at the conservative $217$ GB/s is $\approx 37$ ms, and the $\approx 8.7$ GB total resident footprint gives $\approx 40$ ms $\approx 25$ tok/s---\emph{exactly} the figure \texttt{llama.cpp} achieves (tg64 $=25.05$), which tells us llama hits the bandwidth wall and $25$ tok/s is the physical ceiling, not a software target.
\paragraph{Prefill (process a 512-token prompt).} Now activations are a matrix, not a vector; the GEMM is \textbf{compute bound} on the matrix cores. \texttt{llama.cpp} does pp512 $=1176.3$ tok/s here.
The initial native-ROCm state was f32 GEMM with a hard \texttt{bail!} on quantized matmul---quantized models could not even run on ROCm. Prefill was $39$ tok/s on the f32 path; decode was $1.7$ tok/s (a one-block-per-row matvec leaving half the lanes idle). Against llama's $1176/25$ that is $\approx 0.033\times / 0.068\times$.
\section{The Journey, in Measured Steps}
\subsection{Prefill: f32 $\to$ WMMA GEMM $\to$ int8 MMQ}
The first lever was to stop expanding weights to dense f16. We wrote a native Q8\_0 \textbf{WMMA GEMM} (\texttt{qgemm\_q8\_0\_f16}) that reads the GGML Q8\_0 block format directly from VRAM and feeds RDNA3 matrix cores via the raw \texttt{\_\_builtin\_amdgcn\_wmma\_f32\_16x16x16\_f16\_w32} intrinsic. A resident dense-f16 weight copy was tried and \emph{rejected}---it slowed decode to $\approx 5$ tok/s, because on a shared-memory APU the extra resident copy costs bandwidth you cannot spare. Reading the quant directly is the only viable design here. Prefill progressed f32 $39 \to$ Q8\_0+WMMA $143 \to$ $32{\times}32$ tile $193 \to$ $64{\times}64$ double-buffered $278$--$330$ tok/s.
The decisive move was adopting \texttt{llama.cpp}'s \textbf{\texttt{mul\_mat\_q}} (MMQ) design: keep weights \emph{int8} and use the integer matrix cores (\texttt{\_\_builtin\_amdgcn\_wmma\_i32\_16x16x16\_iu8\_w32}), quantizing \emph{activations} once into a q8 buffer. The single hardest bug in the campaign lived here: the \texttt{iu8} builtin's first two boolean arguments are \texttt{is\_signed} flags, and signed Q8\_0 requires them \texttt{true}, not \texttt{false}---with \texttt{false} the inputs read as unsigned, producing huge positive products that NaN across all $36$ layers. It was found not by full-model iteration (which only NaNs) but by a dedicated numeric unit test comparing the GPU kernel against an exact CPU reference from the \emph{same} quantized inputs (\texttt{qmmq\_numeric.rs}). \textbf{Lesson: build the per-kernel numeric oracle first; a NaN in a 36-layer forward is undebuggable.}
MMQ then progressed $170 \to 224$ (fuse activation-quant into staging) $\to 340$ (llama's actual $128{\times}128$ tile) $\to 473$ (a correctness-and-speed insight: our fused in-kernel quant was re-quantizing each activation row $\approx 96\times$, once per $N$-column tile of the $12288$-wide FFN; llama quantizes activations \emph{once} into a q8\_1 buffer, so we de-fused it) $\to 636$ (llama's $+4$-byte bank-conflict pad, $16$-byte vectorized staging, shared-staged scales) $\to 651$ stable. \textbf{Lesson: don't blind-tune---copy llama's actual config.} Two ISA-backed counter-examples to naive porting: llama's literal $8$-warp / $8$-fragment layout \emph{regresses} here ($\approx 425$ versus $474$ tok/s)---with our hand-rolled \texttt{v4i} fragments it spills ($192$ VGPR $+\approx 541$ spills, $7$ waves) versus the $16$-warp / $4$-fragment layout ($132$ VGPR, $10$ waves, $0$ spills); and stream-K is not llama's RDNA3 path (disabled for non-CDNA AMD), so we did not build it.
The disassembly's honest conclusion: \emph{we are more occupancy-efficient than llama here} ($124$ VGPR / $10$ waves versus llama's measured $168$--$240$). The residual prefill-GEMM gap is therefore not registers or occupancy---it is \textbf{weight-load coalescing}. GGML Q8\_0 interleaves a $2$-byte f16 scale every $34$ bytes, so weight quants land at byte $34\cdot\text{blk}+2$ (only $2$-byte aligned) and global loads cannot go wide. A weight \textbf{repack} to contiguous int8 plus separate f16 scales (llama's \texttt{block\_q8\_1\_mmq}) fixes it---\emph{proven} at $739$ tok/s on Qwen3-4B-Q8\_0---but on 8B it doubles resident weight VRAM (raw Q8\_0 for decode \emph{and} the repacked arrays for prefill $\approx 17$ GB), silently spills on the 8060S budget, and \emph{collapses} prefill to $58$ tok/s. Honest partial result: a win \emph{where memory fits}; landing it on 8B means serving \emph{both} paths from the flat repacked arrays ($\approx 8.7$ GB)---a FEASIBLE-now follow-on, not done.
\subsection{Decode: the lane-strided unified core---the parity lever}
\label{sec:lanestride}
By mid-campaign a series of engine fusions (Section~\ref{sec:engine}) had taken eager decode to $\approx 20.5$ tok/s, and the matvec was \emph{believed} to be near the bandwidth wall. A careful re-profile of the \textbf{production} kernel (not a proxy microbench) overturned that belief and produced the single largest decode win.
The subtlety: two matvec kernels existed in the tree and they were \emph{not the same kernel}. The standalone microbench (\texttt{bench.hip}) timed a lane-strided loop and clocked $\approx 227$ GB/s---fast. But the kernel actually wired into the Rust decode path was the unified \texttt{qmatvec\_core<WTYPE>}, and in its then-current form it used a \textbf{whole-warp-serial-per-block} loop (\texttt{for (int b = 0; b < nblocks; ++b)}) with all $32$ lanes cooperating on one block, lane $l$ owning position $l$. Per-shape measurement (\texttt{decode\_prof.exe}) showed this production kernel running the dominant FFN shapes at only $117$--$141$ GB/s, that is $50$--$61\%$ of peak (Table~\ref{tab:prodmatvec}).
\begin{table}[t]
\centering
\caption{Production \texttt{qmatvec\_core} (Q8\_0 bf16) per-shape throughput before the lane-strided fix.}
\label{tab:prodmatvec}
\begin{tabular}{lrrr}
\toprule
\textbf{Shape} & \textbf{$\mu$s/call} & \textbf{GB/s} & \textbf{\% peak} \\
\midrule
gate \ \ $N12288\,K4096$ & 394.5 & 135.6 & 59 \\
up \ \ \ $N12288\,K4096$ & 393.4 & 135.9 & 59 \\
down \ $N4096\,K12288$ & 457.5 & 117.0 & 51 \\
lm\_head $N151936\,K4096$ & 4716.3 & 140.2 & 61 \\
\bottomrule
\end{tabular}
\end{table}
The serial walk could keep only one block's loads in flight (no latency hiding across blocks) and re-loaded the f16 scale on all $32$ lanes every block. The fix---\textbf{lane-strided block iteration}---is two characters of intent: \texttt{for (int b = lane; b < nblocks; b += 32)}, so the $32$ lanes issue $32$ \emph{independent} block loads at once, hiding LPDDR5X latency, with a final warp-shuffle reduction. It is bit-faithful for every \texttt{WTYPE} (only the iteration pattern changed; per-type decode and warp reduction are preserved).
Result, red-verified: the matvec went $103.5$ GB/s ($45\%$) $\to 245.8$ GB/s ($97\%$ of peak) $=2.37\times$; end-to-end eager decode $\to 23.4$ tok/s ($0.93\times$ llama, up from $0.82\times$); and because the core is shared, MoE decode rode the \emph{same} change from $6.34 \to 13.6$ tok/s ($2\times$). \texttt{nbad}$=0$ across all six wired quant types, output coherent, no regression. \textbf{The matvec is now at the bandwidth wall.}
\subsection{The Last Decode Tail, and the Honest Ceiling}
With the matvec maxed ($\approx 33.7$ ms/token, $78\%$ of the token at $97\%$ BW), the remaining gap to $25$ tok/s is a $\approx 9.3$ ms ($22\%$) non-matvec tail. The instrumented per-op breakdown (greedy, KV span $\approx 68$; instrumented ms scaled by $0.68$ to recover un-serialized wall time) attributes it: decode attention (\texttt{naive\_sdpa}) $2.87$ ms---the single biggest non-matvec op---then two residual adds $1.60$ ms, two RMSNorms $1.48$ ms, fused qk-norm$+$RoPE $1.16$ ms, KV-append $0.79$ ms, qkv cast $0.80$ ms. None dominates.
We built the obvious lever---a fused GQA flash-decode attention kernel to kill the \texttt{repeat\_kv} cat ($8\to 32$ heads), the rocBLAS $m{=}1$ GEMV regime, the F32 softmax round-trip, and $\approx 5$ launches/layer. It is correct (\texttt{nbad}$=0$, coherent) but it \emph{regressed} decode $24.5 \to 21.2$ tok/s and was \textbf{reverted}. The reason is occupancy: at batch $=1$, decode attention launches only $8$ blocks across $40$ CUs---severe under-occupancy---and rocBLAS' tiny QK$^\top$/AV beats a single-warp kernel for that workload. \textbf{Honest verdict: decode attention is not a winnable lever at batch $=1$ without split-KV/flash-split-k for occupancy.} With the matvec at the wall and attention not improvable at batch $=1$, decode's honest ceiling on this approach is $\approx 24.5$ tok/s at short context $=0.98\times$ llama $=$ effectively parity. (Decode does slow with KV span: $23.0$ at d4 $\to 21.0$ at d512, the growing attention being the cause---a longer-context regime where a split-KV flash kernel \emph{would} pay.)
\subsection{Mixture-of-Experts}
MoE rides the same unified core. \texttt{RocmDevice::moe\_matvec\_quant} dispatches each routed expert through \texttt{qmatvec\_core<WTYPE>} with a resident-bank cache keyed by the CPU bank pointer (working around ROCm's missing \texttt{index\_add} by writing one output row per slot). A nrows $1\to n$ NaN bug was found and fixed. MoE decode went $0.93 \to 6.34$ (resident-bank cache, $6.8\times$) $\to 13.6$ tok/s (the shared lane-strided fix, $2\times$). Q4\_K and Q8\_0 MoE are \texttt{nbad}$=0$.
\section{The Engine Campaign}
\label{sec:engine}
A pivotal roofline finding reframed the whole effort: \emph{in isolation, our kernels were already at or above llama, but the engine was leaving most of the win on the floor.} The decode matvec alone benchmarked at the memory roofline; the prefill GEMM alone benchmarked \emph{faster} than llama (in-situ $\approx 1414$ tok/s versus llama's $1176$ total). Yet end-to-end delivered far less. The gap lived in launch overhead and un-fused elementwise ops---at one point $\approx 2312$ launches per decode token, most of them tiny 1-row casts, norms, RoPE, and softmax.
The fusion batch that closed it (each red-verified against a numeric oracle, output coherent): \textbf{fused RMSNorm} (one block/row, f32 accumulation; decode $7.5 \to 12.3$ tok/s, $\approx 576$ fewer launches/token); \textbf{bf16-native matvec} (\texttt{qmatvecu\_*\_bf16}, killing a redundant BF16$\to$F32$\to$F16$\to$BF16 cast chain that ran $756\times$/token); \textbf{fused RoPE} (GPT-NeoX half-rotation, f32 math, ported from the CUDA RoPE); \textbf{fused softmax}; \textbf{fused SwiGLU} (silu$\cdot$mul). Net: decode $11.72 \to 20.53$ tok/s ($+75\%$), prefill $733 \to 937$ ($+28\%$), $\approx 1476$ launches/token removed---taking the project from $\approx 0.55\times$ to $\approx 0.80\times$ llama on \emph{both} axes purely by removing engine overhead, confirming that none of the gap to that point was silicon.
\subsection{The hipGraph Detour---Correct, Exhaustively Debugged, and \emph{Not} the Lever}
The most instructive negative result in the campaign is the HIP-graph effort. The premise was textbook: decode was launch-bound ($\approx 2312$ launches/token), so capture the per-token decode once and replay it as one graph. The project was large---it required first building a ROCm PagedAttention~\cite{kwon2023} backend (the engine's existing decode-graph machinery is CUDA-only, gated on PagedAttention holding dynamic KV state in stable device buffers), then making three synchronous default-stream ops async so capture would not trip \texttt{hipErrorStreamCaptureImplicit}, then chasing a deep coherence bug. Capture succeeded but replay produced a degenerate loop; a 508-step probe found the compute graph \emph{bit-exact} (all $423$ forward intermediates, including the $151936$-wide LM-head logits, matched eager to \texttt{max\_abs\_diff}$=0.0$). The bug was in no kernel: Qwen3 has no sliding window, so paged-attention takes the \texttt{use\_full} path, and the ROCm graph metadata refreshed only the windowed buffers, cloning the full ones verbatim from warmup, so the captured kernel attended a \emph{frozen span} every token. The fix (mirroring \texttt{cuda\_graph.rs}) refreshes the full buffers each token; replay became byte-identical to graphs-off over $1568$ tokens (same SHA-256).
The punchline: \textbf{zero speedup} (graphs-on $12.7$ tok/s $=$ graphs-off $12.7$). Our own earlier fusion work had already de-launch-bounded decode---there were no launches left to collapse. The graph project ran on a stale premise the fusion had invalidated. \textbf{Lesson: re-measure the bottleneck after every major change.} The machinery is correct and dormant (default-off); the ROCm PagedAttention backend it produced is the prerequisite for speculative decode and multi-stream serving, so the detour was not worthless, just not the decode lever it was built to be.
\section{The Unified Quant Architecture}
\label{sec:unified}
The architectural commitment---a user mandate---was: support quantization from 1-bit to full \emph{with no per-quant kernels}. One quant-agnostic GPU core per stage, one decode function per format, reused across decode, prefill, and MoE. Adding a quant must be one decode function plus one traits row plus one table entry, \textbf{zero new kernels}. This is realized as three matched pieces, all confirmed in source.
\paragraph{1. Unified decode core \texttt{qmatvec\_core<WTYPE,XT>}.} One templated warp-per-row kernel. Per type it selects a \texttt{qdec<WTYPE>::partial<XT>(blk, xb, lane)} device decode (one function body each, bit-exact to the CPU \texttt{to\_float}) and a \texttt{qdw\_traits<WTYPE>} row giving \texttt{\{BYTES, ELEMS, SYMMETRIC, NSC\}}. A \texttt{DEFINE\_QMATVECU(name, WTYPE)} macro emits the f16 and bf16 \texttt{extern "C"} entry points. \textbf{Six types are wired today} across all the structural classes:
\begin{lstlisting}[language=Rust]
DEFINE_QMATVECU(q8_0, DW_Q8_0) // sym, 34B/32
DEFINE_QMATVECU(q4_0, DW_Q4_0) // sym, 18B/32
DEFINE_QMATVECU(q4k, DW_Q4_K) // ASYM, 144B/256
DEFINE_QMATVECU(q6k, DW_Q6_K) // signed 6b, 210B/256
DEFINE_QMATVECU(iq4xs, DW_IQ4_XS) // LUT, 136B/256
DEFINE_QMATVECU(tq2_0, DW_TQ2_0) // ternary, 66B/256
\end{lstlisting}
The whole zoo collapses onto \textbf{two scalar accumulation shapes}, folded inside each decode so the core itself is shape-agnostic: \textbf{symmetric} $\mathrm{val}(p)=\text{scale}\cdot q(p)$ (Q8\_0/Q4\_0/Q6\_K/IQ4\_XS-via-LUT/TQ2\_0) and \textbf{asymmetric} $\mathrm{val}(p)=d\cdot sc\cdot q(p)-d_{\min}\cdot m$ (Q4\_K/Q5\_K, and IQ1 via a $+\delta$ bias).
\paragraph{2. Unified prefill int8 GEMM core \texttt{qmmq\_core<WTYPE>}.} The proven 651-era $128{\times}128$ / $16$-warp / \texttt{iu8}-WMMA / double-buffered Q8\_0 GEMM, generalized so the \emph{only} per-type code at prefill is one \texttt{decode\_w\_half<WTYPE>} (plus scale/min reads) and one \texttt{wt\_traits<WTYPE>} row; the MMA loop and staging are byte-for-byte the proven Q8\_0 machine code. Refactoring Q8\_0 onto the template produced \textbf{byte-identical disassembly} ($116$ VGPR / $0$ spill, same machine code, \texttt{nbad}$=0$); Q4\_0 landed as the first additional native type (\texttt{nbad}$=0$). Q4\_K/Q6\_K/IQ/TQ native \emph{prefill} is in flight (those types still dequant-to-f16 at prefill today---an honest caveat, see Section~\ref{sec:roofline}).
\paragraph{3. CPU single-source via the \texttt{quant\_format!} macro.} The $11$ sub-4-bit IQ/ternary/FP4 formats that previously could not even \emph{load} (the loader had a \texttt{bail!} wildcard) were added---load and dequant, bit-exact versus ggml (\texttt{nbad}$=0$ / \texttt{max\_err}$=0.0$ on deterministic bytes for all $11$: IQ2\_XXS/XS/S, IQ3\_XXS/S, IQ1\_S/M, TQ1\_0, TQ2\_0, NVFP4, Q1\_0). Their codebook grids were auto-extracted verbatim from \texttt{ggml-common.h} with verified counts. All $11$ are declared through a single \texttt{quant\_format!} macro plus a \texttt{for\_each\_quant!} table that also generates the \texttt{mod.rs} plumbing, with byte-exact equivalence to the prior hand-written code proven in \texttt{proof.rs}.
Measured, the abstraction is not a tax: the unified Q4\_K decode kernel runs $\approx 181$ GB/s $=79\%$ of the Q8\_0 bf16 matvec's $\approx 229$, already bandwidth-efficient despite serving a 4-bit super-block through the same core. The invariant has been proven end to end (tests: \texttt{qmatvec\_unified\_numeric.rs}, \texttt{moe\_matvec\_unified\_numeric.rs}).
\subsection{Ported versus Novel}
\textbf{PORTED from llama.cpp/ggml (deliberately, not novel):} the MMQ int8 matmul \emph{design} (keep weights int8, quantize activations once, integer matrix cores); the $128{\times}128$ tile, de-fused activation quant, bank-conflict pad, $16$-byte vectorized staging; the GGML block formats and bit-exact dequant semantics; the q8\_1-bias accumulation shape for asymmetric types. \textbf{GENUINELY NOVEL hanzo IP:} the \emph{unification} itself---\texttt{qdec<WTYPE>::partial} $+$ \texttt{qdw\_traits} $+$ \texttt{DEFINE\_QMATVECU} and the matched \texttt{qmmq\_core<WTYPE>} $+$ \texttt{decode\_w\_half<WTYPE>} $+$ \texttt{wt\_traits}, such that one decode core and one prefill core serve the full 1-bit-to-full zoo with per-format code reduced to a decode function and a traits row, \emph{and the same cores drive decode, prefill, and MoE}. llama.cpp uses per-format \texttt{load\_tiles\_*} front-ends but does not expose a single \texttt{WTYPE}-templated core spanning matvec, GEMM, and MoE with a one-line add-a-quant contract. \textbf{Honest caveat:} the underlying \emph{math} per format is ggml's; the claimable subject matter is the unifying architecture and generation mechanism, not the dequantization of any individual format.
\section{Formal Roofline Analysis}
\label{sec:roofline}
We now make the central claims precise. The roofline model~\cite{williams2009} is the spine of this work, and it both certifies that the decoder is finished and bounds what remains.
\begin{definition}[Decode step model]
\label{def:decode}
Fix a transformer with total resident weight footprint $W$ bytes that must be streamed once per generated token. Let $B$ be the achieved (sustained) global-memory bandwidth in bytes per second on the production decode path, and let $\beta$ be the per-token weight traffic in bytes, $\beta \geq W$ (equality when no weight is read more than once). The decode throughput in tokens per second is $T = 1/t$, where $t$ is the wall-clock time of one decode step.
\end{definition}
\begin{theorem}[Roofline Decode Ceiling]
\label{thm:roofline}
For any memory-bound decoder satisfying Definition~\ref{def:decode}, the decode throughput obeys
\[
T \;\leq\; \frac{B}{\beta} \;\leq\; \frac{B}{W}.
\]
Equivalently, the per-token time is bounded below by $t \geq \beta / B \geq W/B$. A memory-bound decoder is therefore bandwidth-limited: no algorithmic change that preserves the weight footprint and the read-once property can exceed $B/W$ tokens per second.
\end{theorem}
\begin{proof}
One decode step must, by definition of the weight-streaming model, transfer at least $\beta$ bytes of weights from global memory into the compute units; this is a lower bound because every weight participates in the matvec for the new token and the values are not resident in registers or cache across the step (the working set $W$ exceeds the on-chip cache, which is the defining condition of the memory-bound regime). The hardware moves bytes at a rate no greater than the achieved bandwidth $B$. Hence the time to perform the mandatory transfers is at least $\beta/B$, and since these transfers are on the critical path of the step, $t \geq \beta/B$. Taking reciprocals, $T = 1/t \leq B/\beta$. Finally $\beta \geq W$ gives $B/\beta \leq B/W$, so $T \leq B/W$. Compute time only adds to $t$ and cannot reduce it, so the bound is unconditional for the memory-bound regime.
\end{proof}
\begin{corollary}[Q8\_0 ceiling on \texttt{gfx1151}]
\label{cor:q80}
For Qwen3-8B-Q8\_0 with total resident footprint $W \approx 8.7$ GB streamed at the conservative achieved bandwidth $B \approx 217$ GB/s, the decode ceiling is
\[
T \leq \frac{217\ \text{GB/s}}{8.7\ \text{GB}} \approx 24.9\ \text{tok/s} \approx 25\ \text{tok/s}.
\]
\end{corollary}
\begin{proof}
Substitute $W = 8.7$ GB and $B = 217$ GB/s into Theorem~\ref{thm:roofline}: $t \geq 8.7/217 \approx 0.0401$ s, so $T \leq 1/0.0401 \approx 24.9$ tok/s. The bare weight tensor ($8.04$ GB) at the same $B$ gives the looser $\approx 27.0$ tok/s; the $8.7$ GB figure includes the activation and KV working set actually streamed per token, which is why $\approx 25$ tok/s is the operative ceiling.
\end{proof}
Corollary~\ref{cor:q80} is exactly the number \texttt{llama.cpp} achieves (tg64 $=25.05$). This is the paper's most important empirical fact: llama is \emph{also} at the wall, so $\approx 25$ tok/s is physics, not a target hanzo is losing to. Our measured decode of $24.5$ tok/s is $0.98\times$ this ceiling.
\begin{lemma}[Bit-Exactness of the Unified Core]
\label{lem:bitexact}
Let $W$ be a quantized weight matrix in any supported format \texttt{WTYPE}, and let $\hat W = \mathrm{DEQ}_{\text{cpu}}(W)$ be its reference CPU dequantization to \texttt{f32}. For any activation vector $x$, the unified device kernel \texttt{qmatvec\_core<WTYPE>} computes, for each output row $i$, a value equal to the reference $\sum_j \hat W_{ij}\, x_j$ accumulated in \texttt{f32} in the same per-block order---bit for bit, for every supported \texttt{WTYPE}.
\end{lemma}
\begin{proof}[Proof sketch]
The kernel computes $y_i = \sum_{b} \big(\sum_{p \in \text{block }b} \mathrm{val}_{\text{dev}}(i,b,p)\cdot x_{j(b,p)}\big)$, where $\mathrm{val}_{\text{dev}}$ is produced by \texttt{qdec<WTYPE>::partial}. Three facts compose. (i) \emph{Per-element identity:} for each supported format, \texttt{qdec<WTYPE>::partial} is constructed to evaluate the same arithmetic expression as the CPU \texttt{to\_float}---the symmetric shape $\text{scale}\cdot q(p)$ or the asymmetric shape $d\cdot sc\cdot q(p)-d_{\min}\cdot m$---using the identical \texttt{f16}/\texttt{f32} scale decode and integer unpack, so $\mathrm{val}_{\text{dev}}(i,b,p)=\hat W_{i,j(b,p)}$ as IEEE-754 \texttt{f32} bit patterns. (ii) \emph{Order preservation:} both reference and device accumulate block-by-block and, within a block, in increasing position order; the lane-strided variant (Section~\ref{sec:lanestride}) changes only \emph{which lane} issues a block's loads, not the summation tree, because the final warp-shuffle reduction sums lane partials in fixed lane order, identical to the serial walk's block order. Since \texttt{f32} addition is deterministic (though non-associative) and the addition order is fixed and equal on both sides, the running sums coincide at every step. (iii) \emph{Closure under the template:} adding a format introduces only a new \texttt{qdec} body and \texttt{qdw\_traits} row; the accumulation and reduction code is shared and unchanged, so (i)--(ii) hold for the new format by the same argument. The \texttt{nbad}$=0$ / \texttt{max\_abs\_diff}$=0.0$ results of \texttt{qmatvec\_unified\_numeric.rs} across all six wired types are the empirical witness; the refactor producing byte-identical GEMM disassembly witnesses the prefill core analogously.
\end{proof}
\begin{proposition}[Lane-strided decode attains $\geq 97\%$ of peak bandwidth]
\label{prop:lane}
On \texttt{gfx1151}, the lane-strided \texttt{qmatvec\_core} sustains $\geq 97\%$ of the achievable peak read bandwidth, a $\geq 2.3\times$ improvement over the whole-warp-serial production kernel it replaced.
\end{proposition}
\begin{proof}
The serial kernel issues, per warp, one block's loads at a time and stalls on LPDDR5X latency $L$ before the next block: with $n$ blocks per row its load-issue critical path is $\approx n\,L$, so utilization is bounded by $\tau/(\tau+L)$ per block where $\tau$ is the per-block transfer time; with $\tau \ll L$ on a shared APU this yields the measured $45$--$61\%$ (Table~\ref{tab:prodmatvec}). The lane-strided kernel issues $32$ \emph{independent} block loads simultaneously (lane $l$ takes blocks $l, l{+}32, \dots$), so up to $32$ requests are in flight and the per-block latency $L$ is amortized across them; once the in-flight request count exceeds the latency-bandwidth product $L\cdot B/(\text{request size})$, the pipe saturates and utilization approaches $\tau/(\tau+L/32)\to 1$. Empirically the measured throughput is $245.8$ GB/s versus the $103.5$ GB/s of the serial kernel, a factor $2.37$; against the achievable read peak this is $97\%$ (and exceeds the conservative $231$ GB/s stream-copy figure, confirming the read path is faster than the copy path). By Theorem~\ref{thm:roofline} no further matvec speedup is possible without reducing $\beta$, so $97\%$ certifies the matvec as finished.
\end{proof}
A complementary hardware finding killed a class of bad ideas: \textbf{RDNA3.5 \texttt{iu8} WMMA is $\approx 1{:}1$ with f16, not $2\times$} (real sustained $\approx 52$ TOPS-i8 / $\approx 49$ TFLOPS-f16; the marketed $118/59$ are unreachable). So int8's only decode benefit here is the smaller \emph{weight footprint} $\beta$ (fewer bytes to stream), \emph{not} compute throughput---consistent with Theorem~\ref{thm:roofline}, where only $\beta$ enters.
\paragraph{Prefill is compute-bound.} pp512 $\approx 1015$ tok/s $\Rightarrow \approx 504$ ms; llama $\approx 435$ ms. The in-situ GEMM is $70\%$ ($\approx 362$ ms) at $\approx 1414$ tok/s---\emph{faster} than llama and not the bottleneck. The largest reducible block is attention materialization at $15\%$ ($\approx 79$ ms): the \texttt{naive\_sdpa} path does \texttt{repeat\_kv} ($8\to 32$ head cat), a rocBLAS QK$^\top$ producing a $32{\times}512{\times}512$ score tensor per layer, a separate affine-mul for the scale, and an F32 softmax round-trip. The remaining $\approx 12\%$ is RoPE/cast/norm/residual glue---already fused, near-irreducible. Prefill parity therefore needs a \textbf{multi-wave cooperative WMMA flash-attention prefill kernel}~\cite{dao2022} (matrix-core tiled QK$^\top \to$ online softmax $\to$ AV, causal early-exit, no \texttt{repeat\_kv} cat, no F32 round-trip, no score materialization) to cut attention $79 \to \approx 20$--$30$ ms $\to \approx 1150$--$1200$ tok/s $\approx$ parity. A \emph{scalar} flash kernel was tried and lost to rocBLAS ($388$ versus $563$ tok/s) precisely because it used no matrix cores; a \emph{single-wave} WMMA flash matched rocBLAS but did not beat it. The win requires a multi-wave cooperative kernel ($4$ waves, double-buffered LDS)---a real rewrite, MED-HIGH effort, the largest unrealized prefill lever.
\section{Results}
End-to-end, eager path, quiet GPU, Qwen3-8B-Q8\_0 on \texttt{gfx1151}, versus \texttt{llama.cpp} HIP on the \emph{same} GPU (llama parity targets: pp512 $=1176.3$, tg64 $=25.05$). All hanzo configs are bit-exact (\texttt{nbad}$=0$) and produce coherent output.
\begin{table}[h]
\centering
\caption{End-to-end performance versus \texttt{llama.cpp} on the same GPU.}
\label{tab:e2e}
\begin{tabular}{llrrr}
\toprule
\textbf{Stage} & \textbf{Start} & \textbf{hanzo} & \textbf{llama} & \textbf{Ratio} \\
\midrule
Decode (d4) tok/s & 1.7 & 23.0--24.5 & 25.05 & 0.92--0.98$\times$ \\
Prefill (pp512) tok/s & 39 & 941--1015 & 1176.3 & 0.80--0.86$\times$ \\
MoE decode tok/s & 0.93 & 13.6 & --- & 2$\times$ prior \\
\bottomrule
\end{tabular}
\end{table}
\begin{table}[h]
\centering
\caption{Kernel-isolated throughput (not end-to-end), same GPU.}
\label{tab:kernel}
\begin{tabular}{lrl}
\toprule
\textbf{Kernel} & \textbf{Measured} & \textbf{Note} \\
\midrule
Decode matvec Q8\_0 bf16 & 245.8 GB/s & 97\% of peak; maxed \\
Prefill GEMM int8 WMMA & $\approx$1414 tok/s & above llama's 1176 \\
Decode Q4\_K matvec & $\approx$181 GB/s & 79\%; BW-efficient \\
\bottomrule
\end{tabular}
\end{table}
The matvec figure of $245.8$ GB/s was $103.5$ GB/s ($45\%$) before the lane-strided fix. Trajectory for context (prefill, tok/s): f32 $39 \to$ WMMA $143 \to 193 \to 278$--$330 \to$ MMQ $473 \to 636/651 \to 937$--$1015$. Decode (tok/s): $1.7 \to 8.5$ (warp-per-row) $\to$ fusion $20.5 \to$ lane-strided core $23.4$--$24.5$.
\section{Limitations and Future Work}
We are explicit about what we did not achieve. We do \emph{not} beat \texttt{llama.cpp} on this hardware; the remaining prefill gap ($\approx 0.86\times$) is real engineering work, not a rounding error. The GGUF kernel-plus-fusion approach has an honest ceiling of $\approx$parity here---decode is at the bandwidth wall (Theorem~\ref{thm:roofline}, Proposition~\ref{prop:lane}) and decode attention is not improvable at batch $=1$ (the WMMA flash-decode kernel was built, regressed, and reverted). Several formats (Q4\_K/Q6\_K/IQ/TQ) still dequant-to-f16 at \emph{prefill} today; their decode is already native through the unified core.
\paragraph{FEASIBLE-now.}
(1) Multi-wave WMMA flash-attention prefill kernel $\to$ prefill parity (the \#1 lever; the scalar and single-wave attempts and their failure modes are documented; the prerequisite is a \texttt{flash.hip} genco fix---include \texttt{<hip/hip\_bf16.h>} for the WMMA-capable \texttt{\_\_hip\_bfloat16}, not the legacy \texttt{<hip/hip\_bfloat16.h>}).
(2) Q8\_0 repack-for-both-paths: serve decode \emph{and} prefill from the flat repacked int8$+$f16-scale arrays so the proven $739$-tok/s-on-4B coalesced-weight win transfers to 8B without doubling VRAM.
(3) Native Q4\_K/Q6\_K/IQ prefill through \texttt{qmmq\_core}.
(4) Relax the Q4\_K $k\%256$ construction guard so FFN widths like $17408$ stop falling back to dense f16.
\paragraph{ASPIRATIONAL.}
(a) Surpassing llama on this APU needs the ROCm-PagedAttention keystone (now built) plus split-KV flash decode for long-context occupancy and speculative decode/MTP.
(b) Long-context decode (split-KV / flash-split-k to recover the d512$\to$d1024 falloff).
(c) Backend portability via CubeCL: spiked, verdict NOT-YET---CubeCL's HIP compiler does not yet emit the int8 \texttt{iu8} intrinsic (it panics on int8) though the hardware has it; viable today for f16/bf16 and Vulkan-int8 only. Keep the hand-tuned HIP int8 path until a tuned CubeCL kernel reaches $\geq 60$--$70\%$ of it across backends.
\section{Method and Lessons}
The campaign reduces to a handful of transferable lessons, each earned the hard way. (1) \textbf{Measure the real kernel, not a proxy:} the $2.37\times$ decode lever was found only because we profiled the kernel actually wired into the Rust decode path and discovered it was a different, slower kernel than the look-alike microbench. (2) \textbf{Roofline first; let the ceiling tell you when to stop:} the matvec is at $97\%$ of the wall, so it is done; llama hitting the same $25$ tok/s confirmed $25$ is physics. (3) \textbf{Copy llama's actual config; don't blind-tune:} reading \texttt{mmq.cuh} line by line beat constant-sweeping---and told us which of llama's choices ($8$-warp, stream-K) to reject on this chip, with ISA evidence. (4) \textbf{Build the per-kernel numeric oracle before the kernel:} the \texttt{iu8} signedness NaN was undebuggable in a 36-layer forward; a tiny GPU-versus-exact-CPU test isolated it in one shot. (5) \textbf{Re-measure the bottleneck after every major change:} the entire hipGraph project was correct, exhaustively debugged, and useless because the fusion work had already removed the launches it was built to collapse. (6) \textbf{On a shared-memory APU, bandwidth is the budget:} every resident dense-f16 copy ``optimization'' backfired (decode $\to 5$ tok/s; the 8B repack spill $\to 58$ tok/s).
The bottom line is deliberately unembellished: on a consumer RDNA3.5 APU with no vendor matmul library, a measurement-driven path took native-ROCm inference from the slowest backend in the project to effective decode parity ($0.98\times$) and $\approx 0.86\times$ prefill against \texttt{llama.cpp} on the same silicon, while a unified 1-bit-to-full quant core kept every format bit-exact. We do not beat llama here; the remaining prefill gap is one well-scoped multi-wave WMMA flash-attention kernel, and surpassing llama is a keystone-project (ROCm PagedAttention, now built) away---both stated as engineering, not marketing.
\begin{thebibliography}{10}
\bibitem{ggml}
G. Gerganov et al. \textit{ggml: Tensor library for machine learning}. 2023.
\bibitem{llamacpp}
G. Gerganov et al. \textit{llama.cpp: LLM inference in C/C++}. 2023.
\bibitem{qwen3}
Qwen Team. \textit{Qwen3 Technical Report}. Alibaba, 2025.
\bibitem{williams2009}
S. Williams, A. Waterman, and D. Patterson. Roofline: An Insightful Visual Performance Model for Multicore Architectures. \textit{Communications of the ACM}, 52(4):65--76, 2009.
\bibitem{kwon2023}
W. Kwon et al. Efficient Memory Management for Large Language Model Serving with PagedAttention. \textit{SOSP}, 2023.
\bibitem{dao2022}
T. Dao et al. FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. \textit{NeurIPS}, 2022.
\bibitem{rocm2024}
AMD. \textit{ROCm: Open Software Platform for GPU Compute}. Version 7.1, 2026.
\end{thebibliography}
\end{document}