spb/forge Public MIT

1% Author: Simon-Pierre Boucher — contact@spboucher.ai2\documentclass[11pt,a4paper]{article}34\usepackage[utf8]{inputenc}5\usepackage[T1]{fontenc}6\usepackage{lmodern}7\usepackage{amsmath,amssymb}8\usepackage{booktabs}9\usepackage{listings}10\usepackage{xcolor}11\usepackage{graphicx}12\usepackage[margin=1in]{geometry}13\usepackage{hyperref}14\hypersetup{colorlinks=true,linkcolor=blue,citecolor=blue,urlcolor=blue}15\usepackage{microtype}1617\lstdefinestyle{msl}{18  basicstyle=\ttfamily\footnotesize,19  keywordstyle=\color{blue!70!black}\bfseries,20  commentstyle=\color{green!40!black}\itshape,21  stringstyle=\color{red!60!black},22  numbers=none, breaklines=true, frame=single, framesep=4pt,23  backgroundcolor=\color{gray!5}, showstringspaces=false,24  morekeywords={kernel,threadgroup,device,constant,constexpr,simdgroup_float8x8,25    simdgroup_matrix,enum,uint,template,typename,tensor,dextents,tensor_inline,half,bfloat}26}27\lstset{style=msl}2829\title{\textbf{Forge}: Building an LLM Training Framework from Scratch\\30in C++ and Metal on Apple Silicon\\31\large Measured Lessons on Compiler Traps, Register Pressure,\\32and the M5 Neural Accelerators}3334\author{Simon-Pierre Boucher\\35\texttt{contact@spboucher.ai}}3637\date{31 July 2026}3839\begin{document}40\maketitle4142\begin{abstract}43We report engineering findings from \emph{Forge}, a transformer training framework44written from scratch in C++20 with hand-written Metal compute kernels, with no45PyTorch, MLX, or other machine-learning dependency. Every GPU kernel is validated46against a CPU reference to $\leq 10^{-4}$, and every optimization reported here was47accepted only after the training loss trajectory remained numerically unchanged.4849Four findings are, we believe, of general interest to practitioners writing Metal50compute kernels. First, the idiom \texttt{constant constexpr} for tile constants in51the Metal Shading Language silently declares an \emph{address-space variable} rather52than a compile-time constant; the resulting loss of loop unrolling spills every53matrix accumulator to the stack and cost a factor of \textbf{12$\times$} on our GEMM.54Second, this class of defect is invisible in the AIR intermediate representation at55every optimization level, because unrolling and fragment promotion occur in the56driver's back end --- practitioners must benchmark rather than read the IR. Third,57per-thread register pressure, not bandwidth, is the dominant constraint in58attention backward kernels; measured spill counts guided three successive59restructurings that took the backward pass from 107\,ms to607.05\,ms, a \textbf{15.2$\times$} improvement. Fourth, on61M5-generation hardware the Metal Performance Primitives cooperative-tensor62\texttt{matmul2d} reaches \textbf{51.5\,TFLOP/s} with \texttt{half} operands against6310.6\,TFLOP/s for a well-tuned hand-written \texttt{simdgroup\_matrix} kernel,64a \textbf{4.9$\times$} gap that inverts the usual cost/benefit assessment of mixed65precision.6667We also report a negative result that we consider equally useful: \texttt{half} and68\texttt{bfloat} operands buy only 18--22\% over \texttt{float} on the69\texttt{simdgroup\_matrix} path, so mixed precision on that path is a memory70optimization and not a compute one. The framework trains models from 12M to 205M71parameters on a single machine.72\end{abstract}7374\tableofcontents7576\section{Introduction}7778Apple Silicon offers a large unified memory pool and substantial GPU throughput, but79the software ecosystem for \emph{training} on it is thin. Inference frameworks are80mature; training-specific kernels, in particular a fused attention backward pass, are81largely absent. A survey we conducted on 31 July 2026 (Section~\ref{sec:survey})82found that of the major open-source frameworks, none ships a fused attention backward83kernel for Metal.8485This paper reports what we learned building one. \emph{Forge} is a decoder-only86transformer training framework in pure C++20 with Metal compute kernels. Its87architecture is entirely configuration-driven, and it has no machine-learning88dependency: only Apple's \texttt{metal-cpp} bindings and a single-header JSON parser.8990Our contribution is not the framework itself but the measurements. Each section below91states a hypothesis, the measurement that tested it, and the outcome --- including92the cases where the hypothesis was wrong.9394\subsection{Methodology and reproducibility}9596All measurements were taken on an Apple M5 Max (40-core GPU) running macOS 27.097(build 26A5388g) with Xcode 26.6 and Metal toolchain 32023.883. Kernels are compiled98at \texttt{-std=metal3.2} except where Metal 4 features are required.99Throughput figures are computed from100\texttt{MTLCommandBuffer.GPUStartTime}/\texttt{GPUEndTime}, which excludes CPU101encoding time.102103Two disciplines govern every result reported here.104105\paragraph{Correctness gates optimization.} Every Metal kernel has a CPU reference106implementation, and 85 parity assertions compare them (maximum absolute error107$\leq 10^{-4}$ in \texttt{float}; most are bit-exact). Gradients are additionally108checked against central finite differences. An optimization is accepted only if the109full suite still passes.110111\paragraph{Numerical equivalence, not just correctness.} For the larger kernel112rewrites we required that the \emph{training loss trajectory} be unchanged. After113replacing the scalar attention kernels with tiled matrix-multiply versions, the loss114at every logged step over 250 optimizer steps was identical115(8.3618, 7.6130, 6.6455, 5.0902, 4.9845; validation 4.9018). This is a stronger116check than parity on a single call, since it accumulates any discrepancy through the117optimizer.118119\section{System overview}120121The framework comprises roughly 6{,}500 lines of C++ and Metal. A \texttt{Tensor}122type provides shared-storage views over pooled \texttt{MTLBuffer} allocations in123\texttt{MTLStorageModeShared}, so the same memory is addressable from CPU and GPU124with no copies. A dynamic autograd tape records backward closures. Operations125dispatch to either a CPU reference backend or the Metal backend behind one interface,126which is what makes the parity testing possible.127128Execution follows the batching discipline recommended for Apple GPUs: the whole129training micro-batch is encoded into one command buffer with a small number of130long-lived compute encoders, and the GPU is synchronized only at the loss-readback131boundary.132133\section{Finding 1: \texttt{constant constexpr} is not a compile-time constant}134\label{sec:constexpr}135136\subsection{The trap}137138Tile geometry in a GEMM kernel is naturally expressed as named constants. The139apparently idiomatic C++ spelling is:140141\begin{lstlisting}142constant constexpr uint TM = 4;   // WRONG in MSL143\end{lstlisting}144145In the Metal Shading Language, however, \texttt{constant} is an \emph{address-space146qualifier}, and program-scope variables are \emph{required} to live in it. The147declaration above therefore creates a variable in the constant address space, not a148compile-time constant. Loop bounds derived from it are not known at compile time, so149the fragment loops do not unroll; \texttt{acc[i][j]} becomes dynamic indexing into an150array of opaque \texttt{simdgroup\_matrix} values; and the compiler is forced to151place all sixteen 8$\times$8 accumulators (256 bytes each) on the stack. Every152matrix-multiply-accumulate then pays roughly a kilobyte of stack traffic.153154The correct spelling uses enumerators, which are true integral constant expressions:155156\begin{lstlisting}157enum : uint { BM = 64, BN = 64, BK = 16, WM = 2, WN = 2,158              TM = BM / (8 * WM), TN = BN / (8 * WN) };159\end{lstlisting}160161together with \verb|#pragma clang loop unroll(full)| on the fragment loops. Note also162that \verb|#pragma unroll full|, which appears in Apple's own header documentation,163is not valid here.164165\subsection{Measurement}166167The change is textual and semantically neutral. Its effect is not:168169\begin{table}[h]170\centering171\begin{tabular}{lrr}172\toprule173Shape & \texttt{constant constexpr} & \texttt{enum} + unroll pragma \\174\midrule175$2048^3$          & 0.82 & 10.21 \\176$4096^3$          & 0.85 & \phantom{0}9.83 \\177$65536\times512\times1408$ & 0.83 & 10.68 \\178\bottomrule179\end{tabular}180\caption{f32 GEMM throughput (TFLOP/s), M5 Max. A 12$\times$ difference from the181declaration form of tile constants.}182\end{table}183184Before this fix the ``optimized'' \texttt{simdgroup\_matrix} kernel was185\emph{three times slower} than a naive 16$\times$16 tiled kernel --- a result that186invites the wrong conclusion, namely that the matrix instructions are not worth187using.188189\section{Finding 2: the intermediate representation cannot diagnose this}190\label{sec:ir}191192Having observed the anomaly, our first instinct was to inspect the compiler output:193194\begin{lstlisting}195xcrun metal -std=metal3.2 -O2 -S -emit-llvm -c matmul_simd.metal -o out.ll196\end{lstlisting}197198The AIR showed three \texttt{alloca} instructions for the fragment arrays and only199two matrix-multiply intrinsics where thirty-two were expected --- apparently200confirming the spill hypothesis. It also, however, showed \emph{exactly the same201thing} at \texttt{-O0}, \texttt{-O2} and \texttt{-O3}, and continued to show it after202the fix that produced the 12$\times$ speedup.203204The explanation is that loop unrolling and promotion of \texttt{simdgroup\_matrix}205values into registers happen in the driver's AIR-to-ISA back end, at pipeline206creation time, not in the front end that \texttt{-emit-llvm} exposes. The IR is207therefore uninformative for exactly the class of question one is most tempted to ask208of it.209210\paragraph{Practical rule.} On Metal, benchmark; do not read the IR. We lost time to211this and record it because the failure mode is silent: the IR looks like a212confirmation.213214\section{Finding 3: register pressure dominates attention backward}215\label{sec:registers}216217\subsection{Three restructurings, each guided by measurement}218219Our first fused attention implementation assigned one thread per query row and held220$q$, $o$ (forward) or $k$, $v$, $dk$, $dv$ (backward) in per-thread arrays. At head221dimension 64 the backward variant holds $4 \times 64 = 256$ floats, or one kilobyte222per thread.223224We tested the hypothesis that this spills, by progressively reducing what each thread225holds:226227\begin{table}[h]228\centering229\begin{tabular}{lrr}230\toprule231Attention backward, per layer & gpt-10m shapes & gpt-25m shapes \\232\midrule233$k,v,dk,dv$ all in registers        & 150 & --- \\234read-only $k,v$ from device         & 118 & 611 \\235$dK$ and $dV$ split into two kernels & 105 & 549 \\236\midrule237tiled with \texttt{simdgroup\_matrix} & \phantom{0}9.66 & \phantom{0}45.5 \\238\quad + $dK$/$dV$ split again        & \textbf{\phantom{0}7.05} & \textbf{\phantom{0}32.9} \\239\bottomrule240\end{tabular}241\caption{Time in milliseconds. gpt-10m shapes are $B{=}64$, $T{=}512$, $H{=}6$,242$d_h{=}64$; gpt-25m shapes are $B{=}64$, $T{=}1024$, $H{=}8$. Total improvement24315.2$\times$ and 17.0$\times$ respectively.}244\end{table}245246The read-only operands $k$ and $v$ are re-read from device memory on every iteration247in the second row, which sounds wasteful; because a thread reads the same address248every iteration, the L1 cache serves it, and removing the two arrays from registers249is the larger effect.250251\subsection{Confirming the mechanism directly}252253The reasoning above was indirect. Xcode~26 ships \texttt{gpudebug}, a headless254command-line GPU debugger that reports compiler statistics per kernel. It requires255shader sources to be embedded in the library, which in turn requires256\texttt{-frecord-sources} at \emph{both} the compile and link steps. (The link must257be performed by \texttt{metal}; \texttt{metallib} rejects the flag.) With that fixed,258the mechanism is directly visible:259260\begin{table}[h]261\centering262\begin{tabular}{lrrr}263\toprule264Kernel & Temp registers & Spilled bytes & Cost \\265\midrule266flash forward, scalar        & 126 & 368 & 38.9\% \\267flash forward, MMA-tiled     & \phantom{0}85 & \textbf{0} & 29.6\% \\268flash backward $dQ$, MMA     & \phantom{0}95 & \textbf{0} & \phantom{0}8.5\% \\269flash backward $dKV$, MMA (fused) & 111 & \textbf{4352} & 16.4\% \\270\bottomrule271\end{tabular}272\caption{Compiler statistics from \texttt{gpudebug}, M5 Max.}273\end{table}274275Two things follow. The matrix-tiled forward eliminates the spill entirely276($368 \rightarrow 0$ bytes), which is the mechanism behind its speedup: an2778$\times$8 fragment occupies two floats per lane, where the scalar formulation held278whole arrays of head-dimension length. And the fused $dK/dV$ backward kernel279\emph{still spilled} 4352 bytes --- a defect we had not suspected, and which made it280the most expensive backward kernel. Splitting it recovered a further 27\%.281282We regard the tooling lesson as the transferable one: register pressure is283measurable, cheaply and without a GUI, and it is worth measuring before restructuring284a kernel on intuition.285286\section{Finding 4: cooperative tensors and the M5 neural accelerators}287\label{sec:mpp}288289\subsection{Mixed precision on the classical path is a memory optimization}290291The literature disagrees about \texttt{half} on Apple GPUs. Published292microbenchmarks report that \texttt{half} and \texttt{float} fused multiply-add293execute at the same rate, the benefit arising from register and bandwidth pressure;294Apple's material for the M3 generation describes up to twice the arithmetic295throughput from co-issue. Since plumbing mixed precision through a training framework296is a substantial amount of work, we measured before committing. Holding the tiling297and the \texttt{float} accumulator fixed and varying only the operand type:298299\begin{table}[h]300\centering301\begin{tabular}{lrrr}302\toprule303Shape & \texttt{float} & \texttt{half} & \texttt{bfloat} \\304\midrule305$2048^3$ & \phantom{0}9.29 & 10.64 & 10.82 \\306$4096^3$ & \phantom{0}9.85 & 11.81 & 12.07 \\307$65536\times512\times1408$ & \phantom{0}9.90 & 12.45 & 12.44 \\308\bottomrule309\end{tabular}310\caption{\texttt{simdgroup\_matrix} throughput (TFLOP/s) by operand precision,311\texttt{float} accumulator throughout.}312\end{table}313314The gain is 18--22\%, consistent with the register/bandwidth explanation rather than315with doubled arithmetic rate. On this path, mixed precision halves activation memory316--- valuable, since activation memory bounds trainable model size --- but it is not a317compute optimization.318319\subsection{The same operands through \texttt{matmul2d}}320321\texttt{MetalPerformancePrimitives.framework} ships in the macOS 26.5 SDK and exposes322\texttt{mpp::tensor\_ops::matmul2d}, a cooperative-tensor matrix multiply that targets323the per-core neural accelerators introduced with the M5 generation. Compiling at324\texttt{-std=metal4.0}, with a 64$\times$32 tile over four SIMD groups and a325\texttt{float} accumulator:326327\begin{table}[h]328\centering329\begin{tabular}{lrrrr}330\toprule331 & \multicolumn{2}{c}{\texttt{simdgroup\_matrix}} & \multicolumn{2}{c}{\texttt{matmul2d}} \\332\cmidrule(lr){2-3}\cmidrule(lr){4-5}333Shape & \texttt{float} & \texttt{half} & \texttt{float} & \texttt{half} \\334\midrule335$2048^3$ & \phantom{0}9.3 & 10.6 & 14.9 & \textbf{51.5} \\336$4096^3$ & \phantom{0}9.9 & 11.8 & 14.6 & \textbf{44.3} \\337$65536\times512\times1408$ & \phantom{0}9.9 & 12.4 & 14.4 & \textbf{23.8} \\338\bottomrule339\end{tabular}340\caption{Throughput (TFLOP/s), M5 Max. A 1.5$\times$ gain in \texttt{float} and3414.3--4.9$\times$ in \texttt{half}.}342\end{table}343344Because a 4.9$\times$ claim invites scepticism, we verified numerically rather than345only timing: against the CPU reference the \texttt{float} path is bit-exact, the346\texttt{half} path differs by $3.8 \times 10^{-6}$ from a reference fed identically347rounded inputs, and every output element is non-zero, confirming that the full348reduction over $K$ occurs. The transposed variants required for training349($X W^{\!\top}$ in the forward, $dX = dY W$ and $dW = dY^{\!\top} X$ in the backward)350are likewise bit-exact.351352\paragraph{Generational caveat.} Published work on M4 Max hardware found353\texttt{matmul2d} still executing on the shader cores and losing to a hand-fused354GEMM. Our M5 result is the opposite. The accelerator hardware differs by generation;355this measurement should be repeated per target device rather than assumed.356357\paragraph{Consequence.} This inverts the priority we had assigned to mixed358precision. Judged on the classical path it is worth 20\% and is chiefly a memory359feature; as the entry condition for the accelerator path it is worth 4.9$\times$.360361\subsection{Interface notes}362363Several details cost us time and are not documented accurately.364365\begin{itemize}366  \item Bind ordinary buffers and construct tensors inside the kernel with the367        \texttt{tensor\_inline} descriptor. The default \texttt{tensor\_handle}368        descriptor wraps an opaque handle obtainable only from a host-side369        \texttt{MTLTensor}, which would require substantial host plumbing.370  \item Extents are ordered (columns, rows), and \texttt{slice()} takes column then371        row; a transposed operand must be sliced the other way and its extents372        swapped.373  \item The element type must be non-\texttt{const}; a \texttt{const float} tensor374        fails a static assertion.375  \item Zero the destination cooperative tensor using376        \texttt{is\_valid\_element(i)}. Apple's own header example calls377        \texttt{get\_mask(i)}, which does not exist in this SDK.378\end{itemize}379380\section{Finding 5: concurrency, not Metal 4, is the encoding win}381382Metal's default compute encoder orders every dispatch against its predecessor. For383the optimizer step this is pure loss: AdamW issues one dispatch per parameter tensor384--- of order one hundred for a 100M-parameter model --- and each gradient-norm385reduction is a single threadgroup. These are mutually independent.386387We introduced a scoped concurrent-dispatch region and applied it to both optimizer388passes. On the 100M configuration end-to-end throughput rose from 7.9k to389\textbf{9.6k tokens/s}, a 22\% improvement, with identical losses. The change is390about twenty lines and requires no Metal 4 adoption.391392An independent investigation confirmed the general shape of this result --- roughly39315$\times$ on a synthetic batch of small independent dispatches --- while also394establishing that Metal 4's command-encoding path offers essentially no CPU-side395advantage once invariant bindings are hoisted out of the dispatch loop. The396concurrency, available since macOS 10.14, is the entire effect.397398\section{Ecosystem survey}399\label{sec:survey}400401We surveyed the major open-source frameworks on 31 July 2026 for a fused attention402backward pass on Metal:403404\begin{table}[h]405\centering406\begin{tabular}{ll}407\toprule408Framework & Fused attention backward on Metal \\409\midrule410MLX & No. \texttt{use\_fallback} returns \texttt{true}; \texttt{eval\_gpu} throws411      \texttt{"NYI"}. The fused backward is CUDA-only. \\412llama.cpp & No. The Metal backend does not support413            \texttt{GGML\_OP\_FLASH\_ATTN\_BACK}. \\414PyTorch MPS & No. Hand-written Metal forward kernels only. \\415Candle & No. Vendored MLX forward only. \\416tinygrad & No. Its flash-attention backward targets AMD. \\417Burn / CubeCL & Kernels exist but are a scalar scaffold, not wired to autodiff. \\418\bottomrule419\end{tabular}420\end{table}421422MLX's Metal forward additionally declines the fused path under gradient tracing, with423the comment that unfused is faster for training on Metal. Production prior art is424limited to Philip Turner's \emph{metal-flash-attention} and one third-party package.425Both use the same three-stage, atomic-free structure we arrived at independently:426a preprocessing kernel for $D = \mathrm{rowsum}(dO \circ O)$ in \texttt{float}, a427$dQ$ kernel parallel over queries, and a $dK/dV$ kernel parallel over key/value rows.428429\section{Attention kernel design}430431Our tiled forward kernel assigns 32 query rows per threadgroup across four SIMD432groups, holding $Q$ and the output accumulator in registers as 8$\times$8 fragments433and staging $K$ and $V$ through threadgroup memory. The score tile is round-tripped434through threadgroup memory so that the softmax reductions run on ordinary threads.435436\paragraph{A scoped problem worth recording.} Online softmax requires rescaling the437output accumulator by a per-row factor at every key/value block,438$O \leftarrow \mathrm{diag}(\mathrm{corr})\, O$. But the Metal specification leaves439the mapping from \texttt{simdgroup\_matrix} elements to lanes \emph{unspecified}, so a440lane cannot determine which row its registers correspond to. MLX resolves this by441reverse-engineering the mapping. We instead construct the 8$\times$8 diagonal matrix442in threadgroup memory and apply it with an ordinary matrix multiply. This is443specification-clean, costs $d_h/8$ additional multiplies per block (about 25\% more444matrix work), and keeps the accumulator in registers --- staging it in threadgroup445memory would have cost 8\,KB and halved residency. The backward pass needs none of446this, since the saved logsumexp already fixes the normalization.447448A second technique worth noting: the $dK/dV$ kernel requires transposed score449matrices, which are obtained not by moving data but by exchanging the operand roles,450$S^{\!\top} = K Q^{\!\top}$ and $dP^{\!\top} = V\, dO^{\!\top}$, with451\texttt{simdgroup\_load(..., transpose=true)} supplying the transposed operands452directly from the staged tiles.453454\section{Results}455456\begin{table}[h]457\centering458\begin{tabular}{lrrr}459\toprule460Kernel & Naive & Tiled & \texttt{simdgroup\_matrix} \\461\midrule462$4096^3$                    & 1.29 & 2.53 & \phantom{0}9.49 \\463Forward MLP $XW_1^{\!\top}$ & 1.46 & 2.53 & 10.68 \\464Backward $dX = dY W$        & 1.51 & 2.54 & 10.58 \\465Backward $dW = dY^{\!\top} X$ & 0.76 & 2.15 & \phantom{0}4.69 \\466\bottomrule467\end{tabular}468\caption{GEMM throughput (TFLOP/s). The $dW$ case lags because $K = BT$ is very large469with few threadgroups; split-K would address it and is not implemented.}470\end{table}471472End-to-end training throughput, all configurations on one M5 Max:473474\begin{table}[h]475\centering476\begin{tabular}{lrrr}477\toprule478Configuration & Parameters & Context & Tokens/s \\479\midrule480gpt-10m  & 12.2M  & \phantom{0}512 & 38.2k \\481gpt-25m  & 29.9M  & 1024 & 22.1k \\482gpt-100m & 97.5M  & 1024 & \phantom{0}9.7k \\483gpt-200m & 205.5M & 1024 & \phantom{0}1.8k\rlap{$^\dagger$} \\484\bottomrule485\end{tabular}486\caption{$^\dagger$ measured before the matrix-tiled backward and concurrent487optimizer landed; this figure is pessimistic.}488\end{table}489490\subsection{A memory bug found by scaling}491492The 100M configuration exhausted GPU memory even at a micro-batch of eight. The cause493was an interaction between two individually reasonable decisions: pooled buffers494released while a command buffer is open are parked on a retire list until the next495synchronization, and the trainer synchronized once per \emph{optimizer} step. With496sixteen gradient-accumulation micro-batches, none of the sixteen sets of activations497were ever recycled, so peak memory scaled with the accumulation factor.498Synchronizing once per micro-batch fixed it. We note this because both decisions are499defensible in isolation and the failure appears only at scale.500501\subsection{End-to-end training run}502\label{sec:epoch}503504To confirm that the optimized kernels train a model and not merely a benchmark, we505trained the 12.2M-parameter configuration for exactly one epoch over the50619.14M-token TinyStories corpus: 584 steps of 64 sequences $\times$ 512 tokens.507Learning rate follows linear warmup over the first 10\% of the run into a cosine508decay to one tenth of peak; AdamW with decoupled weight decay 0.1 on parameters of509rank at least two, gradient clipping at global norm 1.0.510511RESULTS_PLACEHOLDER512513\section{Related work}514515Our GEMM structure follows MLX's STEEL kernels, and our attention design follows516FlashAttention-2 as adapted to Apple GPUs by \emph{metal-flash-attention}. The517training-loop details --- decoupled weight decay with $\varepsilon$ outside the square518root, weight decay applied only to parameters of rank at least two, the gradient-clip519factor folded into the optimizer's gradient read, and the fused classifier that520writes the logit gradient in place --- follow \texttt{llm.c} and \texttt{nanoGPT}.521522\section{Limitations}523524All measurements come from a single M5 Max. The neural-accelerator result in525particular is generation-specific and we expect it not to transfer to the M3526generation. Mixed precision is not implemented: every kernel is \texttt{float}, so527the 4.9$\times$ accelerator path is measured but not yet exploited by the training528loop. Activation checkpointing is absent, which is what bounds model size. Split-K529for the $dW$ GEMM is not implemented. Finally, brief benchmarks on Apple Silicon can530execute in a reduced GPU performance state; long-running configurations should be531cross-checked against the performance-state trace.532533\section{Conclusion}534535The largest factors we encountered were not algorithmic. A twelvefold loss came from536a declaration keyword whose meaning differs between C++ and the Metal Shading537Language. A fifteenfold gain in the attention backward came from reducing per-thread538register pressure, measurable directly once shader sources were embedded in the539library. A fivefold opportunity sits behind a Metal 4 interface whose own header540documentation is stale.541542The methodological lesson is narrower and firmer than any single number. Twice in543this work a confident inference from static artifacts --- the intermediate544representation in Section~\ref{sec:ir}, a published tile configuration in545Section~\ref{sec:registers} --- pointed the wrong way, and was corrected only by546measurement. On this platform the profiler is cheap, scriptable, and headless; the547intuitions are not reliable.548549\section*{Availability}550551\emph{Forge} comprises approximately 6{,}500 lines of C++20 and Metal. The552benchmarks reported here are \texttt{tests/bench\_matmul.cpp},553\texttt{tests/bench\_attention.cpp} and \texttt{tests/bench\_precision.cpp}; the554numerical verification of the cooperative-tensor path is \texttt{tests/mppcheck.cpp}.555Research notes with exact interface signatures and a pitfalls list are maintained in556\texttt{RESEARCH.md}.557558\end{document}559