% Author: Simon-Pierre Boucher — contact@spboucher.ai
\documentclass[11pt,a4paper]{article}

\usepackage[utf8]{inputenc}
\usepackage[T1]{fontenc}
\usepackage{lmodern}
\usepackage{amsmath,amssymb}
\usepackage{booktabs}
\usepackage{listings}
\usepackage{xcolor}
\usepackage{graphicx}
\usepackage[margin=1in]{geometry}
\usepackage{hyperref}
\hypersetup{colorlinks=true,linkcolor=blue,citecolor=blue,urlcolor=blue}
\usepackage{microtype}

\lstdefinestyle{msl}{
  basicstyle=\ttfamily\footnotesize,
  keywordstyle=\color{blue!70!black}\bfseries,
  commentstyle=\color{green!40!black}\itshape,
  stringstyle=\color{red!60!black},
  numbers=none, breaklines=true, frame=single, framesep=4pt,
  backgroundcolor=\color{gray!5}, showstringspaces=false,
  morekeywords={kernel,threadgroup,device,constant,constexpr,simdgroup_float8x8,
    simdgroup_matrix,enum,uint,template,typename,tensor,dextents,tensor_inline,half,bfloat}
}
\lstset{style=msl}

\title{\textbf{Forge}: Building an LLM Training Framework from Scratch\\
in C++ and Metal on Apple Silicon\\
\large Measured Lessons on Compiler Traps, Register Pressure,\\
and the M5 Neural Accelerators}

\author{Simon-Pierre Boucher\\
\texttt{contact@spboucher.ai}}

\date{31 July 2026}

\begin{document}
\maketitle

\begin{abstract}
We report engineering findings from \emph{Forge}, a transformer training framework
written from scratch in C++20 with hand-written Metal compute kernels, with no
PyTorch, MLX, or other machine-learning dependency. Every GPU kernel is validated
against a CPU reference to $\leq 10^{-4}$, and every optimization reported here was
accepted only after the training loss trajectory remained numerically unchanged.

Four findings are, we believe, of general interest to practitioners writing Metal
compute kernels. First, the idiom \texttt{constant constexpr} for tile constants in
the Metal Shading Language silently declares an \emph{address-space variable} rather
than a compile-time constant; the resulting loss of loop unrolling spills every
matrix accumulator to the stack and cost a factor of \textbf{12$\times$} on our GEMM.
Second, this class of defect is invisible in the AIR intermediate representation at
every optimization level, because unrolling and fragment promotion occur in the
driver's back end --- practitioners must benchmark rather than read the IR. Third,
per-thread register pressure, not bandwidth, is the dominant constraint in
attention backward kernels; measured spill counts guided three successive
restructurings that took the backward pass from 107\,ms to
7.05\,ms, a \textbf{15.2$\times$} improvement. Fourth, on
M5-generation hardware the Metal Performance Primitives cooperative-tensor
\texttt{matmul2d} reaches \textbf{51.5\,TFLOP/s} with \texttt{half} operands against
10.6\,TFLOP/s for a well-tuned hand-written \texttt{simdgroup\_matrix} kernel,
a \textbf{4.9$\times$} gap that inverts the usual cost/benefit assessment of mixed
precision.

We also report a negative result that we consider equally useful: \texttt{half} and
\texttt{bfloat} operands buy only 18--22\% over \texttt{float} on the
\texttt{simdgroup\_matrix} path, so mixed precision on that path is a memory
optimization and not a compute one. The framework trains models from 12M to 205M
parameters on a single machine.
\end{abstract}

\tableofcontents

\section{Introduction}

Apple Silicon offers a large unified memory pool and substantial GPU throughput, but
the software ecosystem for \emph{training} on it is thin. Inference frameworks are
mature; training-specific kernels, in particular a fused attention backward pass, are
largely absent. A survey we conducted on 31 July 2026 (Section~\ref{sec:survey})
found that of the major open-source frameworks, none ships a fused attention backward
kernel for Metal.

