Forge — LLM training from scratch in pure C++20 + Metal on Apple Silicon.
-
Fri, Jul 31, 2026 3
-
200M cluster run: one epoch at micro-batch 8
…
8 x 32 x 1024 = 262144 tokens/step (effective batch unchanged, so the 3e-4 LR still applies) x 1553 steps = one epoch over the 407,344,713-token corpus. Micro-batch 8 rather than 16 halves activation memory, which was sitting at 65 of 77.8 GB working set.
-
Add 205.6M-parameter config for the M3 Ultra cluster run
…
vocab 4096 (matches the proven BPE pipeline; the trainer's pair-count array is O(V^2) and the encoder O(V*N), which does not scale to 16384 over 400M tokens), with d_ff 3072 restoring the parameter count to 205.6M. Micro-batch 16 rather than 4 since 96 GB affords it and it cuts sync points 4x.
-
Forge: LLM training from scratch in C++20 + Metal on Apple Silicon
…
A complete transformer training stack with no ML dependencies: tensors, autograd, hand-written Metal kernels, flash attention (forward and backward), AdamW, BPE tokenizer, checkpointing and generation. Architecture is fully config-driven — the same binary trains 12M to 205M parameter models. Every Metal kernel is validated against a CPU reference (85 parity checks, <=1e-4, most bit-exact), gradients against central finite differences, and each optimization was accepted only after the training loss trajectory stayed numerically unchanged. Measured findings (M5 Max, documented in RESEARCH.md and paper/forge.tex): - `constant constexpr` for MSL tile constants declares an address-space variable, not a compile-time constant. Loops stop unrolling and every matrix accumulator spills: 0.82 -> 10.21 TFLOPS once switched to enums. - That defect is invisible in the AIR at every -O level, because unrolling happens in the driver back end. Benchmark; do not read the IR. - Register pressure, not bandwidth, dominates attention backward. Guided by measured spill counts, three restructurings took it 107 -> 7.05 ms (15.2x). - On M5, mpp::tensor_ops::matmul2d reaches 51.5 TFLOPS with f16 operands vs 10.6 for a tuned simdgroup_matrix kernel (4.9x), verified numerically. f16 on the simdgroup path alone is worth only +18-22%. - Concurrent dispatch for the optimizer sweep: +22% on the 100M config. Trained the 12.2M config for one epoch over 19.14M TinyStories tokens: loss 8.40 -> 2.99, validation 3.009, perplexity 20.27, ~38.2k tokens/sec.
-