From d217bf5917bf68c6142277bf5bbbe195b7ed47b1 Mon Sep 17 00:00:00 2001 From: translatingthename Date: Sat, 11 Apr 2026 15:32:19 -0400 Subject: [PATCH] =?UTF-8?q?Non-record:=20Pre-Quant=20AdamW=20TTT=20(Compil?= =?UTF-8?q?ed)=20+=20SP8192=20+=20Depth=20Recurrence=20=E2=80=94=20val=5Fb?= =?UTF-8?q?pb=201.0587=20(3-seed=20mean)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Non-record submission. Pre-quant TTT violates Condition 3 of Issue #1017 (score-before-update). Submitted as technique study documenting: - Condition 3 boundary quantification (illegal TTT -0.044 vs legal -0.002) - Compiled TTT (torch.compile 2x speedup, applicable to legal TTT) - Artifact budget engineering (VE dim optimization, pruning analysis) 3-seed mean sliding BPB: 1.05869 (std 0.00038) All artifacts under 16,000,000 bytes. Zero pruning needed. --- .../README.md | 104 + .../submission.json | 28 + .../train_gpt.py | 2808 +++++++++++++++++ .../train_seed1337.log | 97 + .../train_seed2024.log | 98 + .../train_seed42.log | 95 + 6 files changed, 3230 insertions(+) create mode 100644 records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/README.md create mode 100644 records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/submission.json create mode 100644 records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_gpt.py create mode 100644 records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_seed1337.log create mode 100644 records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_seed2024.log create mode 100644 records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_seed42.log diff --git a/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/README.md b/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/README.md new file mode 100644 index 0000000000..4b0284a675 --- /dev/null +++ b/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/README.md @@ -0,0 +1,104 @@ +# Non-record: Pre-Quant AdamW TTT (Compiled) + SP8192 + Depth Recurrence + +> **Compliance note:** This submission violates Condition 3 of Issue #1017 (score-before-update). Pre-quant TTT fine-tunes on val tokens before scoring them. Submitted as a technique study, not a leaderboard claim. + +**val_bpb = 1.0587** (3-seed mean, std 0.0004) | **~15.5 MB** | 8xH100 SXM + +## 3-Seed Results + +| Seed | Sliding BPB | Roundtrip BPB | Artifact | +|------|-------------|---------------|----------| +| 42 | 1.05840 | 1.06847 | 15,477,275 | +| 1337 | 1.05856 | 1.06904 | 15,439,370 | +| 2024 | 1.05912 | 1.06921 | 15,480,770 | +| **Mean** | **1.05869** | **1.06891** | **15,465,805** | +| **Std** | **0.00038** | **0.00037** | | + +## Why this is a useful non-record + +### 1. Quantifying the Condition 3 boundary + +This submission provides a controlled measurement of how much BPB improvement comes from violating Condition 3: + +| Configuration | BPB | Source | Measured? | +|---|---:|---|---| +| Post-EMA (no TTT, no GPTQ) | 1.1028 | This submission | Yes | +| **Post-GPTQ sliding (illegal 6-epoch TTT)** | **1.0587** | This submission | Yes | +| Post-GPTQ sliding (no TTT) | ~1.106 | Estimated: post-EMA + ~0.003 GPTQ gap | No | +| Post-GPTQ sliding (legal score-first TTT) | ~1.104 | Estimated from PR #1493 delta (-0.002) | No | + +The two measured points bound the illegal TTT contribution at **-0.044 BPB** (post-EMA 1.103 → post-GPTQ sliding 1.059). For comparison, the legal score-first TTT in merged PR #1493 contributes approximately -0.002 BPB (sliding 1.083 → TTT 1.081). This is not an apples-to-apples comparison — the illegal variant uses AdamW for 6 full epochs while the legal variant uses SGD for 3 epochs per chunk, on a different base model — but the order-of-magnitude gap illustrates why Condition 3 is load-bearing. + +**On the theoretical ceiling:** Issue #1017 states: *"Corpus-level TTT has a ceiling of approximately 0.0003 bits"* — this refers specifically to the gain from closing the train-val distribution gap, which the author measured as negligible for FineWeb. However, the author also notes that *"a model that undertrained on the training distribution can still benefit from additional learning at test time."* This means legal TTT can legitimately exceed the 0.0003 ceiling if the model hasn't fully converged during training (our 600s-capped model is certainly in this regime). The merged #1493's legal TTT gain of -0.002 BPB is consistent with this — it reflects real undertraining compensation, not memorization. + +Our illegal TTT's -0.044 BPB gain, however, is 22x larger than legal TTT on a similar architecture. This magnitude is not explainable by undertraining compensation alone and is consistent with memorization of the validation set. A per-epoch ablation (not performed in this submission) would strengthen this argument: if the gain scales roughly linearly with epoch count rather than saturating quickly, that would be a direct memorization signature. + +### 2. Compiled TTT: torch.compile for test-time training + +We demonstrate that `torch.compile(dynamic=False, fullgraph=True)` can be applied to TTT models for a **2x speedup** (860s → 426s for 6 epochs). This is safe because: + +- TTT operates in train mode with `torch.autocast` +- No `torch.inference_mode()` — avoids rotary cache poisoning (a $60+ lesson from our development) +- Fresh model instance created before TTT (deletes compiled training model, resets dynamo) +- Compilation overhead (~20s) amortized over multiple epochs + +This technique applies equally to legal score-first TTT and would reduce eval-time TTT costs. + +### 3. Artifact budget engineering under 16MB + +With SP8192, fitting under 16MB required careful component analysis: + +| Component | Compressed Cost | BPB Benefit | Decision | +|---|---:|---|---| +| BigramHash 2048×128 | +109KB | ~0.001 at SP8192 | **Dropped** — marginal at large vocab | +| VE dim=128 → dim=44 | -340KB | -0.001 | **Shrunk** — optimal via EV analysis | +| VE dim=44 → dim=0 | -150KB | -0.001 | Kept — positive expected value | + +We optimized VE dimension by sweeping dims 0-128, measuring compressed artifact size at each, computing pruning probability vs BPB tradeoff, and selecting the dimension that minimized expected BPB accounting for pruning risk. dim=44 gives 0% pruning risk with 39KB margin. + +## Compliance Statement + +**This submission violates Condition 3 of Issue #1017.** Pre-quant TTT (lines 2417-2455 of `train_gpt.py`) runs 6 AdamW epochs on the full val stream before GPTQ quantization. The same tokens are then scored via sliding window evaluation. No score-before-adapt discipline is implemented. This pattern is structurally identical to the closed PR #1376 and the withdrawn PR #1485 (@ndokutovich acknowledged the violation). + +## Key Techniques + +1. **SP8192 + GPTQ SDClip** — int6 matrices (k=12.85), int8 embeddings (k=20.0) (PR #1394 @clarkkev) +2. **3-Layer Depth Recurrence** (L3-5, 14 virtual from 11 physical) (PR #1493 @bigbag) +3. **Parallel Residuals** (L7+, GPT-J style) (PR #1412 @Robby955, PR #1204 @msisovic) +4. **Pre-Quant AdamW TTT** — 6 epochs, `torch.compile` 2x speedup, freeze 2 blocks (PR #1485 @ndokutovich) +5. **QK-Gain 5.25** + MuonEq-R (Polar Express) + EMA 0.9965 + warmdown 72% (PR #1493 @bigbag) + +## Architecture + +11L × 512d × 8H/4KV, MLP 4× (2048), LeakyReLU(0.5)², Partial RoPE (16/64), LN scale, tied embeddings, softcap=30. Depth recurrence [0,1,2,3,4,5,3,4,5,6,7,8,9,10] = 14 virtual layers. Parallel residuals L7+. XSA all layers. VE dim=44 L9-10. SmearGate. + +## Training + +MuonEq-R (Polar Express, 4 NS steps) + AdamW. ~5160 steps in 600s on 8×H100 SXM. Linear warmdown to 0 over final 72%. EMA 0.9965. Late QAT at LR scale < 15%. + +## Pre-Quant AdamW TTT (VIOLATES CONDITION 3) + +Fine-tunes the EMA model on the full validation token stream before GPTQ: + +- `torch.compile(dynamic=False, fullgraph=True)` for 2x speedup (426s vs 860s) +- AdamW, lr=0.0005, weight_decay=0.0, cosine decay to lr×0.1 +- 6 epochs, freeze first 2 transformer blocks +- Batch: 32 sequences × 2048 tokens, grad clip 1.0 +- Fresh model instance (avoids inference_mode rotary cache poisoning) + +## Quantization + +GPTQ int6 SDClip (k=12.85) + int8 embeddings (k=20.0). 32 AR self-gen calibration sequences. Brotli-11 compression. Zero pruning on all seeds. + +## Reproduction + +```bash +pip install brotli sentencepiece +MATCHED_FINEWEB_REPO_ID=kevclark/parameter-golf python3 data/cached_challenge_fineweb.py --variant sp8192 +VOCAB_SIZE=8192 BIGRAM_VOCAB_SIZE=0 VE_DIM=44 SEED=42 \ + torchrun --standalone --nproc_per_node=8 train_gpt.py +``` + +## Credits + +PR #1394 @clarkkev, PR #1493 @bigbag, PR #1485 @ndokutovich, PR #1412 @Robby955, PR #1204 @msisovic, PR #1285 @dexhunter, PR #549 @abaybektursun diff --git a/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/submission.json b/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/submission.json new file mode 100644 index 0000000000..38a69b739c --- /dev/null +++ b/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/submission.json @@ -0,0 +1,28 @@ +{ + "score": 1.05869, + "score_metric": "val_bpb", + "seeds": { + "42": { + "sliding_bpb": 1.0584, + "roundtrip_bpb": 1.06847, + "artifact_bytes": 15477275 + }, + "1337": { + "sliding_bpb": 1.05856, + "roundtrip_bpb": 1.06904, + "artifact_bytes": 15439370 + }, + "2024": { + "sliding_bpb": 1.05912, + "roundtrip_bpb": 1.06921, + "artifact_bytes": 15480770 + } + }, + "mean_sliding_bpb": 1.05869, + "std_sliding_bpb": 0.00038, + "mean_roundtrip_bpb": 1.06891, + "max_artifact_bytes": 15480770, + "hardware": "8xH100 SXM", + "train_wallclock_seconds": 600, + "track": "10min_16mb" +} \ No newline at end of file diff --git a/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_gpt.py b/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_gpt.py new file mode 100644 index 0000000000..2302920375 --- /dev/null +++ b/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_gpt.py @@ -0,0 +1,2808 @@ +from __future__ import annotations +import copy +import glob +import io +import lzma +import struct +try: + import brotli + _HAS_BROTLI = True +except ImportError: + _HAS_BROTLI = False +import math +import os +import random +import subprocess +import sys +import time +import uuid +import zlib +from pathlib import Path +try: + import zstandard + _COMPRESSOR = "zstd" +except ImportError: + _COMPRESSOR = "zlib" +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 +try: + from flash_attn_interface import flash_attn_func as _fa3_func + _HAS_FA3 = True +except ImportError: + _HAS_FA3 = False + +def flash_attn_3_func(q, k, v, causal=True): + """FA3 with SDPA fallback. q/k/v: [B, T, H, D].""" + if _HAS_FA3: + return _fa3_func(q, k, v, causal=causal) + # SDPA fallback: [B,T,H,D] -> [B,H,T,D] + q2 = q.transpose(1, 2) + k2 = k.transpose(1, 2) + v2 = v.transpose(1, 2) + # GQA: expand KV heads to match Q heads + if k2.size(1) != q2.size(1): + rep = q2.size(1) // k2.size(1) + k2 = k2.repeat_interleave(rep, dim=1) + v2 = v2.repeat_interleave(rep, dim=1) + o = F.scaled_dot_product_attention(q2, k2, v2, is_causal=causal) + return o.transpose(1, 2) # [B,H,T,D] -> [B,T,H,D] +class Hyperparameters: + data_path = os.environ.get("DATA_PATH", "./data/datasets/fineweb10B_sp8192") + 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_8192_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", 4000)) + train_log_every = int(os.environ.get("TRAIN_LOG_EVERY", 500)) + iterations = int(os.environ.get("ITERATIONS", 20000)) + warmdown_iters = int(os.environ.get("WARMDOWN_ITERS", 3500)) + warmdown_frac = float(os.environ.get("WARMDOWN_FRAC", 0.72)) # PR #1493: 72% warmdown (merged SOTA) + lr_floor = float(os.environ.get("LR_FLOOR", 0.0)) # PR #1395: linear decay to 0 reduces quant gap by 61% + warmup_steps = int(os.environ.get("WARMUP_STEPS", 20)) + train_batch_tokens = int(os.environ.get("TRAIN_BATCH_TOKENS", 786_432)) + train_seq_len = int(os.environ.get("TRAIN_SEQ_LEN", 2048)) + eval_seq_len = int(os.environ.get("EVAL_SEQ_LEN", 2048)) + max_wallclock_seconds = float(os.environ.get("MAX_WALLCLOCK_SECONDS", 600.0)) + qk_gain_init = float(os.environ.get("QK_GAIN_INIT", 5.25)) # PR #1493: QK-Gain 5.25 (merged SOTA) + vocab_size = int(os.environ.get("VOCAB_SIZE", 8192)) + num_layers = int(os.environ.get("NUM_LAYERS", 11)) + 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 = float(os.environ.get("MLP_MULT", 4.0)) # 4x matches all top PRs (2048 hidden at dim=512) + 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.03)) # PR #1493: 0.03 + tied_embed_init_std = float(os.environ.get("TIED_EMBED_INIT_STD", 0.005)) + matrix_lr = float(os.environ.get("MATRIX_LR", 0.022)) # PR #1493: 0.022 + scalar_lr = float(os.environ.get("SCALAR_LR", 0.02)) # PR #1493: 0.02 + muon_momentum = float(os.environ.get("MUON_MOMENTUM", 0.99)) + muon_backend_steps = int(os.environ.get("MUON_BACKEND_STEPS", 4)) # 4 with Polar Express, 5 with standard + 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.3)) + eval_stride = int(os.environ.get("EVAL_STRIDE", 64)) + mtp_num_heads = int(os.environ.get("MTP_NUM_HEADS", 0)) + mtp_loss_weight = float(os.environ.get("MTP_LOSS_WEIGHT", 0.2)) + muon_beta2 = float(os.environ.get("MUON_BETA2", 0.95)) + swa_enabled = bool(int(os.environ.get("SWA_ENABLED", "1"))) + swa_every = int(os.environ.get("SWA_EVERY", 50)) + lawa_enabled = bool(int(os.environ.get("LAWA_ENABLED", "0"))) + lawa_k = int(os.environ.get("LAWA_K", 10)) + lawa_freq = int(os.environ.get("LAWA_FREQ", 100)) + muon_wd = float(os.environ.get("MUON_WD", 0.095)) # PR #1493: 0.095 (merged SOTA) + embed_wd = float(os.environ.get("EMBED_WD", 0.095)) # PR #1493: embed WD matches muon WD + adam_wd = float(os.environ.get("ADAM_WD", 0.02)) # PR #1493 default + muoneq_r = bool(int(os.environ.get("MUONEQ_R", "1"))) # MuonEq-R row normalization (arXiv:2603.28254) + qat_enabled = bool(int(os.environ.get("QAT_ENABLED", "0"))) + bigram_vocab_size = int(os.environ.get("BIGRAM_VOCAB_SIZE", 2048)) + bigram_dim = int(os.environ.get("BIGRAM_DIM", 128)) + trigram_enabled = bool(int(os.environ.get("TRIGRAM", "0"))) # TrigramHash (off by default, risky) + xsa_last_n = int(os.environ.get("XSA_LAST_N", 11)) # XSA on ALL layers (our novel contribution) + rope_dims = int(os.environ.get("ROPE_DIMS", 16)) + ln_scale = bool(int(os.environ.get("LN_SCALE", "1"))) + dtg_enabled = bool(int(os.environ.get("DTG_ENABLED", "0"))) + late_qat_threshold = float(os.environ.get("LATE_QAT_THRESHOLD", 0.15)) + ve_enabled = bool(int(os.environ.get("VE_ENABLED", "1"))) + ve_dim = int(os.environ.get("VE_DIM", 44)) # 44 optimal: 0% prune risk, 39KB margin, EV BPB 1.0580 + ve_layers = os.environ.get("VE_LAYERS", "9,10") + gated_attention = bool(int(os.environ.get("GATED_ATTENTION", "0"))) + value_residual = bool(int(os.environ.get("VALUE_RESIDUAL", "0"))) # VRL with sigmoid gates (off by default, risky) + # Depth recurrence (PR #1296, #1204, #1260): reuse layers for free depth + recur_layers = os.environ.get("RECUR_LAYERS", "3,4,5") # PR #1493: 3-layer recurrence (merged SOTA) + recur_start_step = int(os.environ.get("RECUR_START_STEP", 3000)) + # Parallel residuals (PR #1296, #1204, #1289): split attn/MLP into lanes + parallel_start_layer = int(os.environ.get("PARALLEL_START_LAYER", "7")) + # GPTQ calibration + gptq_calib_batches = int(os.environ.get("GPTQ_CALIB_BATCHES", 256)) + gptq_block_size = int(os.environ.get("GPTQ_BLOCK_SIZE", 128)) + # SDClip: clip = k * std(row) instead of percentile search. PR #1394. + # For int6 (b=6): k=12.85 for weight matrices, k=20 for int8 embeddings. + # Entropy ≈ b - log2(k) + const, so larger k → lower entropy → better compression. + sdclip_enabled = bool(int(os.environ.get("SDCLIP_ENABLED", "1"))) + sdclip_k_int6 = float(os.environ.get("SDCLIP_K_INT6", 12.85)) + sdclip_k_int8 = float(os.environ.get("SDCLIP_K_INT8", 20.0)) + # Compression: 'auto' picks best of rANS vs Brotli. 'rans', 'brotli', 'lzma' force a method. + compress_method = os.environ.get("COMPRESS_METHOD", "auto") + # SLOT (per-batch delta optimization) + slot_enabled = bool(int(os.environ.get("SLOT_ENABLED", "1"))) # On by default + slot_lr = float(os.environ.get("SLOT_LR", 0.005)) # AdamW SLOT only + slot_steps = int(os.environ.get("SLOT_STEPS", 8)) # AdamW SLOT only + slot_causal = bool(int(os.environ.get("SLOT_CAUSAL", "1"))) # Context-only (default: safe) + slot_lbfgs = bool(int(os.environ.get("SLOT_LBFGS", "1"))) # L-BFGS logit-space (PR #1350) + slot_lbfgs_max_iter = int(os.environ.get("SLOT_LBFGS_MAX_ITER", 25)) + slot_lbfgs_history = int(os.environ.get("SLOT_LBFGS_HISTORY", 20)) + slot_focal_tokens = int(os.environ.get("SLOT_FOCAL_TOKENS", 128)) + slot_delta_clip = float(os.environ.get("SLOT_DELTA_CLIP", 5.0)) + # Pre-quant AdamW TTT (PR #1485 style): fine-tune on val data BEFORE quantization + # Adapted weights get baked into the artifact — no eval-time adaptation needed (Track A) + ttt_enabled = bool(int(os.environ.get("TTT_ENABLED", "1"))) # ON by default for v9 + ttt_epochs = int(os.environ.get("TTT_EPOCHS", 6)) # PR #1485: 6 epochs + ttt_lr = float(os.environ.get("TTT_LR", 0.0005)) # PR #1485: 0.0005 + ttt_freeze_blocks = int(os.environ.get("TTT_FREEZE_BLOCKS", 2)) # PR #1485: freeze first 2 blocks + ttt_batch_seqs = int(os.environ.get("TTT_BATCH_SEQS", 32)) # PR #1485: 32 sequences per batch + +# --- Batched Newton-Schulz orthogonalization --- + +# Polar Express: minimax-optimal per-step coefficients (arXiv:2505.16932, PR #1344) +# 4 steps with these coefficients matches or beats 5 steps with standard (3.4445, -4.7750, 2.0315) +_PE_COEFFS = [ + (8.156554524902461, -22.48329292557795, 15.878769915207462), + (4.042929935166739, -2.808917465908714, 0.5000178451051316), + (3.8916678022926607, -2.772484153217685, 0.5060648178503393), + (3.285753657755655, -2.3681294933425376, 0.46449024233003106), + (2.3465413258596377, -1.7097828382687081, 0.42323551169305323), +] +_STANDARD_COEFFS = (3.4445, -4.7750, 2.0315) + +def zeropower_via_newtonschulz5(G: Tensor, steps: int = 4, eps: float = 1e-7, + muoneq_r: bool = False, + polar_express: bool = True) -> Tensor: + """Batched Newton-Schulz orthogonalization with Polar Express coefficients. + polar_express=True: minimax-optimal per-step coefficients (4 steps recommended). + polar_express=False: standard fixed coefficients (5 steps recommended). + MuonEq-R (arXiv:2603.28254): row-normalize before NS for better conditioning.""" + was_2d = G.ndim == 2 + if was_2d: + G = G.unsqueeze(0) + X = G.bfloat16() + transposed = X.size(-2) > X.size(-1) + if transposed: + X = X.mT + if muoneq_r: + row_norms = X.norm(dim=-1, keepdim=True).clamp_min(eps) + X = X / row_norms + X = X / (X.norm(dim=(-2, -1), keepdim=True) + eps) + if polar_express: + for a, b, c in _PE_COEFFS[:steps]: + A = X @ X.mT + B = b * A + c * (A @ A) + X = a * X + B @ X + else: + a, b, c = _STANDARD_COEFFS + for _ in range(steps): + A = X @ X.mT + B = b * A + c * (A @ A) + X = a * X + B @ X + if transposed: + X = X.mT + if was_2d: + X = X.squeeze(0) + return X + +# --- Parallel Muon optimizer --- + +class Muon(torch.optim.Optimizer): + """Parallel Muon: post-backward reduce-scatter -> local NS5 -> all-gather. + + No DDP for bank params. After backward, this optimizer: + 1. Launches async reduce-scatter for all banks (biggest first) + 2. Returns control so Adam can step on small params while RS is in-flight + 3. Waits for each RS, runs local NS5 on the shard, launches async all-gather + 4. Each all-gather overlaps with next bank's NS5 + """ + def __init__(self, params, lr: float, momentum: float, backend_steps: int, + nesterov: bool = True, weight_decay: float = 0.0, muoneq_r: bool = False): + super().__init__( + params, + dict(lr=lr, momentum=momentum, backend_steps=backend_steps, + nesterov=nesterov, weight_decay=weight_decay), + ) + self._muoneq_r = muoneq_r + self._built = False + + def _build(self): + self._distributed = dist.is_available() and dist.is_initialized() + self._world_size = dist.get_world_size() if self._distributed else 1 + self._rank = dist.get_rank() if self._distributed else 0 + ws = self._world_size + + self._bank_meta = [] + for group in self.param_groups: + for p in group["params"]: + B = p.shape[0] + padded_B = ((B + ws - 1) // ws) * ws + shard_B = padded_B // ws + tail = p.shape[1:] + dev = p.device + self._bank_meta.append({ + 'p': p, + 'B': B, + 'padded_grad': torch.zeros(padded_B, *tail, device=dev, dtype=torch.bfloat16), + 'shard': torch.zeros(shard_B, *tail, device=dev, dtype=torch.bfloat16), + 'shard_mom': torch.zeros(shard_B, *tail, device=dev, dtype=torch.bfloat16), + 'full_update': torch.zeros(padded_B, *tail, device=dev, dtype=torch.bfloat16), + 'scale': max(1, p.shape[-2] / p.shape[-1]) ** 0.5, + }) + # Sort by size descending -- launch biggest reduce-scatters first + self._bank_meta.sort(key=lambda m: -m['p'].numel()) + self._built = True + + def launch_reduce_scatters(self): + """Phase 1: launch async reduce-scatter for all banks. Call right after backward.""" + if not self._built: + self._build() + if not self._distributed: + return + self._rs_futures = [] + for m in self._bank_meta: + p = m['p'] + if p.grad is None: + self._rs_futures.append(None) + continue + pg = m['padded_grad'] + pg[:m['B']].copy_(p.grad.bfloat16()) + if pg.shape[0] > m['B']: + pg[m['B']:].zero_() + fut = dist.reduce_scatter_tensor(m['shard'], pg, op=dist.ReduceOp.AVG, async_op=True) + self._rs_futures.append(fut) + + @torch.no_grad() + def step(self, closure=None): + """Phase 3: wait for RS, local NS5, all-gather. Call AFTER Adam steps.""" + loss = None + if closure is not None: + with torch.enable_grad(): + loss = closure() + + if not self._built: + self._build() + + for group in self.param_groups: + lr = group["lr"] + momentum = group["momentum"] + backend_steps = group["backend_steps"] + nesterov = group["nesterov"] + wd = group.get("weight_decay", 0.0) + + prev_ag_handle = None + prev_m = None + + sharded = self._distributed and hasattr(self, '_rs_futures') + + for i, m in enumerate(self._bank_meta): + p = m['p'] + if p.grad is None: + continue + + if prev_ag_handle is not None: + prev_ag_handle.wait() + pp = prev_m['p'] + upd = prev_m['full_update'][:prev_m['B']] + if wd > 0.0: + pp.data.mul_(1.0 - lr * wd) + pp.add_(upd.to(dtype=pp.dtype), alpha=-lr * prev_m['scale']) + + if sharded and self._rs_futures[i] is not None: + self._rs_futures[i].wait() + g = m['shard'] + buf = m['shard_mom'] + else: + g = p.grad.bfloat16() + 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: + update = g.add(buf, alpha=momentum) + else: + update = buf + + update = zeropower_via_newtonschulz5(update, steps=backend_steps, muoneq_r=self._muoneq_r) + + if sharded: + prev_ag_handle = dist.all_gather_into_tensor( + m['full_update'], update, async_op=True) + prev_m = m + else: + if wd > 0.0: + p.data.mul_(1.0 - lr * wd) + p.add_(update.to(dtype=p.dtype), alpha=-lr * m['scale']) + + if prev_ag_handle is not None: + prev_ag_handle.wait() + pp = prev_m['p'] + upd = prev_m['full_update'][:prev_m['B']] + if wd > 0.0: + pp.data.mul_(1.0 - lr * wd) + pp.add_(upd.to(dtype=pp.dtype), alpha=-lr * prev_m['scale']) + + if hasattr(self, '_rs_futures'): + del self._rs_futures + + return loss + +# --- Tokenizer evaluation helpers --- + +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("\u2581"): + 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}") + 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, + eval_seq_len: int | None = None, +) -> tuple[float, float]: + seq_len = eval_seq_len or args.train_seq_len + local_batch_tokens = args.val_batch_size // (world_size * grad_accum_steps) + if local_batch_tokens < 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}, seq_len={seq_len}" + ) + local_batch_seqs = local_batch_tokens // seq_len + total_seqs = (val_tokens.numel() - 1) // 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 * seq_len + raw_end = batch_seq_end * seq_len + 1 + local = val_tokens[raw_start:raw_end].to(device=device, dtype=torch.int64, non_blocking=True) + x = local[:-1].reshape(-1, seq_len) + y = local[1:].reshape(-1, 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) + +# --- Quantization helpers --- + +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,smear,dtg_gate,ve_layer_scales,ve_shared.scale,attn_gate,vr_lambda,lane_merge", + ).split(",") + if pattern +) +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 +) +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: + 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() + 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_state_dict_int8(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: + 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) + 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(): + 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: + 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 CastedLinear(nn.Linear): + _qat_enabled: bool = False + def forward(self, x: Tensor) -> Tensor: + w = self.weight.to(x.dtype) + if CastedLinear._qat_enabled and self.training and w.ndim == 2: + with torch.no_grad(): + w32 = self.weight.float() + row_max = w32.abs().amax(dim=1) + scale = (row_max / 31.0).clamp_min(1.0 / 31.0) + w_q = (torch.clamp(torch.round(w32 / scale[:, None]), -32, 31) * scale[:, None]).to(x.dtype) + w = w + (w_q - w).detach() + bias = self.bias.to(x.dtype) if self.bias is not None else None + return F.linear(x, w, bias) +def restore_low_dim_params_to_fp32(module: nn.Module) -> None: + 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, train_seq_len: int = 1024, rope_dims: int = 0): + super().__init__() + self.dim = dim + self.base = base + self.train_seq_len = train_seq_len + self.rope_dims = rope_dims if rope_dims > 0 else dim + inv_freq = 1.0 / (base ** (torch.arange(0, self.rope_dims, 2, dtype=torch.float32) / self.rope_dims)) + 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 + ): + rd = self.rope_dims + if seq_len > self.train_seq_len: + scale = seq_len / self.train_seq_len + new_base = self.base * (scale ** (rd / (rd - 2))) + inv_freq = 1.0 / (new_base ** (torch.arange(0, rd, 2, dtype=torch.float32, device=device) / rd)) + else: + inv_freq = self.inv_freq.to(device) + t = torch.arange(seq_len, device=device, dtype=inv_freq.dtype) + freqs = torch.outer(t, inv_freq) + 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) +def apply_rotary_emb(x: Tensor, cos: Tensor, sin: Tensor, rope_dims: int = 0) -> Tensor: + if rope_dims > 0 and rope_dims < x.size(-1): + x_rope, x_pass = x[..., :rope_dims], x[..., rope_dims:] + half = rope_dims // 2 + x1, x2 = x_rope[..., :half], x_rope[..., half:] + x_rope = torch.cat((x1 * cos + x2 * sin, x1 * (-sin) + x2 * cos), 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) + +class CausalSelfAttention(nn.Module): + def __init__( + self, + dim: int, + num_heads: int, + num_kv_heads: int, + rope_base: float, + qk_gain_init: float, + gated_attention: bool = False, + value_residual: 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") + # No CastedLinear -- weights come from banks + self.q_gain = nn.Parameter(torch.full((num_heads,), qk_gain_init, dtype=torch.float32)) + self.rope_dims = 0 # set by GPT.__init__ for partial RoPE + self.rotary = Rotary(self.head_dim, base=rope_base, train_seq_len=1024) + self.use_xsa = False # set by GPT.__init__ for deep layers only + # Gated attention and value residual (non-banked small params) + self.gated_attention = gated_attention + if gated_attention: + self.attn_gate = nn.Linear(dim, num_heads, bias=True) + nn.init.zeros_(self.attn_gate.weight) + nn.init.constant_(self.attn_gate.bias, 4.0) + self.value_residual = value_residual + if value_residual: + self.vrl_alpha = nn.Parameter(torch.zeros(1, dtype=torch.float32)) # sigmoid gate (PR #569 style) + def _xsa_efficient(self, y: Tensor, v: Tensor) -> Tensor: + """Efficient XSA: subtract self-value projection via GQA-aware reshape (no repeat_interleave). + y: [B, T, H, D], v: [B, T, Hkv, D]. H must be divisible by Hkv.""" + B, T, H, D = y.shape + Hkv = v.size(-2) + group = H // Hkv + y_g = y.reshape(B, T, Hkv, group, D) # [B, T, Hkv, group, D] + vn = F.normalize(v, dim=-1).unsqueeze(-2) # [B, T, Hkv, 1, D] -- broadcast ready + proj = (y_g * vn).sum(dim=-1, keepdim=True) * vn + return (y_g - proj).reshape(B, T, H, D) + def forward(self, x: Tensor, q_w: Tensor, k_w: Tensor, v_w: Tensor, out_w: Tensor, v_embed: Tensor | None = None, v0: Tensor | None = None) -> tuple[Tensor, Tensor | None]: + bsz, seqlen, dim = x.shape + q = F.linear(x, q_w.to(x.dtype)).reshape(bsz, seqlen, self.num_heads, self.head_dim) + k = F.linear(x, k_w.to(x.dtype)).reshape(bsz, seqlen, self.num_kv_heads, self.head_dim) + v = F.linear(x, v_w.to(x.dtype)) + if v_embed is not None: + v = v + v_embed + v = v.reshape(bsz, seqlen, self.num_kv_heads, self.head_dim) + raw_v = v if self.value_residual else None + if self.value_residual and v0 is not None: + alpha = torch.sigmoid(self.vrl_alpha.to(dtype=v.dtype)) + v = v + alpha * v0 # sigmoid-gated residual (PR #569 style) + 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, self.rope_dims) + k = apply_rotary_emb(k, cos, sin, self.rope_dims) + q = q * self.q_gain.to(dtype=q.dtype)[None, None, :, None] + y = flash_attn_3_func(q, k, v, causal=True) + if self.use_xsa: + y = self._xsa_efficient(y, v) + if self.gated_attention: + # gate shape: (bsz, seqlen, num_heads) -> (bsz, seqlen, num_heads, 1) for B,T,H,D layout + gate = torch.sigmoid(self.attn_gate(x)).unsqueeze(-1) + y = y * gate + y = y.reshape(bsz, seqlen, dim) + return F.linear(y, out_w.to(x.dtype)), raw_v + +class SmearGate(nn.Module): + def __init__(self, dim: int): + super().__init__() + self.gate = nn.Parameter(torch.zeros(dim, dtype=torch.float32)) + def forward(self, x: Tensor) -> Tensor: + g = torch.sigmoid(self.gate.to(dtype=x.dtype))[None, None, :] + x_prev = torch.cat([torch.zeros_like(x[:, :1]), x[:, :-1]], dim=1) + return (1 - g) * x + g * x_prev + +class BigramHashEmbedding(nn.Module): + def __init__(self, bigram_vocab_size: int, bigram_dim: int, model_dim: int, trigram: bool = False): + super().__init__() + self.bigram_vocab_size = bigram_vocab_size + self._trigram = trigram + self.embed = nn.Embedding(bigram_vocab_size, bigram_dim) + nn.init.zeros_(self.embed.weight) + self.proj = CastedLinear(bigram_dim, model_dim, bias=False) if bigram_dim != model_dim else None + if self.proj is not None: + nn.init.zeros_(self.proj.weight) + self.scale = nn.Parameter(torch.tensor(0.05, dtype=torch.float32)) + def bigram_hash(self, tokens: Tensor) -> Tensor: + t = tokens.to(torch.int32) + mod = self.bigram_vocab_size - 1 + out = torch.empty_like(t) + out[..., 0] = mod + out[..., 1:] = torch.bitwise_xor(36313 * t[..., 1:], 27191 * t[..., :-1]) % mod + return out.long() + def trigram_hash(self, tokens: Tensor) -> Tensor: + """Hash (t-2, t-1, t) trigrams into same embedding table. Zero extra params.""" + t = tokens.to(torch.int32) + mod = self.bigram_vocab_size - 1 + out = torch.empty_like(t) + out[..., :2] = mod + out[..., 2:] = (36313 * t[..., 2:] ^ 27191 * t[..., 1:-1] ^ 51497 * t[..., :-2]) % mod + return out.long() + def forward(self, token_ids: Tensor) -> Tensor: + h = self.embed(self.bigram_hash(token_ids)) + if self._trigram: + h = h + self.embed(self.trigram_hash(token_ids)) + if self.proj is not None: + h = self.proj(h) + return h * self.scale.to(dtype=h.dtype) + +class ValueEmbedding(nn.Module): + """Reinject token identity into attention values at specific layers. + Each table maps vocab tokens to a low-dim embedding, projected to model_dim.""" + def __init__(self, vocab_size: int, ve_dim: int, model_dim: int): + super().__init__() + self.embed = nn.Embedding(vocab_size, ve_dim) + nn.init.normal_(self.embed.weight, std=0.01) + self.proj = CastedLinear(ve_dim, model_dim, bias=False) if ve_dim != model_dim else None + if self.proj is not None: + nn.init.zeros_(self.proj.weight) + self.scale = nn.Parameter(torch.tensor(0.1, dtype=torch.float32)) + def forward(self, token_ids: Tensor) -> Tensor: + h = self.embed(token_ids) + if self.proj is not None: + h = self.proj(h) + return h * self.scale.to(dtype=h.dtype) + +class MLP(nn.Module): + def __init__(self, dim: int, mlp_mult: int): + super().__init__() + # No CastedLinear -- weights come from banks + def forward(self, x: Tensor, up_w: Tensor, down_w: Tensor) -> Tensor: + x = F.leaky_relu(F.linear(x, up_w.to(x.dtype)), negative_slope=0.5) + return F.linear(x.square(), down_w.to(x.dtype)) + +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, + layer_idx: int = 0, + ln_scale: bool = False, + dtg: bool = False, + gated_attention: bool = False, + value_residual: bool = False, + ): + super().__init__() + self.attn_norm = RMSNorm() + self.mlp_norm = RMSNorm() + self.attn = CausalSelfAttention(dim, num_heads, num_kv_heads, rope_base, qk_gain_init, + gated_attention=gated_attention, value_residual=value_residual) + self.mlp = MLP(dim, mlp_mult) + 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_factor = 1.0 / math.sqrt(layer_idx + 1) if ln_scale else 1.0 + if dtg: + self.dtg_gate = nn.Linear(dim, 1, bias=True) + nn.init.zeros_(self.dtg_gate.weight) + nn.init.constant_(self.dtg_gate.bias, 2.0) + else: + self.dtg_gate = None + def forward(self, x: Tensor, x0: Tensor, q_w: Tensor, k_w: Tensor, v_w: Tensor, out_w: Tensor, up_w: Tensor, down_w: Tensor, v_embed: Tensor | None = None, v0: Tensor | None = None) -> tuple[Tensor, Tensor | None]: + mix = self.resid_mix.to(dtype=x.dtype) + x_in = mix[0][None, None, :] * x + mix[1][None, None, :] * x0 + attn_out, raw_v = self.attn(self.attn_norm(x_in) * self.ln_scale_factor, q_w, k_w, v_w, out_w, v_embed=v_embed, v0=v0) + x_out = x_in + self.attn_scale.to(dtype=x_in.dtype)[None, None, :] * attn_out + x_out = x_out + self.mlp_scale.to(dtype=x_out.dtype)[None, None, :] * self.mlp(self.mlp_norm(x_out) * self.ln_scale_factor, up_w, down_w) + if self.dtg_gate is not None: + gate = torch.sigmoid(self.dtg_gate(x_in.detach())) + x_out = x_in + gate * (x_out - x_in) + return x_out, raw_v + +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, + mtp_num_heads: int = 0, + mtp_loss_weight: float = 0.1, + bigram_vocab_size: int = 0, + bigram_dim: int = 128, + xsa_last_n: int = 0, + rope_dims: int = 0, + ln_scale: bool = False, + dtg: bool = False, + ve_enabled: bool = False, + ve_dim: int = 128, + ve_layers: str = "9,10", + gated_attention: bool = False, + value_residual: bool = False, + recur_layers: str = "", + parallel_start_layer: int = -1, + ): + super().__init__() + self._ve_target_dim = num_kv_heads * (model_dim // num_heads) # kv_dim for value projection + 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.value_residual = value_residual + self.mtp_num_heads = mtp_num_heads + self.mtp_loss_weight = mtp_loss_weight + self.tok_emb = nn.Embedding(vocab_size, model_dim) + self.bigram = BigramHashEmbedding(bigram_vocab_size, bigram_dim, model_dim, trigram=bool(int(os.environ.get("TRIGRAM", "0")))) if bigram_vocab_size > 0 else None + self.smear = SmearGate(model_dim) + self.num_encoder_layers = num_layers // 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.sigmoid_skips = bool(int(os.environ.get("SIGMOID_SKIPS", "0"))) + # Init: sigmoid(2.2) ≈ 0.9, close to the original ones(). Learned gate in [0,1]. + skip_init = 2.2 if self.sigmoid_skips else 1.0 + self.skip_weights = nn.Parameter(torch.full((self.num_skip_weights, model_dim), skip_init, dtype=torch.float32)) + # Parameter banks: contiguous 3D tensors for batched optimizer + head_dim = model_dim // num_heads + kv_dim = num_kv_heads * head_dim + mlp_dim = int(mlp_mult * model_dim) + self.num_layers = num_layers + self.qo_bank = nn.Parameter(torch.empty(2 * num_layers, model_dim, model_dim)) + self.kv_bank = nn.Parameter(torch.empty(2 * num_layers, kv_dim, model_dim)) + self.mlp_up_bank = nn.Parameter(torch.empty(num_layers, mlp_dim, model_dim)) + self.mlp_down_bank = nn.Parameter(torch.empty(num_layers, model_dim, mlp_dim)) + self.blocks = nn.ModuleList( + [ + Block( + model_dim, + num_heads, + num_kv_heads, + mlp_mult, + rope_base, + qk_gain_init, + layer_idx=i, + ln_scale=ln_scale, + dtg=dtg, + gated_attention=gated_attention, + value_residual=value_residual, + ) + for i in range(num_layers) + ] + ) + if rope_dims > 0: + head_dim = model_dim // num_heads + for block in self.blocks: + block.attn.rope_dims = rope_dims + block.attn.rotary = Rotary(head_dim, base=rope_base, train_seq_len=1024, rope_dims=rope_dims) + self.ve_layer_indices = [int(x) for x in ve_layers.split(",") if x.strip()] if ve_enabled else [] + kv_dim_ve = self._ve_target_dim + if self.ve_layer_indices: + self.ve_shared = ValueEmbedding(vocab_size, ve_dim, kv_dim_ve) + self.ve_layer_scales = nn.ParameterList( + [nn.Parameter(torch.ones(1, dtype=torch.float32)) for _ in self.ve_layer_indices] + ) + else: + self.ve_shared = None + self.ve_layer_scales = nn.ParameterList() + self.value_embeds = nn.ModuleList() # keep empty for compat + 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.mtp_heads = nn.ModuleList( + [CastedLinear(model_dim, vocab_size, bias=False) for _ in range(mtp_num_heads)] + ) + for head in self.mtp_heads: + head._zero_init = True + if xsa_last_n > 0: + for i in range(max(0, num_layers - xsa_last_n), num_layers): + self.blocks[i].attn.use_xsa = True + # Depth recurrence: reuse specified layers for free depth (PR #1296, #1204) + self.recur_layers = [int(x) for x in recur_layers.split(",") if x.strip()] + self._recurrence_active = False + # Parallel residuals: split attn/MLP into separate lanes (PR #1296, #1204) + self.parallel_start_layer = parallel_start_layer + if 0 < parallel_start_layer < num_layers: + self.lane_merge = nn.Parameter(torch.tensor(0.5, dtype=torch.float32)) + else: + self.lane_merge = None + self._init_weights() + + def set_recurrence_active(self, active: bool) -> None: + self._recurrence_active = active + + def _get_virtual_layers(self) -> list[int]: + """Return virtual->physical block mapping. + When recurrence is active, the recur_layers are repeated once, + e.g. with num_layers=11 and recur_layers=[4,5]: + [0,1,2,3, 4,5, 4,5, 6,7,8,9,10] + When inactive: [0,1,2,...,num_layers-1] + """ + n = len(self.blocks) + if not self._recurrence_active or not self.recur_layers: + return list(range(n)) + virtual = [] + inserted = False + for i in range(n): + virtual.append(i) + if not inserted and i == self.recur_layers[-1]: + for rl in self.recur_layers: + virtual.append(rl) + inserted = True + return virtual + + 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) + n = self.num_layers + proj_scale = 1.0 / math.sqrt(2 * n) + # Init banks: orthogonal, with proj layers scaled down and out/down zero-init + for i in range(n): + nn.init.orthogonal_(self.qo_bank.data[i], gain=1.0) # Q + nn.init.zeros_(self.qo_bank.data[n + i]) # Out (zero init) + nn.init.orthogonal_(self.kv_bank.data[i], gain=1.0) # K + nn.init.orthogonal_(self.kv_bank.data[n + i], gain=1.0) # V + nn.init.orthogonal_(self.mlp_up_bank.data[i], gain=1.0) # MLP up + nn.init.zeros_(self.mlp_down_bank.data[i]) # MLP down (zero init) + # Scale proj layers (out_proj and mlp_down are "proj" layers) + self.qo_bank.data[n + i].mul_(proj_scale) + self.mlp_down_bank.data[i].mul_(proj_scale) + # Init remaining nn.Linear modules (bigram proj, mtp heads, lm_head) + for name, module in self.named_modules(): + if isinstance(module, nn.Linear): + if getattr(module, "_zero_init", False): + nn.init.zeros_(module.weight) + elif module.weight.ndim == 2 and module.weight.shape[0] >= 64 and module.weight.shape[1] >= 64: + nn.init.orthogonal_(module.weight, gain=1.0) + def _get_ve(self, layer_idx: int, input_ids: Tensor, ve_cache: dict | None = None) -> Tensor | None: + """Get value embedding for a specific layer using shared table + per-layer scale.""" + if self.ve_shared is None or layer_idx not in self.ve_layer_indices: + return None + if ve_cache is not None and 've' not in ve_cache: + ve_cache['ve'] = self.ve_shared(input_ids) + ve_base = ve_cache['ve'] if ve_cache is not None else self.ve_shared(input_ids) + ve_idx = self.ve_layer_indices.index(layer_idx) + return ve_base * self.ve_layer_scales[ve_idx].to(dtype=ve_base.dtype) + def _forward_layers(self, input_ids: Tensor) -> Tensor: + """Shared layer forward pass for forward() and forward_hidden(). + Handles depth recurrence, parallel residuals, U-Net skips, and parameter banks.""" + n = self.num_layers + x = self.tok_emb(input_ids) + if self.bigram is not None: + x = x + self.bigram(input_ids) + x = F.rms_norm(x, (x.size(-1),)) + x = self.smear(x) + x0 = x + v0 = None + skips: list[Tensor] = [] + ve_cache: dict = {} + + virtual_layers = self._get_virtual_layers() + num_virtual = len(virtual_layers) + num_enc = num_virtual // 2 + num_dec = num_virtual - num_enc + + # Determine parallel residuals threshold + parallel_phys = self.parallel_start_layer if self.lane_merge is not None else n + 1 + is_parallel = False + lane0 = None # attention lane + lane1 = None # MLP lane + + # Encoder phase + for vi in range(num_enc): + phys = virtual_layers[vi] + ve = self._get_ve(phys, input_ids, ve_cache) + x, raw_v = self.blocks[phys](x, x0, + self.qo_bank[phys], self.kv_bank[phys], self.kv_bank[n + phys], + self.qo_bank[n + phys], self.mlp_up_bank[phys], self.mlp_down_bank[phys], + v_embed=ve, v0=v0) + if v0 is None and raw_v is not None: + v0 = raw_v + skips.append(x) + + # Decoder phase with U-Net skip connections + parallel residuals + for vi in range(num_dec): + phys = virtual_layers[num_enc + vi] + # U-Net skip connection (only as many as we have skip weights) + if skips and vi < self.num_skip_weights: + sw = torch.sigmoid(self.skip_weights[vi]) if self.sigmoid_skips else self.skip_weights[vi] + skip_val = sw.to(dtype=x.dtype)[None, None, :] * skips.pop() + if is_parallel: + lane0 = lane0 + skip_val + else: + x = x + skip_val + + # Enter parallel mode if physical layer >= threshold + if phys >= parallel_phys and not is_parallel: + lane0 = x + lane1 = x + is_parallel = True + + if is_parallel: + block = self.blocks[phys] + ve = self._get_ve(phys, input_ids, ve_cache) + # Attention on lane0 + mix = block.resid_mix.to(dtype=lane0.dtype) + attn_in = mix[0][None, None, :] * lane0 + mix[1][None, None, :] * x0 + attn_out, _ = block.attn(block.attn_norm(attn_in) * block.ln_scale_factor, + self.qo_bank[phys], self.kv_bank[phys], self.kv_bank[n + phys], + self.qo_bank[n + phys], v_embed=ve, v0=v0) + lane0 = attn_in + block.attn_scale.to(dtype=attn_in.dtype)[None, None, :] * attn_out + # MLP on lane1 + mlp_in = block.mlp_norm(lane1) * block.ln_scale_factor + mlp_out = block.mlp(mlp_in, self.mlp_up_bank[phys], self.mlp_down_bank[phys]) + lane1 = lane1 + block.mlp_scale.to(dtype=lane1.dtype)[None, None, :] * mlp_out + else: + ve = self._get_ve(phys, input_ids, ve_cache) + x, _ = self.blocks[phys](x, x0, + self.qo_bank[phys], self.kv_bank[phys], self.kv_bank[n + phys], + self.qo_bank[n + phys], self.mlp_up_bank[phys], self.mlp_down_bank[phys], + v_embed=ve, v0=v0) + + # Merge parallel lanes + if is_parallel: + m = self.lane_merge.to(dtype=lane0.dtype) + x = m * lane0 + (1 - m) * lane1 + + return self.final_norm(x) + + def forward(self, input_ids: Tensor, target_ids: Tensor) -> Tensor: + x = self._forward_layers(input_ids) + x_flat = x.reshape(-1, x.size(-1)) + targets = target_ids.reshape(-1) + if self.tie_embeddings: + logits_proj = F.linear(x_flat, 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_flat) + logits = self.logit_softcap * torch.tanh(logits_proj / self.logit_softcap) + main_loss = F.cross_entropy(logits.float(), targets, reduction="mean") + if self.training and self.mtp_num_heads > 0 and self.mtp_loss_weight > 0.0: + _, seqlen, dim = x.shape + mtp_loss_sum = x.new_zeros(()) + mtp_loss_count = 0 + for k, mtp_head in enumerate(self.mtp_heads): + valid_t = seqlen - (k + 1) + if valid_t <= 0: + continue + mtp_hidden = x[:, :valid_t, :].reshape(-1, dim) + mtp_targets = target_ids[:, k + 1 :].reshape(-1) + mtp_logits_proj = mtp_head(mtp_hidden) + mtp_logits = self.logit_softcap * torch.tanh(mtp_logits_proj / self.logit_softcap) + mtp_loss_sum = mtp_loss_sum + F.cross_entropy(mtp_logits.float(), mtp_targets, reduction="mean") + mtp_loss_count += 1 + if mtp_loss_count > 0: + main_loss = main_loss + self.mtp_loss_weight * (mtp_loss_sum / mtp_loss_count) + return main_loss + + def forward_hidden(self, input_ids: Tensor) -> Tensor: + """Return final hidden states before lm_head (for SLOT).""" + return self._forward_layers(input_ids) + + def compute_logits(self, hidden_states: Tensor) -> Tensor: + """Project hidden states to logits with softcap.""" + if self.tie_embeddings: + lp = F.linear(hidden_states, self.tok_emb.weight) + else: + lp = self.lm_head(hidden_states) + return self.logit_softcap * torch.tanh(lp / self.logit_softcap) + + def forward_logits(self, input_ids: Tensor) -> Tensor: + """Return logits (bsz, seq_len, vocab) without computing loss.""" + return self.compute_logits(self.forward_hidden(input_ids)) + +# --- Sliding window evaluation --- + +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, + batch_seqs: int = 32, + eval_seq_len: int | None = None, +) -> tuple[float, float]: + """Sliding window evaluation: each token scored with maximum context.""" + seq_len = eval_seq_len or args.train_seq_len + total_tokens = val_tokens.numel() - 1 + window_starts = [ws for ws in range(0, total_tokens, stride) + if min(ws + seq_len, total_tokens) - ws >= 1] + total_windows = len(window_starts) + my_s = (total_windows * rank) // world_size + my_e = (total_windows * (rank + 1)) // world_size + my_windows = window_starts[my_s:my_e] + loss_sum = torch.zeros((), device=device, dtype=torch.float64) + token_count = torch.zeros((), device=device, dtype=torch.float64) + byte_count = torch.zeros((), device=device, dtype=torch.float64) + base_model.eval() + compiled_logits = torch.compile(base_model.forward_logits, dynamic=False, fullgraph=True) + with torch.inference_mode(): + for bi in range(0, len(my_windows), batch_seqs): + batch_ws = my_windows[bi:bi + batch_seqs] + bsz = len(batch_ws) + x_batch = torch.zeros(bsz, seq_len, dtype=torch.int64, device=device) + y_batch = torch.zeros(bsz, seq_len, dtype=torch.int64, device=device) + wlens: list[int] = [] + for i, ws in enumerate(batch_ws): + end = min(ws + seq_len, total_tokens) + wlen = end - ws + wlens.append(wlen) + chunk = val_tokens[ws:end + 1].to(dtype=torch.int64, device=device) + x_batch[i, :wlen] = chunk[:-1] + y_batch[i, :wlen] = chunk[1:] + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + logits = compiled_logits(x_batch) + nll = F.cross_entropy( + logits.reshape(-1, logits.size(-1)).float(), + y_batch.reshape(-1), + reduction="none", + ).reshape(bsz, seq_len) + for i, ws in enumerate(batch_ws): + wlen = wlens[i] + s = 0 if ws == 0 else max(wlen - stride, 0) + scored_nll = nll[i, s:wlen].to(torch.float64) + loss_sum += scored_nll.sum() + token_count += float(wlen - s) + tgt = y_batch[i, s:wlen] + prev = x_batch[i, s:wlen] + tb = base_bytes_lut[tgt].to(torch.float64) + tb += (has_leading_space_lut[tgt] & ~is_boundary_token_lut[prev]).to(torch.float64) + byte_count += tb.sum() + if dist.is_available() and dist.is_initialized(): + dist.all_reduce(loss_sum, op=dist.ReduceOp.SUM) + dist.all_reduce(token_count, op=dist.ReduceOp.SUM) + dist.all_reduce(byte_count, op=dist.ReduceOp.SUM) + val_loss = (loss_sum / token_count).item() + bits_per_token = val_loss / math.log(2.0) + tokens_per_byte = token_count.item() / byte_count.item() + base_model.train() + return val_loss, bits_per_token * tokens_per_byte + + +def generate_autoregressive_calib(model, device, num_seqs=64, seq_len=2048, + vocab_size=1024, temperature=0.8, batch_size=8, seed=42): + """Generate sequences autoregressively from the model for GPTQ calibration. + No external data accessed — fully self-contained.""" + model.eval() + rng = torch.Generator(device=device) + rng.manual_seed(seed) + all_tokens = [] + with torch.inference_mode(), torch.autocast(device_type="cuda", dtype=torch.bfloat16): + for batch_start in range(0, num_seqs, batch_size): + bs = min(batch_size, num_seqs - batch_start) + tokens = torch.randint(0, vocab_size, (bs, 1), device=device, generator=rng) + for pos in range(seq_len - 1): + logits = model.forward_logits(tokens) + next_logit = logits[:, -1, :] + probs = torch.softmax(next_logit / temperature, dim=-1) + next_tok = torch.multinomial(probs, 1, generator=rng) + tokens = torch.cat([tokens, next_tok], dim=1) + for i in range(bs): + all_tokens.append(tokens[i:i+1]) + return all_tokens + + +def collect_hessians_from_tokens(hessian_model, token_seqs, device): + """Collect H = X^T X from pre-generated token sequences.""" + hessians = {} + hooks = [] + for name, module in hessian_model.named_modules(): + if isinstance(module, CastedLinear): + param_name = name + ".weight" + cols = module.weight.shape[1] + hessians[param_name] = torch.zeros(cols, cols, dtype=torch.float32, device='cpu') + def make_hook(pname): + def hook_fn(module, input, output): + x = input[0].detach().float() + if x.ndim == 3: + x = x.reshape(-1, x.shape[-1]) + hessians[pname] += (x.T @ x).cpu() + return hook_fn + h = module.register_forward_hook(make_hook(param_name)) + hooks.append(h) + hessian_model.eval() + with torch.inference_mode(), torch.autocast(device_type="cuda", dtype=torch.bfloat16): + for seq in token_seqs: + x = seq[:, :-1].to(device) + y = seq[:, 1:].to(device) + hessian_model(x, y) + for h in hooks: + h.remove() + num_batches = len(token_seqs) + for name in hessians: + H = hessians[name] + H /= num_batches + damp = 0.01 * torch.diag(H).mean().clamp_min(1e-6) + H += damp * torch.eye(H.shape[0]) + hessians[name] = H + return hessians + + +# --- Brotli + byte-shuffle compression --- + +def _byte_shuffle(data: bytes) -> bytes: + """Reorder bytes by significance position for better compression. + Groups byte 0 of all elements, then byte 1, etc.""" + arr = np.frombuffer(data, dtype=np.uint8) + # Treat as 4-byte groups (torch.save uses 4-byte aligned data) + pad = (4 - len(arr) % 4) % 4 + if pad: + arr = np.concatenate([arr, np.zeros(pad, dtype=np.uint8)]) + grouped = arr.reshape(-1, 4) + shuffled = grouped.T.ravel() + return bytes(shuffled) + pad.to_bytes(1, 'little') + +def _byte_unshuffle(data: bytes) -> bytes: + """Reverse byte-shuffle.""" + pad = data[-1] + arr = np.frombuffer(data[:-1], dtype=np.uint8) + grouped = arr.reshape(4, -1) + unshuffled = grouped.T.ravel() + if pad: + unshuffled = unshuffled[:-pad] + return bytes(unshuffled) + +def compress_artifact(raw: bytes, method: str = "auto") -> bytes: + """Compress artifact. Methods: 'rans', 'brotli', 'lzma', 'auto'. + 'auto' tries rANS first (best for i.i.d. quantized weights), falls back to Brotli/LZMA.""" + if method == "auto": + # Try rANS, compare with Brotli, pick smaller + rans_out = _rans_compress(raw) + if _HAS_BROTLI: + shuffled = _byte_shuffle(raw) + brotli_out = b'BROT' + brotli.compress(shuffled, quality=11) + if len(rans_out) <= len(brotli_out): + return rans_out + return brotli_out + return rans_out + elif method == "rans": + return _rans_compress(raw) + elif method == "brotli": + if not _HAS_BROTLI: + raise ImportError("brotli required") + shuffled = _byte_shuffle(raw) + return b'BROT' + brotli.compress(shuffled, quality=11) + else: + return b'LZMA' + lzma.compress(raw, preset=9) + +def decompress_artifact(blob: bytes) -> bytes: + """Decompress artifact (auto-detects format from 4-byte tag).""" + tag = blob[:4] + payload = blob[4:] + if tag == b'RANS': + return _rans_decompress(payload) + elif tag == b'BROT': + if not _HAS_BROTLI: + raise ImportError("brotli required to decompress this artifact") + return _byte_unshuffle(brotli.decompress(payload)) + elif tag == b'LZMA': + return lzma.decompress(payload) + else: + # Legacy: try LZMA without tag + return lzma.decompress(blob) + + +# --- rANS entropy coding --- +# Near-Shannon-optimal compression for i.i.d. quantized weights. +# 12-27% better than Brotli on peaked int6 distributions. + +_RANS_L = 1 << 23 # Renormalization lower bound +_RANS_SCALE = 1 << 16 # Frequency table precision + +def _build_freq_table(data: np.ndarray) -> dict[int, int]: + """Build rANS frequency table from data, scaled to _RANS_SCALE.""" + counts = {} + for v in data.ravel(): + v = int(v) + counts[v] = counts.get(v, 0) + 1 + total = sum(counts.values()) + symbols = sorted(counts.keys()) + freq = {} + for s in symbols: + freq[s] = max(1, int(counts[s] / total * _RANS_SCALE)) + # Adjust to sum exactly to _RANS_SCALE + diff = _RANS_SCALE - sum(freq.values()) + most_common = max(freq, key=freq.get) + freq[most_common] += diff + return freq + +def _rans_encode(data: np.ndarray, freq: dict[int, int]) -> bytes: + """rANS encode array of symbols. Returns compressed bytes. + Uses 32-bit state with byte-level output (radix=256). + Invariant: state in [RANS_L, RANS_L << 8) after each symbol.""" + symbols = sorted(freq.keys()) + cum = {} + c = 0 + for s in symbols: + cum[s] = (c, freq[s]) + c += freq[s] + + state = _RANS_L + out = bytearray() + flat = data.ravel().tolist() + + for s in reversed(flat): + start, f = cum[s] + # Renormalize: emit bytes to keep state bounded after encode step. + # After encode, state_new = (state/f)*M + (state%f) + start + # For state_new < RANS_L << 8, need state < (RANS_L >> 16) << 8 * f + max_state = ((_RANS_L >> 16) << 8) * f + while state >= max_state: + out.append(state & 0xFF) + state >>= 8 + state = ((state // f) << 16) + (state % f) + start + + # Flush state (4 bytes) + for _ in range(4): + out.append(state & 0xFF) + state >>= 8 + return bytes(out) + +def _rans_decode(data: bytes, num_symbols: int, freq: dict[int, int]) -> np.ndarray: + """rANS decode bytes back to symbol array.""" + symbols = sorted(freq.keys()) + cum = {} + c = 0 + for s in symbols: + cum[s] = (c, freq[s]) + c += freq[s] + + # Build decode LUT + decode_lut = [0] * _RANS_SCALE + for s in symbols: + start, f = cum[s] + for i in range(start, start + f): + decode_lut[i] = s + + pos = len(data) - 1 + state = 0 + for _ in range(4): + state = (state << 8) | data[pos] + pos -= 1 + + result = [] + for _ in range(num_symbols): + slot = state & (_RANS_SCALE - 1) + s = decode_lut[slot] + start, f = cum[s] + state = f * (state >> 16) + slot - start + while state < _RANS_L: + state = (state << 8) | data[pos] + pos -= 1 + result.append(s) + + return np.array(result, dtype=np.int8) + +def _rans_compress(raw: bytes) -> bytes: + """Compress raw artifact bytes with rANS. Format: + b'RANS' + [4: num_bytes] + [4: num_symbols] + [2: freq_table_len] + + freq_table + rans_data""" + arr = np.frombuffer(raw, dtype=np.uint8) + # Treat as signed int8 for better frequency distribution (weights are signed) + signed = arr.view(np.int8) + freq = _build_freq_table(signed) + rans_data = _rans_encode(signed, freq) + + # Serialize: header + freq table + rans data + symbols = sorted(freq.keys()) + header = struct.pack(' bytes: + """Decompress rANS payload back to raw bytes.""" + offset = 0 + orig_size, num_symbols, num_freq = struct.unpack_from(' str: + if "tok_emb" in name or "lm_head" in name: + return "embed" + if ".mlp." in name: + return "mlp" + if ".attn." in name or (".proj." in name and ".mlp." not in name): + return "attn" + return "other" +def _sdclip_scale(t32: Tensor, clip_range: int, k: float) -> Tensor: + """SDClip (PR #1394): clip = k * std(row). Entropy ≈ b - log2(k) + const. + Larger k → values cluster at center → lower entropy → better compression.""" + if t32.ndim == 2: + row_std = t32.std(dim=1) + row_clip = (k * row_std).clamp_min(1e-7) + return (row_clip / clip_range).clamp_min(1.0 / clip_range).to(torch.float16) + std = t32.std().item() + clip = max(k * std, 1e-7) + return torch.tensor(clip / clip_range, dtype=torch.float16) + +def quantize_int6_per_row(t: Tensor, clip_range: int = 31, sdclip_k: float = 0.0) -> tuple[Tensor, Tensor]: + t32 = t.float() + if sdclip_k > 0: + s = _sdclip_scale(t32, clip_range, sdclip_k) + if t32.ndim == 2: + q = torch.clamp(torch.round(t32 / s.float()[:, None]), -clip_range, clip_range).to(torch.int8) + else: + q = torch.clamp(torch.round(t32 / s.float()), -clip_range, clip_range).to(torch.int8) + return q, s + if t32.ndim == 2: + best_q, best_s, best_err = None, None, float('inf') + for pct in [0.9990, 0.9995, 0.9999, 0.99999, 1.0]: + if pct < 1.0: + row_clip = torch.quantile(t32.abs(), pct, dim=1) + else: + row_clip = t32.abs().amax(dim=1) + s = (row_clip / clip_range).clamp_min(1.0 / clip_range).to(torch.float16) + q = torch.clamp(torch.round(t32 / s.float()[:, None]), -clip_range, clip_range).to(torch.int8) + recon = q.float() * s.float()[:, None] + err = (t32 - recon).pow(2).mean().item() + if err < best_err: + best_q, best_s, best_err = q, s, err + return best_q, best_s + amax = t32.abs().max().item() + scale = torch.tensor(amax / clip_range if amax > 0 else 1.0, dtype=torch.float16) + q = torch.clamp(torch.round(t32 / scale.float()), -clip_range, clip_range).to(torch.int8) + return q, scale + +def quantize_int6_gptq(weight, hessian=None, clip_range=31, block_size=128, sdclip_k: float = 0.0): + """Full GPTQ: Hessian-aware int6 quantization with Cholesky error compensation. + If sdclip_k > 0, uses SDClip (one pass) instead of percentile search (5 passes). + If hessian is None, falls back to percentile/SDClip per-row.""" + t32 = weight.float() + if t32.ndim != 2 or hessian is None: + if sdclip_k > 0: + return quantize_int6_per_row(t32, clip_range, sdclip_k=sdclip_k) + return _quantize_int6_percentile(t32, clip_range) + rows, cols = t32.shape + H = hessian.float().clone() + dead = torch.diag(H) == 0 + H[dead, dead] = 1 + damp = 0.01 * torch.mean(torch.diag(H)) + H[torch.arange(cols), torch.arange(cols)] += damp + perm = torch.argsort(torch.diag(H), descending=True) + inv_perm = torch.argsort(perm) + W = t32[:, perm].clone() + W[:, dead[perm]] = 0 + H = H[perm][:, perm] + Hinv = torch.linalg.cholesky(H) + Hinv = torch.cholesky_inverse(Hinv) + Hinv = torch.linalg.cholesky(Hinv, upper=True) + + def _gptq_one_pass(s: Tensor) -> tuple[Tensor, Tensor, float]: + sf = s.float() + Q = torch.zeros_like(W, dtype=torch.int8) + W_work = W.clone() + for i1 in range(0, cols, block_size): + i2 = min(i1 + block_size, cols) + count = i2 - i1 + W1 = W_work[:, i1:i2].clone() + Q1 = torch.zeros(rows, count, dtype=torch.int8) + Err1 = torch.zeros(rows, count) + Hinv1 = Hinv[i1:i2, i1:i2] + for i in range(count): + w = W1[:, i] + d = Hinv1[i, i] + q = torch.clamp(torch.round(w / sf), -clip_range, clip_range).to(torch.int8) + Q1[:, i] = q + err = (w - q.float() * sf) / d + W1[:, i:] -= err.unsqueeze(1) * Hinv1[i, i:].unsqueeze(0) + Err1[:, i] = err + Q[:, i1:i2] = Q1 + if i2 < cols: + W_work[:, i2:] -= Err1 @ Hinv[i1:i2, i2:] + recon = Q.float() * sf[:, None] + mse = (W - recon).pow(2).mean().item() + return Q, s, mse + + if sdclip_k > 0: + # SDClip: single pass with k * std(row) + s = _sdclip_scale(t32, clip_range, sdclip_k) + best_q, best_scale, _ = _gptq_one_pass(s) + else: + # Legacy: percentile search (5 passes) + best_q = None; best_scale = None; best_err = float('inf') + for pct in [0.9990, 0.9995, 0.9999, 0.99999, 1.0]: + if pct < 1.0: + row_clip = torch.quantile(t32.abs(), pct, dim=1) + else: + row_clip = t32.abs().amax(dim=1) + s = (row_clip / clip_range).clamp_min(1.0 / clip_range).to(torch.float16) + Q, sc, mse = _gptq_one_pass(s) + if mse < best_err: + best_q, best_scale, best_err = Q, sc, mse + best_q = best_q[:, inv_perm] + return best_q, best_scale + +def _quantize_int6_percentile(t32, clip_range=31): + """Fallback: percentile search (for 1D or no-Hessian cases).""" + if t32.ndim == 2: + best_q, best_s, best_err = None, None, float('inf') + for pct in [0.9990, 0.9995, 0.9999, 0.99999, 1.0]: + if pct < 1.0: + row_clip = torch.quantile(t32.abs(), pct, dim=1) + else: + row_clip = t32.abs().amax(dim=1) + s = (row_clip / clip_range).clamp_min(1.0 / clip_range).to(torch.float16) + q = torch.clamp(torch.round(t32 / s.float()[:, None]), -clip_range, clip_range).to(torch.int8) + recon = q.float() * s.float()[:, None] + err = (t32 - recon).pow(2).mean().item() + if err < best_err: + best_q, best_s, best_err = q, s, err + return best_q, best_s + amax = t32.abs().max().item() + scale = torch.tensor(amax / clip_range if amax > 0 else 1.0, dtype=torch.float16) + q = torch.clamp(torch.round(t32 / scale.float()), -clip_range, clip_range).to(torch.int8) + return q, scale + +def _unbank_state_dict(sd: dict[str, Tensor], num_layers: int) -> dict[str, Tensor]: + """Convert 3D bank tensors into individual 2D tensors with standard names.""" + out: dict[str, Tensor] = {} + n = num_layers + for name, tensor in sd.items(): + if name == "qo_bank": + for i in range(n): + out[f"blocks.{i}.attn.c_q.weight"] = tensor[i] + out[f"blocks.{i}.attn.proj.weight"] = tensor[n + i] + elif name == "kv_bank": + for i in range(n): + out[f"blocks.{i}.attn.c_k.weight"] = tensor[i] + out[f"blocks.{i}.attn.c_v.weight"] = tensor[n + i] + elif name == "mlp_up_bank": + for i in range(n): + out[f"blocks.{i}.mlp.fc.weight"] = tensor[i] + elif name == "mlp_down_bank": + for i in range(n): + out[f"blocks.{i}.mlp.proj.weight"] = tensor[i] + else: + out[name] = tensor + return out + +def _rebank_state_dict(sd: dict[str, Tensor], num_layers: int, template_sd: dict[str, Tensor]) -> dict[str, Tensor]: + """Convert individual 2D tensors back into 3D bank tensors.""" + out: dict[str, Tensor] = {} + n = num_layers + # Reconstruct banks from individual weight keys + qo_slices = [None] * (2 * n) + kv_slices = [None] * (2 * n) + up_slices = [None] * n + down_slices = [None] * n + consumed = set() + for i in range(n): + qk = f"blocks.{i}.attn.c_q.weight" + if qk in sd: + qo_slices[i] = sd[qk] + consumed.add(qk) + ok = f"blocks.{i}.attn.proj.weight" + if ok in sd: + qo_slices[n + i] = sd[ok] + consumed.add(ok) + kk = f"blocks.{i}.attn.c_k.weight" + if kk in sd: + kv_slices[i] = sd[kk] + consumed.add(kk) + vk = f"blocks.{i}.attn.c_v.weight" + if vk in sd: + kv_slices[n + i] = sd[vk] + consumed.add(vk) + fk = f"blocks.{i}.mlp.fc.weight" + if fk in sd: + up_slices[i] = sd[fk] + consumed.add(fk) + dk = f"blocks.{i}.mlp.proj.weight" + if dk in sd: + down_slices[i] = sd[dk] + consumed.add(dk) + out["qo_bank"] = torch.stack(qo_slices).to(dtype=template_sd["qo_bank"].dtype) + out["kv_bank"] = torch.stack(kv_slices).to(dtype=template_sd["kv_bank"].dtype) + out["mlp_up_bank"] = torch.stack(up_slices).to(dtype=template_sd["mlp_up_bank"].dtype) + out["mlp_down_bank"] = torch.stack(down_slices).to(dtype=template_sd["mlp_down_bank"].dtype) + for name, tensor in sd.items(): + if name not in consumed: + out[name] = tensor + return out + +# --- Non-banked model for Hessian collection --- +# This mirrors the unbanked state dict keys: blocks.{i}.attn.c_q/c_k/c_v/proj, blocks.{i}.mlp.fc/proj + +class _HessianAttn(nn.Module): + """Non-banked attention with CastedLinear layers for Hessian hooks.""" + def __init__(self, dim, num_heads, num_kv_heads, rope_base, qk_gain_init): + super().__init__() + self.num_heads, self.num_kv_heads = num_heads, num_kv_heads + self.head_dim = dim // num_heads + kv_dim = 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.q_gain = nn.Parameter(torch.full((num_heads,), qk_gain_init, dtype=torch.float32)) + self.rope_dims = 0 + self.rotary = Rotary(self.head_dim, base=rope_base, train_seq_len=1024) + self.use_xsa = False + def _xsa_efficient(self, y, v): + B, T, H, D = y.shape; Hkv = v.size(-2); group = H // Hkv + y_g = y.reshape(B, T, Hkv, group, D) + vn = F.normalize(v, dim=-1).unsqueeze(-2) + proj = (y_g * vn).sum(dim=-1, keepdim=True) * vn + return (y_g - proj).reshape(B, T, H, D) + def forward(self, x, v_embed=None): + bsz, seqlen, dim = x.shape + q = self.c_q(x).reshape(bsz, seqlen, self.num_heads, self.head_dim) + k = self.c_k(x).reshape(bsz, seqlen, self.num_kv_heads, self.head_dim) + v = self.c_v(x) + if v_embed is not None: + v = v + v_embed + v = v.reshape(bsz, seqlen, self.num_kv_heads, self.head_dim) + 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, self.rope_dims) + k = apply_rotary_emb(k, cos, sin, self.rope_dims) + q = q * self.q_gain.to(dtype=q.dtype)[None, None, :, None] + y = flash_attn_3_func(q, k, v, causal=True) + if self.use_xsa: + y = self._xsa_efficient(y, v) + return self.proj(y.reshape(bsz, seqlen, dim)) + +class _HessianMLP(nn.Module): + """Non-banked MLP with CastedLinear layers for Hessian hooks.""" + def __init__(self, dim, mlp_mult): + super().__init__() + self.fc = CastedLinear(dim, int(mlp_mult * dim), bias=False) + self.proj = CastedLinear(int(mlp_mult * dim), dim, bias=False) + def forward(self, x): + return self.proj(F.leaky_relu(self.fc(x), negative_slope=0.5).square()) + +class _HessianBlock(nn.Module): + def __init__(self, dim, num_heads, num_kv_heads, mlp_mult, rope_base, qk_gain_init, layer_idx=0, ln_scale=False): + super().__init__() + self.attn_norm = RMSNorm() + self.mlp_norm = RMSNorm() + self.attn = _HessianAttn(dim, num_heads, num_kv_heads, rope_base, qk_gain_init) + self.mlp = _HessianMLP(dim, mlp_mult) + 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_factor = 1.0 / math.sqrt(layer_idx + 1) if ln_scale else 1.0 + def forward(self, x, x0, v_embed=None): + mix = self.resid_mix.to(dtype=x.dtype) + x_in = mix[0][None, None, :] * x + mix[1][None, None, :] * x0 + attn_out = self.attn(self.attn_norm(x_in) * self.ln_scale_factor, v_embed=v_embed) + x_out = x_in + self.attn_scale.to(dtype=x_in.dtype)[None, None, :] * attn_out + x_out = x_out + self.mlp_scale.to(dtype=x_out.dtype)[None, None, :] * self.mlp(self.mlp_norm(x_out) * self.ln_scale_factor) + return x_out + +class _HessianGPT(nn.Module): + """Non-banked GPT model matching unbanked state dict keys for Hessian collection.""" + def __init__(self, vocab_size, num_layers, model_dim, num_heads, num_kv_heads, + mlp_mult, tie_embeddings, logit_softcap, rope_base, qk_gain_init, + bigram_vocab_size=0, bigram_dim=128, xsa_last_n=0, + rope_dims=0, ln_scale=False, + ve_enabled=False, ve_dim=128, ve_layers="9,10"): + super().__init__() + self.tie_embeddings = tie_embeddings + self.logit_softcap = logit_softcap + self.num_layers = num_layers + self.tok_emb = nn.Embedding(vocab_size, model_dim) + self.bigram = BigramHashEmbedding(bigram_vocab_size, bigram_dim, model_dim, trigram=bool(int(os.environ.get("TRIGRAM", "0")))) if bigram_vocab_size > 0 else None + self.smear = SmearGate(model_dim) + self.num_encoder_layers = num_layers // 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.sigmoid_skips = bool(int(os.environ.get("SIGMOID_SKIPS", "0"))) + skip_init = 2.2 if self.sigmoid_skips else 1.0 + self.skip_weights = nn.Parameter(torch.full((self.num_skip_weights, model_dim), skip_init, dtype=torch.float32)) + self.blocks = nn.ModuleList([ + _HessianBlock(model_dim, num_heads, num_kv_heads, mlp_mult, rope_base, qk_gain_init, + layer_idx=i, ln_scale=ln_scale) + for i in range(num_layers) + ]) + if rope_dims > 0: + head_dim = model_dim // num_heads + for block in self.blocks: + block.attn.rope_dims = rope_dims + block.attn.rotary = Rotary(head_dim, base=rope_base, train_seq_len=1024, rope_dims=rope_dims) + if xsa_last_n > 0: + for i in range(max(0, num_layers - xsa_last_n), num_layers): + self.blocks[i].attn.use_xsa = True + kv_dim = num_kv_heads * (model_dim // num_heads) + self.ve_layer_indices = [int(x) for x in ve_layers.split(",") if x.strip()] if ve_enabled else [] + if self.ve_layer_indices: + self.ve_shared = ValueEmbedding(vocab_size, ve_dim, kv_dim) + self.ve_layer_scales = nn.ParameterList([nn.Parameter(torch.ones(1, dtype=torch.float32)) for _ in self.ve_layer_indices]) + else: + self.ve_shared = None + self.ve_layer_scales = nn.ParameterList() + self.final_norm = RMSNorm() + self.lm_head = None if tie_embeddings else CastedLinear(model_dim, vocab_size, bias=False) + def _get_ve(self, layer_idx, input_ids, ve_cache): + if self.ve_shared is None or layer_idx not in self.ve_layer_indices: + return None + if 've' not in ve_cache: + ve_cache['ve'] = self.ve_shared(input_ids) + ve_idx = self.ve_layer_indices.index(layer_idx) + return ve_cache['ve'] * self.ve_layer_scales[ve_idx].to(dtype=ve_cache['ve'].dtype) + def forward(self, input_ids, target_ids): + x = self.tok_emb(input_ids) + if self.bigram is not None: + x = x + self.bigram(input_ids) + x = F.rms_norm(x, (x.size(-1),)) + x = self.smear(x) + x0 = x + skips = [] + ve_cache = {} + for i in range(self.num_encoder_layers): + ve = self._get_ve(i, input_ids, ve_cache) + x = self.blocks[i](x, x0, v_embed=ve) + skips.append(x) + for i in range(self.num_decoder_layers): + bi = self.num_encoder_layers + i + if skips: + sw = torch.sigmoid(self.skip_weights[i]) if self.sigmoid_skips else self.skip_weights[i] + x = x + sw.to(dtype=x.dtype)[None, None, :] * skips.pop() + ve = self._get_ve(bi, input_ids, ve_cache) + x = self.blocks[bi](x, x0, v_embed=ve) + x = self.final_norm(x) + x_flat = x.reshape(-1, x.size(-1)) + targets = target_ids.reshape(-1) + logits_proj = F.linear(x_flat, self.tok_emb.weight) if self.tie_embeddings else self.lm_head(x_flat) + logits = self.logit_softcap * torch.tanh(logits_proj / self.logit_softcap) + return F.cross_entropy(logits.float(), targets, reduction="mean") + +def collect_hessians(hessian_model, train_loader, args, device, grad_accum_steps, num_batches=256): + """Run calibration batches through a non-banked model, collecting H = X^T X for each CastedLinear.""" + hessians = {} + hooks = [] + for name, module in hessian_model.named_modules(): + if isinstance(module, CastedLinear): + param_name = name + ".weight" + cols = module.weight.shape[1] + hessians[param_name] = torch.zeros(cols, cols, dtype=torch.float32, device='cpu') + def make_hook(pname): + def hook_fn(module, input, output): + x = input[0].detach().float() + if x.ndim == 3: + x = x.reshape(-1, x.shape[-1]) + hessians[pname] += (x.T @ x).cpu() + return hook_fn + h = module.register_forward_hook(make_hook(param_name)) + hooks.append(h) + hessian_model.eval() + with torch.inference_mode(), torch.autocast(device_type="cuda", dtype=torch.bfloat16): + for _ in range(num_batches): + x, y = train_loader.next_batch(args.train_batch_tokens, args.train_seq_len, grad_accum_steps) + hessian_model(x, y) + for h in hooks: + h.remove() + for name in hessians: + H = hessians[name] + H /= num_batches + damp = 0.01 * torch.diag(H).mean().clamp_min(1e-6) + H += damp * torch.eye(H.shape[0]) + hessians[name] = H + hessian_model.train() + return hessians + +def mixed_quantize_int6(state_dict: dict[str, Tensor], int6_cats: set[str], hessians: dict[str, Tensor] | None = None, sdclip_k: float = 0.0, sdclip_k_int8: float = 20.0): + num_layers_total = max( + (int(k.split(".")[1]) for k in state_dict if k.startswith("blocks.")), + default=0, + ) + 1 + late_k_layers = set(range(num_layers_total - 2, num_layers_total)) + result: dict[str, Tensor] = {} + meta: dict[str, object] = {} + for name, tensor in state_dict.items(): + t = tensor.detach().cpu().contiguous() + cat = _classify_param(name) + if not t.is_floating_point() or t.numel() <= 65536: + result[name] = t.to(torch.float16) if t.is_floating_point() else t + meta[name] = "passthrough" + continue + if any(p in name for p in CONTROL_TENSOR_NAME_PATTERNS): + result[name] = t.float() + meta[name] = "passthrough_ctrl" + continue + if cat in int6_cats and t.ndim >= 1: + cr = 31 # int6 for all weights + H = hessians.get(name) if hessians else None + if H is not None: + q, s = quantize_int6_gptq(t, hessian=H, clip_range=cr, sdclip_k=sdclip_k) + else: + q, s = quantize_int6_per_row(t, clip_range=cr, sdclip_k=sdclip_k) + result[name + ".q"] = q + result[name + ".scale"] = s + meta[name] = {"type": "int6"} + else: + # Embeddings and other large tensors: int8 with SDClip (PR #1493) + if sdclip_k_int8 > 0 and t.ndim == 2 and cat == "embed": + s = _sdclip_scale(t.float(), 127, sdclip_k_int8) + if s.ndim > 0: + q = torch.clamp(torch.round(t.float() / s.float()[:, None]), -127, 127).to(torch.int8) + else: + q = torch.clamp(torch.round(t.float() / s.float()), -127, 127).to(torch.int8) + else: + q, s = quantize_float_tensor(t) + result[name + ".q"] = q + result[name + ".scale"] = s + meta[name] = {"type": "int8"} + return result, meta +def dequantize_mixed_int6(result: dict[str, Tensor], meta: dict[str, object], + template_sd: dict[str, Tensor]) -> dict[str, Tensor]: + out: dict[str, Tensor] = {} + for name, orig in template_sd.items(): + info = meta.get(name) + if info is None: + continue + orig_dtype = orig.dtype + if info in ("passthrough", "passthrough_ctrl", "passthrough_fp16"): + t = result[name] + if t.dtype == torch.float16 and orig_dtype in (torch.float32, torch.bfloat16): + t = t.to(orig_dtype) + out[name] = t + continue + q, s = result[name + ".q"], result[name + ".scale"] + if s.ndim > 0: + out[name] = (q.float() * s.float().view(q.shape[0], *([1] * (q.ndim - 1)))).to(orig_dtype) + else: + out[name] = (q.float() * float(s.item())).to(orig_dtype) + return out + +# --- Training --- + +def main() -> None: + code = Path(__file__).read_text(encoding="utf-8") + args = Hyperparameters() + # zeropower_via_newtonschulz5 runs eagerly with bmm -- do NOT compile + 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 + 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(False) + 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) + 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"))) + effective_eval_seq_len = args.eval_seq_len if args.eval_seq_len > 0 else args.train_seq_len + val_seq_len = max(args.train_seq_len, effective_eval_seq_len) + val_tokens = load_validation_tokens(args.val_files, val_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}") + CastedLinear._qat_enabled = args.qat_enabled + 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, + mtp_num_heads=args.mtp_num_heads, + mtp_loss_weight=args.mtp_loss_weight, + bigram_vocab_size=args.bigram_vocab_size, + bigram_dim=args.bigram_dim, + xsa_last_n=args.xsa_last_n, + rope_dims=args.rope_dims, + ln_scale=args.ln_scale, + dtg=args.dtg_enabled, + ve_enabled=args.ve_enabled, + ve_dim=args.ve_dim, + ve_layers=args.ve_layers, + gated_attention=args.gated_attention, + value_residual=args.value_residual, + recur_layers=args.recur_layers, + parallel_start_layer=args.parallel_start_layer, + ).to(device).bfloat16() + # Banks stay FP32 (like CastedLinear weights), cast to BF16 in forward + base_model.qo_bank.data = base_model.qo_bank.data.float() + base_model.kv_bank.data = base_model.kv_bank.data.float() + base_model.mlp_up_bank.data = base_model.mlp_up_bank.data.float() + base_model.mlp_down_bank.data = base_model.mlp_down_bank.data.float() + for module in base_model.modules(): + if isinstance(module, CastedLinear): + module.float() + restore_low_dim_params_to_fp32(base_model) + # No DDP -- Parallel Muon handles bank grad communication via reduce-scatter, + # and non-bank grads are manually all-reduced before Adam steps. + # Increase dynamo cache for depth recurrence (forward changes at recur_start_step) + torch._dynamo.config.cache_size_limit = 32 + compiled_model = torch.compile(base_model, dynamic=False, fullgraph=True) + model = compiled_model + + # Optimizer split: + # - 4 parameter banks -> Muon (batched Newton-Schulz) + # - token embedding -> Adam + # - scalars/control tensors -> Adam + # - bigram proj, mtp heads, VE proj -> Adam (small matrix params not worth banking) + matrix_params = [ + base_model.qo_bank, base_model.kv_bank, + base_model.mlp_up_bank, base_model.mlp_down_bank, + ] + block_named_params = list(base_model.blocks.named_parameters()) + 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) + if base_model.bigram is not None: + scalar_params.append(base_model.bigram.scale) + token_lr = args.tied_embed_lr if args.tie_embeddings else args.embed_lr + tok_params = [{"params": [base_model.tok_emb.weight], "lr": token_lr, "base_lr": token_lr}] + if base_model.bigram is not None: + tok_params.append({"params": [base_model.bigram.embed.weight], "lr": token_lr, "base_lr": token_lr}) + if base_model.bigram.proj is not None: + scalar_params.append(base_model.bigram.proj.weight) + if base_model.ve_shared is not None: + tok_params.append({"params": [base_model.ve_shared.embed.weight], "lr": token_lr, "base_lr": token_lr}) + if base_model.ve_shared.proj is not None: + scalar_params.append(base_model.ve_shared.proj.weight) + scalar_params.append(base_model.ve_shared.scale) + for s in base_model.ve_layer_scales: + scalar_params.append(s) + # Parallel residuals lane merge (learnable scalar) + if base_model.lane_merge is not None: + scalar_params.append(base_model.lane_merge) + optimizer_tok = torch.optim.AdamW( + tok_params, + betas=(args.beta1, args.beta2), + eps=args.adam_eps, + weight_decay=args.embed_wd, # PR #1493: embed WD = 0.095 (higher than adam_wd) + fused=True, + ) + optimizer_muon = Muon( + matrix_params, + lr=args.matrix_lr, + momentum=args.muon_momentum, + backend_steps=args.muon_backend_steps, + weight_decay=args.muon_wd, + muoneq_r=args.muoneq_r, + ) + 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, + ) + # Non-bank params that need manual all-reduce (replicated across GPUs) + replicated_params = list(optimizer_tok.param_groups[0]["params"]) + for pg in optimizer_tok.param_groups[1:]: + replicated_params.extend(pg["params"]) + replicated_params.extend(scalar_params) + + optimizer_head = None + 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, + ) + replicated_params.append(base_model.lm_head.weight) + optimizers: list[torch.optim.Optimizer] = [optimizer_tok, optimizer_muon, optimizer_scalar] + if optimizer_head is not None: + optimizers.append(optimizer_head) + n_params = sum(p.numel() for p in base_model.parameters()) + mtp_params = sum(p.numel() for p in base_model.mtp_heads.parameters()) + log0(f"model_params:{n_params}") + log0(f"mtp_num_heads:{args.mtp_num_heads} mtp_loss_weight:{args.mtp_loss_weight} mtp_params:{mtp_params}") + xsa_layers = [i for i, b in enumerate(base_model.blocks) if b.attn.use_xsa] + log0(f"XSA:last_{args.xsa_last_n} active_layers:{xsa_layers}") + 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(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}") + 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: + """Linear LR decay to lr_floor (default 0). PR #1395: letting weights fully settle + before GPTQ reduces quantization gap by 61% (0.038 -> 0.014 BPB).""" + if args.warmdown_frac <= 0: + return 1.0 + # Determine progress fraction + if max_wallclock_ms is not None and max_wallclock_ms > 0: + frac = elapsed_ms / max_wallclock_ms + elif args.iterations > 0: + frac = step / args.iterations + else: + return 1.0 + warmdown_start = 1.0 - args.warmdown_frac + if frac <= warmdown_start: + return 1.0 + # Linear decay from 1.0 to lr_floor over the warmdown period + progress = (frac - warmdown_start) / args.warmdown_frac + return args.lr_floor + (1.0 - args.lr_floor) * max(1.0 - progress, 0.0) + if 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): + 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() + # All-reduce all grads for warmup (simple, not optimized) + if distributed: + for p in base_model.parameters(): + if p.grad is not None: + dist.all_reduce(p.grad, op=dist.ReduceOp.AVG) + 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() + train_loader = DistributedTokenLoader(args.train_files, rank, world_size, device) + swa_state: dict[str, Tensor] | None = None + swa_count = 0 + from collections import deque + lawa_queue: deque[dict[str, Tensor]] = deque(maxlen=args.lawa_k) + ema_state = {name: t.detach().float().clone() for name, t in base_model.state_dict().items()} + ema_decay = 0.9965 # PR #1493: 0.9965 (merged SOTA) + training_time_ms = 0.0 + stop_after_step: int | None = None + torch.cuda.synchronize() + t0 = time.perf_counter() + step = 0 + while True: + 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) + if args.late_qat_threshold > 0 and scale < args.late_qat_threshold and not CastedLinear._qat_enabled: + CastedLinear._qat_enabled = True + log0(f"late_qat:enabled step:{step} scale:{scale:.4f}") + zero_grad_all() + train_loss = torch.zeros((), device=device) + for micro_step in range(grad_accum_steps): + 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) + # === 3-phase overlapped optimizer step === + # Phase 1: Launch async reduce-scatter for banks (biggest first) + optimizer_muon.launch_reduce_scatters() + # Phase 2: All-reduce non-bank grads + step Adam (while bank RS is in-flight) + if distributed: + for p in replicated_params: + if p.grad is not None: + dist.all_reduce(p.grad, op=dist.ReduceOp.AVG) + optimizer_tok.step() + optimizer_scalar.step() + if optimizer_head is not None: + optimizer_head.step() + # Phase 3: Wait for RS, local NS5, all-gather (banks processed last) + optimizer_muon.step() + zero_grad_all() + # EMA update + with torch.no_grad(): + for name, t in base_model.state_dict().items(): + ema_state[name].mul_(ema_decay).add_(t.detach().float(), alpha=1.0 - ema_decay) + step += 1 + # Activate depth recurrence at configured step + if step == args.recur_start_step and not base_model._recurrence_active and base_model.recur_layers: + base_model.set_recurrence_active(True) + log0(f"recurrence:activated at step {step}, virtual_layers={base_model._get_virtual_layers()}") + approx_training_time_ms = training_time_ms + 1000.0 * (time.perf_counter() - t0) + if args.swa_enabled and scale < 0.2 and step % args.swa_every == 0: + if swa_state is None: + swa_state = {name: t.detach().cpu().clone() for name, t in base_model.state_dict().items()} + swa_count = 1 + log0(f"swa:start step:{step}") + else: + for name, t in base_model.state_dict().items(): + swa_state[name] += t.detach().cpu() + swa_count += 1 + if args.lawa_enabled and step % args.lawa_freq == 0: + lawa_queue.append({name: t.detach().cpu().clone() for name, t in base_model.state_dict().items()}) + 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" + ) + 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 + log0( + f"peak memory allocated: {torch.cuda.max_memory_allocated() // 1024 // 1024} MiB " + f"reserved: {torch.cuda.max_memory_reserved() // 1024 // 1024} MiB" + ) + # Apply weight averaging + if args.lawa_enabled and len(lawa_queue) > 1: + log0(f"lawa:applying LAWA averaging k={len(lawa_queue)}") + current_state = base_model.state_dict() + avg_state = {name: torch.zeros(t.shape, dtype=torch.float32, device='cpu') for name, t in current_state.items()} + for snap in lawa_queue: + for name in avg_state: + avg_state[name] += snap[name].float() + for name in avg_state: + avg_state[name] /= len(lawa_queue) + avg_state[name] = avg_state[name].to(dtype=current_state[name].dtype) + base_model.load_state_dict(avg_state, strict=True) + else: + log0("ema:applying EMA weights") + current_state = base_model.state_dict() + avg_state = {name: t.to(dtype=current_state[name].dtype) for name, t in ema_state.items()} + base_model.load_state_dict(avg_state, strict=True) + torch.cuda.synchronize() + t_diag = time.perf_counter() + diag_val_loss, diag_val_bpb = eval_val( + args, compiled_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"DIAGNOSTIC post_ema val_loss:{diag_val_loss:.4f} val_bpb:{diag_val_bpb:.4f} " + f"eval_time:{1000.0 * (time.perf_counter() - t_diag):.0f}ms" + ) + full_state_dict = base_model.state_dict() + export_sd = {k: v for k, v in full_state_dict.items() if "mtp_heads" not in k} + excluded_mtp = sum(int(t.numel()) for k, t in full_state_dict.items() if "mtp_heads" in k) + if excluded_mtp > 0: + log0(f"export_excluding_mtp_params:{excluded_mtp}") + if master_process: + torch.save(export_sd, "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") + # --- Pre-quant AdamW TTT (PR #1485 style) --- + # Fine-tune on val data AFTER EMA, BEFORE quantization. + # Adapted weights get baked into the quantized artifact — Track A compliant. + if args.ttt_enabled: + log0(f"ttt:starting epochs={args.ttt_epochs} lr={args.ttt_lr} freeze_blocks={args.ttt_freeze_blocks} batch_seqs={args.ttt_batch_seqs}") + t_ttt = time.perf_counter() + # CRITICAL: Fresh model instance to avoid inference_mode tensor poisoning + del compiled_model + torch._dynamo.reset() + torch.cuda.empty_cache() + ttt_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, + mtp_num_heads=0, mtp_loss_weight=0.0, + bigram_vocab_size=args.bigram_vocab_size, bigram_dim=args.bigram_dim, + xsa_last_n=args.xsa_last_n, rope_dims=args.rope_dims, ln_scale=args.ln_scale, + dtg=args.dtg_enabled, ve_enabled=args.ve_enabled, ve_dim=args.ve_dim, ve_layers=args.ve_layers, + gated_attention=args.gated_attention, value_residual=args.value_residual, + recur_layers=args.recur_layers, parallel_start_layer=args.parallel_start_layer, + ).to(device).bfloat16() + ttt_model.qo_bank.data = ttt_model.qo_bank.data.float() + ttt_model.kv_bank.data = ttt_model.kv_bank.data.float() + ttt_model.mlp_up_bank.data = ttt_model.mlp_up_bank.data.float() + ttt_model.mlp_down_bank.data = ttt_model.mlp_down_bank.data.float() + for m in ttt_model.modules(): + if isinstance(m, CastedLinear): + m.float() + restore_low_dim_params_to_fp32(ttt_model) + ttt_model.load_state_dict(export_sd, strict=True) + ttt_model.set_recurrence_active(base_model._recurrence_active) + # Freeze first N blocks (PR #1485: freeze 2) + frozen_count = 0 + trainable_count = 0 + if args.ttt_freeze_blocks > 0: + for i, block in enumerate(ttt_model.blocks): + if i < args.ttt_freeze_blocks: + for p in block.parameters(): + p.requires_grad_(False) + frozen_count += p.numel() + ttt_params = [p for p in ttt_model.parameters() if p.requires_grad] + trainable_count = sum(p.numel() for p in ttt_params) + log0(f"ttt:trainable={trainable_count} frozen={frozen_count}") + # AdamW with NO weight decay (PR #1485) + ttt_opt = torch.optim.AdamW(ttt_params, lr=args.ttt_lr, weight_decay=0.0) + # Cosine annealing schedule (PR #1485: eta_min = lr * 0.1) + ttt_scheduler = torch.optim.lr_scheduler.CosineAnnealingLR( + ttt_opt, T_max=args.ttt_epochs, eta_min=args.ttt_lr * 0.1) + # TTT on validation data in batched chunks + ttt_seq_len = args.train_seq_len + total_val_tokens = val_tokens.numel() - 1 + ttt_model.train() + CastedLinear._qat_enabled = False # No QAT during TTT + # Compile TTT model for ~3-4x speedup (14 min -> 4 min) + # Safe: TTT uses train mode + autocast, no inference_mode poisoning + compiled_ttt = torch.compile(ttt_model, dynamic=False, fullgraph=True) + log0("ttt:compiled model for speedup") + for epoch in range(args.ttt_epochs): + ttt_loss_sum = 0.0; ttt_count = 0 + # Iterate through val data in batches of ttt_batch_seqs sequences + for start in range(0, total_val_tokens - ttt_seq_len, ttt_seq_len * args.ttt_batch_seqs): + batch_end = min(start + ttt_seq_len * args.ttt_batch_seqs, total_val_tokens - ttt_seq_len) + batch_starts = list(range(start, batch_end + 1, ttt_seq_len)) + if not batch_starts: + continue + bsz = len(batch_starts) + x_batch = torch.zeros(bsz, ttt_seq_len, dtype=torch.int64, device=device) + y_batch = torch.zeros(bsz, ttt_seq_len, dtype=torch.int64, device=device) + for i, s in enumerate(batch_starts): + chunk = val_tokens[s:s + ttt_seq_len + 1].to(dtype=torch.int64, device=device) + x_batch[i] = chunk[:-1] + y_batch[i] = chunk[1:] + ttt_opt.zero_grad() + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + loss = compiled_ttt(x_batch, y_batch) + loss.backward() + torch.nn.utils.clip_grad_norm_(ttt_params, 1.0) + # All-reduce gradients across GPUs + if distributed: + for p in ttt_params: + if p.grad is not None: + dist.all_reduce(p.grad, op=dist.ReduceOp.AVG) + ttt_opt.step() + ttt_loss_sum += loss.item(); ttt_count += 1 + ttt_scheduler.step() + current_lr = ttt_scheduler.get_last_lr()[0] + log0(f"ttt:epoch {epoch+1}/{args.ttt_epochs} loss={ttt_loss_sum/max(ttt_count,1):.4f} lr={current_lr:.6f} time={time.perf_counter()-t_ttt:.1f}s") + # Replace export_sd with TTT-adapted weights (from base model, not compiled wrapper) + export_sd = {k: v for k, v in ttt_model.state_dict().items() if "mtp_heads" not in k} + del compiled_ttt, ttt_model, ttt_opt, ttt_scheduler + torch._dynamo.reset() + torch.cuda.empty_cache() + log0(f"ttt:done time={time.perf_counter()-t_ttt:.1f}s") + # Unbank 3D tensors into individual 2D tensors for quantization + sd_cpu = {k: v.detach().cpu() for k, v in export_sd.items()} + unbanked_sd = _unbank_state_dict(sd_cpu, args.num_layers) + # Full GPTQ: collect Hessians via a temporary non-banked model + log0(f"gptq:building non-banked model for Hessian collection...") + hessian_model = _HessianGPT( + 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, logit_softcap=args.logit_softcap, + rope_base=args.rope_base, qk_gain_init=args.qk_gain_init, + bigram_vocab_size=args.bigram_vocab_size, bigram_dim=args.bigram_dim, + xsa_last_n=args.xsa_last_n, rope_dims=args.rope_dims, ln_scale=args.ln_scale, + ve_enabled=args.ve_enabled, ve_dim=args.ve_dim, ve_layers=args.ve_layers, + ).to(device).bfloat16() + for m in hessian_model.modules(): + if isinstance(m, CastedLinear): + m.float() + restore_low_dim_params_to_fp32(hessian_model) + # Load unbanked weights into the non-banked model + hessian_model.load_state_dict( + {k: v.to(device) for k, v in unbanked_sd.items() if k in hessian_model.state_dict()}, + strict=False, + ) + # Autoregressive self-generated calibration (no external data) + gptq_calib_seqs = int(os.environ.get("GPTQ_CALIB_SEQS", "32")) + log0(f"gptq:generating autoregressive calibration data ({gptq_calib_seqs} seqs x {args.train_seq_len} tokens, temp=0.8)...") + base_model.load_state_dict(export_sd, strict=False) + t_gen = time.perf_counter() + ar_tokens = generate_autoregressive_calib( + base_model, device, num_seqs=gptq_calib_seqs, seq_len=args.train_seq_len, + vocab_size=args.vocab_size, temperature=0.8, batch_size=8, seed=args.seed, + ) + log0(f"gptq:generated {len(ar_tokens)} sequences in {time.perf_counter()-t_gen:.1f}s") + log0("gptq:collecting hessians from autoregressive data...") + hessians = collect_hessians_from_tokens(hessian_model, ar_tokens, device) + log0(f"gptq:collected hessians for {len(hessians)} layers (AR self-gen)") + # === SAVE RESEARCH DATA (skip for submission runs to save ~2 min) === + if master_process and bool(int(os.environ.get("SAVE_RESEARCH_DATA", "0"))): + research_dir = "/workspace/research_data" + os.makedirs(research_dir, exist_ok=True) + torch.save(unbanked_sd, f"{research_dir}/weights_float32.pt") + log0(f"research:saved float32 weights ({os.path.getsize(f'{research_dir}/weights_float32.pt') / 1024 / 1024:.1f} MB)") + hessian_diags = {k: torch.diag(v).clone() for k, v in hessians.items()} + torch.save(hessian_diags, f"{research_dir}/hessian_diags.pt") + log0(f"research:saved hessian diags ({os.path.getsize(f'{research_dir}/hessian_diags.pt') / 1024 / 1024:.1f} MB)") + top5 = sorted(hessians.keys(), key=lambda k: hessians[k].numel(), reverse=True)[:5] + torch.save({k: hessians[k] for k in top5}, f"{research_dir}/hessians_top5.pt") + log0(f"research:saved top-5 full hessians ({os.path.getsize(f'{research_dir}/hessians_top5.pt') / 1024 / 1024:.1f} MB)") + else: + log0("research:skipped (SAVE_RESEARCH_DATA=0)") + del ar_tokens + del hessian_model + torch.cuda.empty_cache() + sdclip_k = args.sdclip_k_int6 if args.sdclip_enabled else 0.0 + sdclip_k_int8 = args.sdclip_k_int8 if args.sdclip_enabled else 0.0 + quant_result, quant_meta = mixed_quantize_int6(unbanked_sd, {"mlp", "attn"}, hessians=hessians, sdclip_k=sdclip_k, sdclip_k_int8=sdclip_k_int8) + # Save raw quantized values for compression research (only if research data enabled) + if master_process and bool(int(os.environ.get("SAVE_RESEARCH_DATA", "0"))): + torch.save({"w": quant_result, "m": quant_meta}, f"{research_dir}/quant_raw.pt") + log0(f"research:saved raw quant data ({os.path.getsize(f'{research_dir}/quant_raw.pt') / 1024 / 1024:.1f} MB)") + # NOVEL: Selective ±1 pruning by reconstruction error + # Sort ±1 quantized values by their reconstruction error (scale²), + # prune least-impactful first until artifact fits target size. + target_mb = float(os.environ.get("TARGET_MB", "15.22")) # 16MB = 15.2588 MiB; 15.22 leaves 40KB margin + code_bytes_est = len(code.encode("utf-8")) + ones_info = [] # (tensor_key, flat_idx, error) + for name, info in quant_meta.items(): + if not (isinstance(info, dict) and info.get("type") == "int6"): continue + qk, sk = name + ".q", name + ".scale" + if qk not in quant_result or sk not in quant_result: continue + q, s = quant_result[qk], quant_result[sk] + if s.ndim > 0: + ones_mask = (q.abs() == 1) + if ones_mask.any(): + row_idx = torch.arange(q.shape[0]).unsqueeze(1).expand_as(q)[ones_mask] + flat_idx = torch.arange(q.numel()).reshape(q.shape)[ones_mask] + errors = s.float()[row_idx].pow(2) + for fi, err in zip(flat_idx.tolist(), errors.tolist()): + ones_info.append((qk, fi, err)) + if ones_info: + ones_info.sort(key=lambda x: x[2]) + def _apply_prune(n): + """Apply pruning in-place and return compressed size + code.""" + for i in range(min(n, len(ones_info))): + quant_result[ones_info[i][0]].view(-1)[ones_info[i][1]] = 0 + buf = io.BytesIO(); torch.save({"w": quant_result, "m": quant_meta}, buf) + return len(compress_artifact(buf.getvalue(), method=args.compress_method)) + code_bytes_est + def _check_size(): + """Check current compressed size without modifying.""" + buf = io.BytesIO(); torch.save({"w": quant_result, "m": quant_meta}, buf) + return len(compress_artifact(buf.getvalue(), method=args.compress_method)) + code_bytes_est + target_bytes = int(target_mb * 1024 * 1024) + no_sz = _check_size() + log0(f"selective_prune: {len(ones_info)} ±1 candidates, unpruned={no_sz/(1024*1024):.2f}MB target={target_mb}MB") + if no_sz <= target_bytes: + log0("selective_prune: already fits, no pruning needed") + else: + # Linear interpolation: estimate prune count from overshoot ratio + 20% safety margin + overshoot = no_sz - target_bytes + # Estimate bytes saved per pruned value from the distribution + # Each ±1→0 saves ~1 bit in entropy coding, so ~0.125 bytes per value after compression + # But also changes scale utilization. Use conservative 0.05 bytes/value. + est_bytes_per_value = max(overshoot / (len(ones_info) * 0.15), 0.05) + est_prune = int(overshoot / est_bytes_per_value * 1.2) # 20% safety margin + est_prune = min(est_prune, len(ones_info)) + log0(f"selective_prune: overshoot={overshoot/(1024):.0f}KB, est_prune={est_prune} ({100*est_prune/len(ones_info):.1f}%)") + sz = _apply_prune(est_prune) + log0(f"selective_prune: after est_prune={est_prune}: {sz/(1024*1024):.2f}MB {'<=' if sz <= target_bytes else '>'} target") + if sz > target_bytes: + # Still over — apply 50% more incrementally until it fits + while sz > target_bytes and est_prune < len(ones_info): + extra = max(int(est_prune * 0.5), 10000) + new_prune = min(est_prune + extra, len(ones_info)) + _apply_prune(new_prune) # prune from est_prune to new_prune (already sorted) + est_prune = new_prune + sz = _check_size() + log0(f"selective_prune: retry prune={est_prune}: {sz/(1024*1024):.2f}MB {'<=' if sz <= target_bytes else '>'} target") + log0(f"selective_prune: final {est_prune}/{len(ones_info)} ±1 values pruned ({100*est_prune/len(ones_info):.1f}%)") + quant_buf = io.BytesIO() + torch.save({"w": quant_result, "m": quant_meta}, quant_buf) + quant_raw = quant_buf.getvalue() + quant_blob = compress_artifact(quant_raw, method=args.compress_method) + comp_name = quant_blob[:4].decode('ascii', errors='replace') + if master_process: + with open("final_model.int6.ptz", "wb") as f: + f.write(quant_blob) + quant_file_bytes = len(quant_blob) + code_bytes = len(code.encode("utf-8")) + log0(f"Serialized model int6+{comp_name}: {quant_file_bytes} bytes") + log0(f"Total submission size int6+{comp_name}: {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(decompress_artifact(quant_blob_disk)), + map_location="cpu", + ) + deq_unbanked = dequantize_mixed_int6(quant_state["w"], quant_state["m"], unbanked_sd) + # Re-bank the dequantized tensors + deq_state = _rebank_state_dict(deq_unbanked, args.num_layers, sd_cpu) + eval_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, + mtp_num_heads=0, mtp_loss_weight=0.0, + bigram_vocab_size=args.bigram_vocab_size, bigram_dim=args.bigram_dim, + xsa_last_n=args.xsa_last_n, + rope_dims=args.rope_dims, ln_scale=args.ln_scale, dtg=args.dtg_enabled, + ve_enabled=args.ve_enabled, ve_dim=args.ve_dim, ve_layers=args.ve_layers, + gated_attention=args.gated_attention, value_residual=args.value_residual, + recur_layers=args.recur_layers, parallel_start_layer=args.parallel_start_layer, + ).to(device).bfloat16() + eval_model.qo_bank.data = eval_model.qo_bank.data.float() + eval_model.kv_bank.data = eval_model.kv_bank.data.float() + eval_model.mlp_up_bank.data = eval_model.mlp_up_bank.data.float() + eval_model.mlp_down_bank.data = eval_model.mlp_down_bank.data.float() + for m in eval_model.modules(): + if isinstance(m, CastedLinear): + m.float() + restore_low_dim_params_to_fp32(eval_model) + eval_model.load_state_dict(deq_state, strict=True) + # Copy recurrence state from training model to eval model + eval_model.set_recurrence_active(base_model._recurrence_active) + compiled_eval = torch.compile(eval_model, dynamic=False, fullgraph=True) + torch.cuda.synchronize() + t_qeval = time.perf_counter() + q_val_loss, q_val_bpb = eval_val( + args, compiled_eval, rank, world_size, device, grad_accum_steps, + val_tokens, base_bytes_lut, has_leading_space_lut, is_boundary_token_lut, + eval_seq_len=effective_eval_seq_len, + ) + torch.cuda.synchronize() + log0( + f"final_int6_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_int6_roundtrip_exact val_loss:{q_val_loss:.8f} val_bpb:{q_val_bpb:.8f}") + sw_seq_len = effective_eval_seq_len + if args.eval_stride > 0 and args.eval_stride < sw_seq_len: + torch.cuda.synchronize() + t_slide = time.perf_counter() + sw_val_loss, sw_val_bpb = eval_val_sliding( + args, eval_model, rank, world_size, device, + val_tokens, base_bytes_lut, has_leading_space_lut, is_boundary_token_lut, + stride=args.eval_stride, + eval_seq_len=sw_seq_len, + ) + torch.cuda.synchronize() + log0( + f"final_int6_sliding_window val_loss:{sw_val_loss:.4f} val_bpb:{sw_val_bpb:.4f} " + f"stride:{args.eval_stride} eval_time:{1000.0 * (time.perf_counter() - t_slide):.0f}ms" + ) + log0(f"final_int6_sliding_window_exact val_loss:{sw_val_loss:.8f} val_bpb:{sw_val_bpb:.8f}") + log0(f"final_int8_zlib_roundtrip_exact val_loss:{sw_val_loss:.8f} val_bpb:{sw_val_bpb:.8f}") + if args.eval_stride != 64 and 64 < sw_seq_len: + torch.cuda.synchronize() + t_slide64 = time.perf_counter() + sw64_val_loss, sw64_val_bpb = eval_val_sliding( + args, eval_model, rank, world_size, device, + val_tokens, base_bytes_lut, has_leading_space_lut, is_boundary_token_lut, + stride=64, + eval_seq_len=sw_seq_len, + ) + torch.cuda.synchronize() + log0( + f"final_int6_sliding_window_s64 val_loss:{sw64_val_loss:.4f} val_bpb:{sw64_val_bpb:.4f} " + f"stride:64 eval_time:{1000.0 * (time.perf_counter() - t_slide64):.0f}ms" + ) + log0(f"final_int6_sliding_window_s64_exact val_loss:{sw64_val_loss:.8f} val_bpb:{sw64_val_bpb:.8f}") + log0(f"final_int8_zlib_roundtrip_exact val_loss:{sw64_val_loss:.8f} val_bpb:{sw64_val_bpb:.8f}") + # --- SLOT: L-BFGS logit-space (PR #1350) or AdamW hidden-space --- + if args.slot_enabled and args.eval_stride > 0: + try: + slot_stride = args.eval_stride + seq_s = effective_eval_seq_len + total_tok = val_tokens.numel() - 1 + ws_list = [ws for ws in range(0, total_tok, slot_stride) if min(ws + seq_s, total_tok) - ws >= 1] + my_s = (len(ws_list) * rank) // world_size + my_e = (len(ws_list) * (rank + 1)) // world_size + my_ws = ws_list[my_s:my_e] + num_batches = (len(my_ws) + 31) // 32 + sl_loss = torch.zeros((), device=device, dtype=torch.float64) + sl_tc = torch.zeros((), device=device, dtype=torch.float64) + sl_bc = torch.zeros((), device=device, dtype=torch.float64) + V = args.vocab_size + FOCAL_TOKENS = args.slot_focal_tokens + DELTA_CLIP = args.slot_delta_clip + focal_start = max(seq_s - FOCAL_TOKENS, 0) + use_lbfgs = args.slot_lbfgs and args.slot_causal # L-BFGS only with causal mode + slot_mode = "lbfgs-logit-causal" if use_lbfgs else ("adamw-hidden-causal" if args.slot_causal else "adamw-hidden-standard") + log0(f"slot:starting mode={slot_mode} windows={len(my_ws)} batches={num_batches}" + + (f" max_iter={args.slot_lbfgs_max_iter} history={args.slot_lbfgs_history} focal={FOCAL_TOKENS} clip={DELTA_CLIP}" if use_lbfgs else f" lr={args.slot_lr} steps={args.slot_steps}")) + _delta_warmstart = None + torch.cuda.synchronize() + t_slot = time.perf_counter() + eval_model.eval() + for batch_idx, bi in enumerate(range(0, len(my_ws), 32)): + bws = my_ws[bi:bi+32]; bsz = len(bws) + xb = torch.zeros(bsz, seq_s, dtype=torch.int64, device=device) + yb = torch.zeros(bsz, seq_s, dtype=torch.int64, device=device) + wls = [] + for i, ws in enumerate(bws): + end = min(ws + seq_s, total_tok); wl = end - ws; wls.append(wl) + ct = val_tokens[ws:end+1].to(dtype=torch.int64, device=device) + xb[i,:wl] = ct[:-1]; yb[i,:wl] = ct[1:] + # Frozen forward pass + with torch.no_grad(), torch.autocast(device_type="cuda", dtype=torch.bfloat16): + H = eval_model.forward_hidden(xb) + logits_base = eval_model.compute_logits(H).float() + del H + if use_lbfgs: + # --- L-BFGS Causal SLOT in logit space (PR #1350) --- + # Build causal+focal mask: optimize on [focal_start, score_start) per window + opt_mask = torch.zeros(bsz, seq_s, dtype=torch.bool, device=device) + has_opt = False + for i, ws in enumerate(bws): + wl = wls[i] + s = 0 if ws == 0 else max(wl - slot_stride, 0) + if s > focal_start: + opt_mask[i, focal_start:s] = True + has_opt = True + delta = torch.zeros(1, 1, V, device=device, dtype=torch.float32, requires_grad=True) + if _delta_warmstart is not None: + with torch.no_grad(): + delta.data.copy_(_delta_warmstart) + if has_opt: + lbfgs = torch.optim.LBFGS( + [delta], lr=1.0, max_iter=args.slot_lbfgs_max_iter, + history_size=args.slot_lbfgs_history, line_search_fn='strong_wolfe', + tolerance_change=1e-9, tolerance_grad=1e-7, + ) + def _closure(): + lbfgs.zero_grad() + lg = logits_base + delta + nll_all = F.cross_entropy( + lg.reshape(-1, lg.size(-1)), yb.reshape(-1), reduction="none" + ).reshape(bsz, seq_s) + loss = nll_all[opt_mask].mean() + loss.backward() + return loss + lbfgs.step(_closure) + with torch.no_grad(): + delta.data.clamp_(-DELTA_CLIP, DELTA_CLIP) + _delta_warmstart = delta.detach().clone() + # Score new positions with optimized logit delta + with torch.no_grad(): + lg = logits_base + delta.detach() + else: + # --- AdamW hidden-space SLOT (original v6/v7) --- + H_float = logits_base # reuse variable, but we need hidden states + # Re-compute hidden for AdamW mode (needs hidden, not logits) + with torch.no_grad(), torch.autocast(device_type="cuda", dtype=torch.bfloat16): + H_float = eval_model.forward_hidden(xb).detach().float() + del logits_base + if args.slot_causal: + ctx_mask = torch.zeros(bsz, seq_s, device=device, dtype=torch.float32) + for i, ws in enumerate(bws): + wl = wls[i]; score_start = 0 if ws == 0 else max(wl - slot_stride, 0) + if score_start > 0: ctx_mask[i, :score_start] = 1.0 + ctx_mask_sum = ctx_mask.sum().clamp_min(1.0) + delta = torch.zeros(1, 1, H_float.shape[-1], device=device, dtype=H_float.dtype, requires_grad=True) + sopt = torch.optim.AdamW([delta], lr=args.slot_lr, weight_decay=1e-8, eps=1e-5) + for _ in range(args.slot_steps): + sopt.zero_grad() + lg = eval_model.compute_logits((H_float + delta).to(torch.bfloat16)).float() + if args.slot_causal: + nll_all = F.cross_entropy(lg.reshape(-1, lg.size(-1)), yb.reshape(-1), reduction="none").reshape(bsz, seq_s) + loss_s = (nll_all * ctx_mask).sum() / ctx_mask_sum + else: + loss_s = F.cross_entropy(lg.reshape(-1, lg.size(-1)), yb.reshape(-1), reduction="mean") + loss_s.backward() + sopt.step() + with torch.no_grad(): + lg = eval_model.compute_logits((H_float + delta.detach()).to(torch.bfloat16)).float() + nll = F.cross_entropy(lg.reshape(-1, lg.size(-1)), yb.reshape(-1), reduction="none").reshape(bsz, seq_s) + for i, ws in enumerate(bws): + wl = wls[i]; s = 0 if ws == 0 else max(wl - slot_stride, 0) + sl_loss += nll[i, s:wl].to(torch.float64).sum(); sl_tc += float(wl - s) + tgt, prev = yb[i, s:wl], xb[i, s:wl] + tb = base_bytes_lut[tgt].to(torch.float64) + tb += (has_leading_space_lut[tgt] & ~is_boundary_token_lut[prev]).to(torch.float64) + sl_bc += tb.sum() + if not use_lbfgs: + del H_float + if batch_idx % 500 == 0 or batch_idx == num_batches - 1: + log0(f" slot:{slot_mode} batch {batch_idx+1}/{num_batches} time:{time.perf_counter()-t_slot:.1f}s"); sys.stdout.flush() + if dist.is_available() and dist.is_initialized(): + dist.all_reduce(sl_loss, op=dist.ReduceOp.SUM) + dist.all_reduce(sl_tc, op=dist.ReduceOp.SUM) + dist.all_reduce(sl_bc, op=dist.ReduceOp.SUM) + sv_loss = (sl_loss / sl_tc).item() + sv_bpb = sv_loss / math.log(2.0) * (sl_tc.item() / sl_bc.item()) + log0(f"final_slot val_loss:{sv_loss:.4f} val_bpb:{sv_bpb:.4f} mode:{slot_mode} time:{1000*(time.perf_counter()-t_slot):.0f}ms") + log0(f"final_slot_exact val_loss:{sv_loss:.8f} val_bpb:{sv_bpb:.8f}") + log0(f"final_dyneval_exact val_loss:{sv_loss:.8f} val_bpb:{sv_bpb:.8f}") + except Exception as e: + import traceback; log0(f"slot:ERROR {e}"); traceback.print_exc(); sys.stdout.flush() + # Skip dist.destroy_process_group() — can hang on NCCL cleanup after compiled eval, + # causing torchrun to kill the process before the bash script writes SUCCESS marker. + # All results are already flushed to log file at this point. + if distributed: + log0("cleanup:skipping dist.destroy_process_group (NCCL hang risk)") + log0("RUN_COMPLETE") +if __name__ == "__main__": + main() diff --git a/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_seed1337.log b/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_seed1337.log new file mode 100644 index 0000000000..aa7541f0b5 --- /dev/null +++ b/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_seed1337.log @@ -0,0 +1,97 @@ +W0411 07:36:39.938000 3971 torch/distributed/run.py:803] +W0411 07:36:39.938000 3971 torch/distributed/run.py:803] ***************************************** +W0411 07:36:39.938000 3971 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. +W0411 07:36:39.938000 3971 torch/distributed/run.py:803] ***************************************** +logs/066c786e-7fe0-499f-904c-cfae3002b1e5.txt +val_bpb:enabled tokenizer_kind=sentencepiece tokenizer_path=./data/tokenizers/fineweb_8192_bpe.model +train_loader:dataset:fineweb10B_sp8192 train_shards:128 +val_loader:shards pattern=./data/datasets/fineweb10B_sp8192/fineweb_val_*.bin tokens:40540160 +model_params:36311132 +mtp_num_heads:0 mtp_loss_weight:0.2 mtp_params:0 +XSA:last_11 active_layers:[0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10] +world_size:8 grad_accum_steps:1 +sdp_backends:cudnn=False flash=True mem_efficient=False math=False +attention_mode:gqa num_heads:8 num_kv_heads:4 +tie_embeddings:True embed_lr:0.03 head_lr:0.0 matrix_lr:0.022 scalar_lr:0.02 +train_batch_tokens:786432 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:9.0000 val_bpb:3.4842 train_time:0ms step_avg:0.02ms +step:1/20000 train_loss:9.0010 train_time:132ms step_avg:131.87ms +step:2/20000 train_loss:9.1281 train_time:161ms step_avg:80.48ms +step:3/20000 train_loss:9.5563 train_time:257ms step_avg:85.76ms +step:4/20000 train_loss:9.1507 train_time:354ms step_avg:88.59ms +step:5/20000 train_loss:8.6754 train_time:451ms step_avg:90.27ms +step:6/20000 train_loss:8.5371 train_time:548ms step_avg:91.36ms +step:7/20000 train_loss:8.3210 train_time:647ms step_avg:92.40ms +step:8/20000 train_loss:7.8855 train_time:744ms step_avg:93.04ms +step:9/20000 train_loss:7.8673 train_time:842ms step_avg:93.50ms +step:10/20000 train_loss:7.6615 train_time:939ms step_avg:93.87ms +step:500/20000 train_loss:3.3724 train_time:49779ms step_avg:99.56ms +step:1000/20000 train_loss:3.0971 train_time:100060ms step_avg:100.06ms +step:1500/20000 train_loss:3.2994 train_time:150263ms step_avg:100.18ms +step:2000/20000 train_loss:3.1179 train_time:200422ms step_avg:100.21ms +step:2500/20000 train_loss:3.2631 train_time:250581ms step_avg:100.23ms +recurrence:activated at step 3000, virtual_layers=[0, 1, 2, 3, 4, 5, 3, 4, 5, 6, 7, 8, 9, 10] +step:3000/20000 train_loss:2.9732 train_time:300706ms step_avg:100.24ms +step:3500/20000 train_loss:3.0807 train_time:394448ms step_avg:112.70ms +step:4000/20000 train_loss:2.9819 train_time:456100ms step_avg:114.03ms +step:4000/20000 val_loss:2.9667 val_bpb:1.1485 train_time:456170ms step_avg:114.04ms +swa:start step:4500 +step:4500/20000 train_loss:2.9897 train_time:517605ms step_avg:115.02ms +late_qat:enabled step:4640 scale:0.1499 +step:5000/20000 train_loss:2.8100 train_time:580166ms step_avg:116.03ms +step:5159/20000 val_loss:2.8491 val_bpb:1.1030 train_time:600185ms step_avg:116.34ms +stopping_early: wallclock_cap train_time:600185ms step:5159/20000 +peak memory allocated: 32963 MiB reserved: 32968 MiB +ema:applying EMA weights +DIAGNOSTIC post_ema val_loss:2.8463 val_bpb:1.1019 eval_time:2001ms +Serialized model: 136152884 bytes +Code size: 137532 bytes +ttt:starting epochs=6 lr=0.0005 freeze_blocks=2 batch_seqs=32 +ttt:trainable=36307020 frozen=4112 +ttt:compiled model for speedup +ttt:epoch 1/6 loss=2.9077 lr=0.000470 time=156.7s +ttt:epoch 2/6 loss=2.8459 lr=0.000388 time=210.2s +ttt:epoch 3/6 loss=2.7979 lr=0.000275 time=264.0s +ttt:epoch 4/6 loss=2.7459 lr=0.000163 time=317.6s +ttt:epoch 5/6 loss=2.6946 lr=0.000080 time=371.0s +ttt:epoch 6/6 loss=2.6536 lr=0.000050 time=424.5s +ttt:done time=424.8s +gptq:building non-banked model for Hessian collection... +gptq:generating autoregressive calibration data (32 seqs x 2048 tokens, temp=0.8)... +gptq:generated 32 sequences in 128.2s +gptq:collecting hessians from autoregressive data... +gptq:collected hessians for 67 layers (AR self-gen) +research:skipped (SAVE_RESEARCH_DATA=0) +selective_prune: 9943000 ±1 candidates, unpruned=14.72MB target=15.22MB +selective_prune: already fits, no pruning needed +Serialized model int6+BROT: 15301838 bytes +Total submission size int6+BROT: 15439370 bytes +final_int6_roundtrip val_loss:2.7614 val_bpb:1.0690 eval_time:23333ms +final_int6_roundtrip_exact val_loss:2.76144893 val_bpb:1.06904293 +final_int6_sliding_window val_loss:2.7343 val_bpb:1.0586 stride:64 eval_time:109787ms +final_int6_sliding_window_exact val_loss:2.73434671 val_bpb:1.05855527 +final_int8_zlib_roundtrip_exact val_loss:2.73434671 val_bpb:1.05855527 +cleanup:skipping dist.destroy_process_group (NCCL hang risk) +RUN_COMPLETE +[rank0]:[W411 08:04:00.874237743 ProcessGroupNCCL.cpp:1524] Warning: WARNING: destroy_process_group() was not called before program exit, which can leak resources. For more info, please see https://pytorch.org/docs/stable/distributed.html#shutdown (function operator()) diff --git a/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_seed2024.log b/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_seed2024.log new file mode 100644 index 0000000000..9900ba619a --- /dev/null +++ b/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_seed2024.log @@ -0,0 +1,98 @@ +W0411 09:06:39.801000 3967 torch/distributed/run.py:803] +W0411 09:06:39.801000 3967 torch/distributed/run.py:803] ***************************************** +W0411 09:06:39.801000 3967 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. +W0411 09:06:39.801000 3967 torch/distributed/run.py:803] ***************************************** +logs/40bef99f-a60b-4302-801f-10523f70424e.txt +val_bpb:enabled tokenizer_kind=sentencepiece tokenizer_path=./data/tokenizers/fineweb_8192_bpe.model +train_loader:dataset:fineweb10B_sp8192 train_shards:128 +val_loader:shards pattern=./data/datasets/fineweb10B_sp8192/fineweb_val_*.bin tokens:40540160 +model_params:36311132 +mtp_num_heads:0 mtp_loss_weight:0.2 mtp_params:0 +XSA:last_11 active_layers:[0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10] +world_size:8 grad_accum_steps:1 +sdp_backends:cudnn=False flash=True mem_efficient=False math=False +attention_mode:gqa num_heads:8 num_kv_heads:4 +tie_embeddings:True embed_lr:0.03 head_lr:0.0 matrix_lr:0.022 scalar_lr:0.02 +train_batch_tokens:786432 train_seq_len:2048 iterations:20000 warmup_steps:20 max_wallclock_seconds:600.000 +seed:2024 +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:9.0006 val_bpb:3.4844 train_time:0ms step_avg:0.01ms +step:1/20000 train_loss:9.0022 train_time:132ms step_avg:131.76ms +step:2/20000 train_loss:9.1305 train_time:162ms step_avg:80.87ms +step:3/20000 train_loss:9.5925 train_time:258ms step_avg:86.00ms +step:4/20000 train_loss:9.2160 train_time:355ms step_avg:88.70ms +step:5/20000 train_loss:8.7018 train_time:452ms step_avg:90.37ms +step:6/20000 train_loss:8.4469 train_time:549ms step_avg:91.47ms +step:7/20000 train_loss:8.2278 train_time:647ms step_avg:92.49ms +step:8/20000 train_loss:7.7830 train_time:746ms step_avg:93.30ms +step:9/20000 train_loss:7.7574 train_time:847ms step_avg:94.16ms +step:10/20000 train_loss:7.5033 train_time:946ms step_avg:94.61ms +step:500/20000 train_loss:3.3730 train_time:49778ms step_avg:99.56ms +step:1000/20000 train_loss:3.1054 train_time:100058ms step_avg:100.06ms +step:1500/20000 train_loss:3.2960 train_time:150278ms step_avg:100.19ms +step:2000/20000 train_loss:3.1205 train_time:200439ms step_avg:100.22ms +step:2500/20000 train_loss:3.2621 train_time:250527ms step_avg:100.21ms +recurrence:activated at step 3000, virtual_layers=[0, 1, 2, 3, 4, 5, 3, 4, 5, 6, 7, 8, 9, 10] +step:3000/20000 train_loss:2.9741 train_time:300499ms step_avg:100.17ms +step:3500/20000 train_loss:3.0786 train_time:394078ms step_avg:112.59ms +step:4000/20000 train_loss:2.9802 train_time:455628ms step_avg:113.91ms +step:4000/20000 val_loss:2.9686 val_bpb:1.1493 train_time:455698ms step_avg:113.92ms +swa:start step:4500 +step:4500/20000 train_loss:2.9876 train_time:517032ms step_avg:114.90ms +late_qat:enabled step:4645 scale:0.1498 +step:5000/20000 train_loss:2.8099 train_time:579538ms step_avg:115.91ms +step:5164/20000 val_loss:2.8508 val_bpb:1.1036 train_time:600112ms step_avg:116.21ms +stopping_early: wallclock_cap train_time:600112ms step:5164/20000 +peak memory allocated: 32964 MiB reserved: 33018 MiB +ema:applying EMA weights +DIAGNOSTIC post_ema val_loss:2.8478 val_bpb:1.1025 eval_time:1997ms +Serialized model: 136152884 bytes +Code size: 137532 bytes +ttt:starting epochs=6 lr=0.0005 freeze_blocks=2 batch_seqs=32 +ttt:trainable=36307020 frozen=4112 +ttt:compiled model for speedup +ttt:epoch 1/6 loss=2.9076 lr=0.000470 time=159.4s +ttt:epoch 2/6 loss=2.8469 lr=0.000388 time=212.9s +ttt:epoch 3/6 loss=2.7976 lr=0.000275 time=266.6s +ttt:epoch 4/6 loss=2.7455 lr=0.000163 time=320.2s +ttt:epoch 5/6 loss=2.6940 lr=0.000080 time=373.6s +ttt:epoch 6/6 loss=2.6526 lr=0.000050 time=426.9s +ttt:done time=427.2s +gptq:building non-banked model for Hessian collection... +gptq:generating autoregressive calibration data (32 seqs x 2048 tokens, temp=0.8)... +gptq:generated 32 sequences in 128.5s +gptq:collecting hessians from autoregressive data... +gptq:collected hessians for 67 layers (AR self-gen) +research:skipped (SAVE_RESEARCH_DATA=0) +selective_prune: 9996116 ±1 candidates, unpruned=14.76MB target=15.22MB +selective_prune: already fits, no pruning needed +Serialized model int6+BROT: 15343238 bytes +Total submission size int6+BROT: 15480770 bytes +final_int6_roundtrip val_loss:2.7619 val_bpb:1.0692 eval_time:24151ms +final_int6_roundtrip_exact val_loss:2.76187094 val_bpb:1.06920631 +final_int6_sliding_window val_loss:2.7358 val_bpb:1.0591 stride:64 eval_time:109724ms +final_int6_sliding_window_exact val_loss:2.73581780 val_bpb:1.05912478 +final_int8_zlib_roundtrip_exact val_loss:2.73581780 val_bpb:1.05912478 +cleanup:skipping dist.destroy_process_group (NCCL hang risk) +RUN_COMPLETE +[rank0]:[W411 09:34:01.731678710 ProcessGroupNCCL.cpp:1524] Warning: WARNING: destroy_process_group() was not called before program exit, which can leak resources. For more info, please see https://pytorch.org/docs/stable/distributed.html#shutdown (function operator()) +SEED_2024_SUCCESS diff --git a/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_seed42.log b/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_seed42.log new file mode 100644 index 0000000000..bfe63a5c4f --- /dev/null +++ b/records/track_10min_16mb/2026-04-11_SP8192_PreQuantTTT_CompiledTTT/train_seed42.log @@ -0,0 +1,95 @@ +W0411 06:38:33.527000 3959 torch/distributed/run.py:803] +W0411 06:38:33.527000 3959 torch/distributed/run.py:803] ***************************************** +W0411 06:38:33.527000 3959 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. +W0411 06:38:33.527000 3959 torch/distributed/run.py:803] ***************************************** +logs/734c523f-d4b5-4e76-b3d1-2b23345e751e.txt +val_bpb:enabled tokenizer_kind=sentencepiece tokenizer_path=./data/tokenizers/fineweb_8192_bpe.model +train_loader:dataset:fineweb10B_sp8192 train_shards:128 +val_loader:shards pattern=./data/datasets/fineweb10B_sp8192/fineweb_val_*.bin tokens:40540160 +model_params:36311132 +mtp_num_heads:0 mtp_loss_weight:0.2 mtp_params:0 +XSA:last_11 active_layers:[0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10] +world_size:8 grad_accum_steps:1 +sdp_backends:cudnn=False flash=True mem_efficient=False math=False +attention_mode:gqa num_heads:8 num_kv_heads:4 +tie_embeddings:True embed_lr:0.03 head_lr:0.0 matrix_lr:0.022 scalar_lr:0.02 +train_batch_tokens:786432 train_seq_len:2048 iterations:20000 warmup_steps:20 max_wallclock_seconds:600.000 +seed:42 +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:8.9996 val_bpb:3.4840 train_time:0ms step_avg:0.01ms +step:1/20000 train_loss:9.0002 train_time:131ms step_avg:130.87ms +step:2/20000 train_loss:9.1220 train_time:161ms step_avg:80.26ms +step:3/20000 train_loss:9.4771 train_time:257ms step_avg:85.80ms +step:4/20000 train_loss:9.0481 train_time:354ms step_avg:88.56ms +step:5/20000 train_loss:8.5784 train_time:451ms step_avg:90.19ms +step:6/20000 train_loss:8.4059 train_time:548ms step_avg:91.36ms +step:7/20000 train_loss:8.1422 train_time:648ms step_avg:92.57ms +step:8/20000 train_loss:7.7518 train_time:748ms step_avg:93.52ms +step:9/20000 train_loss:7.7380 train_time:847ms step_avg:94.10ms +step:10/20000 train_loss:7.5287 train_time:948ms step_avg:94.81ms +step:500/20000 train_loss:3.3825 train_time:49856ms step_avg:99.71ms +step:1000/20000 train_loss:3.1040 train_time:100100ms step_avg:100.10ms +step:1500/20000 train_loss:3.2985 train_time:150318ms step_avg:100.21ms +step:2000/20000 train_loss:3.1237 train_time:200508ms step_avg:100.25ms +step:2500/20000 train_loss:3.2609 train_time:250657ms step_avg:100.26ms +recurrence:activated at step 3000, virtual_layers=[0, 1, 2, 3, 4, 5, 3, 4, 5, 6, 7, 8, 9, 10] +step:3000/20000 train_loss:2.9775 train_time:300765ms step_avg:100.25ms +step:3500/20000 train_loss:3.0857 train_time:394306ms step_avg:112.66ms +step:4000/20000 train_loss:2.9865 train_time:455967ms step_avg:113.99ms +step:4000/20000 val_loss:2.9692 val_bpb:1.1495 train_time:456035ms step_avg:114.01ms +swa:start step:4500 +step:4500/20000 train_loss:2.9868 train_time:517537ms step_avg:115.01ms +late_qat:enabled step:4641 scale:0.1497 +step:5000/20000 train_loss:2.8141 train_time:580273ms step_avg:116.05ms +step:5158/20000 val_loss:2.8514 val_bpb:1.1039 train_time:600137ms step_avg:116.35ms +stopping_early: wallclock_cap train_time:600137ms step:5158/20000 +peak memory allocated: 32963 MiB reserved: 32968 MiB +ema:applying EMA weights +DIAGNOSTIC post_ema val_loss:2.8486 val_bpb:1.1028 eval_time:2001ms +Serialized model: 136152884 bytes +Code size: 137223 bytes +ttt:starting epochs=6 lr=0.0005 freeze_blocks=2 batch_seqs=32 +ttt:trainable=36307020 frozen=4112 +ttt:compiled model for speedup +ttt:epoch 1/6 loss=2.9106 lr=0.000470 time=157.4s +ttt:epoch 2/6 loss=2.8479 lr=0.000388 time=211.0s +ttt:epoch 3/6 loss=2.7998 lr=0.000275 time=264.9s +ttt:epoch 4/6 loss=2.7478 lr=0.000163 time=318.6s +ttt:epoch 5/6 loss=2.6964 lr=0.000080 time=372.3s +ttt:epoch 6/6 loss=2.6552 lr=0.000050 time=425.9s +ttt:done time=426.2s +gptq:building non-banked model for Hessian collection... +gptq:generating autoregressive calibration data (32 seqs x 2048 tokens, temp=0.8)... +gptq:generated 32 sequences in 129.6s +gptq:collecting hessians from autoregressive data... +gptq:collected hessians for 67 layers (AR self-gen) +research:skipped (SAVE_RESEARCH_DATA=0) +selective_prune: 9985345 ±1 candidates, unpruned=14.76MB target=15.22MB +selective_prune: already fits, no pruning needed +Serialized model int6+BROT: 15340052 bytes +Total submission size int6+BROT: 15477275 bytes +final_int6_roundtrip val_loss:2.7600 val_bpb:1.0685 eval_time:23904ms +final_int6_roundtrip_exact val_loss:2.75997805 val_bpb:1.06847351 +final_int6_sliding_window val_loss:2.7339 val_bpb:1.0584 stride:64 +final_int6_sliding_window_exact val_loss:2.73394084 val_bpb:1.05839815 +final_int8_zlib_roundtrip_exact val_loss:2.73394084 val_bpb:1.05839815 +# NOTE: sliding window lines appended from runner polling data (pod died before log backup)