This paper reports what we learned building one. \emph{Forge} is a decoder-only
transformer training framework in pure C++20 with Metal compute kernels. Its
architecture is entirely configuration-driven, and it has no machine-learning
dependency: only Apple's \texttt{metal-cpp} bindings and a single-header JSON parser.

Our contribution is not the framework itself but the measurements. Each section below
states a hypothesis, the measurement that tested it, and the outcome --- including
the cases where the hypothesis was wrong.

\subsection{Methodology and reproducibility}

All measurements were taken on an Apple M5 Max (40-core GPU) running macOS 27.0
(build 26A5388g) with Xcode 26.6 and Metal toolchain 32023.883. Kernels are compiled
at \texttt{-std=metal3.2} except where Metal 4 features are required.
Throughput figures are computed from
\texttt{MTLCommandBuffer.GPUStartTime}/\texttt{GPUEndTime}, which excludes CPU
encoding time.

Two disciplines govern every result reported here.

\paragraph{Correctness gates optimization.} Every Metal kernel has a CPU reference
implementation, and 85 parity assertions compare them (maximum absolute error
$\leq 10^{-4}$ in \texttt{float}; most are bit-exact). Gradients are additionally
checked against central finite differences. An optimization is accepted only if the
full suite still passes.

\paragraph{Numerical equivalence, not just correctness.} For the larger kernel
rewrites we required that the \emph{training loss trajectory} be unchanged. After
replacing the scalar attention kernels with tiled matrix-multiply versions, the loss
at every logged step over 250 optimizer steps was identical
(8.3618, 7.6130, 6.6455, 5.0902, 4.9845; validation 4.9018). This is a stronger
check than parity on a single call, since it accumulates any discrepancy through the
optimizer.

\section{System overview}

The framework comprises roughly 6{,}500 lines of C++ and Metal. A \texttt{Tensor}
type provides shared-storage views over pooled \texttt{MTLBuffer} allocations in
\texttt{MTLStorageModeShared}, so the same memory is addressable from CPU and GPU
with no copies. A dynamic autograd tape records backward closures. Operations
dispatch to either a CPU reference backend or the Metal backend behind one interface,
which is what makes the parity testing possible.

Execution follows the batching discipline recommended for Apple GPUs: the whole
training micro-batch is encoded into one command buffer with a small number of
long-lived compute encoders, and the GPU is synchronized only at the loss-readback
boundary.

\section{Finding 1: \texttt{constant constexpr} is not a compile-time constant}
\label{sec:constexpr}

\subsection{The trap}

Tile geometry in a GEMM kernel is naturally expressed as named constants. The
apparently idiomatic C++ spelling is:

\begin{lstlisting}
constant constexpr uint TM = 4;   // WRONG in MSL
\end{lstlisting}

In the Metal Shading Language, however, \texttt{constant} is an \emph{address-space
qualifier}, and program-scope variables are \emph{required} to live in it. The
declaration above therefore creates a variable in the constant address space, not a
compile-time constant. Loop bounds derived from it are not known at compile time, so
the fragment loops do not unroll; \texttt{acc[i][j]} becomes dynamic indexing into an
array of opaque \texttt{simdgroup\_matrix} values; and the compiler is forced to
place all sixteen 8$\times$8 accumulators (256 bytes each) on the stack. Every
matrix-multiply-accumulate then pays roughly a kilobyte of stack traffic.

The correct spelling uses enumerators, which are true integral constant expressions:

\begin{lstlisting}
enum : uint { BM = 64, BN = 64, BK = 16, WM = 2, WN = 2,
              TM = BM / (8 * WM), TN = BN / (8 * WN) };
\end{lstlisting}

together with \verb|#pragma clang loop unroll(full)| on the fragment loops. Note also
that \verb|#pragma unroll full|, which appears in Apple's own header documentation,
is not valid here.

