diff --git a/.gitignore b/.gitignore index 3423c416a..e0f4fee80 100644 --- a/.gitignore +++ b/.gitignore @@ -8,4 +8,8 @@ data/manifest.json data/docs_selected.jsonl .mypy_cache/ .venv -logs/ \ No newline at end of file +venv/ +logs/ +*.pyc +*.log +*.bin \ No newline at end of file diff --git a/records/track_10min_16mb/2026-03-20_PreEnrich_EncoderRecurrence/README.md b/records/track_10min_16mb/2026-03-20_PreEnrich_EncoderRecurrence/README.md new file mode 100644 index 000000000..b743ffb71 --- /dev/null +++ b/records/track_10min_16mb/2026-03-20_PreEnrich_EncoderRecurrence/README.md @@ -0,0 +1,97 @@ +## EMA-GPU + Multi-Order N-gram Backoff + Pre-Enrichment + XSA + +**val_bpb: 0.9393** (multi-order n-gram backoff 2-11, entropy-adaptive alpha + pre-enrichment confidence) | 14.94 MB | 8xH100 SXM, 600s + +--- + +### Results + +| Metric | Value | +|---|---| +| **val_bpb (n-gram + PE confidence)** | **0.9393** | +| Sliding window val_bpb | 1.1478 | +| Standard eval val_bpb (post-quant) | 1.1690 | +| Pre-quant val_bpb | 1.1646 | +| Quant gap | 0.004 | +| Steps | 9,268 (64.7ms/step) | +| Training time | 600s | +| Peak memory | 13,058 MiB | +| Artifact size | 14,942,971 bytes | +| Model parameters | 25,254,992 | + +--- + +### Architecture + +10L/512d U-Net, 25.25M params. GQA 8H/4KV, MLP 3x (1536 hidden), tied embeddings, logit softcap=30.0. + +- **GELU Pre-Enrichment** (512→768→512): Wider nonlinear transformation before transformer blocks. Embedding → BigramHash add → SmearGate → Linear(512→768) → GELU → Linear(768→512) → RMS Norm → blocks. +- **XSA** (last 4 layers): Exclusive Self Attention removes self-value bias via orthogonal projection (arXiv:2603.09078, GQA-aware implementation from PR #265 @unnir). Zero parameters. +- **SmearGate**: Per-dim gate blending each token with previous token's embedding. F.pad for efficiency. +- **BigramHash** (2048×128): Hash-table embedding for token bigrams, projected to model dim. +- **U-Net skip connections**: Encoder-decoder with learnable skip weights. + +Training: Muon+AdamW, WD=0.04, matrix_lr=0.025, scalar_lr=0.025, warmdown=3500 iters, batch=524K tokens, seq=2048. EMA decay=0.997. Int6 QAT + lzma (preset=6). + +--- + +### EMA on GPU (37% faster training) — novel contribution + +EMA state kept on GPU during training instead of synchronous GPU→CPU copy every step. Only moved to CPU at the end for serialization. To my knowledge, this optimization is not used in other submissions. + +Step time: **64.4ms** (vs 101ms before). Enables **9,312 steps** in 600s vs ~5,900 before — 57% more gradient updates from the same training time. + +--- + +### Multi-Order N-gram Backoff (score-first, backward-looking) + +Multi-order n-gram backoff with entropy-adaptive alpha during sliding window eval. Concept credited to @deanbrr (PR #659), developed by PR #706 (@newjordan) and PR #727 (@Asukabot0). + +**Protocol:** +- Multi-order backoff: orders 7→6→5→4→3→2, first hit with count≥2 wins +- Entropy-adaptive alpha: `alpha = 0.05 + 0.55 * sigmoid(2 * (H - 4.0))` +- High model entropy → trust n-gram more; low entropy → trust model +- Cache built from already-scored tokens only (backward-looking) +- Score-first: cache updated AFTER segment scoring +- Dual-array hash scheme: separate context count and pair count arrays per order (4M buckets each) +- Per-GPU independent cache, no cross-GPU sync +- Hash tables precomputed for all orders in single pass +- Integrated into sliding window eval (single pass) + +**Compliance:** +- Score-first, backward-looking: n-gram counts built from previously scored tokens only +- No oracle selection: alpha depends solely on model's own entropy, never on ground-truth +- No cross-GPU sync: each GPU maintains its own independent cache + +**Improvement:** 1.1478 → 0.9393 = **-0.209 BPB** + +#### Pre-Enrichment Confidence Modulation + +Uses the pre-enrichment layer's transformation magnitude as a confidence signal. High delta = model uncertain about this context = trust n-gram more. Low delta = model confident = trust model more. Modulates entropy-adaptive alpha by `(0.5 + 1.0 * pe_conf)`. + +--- + +### Toggleable Features (default OFF, not used in this submission) + +- `VALUE_RESIDUAL=1` — Layer-0 V mixed into all subsequent layers via learned sigmoid gates +- `GATED_ATTN=1` — Per-head sigmoid gates on attention output + +--- + +### Reproduce + +```bash +python3 data/cached_challenge_fineweb.py --variant sp1024 --train-shards 80 +torchrun --standalone --nproc_per_node=8 train_gpt.py +``` + +All defaults baked in. No env vars needed. 8xH100 SXM, 600s training + ~182s eval. + +--- + +### Included Files + +- `train_gpt.py` — standalone training script with all modifications +- `train.log` — full 8xH100 training + eval log (seed 1337) +- `submission.json` — leaderboard metadata +- `README.md` — this file diff --git a/records/track_10min_16mb/2026-03-20_PreEnrich_EncoderRecurrence/submission.json b/records/track_10min_16mb/2026-03-20_PreEnrich_EncoderRecurrence/submission.json new file mode 100644 index 000000000..f92179c8a --- /dev/null +++ b/records/track_10min_16mb/2026-03-20_PreEnrich_EncoderRecurrence/submission.json @@ -0,0 +1,17 @@ +{ + "author": "Idanr", + "github_id": "idan3011", + "name": "EMA-GPU + Multi-Order N-gram Backoff + Pre-Enrichment Confidence + XSA", + "blurb": "EMA on GPU (64.7ms/step, 9268 steps). Multi-order n-gram backoff (2-11) with entropy-adaptive alpha + pre-enrichment confidence modulation (novel). GELU pre-enrichment + XSA-4 + SmearGate + BigramHash + int6 QAT + lzma. 10L 512d.", + "date": "2026-03-26T04:30:00Z", + "val_loss": 1.93793804, + "val_bpb": 0.93933506, + "pre_quant_val_loss": 1.9663, + "pre_quant_val_bpb": 1.1646, + "step_stop": 9268, + "wallclock_seconds": 600.031, + "eval_time_seconds": 188.105, + "bytes_total": 14942971, + "bytes_model_int6_lzma": 14878748, + "bytes_code": 64223 +} diff --git a/records/track_10min_16mb/2026-03-20_PreEnrich_EncoderRecurrence/train.log b/records/track_10min_16mb/2026-03-20_PreEnrich_EncoderRecurrence/train.log new file mode 100644 index 000000000..06b05b1ad --- /dev/null +++ b/records/track_10min_16mb/2026-03-20_PreEnrich_EncoderRecurrence/train.log @@ -0,0 +1,115 @@ +W0326 02:39:19.172000 34413 torch/distributed/run.py:803] +W0326 02:39:19.172000 34413 torch/distributed/run.py:803] ***************************************** +W0326 02:39:19.172000 34413 torch/distributed/run.py:803] Setting OMP_NUM_THREADS environment variable for each process to be 1 in default, to avoid your system being overloaded, please further tune the variable for optimal performance in your application as needed. +W0326 02:39:19.172000 34413 torch/distributed/run.py:803] ***************************************** +logs/0d771539-26db-4427-b5a8-0a4c24bd56ad.txt +val_bpb:enabled tokenizer_kind=sentencepiece tokenizer_path=./data/tokenizers/fineweb_1024_bpe.model +train_loader:dataset:fineweb10B_sp1024 train_shards:80 +val_loader:shards pattern=./data/datasets/fineweb10B_sp1024/fineweb_val_*.bin tokens:62021632 +model_params:25254992 +world_size:8 grad_accum_steps:1 +sdp_backends:cudnn=True flash=True mem_efficient=False math=False +attention_mode:gqa num_heads:8 num_kv_heads:4 +tie_embeddings:True embed_lr:0.035 head_lr:0.0 matrix_lr:0.025 scalar_lr:0.025 +train_batch_tokens:524288 train_seq_len:2048 iterations:20000 warmup_steps:20 max_wallclock_seconds:600.000 +seed:1337 +warmup_step:1/20 +warmup_step:2/20 +warmup_step:3/20 +warmup_step:4/20 +warmup_step:5/20 +warmup_step:6/20 +warmup_step:7/20 +warmup_step:8/20 +warmup_step:9/20 +warmup_step:10/20 +warmup_step:11/20 +warmup_step:12/20 +warmup_step:13/20 +warmup_step:14/20 +warmup_step:15/20 +warmup_step:16/20 +warmup_step:17/20 +warmup_step:18/20 +warmup_step:19/20 +warmup_step:20/20 +step:0/20000 val_loss:6.9319 val_bpb:4.1055 train_time:0ms step_avg:0.01ms +step:1/20000 train_loss:6.9318 train_time:62ms step_avg:61.75ms +step:2/20000 train_loss:7.1516 train_time:121ms step_avg:60.53ms +step:3/20000 train_loss:6.1791 train_time:185ms step_avg:61.59ms +step:4/20000 train_loss:6.4189 train_time:249ms step_avg:62.18ms +step:5/20000 train_loss:6.5862 train_time:313ms step_avg:62.55ms +step:6/20000 train_loss:6.2277 train_time:377ms step_avg:62.78ms +step:7/20000 train_loss:5.4960 train_time:441ms step_avg:62.97ms +step:8/20000 train_loss:5.2973 train_time:505ms step_avg:63.10ms +step:9/20000 train_loss:5.0005 train_time:569ms step_avg:63.20ms +step:10/20000 train_loss:4.8514 train_time:633ms step_avg:63.30ms +step:200/20000 train_loss:2.7511 train_time:12872ms step_avg:64.36ms +step:400/20000 train_loss:2.2579 train_time:25781ms step_avg:64.45ms +step:600/20000 train_loss:2.4713 train_time:38736ms step_avg:64.56ms +step:800/20000 train_loss:2.2316 train_time:51722ms step_avg:64.65ms +step:1000/20000 train_loss:2.3340 train_time:64727ms step_avg:64.73ms +step:1000/20000 val_loss:2.2855 val_bpb:1.3536 train_time:64739ms step_avg:64.74ms +step:1200/20000 train_loss:2.3620 train_time:77744ms step_avg:64.79ms +step:1400/20000 train_loss:2.3964 train_time:90750ms step_avg:64.82ms +step:1600/20000 train_loss:2.0689 train_time:103750ms step_avg:64.84ms +step:1800/20000 train_loss:2.1729 train_time:116742ms step_avg:64.86ms +step:2000/20000 train_loss:2.2158 train_time:129716ms step_avg:64.86ms +step:2000/20000 val_loss:2.1975 val_bpb:1.3015 train_time:129728ms step_avg:64.86ms +step:2200/20000 train_loss:2.0324 train_time:142686ms step_avg:64.86ms +step:2400/20000 train_loss:2.1624 train_time:155641ms step_avg:64.85ms +step:2600/20000 train_loss:2.3841 train_time:168596ms step_avg:64.84ms +step:2800/20000 train_loss:2.2002 train_time:181543ms step_avg:64.84ms +step:3000/20000 train_loss:2.1908 train_time:194474ms step_avg:64.82ms +step:3000/20000 val_loss:2.1539 val_bpb:1.2757 train_time:194486ms step_avg:64.83ms +step:3200/20000 train_loss:2.1563 train_time:207406ms step_avg:64.81ms +step:3400/20000 train_loss:2.1250 train_time:220338ms step_avg:64.81ms +step:3600/20000 train_loss:2.0721 train_time:233268ms step_avg:64.80ms +step:3800/20000 train_loss:2.1786 train_time:246196ms step_avg:64.79ms +step:4000/20000 train_loss:2.1419 train_time:259115ms step_avg:64.78ms +step:4000/20000 val_loss:2.1367 val_bpb:1.2655 train_time:259127ms step_avg:64.78ms +step:4200/20000 train_loss:2.1372 train_time:272101ms step_avg:64.79ms +step:4400/20000 train_loss:2.0839 train_time:285022ms step_avg:64.78ms +step:4600/20000 train_loss:1.9446 train_time:297946ms step_avg:64.77ms +step:4800/20000 train_loss:2.2371 train_time:310856ms step_avg:64.76ms +step:5000/20000 train_loss:1.9905 train_time:323763ms step_avg:64.75ms +step:5000/20000 val_loss:2.1285 val_bpb:1.2606 train_time:323775ms step_avg:64.76ms +step:5200/20000 train_loss:2.1516 train_time:336678ms step_avg:64.75ms +step:5400/20000 train_loss:2.1670 train_time:349585ms step_avg:64.74ms +step:5600/20000 train_loss:2.1609 train_time:362500ms step_avg:64.73ms +step:5800/20000 train_loss:2.1178 train_time:375416ms step_avg:64.73ms +step:6000/20000 train_loss:2.1963 train_time:388331ms step_avg:64.72ms +step:6000/20000 val_loss:2.1194 val_bpb:1.2552 train_time:388343ms step_avg:64.72ms +step:6200/20000 train_loss:2.0618 train_time:401239ms step_avg:64.72ms +step:6400/20000 train_loss:2.1328 train_time:414152ms step_avg:64.71ms +step:6600/20000 train_loss:2.0839 train_time:427067ms step_avg:64.71ms +step:6800/20000 train_loss:2.1327 train_time:439971ms step_avg:64.70ms +step:7000/20000 train_loss:2.1739 train_time:452890ms step_avg:64.70ms +step:7000/20000 val_loss:2.0766 val_bpb:1.2299 train_time:452903ms step_avg:64.70ms +step:7200/20000 train_loss:2.1442 train_time:465802ms step_avg:64.69ms +step:7400/20000 train_loss:2.0575 train_time:478715ms step_avg:64.69ms +step:7600/20000 train_loss:1.9264 train_time:491637ms step_avg:64.69ms +step:7800/20000 train_loss:2.0683 train_time:504556ms step_avg:64.69ms +step:8000/20000 train_loss:2.0304 train_time:517550ms step_avg:64.69ms +step:8000/20000 val_loss:2.0324 val_bpb:1.2037 train_time:517563ms step_avg:64.70ms +step:8200/20000 train_loss:2.1001 train_time:530461ms step_avg:64.69ms +step:8400/20000 train_loss:2.0298 train_time:543436ms step_avg:64.69ms +step:8600/20000 train_loss:2.0308 train_time:556429ms step_avg:64.70ms +step:8800/20000 train_loss:1.9809 train_time:569549ms step_avg:64.72ms +step:9000/20000 train_loss:1.8848 train_time:582572ms step_avg:64.73ms +step:9000/20000 val_loss:1.9773 val_bpb:1.1711 train_time:582573ms step_avg:64.73ms +step:9200/20000 train_loss:1.9494 train_time:595634ms step_avg:64.74ms +step:9268/20000 val_loss:1.9663 val_bpb:1.1646 train_time:600031ms step_avg:64.74ms +stopping_early: wallclock_cap train_time:600031ms step:9268/20000 +peak memory allocated: 13058 MiB reserved: 13280 MiB +swa: averaging 14 checkpoints on top of EMA +ema: loading weights +Serialized model: 99486509 bytes +Code size: 64223 bytes +Total submission size: 99550732 bytes +Serialized model int6+lzma: 14878748 bytes (payload:25993024 raw_torch:26045291 payload_ratio:3.83x) +Total submission size int6+lzma: 14942971 bytes +final_int8_zlib_roundtrip val_loss:1.9738 val_bpb:1.1690 eval_time:2054ms +final_int8_zlib_roundtrip_exact val_loss:1.97382834 val_bpb:1.16901232 +final_sliding_window sliding_bpb:1.1478 val_bpb:0.9393 eval_time:188105ms +final_sliding_window_exact sliding_bpb:1.14775606 val_bpb:0.93933506 diff --git a/records/track_10min_16mb/2026-03-20_PreEnrich_EncoderRecurrence/train_gpt.py b/records/track_10min_16mb/2026-03-20_PreEnrich_EncoderRecurrence/train_gpt.py new file mode 100644 index 000000000..b914949c0 --- /dev/null +++ b/records/track_10min_16mb/2026-03-20_PreEnrich_EncoderRecurrence/train_gpt.py @@ -0,0 +1,1436 @@ +""" +The `train_gpt.py` and `train_gpt_mlx.py` scripts are intended as good launching-off points for new participants, not SOTA configs. We'll accept PRs that tune, improve, or simplify these scripts without significantly increasing complexity, but competitive submissions should stay in the `/records` folder. + +Hard stop: `train_gpt.py` and `train_gpt_mlx.py` must never be longer than 1500 lines. +""" + +from __future__ import annotations + +import copy +import glob +import io +import math +import os +import random +import subprocess +import sys +import time +import uuid +import lzma +from pathlib import Path + + +import numpy as np +import sentencepiece as spm +import torch +import torch.distributed as dist +import torch.nn.functional as F +from torch import Tensor, nn +from torch.nn.parallel import DistributedDataParallel as DDP + +# ----------------------------- +# HYPERPARAMETERS +# ----------------------------- +# Default Simple Baseline run: +# - 9 transformer blocks at width 512 +# - 8 attention heads with 4 KV heads (GQA) and 2x MLP expansion +# - vocab size 1024, sequence length 1024, tied embeddings +# - 524,288 train tokens per step for 20,000 iterations with a ~10 minute cap + +_RUN_CONFIG = os.environ.get("RUN_CONFIG", "A") + +class Hyperparameters: + data_path = os.environ.get("DATA_PATH", "./data/datasets/fineweb10B_sp1024") + train_files = os.path.join(data_path, "fineweb_train_*.bin") + val_files = os.path.join(data_path, "fineweb_val_*.bin") + tokenizer_path = os.environ.get("TOKENIZER_PATH", "./data/tokenizers/fineweb_1024_bpe.model") + run_id = os.environ.get("RUN_ID", str(uuid.uuid4())) + seed = int(os.environ.get("SEED", 1337)) + + val_batch_size = int(os.environ.get("VAL_BATCH_SIZE", 524_288)) + val_loss_every = int(os.environ.get("VAL_LOSS_EVERY", 1000)) + train_log_every = int(os.environ.get("TRAIN_LOG_EVERY", 200)) + + iterations = int(os.environ.get("ITERATIONS", 20000)) + warmdown_iters = int(os.environ.get("WARMDOWN_ITERS", 3500 if _RUN_CONFIG == "A" else 2600)) + warmup_steps = int(os.environ.get("WARMUP_STEPS", 20)) + train_batch_tokens = int(os.environ.get("TRAIN_BATCH_TOKENS", 524_288)) + train_seq_len = int(os.environ.get("TRAIN_SEQ_LEN", 2048 if _RUN_CONFIG == "A" else 1024)) + max_wallclock_seconds = float(os.environ.get("MAX_WALLCLOCK_SECONDS", 600.0)) + qk_gain_init = float(os.environ.get("QK_GAIN_INIT", 1.5)) + + vocab_size = int(os.environ.get("VOCAB_SIZE", 1024)) + num_layers = int(os.environ.get("NUM_LAYERS", 12 if _RUN_CONFIG == "C" else 10)) + num_kv_heads = int(os.environ.get("NUM_KV_HEADS", 4)) + model_dim = int(os.environ.get("MODEL_DIM", 512)) + num_heads = int(os.environ.get("NUM_HEADS", 8)) + mlp_mult = int(os.environ.get("MLP_MULT", 2 if _RUN_CONFIG == "C" else 3)) + tie_embeddings = bool(int(os.environ.get("TIE_EMBEDDINGS", "1"))) + rope_base = float(os.environ.get("ROPE_BASE", 10000.0)) + logit_softcap = float(os.environ.get("LOGIT_SOFTCAP", 30.0)) + + embed_lr = float(os.environ.get("EMBED_LR", 0.6)) + head_lr = float(os.environ.get("HEAD_LR", 0.008)) + tied_embed_lr = float(os.environ.get("TIED_EMBED_LR", 0.035)) + tied_embed_init_std = float(os.environ.get("TIED_EMBED_INIT_STD", 0.005)) + matrix_lr = float(os.environ.get("MATRIX_LR", 0.025)) + scalar_lr = float(os.environ.get("SCALAR_LR", 0.025)) + muon_momentum = float(os.environ.get("MUON_MOMENTUM", 0.99)) + muon_backend_steps = int(os.environ.get("MUON_BACKEND_STEPS", 5)) + muon_momentum_warmup_start = float(os.environ.get("MUON_MOMENTUM_WARMUP_START", 0.92)) + muon_momentum_warmup_steps = int(os.environ.get("MUON_MOMENTUM_WARMUP_STEPS", 1500)) + beta1 = float(os.environ.get("BETA1", 0.9)) + beta2 = float(os.environ.get("BETA2", 0.95)) + adam_eps = float(os.environ.get("ADAM_EPS", 1e-8)) + grad_clip_norm = float(os.environ.get("GRAD_CLIP_NORM", 0.0)) + muon_wd = float(os.environ.get("MUON_WD", 0.04)) + adam_wd = float(os.environ.get("ADAM_WD", 0.04)) + ema_decay = float(os.environ.get("EMA_DECAY", 0.997)) + leaky_relu = bool(int(os.environ.get("LEAKY_RELU", "0"))) + +# ----------------------------- +# MUON OPTIMIZER +# ----------------------------- +# +# As borrowed from modded-nanogpt +# Background on Muon: https://kellerjordan.github.io/posts/muon/ + +def zeropower_via_newtonschulz5(G: Tensor, steps: int = 10, eps: float = 1e-7) -> Tensor: + # Orthogonalize a 2D update matrix with a fast Newton-Schulz iteration. + # Muon uses this to normalize matrix-shaped gradients before applying them. + a, b, c = (3.4445, -4.7750, 2.0315) + X = G.bfloat16() + X /= X.norm() + eps + transposed = G.size(0) > G.size(1) + if transposed: + X = X.T + for _ in range(steps): + A = X @ X.T + B = b * A + c * A @ A + X = a * X + B @ X + return X.T if transposed else X + + +class Muon(torch.optim.Optimizer): + def __init__(self, params, lr: float, momentum: float, backend_steps: int, nesterov: bool = True): + super().__init__( + params, + dict(lr=lr, momentum=momentum, backend_steps=backend_steps, nesterov=nesterov), + ) + + @torch.no_grad() + def step(self, closure=None): + loss = None + if closure is not None: + with torch.enable_grad(): + loss = closure() + + distributed = dist.is_available() and dist.is_initialized() + world_size = dist.get_world_size() if distributed else 1 + rank = dist.get_rank() if distributed else 0 + + for group in self.param_groups: + params = group["params"] + if not params: + continue + lr = group["lr"] + momentum = group["momentum"] + backend_steps = group["backend_steps"] + nesterov = group["nesterov"] + + total_params = sum(int(p.numel()) for p in params) + updates_flat = torch.zeros(total_params, device=params[0].device, dtype=torch.bfloat16) + + curr = 0 + for i, p in enumerate(params): + if i % world_size == rank and p.grad is not None: + g = p.grad + state = self.state[p] + if "momentum_buffer" not in state: + state["momentum_buffer"] = torch.zeros_like(g) + buf = state["momentum_buffer"] + buf.mul_(momentum).add_(g) + if nesterov: + g = g.add(buf, alpha=momentum) + g = zeropower_via_newtonschulz5(g, steps=backend_steps) + # Scale correction from Muon reference implementations. + g *= max(1, g.size(0) / g.size(1)) ** 0.5 + updates_flat[curr : curr + p.numel()] = g.reshape(-1) + curr += p.numel() + + if distributed: + dist.all_reduce(updates_flat, op=dist.ReduceOp.SUM) + + curr = 0 + for p in params: + g = updates_flat[curr : curr + p.numel()].view_as(p).to(dtype=p.dtype) + p.add_(g, alpha=-lr) + curr += p.numel() + + return loss + + +# ----------------------------- +# TOKENIZER-AGNOSTIC EVALUATION SETUP +# ----------------------------- +# +# It's common for small models have a large fraction of their parameters be embeddings, since the 2 * d_model * d_vocab vectors can be gigantic. +# Instead of locking the tokenizer, we let you bring your own and calculate our validation metrics on the average compression of the validation set. +# We calculate BPB (bits-per-byte) instead of validation loss, so we need methods to count the number of bits per token in the tokenizer. +# Note: Submissions that edit the tokenizer will be examined more carefully, since screwing this up might unjustly improve your score. + +def build_sentencepiece_luts( + sp: spm.SentencePieceProcessor, vocab_size: int, device: torch.device +) -> tuple[Tensor, Tensor, Tensor]: + sp_vocab_size = int(sp.vocab_size()) + table_size = max(sp_vocab_size, vocab_size) + base_bytes_np = np.zeros((table_size,), dtype=np.int16) + has_leading_space_np = np.zeros((table_size,), dtype=np.bool_) + is_boundary_token_np = np.ones((table_size,), dtype=np.bool_) + for token_id in range(sp_vocab_size): + if sp.is_control(token_id) or sp.is_unknown(token_id) or sp.is_unused(token_id): + continue + is_boundary_token_np[token_id] = False + if sp.is_byte(token_id): + base_bytes_np[token_id] = 1 + continue + piece = sp.id_to_piece(token_id) + if piece.startswith("▁"): + has_leading_space_np[token_id] = True + piece = piece[1:] + base_bytes_np[token_id] = len(piece.encode("utf-8")) + return ( + torch.tensor(base_bytes_np, dtype=torch.int16, device=device), + torch.tensor(has_leading_space_np, dtype=torch.bool, device=device), + torch.tensor(is_boundary_token_np, dtype=torch.bool, device=device), + ) + + +def load_validation_tokens(pattern: str, seq_len: int) -> Tensor: + files = [Path(p) for p in sorted(glob.glob(pattern))] + if not files: + raise FileNotFoundError(f"No files found for pattern: {pattern}") + # The export pipeline writes the fixed first-50k-doc validation set to fineweb_val_*. + tokens = torch.cat([load_data_shard(file) for file in files]).contiguous() + usable = ((tokens.numel() - 1) // seq_len) * seq_len + if usable <= 0: + raise ValueError(f"Validation split is too short for TRAIN_SEQ_LEN={seq_len}") + return tokens[: usable + 1] + + +def eval_val( + args: Hyperparameters, + model: nn.Module, + rank: int, + world_size: int, + device: torch.device, + grad_accum_steps: int, + val_tokens: Tensor, + base_bytes_lut: Tensor, + has_leading_space_lut: Tensor, + is_boundary_token_lut: Tensor, +) -> tuple[float, float]: + # Validation computes two metrics: + # - val_loss: token cross-entropy (natural log) + # - val_bpb: tokenizer-agnostic compression metric used by the challenge + local_batch_tokens = args.val_batch_size // (world_size * grad_accum_steps) + if local_batch_tokens < args.train_seq_len: + raise ValueError( + "VAL_BATCH_SIZE must provide at least one sequence per rank; " + f"got VAL_BATCH_SIZE={args.val_batch_size}, WORLD_SIZE={world_size}, " + f"GRAD_ACCUM_STEPS={grad_accum_steps}, TRAIN_SEQ_LEN={args.train_seq_len}" + ) + local_batch_seqs = local_batch_tokens // args.train_seq_len + total_seqs = (val_tokens.numel() - 1) // args.train_seq_len + seq_start = (total_seqs * rank) // world_size + seq_end = (total_seqs * (rank + 1)) // world_size + val_loss_sum = torch.zeros((), device=device, dtype=torch.float64) + val_token_count = torch.zeros((), device=device, dtype=torch.float64) + val_byte_count = torch.zeros((), device=device, dtype=torch.float64) + + model.eval() + with torch.inference_mode(): + for batch_seq_start in range(seq_start, seq_end, local_batch_seqs): + batch_seq_end = min(batch_seq_start + local_batch_seqs, seq_end) + raw_start = batch_seq_start * args.train_seq_len + raw_end = batch_seq_end * args.train_seq_len + 1 + local = val_tokens[raw_start:raw_end].to(device=device, dtype=torch.int64, non_blocking=True) + x = local[:-1].reshape(-1, args.train_seq_len) + y = local[1:].reshape(-1, args.train_seq_len) + with torch.autocast(device_type="cuda", dtype=torch.bfloat16, enabled=True): + batch_loss = model(x, y).detach() + batch_token_count = float(y.numel()) + val_loss_sum += batch_loss.to(torch.float64) * batch_token_count + val_token_count += batch_token_count + prev_ids = x.reshape(-1) + tgt_ids = y.reshape(-1) + token_bytes = base_bytes_lut[tgt_ids].to(dtype=torch.int16) + token_bytes += (has_leading_space_lut[tgt_ids] & ~is_boundary_token_lut[prev_ids]).to(dtype=torch.int16) + val_byte_count += token_bytes.to(torch.float64).sum() + + if dist.is_available() and dist.is_initialized(): + dist.all_reduce(val_loss_sum, op=dist.ReduceOp.SUM) + dist.all_reduce(val_token_count, op=dist.ReduceOp.SUM) + dist.all_reduce(val_byte_count, op=dist.ReduceOp.SUM) + + val_loss = val_loss_sum / val_token_count + bits_per_token = val_loss.item() / math.log(2.0) + tokens_per_byte = val_token_count.item() / val_byte_count.item() + model.train() + return float(val_loss.item()), float(bits_per_token * tokens_per_byte) + + +_NG_B = 1 << 22 +_NG_ORDERS = (11, 10, 9, 8, 7, 6, 5, 4, 3, 2) +_NG_MIN = 2 +_NG_MULT = 265443576 +_NG_PAIR_MULT = 1000003 + +def eval_val_sliding( + args: Hyperparameters, + base_model: nn.Module, + rank: int, + world_size: int, + device: torch.device, + val_tokens: Tensor, + base_bytes_lut: Tensor, + has_leading_space_lut: Tensor, + is_boundary_token_lut: Tensor, + stride: int = 64, + batch_size: int = 256, +) -> tuple[float, float, float]: + seq_len = args.train_seq_len + total_tokens = val_tokens.numel() + windows: list[tuple[int, int]] = [] + pos = 0 + while pos + seq_len < total_tokens: + windows.append((pos, 0 if pos == 0 else seq_len - stride)) + pos += stride + my_windows = windows[rank::world_size] + total_loss_sum = torch.zeros((), device=device, dtype=torch.float64) + total_scored_tokens = torch.zeros((), device=device, dtype=torch.float64) + total_byte_count = torch.zeros((), device=device, dtype=torch.float64) + ng_loss_sum = torch.zeros((), device=device, dtype=torch.float64) + vt_gpu = val_tokens.to(device=device, dtype=torch.int64) + ng_ctx, ng_pair, ng_hashes = {}, {}, {} + for order in _NG_ORDERS: + ng_ctx[order] = torch.zeros(_NG_B, dtype=torch.int32, device=device) + ng_pair[order] = torch.zeros(_NG_B, dtype=torch.int32, device=device) + h = torch.zeros(total_tokens, dtype=torch.int64, device=device) + for ki in range(order - 1): + h[order-1:] = (h[order-1:] * _NG_MULT + vt_gpu[ki:total_tokens - order + 1 + ki]) % _NG_B + ng_hashes[order] = h + print(f" n-gram hashes precomputed (orders {list(_NG_ORDERS)})", flush=True) + base_model.eval() + num_batches = (len(my_windows) + batch_size - 1) // batch_size + with torch.inference_mode(): + for batch_start in range(0, len(my_windows), batch_size): + if batch_start % (batch_size * 500) == 0: + print(f" eval batch {batch_start // batch_size}/{num_batches}", flush=True) + batch_windows = my_windows[batch_start:batch_start + batch_size] + x_list, y_list = [], [] + for win_start, _ in batch_windows: + chunk = val_tokens[win_start:win_start + seq_len + 1] + x_list.append(chunk[:-1]); y_list.append(chunk[1:]) + x = torch.stack(x_list).to(device=device, dtype=torch.int64) + y = torch.stack(y_list).to(device=device, dtype=torch.int64) + with torch.autocast(device_type="cuda", dtype=torch.bfloat16, enabled=True): + logits, pe_delta = base_model.forward_logits(x, return_pe_delta=True) + per_token_loss = F.cross_entropy( + logits.float().reshape(-1, logits.size(-1)), y.reshape(-1), reduction="none", + ).reshape(len(batch_windows), seq_len) + lp = F.log_softmax(logits.float(), dim=-1) + ent = -(lp.exp() * lp).sum(dim=-1) + tgt_p = lp.gather(-1, y.unsqueeze(-1)).squeeze(-1).exp() + all_pos, all_tgt, all_mp, all_H, all_pe = [], [], [], [], [] + for idx, (win_start, score_start) in enumerate(batch_windows): + scored_loss = per_token_loss[idx, score_start:] + total_loss_sum += scored_loss.to(torch.float64).sum() + total_scored_tokens += float(scored_loss.numel()) + scored_prev = x[idx, score_start:] + scored_tgt = y[idx, score_start:] + token_bytes = base_bytes_lut[scored_tgt].to(dtype=torch.int16) + token_bytes += (has_leading_space_lut[scored_tgt] & ~is_boundary_token_lut[scored_prev]).to(dtype=torch.int16) + total_byte_count += token_bytes.to(torch.float64).sum() + pos = torch.arange(score_start, seq_len, dtype=torch.int64, device=device) + win_start + 1 + all_pos.append(pos); all_tgt.append(vt_gpu[pos]); all_mp.append(tgt_p[idx, score_start:]) + all_H.append(ent[idx, score_start:]); all_pe.append(pe_delta[idx, score_start:]) + ap = torch.cat(all_pos); at = torch.cat(all_tgt); amp = torch.cat(all_mp) + aH = torch.cat(all_H) + n = ap.shape[0] + EPS = 1e-8 + best_ng = torch.zeros(n, device=device); found = torch.zeros(n, dtype=torch.bool, device=device) + for order in _NG_ORDERS: + m = (ap >= order) & (~found) + if not m.any(): continue + ch = ng_hashes[order][ap[m]] + cc = ng_ctx[order][ch]; has = cc >= _NG_MIN + if not has.any(): continue + ph = (ch * _NG_PAIR_MULT + at[m]) % _NG_B + ng_p = (ng_pair[order][ph].float() / cc.float().clamp(min=1)).clamp(EPS, 1 - EPS) + ix = m.nonzero(as_tuple=True)[0]; best_ng[ix[has]] = ng_p[has]; found[ix[has]] = True + alpha = 0.05 + 0.55 / (1.0 + torch.exp(-3.0 * (aH - 3.5))) + aPE = torch.cat(all_pe) + pe_conf = aPE / aPE.max().clamp(min=1e-8) + alpha = alpha * (0.5 + 1.0 * pe_conf) + mixed = torch.where(found, (1 - alpha) * amp + alpha * best_ng, amp) + ng_loss_sum -= torch.log(mixed.clamp(min=1e-20)).to(torch.float64).sum() + for order in _NG_ORDERS: + v = ap >= order + if not v.any(): continue + ch = ng_hashes[order][ap[v]] + ng_ctx[order].scatter_add_(0, ch, torch.ones_like(ch, dtype=torch.int32)) + ph = (ch * _NG_PAIR_MULT + at[v]) % _NG_B + ng_pair[order].scatter_add_(0, ph, torch.ones_like(ph, dtype=torch.int32)) + ng_loss_t = ng_loss_sum + if dist.is_available() and dist.is_initialized(): + dist.all_reduce(total_loss_sum, op=dist.ReduceOp.SUM) + dist.all_reduce(total_scored_tokens, op=dist.ReduceOp.SUM) + dist.all_reduce(total_byte_count, op=dist.ReduceOp.SUM) + dist.all_reduce(ng_loss_t, op=dist.ReduceOp.SUM) + val_loss = (total_loss_sum / total_scored_tokens).item() + bpb = (total_loss_sum / (total_byte_count * math.log(2.0))).item() + ng_bpb = (ng_loss_t / (total_byte_count * math.log(2.0))).item() + base_model.train() + return float(val_loss), float(bpb), float(ng_bpb) + + + +# ----------------------------- +# POST-TRAINING QUANTIZATION +# ----------------------------- +# +# It's silly to export our model, which is trained in bf16 and fp32, at that same precision. +# Instead, we get approximately the same model (with a small hit) by quantizing the model to int8 & zlib compressing. +# We can then decompress the model and run in higher precision for evaluation, after closing in under the size limit. + +_ctrl_default = "attn_scale,attn_scales,mlp_scale,mlp_scales,resid_mix,resid_mixes,q_gain,skip_weight,skip_weights" +CONTROL_TENSOR_NAME_PATTERNS = tuple( + p for p in os.environ.get("CONTROL_TENSOR_NAME_PATTERNS", _ctrl_default).split(",") if p) +INT8_KEEP_FLOAT_FP32_NAME_PATTERNS = tuple( + p for p in os.environ.get("INT8_KEEP_FLOAT_FP32_NAME_PATTERNS", + ",".join(CONTROL_TENSOR_NAME_PATTERNS)).split(",") if p) +INT8_KEEP_FLOAT_MAX_NUMEL = 65_536 +INT8_KEEP_FLOAT_STORE_DTYPE = torch.float16 +INT8_PER_ROW_SCALE_DTYPE = torch.float16 +INT8_CLIP_PERCENTILE = 99.99984 +INT8_CLIP_Q = INT8_CLIP_PERCENTILE / 100.0 + +def tensor_nbytes(t: Tensor) -> int: + return int(t.numel()) * int(t.element_size()) + +def keep_float_tensor(name: str, t: Tensor, passthrough_orig_dtypes: dict[str, str]) -> Tensor: + if any(pattern in name for pattern in INT8_KEEP_FLOAT_FP32_NAME_PATTERNS): + return t.float().contiguous() + if t.dtype in {torch.float32, torch.bfloat16}: + passthrough_orig_dtypes[name] = str(t.dtype).removeprefix("torch.") + return t.to(dtype=INT8_KEEP_FLOAT_STORE_DTYPE).contiguous() + return t + +def quantize_float_tensor(t: Tensor) -> tuple[Tensor, Tensor]: + t32 = t.float() + if t32.ndim == 2: + # Matrices get one scale per row, which usually tracks output-channel + # ranges much better than a single tensor-wide scale. + clip_abs = ( + torch.quantile(t32.abs(), INT8_CLIP_Q, dim=1) + if t32.numel() + else torch.empty((t32.shape[0],), dtype=torch.float32) + ) + clipped = torch.maximum(torch.minimum(t32, clip_abs[:, None]), -clip_abs[:, None]) + scale = (clip_abs / 127.0).clamp_min(1.0 / 127.0) + q = torch.clamp(torch.round(clipped / scale[:, None]), -127, 127).to(torch.int8).contiguous() + return q, scale.to(dtype=INT8_PER_ROW_SCALE_DTYPE).contiguous() + + # Vectors / scalars use a simpler per-tensor scale. + clip_abs = float(torch.quantile(t32.abs().flatten(), INT8_CLIP_Q).item()) if t32.numel() else 0.0 + scale = torch.tensor(clip_abs / 127.0 if clip_abs > 0 else 1.0, dtype=torch.float32) + q = torch.clamp(torch.round(torch.clamp(t32, -clip_abs, clip_abs) / scale), -127, 127).to(torch.int8).contiguous() + return q, scale + +def quantize_float_tensor_int6(t: Tensor) -> tuple[Tensor, Tensor]: + t32 = t.float() + if t32.ndim == 2: + best_q, best_s, best_mse = None, None, float("inf") + for pct in [0.999, 0.9999, 0.99999, 0.999999, 0.9999999]: + ca = torch.quantile(t32.abs(), pct, dim=1) if t32.numel() else torch.empty((t32.shape[0],), dtype=torch.float32) + s = (ca / 31.0).clamp_min(1.0 / 31.0) + q = torch.clamp(torch.round(torch.clamp(t32, -ca[:, None], ca[:, None]) / s[:, None]), -31, 31) + mse = ((q * s[:, None] - t32) ** 2).mean().item() + if mse < best_mse: best_q, best_s, best_mse = q.to(torch.int8).contiguous(), s.to(dtype=INT8_PER_ROW_SCALE_DTYPE).contiguous(), mse + return best_q, best_s + clip_abs = float(torch.quantile(t32.abs().flatten(), INT8_CLIP_Q).item()) if t32.numel() else 0.0 + scale = torch.tensor(clip_abs / 31.0 if clip_abs > 0 else 1.0, dtype=torch.float32) + q = torch.clamp(torch.round(torch.clamp(t32, -clip_abs, clip_abs) / scale), -31, 31).to(torch.int8).contiguous() + return q, scale + +def quantize_state_dict_int6(state_dict: dict[str, Tensor]): + quantized: dict[str, Tensor] = {} + scales: dict[str, Tensor] = {} + dtypes: dict[str, str] = {} + passthrough: dict[str, Tensor] = {} + passthrough_orig_dtypes: dict[str, str] = {} + qmeta: dict[str, dict[str, object]] = {} + stats = dict.fromkeys( + ("param_count", "num_tensors", "num_float_tensors", "num_nonfloat_tensors", "baseline_tensor_bytes", "int8_payload_bytes"), + 0, + ) + for name, tensor in state_dict.items(): + t = tensor.detach().to("cpu").contiguous() + stats["param_count"] += int(t.numel()) + stats["num_tensors"] += 1 + stats["baseline_tensor_bytes"] += tensor_nbytes(t) + if not t.is_floating_point(): + stats["num_nonfloat_tensors"] += 1 + passthrough[name] = t + stats["int8_payload_bytes"] += tensor_nbytes(t) + continue + if t.numel() <= INT8_KEEP_FLOAT_MAX_NUMEL or "tok_emb.weight" in name: + kept = keep_float_tensor(name, t, passthrough_orig_dtypes) + passthrough[name] = kept + stats["int8_payload_bytes"] += tensor_nbytes(kept) + continue + stats["num_float_tensors"] += 1 + q, s = quantize_float_tensor_int6(t) + if s.ndim > 0: + qmeta[name] = {"scheme": "per_row", "axis": 0} + quantized[name] = q + scales[name] = s + dtypes[name] = str(t.dtype).removeprefix("torch.") + stats["int8_payload_bytes"] += tensor_nbytes(q) + tensor_nbytes(s) + obj: dict[str, object] = { + "__quant_format__": "int6_per_row_v1", + "quantized": quantized, + "scales": scales, + "dtypes": dtypes, + "passthrough": passthrough, + } + if qmeta: + obj["qmeta"] = qmeta + if passthrough_orig_dtypes: + obj["passthrough_orig_dtypes"] = passthrough_orig_dtypes + return obj, stats + +def quantize_state_dict_int8(state_dict: dict[str, Tensor]): + # Single supported clean-script export format: + # - per-row int8 for 2D float tensors + # - per-tensor int8 for other float tensors + # - exact passthrough for non-floats + # - passthrough for small float tensors, stored as fp16 to save bytes + quantized: dict[str, Tensor] = {} + scales: dict[str, Tensor] = {} + dtypes: dict[str, str] = {} + passthrough: dict[str, Tensor] = {} + passthrough_orig_dtypes: dict[str, str] = {} + qmeta: dict[str, dict[str, object]] = {} + stats = dict.fromkeys( + ("param_count", "num_tensors", "num_float_tensors", "num_nonfloat_tensors", "baseline_tensor_bytes", "int8_payload_bytes"), + 0, + ) + + for name, tensor in state_dict.items(): + t = tensor.detach().to("cpu").contiguous() + stats["param_count"] += int(t.numel()) + stats["num_tensors"] += 1 + stats["baseline_tensor_bytes"] += tensor_nbytes(t) + + if not t.is_floating_point(): + stats["num_nonfloat_tensors"] += 1 + passthrough[name] = t + stats["int8_payload_bytes"] += tensor_nbytes(t) + continue + + # Small float tensors are cheap enough to keep directly. We still downcast + # fp32/bf16 passthrough tensors to fp16 so metadata does not dominate size. + if t.numel() <= INT8_KEEP_FLOAT_MAX_NUMEL: + kept = keep_float_tensor(name, t, passthrough_orig_dtypes) + passthrough[name] = kept + stats["int8_payload_bytes"] += tensor_nbytes(kept) + continue + + stats["num_float_tensors"] += 1 + q, s = quantize_float_tensor(t) + if s.ndim > 0: + qmeta[name] = {"scheme": "per_row", "axis": 0} + quantized[name] = q + scales[name] = s + dtypes[name] = str(t.dtype).removeprefix("torch.") + stats["int8_payload_bytes"] += tensor_nbytes(q) + tensor_nbytes(s) + + obj: dict[str, object] = { + "__quant_format__": "int8_clean_per_row_v1", + "quantized": quantized, + "scales": scales, + "dtypes": dtypes, + "passthrough": passthrough, + } + if qmeta: + obj["qmeta"] = qmeta + if passthrough_orig_dtypes: + obj["passthrough_orig_dtypes"] = passthrough_orig_dtypes + return obj, stats + +def dequantize_state_dict_int8(obj: dict[str, object]) -> dict[str, Tensor]: + out: dict[str, Tensor] = {} + qmeta = obj.get("qmeta", {}) + passthrough_orig_dtypes = obj.get("passthrough_orig_dtypes", {}) + for name, q in obj["quantized"].items(): + dtype = getattr(torch, obj["dtypes"][name]) + s = obj["scales"][name] + if qmeta.get(name, {}).get("scheme") == "per_row" or s.ndim > 0: + s = s.to(dtype=torch.float32) + # Broadcast the saved row scale back across trailing dimensions. + out[name] = (q.float() * s.view(q.shape[0], *([1] * (q.ndim - 1)))).to(dtype=dtype).contiguous() + else: + scale = float(s.item()) + out[name] = (q.float() * scale).to(dtype=dtype).contiguous() + for name, t in obj["passthrough"].items(): + # Restore small tensors, undoing the temporary fp16 storage cast if needed. + out_t = t.detach().to("cpu").contiguous() + orig_dtype = passthrough_orig_dtypes.get(name) + if isinstance(orig_dtype, str): + out_t = out_t.to(dtype=getattr(torch, orig_dtype)).contiguous() + out[name] = out_t + return out + + +# ----------------------------- +# DATA LOADING +# ----------------------------- + +def load_data_shard(file: Path) -> Tensor: + header_bytes = 256 * np.dtype(" None: + self.file_idx = (self.file_idx + 1) % len(self.files) + self.tokens = load_data_shard(self.files[self.file_idx]) + self.pos = 0 + + def take(self, n: int) -> Tensor: + chunks: list[Tensor] = [] + remaining = n + while remaining > 0: + avail = self.tokens.numel() - self.pos + if avail <= 0: + self._advance_file() + continue + k = min(remaining, avail) + chunks.append(self.tokens[self.pos : self.pos + k]) + self.pos += k + remaining -= k + return chunks[0] if len(chunks) == 1 else torch.cat(chunks) + + +class DistributedTokenLoader: + # Each call consumes a contiguous chunk from the shared token stream, then slices out + # one disjoint span per rank. The extra "+1" token lets us build (x, y) by shifting. + def __init__(self, pattern: str, rank: int, world_size: int, device: torch.device): + self.rank = rank + self.world_size = world_size + self.device = device + self.stream = TokenStream(pattern) + + def next_batch(self, global_tokens: int, seq_len: int, grad_accum_steps: int) -> tuple[Tensor, Tensor]: + local_tokens = global_tokens // (self.world_size * grad_accum_steps) + per_rank_span = local_tokens + 1 + chunk = self.stream.take(per_rank_span * self.world_size) + start = self.rank * per_rank_span + local = chunk[start : start + per_rank_span].to(dtype=torch.int64) + x = local[:-1].reshape(-1, seq_len) + y = local[1:].reshape(-1, seq_len) + return x.to(self.device, non_blocking=True), y.to(self.device, non_blocking=True) + +# ----------------------------- +# TRANSFORMER MODULES +# ----------------------------- + +class RMSNorm(nn.Module): + def __init__(self, eps: float | None = None): + super().__init__() + self.eps = eps + + def forward(self, x: Tensor) -> Tensor: + return F.rms_norm(x, (x.size(-1),), eps=self.eps) + + +class _FakeQuantInt6(torch.autograd.Function): + @staticmethod + def forward(ctx, w: Tensor) -> Tensor: + if w.ndim != 2: + return w + row_max = w.abs().amax(dim=1, keepdim=True).clamp_min(1e-12) + scale = row_max / 31.0 + q = (w / scale).round().clamp(-31, 31) + return q * scale + + @staticmethod + def backward(ctx, grad: Tensor) -> Tensor: + return grad + +def fake_quant_int6(w: Tensor) -> Tensor: + return _FakeQuantInt6.apply(w) + +class CastedLinear(nn.Linear): + def __init__(self, *args, **kwargs): + super().__init__(*args, **kwargs) + self.use_qat = False + + def forward(self, x: Tensor) -> Tensor: + w = self.weight + if self.use_qat and self.training: + w = fake_quant_int6(w) + bias = self.bias.to(x.dtype) if self.bias is not None else None + return F.linear(x, w.to(x.dtype), bias) + + +def restore_low_dim_params_to_fp32(module: nn.Module) -> None: + # Keep small/control parameters in fp32 even when the model body runs in bf16. + with torch.no_grad(): + for name, param in module.named_parameters(): + if (param.ndim < 2 or any(pattern in name for pattern in CONTROL_TENSOR_NAME_PATTERNS)) and param.dtype != torch.float32: + param.data = param.data.float() + + +class Rotary(nn.Module): + def __init__(self, dim: int, base: float = 10000.0): + super().__init__() + rdim = _ROPE_DIMS if _ROPE_DIMS > 0 else dim + inv_freq = 1.0 / (base ** (torch.arange(0, rdim, 2, dtype=torch.float32) / rdim)) + self.register_buffer("inv_freq", inv_freq, persistent=False) + self._seq_len_cached = 0 + self._cos_cached: Tensor | None = None + self._sin_cached: Tensor | None = None + + def forward(self, seq_len: int, device: torch.device, dtype: torch.dtype) -> tuple[Tensor, Tensor]: + if ( + self._cos_cached is None + or self._sin_cached is None + or self._seq_len_cached != seq_len + or self._cos_cached.device != device + ): + t = torch.arange(seq_len, device=device, dtype=self.inv_freq.dtype) + freqs = torch.outer(t, self.inv_freq.to(device)) + self._cos_cached = freqs.cos()[None, None, :, :] + self._sin_cached = freqs.sin()[None, None, :, :] + self._seq_len_cached = seq_len + return self._cos_cached.to(dtype=dtype), self._sin_cached.to(dtype=dtype) + + +_ROPE_DIMS = int(os.environ.get("ROPE_DIMS", 0)) + +def apply_rotary_emb(x: Tensor, cos: Tensor, sin: Tensor) -> Tensor: + rd = _ROPE_DIMS + if rd > 0 and rd < x.size(-1): + x_rope, x_pass = x[..., :rd], x[..., rd:] + half = rd // 2 + x1, x2 = x_rope[..., :half], x_rope[..., half:] + x_rope = torch.cat((x1 * cos[..., :half] + x2 * sin[..., :half], x1 * (-sin[..., :half]) + x2 * cos[..., :half]), dim=-1) + return torch.cat((x_rope, x_pass), dim=-1) + half = x.size(-1) // 2 + x1, x2 = x[..., :half], x[..., half:] + return torch.cat((x1 * cos + x2 * sin, x1 * (-sin) + x2 * cos), dim=-1) + + +_GATED_ATTN = bool(int(os.environ.get("GATED_ATTN", "0"))) +_VALUE_RESIDUAL = bool(int(os.environ.get("VALUE_RESIDUAL", "0"))) + +class CausalSelfAttention(nn.Module): + def __init__( + self, + dim: int, + num_heads: int, + num_kv_heads: int, + rope_base: float, + qk_gain_init: float, + use_xsa: bool = False, + ): + super().__init__() + if dim % num_heads != 0: + raise ValueError("model_dim must be divisible by num_heads") + if num_heads % num_kv_heads != 0: + raise ValueError("num_heads must be divisible by num_kv_heads") + self.num_heads = num_heads + self.num_kv_heads = num_kv_heads + self.head_dim = dim // num_heads + if self.head_dim % 2 != 0: + raise ValueError("head_dim must be even for RoPE") + kv_dim = self.num_kv_heads * self.head_dim + self.c_q = CastedLinear(dim, dim, bias=False) + self.c_k = CastedLinear(dim, kv_dim, bias=False) + self.c_v = CastedLinear(dim, kv_dim, bias=False) + self.proj = CastedLinear(dim, dim, bias=False) + self.proj._zero_init = True + self.q_gain = nn.Parameter(torch.full((num_heads,), qk_gain_init, dtype=torch.float32)) + if _GATED_ATTN: + self.attn_gate = nn.Parameter(torch.ones(num_heads, dtype=torch.float32)) + self.rotary = Rotary(self.head_dim, base=rope_base) + self.use_xsa = use_xsa + if _VALUE_RESIDUAL: + self.vr_lambda = nn.Parameter(torch.tensor([0.5, 0.5], dtype=torch.float32)) + + def forward(self, x: Tensor, v0: Tensor | None = None) -> Tensor: + bsz, seqlen, dim = x.shape + q = self.c_q(x).reshape(bsz, seqlen, self.num_heads, self.head_dim).transpose(1, 2) + k = self.c_k(x).reshape(bsz, seqlen, self.num_kv_heads, self.head_dim).transpose(1, 2) + v = self.c_v(x).reshape(bsz, seqlen, self.num_kv_heads, self.head_dim).transpose(1, 2) + if _VALUE_RESIDUAL and v0 is not None: + lam = torch.sigmoid(self.vr_lambda).to(dtype=v.dtype) + v = lam[0] * v0 + lam[1] * v + q = F.rms_norm(q, (q.size(-1),)) + k = F.rms_norm(k, (k.size(-1),)) + cos, sin = self.rotary(seqlen, x.device, q.dtype) + q = apply_rotary_emb(q, cos, sin) + k = apply_rotary_emb(k, cos, sin) + q = q * self.q_gain.to(dtype=q.dtype)[None, :, None, None] + y = F.scaled_dot_product_attention(q, k, v, attn_mask=None, is_causal=True, enable_gqa=(self.num_kv_heads != self.num_heads)) + if self.use_xsa: + vn = F.normalize(v.repeat_interleave(self.num_heads // self.num_kv_heads, dim=1), dim=-1) + y = y - (y * vn).sum(dim=-1, keepdim=True) * vn + if _GATED_ATTN: + y = y * torch.sigmoid(self.attn_gate).to(dtype=y.dtype)[None, :, None, None] + y = y.transpose(1, 2).contiguous().reshape(bsz, seqlen, dim) + return self.proj(y), v + + +class MLP(nn.Module): + def __init__(self, dim: int, mlp_mult: int, leaky: bool = False): + super().__init__() + hidden = mlp_mult * dim + self.fc = CastedLinear(dim, hidden, bias=False) + self.proj = CastedLinear(hidden, dim, bias=False) + self.proj._zero_init = True + self._leaky = leaky + + def forward(self, x: Tensor) -> Tensor: + x = F.leaky_relu(self.fc(x), 0.5) if self._leaky else torch.relu(self.fc(x)) + return self.proj(x.square()) + + +_LN_SCALE = bool(int(os.environ.get("LN_SCALE", "0"))) + +class Block(nn.Module): + def __init__(self, dim: int, num_heads: int, num_kv_heads: int, mlp_mult: int, + rope_base: float, qk_gain_init: float, use_xsa: bool = False, leaky: bool = False, layer_idx: int = 0): + super().__init__() + self.attn_norm = RMSNorm() + self.mlp_norm = RMSNorm() + self.attn = CausalSelfAttention(dim, num_heads, num_kv_heads, rope_base, qk_gain_init, use_xsa=use_xsa) + self.mlp = MLP(dim, mlp_mult, leaky=leaky) + self.attn_scale = nn.Parameter(torch.ones(dim, dtype=torch.float32)) + self.mlp_scale = nn.Parameter(torch.ones(dim, dtype=torch.float32)) + self.resid_mix = nn.Parameter(torch.stack((torch.ones(dim), torch.zeros(dim))).float()) + self._ln_scale = 1.0 / math.sqrt(layer_idx + 1) if _LN_SCALE else 1.0 + + def forward(self, x: Tensor, x0: Tensor, v0: Tensor | None = None) -> tuple[Tensor, Tensor]: + mix = self.resid_mix.to(dtype=x.dtype) + x = mix[0][None, None, :] * x + mix[1][None, None, :] * x0 + s = self._ln_scale + attn_out, v = self.attn(self.attn_norm(x), v0 if _VALUE_RESIDUAL else None) + x = x + s * self.attn_scale.to(dtype=x.dtype)[None, None, :] * attn_out + x = x + s * self.mlp_scale.to(dtype=x.dtype)[None, None, :] * self.mlp(self.mlp_norm(x)) + return x, v + + + +class SmearGate(nn.Module): + def __init__(self, dim: int): + super().__init__() + self.gate = nn.Parameter(torch.full((dim,), 3.0, dtype=torch.float32)) + + def forward(self, x: Tensor) -> Tensor: + g = torch.sigmoid(self.gate).to(dtype=x.dtype) + x_prev = F.pad(x[:, :-1], (0, 0, 1, 0)) + return g * x + (1.0 - g) * x_prev + + +class BigramHash(nn.Module): + def __init__(self, num_buckets: int, hash_dim: int, model_dim: int): + super().__init__() + self.num_buckets = num_buckets + self.table = nn.Embedding(num_buckets, hash_dim) + self.proj = CastedLinear(hash_dim, model_dim, bias=False) + nn.init.normal_(self.table.weight, std=0.01) + + def forward(self, input_ids: Tensor) -> Tensor: + prev_ids = torch.cat([torch.zeros_like(input_ids[:, :1]), input_ids[:, :-1]], dim=1) + h = ((prev_ids.long() * 92821 + input_ids.long()) % self.num_buckets).long() + return self.proj(self.table(h)) + + +class GPT(nn.Module): + def __init__( + self, + vocab_size: int, + num_layers: int, + model_dim: int, + num_heads: int, + num_kv_heads: int, + mlp_mult: int, + tie_embeddings: bool, + tied_embed_init_std: float, + logit_softcap: float, + rope_base: float, + qk_gain_init: float, + ): + super().__init__() + if logit_softcap <= 0.0: + raise ValueError(f"logit_softcap must be positive, got {logit_softcap}") + self.tie_embeddings = tie_embeddings + self.tied_embed_init_std = tied_embed_init_std + self.logit_softcap = logit_softcap + self.tok_emb = nn.Embedding(vocab_size, model_dim) + self.bigram_hash = BigramHash(2048, 128, model_dim) + self.smear_gate = SmearGate(model_dim) + pre_enrich_hidden = model_dim * 3 // 2 + self.pre_enrich = nn.Sequential( + CastedLinear(model_dim, pre_enrich_hidden, bias=False), + nn.GELU(), + CastedLinear(pre_enrich_hidden, model_dim, bias=False), + ) + self.num_encoder_layers = (num_layers + 1) // 2 + self.num_decoder_layers = num_layers - self.num_encoder_layers + self.num_skip_weights = min(self.num_encoder_layers, self.num_decoder_layers) + self.skip_weights = nn.Parameter(torch.ones(self.num_skip_weights, model_dim, dtype=torch.float32)) + xsa_last_n = int(os.environ.get("XSA_LAST_N", 4)) + mlp_mult_enc = int(os.environ.get("MLP_MULT_ENCODER", mlp_mult)) + mlp_mult_dec = int(os.environ.get("MLP_MULT_DECODER", mlp_mult)) + leaky = bool(int(os.environ.get("LEAKY_RELU", "0"))) + self.blocks = nn.ModuleList( + [ + Block(model_dim, num_heads, num_kv_heads, + mlp_mult_enc if i < self.num_encoder_layers else mlp_mult_dec, + rope_base, qk_gain_init, use_xsa=(i >= num_layers - xsa_last_n), leaky=leaky, layer_idx=i) + for i in range(num_layers) + ] + ) + self.final_norm = RMSNorm() + self.lm_head = None if tie_embeddings else CastedLinear(model_dim, vocab_size, bias=False) + if self.lm_head is not None: + self.lm_head._zero_init = True + self._init_weights() + + def _init_weights(self) -> None: + if self.tie_embeddings: + nn.init.normal_(self.tok_emb.weight, mean=0.0, std=self.tied_embed_init_std) + with torch.no_grad(): + U, S, V = torch.linalg.svd(self.tok_emb.weight.data, full_matrices=False) + target_S = S[0] * (1.0 / torch.arange(1, S.shape[0] + 1, dtype=S.dtype)) ** 0.5 + self.tok_emb.weight.data = (U * target_S[None, :]) @ V + for module in self.modules(): + if isinstance(module, nn.Linear) and getattr(module, "_zero_init", False): + nn.init.zeros_(module.weight) + + def _run_blocks(self, x: Tensor, x0: Tensor) -> Tensor: + v0 = None + skips: list[Tensor] = [] + for i in range(self.num_encoder_layers): + x, v = self.blocks[i](x, x0, v0) + if v0 is None: v0 = v + skips.append(x) + for i in range(self.num_decoder_layers): + if skips: + x = x + self.skip_weights[i].to(dtype=x.dtype)[None, None, :] * skips.pop() + x, v = self.blocks[self.num_encoder_layers + i](x, x0, v0) + return x + + def _compute_logits(self, x: Tensor) -> Tensor: + if self.tie_embeddings: + logits_proj = F.linear(x, self.tok_emb.weight) + else: + if self.lm_head is None: + raise RuntimeError("lm_head is required when tie_embeddings=False") + logits_proj = self.lm_head(x) + return self.logit_softcap * torch.tanh(logits_proj / self.logit_softcap) + + def forward(self, input_ids: Tensor, target_ids: Tensor) -> Tensor: + x = self.tok_emb(input_ids) + self.bigram_hash(input_ids) + x = self.smear_gate(x) + x = self.pre_enrich(x) + x = F.rms_norm(x, (x.size(-1),)) + x0 = x + x = self._run_blocks(x, x0) + x = self.final_norm(x).reshape(-1, x.size(-1)) + targets = target_ids.reshape(-1) + logits = self._compute_logits(x) + return F.cross_entropy(logits.float(), targets, reduction="mean") + + def forward_logits(self, input_ids: Tensor, return_pe_delta: bool = False) -> Tensor | tuple[Tensor, Tensor]: + x = self.tok_emb(input_ids) + self.bigram_hash(input_ids) + x = self.smear_gate(x) + x_pre = x + x = self.pre_enrich(x) + pe_delta = (x - x_pre).norm(dim=-1) if return_pe_delta else None + x = F.rms_norm(x, (x.size(-1),)) + x0 = x + x = self._run_blocks(x, x0) + x = self.final_norm(x) + logits = self._compute_logits(x) + return (logits, pe_delta) if return_pe_delta else logits + + +# ----------------------------- +# TRAINING +# ----------------------------- + +def main() -> None: + global zeropower_via_newtonschulz5 + + code = Path(__file__).read_text(encoding="utf-8") + eval_only = bool(int(os.environ.get("EVAL_ONLY", "0"))) + args = Hyperparameters() + zeropower_via_newtonschulz5 = torch.compile(zeropower_via_newtonschulz5) + + # ----------------------------- + # DISTRIBUTED + CUDA SETUP + # ----------------------------- + + distributed = "RANK" in os.environ and "WORLD_SIZE" in os.environ + rank = int(os.environ.get("RANK", "0")) + world_size = int(os.environ.get("WORLD_SIZE", "1")) + local_rank = int(os.environ.get("LOCAL_RANK", "0")) + if world_size <= 0: + raise ValueError(f"WORLD_SIZE must be positive, got {world_size}") + if 8 % world_size != 0: + raise ValueError(f"WORLD_SIZE={world_size} must divide 8 so grad_accum_steps stays integral") + grad_accum_steps = 8 // world_size + grad_scale = 1.0 / grad_accum_steps + if not torch.cuda.is_available(): + raise RuntimeError("CUDA is required") + device = torch.device("cuda", local_rank) + torch.cuda.set_device(device) + if distributed: + dist.init_process_group(backend="nccl", device_id=device) + dist.barrier() + master_process = rank == 0 + + # Fast math knobs + torch.backends.cuda.matmul.allow_tf32 = True + torch.backends.cudnn.allow_tf32 = True + from torch.backends.cuda import enable_cudnn_sdp, enable_flash_sdp, enable_math_sdp, enable_mem_efficient_sdp + + enable_cudnn_sdp(True) + enable_flash_sdp(True) + enable_mem_efficient_sdp(False) + enable_math_sdp(False) + + logfile = None + if master_process: + os.makedirs("logs", exist_ok=True) + logfile = f"logs/{args.run_id}.txt" + print(logfile) + + def log0(msg: str, console: bool = True) -> None: + if not master_process: + return + if console: + print(msg) + if logfile is not None: + with open(logfile, "a", encoding="utf-8") as f: + print(msg, file=f) + + log0(code, console=False) + log0("=" * 100, console=False) + log0(f"Running Python {sys.version}", console=False) + log0(f"Running PyTorch {torch.__version__}", console=False) + log0( + subprocess.run(["nvidia-smi"], stdout=subprocess.PIPE, stderr=subprocess.PIPE, text=True, check=False).stdout, + console=False, + ) + log0("=" * 100, console=False) + + # ----------------------------- + # TOKENIZER + VALIDATION METRIC SETUP + # ----------------------------- + + random.seed(args.seed) + np.random.seed(args.seed) + torch.manual_seed(args.seed) + torch.cuda.manual_seed_all(args.seed) + + if not args.tokenizer_path.endswith(".model"): + raise ValueError(f"Script only setup for SentencePiece .model file: {args.tokenizer_path}") + sp = spm.SentencePieceProcessor(model_file=args.tokenizer_path) + if int(sp.vocab_size()) != args.vocab_size: + raise ValueError( + f"VOCAB_SIZE={args.vocab_size} does not match tokenizer vocab_size={int(sp.vocab_size())}" + ) + dataset_dir = Path(args.data_path).resolve() + actual_train_files = len(list(dataset_dir.glob("fineweb_train_*.bin"))) + val_tokens = load_validation_tokens(args.val_files, args.train_seq_len) + base_bytes_lut, has_leading_space_lut, is_boundary_token_lut = build_sentencepiece_luts( + sp, args.vocab_size, device + ) + log0(f"val_bpb:enabled tokenizer_kind=sentencepiece tokenizer_path={args.tokenizer_path}") + log0(f"train_loader:dataset:{dataset_dir.name} train_shards:{actual_train_files}") + log0(f"val_loader:shards pattern={args.val_files} tokens:{val_tokens.numel() - 1}") + + # ----------------------------- + # MODEL + OPTIMIZER SETUP + # ----------------------------- + + base_model = GPT( + vocab_size=args.vocab_size, + num_layers=args.num_layers, + model_dim=args.model_dim, + num_heads=args.num_heads, + num_kv_heads=args.num_kv_heads, + mlp_mult=args.mlp_mult, + tie_embeddings=args.tie_embeddings, + tied_embed_init_std=args.tied_embed_init_std, + logit_softcap=args.logit_softcap, + rope_base=args.rope_base, + qk_gain_init=args.qk_gain_init, + ).to(device).bfloat16() + for module in base_model.modules(): + if isinstance(module, CastedLinear): + module.float() + restore_low_dim_params_to_fp32(base_model) + for module in base_model.modules(): + if isinstance(module, CastedLinear): + module.use_qat = True + compiled_model = torch.compile(base_model, dynamic=False, fullgraph=True) + model: nn.Module = DDP(compiled_model, device_ids=[local_rank], broadcast_buffers=False) if distributed else compiled_model + + # Optimizer split: + # - token embedding (Adam) uses EMBED_LR + # - untied lm_head (Adam) uses HEAD_LR + # - matrix params in transformer blocks use MATRIX_LR via Muon + # - vectors/scalars use SCALAR_LR via Adam + block_named_params = list(base_model.blocks.named_parameters()) + matrix_params = [ + p + for name, p in block_named_params + if p.ndim == 2 and not any(pattern in name for pattern in CONTROL_TENSOR_NAME_PATTERNS) + ] + matrix_params.extend(p for p in base_model.pre_enrich.parameters() if p.ndim == 2) + matrix_params.extend(p for p in base_model.bigram_hash.parameters() if p.ndim == 2) + scalar_params = [ + p + for name, p in block_named_params + if p.ndim < 2 or any(pattern in name for pattern in CONTROL_TENSOR_NAME_PATTERNS) + ] + if base_model.skip_weights.numel() > 0: + scalar_params.append(base_model.skip_weights) + scalar_params.append(base_model.smear_gate.gate) + token_lr = args.tied_embed_lr if args.tie_embeddings else args.embed_lr + optimizer_tok = torch.optim.AdamW( + [{"params": [base_model.tok_emb.weight], "lr": token_lr, "base_lr": token_lr}], + betas=(args.beta1, args.beta2), + eps=args.adam_eps, + weight_decay=args.adam_wd, + fused=True, + ) + optimizer_muon = Muon( + matrix_params, + lr=args.matrix_lr, + momentum=args.muon_momentum, + backend_steps=args.muon_backend_steps, + ) + for group in optimizer_muon.param_groups: + group["base_lr"] = args.matrix_lr + optimizer_scalar = torch.optim.AdamW( + [{"params": scalar_params, "lr": args.scalar_lr, "base_lr": args.scalar_lr}], + betas=(args.beta1, args.beta2), + eps=args.adam_eps, + weight_decay=args.adam_wd, + fused=True, + ) + optimizers: list[torch.optim.Optimizer] = [optimizer_tok, optimizer_muon, optimizer_scalar] + if base_model.lm_head is not None: + optimizer_head = torch.optim.Adam( + [{"params": [base_model.lm_head.weight], "lr": args.head_lr, "base_lr": args.head_lr}], + betas=(args.beta1, args.beta2), + eps=args.adam_eps, + fused=True, + ) + optimizers.insert(1, optimizer_head) + + n_params = sum(p.numel() for p in base_model.parameters()) + log0(f"model_params:{n_params}") + log0(f"world_size:{world_size} grad_accum_steps:{grad_accum_steps}") + log0("sdp_backends:cudnn=True flash=True mem_efficient=False math=False") + log0(f"attention_mode:gqa num_heads:{args.num_heads} num_kv_heads:{args.num_kv_heads}") + log0( + f"tie_embeddings:{args.tie_embeddings} embed_lr:{token_lr} " + f"head_lr:{args.head_lr if base_model.lm_head is not None else 0.0} " + f"matrix_lr:{args.matrix_lr} scalar_lr:{args.scalar_lr}" + ) + log0( + f"train_batch_tokens:{args.train_batch_tokens} train_seq_len:{args.train_seq_len} " + f"iterations:{args.iterations} warmup_steps:{args.warmup_steps} " + f"max_wallclock_seconds:{args.max_wallclock_seconds:.3f}" + ) + log0(f"seed:{args.seed}") + + # ----------------------------- + # DATA LOADER & MODEL WARMUP + # ----------------------------- + + train_loader = DistributedTokenLoader(args.train_files, rank, world_size, device) + + def zero_grad_all() -> None: + for opt in optimizers: + opt.zero_grad(set_to_none=True) + + max_wallclock_ms = 1000.0 * args.max_wallclock_seconds if args.max_wallclock_seconds > 0 else None + + def lr_mul(step: int, elapsed_ms: float) -> float: + if args.warmdown_iters <= 0: + return 1.0 + if max_wallclock_ms is None: + warmdown_start = max(args.iterations - args.warmdown_iters, 0) + return max((args.iterations - step) / max(args.warmdown_iters, 1), 0.0) if warmdown_start <= step < args.iterations else 1.0 + step_ms = elapsed_ms / max(step, 1) + warmdown_ms = args.warmdown_iters * step_ms + remaining_ms = max(max_wallclock_ms - elapsed_ms, 0.0) + return remaining_ms / max(warmdown_ms, 1e-9) if remaining_ms <= warmdown_ms else 1.0 + + if eval_only: + log0("eval_only: loading final_model.int6.ptz") + with open("final_model.int6.ptz", "rb") as f: + base_model.load_state_dict(dequantize_state_dict_int8( + torch.load(io.BytesIO(lzma.decompress(f.read())), map_location="cpu")), strict=True) + elif args.warmup_steps > 0: + initial_model_state = {name: tensor.detach().cpu().clone() for name, tensor in base_model.state_dict().items()} + initial_optimizer_states = [copy.deepcopy(opt.state_dict()) for opt in optimizers] + model.train() + for warmup_step in range(args.warmup_steps): + zero_grad_all() + for micro_step in range(grad_accum_steps): + if distributed: + model.require_backward_grad_sync = micro_step == grad_accum_steps - 1 + x, y = train_loader.next_batch(args.train_batch_tokens, args.train_seq_len, grad_accum_steps) + with torch.autocast(device_type="cuda", dtype=torch.bfloat16, enabled=True): + warmup_loss = model(x, y) + (warmup_loss * grad_scale).backward() + for opt in optimizers: + opt.step() + zero_grad_all() + if args.warmup_steps <= 20 or (warmup_step + 1) % 10 == 0 or warmup_step + 1 == args.warmup_steps: + log0(f"warmup_step:{warmup_step + 1}/{args.warmup_steps}") + base_model.load_state_dict(initial_model_state, strict=True) + for opt, state in zip(optimizers, initial_optimizer_states, strict=True): + opt.load_state_dict(state) + zero_grad_all() + if distributed: + model.require_backward_grad_sync = True + train_loader = DistributedTokenLoader(args.train_files, rank, world_size, device) + + # ----------------------------- + # MAIN TRAINING LOOP + # ----------------------------- + + training_time_ms = 0.0 + if not eval_only: + stop_after_step: int | None = None + ema_state = {k: v.detach().clone().float() for k, v in base_model.state_dict().items()} + swa_state: dict[str, Tensor] | None = None + swa_count = 0 + torch.cuda.synchronize() + t0 = time.perf_counter() + + step = 0 + while not eval_only: + last_step = step == args.iterations or (stop_after_step is not None and step >= stop_after_step) + + should_validate = last_step or (args.val_loss_every > 0 and step % args.val_loss_every == 0) + if should_validate: + torch.cuda.synchronize() + training_time_ms += 1000.0 * (time.perf_counter() - t0) + val_loss, val_bpb = eval_val( + args, + model, + rank, + world_size, + device, + grad_accum_steps, + val_tokens, + base_bytes_lut, + has_leading_space_lut, + is_boundary_token_lut, + ) + log0( + f"step:{step}/{args.iterations} val_loss:{val_loss:.4f} val_bpb:{val_bpb:.4f} " + f"train_time:{training_time_ms:.0f}ms step_avg:{training_time_ms / max(step, 1):.2f}ms" + ) + torch.cuda.synchronize() + t0 = time.perf_counter() + + if last_step: + if stop_after_step is not None and step < args.iterations: + log0( + f"stopping_early: wallclock_cap train_time:{training_time_ms:.0f}ms " + f"step:{step}/{args.iterations}" + ) + break + + elapsed_ms = training_time_ms + 1000.0 * (time.perf_counter() - t0) + scale = lr_mul(step, elapsed_ms) + zero_grad_all() + train_loss = torch.zeros((), device=device) + for micro_step in range(grad_accum_steps): + if distributed: + model.require_backward_grad_sync = micro_step == grad_accum_steps - 1 + x, y = train_loader.next_batch(args.train_batch_tokens, args.train_seq_len, grad_accum_steps) + with torch.autocast(device_type="cuda", dtype=torch.bfloat16, enabled=True): + loss = model(x, y) + train_loss += loss.detach() + (loss * grad_scale).backward() + train_loss /= grad_accum_steps + + frac = min(step / args.muon_momentum_warmup_steps, 1.0) if args.muon_momentum_warmup_steps > 0 else 1.0 + muon_momentum = (1 - frac) * args.muon_momentum_warmup_start + frac * args.muon_momentum + for group in optimizer_muon.param_groups: + group["momentum"] = muon_momentum + + for opt in optimizers: + for group in opt.param_groups: + group["lr"] = group["base_lr"] * scale + + if args.grad_clip_norm > 0: + torch.nn.utils.clip_grad_norm_(base_model.parameters(), args.grad_clip_norm) + for opt in optimizers: + opt.step() + with torch.no_grad(): + muon_lr = optimizer_muon.param_groups[0]["lr"] + for p in matrix_params: + p.mul_(1.0 - args.muon_wd * muon_lr) + zero_grad_all() + + step += 1 + with torch.no_grad(): + for k, v in base_model.state_dict().items(): + ema_state[k].mul_(args.ema_decay).add_(v.detach().float(), alpha=1.0 - args.ema_decay) + if scale < 0.2 and step % 50 == 0: + sd = {k: v.detach().cpu().float() for k, v in base_model.state_dict().items()} + if swa_state is None: swa_state, swa_count = sd, 1 + else: + for k in swa_state: swa_state[k] += sd[k] + swa_count += 1 + approx_training_time_ms = training_time_ms + 1000.0 * (time.perf_counter() - t0) + should_log_train = ( + args.train_log_every > 0 + and (step <= 10 or step % args.train_log_every == 0 or stop_after_step is not None) + ) + if should_log_train: + log0( + f"step:{step}/{args.iterations} train_loss:{train_loss.item():.4f} " + f"train_time:{approx_training_time_ms:.0f}ms step_avg:{approx_training_time_ms / step:.2f}ms" + ) + + # Needed to sync whether we've reached the wallclock cap. + reached_cap = max_wallclock_ms is not None and approx_training_time_ms >= max_wallclock_ms + if distributed and max_wallclock_ms is not None: + reached_cap_tensor = torch.tensor(int(reached_cap), device=device) + dist.all_reduce(reached_cap_tensor, op=dist.ReduceOp.MAX) + reached_cap = bool(reached_cap_tensor.item()) + if stop_after_step is None and reached_cap: + stop_after_step = step + + if not eval_only: + log0( + f"peak memory allocated: {torch.cuda.max_memory_allocated() // 1024 // 1024} MiB " + f"reserved: {torch.cuda.max_memory_reserved() // 1024 // 1024} MiB" + ) + ema_state = {k: v.cpu() for k, v in ema_state.items()} + if swa_state is not None and swa_count > 0: + log0(f"swa: averaging {swa_count} checkpoints on top of EMA") + for k in swa_state: + swa_state[k] /= swa_count + ema_state[k] = 0.5 * ema_state[k] + 0.5 * swa_state[k] + del swa_state + log0("ema: loading weights") + base_model.load_state_dict(ema_state, strict=True) + for module in base_model.modules(): + if isinstance(module, CastedLinear): + module.float() + restore_low_dim_params_to_fp32(base_model) + del ema_state + if master_process: + torch.save(base_model.state_dict(), "final_model.pt") + model_bytes = os.path.getsize("final_model.pt") + code_bytes = len(code.encode("utf-8")) + log0(f"Serialized model: {model_bytes} bytes") + log0(f"Code size: {code_bytes} bytes") + log0(f"Total submission size: {model_bytes + code_bytes} bytes") + quant_obj, quant_stats = quantize_state_dict_int6(base_model.state_dict()) + quant_buf = io.BytesIO() + torch.save(quant_obj, quant_buf) + quant_raw = quant_buf.getvalue() + quant_blob = lzma.compress(quant_raw, preset=6) + quant_raw_bytes = len(quant_raw) + if master_process: + with open("final_model.int6.ptz", "wb") as f: + f.write(quant_blob) + quant_file_bytes = os.path.getsize("final_model.int6.ptz") + code_bytes = len(code.encode("utf-8")) + ratio = quant_stats["baseline_tensor_bytes"] / max(quant_stats["int8_payload_bytes"], 1) + log0( + f"Serialized model int6+lzma: {quant_file_bytes} bytes " + f"(payload:{quant_stats['int8_payload_bytes']} raw_torch:{quant_raw_bytes} payload_ratio:{ratio:.2f}x)" + ) + log0(f"Total submission size int6+lzma: {quant_file_bytes + code_bytes} bytes") + if distributed: + dist.barrier() + + with open("final_model.int6.ptz", "rb") as f: + quant_blob_disk = f.read() + quant_state = torch.load(io.BytesIO(lzma.decompress(quant_blob_disk)), map_location="cpu") + base_model.load_state_dict(dequantize_state_dict_int8(quant_state), strict=True) + torch.cuda.synchronize() + t_qeval = time.perf_counter() + q_val_loss, q_val_bpb = eval_val( + args, model, rank, world_size, device, grad_accum_steps, + val_tokens, base_bytes_lut, has_leading_space_lut, is_boundary_token_lut, + ) + torch.cuda.synchronize() + log0( + f"final_int8_zlib_roundtrip val_loss:{q_val_loss:.4f} val_bpb:{q_val_bpb:.4f} " + f"eval_time:{1000.0 * (time.perf_counter() - t_qeval):.0f}ms" + ) + log0(f"final_int8_zlib_roundtrip_exact val_loss:{q_val_loss:.8f} val_bpb:{q_val_bpb:.8f}") + + torch.cuda.synchronize() + t_slide = time.perf_counter() + sw_val_loss, sw_val_bpb, ng_bpb = eval_val_sliding( + args, base_model, rank, world_size, device, + val_tokens, base_bytes_lut, has_leading_space_lut, is_boundary_token_lut, + ) + torch.cuda.synchronize() + log0( + f"final_sliding_window sliding_bpb:{sw_val_bpb:.4f} val_bpb:{ng_bpb:.4f} " + f"eval_time:{1000.0 * (time.perf_counter() - t_slide):.0f}ms" + ) + log0(f"final_sliding_window_exact sliding_bpb:{sw_val_bpb:.8f} val_bpb:{ng_bpb:.8f}") + + if distributed: + dist.destroy_process_group() + + +if __name__ == "__main__": + main() diff --git a/train_gpt.py b/train_gpt.py index 85e2cc463..b914949c0 100644 --- a/train_gpt.py +++ b/train_gpt.py @@ -1,7 +1,7 @@ """ The `train_gpt.py` and `train_gpt_mlx.py` scripts are intended as good launching-off points for new participants, not SOTA configs. We'll accept PRs that tune, improve, or simplify these scripts without significantly increasing complexity, but competitive submissions should stay in the `/records` folder. -Hard stop: To keep readable for newcomers, let's make sure `train_gpt.py` and `train_gpt_mlx.py` never are longer than 1500 lines. +Hard stop: `train_gpt.py` and `train_gpt_mlx.py` must never be longer than 1500 lines. """ from __future__ import annotations @@ -16,9 +16,10 @@ import sys import time import uuid -import zlib +import lzma from pathlib import Path + import numpy as np import sentencepiece as spm import torch @@ -36,8 +37,9 @@ # - vocab size 1024, sequence length 1024, tied embeddings # - 524,288 train tokens per step for 20,000 iterations with a ~10 minute cap +_RUN_CONFIG = os.environ.get("RUN_CONFIG", "A") + class Hyperparameters: - # Data paths are shard globs produced by the existing preprocessing pipeline. data_path = os.environ.get("DATA_PATH", "./data/datasets/fineweb10B_sp1024") train_files = os.path.join(data_path, "fineweb_train_*.bin") val_files = os.path.join(data_path, "fineweb_val_*.bin") @@ -45,53 +47,46 @@ class Hyperparameters: run_id = os.environ.get("RUN_ID", str(uuid.uuid4())) seed = int(os.environ.get("SEED", 1337)) - # Validation cadence and batch size. Validation always uses the full fineweb_val split. val_batch_size = int(os.environ.get("VAL_BATCH_SIZE", 524_288)) val_loss_every = int(os.environ.get("VAL_LOSS_EVERY", 1000)) train_log_every = int(os.environ.get("TRAIN_LOG_EVERY", 200)) - # Training length. iterations = int(os.environ.get("ITERATIONS", 20000)) - warmdown_iters = int(os.environ.get("WARMDOWN_ITERS", 1200)) + warmdown_iters = int(os.environ.get("WARMDOWN_ITERS", 3500 if _RUN_CONFIG == "A" else 2600)) warmup_steps = int(os.environ.get("WARMUP_STEPS", 20)) train_batch_tokens = int(os.environ.get("TRAIN_BATCH_TOKENS", 524_288)) - train_seq_len = int(os.environ.get("TRAIN_SEQ_LEN", 1024)) + train_seq_len = int(os.environ.get("TRAIN_SEQ_LEN", 2048 if _RUN_CONFIG == "A" else 1024)) max_wallclock_seconds = float(os.environ.get("MAX_WALLCLOCK_SECONDS", 600.0)) qk_gain_init = float(os.environ.get("QK_GAIN_INIT", 1.5)) - # Model shape. vocab_size = int(os.environ.get("VOCAB_SIZE", 1024)) - num_layers = int(os.environ.get("NUM_LAYERS", 9)) + num_layers = int(os.environ.get("NUM_LAYERS", 12 if _RUN_CONFIG == "C" else 10)) num_kv_heads = int(os.environ.get("NUM_KV_HEADS", 4)) model_dim = int(os.environ.get("MODEL_DIM", 512)) num_heads = int(os.environ.get("NUM_HEADS", 8)) - mlp_mult = int(os.environ.get("MLP_MULT", 2)) + mlp_mult = int(os.environ.get("MLP_MULT", 2 if _RUN_CONFIG == "C" else 3)) tie_embeddings = bool(int(os.environ.get("TIE_EMBEDDINGS", "1"))) rope_base = float(os.environ.get("ROPE_BASE", 10000.0)) logit_softcap = float(os.environ.get("LOGIT_SOFTCAP", 30.0)) - # Optimizer hyperparameters. embed_lr = float(os.environ.get("EMBED_LR", 0.6)) head_lr = float(os.environ.get("HEAD_LR", 0.008)) - tied_embed_lr = float(os.environ.get("TIED_EMBED_LR", 0.05)) + tied_embed_lr = float(os.environ.get("TIED_EMBED_LR", 0.035)) tied_embed_init_std = float(os.environ.get("TIED_EMBED_INIT_STD", 0.005)) - matrix_lr = float(os.environ.get("MATRIX_LR", 0.04)) - scalar_lr = float(os.environ.get("SCALAR_LR", 0.04)) - muon_momentum = float(os.environ.get("MUON_MOMENTUM", 0.95)) + matrix_lr = float(os.environ.get("MATRIX_LR", 0.025)) + scalar_lr = float(os.environ.get("SCALAR_LR", 0.025)) + muon_momentum = float(os.environ.get("MUON_MOMENTUM", 0.99)) muon_backend_steps = int(os.environ.get("MUON_BACKEND_STEPS", 5)) - muon_momentum_warmup_start = float(os.environ.get("MUON_MOMENTUM_WARMUP_START", 0.85)) - muon_momentum_warmup_steps = int(os.environ.get("MUON_MOMENTUM_WARMUP_STEPS", 500)) + muon_momentum_warmup_start = float(os.environ.get("MUON_MOMENTUM_WARMUP_START", 0.92)) + muon_momentum_warmup_steps = int(os.environ.get("MUON_MOMENTUM_WARMUP_STEPS", 1500)) beta1 = float(os.environ.get("BETA1", 0.9)) beta2 = float(os.environ.get("BETA2", 0.95)) adam_eps = float(os.environ.get("ADAM_EPS", 1e-8)) grad_clip_norm = float(os.environ.get("GRAD_CLIP_NORM", 0.0)) - - # Test-time training (LoRA) hyperparameters. - ttt_lora_rank = int(os.environ.get("TTT_LORA_RANK", 8)) - ttt_lora_lr = float(os.environ.get("TTT_LORA_LR", 0.01)) - ttt_chunk_size = int(os.environ.get("TTT_CHUNK_SIZE", 256)) - ttt_eval_seq_len = int(os.environ.get("TTT_EVAL_SEQ_LEN", 1024)) - ttt_batch_size = int(os.environ.get("TTT_BATCH_SIZE", 64)) + muon_wd = float(os.environ.get("MUON_WD", 0.04)) + adam_wd = float(os.environ.get("ADAM_WD", 0.04)) + ema_decay = float(os.environ.get("EMA_DECAY", 0.997)) + leaky_relu = bool(int(os.environ.get("LEAKY_RELU", "0"))) # ----------------------------- # MUON OPTIMIZER @@ -284,6 +279,123 @@ def eval_val( model.train() return float(val_loss.item()), float(bits_per_token * tokens_per_byte) + +_NG_B = 1 << 22 +_NG_ORDERS = (11, 10, 9, 8, 7, 6, 5, 4, 3, 2) +_NG_MIN = 2 +_NG_MULT = 265443576 +_NG_PAIR_MULT = 1000003 + +def eval_val_sliding( + args: Hyperparameters, + base_model: nn.Module, + rank: int, + world_size: int, + device: torch.device, + val_tokens: Tensor, + base_bytes_lut: Tensor, + has_leading_space_lut: Tensor, + is_boundary_token_lut: Tensor, + stride: int = 64, + batch_size: int = 256, +) -> tuple[float, float, float]: + seq_len = args.train_seq_len + total_tokens = val_tokens.numel() + windows: list[tuple[int, int]] = [] + pos = 0 + while pos + seq_len < total_tokens: + windows.append((pos, 0 if pos == 0 else seq_len - stride)) + pos += stride + my_windows = windows[rank::world_size] + total_loss_sum = torch.zeros((), device=device, dtype=torch.float64) + total_scored_tokens = torch.zeros((), device=device, dtype=torch.float64) + total_byte_count = torch.zeros((), device=device, dtype=torch.float64) + ng_loss_sum = torch.zeros((), device=device, dtype=torch.float64) + vt_gpu = val_tokens.to(device=device, dtype=torch.int64) + ng_ctx, ng_pair, ng_hashes = {}, {}, {} + for order in _NG_ORDERS: + ng_ctx[order] = torch.zeros(_NG_B, dtype=torch.int32, device=device) + ng_pair[order] = torch.zeros(_NG_B, dtype=torch.int32, device=device) + h = torch.zeros(total_tokens, dtype=torch.int64, device=device) + for ki in range(order - 1): + h[order-1:] = (h[order-1:] * _NG_MULT + vt_gpu[ki:total_tokens - order + 1 + ki]) % _NG_B + ng_hashes[order] = h + print(f" n-gram hashes precomputed (orders {list(_NG_ORDERS)})", flush=True) + base_model.eval() + num_batches = (len(my_windows) + batch_size - 1) // batch_size + with torch.inference_mode(): + for batch_start in range(0, len(my_windows), batch_size): + if batch_start % (batch_size * 500) == 0: + print(f" eval batch {batch_start // batch_size}/{num_batches}", flush=True) + batch_windows = my_windows[batch_start:batch_start + batch_size] + x_list, y_list = [], [] + for win_start, _ in batch_windows: + chunk = val_tokens[win_start:win_start + seq_len + 1] + x_list.append(chunk[:-1]); y_list.append(chunk[1:]) + x = torch.stack(x_list).to(device=device, dtype=torch.int64) + y = torch.stack(y_list).to(device=device, dtype=torch.int64) + with torch.autocast(device_type="cuda", dtype=torch.bfloat16, enabled=True): + logits, pe_delta = base_model.forward_logits(x, return_pe_delta=True) + per_token_loss = F.cross_entropy( + logits.float().reshape(-1, logits.size(-1)), y.reshape(-1), reduction="none", + ).reshape(len(batch_windows), seq_len) + lp = F.log_softmax(logits.float(), dim=-1) + ent = -(lp.exp() * lp).sum(dim=-1) + tgt_p = lp.gather(-1, y.unsqueeze(-1)).squeeze(-1).exp() + all_pos, all_tgt, all_mp, all_H, all_pe = [], [], [], [], [] + for idx, (win_start, score_start) in enumerate(batch_windows): + scored_loss = per_token_loss[idx, score_start:] + total_loss_sum += scored_loss.to(torch.float64).sum() + total_scored_tokens += float(scored_loss.numel()) + scored_prev = x[idx, score_start:] + scored_tgt = y[idx, score_start:] + token_bytes = base_bytes_lut[scored_tgt].to(dtype=torch.int16) + token_bytes += (has_leading_space_lut[scored_tgt] & ~is_boundary_token_lut[scored_prev]).to(dtype=torch.int16) + total_byte_count += token_bytes.to(torch.float64).sum() + pos = torch.arange(score_start, seq_len, dtype=torch.int64, device=device) + win_start + 1 + all_pos.append(pos); all_tgt.append(vt_gpu[pos]); all_mp.append(tgt_p[idx, score_start:]) + all_H.append(ent[idx, score_start:]); all_pe.append(pe_delta[idx, score_start:]) + ap = torch.cat(all_pos); at = torch.cat(all_tgt); amp = torch.cat(all_mp) + aH = torch.cat(all_H) + n = ap.shape[0] + EPS = 1e-8 + best_ng = torch.zeros(n, device=device); found = torch.zeros(n, dtype=torch.bool, device=device) + for order in _NG_ORDERS: + m = (ap >= order) & (~found) + if not m.any(): continue + ch = ng_hashes[order][ap[m]] + cc = ng_ctx[order][ch]; has = cc >= _NG_MIN + if not has.any(): continue + ph = (ch * _NG_PAIR_MULT + at[m]) % _NG_B + ng_p = (ng_pair[order][ph].float() / cc.float().clamp(min=1)).clamp(EPS, 1 - EPS) + ix = m.nonzero(as_tuple=True)[0]; best_ng[ix[has]] = ng_p[has]; found[ix[has]] = True + alpha = 0.05 + 0.55 / (1.0 + torch.exp(-3.0 * (aH - 3.5))) + aPE = torch.cat(all_pe) + pe_conf = aPE / aPE.max().clamp(min=1e-8) + alpha = alpha * (0.5 + 1.0 * pe_conf) + mixed = torch.where(found, (1 - alpha) * amp + alpha * best_ng, amp) + ng_loss_sum -= torch.log(mixed.clamp(min=1e-20)).to(torch.float64).sum() + for order in _NG_ORDERS: + v = ap >= order + if not v.any(): continue + ch = ng_hashes[order][ap[v]] + ng_ctx[order].scatter_add_(0, ch, torch.ones_like(ch, dtype=torch.int32)) + ph = (ch * _NG_PAIR_MULT + at[v]) % _NG_B + ng_pair[order].scatter_add_(0, ph, torch.ones_like(ph, dtype=torch.int32)) + ng_loss_t = ng_loss_sum + if dist.is_available() and dist.is_initialized(): + dist.all_reduce(total_loss_sum, op=dist.ReduceOp.SUM) + dist.all_reduce(total_scored_tokens, op=dist.ReduceOp.SUM) + dist.all_reduce(total_byte_count, op=dist.ReduceOp.SUM) + dist.all_reduce(ng_loss_t, op=dist.ReduceOp.SUM) + val_loss = (total_loss_sum / total_scored_tokens).item() + bpb = (total_loss_sum / (total_byte_count * math.log(2.0))).item() + ng_bpb = (ng_loss_t / (total_byte_count * math.log(2.0))).item() + base_model.train() + return float(val_loss), float(bpb), float(ng_bpb) + + + # ----------------------------- # POST-TRAINING QUANTIZATION # ----------------------------- @@ -292,22 +404,12 @@ def eval_val( # Instead, we get approximately the same model (with a small hit) by quantizing the model to int8 & zlib compressing. # We can then decompress the model and run in higher precision for evaluation, after closing in under the size limit. +_ctrl_default = "attn_scale,attn_scales,mlp_scale,mlp_scales,resid_mix,resid_mixes,q_gain,skip_weight,skip_weights" CONTROL_TENSOR_NAME_PATTERNS = tuple( - pattern - for pattern in os.environ.get( - "CONTROL_TENSOR_NAME_PATTERNS", - "attn_scale,attn_scales,mlp_scale,mlp_scales,resid_mix,resid_mixes,q_gain,skip_weight,skip_weights", - ).split(",") - if pattern -) + p for p in os.environ.get("CONTROL_TENSOR_NAME_PATTERNS", _ctrl_default).split(",") if p) INT8_KEEP_FLOAT_FP32_NAME_PATTERNS = tuple( - pattern - for pattern in os.environ.get( - "INT8_KEEP_FLOAT_FP32_NAME_PATTERNS", - ",".join(CONTROL_TENSOR_NAME_PATTERNS), - ).split(",") - if pattern -) + p for p in os.environ.get("INT8_KEEP_FLOAT_FP32_NAME_PATTERNS", + ",".join(CONTROL_TENSOR_NAME_PATTERNS)).split(",") if p) INT8_KEEP_FLOAT_MAX_NUMEL = 65_536 INT8_KEEP_FLOAT_STORE_DTYPE = torch.float16 INT8_PER_ROW_SCALE_DTYPE = torch.float16 @@ -346,6 +448,69 @@ def quantize_float_tensor(t: Tensor) -> tuple[Tensor, Tensor]: q = torch.clamp(torch.round(torch.clamp(t32, -clip_abs, clip_abs) / scale), -127, 127).to(torch.int8).contiguous() return q, scale +def quantize_float_tensor_int6(t: Tensor) -> tuple[Tensor, Tensor]: + t32 = t.float() + if t32.ndim == 2: + best_q, best_s, best_mse = None, None, float("inf") + for pct in [0.999, 0.9999, 0.99999, 0.999999, 0.9999999]: + ca = torch.quantile(t32.abs(), pct, dim=1) if t32.numel() else torch.empty((t32.shape[0],), dtype=torch.float32) + s = (ca / 31.0).clamp_min(1.0 / 31.0) + q = torch.clamp(torch.round(torch.clamp(t32, -ca[:, None], ca[:, None]) / s[:, None]), -31, 31) + mse = ((q * s[:, None] - t32) ** 2).mean().item() + if mse < best_mse: best_q, best_s, best_mse = q.to(torch.int8).contiguous(), s.to(dtype=INT8_PER_ROW_SCALE_DTYPE).contiguous(), mse + return best_q, best_s + clip_abs = float(torch.quantile(t32.abs().flatten(), INT8_CLIP_Q).item()) if t32.numel() else 0.0 + scale = torch.tensor(clip_abs / 31.0 if clip_abs > 0 else 1.0, dtype=torch.float32) + q = torch.clamp(torch.round(torch.clamp(t32, -clip_abs, clip_abs) / scale), -31, 31).to(torch.int8).contiguous() + return q, scale + +def quantize_state_dict_int6(state_dict: dict[str, Tensor]): + quantized: dict[str, Tensor] = {} + scales: dict[str, Tensor] = {} + dtypes: dict[str, str] = {} + passthrough: dict[str, Tensor] = {} + passthrough_orig_dtypes: dict[str, str] = {} + qmeta: dict[str, dict[str, object]] = {} + stats = dict.fromkeys( + ("param_count", "num_tensors", "num_float_tensors", "num_nonfloat_tensors", "baseline_tensor_bytes", "int8_payload_bytes"), + 0, + ) + for name, tensor in state_dict.items(): + t = tensor.detach().to("cpu").contiguous() + stats["param_count"] += int(t.numel()) + stats["num_tensors"] += 1 + stats["baseline_tensor_bytes"] += tensor_nbytes(t) + if not t.is_floating_point(): + stats["num_nonfloat_tensors"] += 1 + passthrough[name] = t + stats["int8_payload_bytes"] += tensor_nbytes(t) + continue + if t.numel() <= INT8_KEEP_FLOAT_MAX_NUMEL or "tok_emb.weight" in name: + kept = keep_float_tensor(name, t, passthrough_orig_dtypes) + passthrough[name] = kept + stats["int8_payload_bytes"] += tensor_nbytes(kept) + continue + stats["num_float_tensors"] += 1 + q, s = quantize_float_tensor_int6(t) + if s.ndim > 0: + qmeta[name] = {"scheme": "per_row", "axis": 0} + quantized[name] = q + scales[name] = s + dtypes[name] = str(t.dtype).removeprefix("torch.") + stats["int8_payload_bytes"] += tensor_nbytes(q) + tensor_nbytes(s) + obj: dict[str, object] = { + "__quant_format__": "int6_per_row_v1", + "quantized": quantized, + "scales": scales, + "dtypes": dtypes, + "passthrough": passthrough, + } + if qmeta: + obj["qmeta"] = qmeta + if passthrough_orig_dtypes: + obj["passthrough_orig_dtypes"] = passthrough_orig_dtypes + return obj, stats + def quantize_state_dict_int8(state_dict: dict[str, Tensor]): # Single supported clean-script export format: # - per-row int8 for 2D float tensors @@ -513,11 +678,34 @@ def forward(self, x: Tensor) -> Tensor: return F.rms_norm(x, (x.size(-1),), eps=self.eps) +class _FakeQuantInt6(torch.autograd.Function): + @staticmethod + def forward(ctx, w: Tensor) -> Tensor: + if w.ndim != 2: + return w + row_max = w.abs().amax(dim=1, keepdim=True).clamp_min(1e-12) + scale = row_max / 31.0 + q = (w / scale).round().clamp(-31, 31) + return q * scale + + @staticmethod + def backward(ctx, grad: Tensor) -> Tensor: + return grad + +def fake_quant_int6(w: Tensor) -> Tensor: + return _FakeQuantInt6.apply(w) + class CastedLinear(nn.Linear): - # Keep weights in fp32 for optimizer/state quality, cast at matmul time for bf16 compute. + def __init__(self, *args, **kwargs): + super().__init__(*args, **kwargs) + self.use_qat = False + def forward(self, x: Tensor) -> Tensor: + w = self.weight + if self.use_qat and self.training: + w = fake_quant_int6(w) bias = self.bias.to(x.dtype) if self.bias is not None else None - return F.linear(x, self.weight.to(x.dtype), bias) + return F.linear(x, w.to(x.dtype), bias) def restore_low_dim_params_to_fp32(module: nn.Module) -> None: @@ -529,10 +717,10 @@ def restore_low_dim_params_to_fp32(module: nn.Module) -> None: class Rotary(nn.Module): - # Caches cos/sin tables per sequence length on the current device. def __init__(self, dim: int, base: float = 10000.0): super().__init__() - inv_freq = 1.0 / (base ** (torch.arange(0, dim, 2, dtype=torch.float32) / dim)) + rdim = _ROPE_DIMS if _ROPE_DIMS > 0 else dim + inv_freq = 1.0 / (base ** (torch.arange(0, rdim, 2, dtype=torch.float32) / rdim)) self.register_buffer("inv_freq", inv_freq, persistent=False) self._seq_len_cached = 0 self._cos_cached: Tensor | None = None @@ -553,12 +741,24 @@ def forward(self, seq_len: int, device: torch.device, dtype: torch.dtype) -> tup return self._cos_cached.to(dtype=dtype), self._sin_cached.to(dtype=dtype) +_ROPE_DIMS = int(os.environ.get("ROPE_DIMS", 0)) + def apply_rotary_emb(x: Tensor, cos: Tensor, sin: Tensor) -> Tensor: + rd = _ROPE_DIMS + if rd > 0 and rd < x.size(-1): + x_rope, x_pass = x[..., :rd], x[..., rd:] + half = rd // 2 + x1, x2 = x_rope[..., :half], x_rope[..., half:] + x_rope = torch.cat((x1 * cos[..., :half] + x2 * sin[..., :half], x1 * (-sin[..., :half]) + x2 * cos[..., :half]), dim=-1) + return torch.cat((x_rope, x_pass), dim=-1) half = x.size(-1) // 2 x1, x2 = x[..., :half], x[..., half:] return torch.cat((x1 * cos + x2 * sin, x1 * (-sin) + x2 * cos), dim=-1) +_GATED_ATTN = bool(int(os.environ.get("GATED_ATTN", "0"))) +_VALUE_RESIDUAL = bool(int(os.environ.get("VALUE_RESIDUAL", "0"))) + class CausalSelfAttention(nn.Module): def __init__( self, @@ -567,6 +767,7 @@ def __init__( num_kv_heads: int, rope_base: float, qk_gain_init: float, + use_xsa: bool = False, ): super().__init__() if dim % num_heads != 0: @@ -585,77 +786,100 @@ def __init__( self.proj = CastedLinear(dim, dim, bias=False) self.proj._zero_init = True self.q_gain = nn.Parameter(torch.full((num_heads,), qk_gain_init, dtype=torch.float32)) + if _GATED_ATTN: + self.attn_gate = nn.Parameter(torch.ones(num_heads, dtype=torch.float32)) self.rotary = Rotary(self.head_dim, base=rope_base) + self.use_xsa = use_xsa + if _VALUE_RESIDUAL: + self.vr_lambda = nn.Parameter(torch.tensor([0.5, 0.5], dtype=torch.float32)) - def forward(self, x: Tensor, q_delta=None, v_delta=None) -> Tensor: + def forward(self, x: Tensor, v0: Tensor | None = None) -> Tensor: bsz, seqlen, dim = x.shape - q = self.c_q(x) + (q_delta if q_delta is not None else 0) - k = self.c_k(x) - v = self.c_v(x) + (v_delta if v_delta is not None else 0) - q = q.reshape(bsz, seqlen, self.num_heads, self.head_dim).transpose(1, 2) - k = k.reshape(bsz, seqlen, self.num_kv_heads, self.head_dim).transpose(1, 2) - v = v.reshape(bsz, seqlen, self.num_kv_heads, self.head_dim).transpose(1, 2) + q = self.c_q(x).reshape(bsz, seqlen, self.num_heads, self.head_dim).transpose(1, 2) + k = self.c_k(x).reshape(bsz, seqlen, self.num_kv_heads, self.head_dim).transpose(1, 2) + v = self.c_v(x).reshape(bsz, seqlen, self.num_kv_heads, self.head_dim).transpose(1, 2) + if _VALUE_RESIDUAL and v0 is not None: + lam = torch.sigmoid(self.vr_lambda).to(dtype=v.dtype) + v = lam[0] * v0 + lam[1] * v q = F.rms_norm(q, (q.size(-1),)) k = F.rms_norm(k, (k.size(-1),)) cos, sin = self.rotary(seqlen, x.device, q.dtype) q = apply_rotary_emb(q, cos, sin) k = apply_rotary_emb(k, cos, sin) q = q * self.q_gain.to(dtype=q.dtype)[None, :, None, None] - y = F.scaled_dot_product_attention( - q, - k, - v, - attn_mask=None, - is_causal=True, - enable_gqa=(self.num_kv_heads != self.num_heads), - ) + y = F.scaled_dot_product_attention(q, k, v, attn_mask=None, is_causal=True, enable_gqa=(self.num_kv_heads != self.num_heads)) + if self.use_xsa: + vn = F.normalize(v.repeat_interleave(self.num_heads // self.num_kv_heads, dim=1), dim=-1) + y = y - (y * vn).sum(dim=-1, keepdim=True) * vn + if _GATED_ATTN: + y = y * torch.sigmoid(self.attn_gate).to(dtype=y.dtype)[None, :, None, None] y = y.transpose(1, 2).contiguous().reshape(bsz, seqlen, dim) - return self.proj(y) + return self.proj(y), v class MLP(nn.Module): - # relu^2 MLP from the original modded-nanogpt setup - def __init__(self, dim: int, mlp_mult: int): + def __init__(self, dim: int, mlp_mult: int, leaky: bool = False): super().__init__() hidden = mlp_mult * dim self.fc = CastedLinear(dim, hidden, bias=False) self.proj = CastedLinear(hidden, dim, bias=False) self.proj._zero_init = True + self._leaky = leaky def forward(self, x: Tensor) -> Tensor: - x = torch.relu(self.fc(x)) + x = F.leaky_relu(self.fc(x), 0.5) if self._leaky else torch.relu(self.fc(x)) return self.proj(x.square()) +_LN_SCALE = bool(int(os.environ.get("LN_SCALE", "0"))) + class Block(nn.Module): - def __init__( - self, - dim: int, - num_heads: int, - num_kv_heads: int, - mlp_mult: int, - rope_base: float, - qk_gain_init: float, - ): + def __init__(self, dim: int, num_heads: int, num_kv_heads: int, mlp_mult: int, + rope_base: float, qk_gain_init: float, use_xsa: bool = False, leaky: bool = False, layer_idx: int = 0): super().__init__() self.attn_norm = RMSNorm() self.mlp_norm = RMSNorm() - self.attn = CausalSelfAttention(dim, num_heads, num_kv_heads, rope_base, qk_gain_init) - self.mlp = MLP(dim, mlp_mult) + self.attn = CausalSelfAttention(dim, num_heads, num_kv_heads, rope_base, qk_gain_init, use_xsa=use_xsa) + self.mlp = MLP(dim, mlp_mult, leaky=leaky) self.attn_scale = nn.Parameter(torch.ones(dim, dtype=torch.float32)) self.mlp_scale = nn.Parameter(torch.ones(dim, dtype=torch.float32)) self.resid_mix = nn.Parameter(torch.stack((torch.ones(dim), torch.zeros(dim))).float()) + self._ln_scale = 1.0 / math.sqrt(layer_idx + 1) if _LN_SCALE else 1.0 - def forward(self, x: Tensor, x0: Tensor, q_delta_fn=None, v_delta_fn=None) -> Tensor: + def forward(self, x: Tensor, x0: Tensor, v0: Tensor | None = None) -> tuple[Tensor, Tensor]: mix = self.resid_mix.to(dtype=x.dtype) x = mix[0][None, None, :] * x + mix[1][None, None, :] * x0 - n = self.attn_norm(x) - qd = q_delta_fn(n) if q_delta_fn is not None else None - vd = v_delta_fn(n) if v_delta_fn is not None else None - attn_out = self.attn(n, qd, vd) - x = x + self.attn_scale.to(dtype=x.dtype)[None, None, :] * attn_out - x = x + self.mlp_scale.to(dtype=x.dtype)[None, None, :] * self.mlp(self.mlp_norm(x)) - return x + s = self._ln_scale + attn_out, v = self.attn(self.attn_norm(x), v0 if _VALUE_RESIDUAL else None) + x = x + s * self.attn_scale.to(dtype=x.dtype)[None, None, :] * attn_out + x = x + s * self.mlp_scale.to(dtype=x.dtype)[None, None, :] * self.mlp(self.mlp_norm(x)) + return x, v + + + +class SmearGate(nn.Module): + def __init__(self, dim: int): + super().__init__() + self.gate = nn.Parameter(torch.full((dim,), 3.0, dtype=torch.float32)) + + def forward(self, x: Tensor) -> Tensor: + g = torch.sigmoid(self.gate).to(dtype=x.dtype) + x_prev = F.pad(x[:, :-1], (0, 0, 1, 0)) + return g * x + (1.0 - g) * x_prev + + +class BigramHash(nn.Module): + def __init__(self, num_buckets: int, hash_dim: int, model_dim: int): + super().__init__() + self.num_buckets = num_buckets + self.table = nn.Embedding(num_buckets, hash_dim) + self.proj = CastedLinear(hash_dim, model_dim, bias=False) + nn.init.normal_(self.table.weight, std=0.01) + + def forward(self, input_ids: Tensor) -> Tensor: + prev_ids = torch.cat([torch.zeros_like(input_ids[:, :1]), input_ids[:, :-1]], dim=1) + h = ((prev_ids.long() * 92821 + input_ids.long()) % self.num_buckets).long() + return self.proj(self.table(h)) class GPT(nn.Module): @@ -680,20 +904,27 @@ def __init__( self.tied_embed_init_std = tied_embed_init_std self.logit_softcap = logit_softcap self.tok_emb = nn.Embedding(vocab_size, model_dim) - self.num_encoder_layers = num_layers // 2 + self.bigram_hash = BigramHash(2048, 128, model_dim) + self.smear_gate = SmearGate(model_dim) + pre_enrich_hidden = model_dim * 3 // 2 + self.pre_enrich = nn.Sequential( + CastedLinear(model_dim, pre_enrich_hidden, bias=False), + nn.GELU(), + CastedLinear(pre_enrich_hidden, model_dim, bias=False), + ) + self.num_encoder_layers = (num_layers + 1) // 2 self.num_decoder_layers = num_layers - self.num_encoder_layers self.num_skip_weights = min(self.num_encoder_layers, self.num_decoder_layers) self.skip_weights = nn.Parameter(torch.ones(self.num_skip_weights, model_dim, dtype=torch.float32)) + xsa_last_n = int(os.environ.get("XSA_LAST_N", 4)) + mlp_mult_enc = int(os.environ.get("MLP_MULT_ENCODER", mlp_mult)) + mlp_mult_dec = int(os.environ.get("MLP_MULT_DECODER", mlp_mult)) + leaky = bool(int(os.environ.get("LEAKY_RELU", "0"))) self.blocks = nn.ModuleList( [ - Block( - model_dim, - num_heads, - num_kv_heads, - mlp_mult, - rope_base, - qk_gain_init, - ) + Block(model_dim, num_heads, num_kv_heads, + mlp_mult_enc if i < self.num_encoder_layers else mlp_mult_dec, + rope_base, qk_gain_init, use_xsa=(i >= num_layers - xsa_last_n), leaky=leaky, layer_idx=i) for i in range(num_layers) ] ) @@ -706,253 +937,61 @@ def __init__( def _init_weights(self) -> None: if self.tie_embeddings: nn.init.normal_(self.tok_emb.weight, mean=0.0, std=self.tied_embed_init_std) + with torch.no_grad(): + U, S, V = torch.linalg.svd(self.tok_emb.weight.data, full_matrices=False) + target_S = S[0] * (1.0 / torch.arange(1, S.shape[0] + 1, dtype=S.dtype)) ** 0.5 + self.tok_emb.weight.data = (U * target_S[None, :]) @ V for module in self.modules(): if isinstance(module, nn.Linear) and getattr(module, "_zero_init", False): nn.init.zeros_(module.weight) - def forward(self, input_ids: Tensor, target_ids: Tensor, lora=None) -> Tensor: - x = self.tok_emb(input_ids) - x = F.rms_norm(x, (x.size(-1),)) - x0 = x + def _run_blocks(self, x: Tensor, x0: Tensor) -> Tensor: + v0 = None skips: list[Tensor] = [] - - # First half stores skips; second half reuses them in reverse order. for i in range(self.num_encoder_layers): - qd = lora.q_loras[i] if lora else None - vd = lora.v_loras[i] if lora else None - x = self.blocks[i](x, x0, qd, vd) + x, v = self.blocks[i](x, x0, v0) + if v0 is None: v0 = v skips.append(x) for i in range(self.num_decoder_layers): - bi = self.num_encoder_layers + i if skips: x = x + self.skip_weights[i].to(dtype=x.dtype)[None, None, :] * skips.pop() - qd = lora.q_loras[bi] if lora else None - vd = lora.v_loras[bi] if lora else None - x = self.blocks[bi](x, x0, qd, vd) - x = self.final_norm(x) - if self.tie_embeddings: - logits = F.linear(x, self.tok_emb.weight) - else: - logits = self.lm_head(x) - logits = logits + (lora.lm_head_lora(x) if lora else 0) - logits = self.logit_softcap * torch.tanh(logits / self.logit_softcap) - if lora: - bsz, sl, V = logits.shape - return F.cross_entropy( - logits.float().reshape(-1, V), target_ids.reshape(-1), reduction="none").reshape(bsz, sl) - return F.cross_entropy(logits.float().reshape(-1, logits.size(-1)), target_ids.reshape(-1), reduction="mean") - - -# ----------------------------- -# TEST-TIME TRAINING (LoRA) -# ----------------------------- -# -# At evaluation time, we adapt per-document low-rank adapters on the validation data. -# Each document gets its own adapter, so there is no inter-document dependency. - -BOS_ID = 1 - -class BatchedLinearLoRA(nn.Module): - """LoRA for a linear layer, with independent weights per batch element. - Computes x @ Aᵀ @ Bᵀ = x @ (BA)ᵀ, i.e. the LoRA delta is ΔW = BA.""" - def __init__(self, bsz: int, in_features: int, out_features: int, rank: int): - super().__init__() - self.in_features = in_features - self.A = nn.Parameter(torch.empty(bsz, rank, in_features)) # down-projection - self.B = nn.Parameter(torch.zeros(bsz, out_features, rank)) # up-projection - self.reset() - - def forward(self, x: Tensor) -> Tensor: - return (x @ self.A.transpose(1, 2)) @ self.B.transpose(1, 2) # (bsz, T, out) - - def reset(self) -> None: - bound = 1.0 / math.sqrt(self.in_features) - with torch.no_grad(): - self.A.uniform_(-bound, bound) # kaiming-uniform - self.B.zero_() - -class BatchedTTTLoRA(nn.Module): - """All LoRA adapters for one batch: LM head and Q/V per block.""" - def __init__(self, bsz: int, model: GPT, rank: int): - super().__init__() - dim = model.tok_emb.embedding_dim - vocab = model.tok_emb.num_embeddings - self.lm_head_lora = BatchedLinearLoRA(bsz, dim, vocab, rank) - self.q_loras = nn.ModuleList() - self.v_loras = nn.ModuleList() - for block in model.blocks: - self.q_loras.append(BatchedLinearLoRA(bsz, dim, block.attn.c_q.weight.shape[0], rank)) - self.v_loras.append(BatchedLinearLoRA(bsz, dim, block.attn.c_v.weight.shape[0], rank)) - - def reset(self) -> None: - for m in self.modules(): - if isinstance(m, BatchedLinearLoRA): - m.reset() - -def _reset_ttt_optimizer(opt): - for group in opt.param_groups: - for p in group['params']: - s = opt.state.get(p) - if not s: # Fresh state. - continue - s['exp_avg'].zero_() - s['exp_avg_sq'].zero_() - s['step'].fill_(0) - -def _build_ttt_optimizer(lora, args: Hyperparameters): - return torch.optim.Adam(lora.parameters(), lr=args.ttt_lora_lr, betas=(args.beta1, args.beta2), eps=1e-10) - -def _find_docs(all_tokens: Tensor, include_next_bos: bool = True) -> list[tuple[int, int]]: - """Return (start_offset, length) for each document, identified by BOS boundaries. - - If include_next_bos is True, include next document's BOS (to match continuous-stream - eval token count exactly). - """ - bos_positions = (all_tokens == BOS_ID).nonzero(as_tuple=True)[0].numpy() - docs = [] - for i in range(len(bos_positions)): - start = int(bos_positions[i]) - end = int(bos_positions[i + 1]) if i + 1 < len(bos_positions) else all_tokens.numel() - if include_next_bos and i + 1 < len(bos_positions): - end += 1 - assert end - start >= 2 - docs.append((start, end - start)) - return docs - -def _compute_chunk_window(ci: int, pred_len: int, num_chunks: int, chunk_size: int, eval_seq_len: int): - """Return (win_start, win_len, chunk_offset, chunk_len) for chunk `ci` of a doc.""" - chunk_start = ci * chunk_size - chunk_end = pred_len if ci == num_chunks - 1 else (ci + 1) * chunk_size - win_start = max(0, chunk_end - eval_seq_len) - win_len = chunk_end - win_start - chunk_offset = chunk_start - win_start - chunk_len = chunk_end - chunk_start - return win_start, win_len, chunk_offset, chunk_len - -def _accumulate_bpb( - ptl: Tensor, x: Tensor, y: Tensor, - batch_i: int, chunk_offset: int, chunk_len: int, - base_bytes_lut: Tensor, has_leading_space_lut: Tensor, is_boundary_token_lut: Tensor, - loss_sum: Tensor, byte_sum: Tensor, token_count: Tensor, -): - """Add one doc-chunk's contribution to the running BPB accumulators.""" - lbl = ptl[batch_i, chunk_offset:chunk_offset + chunk_len].to(torch.float64) - prev = x[batch_i, chunk_offset:chunk_offset + chunk_len] - tgt = y[batch_i, chunk_offset:chunk_offset + chunk_len] - tok_bytes = base_bytes_lut[tgt].to(torch.float64) - tok_bytes += has_leading_space_lut[tgt] & ~is_boundary_token_lut[prev] - loss_sum += lbl.sum() - byte_sum += tok_bytes.sum() - token_count += chunk_len - -def eval_val_ttt_lora( - args: Hyperparameters, - base_model: GPT, - rank: int, - world_size: int, - device: torch.device, - base_bytes_lut: Tensor, - has_leading_space_lut: Tensor, - is_boundary_token_lut: Tensor, -) -> tuple[float, float]: - """Evaluate with batched LoRA test-time training. Returns (val_loss, val_bpb).""" - # Load validation tokens and find document boundaries - files = sorted(glob.glob(args.val_files)) - all_tokens = torch.cat([load_data_shard(Path(f)) for f in files]) - docs = _find_docs(all_tokens) - - # Each rank takes a contiguous slice of documents - rank_docs = docs[(len(docs) * rank) // world_size : (len(docs) * (rank + 1)) // world_size] - chunk_size = args.ttt_chunk_size - eval_seq_len = args.ttt_eval_seq_len - batch_size = args.ttt_batch_size - lora_rank = args.ttt_lora_rank - - rank_docs.sort(key=lambda d: (d[1] - 2) // chunk_size) - - base_model.eval() - for p in base_model.parameters(): - p.requires_grad_(False) - - lora = BatchedTTTLoRA(batch_size, base_model, lora_rank).to(device) - opt = _build_ttt_optimizer(lora, args) - - loss_sum = torch.zeros((), device=device, dtype=torch.float64) - byte_sum = torch.zeros((), device=device, dtype=torch.float64) - token_count = torch.zeros((), device=device, dtype=torch.float64) - - for bi in range(0, len(rank_docs), batch_size): - batch = rank_docs[bi:bi + batch_size] - bsz = len(batch) + x, v = self.blocks[self.num_encoder_layers + i](x, x0, v0) + return x - if bsz == batch_size: - cur_lora, cur_opt = lora, opt - cur_lora.reset() - _reset_ttt_optimizer(cur_opt) + def _compute_logits(self, x: Tensor) -> Tensor: + if self.tie_embeddings: + logits_proj = F.linear(x, self.tok_emb.weight) else: - cur_lora = BatchedTTTLoRA(bsz, base_model, lora_rank).to(device) - cur_opt = _build_ttt_optimizer(cur_lora, args) - - pred_lens = [doc_len - 1 for _, doc_len in batch] - num_chunks = [(pl + chunk_size - 1) // chunk_size for pl in pred_lens] - max_nc = max(num_chunks) - - for ci in range(max_nc): - chunk_stats = _compute_chunk_window(ci, (ci + 1) * chunk_size, ci + 1, chunk_size, eval_seq_len) - context_size, chunk_offset = chunk_stats[1], chunk_stats[2] - - active = [ci < nc for nc in num_chunks] - needs_train = any(ci < nc - 1 for nc in num_chunks) - - x = torch.zeros(bsz, context_size, dtype=torch.int64, device=device) - y = torch.zeros(bsz, context_size, dtype=torch.int64, device=device) - doc_info = [] # (chunk_offset, chunk_len) per doc - for b in range(bsz): - if not active[b]: - doc_info.append((0, 0)) - continue - ds, dl = batch[b] - ws, wl, co, cl = _compute_chunk_window(ci, pred_lens[b], num_chunks[b], chunk_size, eval_seq_len) - chunk = all_tokens[ds + ws: ds + ws + wl + 1] - toks = chunk.to(dtype=torch.int64, device=device) - x[b, :wl] = toks[:-1] - y[b, :wl] = toks[1:] - doc_info.append((co, cl)) - - # Forward pass (keep grad graph alive only when we need to train) - if needs_train: - with torch.autocast(device_type="cuda", dtype=torch.bfloat16): - ptl = base_model(x, y, lora=cur_lora) - else: - with torch.no_grad(), torch.autocast(device_type="cuda", dtype=torch.bfloat16): - ptl = base_model(x, y, lora=cur_lora) - - # Score: accumulate loss and byte counts for BPB (before training on chunk) - with torch.no_grad(): - for b in range(bsz): - if not active[b]: - continue - co, cl = doc_info[b] - _accumulate_bpb( - ptl, x, y, b, co, cl, base_bytes_lut, has_leading_space_lut, - is_boundary_token_lut, loss_sum, byte_sum, token_count) - - # Train: one Adam step on the LoRA params using this chunk's loss - if needs_train: - mask = torch.tensor([float(ci < num_chunks[b] - 1) for b in range(bsz)], device=device) - per_doc = ptl[:, chunk_offset:chunk_offset + chunk_size].mean(dim=-1) - cur_opt.zero_grad() - (per_doc * mask).sum().backward() - cur_opt.step() - - if dist.is_available() and dist.is_initialized(): - dist.all_reduce(loss_sum, op=dist.ReduceOp.SUM) - dist.all_reduce(byte_sum, op=dist.ReduceOp.SUM) - dist.all_reduce(token_count, op=dist.ReduceOp.SUM) + if self.lm_head is None: + raise RuntimeError("lm_head is required when tie_embeddings=False") + logits_proj = self.lm_head(x) + return self.logit_softcap * torch.tanh(logits_proj / self.logit_softcap) + + def forward(self, input_ids: Tensor, target_ids: Tensor) -> Tensor: + x = self.tok_emb(input_ids) + self.bigram_hash(input_ids) + x = self.smear_gate(x) + x = self.pre_enrich(x) + x = F.rms_norm(x, (x.size(-1),)) + x0 = x + x = self._run_blocks(x, x0) + x = self.final_norm(x).reshape(-1, x.size(-1)) + targets = target_ids.reshape(-1) + logits = self._compute_logits(x) + return F.cross_entropy(logits.float(), targets, reduction="mean") + + def forward_logits(self, input_ids: Tensor, return_pe_delta: bool = False) -> Tensor | tuple[Tensor, Tensor]: + x = self.tok_emb(input_ids) + self.bigram_hash(input_ids) + x = self.smear_gate(x) + x_pre = x + x = self.pre_enrich(x) + pe_delta = (x - x_pre).norm(dim=-1) if return_pe_delta else None + x = F.rms_norm(x, (x.size(-1),)) + x0 = x + x = self._run_blocks(x, x0) + x = self.final_norm(x) + logits = self._compute_logits(x) + return (logits, pe_delta) if return_pe_delta else logits - val_loss = float(loss_sum.item() / token_count.item()) - val_bpb = float((loss_sum.item() / math.log(2.0)) / byte_sum.item()) - return val_loss, val_bpb # ----------------------------- # TRAINING @@ -962,6 +1001,7 @@ def main() -> None: global zeropower_via_newtonschulz5 code = Path(__file__).read_text(encoding="utf-8") + eval_only = bool(int(os.environ.get("EVAL_ONLY", "0"))) args = Hyperparameters() zeropower_via_newtonschulz5 = torch.compile(zeropower_via_newtonschulz5) @@ -993,7 +1033,7 @@ def main() -> None: torch.backends.cudnn.allow_tf32 = True from torch.backends.cuda import enable_cudnn_sdp, enable_flash_sdp, enable_math_sdp, enable_mem_efficient_sdp - enable_cudnn_sdp(False) + enable_cudnn_sdp(True) enable_flash_sdp(True) enable_mem_efficient_sdp(False) enable_math_sdp(False) @@ -1069,9 +1109,10 @@ def log0(msg: str, console: bool = True) -> None: for module in base_model.modules(): if isinstance(module, CastedLinear): module.float() - if isinstance(module, Rotary): - module.inv_freq.data = module.inv_freq.data.float() restore_low_dim_params_to_fp32(base_model) + for module in base_model.modules(): + if isinstance(module, CastedLinear): + module.use_qat = True compiled_model = torch.compile(base_model, dynamic=False, fullgraph=True) model: nn.Module = DDP(compiled_model, device_ids=[local_rank], broadcast_buffers=False) if distributed else compiled_model @@ -1086,6 +1127,8 @@ def log0(msg: str, console: bool = True) -> None: for name, p in block_named_params if p.ndim == 2 and not any(pattern in name for pattern in CONTROL_TENSOR_NAME_PATTERNS) ] + matrix_params.extend(p for p in base_model.pre_enrich.parameters() if p.ndim == 2) + matrix_params.extend(p for p in base_model.bigram_hash.parameters() if p.ndim == 2) scalar_params = [ p for name, p in block_named_params @@ -1093,11 +1136,13 @@ def log0(msg: str, console: bool = True) -> None: ] if base_model.skip_weights.numel() > 0: scalar_params.append(base_model.skip_weights) + scalar_params.append(base_model.smear_gate.gate) token_lr = args.tied_embed_lr if args.tie_embeddings else args.embed_lr - optimizer_tok = torch.optim.Adam( + optimizer_tok = torch.optim.AdamW( [{"params": [base_model.tok_emb.weight], "lr": token_lr, "base_lr": token_lr}], betas=(args.beta1, args.beta2), eps=args.adam_eps, + weight_decay=args.adam_wd, fused=True, ) optimizer_muon = Muon( @@ -1108,10 +1153,11 @@ def log0(msg: str, console: bool = True) -> None: ) for group in optimizer_muon.param_groups: group["base_lr"] = args.matrix_lr - optimizer_scalar = torch.optim.Adam( + optimizer_scalar = torch.optim.AdamW( [{"params": scalar_params, "lr": args.scalar_lr, "base_lr": args.scalar_lr}], betas=(args.beta1, args.beta2), eps=args.adam_eps, + weight_decay=args.adam_wd, fused=True, ) optimizers: list[torch.optim.Optimizer] = [optimizer_tok, optimizer_muon, optimizer_scalar] @@ -1127,7 +1173,7 @@ def log0(msg: str, console: bool = True) -> None: n_params = sum(p.numel() for p in base_model.parameters()) log0(f"model_params:{n_params}") log0(f"world_size:{world_size} grad_accum_steps:{grad_accum_steps}") - log0("sdp_backends:cudnn=False flash=True mem_efficient=False math=False") + log0("sdp_backends:cudnn=True flash=True mem_efficient=False math=False") log0(f"attention_mode:gqa num_heads:{args.num_heads} num_kv_heads:{args.num_kv_heads}") log0( f"tie_embeddings:{args.tie_embeddings} embed_lr:{token_lr} " @@ -1164,9 +1210,12 @@ def lr_mul(step: int, elapsed_ms: float) -> float: remaining_ms = max(max_wallclock_ms - elapsed_ms, 0.0) return remaining_ms / max(warmdown_ms, 1e-9) if remaining_ms <= warmdown_ms else 1.0 - # Warmup primes the compiled forward/backward/optimizer paths, then we restore the - # initial weights/optimizer state so measured training starts from the true init. - if args.warmup_steps > 0: + if eval_only: + log0("eval_only: loading final_model.int6.ptz") + with open("final_model.int6.ptz", "rb") as f: + base_model.load_state_dict(dequantize_state_dict_int8( + torch.load(io.BytesIO(lzma.decompress(f.read())), map_location="cpu")), strict=True) + elif args.warmup_steps > 0: initial_model_state = {name: tensor.detach().cpu().clone() for name, tensor in base_model.state_dict().items()} initial_optimizer_states = [copy.deepcopy(opt.state_dict()) for opt in optimizers] model.train() @@ -1197,12 +1246,16 @@ def lr_mul(step: int, elapsed_ms: float) -> float: # ----------------------------- training_time_ms = 0.0 - stop_after_step: int | None = None - torch.cuda.synchronize() - t0 = time.perf_counter() + if not eval_only: + stop_after_step: int | None = None + ema_state = {k: v.detach().clone().float() for k, v in base_model.state_dict().items()} + swa_state: dict[str, Tensor] | None = None + swa_count = 0 + torch.cuda.synchronize() + t0 = time.perf_counter() step = 0 - while True: + while not eval_only: last_step = step == args.iterations or (stop_after_step is not None and step >= stop_after_step) should_validate = last_step or (args.val_loss_every > 0 and step % args.val_loss_every == 0) @@ -1263,9 +1316,22 @@ def lr_mul(step: int, elapsed_ms: float) -> float: torch.nn.utils.clip_grad_norm_(base_model.parameters(), args.grad_clip_norm) for opt in optimizers: opt.step() + with torch.no_grad(): + muon_lr = optimizer_muon.param_groups[0]["lr"] + for p in matrix_params: + p.mul_(1.0 - args.muon_wd * muon_lr) zero_grad_all() step += 1 + with torch.no_grad(): + for k, v in base_model.state_dict().items(): + ema_state[k].mul_(args.ema_decay).add_(v.detach().float(), alpha=1.0 - args.ema_decay) + if scale < 0.2 and step % 50 == 0: + sd = {k: v.detach().cpu().float() for k, v in base_model.state_dict().items()} + if swa_state is None: swa_state, swa_count = sd, 1 + else: + for k in swa_state: swa_state[k] += sd[k] + swa_count += 1 approx_training_time_ms = training_time_ms + 1000.0 * (time.perf_counter() - t0) should_log_train = ( args.train_log_every > 0 @@ -1286,62 +1352,61 @@ def lr_mul(step: int, elapsed_ms: float) -> float: if stop_after_step is None and reached_cap: stop_after_step = step - log0( - f"peak memory allocated: {torch.cuda.max_memory_allocated() // 1024 // 1024} MiB " - f"reserved: {torch.cuda.max_memory_reserved() // 1024 // 1024} MiB" - ) - - # ----------------------------- - # SERIALIZATION + ROUNDTRIP VALIDATION - # ----------------------------- - # Save the raw state (useful for debugging/loading in PyTorch directly), then always produce - # the compressed int8+zlib artifact and validate the round-tripped weights. - - if master_process: - torch.save(base_model.state_dict(), "final_model.pt") - model_bytes = os.path.getsize("final_model.pt") - code_bytes = len(code.encode("utf-8")) - log0(f"Serialized model: {model_bytes} bytes") - log0(f"Code size: {code_bytes} bytes") - log0(f"Total submission size: {model_bytes + code_bytes} bytes") - - quant_obj, quant_stats = quantize_state_dict_int8(base_model.state_dict()) - quant_buf = io.BytesIO() - torch.save(quant_obj, quant_buf) - quant_raw = quant_buf.getvalue() - quant_blob = zlib.compress(quant_raw, level=9) - quant_raw_bytes = len(quant_raw) - if master_process: - with open("final_model.int8.ptz", "wb") as f: - f.write(quant_blob) - quant_file_bytes = os.path.getsize("final_model.int8.ptz") - code_bytes = len(code.encode("utf-8")) - ratio = quant_stats["baseline_tensor_bytes"] / max(quant_stats["int8_payload_bytes"], 1) + if not eval_only: log0( - f"Serialized model int8+zlib: {quant_file_bytes} bytes " - f"(payload:{quant_stats['int8_payload_bytes']} raw_torch:{quant_raw_bytes} payload_ratio:{ratio:.2f}x)" + f"peak memory allocated: {torch.cuda.max_memory_allocated() // 1024 // 1024} MiB " + f"reserved: {torch.cuda.max_memory_reserved() // 1024 // 1024} MiB" ) - log0(f"Total submission size int8+zlib: {quant_file_bytes + code_bytes} bytes") + ema_state = {k: v.cpu() for k, v in ema_state.items()} + if swa_state is not None and swa_count > 0: + log0(f"swa: averaging {swa_count} checkpoints on top of EMA") + for k in swa_state: + swa_state[k] /= swa_count + ema_state[k] = 0.5 * ema_state[k] + 0.5 * swa_state[k] + del swa_state + log0("ema: loading weights") + base_model.load_state_dict(ema_state, strict=True) + for module in base_model.modules(): + if isinstance(module, CastedLinear): + module.float() + restore_low_dim_params_to_fp32(base_model) + del ema_state + if master_process: + torch.save(base_model.state_dict(), "final_model.pt") + model_bytes = os.path.getsize("final_model.pt") + code_bytes = len(code.encode("utf-8")) + log0(f"Serialized model: {model_bytes} bytes") + log0(f"Code size: {code_bytes} bytes") + log0(f"Total submission size: {model_bytes + code_bytes} bytes") + quant_obj, quant_stats = quantize_state_dict_int6(base_model.state_dict()) + quant_buf = io.BytesIO() + torch.save(quant_obj, quant_buf) + quant_raw = quant_buf.getvalue() + quant_blob = lzma.compress(quant_raw, preset=6) + quant_raw_bytes = len(quant_raw) + if master_process: + with open("final_model.int6.ptz", "wb") as f: + f.write(quant_blob) + quant_file_bytes = os.path.getsize("final_model.int6.ptz") + code_bytes = len(code.encode("utf-8")) + ratio = quant_stats["baseline_tensor_bytes"] / max(quant_stats["int8_payload_bytes"], 1) + log0( + f"Serialized model int6+lzma: {quant_file_bytes} bytes " + f"(payload:{quant_stats['int8_payload_bytes']} raw_torch:{quant_raw_bytes} payload_ratio:{ratio:.2f}x)" + ) + log0(f"Total submission size int6+lzma: {quant_file_bytes + code_bytes} bytes") + if distributed: + dist.barrier() - if distributed: - dist.barrier() - with open("final_model.int8.ptz", "rb") as f: + with open("final_model.int6.ptz", "rb") as f: quant_blob_disk = f.read() - quant_state = torch.load(io.BytesIO(zlib.decompress(quant_blob_disk)), map_location="cpu") + quant_state = torch.load(io.BytesIO(lzma.decompress(quant_blob_disk)), map_location="cpu") base_model.load_state_dict(dequantize_state_dict_int8(quant_state), strict=True) torch.cuda.synchronize() t_qeval = time.perf_counter() q_val_loss, q_val_bpb = eval_val( - args, - model, - rank, - world_size, - device, - grad_accum_steps, - val_tokens, - base_bytes_lut, - has_leading_space_lut, - is_boundary_token_lut, + args, model, rank, world_size, device, grad_accum_steps, + val_tokens, base_bytes_lut, has_leading_space_lut, is_boundary_token_lut, ) torch.cuda.synchronize() log0( @@ -1350,19 +1415,18 @@ def lr_mul(step: int, elapsed_ms: float) -> float: ) log0(f"final_int8_zlib_roundtrip_exact val_loss:{q_val_loss:.8f} val_bpb:{q_val_bpb:.8f}") - # LoRA test-time training evaluation (the competition score) - torch._dynamo.reset() torch.cuda.synchronize() - t_ttt = time.perf_counter() - ttt_val_loss, ttt_val_bpb = eval_val_ttt_lora( + t_slide = time.perf_counter() + sw_val_loss, sw_val_bpb, ng_bpb = eval_val_sliding( args, base_model, rank, world_size, device, - base_bytes_lut, has_leading_space_lut, is_boundary_token_lut, + val_tokens, base_bytes_lut, has_leading_space_lut, is_boundary_token_lut, ) torch.cuda.synchronize() log0( - f"final_int8_ttt_lora val_loss:{ttt_val_loss:.4f} val_bpb:{ttt_val_bpb:.4f} " - f"eval_time:{1000.0 * (time.perf_counter() - t_ttt):.0f}ms" + f"final_sliding_window sliding_bpb:{sw_val_bpb:.4f} val_bpb:{ng_bpb:.4f} " + f"eval_time:{1000.0 * (time.perf_counter() - t_slide):.0f}ms" ) + log0(f"final_sliding_window_exact sliding_bpb:{sw_val_bpb:.8f} val_bpb:{ng_bpb:.8f}") if distributed: dist.destroy_process_group()