spb/forge Public MIT

Wave 1 (module-level):
- rope refactor: kernels read a host-precomputed inv-freq table; unlocks
  HF-"llama3" rope scaling (rope_scale_*), per-layer theta, and NoPE
  layers (nope_every, SmolLM3)
- attention_bias (Qwen2.5 QKV bias), head_dim decoupled from
  d_model/n_heads (Qwen3), relu2 activation (nanoGPT-speedrun lineage),
  norm_placement pre|post|sandwich (OLMo2/Gemma)

Wave 2 (attention kernels):
- sliding_window + sliding_global_every (Mistral / Gemma3 local:global
  patterns) in the CPU reference, the unfused Metal kernels, and the
  fused scalar flash kernels — out-of-window KV blocks are skipped, so
  cost scales with the window; window > 0 auto-routes off the MMA kernel
- attn_softcap (Gemma2): cap*tanh on scores pre-softmax, unfused path,
  exact tanh' chain in all backwards
- rope_theta_global for dual-theta local/global layers (Gemma3)

Parity suites cover every knob (fused + unfused paths); gradcheck and
overfit stay green. New demo configs: gpt-50m-mistral, gpt-50m-gemma;
ARCHITECTURES.md documents the per-family config matrix.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>

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.