\subsection{Measurement}

The change is textual and semantically neutral. Its effect is not:

\begin{table}[h]
\centering
\begin{tabular}{lrr}
\toprule
Shape & \texttt{constant constexpr} & \texttt{enum} + unroll pragma \\
\midrule
$2048^3$          & 0.82 & 10.21 \\
$4096^3$          & 0.85 & \phantom{0}9.83 \\
$65536\times512\times1408$ & 0.83 & 10.68 \\
\bottomrule
\end{tabular}
\caption{f32 GEMM throughput (TFLOP/s), M5 Max. A 12$\times$ difference from the
declaration form of tile constants.}
\end{table}

Before this fix the ``optimized'' \texttt{simdgroup\_matrix} kernel was
\emph{three times slower} than a naive 16$\times$16 tiled kernel --- a result that
invites the wrong conclusion, namely that the matrix instructions are not worth
using.

\section{Finding 2: the intermediate representation cannot diagnose this}
\label{sec:ir}

Having observed the anomaly, our first instinct was to inspect the compiler output:

\begin{lstlisting}
xcrun metal -std=metal3.2 -O2 -S -emit-llvm -c matmul_simd.metal -o out.ll
\end{lstlisting}

The AIR showed three \texttt{alloca} instructions for the fragment arrays and only
two matrix-multiply intrinsics where thirty-two were expected --- apparently
confirming the spill hypothesis. It also, however, showed \emph{exactly the same
thing} at \texttt{-O0}, \texttt{-O2} and \texttt{-O3}, and continued to show it after
the fix that produced the 12$\times$ speedup.

The explanation is that loop unrolling and promotion of \texttt{simdgroup\_matrix}
values into registers happen in the driver's AIR-to-ISA back end, at pipeline
creation time, not in the front end that \texttt{-emit-llvm} exposes. The IR is
therefore uninformative for exactly the class of question one is most tempted to ask
of it.

\paragraph{Practical rule.} On Metal, benchmark; do not read the IR. We lost time to
this and record it because the failure mode is silent: the IR looks like a
confirmation.

\section{Finding 3: register pressure dominates attention backward}
\label{sec:registers}

\subsection{Three restructurings, each guided by measurement}

Our first fused attention implementation assigned one thread per query row and held
$q$, $o$ (forward) or $k$, $v$, $dk$, $dv$ (backward) in per-thread arrays. At head
dimension 64 the backward variant holds $4 \times 64 = 256$ floats, or one kilobyte
per thread.

We tested the hypothesis that this spills, by progressively reducing what each thread
holds:

\begin{table}[h]
\centering
\begin{tabular}{lrr}
\toprule
Attention backward, per layer & gpt-10m shapes & gpt-25m shapes \\
\midrule
$k,v,dk,dv$ all in registers        & 150 & --- \\
read-only $k,v$ from device         & 118 & 611 \\
$dK$ and $dV$ split into two kernels & 105 & 549 \\
\midrule
tiled with \texttt{simdgroup\_matrix} & \phantom{0}9.66 & \phantom{0}45.5 \\
\quad + $dK$/$dV$ split again        & \textbf{\phantom{0}7.05} & \textbf{\phantom{0}32.9} \\
\bottomrule
\end{tabular}
\caption{Time in milliseconds. gpt-10m shapes are $B{=}64$, $T{=}512$, $H{=}6$,
$d_h{=}64$; gpt-25m shapes are $B{=}64$, $T{=}1024$, $H{=}8$. Total improvement
15.2$\times$ and 17.0$\times$ respectively.}
\end{table}

The read-only operands $k$ and $v$ are re-read from device memory on every iteration
in the second row, which sounds wasteful; because a thread reads the same address
every iteration, the L1 cache serves it, and removing the two arrays from registers
is the larger effect.

\subsection{Confirming the mechanism directly}

The reasoning above was indirect. Xcode~26 ships \texttt{gpudebug}, a headless
command-line GPU debugger that reports compiler statistics per kernel. It requires
shader sources to be embedded in the library, which in turn requires
\texttt{-frecord-sources} at \emph{both} the compile and link steps. (The link must
be performed by \texttt{metal}; \texttt{metallib} rejects the flag.) With that fixed,
the mechanism is directly visible:

\begin{table}[h]
\centering
\begin{tabular}{lrrr}
\toprule
Kernel & Temp registers & Spilled bytes & Cost \\
\midrule
flash forward, scalar        & 126 & 368 & 38.9\% \\
flash forward, MMA-tiled     & \phantom{0}85 & \textbf{0} & 29.6\% \\
flash backward $dQ$, MMA     & \phantom{0}95 & \textbf{0} & \phantom{0}8.5\% \\
flash backward $dKV$, MMA (fused) & 111 & \textbf{4352} & 16.4\% \\
\bottomrule
\end{tabular}
\caption{Compiler statistics from \texttt{gpudebug}, M5 Max.}
\end{table}

Two things follow. The matrix-tiled forward eliminates the spill entirely
($368 \rightarrow 0$ bytes), which is the mechanism behind its speedup: an
8$\times$8 fragment occupies two floats per lane, where the scalar formulation held
whole arrays of head-dimension length. And the fused $dK/dV$ backward kernel
\emph{still spilled} 4352 bytes --- a defect we had not suspected, and which made it
the most expensive backward kernel. Splitting it recovered a further 27\%.

We regard the tooling lesson as the transferable one: register pressure is
measurable, cheaply and without a GUI, and it is worth measuring before restructuring
a kernel on intuition.

\section{Finding 4: cooperative tensors and the M5 neural accelerators}
\label{sec:mpp}

\subsection{Mixed precision on the classical path is a memory optimization}

The literature disagrees about \texttt{half} on Apple GPUs. Published
microbenchmarks report that \texttt{half} and \texttt{float} fused multiply-add
execute at the same rate, the benefit arising from register and bandwidth pressure;
Apple's material for the M3 generation describes up to twice the arithmetic
throughput from co-issue. Since plumbing mixed precision through a training framework
is a substantial amount of work, we measured before committing. Holding the tiling
and the \texttt{float} accumulator fixed and varying only the operand type:

\begin{table}[h]
\centering
\begin{tabular}{lrrr}
\toprule
Shape & \texttt{float} & \texttt{half} & \texttt{bfloat} \\
\midrule
$2048^3$ & \phantom{0}9.29 & 10.64 & 10.82 \\
$4096^3$ & \phantom{0}9.85 & 11.81 & 12.07 \\
$65536\times512\times1408$ & \phantom{0}9.90 & 12.45 & 12.44 \\
\bottomrule
\end{tabular}
\caption{\texttt{simdgroup\_matrix} throughput (TFLOP/s) by operand precision,
\texttt{float} accumulator throughout.}
\end{table}

The gain is 18--22\%, consistent with the register/bandwidth explanation rather than
with doubled arithmetic rate. On this path, mixed precision halves activation memory
--- valuable, since activation memory bounds trainable model size --- but it is not a
compute optimization.

\subsection{The same operands through \texttt{matmul2d}}

\texttt{MetalPerformancePrimitives.framework} ships in the macOS 26.5 SDK and exposes
\texttt{mpp::tensor\_ops::matmul2d}, a cooperative-tensor matrix multiply that targets
the per-core neural accelerators introduced with the M5 generation. Compiling at
\texttt{-std=metal4.0}, with a 64$\times$32 tile over four SIMD groups and a
\texttt{float} accumulator:

\begin{table}[h]
\centering
\begin{tabular}{lrrrr}
\toprule
 & \multicolumn{2}{c}{\texttt{simdgroup\_matrix}} & \multicolumn{2}{c}{\texttt{matmul2d}} \\
\cmidrule(lr){2-3}\cmidrule(lr){4-5}
Shape & \texttt{float} & \texttt{half} & \texttt{float} & \texttt{half} \\
\midrule
$2048^3$ & \phantom{0}9.3 & 10.6 & 14.9 & \textbf{51.5} \\
$4096^3$ & \phantom{0}9.9 & 11.8 & 14.6 & \textbf{44.3} \\
$65536\times512\times1408$ & \phantom{0}9.9 & 12.4 & 14.4 & \textbf{23.8} \\
\bottomrule
\end{tabular}
\caption{Throughput (TFLOP/s), M5 Max. A 1.5$\times$ gain in \texttt{float} and
4.3--4.9$\times$ in \texttt{half}.}
\end{table}

Because a 4.9$\times$ claim invites scepticism, we verified numerically rather than
only timing: against the CPU reference the \texttt{float} path is bit-exact, the
\texttt{half} path differs by $3.8 \times 10^{-6}$ from a reference fed identically
rounded inputs, and every output element is non-zero, confirming that the full
reduction over $K$ occurs. The transposed variants required for training
($X W^{\!\top}$ in the forward, $dX = dY W$ and $dW = dY^{\!\top} X$ in the backward)
are likewise bit-exact.

\paragraph{Generational caveat.} Published work on M4 Max hardware found
\texttt{matmul2d} still executing on the shader cores and losing to a hand-fused
GEMM. Our M5 result is the opposite. The accelerator hardware differs by generation;
this measurement should be repeated per target device rather than assumed.

\paragraph{Consequence.} This inverts the priority we had assigned to mixed
precision. Judged on the classical path it is worth 20\% and is chiefly a memory
feature; as the entry condition for the accelerator path it is worth 4.9$\times$.

\subsection{Interface notes}

Several details cost us time and are not documented accurately.

\begin{itemize}
  \item Bind ordinary buffers and construct tensors inside the kernel with the
        \texttt{tensor\_inline} descriptor. The default \texttt{tensor\_handle}
        descriptor wraps an opaque handle obtainable only from a host-side
        \texttt{MTLTensor}, which would require substantial host plumbing.
  \item Extents are ordered (columns, rows), and \texttt{slice()} takes column then
        row; a transposed operand must be sliced the other way and its extents
        swapped.
  \item The element type must be non-\texttt{const}; a \texttt{const float} tensor
        fails a static assertion.
  \item Zero the destination cooperative tensor using
        \texttt{is\_valid\_element(i)}. Apple's own header example calls
        \texttt{get\_mask(i)}, which does not exist in this SDK.
\end{itemize}

\section{Finding 5: concurrency, not Metal 4, is the encoding win}

Metal's default compute encoder orders every dispatch against its predecessor. For
the optimizer step this is pure loss: AdamW issues one dispatch per parameter tensor
--- of order one hundred for a 100M-parameter model --- and each gradient-norm
reduction is a single threadgroup. These are mutually independent.

We introduced a scoped concurrent-dispatch region and applied it to both optimizer
passes. On the 100M configuration end-to-end throughput rose from 7.9k to
\textbf{9.6k tokens/s}, a 22\% improvement, with identical losses. The change is
about twenty lines and requires no Metal 4 adoption.

An independent investigation confirmed the general shape of this result --- roughly
15$\times$ on a synthetic batch of small independent dispatches --- while also
establishing that Metal 4's command-encoding path offers essentially no CPU-side
advantage once invariant bindings are hoisted out of the dispatch loop. The
concurrency, available since macOS 10.14, is the entire effect.

\section{Ecosystem survey}
\label{sec:survey}

We surveyed the major open-source frameworks on 31 July 2026 for a fused attention
backward pass on Metal:

\begin{table}[h]
\centering
\begin{tabular}{ll}
\toprule
Framework & Fused attention backward on Metal \\
\midrule
MLX & No. \texttt{use\_fallback} returns \texttt{true}; \texttt{eval\_gpu} throws
      \texttt{"NYI"}. The fused backward is CUDA-only. \\
llama.cpp & No. The Metal backend does not support
            \texttt{GGML\_OP\_FLASH\_ATTN\_BACK}. \\
PyTorch MPS & No. Hand-written Metal forward kernels only. \\
Candle & No. Vendored MLX forward only. \\
tinygrad & No. Its flash-attention backward targets AMD. \\
Burn / CubeCL & Kernels exist but are a scalar scaffold, not wired to autodiff. \\
\bottomrule
\end{tabular}
\end{table}

MLX's Metal forward additionally declines the fused path under gradient tracing, with
the comment that unfused is faster for training on Metal. Production prior art is
limited to Philip Turner's \emph{metal-flash-attention} and one third-party package.
Both use the same three-stage, atomic-free structure we arrived at independently:
a preprocessing kernel for $D = \mathrm{rowsum}(dO \circ O)$ in \texttt{float}, a
$dQ$ kernel parallel over queries, and a $dK/dV$ kernel parallel over key/value rows.

\section{Attention kernel design}

Our tiled forward kernel assigns 32 query rows per threadgroup across four SIMD
groups, holding $Q$ and the output accumulator in registers as 8$\times$8 fragments
and staging $K$ and $V$ through threadgroup memory. The score tile is round-tripped
through threadgroup memory so that the softmax reductions run on ordinary threads.

\paragraph{A scoped problem worth recording.} Online softmax requires rescaling the
output accumulator by a per-row factor at every key/value block,
$O \leftarrow \mathrm{diag}(\mathrm{corr})\, O$. But the Metal specification leaves
the mapping from \texttt{simdgroup\_matrix} elements to lanes \emph{unspecified}, so a
lane cannot determine which row its registers correspond to. MLX resolves this by
reverse-engineering the mapping. We instead construct the 8$\times$8 diagonal matrix
in threadgroup memory and apply it with an ordinary matrix multiply. This is
specification-clean, costs $d_h/8$ additional multiplies per block (about 25\% more
matrix work), and keeps the accumulator in registers --- staging it in threadgroup
memory would have cost 8\,KB and halved residency. The backward pass needs none of
this, since the saved logsumexp already fixes the normalization.

A second technique worth noting: the $dK/dV$ kernel requires transposed score
matrices, which are obtained not by moving data but by exchanging the operand roles,
$S^{\!\top} = K Q^{\!\top}$ and $dP^{\!\top} = V\, dO^{\!\top}$, with
\texttt{simdgroup\_load(..., transpose=true)} supplying the transposed operands
directly from the staged tiles.

\section{Results}

\begin{table}[h]
\centering
\begin{tabular}{lrrr}
\toprule
Kernel & Naive & Tiled & \texttt{simdgroup\_matrix} \\
\midrule
$4096^3$                    & 1.29 & 2.53 & \phantom{0}9.49 \\
Forward MLP $XW_1^{\!\top}$ & 1.46 & 2.53 & 10.68 \\
Backward $dX = dY W$        & 1.51 & 2.54 & 10.58 \\
Backward $dW = dY^{\!\top} X$ & 0.76 & 2.15 & \phantom{0}4.69 \\
\bottomrule
\end{tabular}
\caption{GEMM throughput (TFLOP/s). The $dW$ case lags because $K = BT$ is very large
with few threadgroups; split-K would address it and is not implemented.}
\end{table}

End-to-end training throughput, all configurations on one M5 Max:

\begin{table}[h]
\centering
\begin{tabular}{lrrr}
\toprule
Configuration & Parameters & Context & Tokens/s \\
\midrule
gpt-10m  & 12.2M  & \phantom{0}512 & 38.2k \\
gpt-25m  & 29.9M  & 1024 & 22.1k \\
gpt-100m & 97.5M  & 1024 & \phantom{0}9.7k \\
gpt-200m & 205.5M & 1024 & \phantom{0}1.8k\rlap{$^\dagger$} \\
\bottomrule
\end{tabular}
\caption{$^\dagger$ measured before the matrix-tiled backward and concurrent
optimizer landed; this figure is pessimistic.}
\end{table}

\subsection{A memory bug found by scaling}

The 100M configuration exhausted GPU memory even at a micro-batch of eight. The cause
was an interaction between two individually reasonable decisions: pooled buffers
released while a command buffer is open are parked on a retire list until the next
synchronization, and the trainer synchronized once per \emph{optimizer} step. With
sixteen gradient-accumulation micro-batches, none of the sixteen sets of activations
were ever recycled, so peak memory scaled with the accumulation factor.
Synchronizing once per micro-batch fixed it. We note this because both decisions are
defensible in isolation and the failure appears only at scale.

\subsection{End-to-end training run}
\label{sec:epoch}

To confirm that the optimized kernels train a model and not merely a benchmark, we
trained the 12.2M-parameter configuration for exactly one epoch over the
19.14M-token TinyStories corpus: 584 steps of 64 sequences $\times$ 512 tokens.
Learning rate follows linear warmup over the first 10\% of the run into a cosine
decay to one tenth of peak; AdamW with decoupled weight decay 0.1 on parameters of
rank at least two, gradient clipping at global norm 1.0.

RESULTS_PLACEHOLDER

\section{Related work}

Our GEMM structure follows MLX's STEEL kernels, and our attention design follows
FlashAttention-2 as adapted to Apple GPUs by \emph{metal-flash-attention}. The
training-loop details --- decoupled weight decay with $\varepsilon$ outside the square
root, weight decay applied only to parameters of rank at least two, the gradient-clip
factor folded into the optimizer's gradient read, and the fused classifier that
writes the logit gradient in place --- follow \texttt{llm.c} and \texttt{nanoGPT}.

\section{Limitations}

All measurements come from a single M5 Max. The neural-accelerator result in
particular is generation-specific and we expect it not to transfer to the M3
generation. Mixed precision is not implemented: every kernel is \texttt{float}, so
the 4.9$\times$ accelerator path is measured but not yet exploited by the training
loop. Activation checkpointing is absent, which is what bounds model size. Split-K
for the $dW$ GEMM is not implemented. Finally, brief benchmarks on Apple Silicon can
execute in a reduced GPU performance state; long-running configurations should be
cross-checked against the performance-state trace.

\section{Conclusion}

The largest factors we encountered were not algorithmic. A twelvefold loss came from
a declaration keyword whose meaning differs between C++ and the Metal Shading
Language. A fifteenfold gain in the attention backward came from reducing per-thread
register pressure, measurable directly once shader sources were embedded in the
library. A fivefold opportunity sits behind a Metal 4 interface whose own header
documentation is stale.

The methodological lesson is narrower and firmer than any single number. Twice in
this work a confident inference from static artifacts --- the intermediate
representation in Section~\ref{sec:ir}, a published tile configuration in
Section~\ref{sec:registers} --- pointed the wrong way, and was corrected only by
measurement. On this platform the profiler is cheap, scriptable, and headless; the
intuitions are not reliable.

\section*{Availability}

\emph{Forge} comprises approximately 6{,}500 lines of C++20 and Metal. The
benchmarks reported here are \texttt{tests/bench\_matmul.cpp},
\texttt{tests/bench\_attention.cpp} and \texttt{tests/bench\_precision.cpp}; the
numerical verification of the cooperative-tensor path is \texttt{tests/mppcheck.cpp}.
Research notes with exact interface signatures and a pitfalls list are maintained in
\texttt{RESEARCH.md}.

\end{document}