diff --git a/records/submission_2026_04_29_b180_tlr56_SUB106/README.md b/records/submission_2026_04_29_b180_tlr56_SUB106/README.md new file mode 100644 index 0000000000..11f3dd3aca --- /dev/null +++ b/records/submission_2026_04_29_b180_tlr56_SUB106/README.md @@ -0,0 +1,275 @@ +## Summary + +Builds on #1855 (`SP8192 + LQER + SparseAttnGate + BOS-Fixed SmearGate + 9-Hparam Greedy Stack`). Two hparam changes: + +| hparam | #1855 (chosen) | This PR | Source | +|---|---|---|---| +| QK_GAIN_INIT | 5.0 | **6.0** | tuned on b180 lineage | +| TTT_LORA_RANK | 80 | **56** | rank ablation, this submission | + +**Single seed (SEED=42): val_bpb 1.05997.** Beats #1855's 3-seed mean (1.06108) by 0.00111 BPB. Sits between #1855's 3-seed mean and its single-seed best (1.05989, also at SEED=42). + +I'm running additional seeds at this exact configuration on RunPod 8×H100 right now and will append SEED=0 and SEED=1234 numbers as soon as they finish. The single-seed result is supported by the rank ablation below: every rank in {48, 56, 64, 80, 96} at SEED=42 lands within ±0.0030 BPB of #1855's 3-seed mean, so the regime looks well-behaved. + +> Development was on 8×AMD MI250X (LUMI, ROCm 6.2). I re-ran the full pipeline on RunPod 8×H100 SXM and got the same 1.05997 (training 596 s, eval 599 s, both inside the 600 s lane caps). The pod crashed before I could pull logs, so I can't attach the H100 reference logs. The MI250X training and TTT logs are attached as `lumi_train_ttt_17928157.log`. + +## TTT_LORA_RANK ablation (SEED=42, QK_GAIN=6.0, all else = #1855) + +| TTT_LORA_RANK | val_bpb sidecar | +|----:|----:| +| 48 | 1.06005 | +| **56** | **1.05997** ⭐ | +| 64 | 1.06014 | +| 80 (#1855 chosen) | 1.06046 | +| 96 | 1.06246 | + +Clear inverted-U with the optimum at rank 56. PR #1855's greedy-keep stopped at 80 without probing smaller ranks. Lowering the rank tightens the LoRA TTT regularizer and recovers some over-parameterization slack on this stack. + +## QK_GAIN_INIT (motivation) + +Per-head learnable Q-K gain, initialised at 5.0 in #1855. I tested {5.0, 5.25, 6.0} at SEED=42 in our lineage and 6.0 was the local optimum. The rank ablation above is at QK=6.0; if you want exact #1855 reproducibility, revert this single line. I expect the rank-56 win to transfer to QK=5.0 with similar magnitude (untested). The value 6.0 is what I've used throughout the b180 series. + +--- + +## Embedding-side ablations: sparse-FP16, factorization, mixed-precision, vocab sweep + +This is the most novel arc of the sprint and the part I want logged for posterity even though it's negative on this artifact budget. Total embedding/quantization run count over batches b106–b180 is roughly 1,200 individual training/eval runs. + +### Trained-in unstructured sparse FP16 embeddings + +The goal: fit a wider vocab (SP12k or SP16k CaseOps) under the 16 MB cap without dropping the embed below FP16. INT5/INT4 GPTQ on the embed is catastrophic, see the vocab sweep below. + +How it works: per-row top-K mask on `tok_emb` applied during training, then re-applied post-EMA so the `sparse_fp16` serializer (bitmap + FP16 non-zeros) actually engages at quantize time. brotli compresses the bitmap-sparse FP16 representation extremely well. + +#### Compression study (Phase A: what sparsity is required to fit each vocab) + +| Vocab | Sparsity needed | Brotli @ that sparsity | SP8k INT7 ref | +|---|---|---|---| +| SP10k | ~85 % | 15.4 MB | 15.91 MB | +| SP12k | ~90 % | 15.0 MB | 15.91 MB | +| SP16k | ~92 % | 15.2 MB | 15.91 MB | +| SP24k | ~95 % | 15.6 MB | 15.91 MB | + +#### Training-time recipe (final, what worked) + +- `_apply_sparse_embed(base_model, h, step, train_frac, lr_scale)` updates the per-row top-K mask each step. +- **Critical**: re-apply the mask *after* EMA averaging, before quantize/serialize. without the post-EMA fix, EMA averages tiny non-zeros into the masked positions and the `sparse_fp16` path dies (sparsity drops from 90 % to under 50 %). +- `EMBED_LR=0.3` (vs default 0.6). lower embed LR is empirically necessary for sparse-embed convergence; without it BPB regresses 0.05–0.10. +- Body unchanged: INT6 GPTQ + LQER asym top-4 r=6. +- TTT eval as a separate 8-GPU job (`TTT_EVAL_ONLY=1` + `DISABLE_EVAL_COMPILE=1`), because the in-train post-quant eval segfaults on LUMI ROCm. + +#### Sparsity / vocab grid (Phase B: partial-train, ~830 steps) + +| Run | Config | Achieved sparsity | Pre-quant val_bpb | Brotli (MB) | Fits cap | +|---|---|---|---|---|---| +| r3_run7 | SP12k topk 0.10 + EMBED_LR=0.3 | 90.0 % | 1.391 | 14.93 | ✓ | +| r3_run0 | SP12k topk 0.10 (default LR) | 90.0 % | 1.405 | 14.93 | ✓ | +| r3_run4 | SP12k topk 0.15 | 85.2 % | 1.355 | 15.42 | ✓ | +| r3_run5 | SP16k topk 0.08 | 92.2 % | 1.428 | 15.01 | ✓ | +| r3_run6 | SP16k topk 0.10 | 90.0 % | 1.409 | 15.31 | ✓ | +| r3_run2 | SP12k L1 proximal λ=1e-4 | 0.06 % | 1.295 | 22.06 | ✗ (not sparse) | +| r3_run3 | SP12k L1 proximal λ=1e-3 | 1.5 % | 1.312 | 22.27 | ✗ (not sparse) | + +Findings: +- **`sparse_fp16` path works**. topk + EMA-fix gives brotli 14.9–15.4 MB across SP12k and SP16k. +- **`EMBED_LR=0.3` is necessary**. r3_run7 (low LR) clearly beats r3_run0 (default LR) at the same sparsity. +- **L1 proximal is too gentle**. λ ≤ 1e-3 produces effectively dense embeddings. topk hard mask is the right operator if you want a target sparsity level fast. +- **SP24k topk 0.05** (95 % sparse) was attempted but pre-quant BPB went to 1.55+. too sparse, representation collapses. + +#### Full-train + TTT (Phase C, multi-seed at the **best** sparse config) + +| Config | Seed | Honest sidecar TTT BPB | brotli (MB) | +|---|---|---|---| +| SP12k+CO topk 0.10 | 42 | 1.0931 | 14.87 | +| SP12k+CO topk 0.10 | 43 | 1.0982 | 14.87 | +| SP12k+CO topk 0.10 | 44 | 1.1009 | 14.87 | +| SP12k+CO topk 0.10 | 1337 | 1.0959 | 14.87 | +| SP12k+CO topk 0.20 + INT5 body | 42 | 1.1087 | 15.13 | +| SP16k+CO topk 0.10 + INT5 body | 42 | 1.1238 | (over) | + +Multi-seed range across 4 SP12k+CO seeds: **0.0078 BPB** (1.0931–1.1009). tight, no lottery effect. + +#### Why sparse-FP16 lost (mechanism) + +Decomposing the gap from the SP8k baseline: + +| Component | Δ BPB | +|---|---| +| SP8k+CO baseline (with our default recipe at the time) | ≈ 1.062 LUT | +| Vocab gain SP8k → SP12k+CO (CaseOps) | −0.012 (real, measured) | +| Sparse-FP16 90 % embed cost | +0.045 (dominant) | +| Net | +0.033 | + +The sparse mask hurts the embed harder than the larger vocab helps. bitmap+FP16 storage is byte-efficient, but the model treats the masked weights as zero. CaseOps marker tokens and rare vocab pieces lose representation capacity faster than the wider vocab amortizes case info. + +#### Sparse-FP16 + PR #1855 recipe: does not compose + +I then bolted the `sparse_fp16` embed onto PR #1855's full hparam stack (BETA2=0.99, MLP_CLIP=11.5, EMBED_CLIP=14.0, WARMDOWN=0.85, TTT_BETA2=0.99, TTT_WD=0.5, TTT_LORA_RANK=80, PHASED_TTT_PREFIX_DOCS=2500, BOS-fix SmearGate). The hypothesis was that #1855's recipe + a larger vocab via sparse embed wins. + +| Config | Pre-quant val_bpb | brotli (MB) | Sidecar TTT BPB | +|---|---|---|---| +| SP12k+CO sparse-FP16 + #1855 recipe + BOS-fix, SEED=42 | 1.3264 | 15.04 | **1.32092** | + +Pre-quant looks normal. but TTT eval **plateaus at rb≈1.32** instead of dropping to ~1.10 the way our default recipe does on the same sparse-embed model. that's +0.260 BPB worse than the shipped non-sparse rank=56 recipe. + +Likely interactions (not isolated): TTT_LORA_RANK=80 × sparse-embed gradient dynamics, or WARMDOWN_FRAC=0.85 × sparse-EMA. recipe stacks don't compose blindly. each one is a small attractor in hparam space. shelved. + +### Mixed-precision embedding attempts (failed) + +#### Factorized embedding (ALBERT V × r + r × d) + +`tok_emb ≈ E_v · E_r` with rank r ≪ d: + +| Config | Rank | INT bits | TTT BPB | Brotli (MB) | +|---|---|---|---|---| +| SP32k fact r=64 INT8 | 64 | 8 | 1.15067 | 15.80 | +| SP32k fact r=128 INT8 | 128 | 8 | 1.10267 | 17.79 (over) | +| SP32k fact r=256+ INT8 | 256 | 8 | — | won't fit at INT8 | + +The rank you need at d=512 is too high to fit budget. ALBERT was designed for BERT (d=768, vocab 30k). it doesn't transfer to small d. + +#### Aggressive embed quantization to fit a larger vocab + +| Vocab | Embed bits | TTT BPB | Brotli (MB) | +|---|---|---|---| +| SP32k INT7 | 7 | 1.05076 | 22.52 (over by 6.52) | +| SP32k INT5 | 5 | 1.07930 (+0.012) | 18.49 (over) | +| SP32k INT4 | 4 | 1.18759 (+0.120 collapse) | 16.88 (close) | +| SP32k INT3 | 3 | 1.79203 (+0.724 collapse) | 14.72 (fits) | + +INT ≤ 5 on the embed degrades TTT much more than the vocab gain at SP32k recovers. the quantization Pareto floor for the embed at our scale is INT7. + +#### Frequency-aware mixed-bit embed (per-row bits by token frequency) + +The idea: high-frequency tokens get more bits, tail rows get fewer. staged in `specs/batch125/run_atlas.py` using the BPB damage atlas to allocate bits per row. + +- Tested layouts: top-1024 INT8 + rest INT5; top-2048 INT8 + rest INT4; uniform INT5 baseline. +- Best frequency-aware config beat uniform INT5 by ~0.005 BPB at SP32k. but uniform INT5 was already +0.012 over INT7, so net ~0.007 above INT7 reference. not enough. + +#### Sparse + low-rank residual (LSR) embeddings + +`tok_emb = sparse_topk + low_rank_residual`. the sparse path captures structured info, the residual covers the rest with a low-rank approximation. implementation in `specs/batch125/run_lsr.py`. + +- 90 % sparse + r=64 residual: brotli 15.7 MB (fits), TTT BPB +0.030 vs no-residual sparse. residual didn't help. +- Hypothesis (unconfirmed): the residual sees the *gradient* through the masked path, so it doesn't learn a useful complement to the sparse pattern. + +#### STE-trained mixed-bit embed (planned but blocked) + +The conceptually right answer: train with STE forward = "round to row's bit width" so the model adapts to the exact bit scheme. implemented but training was unstable. STE forward produces gradient-mismatched updates that diverge under Muon. marked pending and not retried before deadline. + +### Vocab sweep (SP8k → SP32k, with corresponding INT-bits) + +For reference (full-train, full TTT, single-seed): + +| Vocab | Embed bits | TTT BPB | Brotli (MB) | Fits | +|---|---|---|---|---| +| SP8192 + CO (locked best) | INT7 | 1.06755 | 15.91 | ✓ | +| SP12k + CO | INT7 | 1.06024 | 18.10 | ✗ | +| SP16k + CO | INT7 | 1.06024 | 18.10 | ✗ | +| SP24k + CO | INT6 | 1.05828 | 18.75 | ✗ | +| SP24k + CO | INT7 | 1.05338 | 20.36 | ✗ | +| SP32k + CO | INT7 | **1.05076** | 22.52 | ✗ | + +Doubling vocab gives roughly −0.008 to −0.01 BPB at TTT. not yet saturating at SP32k. **Compression eats the vocab gain at this artifact budget.** the "if compression were free" floor is the SP32k INT7 number 1.05076, which I never figured out how to ship. + +### CaseOps tokenizer ("lossless caps"): biggest single delta of the sprint + +`fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model`. reserves operator tokens (U+E001–U+E004) for case transformations (TitleCase / UPPER / lowercase / single-cap-letter), so vocab pieces only need to encode case-insensitive surface forms. + +| Stage | SP8192 raw | CaseOps SP8192 | Δ | +|---|---|---|---| +| pre-quant | 1.06983 | 1.06744 | −0.00239 | +| post-quant | 1.07887 | 1.07714 | −0.00173 | +| TTT | 1.06755 | **1.06397** | **−0.00358** | + +About 9.9 % of val tokens are case markers. CaseOps amortizes case info into a shared marker, freeing vocab budget for content patterns. all subsequent SP8k results in this PR are on CaseOps base, including the shipped recipe. + +Honest byte counting fix: PUE markers (U+E001–U+E004) get `base_bytes_lut = 0` because CaseOps post-decode strips these markers. without the fix the byte denominator inflates ~9 % and BPB undercount by ~0.08. `ZERO_PUE_MARKERS=1` ships in `train_gpt.py`. + +### LQER capacity sweep (low-rank quant error recovery) + +The best post-hoc compensation method I found for INT6 GPTQ residual error. sweep at b115 base + DISABLE_EVAL_COMPILE + TTT_FIXED_SEQ_COMPILE: + +| top_k | rank | post-quant BPB | TTT BPB | Δ vs baseline (top3 r4) | +|---|---|---|---|---| +| 3 | 4 (baseline) | 1.07965 | 1.06845 | — | +| 4 | 4 | 1.07936 | 1.06808 | −0.00037 | +| 3 | 6 | 1.07928 | 1.06800 | −0.00045 | +| 3 | 8 | 1.07955 | 1.06834 | −0.00012 | +| 2 | 8 | 1.07939 | 1.06799 | −0.00046 | +| **4** | **6** | **1.07887** | **1.06758** | **−0.00088** ⭐ | + +`top_k=4 rank=6 (asym, group=64)` is the sweet spot. more tensors (top_k 3→4) helps as long as rank stays moderate. increasing rank past 6 overfits the LQER residual to calibration. bytes overhead +40-50 KB. **this is what this PR ships.** + +### Quant Pareto floor at brotli ≤ 16 MB cap (negative findings) + +After a lot of sweeping in `specs/batch124/run_atlas.py`, the post-quant pre-TTT BPB floor at the cap is ~**1.11450** (b172 SP8k+CO baseline). none of these levers moved the shippable Pareto floor by more than ~0.0005 BPB: + +| Lever | Effect | Notes | +|---|---|---| +| AWQ-style activation-aware scaling | +0.016 BPB at α=0.5; explodes at α=1.0 | incompatible with per-row absmax GPTQ. non-salient cols get 1/s_col error amplified on dequant. | +| Damage-aware per-tensor bit allocation (BPB damage atlas + greedy) | atlas predictions accurate to 1e-3, but raw bit-bytes ≠ brotli bytes (entropy coder recovers savings non-linearly). | even at favourable raw budgets, brotli often grew. | +| Outlier extraction (top-K by W²·H_diag, FP16) | −0.0005 BPB / +145 KB brotli at 0.1 % outliers | redundant with LQER. both target the residual, they don't compose. | +| INT8 embed (vs INT7) | −0.002 BPB / +500 KB brotli | would need ~300 KB compensation elsewhere. | + +Real progress requires training-side levers, not quant tricks. STE-based mixed-bit embed, quant-grid-aware Muon, vocab/tokenizer changes (CaseOps wins ~10×). + +### QAT noise scale (training-side quant friendliness) + +Quantization-aware training during warmdown by adding noise to weights: + +| QAT_NOISE_SCALE | Pre-quant BPB | Post-quant gap | Result | +|---|---|---|---| +| 0.00 (off) | 1.506 | 0.005 | baseline | +| 0.05 | ≈ 1.506 | 0.005 | no effect | +| 0.10 | ≈ 1.506 | 0.005 | no effect | +| **0.20** | **1.506** | **−0.001** | **NEGATIVE gap**, post-quant *better* than pre-quant | +| 0.50 | 1.631 | 0.005 | catastrophic pre-quant cost | + +Sweet spot at 0.20. didn't ship in this PR because it doesn't compose with #1855's `WARMDOWN_FRAC=0.85` recipe at our wallclock budget. + +### GPTQ block size + +| block_size | Quant gap | Notes | +|---|---|---| +| 64 | 0.050 (10× worse) | never use < 128 | +| **128** (default) | **0.005** | optimal | +| 256 | 0.007 | slightly worse | + +Default 128 is correct. don't change. + +--- + +## Per-group lrzip artifact + +Composition follows #1855 exactly: per-group bucketing + L1 sim-sort on hot 2D groups + lrzip ZPAQ on each group blob. + +| Component | Bytes | +|---|---:| +| Per-group lrzip quant blob | 15,920,473 | +| Pyminified `train_gpt.py` wrapper | 33,270 | +| **Total** | **15,953,743** | +| Cap | 16,000,000 | +| **Margin** | **46,257 (under cap)** | + +Roundtrip verified lossless: 275 quant tensors decompress byte-exact via Docker amd64 lrzip 0.651 (`scripts/pergroup_lrzip_recompress.py --roundtrip`). + +Eval host needs `apt install lrzip` (same as #1855). + +## Test plan + +- [x] TTT_LORA_RANK ablation at SEED=42 (5 ranks): 48 / 56 / 64 / 80 / 96. +- [x] Per-group lrzip artifact roundtrip verified lossless (275 tensors, byte-exact). +- [x] Total submission ≤ 16 MB (15,953,743 / 16,000,000 = 99.71 %). +- [x] H100 cross-hardware re-run matches MI250X result (596 s train + 599 s eval, both under 600 s lane caps; logs lost to pod crash). +- [ ] **SEED=0 + SEED=1234 at rank=56**: running on RunPod 8×H100, will append. + +## Files changed + +- `records/submission_2026_04_29_b180_tlr56_SUB106/` (new) + - `final_model.int6.ptz`: per-group lrzip quant blob (15,920,473 B) + - `train_gpt.py`: recipe. behavioural diff vs #1855 is `QK_GAIN_INIT=6.0` and `TTT_LORA_RANK=56`. source-level diff is bigger because this is our ROCm-ported lineage with FA fallback / inductor shims, but the H100 path is structurally equivalent. + - `submission.json`: metadata (val_bpb 1.05997, ablation tables, artifact byte breakdown) + - `lossless_caps.py`, `prepare_caseops_data.py`: CaseOps preprocessing (unchanged from #1855 lineage) + - `fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model`: tokenizer + - `lumi_train_ttt_17928157.log`: combined MI250X training + TTT eval log (rank=56 SEED=42, single SLURM job, ~37 KB) diff --git a/records/submission_2026_04_29_b180_tlr56_SUB106/final_model.int6.ptz b/records/submission_2026_04_29_b180_tlr56_SUB106/final_model.int6.ptz new file mode 100644 index 0000000000..ab1bdc9e91 Binary files /dev/null and b/records/submission_2026_04_29_b180_tlr56_SUB106/final_model.int6.ptz differ diff --git a/records/submission_2026_04_29_b180_tlr56_SUB106/fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model b/records/submission_2026_04_29_b180_tlr56_SUB106/fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model new file mode 100644 index 0000000000..fffc8bb306 Binary files /dev/null and b/records/submission_2026_04_29_b180_tlr56_SUB106/fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model differ diff --git a/records/submission_2026_04_29_b180_tlr56_SUB106/lossless_caps.py b/records/submission_2026_04_29_b180_tlr56_SUB106/lossless_caps.py new file mode 100644 index 0000000000..98e472f824 --- /dev/null +++ b/records/submission_2026_04_29_b180_tlr56_SUB106/lossless_caps.py @@ -0,0 +1,833 @@ +"""Lossless capitalization pre-encoding helpers. + +This module provides a narrow, reversible transform that only touches +ASCII capital letters `A-Z`. Each uppercase ASCII letter is rewritten as +``, where `sentinel` is a private-use Unicode +character that is escaped by doubling if it appears literally in the +input text. + +Example with the default sentinel `\\uE000`: + + "The NASA Launch" -> "\\uE000the \\uE000n\\uE000a\\uE000s\\uE000a \\uE000launch" + +The transform is intentionally simple for v1: + +- lowercase ASCII letters are unchanged +- uppercase ASCII letters become sentinel + lowercase letter +- non-ASCII characters are left untouched +- literal sentinel characters are escaped as sentinel + sentinel + +This makes the transform exactly invertible while allowing a downstream +tokenizer to reuse lowercase subwords across case variants. +""" + +from __future__ import annotations + +import json +from pathlib import Path +from typing import Callable, Iterable + +LOSSLESS_CAPS_V1 = "lossless_caps_v1" +LOSSLESS_CAPS_V2 = "lossless_caps_v2" +LOSSLESS_CAPS_V3 = "lossless_caps_v3" +LOSSLESS_CAPS_V4 = "lossless_caps_v4" +LOSSLESS_CAPS_V5 = "lossless_caps_v5" +LOSSLESS_CAPS_V6 = "lossless_caps_v6" +LOSSLESS_CAPS_V7 = "lossless_caps_v7" +LOSSLESS_CAPS_CASEOPS_V1 = "lossless_caps_caseops_v1" +IDENTITY = "identity" +DEFAULT_SENTINEL = "\uE000" +DEFAULT_V2_TITLE = "\uE001" +DEFAULT_V2_ALLCAPS = "\uE002" +DEFAULT_V2_CAPNEXT = "\uE003" +DEFAULT_V2_ESC = "\uE004" +DEFAULT_V5_TITLE_MIN_LEN = 7 +DEFAULT_V6_ALLCAPS_MIN_LEN = 3 +DEFAULT_V7_ALLCAPS_MIN_LEN = 4 + + +class LosslessCapsError(ValueError): + """Raised when a transformed string is malformed.""" + + +def _is_ascii_upper(ch: str) -> bool: + return "A" <= ch <= "Z" + + +def _is_ascii_lower(ch: str) -> bool: + return "a" <= ch <= "z" + + +def _is_ascii_alpha(ch: str) -> bool: + return _is_ascii_lower(ch) or _is_ascii_upper(ch) + + +def _validate_distinct_single_chars(*chars: str) -> None: + if any(len(ch) != 1 for ch in chars): + raise ValueError("all control characters must be exactly one character") + if len(set(chars)) != len(chars): + raise ValueError("control characters must be distinct") + + +def encode_lossless_caps_v1(text: str, *, sentinel: str = DEFAULT_SENTINEL) -> str: + """Encode ASCII capitals reversibly using a one-character sentinel.""" + if len(sentinel) != 1: + raise ValueError("sentinel must be exactly one character") + out: list[str] = [] + for ch in text: + if ch == sentinel: + out.append(sentinel) + out.append(sentinel) + elif _is_ascii_upper(ch): + out.append(sentinel) + out.append(ch.lower()) + else: + out.append(ch) + return "".join(out) + + +def decode_lossless_caps_v1(text: str, *, sentinel: str = DEFAULT_SENTINEL) -> str: + """Decode the `lossless_caps_v1` transform back to the original text.""" + if len(sentinel) != 1: + raise ValueError("sentinel must be exactly one character") + out: list[str] = [] + i = 0 + n = len(text) + while i < n: + ch = text[i] + if ch != sentinel: + out.append(ch) + i += 1 + continue + if i + 1 >= n: + raise LosslessCapsError("dangling capitalization sentinel at end of string") + nxt = text[i + 1] + if nxt == sentinel: + out.append(sentinel) + elif _is_ascii_lower(nxt): + out.append(nxt.upper()) + else: + raise LosslessCapsError( + f"invalid sentinel escape sequence {sentinel + nxt!r}; " + "expected doubled sentinel or sentinel + lowercase ASCII letter" + ) + i += 2 + return "".join(out) + + +def encode_lossless_caps_v2( + text: str, + *, + title: str = DEFAULT_V2_TITLE, + allcaps: str = DEFAULT_V2_ALLCAPS, + capnext: str = DEFAULT_V2_CAPNEXT, + esc: str = DEFAULT_V2_ESC, +) -> str: + """Encode ASCII word capitalization with cheap word-level markers. + + Rules over maximal ASCII alphabetic runs: + - lowercase words stay unchanged + - TitleCase words become `title + lowercase(word)` + - ALLCAPS words become `allcaps + lowercase(word)` + - mixed-case words use: + - optional `title` when the first letter is uppercase + - `capnext + lowercase(letter)` for subsequent uppercase letters + - literal control characters are escaped as `esc + literal` + """ + _validate_distinct_single_chars(title, allcaps, capnext, esc) + controls = {title, allcaps, capnext, esc} + out: list[str] = [] + i = 0 + n = len(text) + while i < n: + ch = text[i] + if ch in controls: + out.append(esc) + out.append(ch) + i += 1 + continue + if not _is_ascii_alpha(ch): + out.append(ch) + i += 1 + continue + + j = i + 1 + while j < n and _is_ascii_alpha(text[j]): + j += 1 + word = text[i:j] + lower_word = word.lower() + + if word.islower(): + out.append(word) + elif len(word) >= 2 and word.isupper(): + out.append(allcaps) + out.append(lower_word) + elif _is_ascii_upper(word[0]) and word[1:].islower(): + out.append(title) + out.append(lower_word) + else: + if _is_ascii_upper(word[0]): + out.append(title) + out.append(lower_word[0]) + for orig_ch, lower_ch in zip(word[1:], lower_word[1:], strict=True): + if _is_ascii_upper(orig_ch): + out.append(capnext) + out.append(lower_ch) + i = j + return "".join(out) + + +def decode_lossless_caps_v2( + text: str, + *, + title: str = DEFAULT_V2_TITLE, + allcaps: str = DEFAULT_V2_ALLCAPS, + capnext: str = DEFAULT_V2_CAPNEXT, + esc: str = DEFAULT_V2_ESC, +) -> str: + """Decode the `lossless_caps_v2` transform back to the original text.""" + _validate_distinct_single_chars(title, allcaps, capnext, esc) + out: list[str] = [] + pending_escape = False + pending_word_mode: str | None = None + active_allcaps = False + pending_capnext = False + in_ascii_word = False + + for ch in text: + if pending_escape: + if pending_word_mode is not None and not _is_ascii_alpha(ch): + raise LosslessCapsError("escaped control char cannot satisfy pending word capitalization mode") + out.append(ch) + pending_escape = False + if _is_ascii_alpha(ch): + in_ascii_word = True + else: + in_ascii_word = False + active_allcaps = False + continue + + if ch == esc: + pending_escape = True + continue + if ch == title: + if pending_word_mode is not None or in_ascii_word or pending_capnext: + raise LosslessCapsError("invalid title marker placement") + pending_word_mode = "title" + continue + if ch == allcaps: + if pending_word_mode is not None or in_ascii_word or pending_capnext: + raise LosslessCapsError("invalid allcaps marker placement") + pending_word_mode = "allcaps" + continue + if ch == capnext: + if pending_capnext: + raise LosslessCapsError("duplicate capnext marker") + pending_capnext = True + continue + + if _is_ascii_alpha(ch): + at_word_start = not in_ascii_word + if at_word_start: + if pending_word_mode == "allcaps": + out.append(ch.upper()) + active_allcaps = True + elif pending_word_mode == "title": + out.append(ch.upper()) + elif pending_capnext: + out.append(ch.upper()) + else: + out.append(ch) + pending_word_mode = None + pending_capnext = False + in_ascii_word = True + continue + + if pending_word_mode is not None: + raise LosslessCapsError("word capitalization marker leaked into the middle of a word") + if active_allcaps: + out.append(ch.upper()) + elif pending_capnext: + out.append(ch.upper()) + else: + out.append(ch) + pending_capnext = False + continue + + if pending_word_mode is not None or pending_capnext: + raise LosslessCapsError("capitalization marker not followed by an ASCII letter") + out.append(ch) + in_ascii_word = False + active_allcaps = False + + if pending_escape: + raise LosslessCapsError("dangling escape marker at end of string") + if pending_word_mode is not None or pending_capnext: + raise LosslessCapsError("dangling capitalization marker at end of string") + return "".join(out) + + +def encode_lossless_caps_v3( + text: str, + *, + title: str = DEFAULT_V2_TITLE, + allcaps: str = DEFAULT_V2_ALLCAPS, + esc: str = DEFAULT_V2_ESC, +) -> str: + """Encode only common word-level capitalization patterns. + + Rules over maximal ASCII alphabetic runs: + - lowercase words stay unchanged + - TitleCase words become `title + lowercase(word)` + - ALLCAPS words become `allcaps + lowercase(word)` + - all other mixed-case words are left unchanged + - literal control characters are escaped as `esc + literal` + """ + _validate_distinct_single_chars(title, allcaps, esc) + controls = {title, allcaps, esc} + out: list[str] = [] + i = 0 + n = len(text) + while i < n: + ch = text[i] + if ch in controls: + out.append(esc) + out.append(ch) + i += 1 + continue + if not _is_ascii_alpha(ch): + out.append(ch) + i += 1 + continue + + j = i + 1 + while j < n and _is_ascii_alpha(text[j]): + j += 1 + word = text[i:j] + + if word.islower(): + out.append(word) + elif len(word) >= 2 and word.isupper(): + out.append(allcaps) + out.append(word.lower()) + elif _is_ascii_upper(word[0]) and word[1:].islower(): + out.append(title) + out.append(word.lower()) + else: + out.append(word) + i = j + return "".join(out) + + +def decode_lossless_caps_v3( + text: str, + *, + title: str = DEFAULT_V2_TITLE, + allcaps: str = DEFAULT_V2_ALLCAPS, + esc: str = DEFAULT_V2_ESC, +) -> str: + """Decode the `lossless_caps_v3` transform back to the original text.""" + _validate_distinct_single_chars(title, allcaps, esc) + out: list[str] = [] + pending_escape = False + pending_word_mode: str | None = None + active_allcaps = False + in_ascii_word = False + + for ch in text: + if pending_escape: + if pending_word_mode is not None and not _is_ascii_alpha(ch): + raise LosslessCapsError("escaped control char cannot satisfy pending word capitalization mode") + out.append(ch) + pending_escape = False + if _is_ascii_alpha(ch): + in_ascii_word = True + else: + in_ascii_word = False + active_allcaps = False + continue + + if ch == esc: + pending_escape = True + continue + if ch == title: + if pending_word_mode is not None or in_ascii_word: + raise LosslessCapsError("invalid title marker placement") + pending_word_mode = "title" + continue + if ch == allcaps: + if pending_word_mode is not None or in_ascii_word: + raise LosslessCapsError("invalid allcaps marker placement") + pending_word_mode = "allcaps" + continue + + if _is_ascii_alpha(ch): + at_word_start = not in_ascii_word + if at_word_start: + if pending_word_mode == "allcaps": + out.append(ch.upper()) + active_allcaps = True + elif pending_word_mode == "title": + out.append(ch.upper()) + else: + out.append(ch) + pending_word_mode = None + in_ascii_word = True + continue + + if pending_word_mode is not None: + raise LosslessCapsError("word capitalization marker leaked into the middle of a word") + out.append(ch.upper() if active_allcaps else ch) + continue + + if pending_word_mode is not None: + raise LosslessCapsError("capitalization marker not followed by an ASCII letter") + out.append(ch) + in_ascii_word = False + active_allcaps = False + + if pending_escape: + raise LosslessCapsError("dangling escape marker at end of string") + if pending_word_mode is not None: + raise LosslessCapsError("dangling capitalization marker at end of string") + return "".join(out) + + +def encode_lossless_caps_v4( + text: str, + *, + allcaps: str = DEFAULT_V2_ALLCAPS, + esc: str = DEFAULT_V2_ESC, +) -> str: + """Encode only ALLCAPS ASCII words, leaving all other case untouched.""" + _validate_distinct_single_chars(allcaps, esc) + controls = {allcaps, esc} + out: list[str] = [] + i = 0 + n = len(text) + while i < n: + ch = text[i] + if ch in controls: + out.append(esc) + out.append(ch) + i += 1 + continue + if not _is_ascii_alpha(ch): + out.append(ch) + i += 1 + continue + j = i + 1 + while j < n and _is_ascii_alpha(text[j]): + j += 1 + word = text[i:j] + if len(word) >= 2 and word.isupper(): + out.append(allcaps) + out.append(word.lower()) + else: + out.append(word) + i = j + return "".join(out) + + +def decode_lossless_caps_v4( + text: str, + *, + allcaps: str = DEFAULT_V2_ALLCAPS, + esc: str = DEFAULT_V2_ESC, +) -> str: + """Decode the `lossless_caps_v4` transform back to the original text.""" + _validate_distinct_single_chars(allcaps, esc) + out: list[str] = [] + pending_escape = False + pending_allcaps = False + in_ascii_word = False + active_allcaps = False + + for ch in text: + if pending_escape: + if pending_allcaps and not _is_ascii_alpha(ch): + raise LosslessCapsError("escaped control char cannot satisfy pending allcaps mode") + out.append(ch) + pending_escape = False + if _is_ascii_alpha(ch): + in_ascii_word = True + else: + in_ascii_word = False + active_allcaps = False + continue + + if ch == esc: + pending_escape = True + continue + if ch == allcaps: + if pending_allcaps or in_ascii_word: + raise LosslessCapsError("invalid allcaps marker placement") + pending_allcaps = True + continue + + if _is_ascii_alpha(ch): + if not in_ascii_word: + active_allcaps = pending_allcaps + pending_allcaps = False + in_ascii_word = True + out.append(ch.upper() if active_allcaps else ch) + continue + + if pending_allcaps: + raise LosslessCapsError("allcaps marker not followed by an ASCII letter") + out.append(ch) + in_ascii_word = False + active_allcaps = False + + if pending_escape: + raise LosslessCapsError("dangling escape marker at end of string") + if pending_allcaps: + raise LosslessCapsError("dangling allcaps marker at end of string") + return "".join(out) + + +def encode_lossless_caps_v5( + text: str, + *, + title: str = DEFAULT_V2_TITLE, + allcaps: str = DEFAULT_V2_ALLCAPS, + esc: str = DEFAULT_V2_ESC, + title_min_len: int = DEFAULT_V5_TITLE_MIN_LEN, +) -> str: + """Encode ALLCAPS words and only sufficiently long TitleCase words.""" + _validate_distinct_single_chars(title, allcaps, esc) + controls = {title, allcaps, esc} + out: list[str] = [] + i = 0 + n = len(text) + while i < n: + ch = text[i] + if ch in controls: + out.append(esc) + out.append(ch) + i += 1 + continue + if not _is_ascii_alpha(ch): + out.append(ch) + i += 1 + continue + j = i + 1 + while j < n and _is_ascii_alpha(text[j]): + j += 1 + word = text[i:j] + if len(word) >= 2 and word.isupper(): + out.append(allcaps) + out.append(word.lower()) + elif len(word) >= title_min_len and _is_ascii_upper(word[0]) and word[1:].islower(): + out.append(title) + out.append(word.lower()) + else: + out.append(word) + i = j + return "".join(out) + + +def decode_lossless_caps_v5( + text: str, + *, + title: str = DEFAULT_V2_TITLE, + allcaps: str = DEFAULT_V2_ALLCAPS, + esc: str = DEFAULT_V2_ESC, +) -> str: + """Decode the `lossless_caps_v5` transform back to the original text.""" + return decode_lossless_caps_v3(text, title=title, allcaps=allcaps, esc=esc) + + +def encode_lossless_caps_v6( + text: str, + *, + allcaps: str = DEFAULT_V2_ALLCAPS, + esc: str = DEFAULT_V2_ESC, + allcaps_min_len: int = DEFAULT_V6_ALLCAPS_MIN_LEN, +) -> str: + """Encode only ALLCAPS words with length >= allcaps_min_len.""" + _validate_distinct_single_chars(allcaps, esc) + controls = {allcaps, esc} + out: list[str] = [] + i = 0 + n = len(text) + while i < n: + ch = text[i] + if ch in controls: + out.append(esc) + out.append(ch) + i += 1 + continue + if not _is_ascii_alpha(ch): + out.append(ch) + i += 1 + continue + j = i + 1 + while j < n and _is_ascii_alpha(text[j]): + j += 1 + word = text[i:j] + if len(word) >= allcaps_min_len and word.isupper(): + out.append(allcaps) + out.append(word.lower()) + else: + out.append(word) + i = j + return "".join(out) + + +def decode_lossless_caps_v6( + text: str, + *, + allcaps: str = DEFAULT_V2_ALLCAPS, + esc: str = DEFAULT_V2_ESC, +) -> str: + """Decode the `lossless_caps_v6` transform back to the original text.""" + return decode_lossless_caps_v4(text, allcaps=allcaps, esc=esc) + + +def encode_lossless_caps_v7( + text: str, + *, + allcaps: str = DEFAULT_V2_ALLCAPS, + esc: str = DEFAULT_V2_ESC, + allcaps_min_len: int = DEFAULT_V7_ALLCAPS_MIN_LEN, +) -> str: + """Encode only ALLCAPS words with length >= 4.""" + return encode_lossless_caps_v6( + text, + allcaps=allcaps, + esc=esc, + allcaps_min_len=allcaps_min_len, + ) + + +def decode_lossless_caps_v7( + text: str, + *, + allcaps: str = DEFAULT_V2_ALLCAPS, + esc: str = DEFAULT_V2_ESC, +) -> str: + """Decode the `lossless_caps_v7` transform back to the original text.""" + return decode_lossless_caps_v6(text, allcaps=allcaps, esc=esc) + + +def get_text_transform(name: str | None) -> Callable[[str], str]: + """Return the forward text transform for the given config name.""" + normalized = IDENTITY if name in {None, "", IDENTITY} else str(name) + if normalized == IDENTITY: + return lambda text: text + if normalized == LOSSLESS_CAPS_V1: + return encode_lossless_caps_v1 + if normalized == LOSSLESS_CAPS_V2: + return encode_lossless_caps_v2 + if normalized == LOSSLESS_CAPS_V3: + return encode_lossless_caps_v3 + if normalized == LOSSLESS_CAPS_V4: + return encode_lossless_caps_v4 + if normalized == LOSSLESS_CAPS_V5: + return encode_lossless_caps_v5 + if normalized == LOSSLESS_CAPS_V6: + return encode_lossless_caps_v6 + if normalized == LOSSLESS_CAPS_V7: + return encode_lossless_caps_v7 + if normalized == LOSSLESS_CAPS_CASEOPS_V1: + return encode_lossless_caps_v2 + raise ValueError(f"unsupported text_transform={name!r}") + + +def get_text_inverse_transform(name: str | None) -> Callable[[str], str]: + """Return the inverse transform for the given config name.""" + normalized = IDENTITY if name in {None, "", IDENTITY} else str(name) + if normalized == IDENTITY: + return lambda text: text + if normalized == LOSSLESS_CAPS_V1: + return decode_lossless_caps_v1 + if normalized == LOSSLESS_CAPS_V2: + return decode_lossless_caps_v2 + if normalized == LOSSLESS_CAPS_V3: + return decode_lossless_caps_v3 + if normalized == LOSSLESS_CAPS_V4: + return decode_lossless_caps_v4 + if normalized == LOSSLESS_CAPS_V5: + return decode_lossless_caps_v5 + if normalized == LOSSLESS_CAPS_V6: + return decode_lossless_caps_v6 + if normalized == LOSSLESS_CAPS_V7: + return decode_lossless_caps_v7 + if normalized == LOSSLESS_CAPS_CASEOPS_V1: + return decode_lossless_caps_v2 + raise ValueError(f"unsupported text_transform={name!r}") + + +def normalize_text_transform_name(name: str | None) -> str: + """Normalize empty/None transform names to the identity transform.""" + return IDENTITY if name in {None, "", IDENTITY} else str(name) + + +def get_text_transform_control_symbols(name: str | None) -> list[str]: + """Return reserved control symbols used by a transform, if any.""" + normalized = normalize_text_transform_name(name) + if normalized == IDENTITY: + return [] + if normalized == LOSSLESS_CAPS_V1: + return [DEFAULT_SENTINEL] + if normalized == LOSSLESS_CAPS_V2: + return [DEFAULT_V2_TITLE, DEFAULT_V2_ALLCAPS, DEFAULT_V2_CAPNEXT, DEFAULT_V2_ESC] + if normalized == LOSSLESS_CAPS_CASEOPS_V1: + return [DEFAULT_V2_TITLE, DEFAULT_V2_ALLCAPS, DEFAULT_V2_CAPNEXT, DEFAULT_V2_ESC] + if normalized in {LOSSLESS_CAPS_V3, LOSSLESS_CAPS_V5}: + return [DEFAULT_V2_TITLE, DEFAULT_V2_ALLCAPS, DEFAULT_V2_ESC] + if normalized in {LOSSLESS_CAPS_V4, LOSSLESS_CAPS_V6, LOSSLESS_CAPS_V7}: + return [DEFAULT_V2_ALLCAPS, DEFAULT_V2_ESC] + raise ValueError(f"unsupported text_transform={name!r}") + + +def infer_text_transform_from_manifest(tokenizer_path: str | Path) -> str: + """Best-effort lookup of a tokenizer's text transform from a local manifest.""" + tokenizer_path = Path(tokenizer_path).expanduser().resolve() + manifest_candidates = [ + tokenizer_path.parent.parent / "manifest.json", + tokenizer_path.parent / "manifest.json", + ] + for manifest_path in manifest_candidates: + if not manifest_path.is_file(): + continue + try: + payload = json.loads(manifest_path.read_text(encoding="utf-8")) + except (OSError, json.JSONDecodeError): + continue + tokenizers = payload.get("tokenizers") + if not isinstance(tokenizers, list): + continue + for tokenizer_meta in tokenizers: + if not isinstance(tokenizer_meta, dict): + continue + model_path = tokenizer_meta.get("model_path") or tokenizer_meta.get("path") + if not model_path: + continue + candidate = (manifest_path.parent / str(model_path)).resolve() + if candidate == tokenizer_path: + return normalize_text_transform_name(tokenizer_meta.get("text_transform")) + return IDENTITY + + +def surface_piece_original_byte_counts( + surfaces: Iterable[str], + *, + text_transform_name: str | None = None, + sentinel: str = DEFAULT_SENTINEL, +) -> list[int]: + """Return exact original UTF-8 byte counts contributed by each surface piece. + + `surfaces` must be the exact decoded text fragments emitted by SentencePiece + in order, e.g. `piece.surface` from `encode_as_immutable_proto`. + """ + normalized = normalize_text_transform_name(text_transform_name) + if normalized == IDENTITY: + return [len(surface.encode("utf-8")) for surface in surfaces] + if normalized == LOSSLESS_CAPS_V1: + if len(sentinel) != 1: + raise ValueError("sentinel must be exactly one character") + sentinel_bytes = len(sentinel.encode("utf-8")) + pending_sentinel = False + counts: list[int] = [] + for surface in surfaces: + piece_bytes = 0 + for ch in surface: + if pending_sentinel: + if ch == sentinel: + piece_bytes += sentinel_bytes + elif _is_ascii_lower(ch): + piece_bytes += 1 + else: + raise LosslessCapsError( + f"invalid continuation {ch!r} after capitalization sentinel" + ) + pending_sentinel = False + continue + if ch == sentinel: + pending_sentinel = True + else: + piece_bytes += len(ch.encode("utf-8")) + counts.append(piece_bytes) + if pending_sentinel: + raise LosslessCapsError("dangling capitalization sentinel across piece boundary") + return counts + if normalized not in {LOSSLESS_CAPS_V2, LOSSLESS_CAPS_V3, LOSSLESS_CAPS_V4, LOSSLESS_CAPS_V5, LOSSLESS_CAPS_V6, LOSSLESS_CAPS_V7, LOSSLESS_CAPS_CASEOPS_V1}: + raise ValueError(f"unsupported text_transform={text_transform_name!r}") + + title = DEFAULT_V2_TITLE + allcaps = DEFAULT_V2_ALLCAPS + capnext = DEFAULT_V2_CAPNEXT + esc = DEFAULT_V2_ESC + if normalized in {LOSSLESS_CAPS_V2, LOSSLESS_CAPS_CASEOPS_V1}: + _validate_distinct_single_chars(title, allcaps, capnext, esc) + elif normalized in {LOSSLESS_CAPS_V4, LOSSLESS_CAPS_V6, LOSSLESS_CAPS_V7}: + _validate_distinct_single_chars(allcaps, esc) + else: + _validate_distinct_single_chars(title, allcaps, esc) + pending_escape = False + pending_word_mode: str | None = None + active_allcaps = False + pending_capnext = False + in_ascii_word = False + counts: list[int] = [] + for surface in surfaces: + piece_bytes = 0 + for ch in surface: + if pending_escape: + if pending_word_mode is not None and not _is_ascii_alpha(ch): + raise LosslessCapsError("escaped control char cannot satisfy pending word capitalization mode") + piece_bytes += len(ch.encode("utf-8")) + pending_escape = False + if _is_ascii_alpha(ch): + in_ascii_word = True + else: + in_ascii_word = False + active_allcaps = False + continue + if ch == esc: + pending_escape = True + continue + if normalized in {LOSSLESS_CAPS_V2, LOSSLESS_CAPS_V3, LOSSLESS_CAPS_V5, LOSSLESS_CAPS_CASEOPS_V1} and ch == title: + if pending_word_mode is not None or in_ascii_word or pending_capnext: + raise LosslessCapsError("invalid title marker placement") + pending_word_mode = "title" + continue + if ch == allcaps: + if pending_word_mode is not None or in_ascii_word or pending_capnext: + raise LosslessCapsError("invalid allcaps marker placement") + pending_word_mode = "allcaps" + continue + if normalized in {LOSSLESS_CAPS_V2, LOSSLESS_CAPS_CASEOPS_V1} and ch == capnext: + if pending_capnext: + raise LosslessCapsError("duplicate capnext marker") + pending_capnext = True + continue + + if _is_ascii_alpha(ch): + at_word_start = not in_ascii_word + if at_word_start: + piece_bytes += 1 + active_allcaps = pending_word_mode == "allcaps" + pending_word_mode = None + pending_capnext = False + in_ascii_word = True + continue + if pending_word_mode is not None: + raise LosslessCapsError("word capitalization marker leaked into the middle of a word") + piece_bytes += 1 + pending_capnext = False + continue + + if pending_word_mode is not None or pending_capnext: + raise LosslessCapsError("capitalization marker not followed by an ASCII letter") + piece_bytes += len(ch.encode("utf-8")) + in_ascii_word = False + active_allcaps = False + counts.append(piece_bytes) + if pending_escape: + raise LosslessCapsError("dangling escape marker across piece boundary") + if pending_word_mode is not None or pending_capnext: + raise LosslessCapsError("dangling capitalization marker across piece boundary") + return counts diff --git a/records/submission_2026_04_29_b180_tlr56_SUB106/lumi_train_ttt_17928157.log b/records/submission_2026_04_29_b180_tlr56_SUB106/lumi_train_ttt_17928157.log new file mode 100644 index 0000000000..200b6baeef --- /dev/null +++ b/records/submission_2026_04_29_b180_tlr56_SUB106/lumi_train_ttt_17928157.log @@ -0,0 +1,979 @@ +Setting up environment: + CODE_DIR=/users/vtoivone/parameter-golf + SCRATCH_BASE=/scratch/project_462001391/vtoivone/pgolf_cache + PROJAPPL_BASE=/projappl/project_462001391/vtoivone +Environment configured: + PYTHONUSERBASE=/projappl/project_462001391/vtoivone/python_user + HF_HOME=/scratch/project_462001391/vtoivone/pgolf_cache/hf_cache + MIOPEN_USER_DB_PATH=/tmp/vtoivone-miopen-cache-0 + CC=/opt/rocm-6.2.4/llvm/bin/clang + CXX=/opt/rocm-6.2.4/llvm/bin/clang++ +=========================================== +Parameter Golf Training +=========================================== +Job ID: 17928157 +GPUs: 8 +DATA_PATH: /scratch/project_462001391/vtoivone/pgolf_cache/datasets/fineweb10B_caseops +TOKENIZER_PATH: /scratch/project_462001391/vtoivone/pgolf_cache/tokenizers/fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model + +Hyperparameters: + adam_eps: 1e-08 + adam_wd: 0.02 + artifact_dir: /users/vtoivone/parameter-golf/artifacts/b180_tlr56_s42_v4 + attn_clip_sigmas: 13.0 + beta1: 0.9 + beta2: 0.99 + compressor: brotli + data_dir: ./data/ + datasets_dir: /scratch/project_462001391/vtoivone/pgolf_cache/datasets/fineweb10B_caseops + distributed: True + ema_decay: 0.9965 + embed_bits: 7 + embed_clip_sigmas: 14.0 + embed_lr: 0.6 + embed_wd: 0.085 + enable_looping_at: 0.35 + eval_seq_len: 2048 + eval_stride: 64 + gate_attn_out: False + gate_attn_width: 24 + gate_window: 12 + global_ttt_batch_seqs: 32 + global_ttt_chunk_tokens: 32768 + global_ttt_epochs: 1 + global_ttt_grad_clip: 1.0 + global_ttt_lr: 0.001 + global_ttt_momentum: 0.9 + global_ttt_respect_doc_boundaries: True + global_ttt_warmup_chunks: 0 + global_ttt_warmup_start_lr: 0.0 + gptq_calibration_batches: 16 + gptq_reserve_seconds: 4.0 + grad_accum_steps: 1 + grad_clip_norm: 0.3 + is_main_process: True + iterations: 20000 + ln_scale: True + local_rank: 0 + logfile: /users/vtoivone/parameter-golf/artifacts/b180_tlr56_s42_v4/b180_tlr56_s42.txt + logit_softcap: 30.0 + loop_end: 5 + loop_start: 3 + lqer_adaptive: True + lqer_asym_enabled: True + lqer_asym_group: 64 + lqer_budget_bytes: 80000 + lqer_cand_ranks: 4_6 + lqer_enabled: True + lqer_factor_bits: 4 + lqer_rank: 4 + lqer_top_k: 3 + matrix_bits: 6 + matrix_clip_sigmas: 12.85 + matrix_lr: 0.026 + max_wallclock_seconds: 3400.0 + min_lr: 0.1 + mlp_clip_sigmas: 11.5 + mlp_mult: 4.0 + model_dim: 512 + model_path: /users/vtoivone/parameter-golf/artifacts/b180_tlr56_s42_v4/final_model.pt + muon_backend_steps: 5 + muon_momentum: 0.97 + muon_momentum_warmup_start: 0.92 + muon_momentum_warmup_steps: 1500 + muon_row_normalize: True + muon_wd: 0.095 + num_heads: 8 + num_kv_heads: 4 + num_layers: 11 + num_loops: 2 + parallel_final_lane: mean + parallel_start_layer: 8 + phased_ttt_enabled: True + phased_ttt_num_phases: 3 + phased_ttt_prefix_docs: 2500 + qk_gain_init: 6.0 + quantized_model_path: /users/vtoivone/parameter-golf/artifacts/b180_tlr56_s42_v4/final_model.int6.ptz + rank: 0 + rope_base: 10000.0 + rope_dims: 16 + rope_train_seq_len: 2048 + rope_yarn: False + run_id: b180_tlr56_s42 + scalar_lr: 0.02 + seed: 42 + skip_gates_enabled: True + sliding_window_enabled: False + smear_gate_enabled: True + smear_gate_width: 12 + sparse_attn_gate_enabled: True + sparse_attn_gate_scale: 0.5 + tie_embeddings: True + tied_embed_init_std: 0.005 + tied_embed_lr: 0.03 + tokenizer_path: /scratch/project_462001391/vtoivone/pgolf_cache/tokenizers/fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model + train_batch_tokens: 786432 + train_files: /scratch/project_462001391/vtoivone/pgolf_cache/datasets/fineweb10B_caseops/fineweb_train_[0-9]*.bin + train_log_every: 500 + train_seq_len: 2048 + ttt_batch_size: 64 + ttt_beta1: 0.0 + ttt_beta2: 0.99 + ttt_chunk_size: 32 + ttt_enabled: True + ttt_eval_batches: + ttt_eval_seq_len: 2048 + ttt_grad_steps: 1 + ttt_k_lora: True + ttt_lora_lr: 0.0001 + ttt_lora_rank: 56 + ttt_mlp_lora: True + ttt_o_lora: True + ttt_optimizer: adam + ttt_weight_decay: 0.5 + val_batch_tokens: 524288 + val_doc_fraction: 1.0 + val_files: /scratch/project_462001391/vtoivone/pgolf_cache/datasets/fineweb10B_caseops/fineweb_val_[0-9]*.bin + val_loss_every: 4000 + vocab_size: 8192 + warmdown_frac: 0.85 + warmup_steps: 20 + world_size: 8 + xsa_last_n: 11 +build_sentencepiece_luts: 4 PUE-marker tokens zeroed out +build_sentencepiece_luts: 4 PUE-marker tokens zeroed out +build_sentencepiece_luts: 4 PUE-marker tokens zeroed out +build_sentencepiece_luts: 4 PUE-marker tokens zeroed out +build_sentencepiece_luts: 4 PUE-marker tokens zeroed out +train_shards: 150 +val_tokens: 47851520 +build_sentencepiece_luts: 4 PUE-marker tokens zeroed out +build_sentencepiece_luts: 4 PUE-marker tokens zeroed out +build_sentencepiece_luts: 4 PUE-marker tokens zeroed out +model_params:35945671 +gptq:reserving 4s, effective=3396000ms +warmup_cu_buckets:64,128,192,256 iters_each:3 +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: 10/20 +warmup_step: 20/20 +loop_warmup:enabled encoder:[0, 1, 2, 3, 4, 5, 3, 4] decoder:[5, 3, 4, 5, 6, 7, 8, 9, 10] +loop_warmup_step: 1/20 +loop_warmup_step: 2/20 +loop_warmup_step: 3/20 +loop_warmup_step: 4/20 +loop_warmup_step: 5/20 +loop_warmup_step: 6/20 +loop_warmup_step: 10/20 +loop_warmup_step: 20/20 +0/20000 val_loss: 9.0076 val_bpb: 4.1159 +1/20000 train_loss: 9.0087 train_time: 0.0m tok/s: 11702455 +2/20000 train_loss: 12.8129 train_time: 0.0m tok/s: 2838208 +3/20000 train_loss: 10.2217 train_time: 0.0m tok/s: 2173916 +4/20000 train_loss: 8.6794 train_time: 0.0m tok/s: 1936572 +5/20000 train_loss: 7.9367 train_time: 0.0m tok/s: 1806449 +500/20000 train_loss: 2.5748 train_time: 4.5m tok/s: 1469725 +1000/20000 train_loss: 2.8010 train_time: 8.9m tok/s: 1468517 +1500/20000 train_loss: 2.6241 train_time: 13.4m tok/s: 1467981 +2000/20000 train_loss: 2.6577 train_time: 17.9m tok/s: 1467467 +layer_loop:enabled step:2217 frac:0.350 encoder:[0, 1, 2, 3, 4, 5, 3, 4] decoder:[5, 3, 4, 5, 6, 7, 8, 9, 10] +2500/20000 train_loss: 2.5439 train_time: 23.5m tok/s: 1392410 +3000/20000 train_loss: 2.5596 train_time: 30.1m tok/s: 1305106 +3500/20000 train_loss: 2.5643 train_time: 36.7m tok/s: 1248716 +4000/20000 train_loss: 2.4106 train_time: 43.3m tok/s: 1209902 +4000/20000 val_loss: 2.4312 val_bpb: 1.1109 +4500/20000 train_loss: 2.2849 train_time: 49.9m tok/s: 1181486 +5000/20000 train_loss: 2.4412 train_time: 56.5m tok/s: 1159728 +5007/20000 val_loss: 2.3546 val_bpb: 1.0759 +stopping_early: wallclock_cap train_time: 3396829ms step: 5007/20000 +peak memory allocated: 41211 MiB reserved: 46604 MiB +ema:applying EMA weights +diagnostic pre-quantization post-ema val_loss:2.32876861 val_bpb:1.06408552 eval_time:18050ms +Serialized model: 135417533 bytes +Code size (uncompressed): 142470 bytes +Code size (compressed): 35985 bytes +GPTQ:collecting Hessians from calibration data... +GPTQ:collected 67 Hessians in 30.9s +LQER adaptive: budget=80000 bytes, used=79944 bytes, tensors=20 + blocks.0.mlp.fc.weight: rank=6 bytes=4706 + blocks.1.mlp.fc.weight: rank=6 bytes=4706 + blocks.10.attn.c_v.weight: rank=6 bytes=2018 + blocks.10.attn.proj.weight: rank=6 bytes=2402 + blocks.10.mlp.fc.weight: rank=4 bytes=3138 + blocks.2.mlp.fc.weight: rank=6 bytes=4706 + blocks.3.mlp.fc.weight: rank=4 bytes=3138 + blocks.4.attn.proj.weight: rank=6 bytes=2402 + blocks.4.mlp.fc.weight: rank=6 bytes=4706 + blocks.5.mlp.fc.weight: rank=6 bytes=4706 + blocks.6.attn.proj.weight: rank=6 bytes=2402 + blocks.6.mlp.fc.weight: rank=6 bytes=4706 + blocks.7.mlp.fc.weight: rank=6 bytes=4706 + blocks.8.attn.proj.weight: rank=6 bytes=2402 + blocks.8.mlp.fc.weight: rank=6 bytes=4706 + blocks.9.attn.c_k.weight: rank=4 bytes=1346 + blocks.9.attn.c_v.weight: rank=6 bytes=2018 + blocks.9.attn.proj.weight: rank=6 bytes=2402 + blocks.9.mlp.fc.weight: rank=6 bytes=4706 + tok_emb.weight: rank=6 bytes=13922 +Quantized weights: + gptq (int6): blocks.attn.c_k.weight, blocks.attn.c_q.weight, blocks.attn.c_v.weight, blocks.attn.proj.weight, blocks.mlp.proj.weight + gptq (int6)+lqer_asym_r4: blocks.attn.c_k.weight, blocks.mlp.fc.weight + gptq (int6)+lqer_asym_r6: blocks.attn.c_v.weight, blocks.attn.proj.weight, blocks.mlp.fc.weight + gptq (int7)+lqer_asym_r6: tok_emb.weight + passthrough (float16): blocks.attn.attn_gate_w, blocks.attn.q_gain, blocks.attn_scale, blocks.mlp_scale, blocks.resid_mix, parallel_post_lambdas, parallel_resid_lambdas, skip_gates, skip_weights, smear_gate.weight, smear_lambda +Serialized model quantized+brotli: 16156886 bytes +Total submission size quantized+brotli: 16192871 bytes +diagnostic quantized val_loss:2.34799135 val_bpb:1.07286898 eval_time:30243ms +ttt_lora:warming up compile (random tokens, no val data) +ttt_lora:compile warmup done (62.9s) + +beginning TTT eval timer +ttt_phased: total_docs:50000 prefix_docs:2500 suffix_docs:47500 num_phases:3 boundaries:[833, 1666, 2500] +ttp: b779/782 bl:2.2249 bb:1.0526 rl:2.2249 rb:1.0526 dl:10442-13079 gd:0 +ttp: b770/782 bl:2.2860 bb:1.0793 rl:2.2442 rb:1.0610 dl:5311-5522 gd:0 +ttp: b764/782 bl:2.2905 bb:1.0729 rl:2.2536 rb:1.0635 dl:4284-4392 gd:0 +ttpp: phase:1/3 pd:1296 gd:833 t:907.3s +tttg: c1/131 lr:0.001000 t:0.1s +tttg: c2/131 lr:0.001000 t:1.6s +tttg: c3/131 lr:0.000999 t:1.7s +tttg: c4/131 lr:0.000999 t:1.9s +tttg: c5/131 lr:0.000998 t:2.0s +tttg: c6/131 lr:0.000996 t:2.1s +tttg: c7/131 lr:0.000995 t:2.3s +tttg: c8/131 lr:0.000993 t:2.4s +tttg: c9/131 lr:0.000991 t:2.5s +tttg: c10/131 lr:0.000988 t:2.7s +tttg: c11/131 lr:0.000985 t:2.8s +tttg: c12/131 lr:0.000982 t:2.9s +tttg: c13/131 lr:0.000979 t:3.1s +tttg: c14/131 lr:0.000976 t:3.2s +tttg: c15/131 lr:0.000972 t:3.3s +tttg: c16/131 lr:0.000968 t:3.5s +tttg: c17/131 lr:0.000963 t:3.6s +tttg: c18/131 lr:0.000958 t:3.7s +tttg: c19/131 lr:0.000953 t:3.9s +tttg: c20/131 lr:0.000948 t:4.0s +tttg: c21/131 lr:0.000943 t:4.1s +tttg: c22/131 lr:0.000937 t:4.3s +tttg: c23/131 lr:0.000931 t:4.4s +tttg: c24/131 lr:0.000925 t:4.5s +tttg: c25/131 lr:0.000918 t:4.7s +tttg: c26/131 lr:0.000911 t:4.8s +tttg: c27/131 lr:0.000905 t:4.9s +tttg: c28/131 lr:0.000897 t:5.0s +tttg: c29/131 lr:0.000890 t:5.2s +tttg: c30/131 lr:0.000882 t:5.3s +tttg: c31/131 lr:0.000874 t:5.4s +tttg: c32/131 lr:0.000866 t:5.6s +tttg: c33/131 lr:0.000858 t:5.7s +tttg: c34/131 lr:0.000849 t:5.8s +tttg: c35/131 lr:0.000841 t:6.0s +tttg: c36/131 lr:0.000832 t:6.1s +tttg: c37/131 lr:0.000822 t:6.2s +tttg: c38/131 lr:0.000813 t:6.4s +tttg: c39/131 lr:0.000804 t:6.5s +tttg: c40/131 lr:0.000794 t:6.6s +tttg: c41/131 lr:0.000784 t:6.8s +tttg: c42/131 lr:0.000774 t:6.9s +tttg: c43/131 lr:0.000764 t:7.0s +tttg: c44/131 lr:0.000753 t:7.2s +tttg: c45/131 lr:0.000743 t:7.3s +tttg: c46/131 lr:0.000732 t:7.4s +tttg: c47/131 lr:0.000722 t:7.6s +tttg: c48/131 lr:0.000711 t:7.7s +tttg: c49/131 lr:0.000700 t:7.8s +tttg: c50/131 lr:0.000689 t:8.0s +tttg: c51/131 lr:0.000677 t:8.1s +tttg: c52/131 lr:0.000666 t:8.2s +tttg: c53/131 lr:0.000655 t:8.4s +tttg: c54/131 lr:0.000643 t:8.5s +tttg: c55/131 lr:0.000631 t:8.6s +tttg: c56/131 lr:0.000620 t:8.7s +tttg: c57/131 lr:0.000608 t:8.9s +tttg: c58/131 lr:0.000596 t:9.0s +tttg: c59/131 lr:0.000584 t:9.1s +tttg: c60/131 lr:0.000572 t:9.3s +tttg: c61/131 lr:0.000560 t:9.4s +tttg: c62/131 lr:0.000548 t:9.5s +tttg: c63/131 lr:0.000536 t:9.7s +tttg: c64/131 lr:0.000524 t:9.8s +tttg: c65/131 lr:0.000512 t:9.9s +tttg: c66/131 lr:0.000500 t:10.1s +tttg: c67/131 lr:0.000488 t:10.2s +tttg: c68/131 lr:0.000476 t:10.3s +tttg: c69/131 lr:0.000464 t:10.5s +tttg: c70/131 lr:0.000452 t:10.6s +tttg: c71/131 lr:0.000440 t:10.7s +tttg: c72/131 lr:0.000428 t:10.9s +tttg: c73/131 lr:0.000416 t:11.0s +tttg: c74/131 lr:0.000404 t:11.1s +tttg: c75/131 lr:0.000392 t:11.3s +tttg: c76/131 lr:0.000380 t:11.4s +tttg: c77/131 lr:0.000369 t:11.5s +tttg: c78/131 lr:0.000357 t:11.7s +tttg: c79/131 lr:0.000345 t:11.8s +tttg: c80/131 lr:0.000334 t:11.9s +tttg: c81/131 lr:0.000323 t:12.1s +tttg: c82/131 lr:0.000311 t:12.2s +tttg: c83/131 lr:0.000300 t:12.3s +tttg: c84/131 lr:0.000289 t:12.5s +tttg: c85/131 lr:0.000278 t:12.6s +tttg: c86/131 lr:0.000268 t:12.7s +tttg: c87/131 lr:0.000257 t:12.8s +tttg: c88/131 lr:0.000247 t:13.0s +tttg: c89/131 lr:0.000236 t:13.1s +tttg: c90/131 lr:0.000226 t:13.2s +tttg: c91/131 lr:0.000216 t:13.4s +tttg: c92/131 lr:0.000206 t:13.5s +tttg: c93/131 lr:0.000196 t:13.6s +tttg: c94/131 lr:0.000187 t:13.8s +tttg: c95/131 lr:0.000178 t:13.9s +tttg: c96/131 lr:0.000168 t:14.0s +tttg: c97/131 lr:0.000159 t:14.2s +tttg: c98/131 lr:0.000151 t:14.3s +tttg: c99/131 lr:0.000142 t:14.4s +tttg: c100/131 lr:0.000134 t:14.6s +tttg: c101/131 lr:0.000126 t:14.7s +tttg: c102/131 lr:0.000118 t:14.8s +tttg: c103/131 lr:0.000110 t:15.0s +tttg: c104/131 lr:0.000103 t:15.1s +tttg: c105/131 lr:0.000095 t:15.2s +tttg: c106/131 lr:0.000089 t:15.3s +tttg: c107/131 lr:0.000082 t:15.5s +tttg: c108/131 lr:0.000075 t:15.6s +tttg: c109/131 lr:0.000069 t:15.7s +tttg: c110/131 lr:0.000063 t:15.9s +tttg: c111/131 lr:0.000057 t:16.0s +tttg: c112/131 lr:0.000052 t:16.1s +tttg: c113/131 lr:0.000047 t:16.3s +tttg: c114/131 lr:0.000042 t:16.4s +tttg: c115/131 lr:0.000037 t:16.5s +tttg: c116/131 lr:0.000032 t:16.7s +tttg: c117/131 lr:0.000028 t:16.8s +tttg: c118/131 lr:0.000024 t:16.9s +tttg: c119/131 lr:0.000021 t:17.1s +tttg: c120/131 lr:0.000018 t:17.2s +tttg: c121/131 lr:0.000015 t:17.3s +tttg: c122/131 lr:0.000012 t:17.4s +tttg: c123/131 lr:0.000009 t:17.6s +tttg: c124/131 lr:0.000007 t:17.7s +tttg: c125/131 lr:0.000005 t:17.8s +tttg: c126/131 lr:0.000004 t:18.0s +tttg: c127/131 lr:0.000002 t:18.1s +tttg: c128/131 lr:0.000001 t:18.2s +tttg: c129/131 lr:0.000001 t:18.4s +tttg: c130/131 lr:0.000000 t:18.5s +ttpr: phase:1/3 t:927.1s +ttp: b754/782 bl:2.2873 bb:1.0581 rl:2.2582 rb:1.0627 dl:3345-3397 gd:0 +ttp: b753/782 bl:2.2158 bb:1.0003 rl:2.2532 rb:1.0551 dl:3284-3344 gd:0 +ttpp: phase:2/3 pd:2128 gd:1666 t:1178.6s +tttg: c1/219 lr:0.001000 t:0.1s +tttg: c2/219 lr:0.001000 t:0.3s +tttg: c3/219 lr:0.001000 t:0.4s +tttg: c4/219 lr:0.001000 t:0.5s +tttg: c5/219 lr:0.000999 t:0.7s +tttg: c6/219 lr:0.000999 t:0.8s +tttg: c7/219 lr:0.000998 t:0.9s +tttg: c8/219 lr:0.000997 t:1.1s +tttg: c9/219 lr:0.000997 t:1.2s +tttg: c10/219 lr:0.000996 t:1.3s +tttg: c11/219 lr:0.000995 t:1.4s +tttg: c12/219 lr:0.000994 t:1.6s +tttg: c13/219 lr:0.000993 t:1.7s +tttg: c14/219 lr:0.000991 t:1.8s +tttg: c15/219 lr:0.000990 t:2.0s +tttg: c16/219 lr:0.000988 t:2.1s +tttg: c17/219 lr:0.000987 t:2.2s +tttg: c18/219 lr:0.000985 t:2.4s +tttg: c19/219 lr:0.000983 t:2.5s +tttg: c20/219 lr:0.000981 t:2.6s +tttg: c21/219 lr:0.000979 t:2.8s +tttg: c22/219 lr:0.000977 t:2.9s +tttg: c23/219 lr:0.000975 t:3.0s +tttg: c24/219 lr:0.000973 t:3.2s +tttg: c25/219 lr:0.000970 t:3.3s +tttg: c26/219 lr:0.000968 t:3.4s +tttg: c27/219 lr:0.000965 t:3.6s +tttg: c28/219 lr:0.000963 t:3.7s +tttg: c29/219 lr:0.000960 t:3.8s +tttg: c30/219 lr:0.000957 t:4.0s +tttg: c31/219 lr:0.000954 t:4.1s +tttg: c32/219 lr:0.000951 t:4.2s +tttg: c33/219 lr:0.000948 t:4.4s +tttg: c34/219 lr:0.000945 t:4.5s +tttg: c35/219 lr:0.000941 t:4.6s +tttg: c36/219 lr:0.000938 t:4.8s +tttg: c37/219 lr:0.000934 t:4.9s +tttg: c38/219 lr:0.000931 t:5.0s +tttg: c39/219 lr:0.000927 t:5.2s +tttg: c40/219 lr:0.000923 t:5.3s +tttg: c41/219 lr:0.000919 t:5.4s +tttg: c42/219 lr:0.000915 t:5.6s +tttg: c43/219 lr:0.000911 t:5.7s +tttg: c44/219 lr:0.000907 t:5.8s +tttg: c45/219 lr:0.000903 t:6.0s +tttg: c46/219 lr:0.000898 t:6.1s +tttg: c47/219 lr:0.000894 t:6.2s +tttg: c48/219 lr:0.000890 t:6.4s +tttg: c49/219 lr:0.000885 t:6.5s +tttg: c50/219 lr:0.000880 t:6.6s +tttg: c51/219 lr:0.000876 t:6.8s +tttg: c52/219 lr:0.000871 t:6.9s +tttg: c53/219 lr:0.000866 t:7.0s +tttg: c54/219 lr:0.000861 t:7.1s +tttg: c55/219 lr:0.000856 t:7.3s +tttg: c56/219 lr:0.000851 t:7.4s +tttg: c57/219 lr:0.000846 t:7.5s +tttg: c58/219 lr:0.000841 t:7.7s +tttg: c59/219 lr:0.000835 t:7.8s +tttg: c60/219 lr:0.000830 t:7.9s +tttg: c61/219 lr:0.000824 t:8.1s +tttg: c62/219 lr:0.000819 t:8.2s +tttg: c63/219 lr:0.000813 t:8.3s +tttg: c64/219 lr:0.000808 t:8.5s +tttg: c65/219 lr:0.000802 t:8.6s +tttg: c66/219 lr:0.000796 t:8.7s +tttg: c67/219 lr:0.000790 t:8.9s +tttg: c68/219 lr:0.000784 t:9.0s +tttg: c69/219 lr:0.000779 t:9.1s +tttg: c70/219 lr:0.000773 t:9.3s +tttg: c71/219 lr:0.000766 t:9.4s +tttg: c72/219 lr:0.000760 t:9.5s +tttg: c73/219 lr:0.000754 t:9.7s +tttg: c74/219 lr:0.000748 t:9.8s +tttg: c75/219 lr:0.000742 t:9.9s +tttg: c76/219 lr:0.000735 t:10.1s +tttg: c77/219 lr:0.000729 t:10.2s +tttg: c78/219 lr:0.000722 t:10.3s +tttg: c79/219 lr:0.000716 t:10.5s +tttg: c80/219 lr:0.000709 t:10.6s +tttg: c81/219 lr:0.000703 t:10.7s +tttg: c82/219 lr:0.000696 t:10.9s +tttg: c83/219 lr:0.000690 t:11.0s +tttg: c84/219 lr:0.000683 t:11.1s +tttg: c85/219 lr:0.000676 t:11.3s +tttg: c86/219 lr:0.000670 t:11.4s +tttg: c87/219 lr:0.000663 t:11.5s +tttg: c88/219 lr:0.000656 t:11.7s +tttg: c89/219 lr:0.000649 t:11.8s +tttg: c90/219 lr:0.000642 t:11.9s +tttg: c91/219 lr:0.000635 t:12.1s +tttg: c92/219 lr:0.000628 t:12.2s +tttg: c93/219 lr:0.000621 t:12.3s +tttg: c94/219 lr:0.000614 t:12.5s +tttg: c95/219 lr:0.000607 t:12.6s +tttg: c96/219 lr:0.000600 t:12.7s +tttg: c97/219 lr:0.000593 t:12.8s +tttg: c98/219 lr:0.000586 t:13.0s +tttg: c99/219 lr:0.000579 t:13.1s +tttg: c100/219 lr:0.000572 t:13.2s +tttg: c101/219 lr:0.000565 t:13.4s +tttg: c102/219 lr:0.000558 t:13.5s +tttg: c103/219 lr:0.000550 t:13.6s +tttg: c104/219 lr:0.000543 t:13.8s +tttg: c105/219 lr:0.000536 t:13.9s +tttg: c106/219 lr:0.000529 t:14.0s +tttg: c107/219 lr:0.000522 t:14.2s +tttg: c108/219 lr:0.000514 t:14.3s +tttg: c109/219 lr:0.000507 t:14.4s +tttg: c110/219 lr:0.000500 t:14.6s +tttg: c111/219 lr:0.000493 t:14.7s +tttg: c112/219 lr:0.000486 t:14.8s +tttg: c113/219 lr:0.000478 t:15.0s +tttg: c114/219 lr:0.000471 t:15.1s +tttg: c115/219 lr:0.000464 t:15.2s +tttg: c116/219 lr:0.000457 t:15.4s +tttg: c117/219 lr:0.000450 t:15.5s +tttg: c118/219 lr:0.000442 t:15.6s +tttg: c119/219 lr:0.000435 t:15.8s +tttg: c120/219 lr:0.000428 t:15.9s +tttg: c121/219 lr:0.000421 t:16.0s +tttg: c122/219 lr:0.000414 t:16.2s +tttg: c123/219 lr:0.000407 t:16.3s +tttg: c124/219 lr:0.000400 t:16.4s +tttg: c125/219 lr:0.000393 t:16.6s +tttg: c126/219 lr:0.000386 t:16.7s +tttg: c127/219 lr:0.000379 t:16.8s +tttg: c128/219 lr:0.000372 t:17.0s +tttg: c129/219 lr:0.000365 t:17.1s +tttg: c130/219 lr:0.000358 t:17.2s +tttg: c131/219 lr:0.000351 t:17.4s +tttg: c132/219 lr:0.000344 t:17.5s +tttg: c133/219 lr:0.000337 t:17.6s +tttg: c134/219 lr:0.000330 t:17.8s +tttg: c135/219 lr:0.000324 t:17.9s +tttg: c136/219 lr:0.000317 t:18.0s +tttg: c137/219 lr:0.000310 t:18.2s +tttg: c138/219 lr:0.000304 t:18.3s +tttg: c139/219 lr:0.000297 t:18.4s +tttg: c140/219 lr:0.000291 t:18.6s +tttg: c141/219 lr:0.000284 t:18.7s +tttg: c142/219 lr:0.000278 t:18.8s +tttg: c143/219 lr:0.000271 t:19.0s +tttg: c144/219 lr:0.000265 t:19.1s +tttg: c145/219 lr:0.000258 t:19.2s +tttg: c146/219 lr:0.000252 t:19.3s +tttg: c147/219 lr:0.000246 t:19.5s +tttg: c148/219 lr:0.000240 t:19.6s +tttg: c149/219 lr:0.000234 t:19.7s +tttg: c150/219 lr:0.000227 t:19.9s +tttg: c151/219 lr:0.000221 t:20.0s +tttg: c152/219 lr:0.000216 t:20.1s +tttg: c153/219 lr:0.000210 t:20.3s +tttg: c154/219 lr:0.000204 t:20.4s +tttg: c155/219 lr:0.000198 t:20.5s +tttg: c156/219 lr:0.000192 t:20.7s +tttg: c157/219 lr:0.000187 t:20.8s +tttg: c158/219 lr:0.000181 t:20.9s +tttg: c159/219 lr:0.000176 t:21.1s +tttg: c160/219 lr:0.000170 t:21.2s +tttg: c161/219 lr:0.000165 t:21.3s +tttg: c162/219 lr:0.000159 t:21.5s +tttg: c163/219 lr:0.000154 t:21.6s +tttg: c164/219 lr:0.000149 t:21.7s +tttg: c165/219 lr:0.000144 t:21.9s +tttg: c166/219 lr:0.000139 t:22.0s +tttg: c167/219 lr:0.000134 t:22.1s +tttg: c168/219 lr:0.000129 t:22.3s +tttg: c169/219 lr:0.000124 t:22.4s +tttg: c170/219 lr:0.000120 t:22.5s +tttg: c171/219 lr:0.000115 t:22.7s +tttg: c172/219 lr:0.000110 t:22.8s +tttg: c173/219 lr:0.000106 t:22.9s +tttg: c174/219 lr:0.000102 t:23.1s +tttg: c175/219 lr:0.000097 t:23.2s +tttg: c176/219 lr:0.000093 t:23.3s +tttg: c177/219 lr:0.000089 t:23.5s +tttg: c178/219 lr:0.000085 t:23.6s +tttg: c179/219 lr:0.000081 t:23.7s +tttg: c180/219 lr:0.000077 t:23.9s +tttg: c181/219 lr:0.000073 t:24.0s +tttg: c182/219 lr:0.000069 t:24.1s +tttg: c183/219 lr:0.000066 t:24.3s +tttg: c184/219 lr:0.000062 t:24.4s +tttg: c185/219 lr:0.000059 t:24.5s +tttg: c186/219 lr:0.000055 t:24.7s +tttg: c187/219 lr:0.000052 t:24.8s +tttg: c188/219 lr:0.000049 t:24.9s +tttg: c189/219 lr:0.000046 t:25.1s +tttg: c190/219 lr:0.000043 t:25.2s +tttg: c191/219 lr:0.000040 t:25.3s +tttg: c192/219 lr:0.000037 t:25.4s +tttg: c193/219 lr:0.000035 t:25.6s +tttg: c194/219 lr:0.000032 t:25.7s +tttg: c195/219 lr:0.000030 t:25.8s +tttg: c196/219 lr:0.000027 t:26.0s +tttg: c197/219 lr:0.000025 t:26.1s +tttg: c198/219 lr:0.000023 t:26.2s +tttg: c199/219 lr:0.000021 t:26.4s +tttg: c200/219 lr:0.000019 t:26.5s +tttg: c201/219 lr:0.000017 t:26.6s +tttg: c202/219 lr:0.000015 t:26.8s +tttg: c203/219 lr:0.000013 t:26.9s +tttg: c204/219 lr:0.000012 t:27.0s +tttg: c205/219 lr:0.000010 t:27.2s +tttg: c206/219 lr:0.000009 t:27.3s +tttg: c207/219 lr:0.000007 t:27.4s +tttg: c208/219 lr:0.000006 t:27.6s +tttg: c209/219 lr:0.000005 t:27.7s +tttg: c210/219 lr:0.000004 t:27.8s +tttg: c211/219 lr:0.000003 t:28.0s +tttg: c212/219 lr:0.000003 t:28.1s +tttg: c213/219 lr:0.000002 t:28.2s +tttg: c214/219 lr:0.000001 t:28.4s +tttg: c215/219 lr:0.000001 t:28.5s +tttg: c216/219 lr:0.000000 t:28.6s +tttg: c217/219 lr:0.000000 t:28.8s +tttg: c218/219 lr:0.000000 t:28.9s +ttpr: phase:2/3 t:1208.8s +ttp: b741/782 bl:2.3189 bb:1.0399 rl:2.2589 rb:1.0537 dl:2686-2730 gd:0 +ttp: b739/782 bl:2.2886 bb:1.0212 rl:2.2613 rb:1.0510 dl:2619-2652 gd:0 +ttpp: phase:3/3 pd:2960 gd:2500 t:1366.8s +tttg: c1/289 lr:0.001000 t:0.2s +tttg: c2/289 lr:0.001000 t:0.3s +tttg: c3/289 lr:0.001000 t:0.5s +tttg: c4/289 lr:0.001000 t:0.6s +tttg: c5/289 lr:0.001000 t:0.8s +tttg: c6/289 lr:0.000999 t:0.9s +tttg: c7/289 lr:0.000999 t:1.1s +tttg: c8/289 lr:0.000999 t:1.2s +tttg: c9/289 lr:0.000998 t:1.3s +tttg: c10/289 lr:0.000998 t:1.5s +tttg: c11/289 lr:0.000997 t:1.6s +tttg: c12/289 lr:0.000996 t:1.7s +tttg: c13/289 lr:0.000996 t:1.9s +tttg: c14/289 lr:0.000995 t:2.0s +tttg: c15/289 lr:0.000994 t:2.1s +tttg: c16/289 lr:0.000993 t:2.3s +tttg: c17/289 lr:0.000992 t:2.4s +tttg: c18/289 lr:0.000991 t:2.5s +tttg: c19/289 lr:0.000990 t:2.7s +tttg: c20/289 lr:0.000989 t:2.8s +tttg: c21/289 lr:0.000988 t:2.9s +tttg: c22/289 lr:0.000987 t:3.1s +tttg: c23/289 lr:0.000986 t:3.2s +tttg: c24/289 lr:0.000984 t:3.3s +tttg: c25/289 lr:0.000983 t:3.5s +tttg: c26/289 lr:0.000982 t:3.6s +tttg: c27/289 lr:0.000980 t:3.7s +tttg: c28/289 lr:0.000978 t:3.9s +tttg: c29/289 lr:0.000977 t:4.0s +tttg: c30/289 lr:0.000975 t:4.1s +tttg: c31/289 lr:0.000973 t:4.3s +tttg: c32/289 lr:0.000972 t:4.4s +tttg: c33/289 lr:0.000970 t:4.5s +tttg: c34/289 lr:0.000968 t:4.7s +tttg: c35/289 lr:0.000966 t:4.8s +tttg: c36/289 lr:0.000964 t:4.9s +tttg: c37/289 lr:0.000962 t:5.1s +tttg: c38/289 lr:0.000960 t:5.2s +tttg: c39/289 lr:0.000958 t:5.3s +tttg: c40/289 lr:0.000955 t:5.4s +tttg: c41/289 lr:0.000953 t:5.6s +tttg: c42/289 lr:0.000951 t:5.7s +tttg: c43/289 lr:0.000948 t:5.8s +tttg: c44/289 lr:0.000946 t:6.0s +tttg: c45/289 lr:0.000944 t:6.1s +tttg: c46/289 lr:0.000941 t:6.2s +tttg: c47/289 lr:0.000938 t:6.4s +tttg: c48/289 lr:0.000936 t:6.5s +tttg: c49/289 lr:0.000933 t:6.6s +tttg: c50/289 lr:0.000930 t:6.8s +tttg: c51/289 lr:0.000927 t:6.9s +tttg: c52/289 lr:0.000925 t:7.0s +tttg: c53/289 lr:0.000922 t:7.2s +tttg: c54/289 lr:0.000919 t:7.3s +tttg: c55/289 lr:0.000916 t:7.4s +tttg: c56/289 lr:0.000913 t:7.6s +tttg: c57/289 lr:0.000910 t:7.7s +tttg: c58/289 lr:0.000906 t:7.8s +tttg: c59/289 lr:0.000903 t:8.0s +tttg: c60/289 lr:0.000900 t:8.1s +tttg: c61/289 lr:0.000897 t:8.2s +tttg: c62/289 lr:0.000893 t:8.4s +tttg: c63/289 lr:0.000890 t:8.5s +tttg: c64/289 lr:0.000887 t:8.6s +tttg: c65/289 lr:0.000883 t:8.8s +tttg: c66/289 lr:0.000879 t:8.9s +tttg: c67/289 lr:0.000876 t:9.0s +tttg: c68/289 lr:0.000872 t:9.2s +tttg: c69/289 lr:0.000869 t:9.3s +tttg: c70/289 lr:0.000865 t:9.4s +tttg: c71/289 lr:0.000861 t:9.5s +tttg: c72/289 lr:0.000857 t:9.7s +tttg: c73/289 lr:0.000854 t:9.8s +tttg: c74/289 lr:0.000850 t:9.9s +tttg: c75/289 lr:0.000846 t:10.1s +tttg: c76/289 lr:0.000842 t:10.2s +tttg: c77/289 lr:0.000838 t:10.3s +tttg: c78/289 lr:0.000834 t:10.5s +tttg: c79/289 lr:0.000830 t:10.6s +tttg: c80/289 lr:0.000826 t:10.7s +tttg: c81/289 lr:0.000821 t:10.9s +tttg: c82/289 lr:0.000817 t:11.0s +tttg: c83/289 lr:0.000813 t:11.1s +tttg: c84/289 lr:0.000809 t:11.3s +tttg: c85/289 lr:0.000804 t:11.4s +tttg: c86/289 lr:0.000800 t:11.5s +tttg: c87/289 lr:0.000796 t:11.7s +tttg: c88/289 lr:0.000791 t:11.8s +tttg: c89/289 lr:0.000787 t:11.9s +tttg: c90/289 lr:0.000782 t:12.1s +tttg: c91/289 lr:0.000778 t:12.2s +tttg: c92/289 lr:0.000773 t:12.3s +tttg: c93/289 lr:0.000769 t:12.5s +tttg: c94/289 lr:0.000764 t:12.6s +tttg: c95/289 lr:0.000759 t:12.7s +tttg: c96/289 lr:0.000755 t:12.9s +tttg: c97/289 lr:0.000750 t:13.0s +tttg: c98/289 lr:0.000745 t:13.1s +tttg: c99/289 lr:0.000740 t:13.3s +tttg: c100/289 lr:0.000736 t:13.4s +tttg: c101/289 lr:0.000731 t:13.5s +tttg: c102/289 lr:0.000726 t:13.6s +tttg: c103/289 lr:0.000721 t:13.8s +tttg: c104/289 lr:0.000716 t:13.9s +tttg: c105/289 lr:0.000711 t:14.0s +tttg: c106/289 lr:0.000706 t:14.2s +tttg: c107/289 lr:0.000701 t:14.3s +tttg: c108/289 lr:0.000696 t:14.4s +tttg: c109/289 lr:0.000691 t:14.6s +tttg: c110/289 lr:0.000686 t:14.7s +tttg: c111/289 lr:0.000681 t:14.8s +tttg: c112/289 lr:0.000676 t:15.0s +tttg: c113/289 lr:0.000671 t:15.1s +tttg: c114/289 lr:0.000666 t:15.2s +tttg: c115/289 lr:0.000661 t:15.4s +tttg: c116/289 lr:0.000656 t:15.5s +tttg: c117/289 lr:0.000650 t:15.6s +tttg: c118/289 lr:0.000645 t:15.8s +tttg: c119/289 lr:0.000640 t:15.9s +tttg: c120/289 lr:0.000635 t:16.0s +tttg: c121/289 lr:0.000629 t:16.2s +tttg: c122/289 lr:0.000624 t:16.3s +tttg: c123/289 lr:0.000619 t:16.4s +tttg: c124/289 lr:0.000614 t:16.6s +tttg: c125/289 lr:0.000608 t:16.7s +tttg: c126/289 lr:0.000603 t:16.8s +tttg: c127/289 lr:0.000598 t:17.0s +tttg: c128/289 lr:0.000592 t:17.1s +tttg: c129/289 lr:0.000587 t:17.2s +tttg: c130/289 lr:0.000581 t:17.4s +tttg: c131/289 lr:0.000576 t:17.5s +tttg: c132/289 lr:0.000571 t:17.6s +tttg: c133/289 lr:0.000565 t:17.7s +tttg: c134/289 lr:0.000560 t:17.9s +tttg: c135/289 lr:0.000554 t:18.0s +tttg: c136/289 lr:0.000549 t:18.1s +tttg: c137/289 lr:0.000544 t:18.3s +tttg: c138/289 lr:0.000538 t:18.4s +tttg: c139/289 lr:0.000533 t:18.5s +tttg: c140/289 lr:0.000527 t:18.7s +tttg: c141/289 lr:0.000522 t:18.8s +tttg: c142/289 lr:0.000516 t:18.9s +tttg: c143/289 lr:0.000511 t:19.1s +tttg: c144/289 lr:0.000505 t:19.2s +tttg: c145/289 lr:0.000500 t:19.3s +tttg: c146/289 lr:0.000495 t:19.5s +tttg: c147/289 lr:0.000489 t:19.6s +tttg: c148/289 lr:0.000484 t:19.7s +tttg: c149/289 lr:0.000478 t:19.9s +tttg: c150/289 lr:0.000473 t:20.0s +tttg: c151/289 lr:0.000467 t:20.1s +tttg: c152/289 lr:0.000462 t:20.3s +tttg: c153/289 lr:0.000456 t:20.4s +tttg: c154/289 lr:0.000451 t:20.5s +tttg: c155/289 lr:0.000446 t:20.7s +tttg: c156/289 lr:0.000440 t:20.8s +tttg: c157/289 lr:0.000435 t:20.9s +tttg: c158/289 lr:0.000429 t:21.1s +tttg: c159/289 lr:0.000424 t:21.2s +tttg: c160/289 lr:0.000419 t:21.3s +tttg: c161/289 lr:0.000413 t:21.5s +tttg: c162/289 lr:0.000408 t:21.6s +tttg: c163/289 lr:0.000402 t:21.7s +tttg: c164/289 lr:0.000397 t:21.8s +tttg: c165/289 lr:0.000392 t:22.0s +tttg: c166/289 lr:0.000386 t:22.1s +tttg: c167/289 lr:0.000381 t:22.2s +tttg: c168/289 lr:0.000376 t:22.4s +tttg: c169/289 lr:0.000371 t:22.5s +tttg: c170/289 lr:0.000365 t:22.6s +tttg: c171/289 lr:0.000360 t:22.8s +tttg: c172/289 lr:0.000355 t:22.9s +tttg: c173/289 lr:0.000350 t:23.0s +tttg: c174/289 lr:0.000344 t:23.2s +tttg: c175/289 lr:0.000339 t:23.3s +tttg: c176/289 lr:0.000334 t:23.4s +tttg: c177/289 lr:0.000329 t:23.6s +tttg: c178/289 lr:0.000324 t:23.7s +tttg: c179/289 lr:0.000319 t:23.8s +tttg: c180/289 lr:0.000314 t:24.0s +tttg: c181/289 lr:0.000309 t:24.1s +tttg: c182/289 lr:0.000304 t:24.2s +tttg: c183/289 lr:0.000299 t:24.4s +tttg: c184/289 lr:0.000294 t:24.5s +tttg: c185/289 lr:0.000289 t:24.6s +tttg: c186/289 lr:0.000284 t:24.8s +tttg: c187/289 lr:0.000279 t:24.9s +tttg: c188/289 lr:0.000274 t:25.0s +tttg: c189/289 lr:0.000269 t:25.2s +tttg: c190/289 lr:0.000264 t:25.3s +tttg: c191/289 lr:0.000260 t:25.4s +tttg: c192/289 lr:0.000255 t:25.6s +tttg: c193/289 lr:0.000250 t:25.7s +tttg: c194/289 lr:0.000245 t:25.8s +tttg: c195/289 lr:0.000241 t:25.9s +tttg: c196/289 lr:0.000236 t:26.1s +tttg: c197/289 lr:0.000231 t:26.2s +tttg: c198/289 lr:0.000227 t:26.3s +tttg: c199/289 lr:0.000222 t:26.5s +tttg: c200/289 lr:0.000218 t:26.6s +tttg: c201/289 lr:0.000213 t:26.7s +tttg: c202/289 lr:0.000209 t:26.9s +tttg: c203/289 lr:0.000204 t:27.0s +tttg: c204/289 lr:0.000200 t:27.1s +tttg: c205/289 lr:0.000196 t:27.3s +tttg: c206/289 lr:0.000191 t:27.4s +tttg: c207/289 lr:0.000187 t:27.5s +tttg: c208/289 lr:0.000183 t:27.7s +tttg: c209/289 lr:0.000179 t:27.8s +tttg: c210/289 lr:0.000174 t:27.9s +tttg: c211/289 lr:0.000170 t:28.1s +tttg: c212/289 lr:0.000166 t:28.2s +tttg: c213/289 lr:0.000162 t:28.3s +tttg: c214/289 lr:0.000158 t:28.5s +tttg: c215/289 lr:0.000154 t:28.6s +tttg: c216/289 lr:0.000150 t:28.7s +tttg: c217/289 lr:0.000146 t:28.9s +tttg: c218/289 lr:0.000143 t:29.0s +tttg: c219/289 lr:0.000139 t:29.1s +tttg: c220/289 lr:0.000135 t:29.3s +tttg: c221/289 lr:0.000131 t:29.4s +tttg: c222/289 lr:0.000128 t:29.5s +tttg: c223/289 lr:0.000124 t:29.7s +tttg: c224/289 lr:0.000121 t:29.8s +tttg: c225/289 lr:0.000117 t:29.9s +tttg: c226/289 lr:0.000113 t:30.0s +tttg: c227/289 lr:0.000110 t:30.2s +tttg: c228/289 lr:0.000107 t:30.3s +tttg: c229/289 lr:0.000103 t:30.4s +tttg: c230/289 lr:0.000100 t:30.6s +tttg: c231/289 lr:0.000097 t:30.7s +tttg: c232/289 lr:0.000094 t:30.8s +tttg: c233/289 lr:0.000090 t:31.0s +tttg: c234/289 lr:0.000087 t:31.1s +tttg: c235/289 lr:0.000084 t:31.2s +tttg: c236/289 lr:0.000081 t:31.4s +tttg: c237/289 lr:0.000078 t:31.5s +tttg: c238/289 lr:0.000075 t:31.6s +tttg: c239/289 lr:0.000073 t:31.8s +tttg: c240/289 lr:0.000070 t:31.9s +tttg: c241/289 lr:0.000067 t:32.0s +tttg: c242/289 lr:0.000064 t:32.2s +tttg: c243/289 lr:0.000062 t:32.3s +tttg: c244/289 lr:0.000059 t:32.4s +tttg: c245/289 lr:0.000056 t:32.6s +tttg: c246/289 lr:0.000054 t:32.7s +tttg: c247/289 lr:0.000052 t:32.8s +tttg: c248/289 lr:0.000049 t:33.0s +tttg: c249/289 lr:0.000047 t:33.1s +tttg: c250/289 lr:0.000045 t:33.2s +tttg: c251/289 lr:0.000042 t:33.4s +tttg: c252/289 lr:0.000040 t:33.5s +tttg: c253/289 lr:0.000038 t:33.6s +tttg: c254/289 lr:0.000036 t:33.8s +tttg: c255/289 lr:0.000034 t:33.9s +tttg: c256/289 lr:0.000032 t:34.0s +tttg: c257/289 lr:0.000030 t:34.2s +tttg: c258/289 lr:0.000028 t:34.3s +tttg: c259/289 lr:0.000027 t:34.4s +tttg: c260/289 lr:0.000025 t:34.5s +tttg: c261/289 lr:0.000023 t:34.7s +tttg: c262/289 lr:0.000022 t:34.8s +tttg: c263/289 lr:0.000020 t:34.9s +tttg: c264/289 lr:0.000018 t:35.1s +tttg: c265/289 lr:0.000017 t:35.2s +tttg: c266/289 lr:0.000016 t:35.3s +tttg: c267/289 lr:0.000014 t:35.5s +tttg: c268/289 lr:0.000013 t:35.6s +tttg: c269/289 lr:0.000012 t:35.7s +tttg: c270/289 lr:0.000011 t:35.9s +tttg: c271/289 lr:0.000010 t:36.0s +tttg: c272/289 lr:0.000009 t:36.1s +tttg: c273/289 lr:0.000008 t:36.3s +tttg: c274/289 lr:0.000007 t:36.4s +tttg: c275/289 lr:0.000006 t:36.5s +tttg: c276/289 lr:0.000005 t:36.7s +tttg: c277/289 lr:0.000004 t:36.8s +tttg: c278/289 lr:0.000004 t:36.9s +tttg: c279/289 lr:0.000003 t:37.1s +tttg: c280/289 lr:0.000002 t:37.2s +tttg: c281/289 lr:0.000002 t:37.3s +tttg: c282/289 lr:0.000001 t:37.5s +tttg: c283/289 lr:0.000001 t:37.6s +tttg: c284/289 lr:0.000001 t:37.7s +tttg: c285/289 lr:0.000000 t:37.9s +tttg: c286/289 lr:0.000000 t:38.0s +tttg: c287/289 lr:0.000000 t:38.1s +tttg: c288/289 lr:0.000000 t:38.3s +ttpr: phase:3/3 t:1406.3s +ttp: b728/782 bl:2.3594 bb:1.0802 rl:2.2676 rb:1.0529 dl:2306-2324 gd:1 +ttp: b727/782 bl:2.2628 bb:1.0428 rl:2.2673 rb:1.0523 dl:2277-2305 gd:1 +ttp: b715/782 bl:2.3559 bb:1.0271 rl:2.2719 rb:1.0510 dl:2036-2053 gd:1 +ttp: b706/782 bl:2.4010 bb:1.0738 rl:2.2777 rb:1.0520 dl:1898-1910 gd:1 +ttp: b699/782 bl:2.4158 bb:1.0546 rl:2.2834 rb:1.0521 dl:1814-1824 gd:1 +ttp: b692/782 bl:2.2889 bb:1.0273 rl:2.2836 rb:1.0512 dl:1737-1746 gd:1 +ttp: b680/782 bl:2.2806 bb:1.0270 rl:2.2835 rb:1.0503 dl:1618-1628 gd:1 +ttp: b679/782 bl:2.3035 bb:1.0577 rl:2.2842 rb:1.0506 dl:1610-1618 gd:1 +ttp: b664/782 bl:2.3398 bb:1.0269 rl:2.2858 rb:1.0498 dl:1493-1499 gd:1 +ttp: b661/782 bl:2.4001 bb:1.0850 rl:2.2891 rb:1.0508 dl:1474-1480 gd:1 +ttp: b648/782 bl:2.2811 bb:1.0066 rl:2.2889 rb:1.0496 dl:1387-1392 gd:1 +ttp: b640/782 bl:2.3056 bb:1.0503 rl:2.2893 rb:1.0497 dl:1337-1343 gd:1 +ttp: b634/782 bl:2.3825 bb:1.0488 rl:2.2915 rb:1.0496 dl:1302-1308 gd:1 +ttp: b627/782 bl:2.3776 bb:1.0705 rl:2.2934 rb:1.0501 dl:1266-1271 gd:1 +ttp: b619/782 bl:2.3197 bb:1.0578 rl:2.2940 rb:1.0503 dl:1221-1226 gd:1 +ttp: b609/782 bl:2.2707 bb:1.0173 rl:2.2935 rb:1.0496 dl:1172-1177 gd:1 +ttp: b604/782 bl:2.3814 bb:1.0453 rl:2.2952 rb:1.0495 dl:1150-1154 gd:1 +ttp: b593/782 bl:2.2906 bb:1.0111 rl:2.2951 rb:1.0488 dl:1103-1107 gd:1 +ttp: b585/782 bl:2.2829 bb:1.0350 rl:2.2949 rb:1.0486 dl:1069-1073 gd:1 +ttp: b577/782 bl:2.2897 bb:1.0306 rl:2.2948 rb:1.0483 dl:1037-1041 gd:1 +ttp: b569/782 bl:2.3049 bb:1.0423 rl:2.2949 rb:1.0482 dl:1007-1010 gd:1 +ttp: b561/782 bl:2.2472 bb:1.0136 rl:2.2942 rb:1.0477 dl:979-983 gd:1 +ttp: b553/782 bl:2.2832 bb:1.0294 rl:2.2941 rb:1.0474 dl:952-955 gd:1 +ttp: b545/782 bl:2.3306 bb:1.0306 rl:2.2946 rb:1.0472 dl:927-930 gd:1 +ttp: b536/782 bl:2.3148 bb:1.0424 rl:2.2948 rb:1.0471 dl:899-902 gd:1 +ttp: b528/782 bl:2.3383 bb:1.0451 rl:2.2954 rb:1.0471 dl:875-878 gd:1 +ttp: b520/782 bl:2.3266 bb:1.0030 rl:2.2958 rb:1.0465 dl:852-854 gd:1 +ttp: b512/782 bl:2.3034 bb:1.0637 rl:2.2959 rb:1.0467 dl:829-832 gd:1 +ttp: b509/782 bl:2.3569 bb:1.0348 rl:2.2966 rb:1.0466 dl:820-823 gd:1 +ttp: b500/782 bl:2.3246 bb:1.0642 rl:2.2969 rb:1.0468 dl:796-799 gd:1 +ttp: b493/782 bl:2.3650 bb:1.0439 rl:2.2976 rb:1.0467 dl:778-780 gd:1 +ttp: b484/782 bl:2.3654 bb:1.0481 rl:2.2983 rb:1.0467 dl:756-759 gd:1 +ttp: b476/782 bl:2.2767 bb:1.0317 rl:2.2981 rb:1.0466 dl:738-740 gd:1 +ttp: b464/782 bl:2.2690 bb:1.0168 rl:2.2978 rb:1.0463 dl:710-712 gd:1 +ttp: b456/782 bl:2.3487 bb:1.0404 rl:2.2982 rb:1.0463 dl:693-695 gd:1 +ttp: b448/782 bl:2.3086 bb:1.0064 rl:2.2983 rb:1.0459 dl:677-678 gd:1 +ttp: b440/782 bl:2.2368 bb:0.9846 rl:2.2978 rb:1.0454 dl:659-662 gd:1 +ttp: b432/782 bl:2.3337 bb:1.0372 rl:2.2981 rb:1.0453 dl:643-645 gd:1 +ttp: b424/782 bl:2.3404 bb:1.0611 rl:2.2984 rb:1.0454 dl:629-630 gd:1 +ttp: b416/782 bl:2.3714 bb:1.0426 rl:2.2990 rb:1.0454 dl:613-615 gd:1 +ttp: b408/782 bl:2.3009 bb:1.0699 rl:2.2990 rb:1.0456 dl:597-598 gd:1 +ttp: b400/782 bl:2.3057 bb:1.0375 rl:2.2991 rb:1.0455 dl:582-584 gd:1 +ttp: b392/782 bl:2.2436 bb:1.0320 rl:2.2987 rb:1.0454 dl:568-570 gd:1 +ttp: b384/782 bl:2.3383 bb:1.0518 rl:2.2990 rb:1.0455 dl:554-555 gd:1 +ttp: b376/782 bl:2.3236 bb:1.0421 rl:2.2991 rb:1.0454 dl:540-542 gd:1 +ttp: b369/782 bl:2.3502 bb:1.0615 rl:2.2994 rb:1.0455 dl:528-530 gd:1 +ttp: b361/782 bl:2.3429 bb:1.0938 rl:2.2997 rb:1.0458 dl:515-517 gd:1 +ttp: b353/782 bl:2.1979 bb:1.0049 rl:2.2991 rb:1.0456 dl:501-503 gd:1 +ttp: b345/782 bl:2.3547 bb:1.0719 rl:2.2994 rb:1.0457 dl:489-491 gd:1 +ttp: b337/782 bl:2.3136 bb:1.0529 rl:2.2995 rb:1.0458 dl:477-478 gd:1 +ttp: b332/782 bl:2.3082 bb:1.0449 rl:2.2995 rb:1.0458 dl:469-471 gd:1 +ttp: b325/782 bl:2.3513 bb:1.0815 rl:2.2998 rb:1.0460 dl:459-461 gd:1 +ttp: b317/782 bl:2.2992 bb:1.0444 rl:2.2998 rb:1.0460 dl:446-448 gd:1 +ttp: b309/782 bl:2.4045 bb:1.1034 rl:2.3003 rb:1.0462 dl:435-437 gd:1 +ttp: b301/782 bl:2.3512 bb:1.0915 rl:2.3006 rb:1.0465 dl:422-424 gd:1 +ttp: b292/782 bl:2.3324 bb:1.1045 rl:2.3007 rb:1.0467 dl:409-410 gd:1 +ttp: b283/782 bl:2.3740 bb:1.1286 rl:2.3011 rb:1.0471 dl:396-398 gd:1 +ttp: b275/782 bl:2.3477 bb:1.0575 rl:2.3013 rb:1.0471 dl:385-386 gd:1 +ttp: b264/782 bl:2.4255 bb:1.1053 rl:2.3018 rb:1.0474 dl:371-372 gd:1 +ttp: b256/782 bl:2.5360 bb:1.1194 rl:2.3027 rb:1.0477 dl:361-362 gd:1 +ttp: b248/782 bl:2.4610 bb:1.1878 rl:2.3034 rb:1.0482 dl:351-352 gd:1 +ttp: b243/782 bl:2.3528 bb:1.0796 rl:2.3035 rb:1.0483 dl:345-346 gd:1 +ttp: b235/782 bl:2.2815 bb:1.0984 rl:2.3035 rb:1.0485 dl:335-336 gd:1 +ttp: b227/782 bl:2.4841 bb:1.1534 rl:2.3041 rb:1.0488 dl:325-327 gd:1 +ttp: b219/782 bl:2.3406 bb:1.1201 rl:2.3042 rb:1.0491 dl:316-317 gd:1 +ttp: b211/782 bl:2.3899 bb:1.0887 rl:2.3045 rb:1.0492 dl:307-308 gd:1 +ttp: b203/782 bl:2.4238 bb:1.1057 rl:2.3049 rb:1.0494 dl:299-300 gd:1 +ttp: b195/782 bl:2.4209 bb:1.1291 rl:2.3053 rb:1.0496 dl:290-291 gd:1 +ttp: b184/782 bl:2.3848 bb:1.1243 rl:2.3055 rb:1.0499 dl:278-279 gd:1 +ttp: b176/782 bl:2.3205 bb:1.1268 rl:2.3056 rb:1.0501 dl:270-271 gd:1 +ttp: b167/782 bl:2.5207 bb:1.1243 rl:2.3062 rb:1.0503 dl:262-263 gd:1 +ttp: b159/782 bl:2.4752 bb:1.1484 rl:2.3066 rb:1.0505 dl:254-255 gd:1 +ttp: b155/782 bl:2.4035 bb:1.1112 rl:2.3069 rb:1.0507 dl:250-251 gd:1 +ttp: b147/782 bl:2.4585 bb:1.1186 rl:2.3073 rb:1.0509 dl:242-243 gd:1 +ttp: b139/782 bl:2.4305 bb:1.1319 rl:2.3076 rb:1.0511 dl:234-235 gd:1 +ttp: b131/782 bl:2.3909 bb:1.1544 rl:2.3078 rb:1.0513 dl:227-228 gd:1 +ttp: b121/782 bl:2.4356 bb:1.1117 rl:2.3081 rb:1.0514 dl:218-219 gd:1 +ttp: b114/782 bl:2.4663 bb:1.1437 rl:2.3084 rb:1.0517 dl:211-212 gd:1 +ttp: b106/782 bl:2.4249 bb:1.1673 rl:2.3087 rb:1.0519 dl:204-205 gd:1 +ttp: b98/782 bl:2.5923 bb:1.2163 rl:2.3093 rb:1.0522 dl:197-198 gd:1 +ttp: b89/782 bl:2.4888 bb:1.1500 rl:2.3096 rb:1.0524 dl:189-190 gd:1 +ttp: b83/782 bl:2.4263 bb:1.1449 rl:2.3099 rb:1.0526 dl:183-184 gd:1 +ttp: b75/782 bl:2.5696 bb:1.1914 rl:2.3103 rb:1.0528 dl:176-177 gd:1 +ttp: b65/782 bl:2.4719 bb:1.1725 rl:2.3106 rb:1.0530 dl:167-169 gd:1 +ttp: b58/782 bl:2.5095 bb:1.2180 rl:2.3110 rb:1.0533 dl:161-162 gd:1 +ttp: b48/782 bl:2.5193 bb:1.2147 rl:2.3113 rb:1.0535 dl:151-152 gd:1 +ttp: b40/782 bl:2.4825 bb:1.1500 rl:2.3115 rb:1.0537 dl:143-144 gd:1 +ttp: b32/782 bl:2.6027 bb:1.2137 rl:2.3120 rb:1.0539 dl:135-136 gd:1 +ttp: b24/782 bl:2.4605 bb:1.1604 rl:2.3122 rb:1.0540 dl:127-128 gd:1 +ttp: b20/782 bl:2.5872 bb:1.2389 rl:2.3125 rb:1.0543 dl:122-123 gd:1 +ttp: b12/782 bl:2.5610 bb:1.1845 rl:2.3128 rb:1.0544 dl:110-112 gd:1 +ttp: b4/782 bl:2.7475 bb:1.2311 rl:2.3132 rb:1.0546 dl:93-96 gd:1 +quantized_ttt_phased val_loss:2.31960602 val_bpb:1.05996771 eval_time:3348399ms +total_eval_time:3348.4s +Training complete diff --git a/records/submission_2026_04_29_b180_tlr56_SUB106/prepare_caseops_data.py b/records/submission_2026_04_29_b180_tlr56_SUB106/prepare_caseops_data.py new file mode 100644 index 0000000000..5c3f13e69c --- /dev/null +++ b/records/submission_2026_04_29_b180_tlr56_SUB106/prepare_caseops_data.py @@ -0,0 +1,177 @@ +"""Prepare CaseOps-tokenized FineWeb shards + per-token byte sidecar. + +CaseOps (``lossless_caps_caseops_v1``) is a bijective, character-level text +transform that introduces four operator tokens in place of explicit +capitalization: TITLE, ALLCAPS, CAPNEXT, ESC. The transform is fully +reversible — no information is lost relative to the untransformed UTF-8 +text, so BPB stays computable on TRUE byte counts. + +Forward pipeline: + 1. Read the canonical FineWeb-10B doc stream (``docs_selected.jsonl`` + produced by ``data/download_hf_docs_and_tokenize.py`` in the root repo). + 2. Apply ``encode_lossless_caps_v2`` (the caseops_v1 alias) to each doc. + 3. Tokenize with the shipped SP model + ``tokenizers/fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model`` + (reserves TITLE/ALLCAPS/CAPNEXT/ESC + sentinel as user_defined_symbols). + 4. Write uint16 train/val shards (``fineweb_{train,val}_XXXXXX.bin``). + 5. For the VAL stream only, emit per-token byte sidecar shards + (``fineweb_val_bytes_XXXXXX.bin``, uint16 parallel arrays) that record + each token's ORIGINAL pre-transform UTF-8 byte count. BPB is computed + from these canonical bytes so the score is on the untransformed text + (not the transformed representation). + +Output layout — matches what ``train_gpt.py`` expects under +``DATA_DIR=./data`` with ``CASEOPS_ENABLED=1``: + + data/datasets/fineweb10B_sp8192_caseops/datasets/ + tokenizers/fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model + datasets/fineweb10B_sp8192_lossless_caps_caseops_v1_reserved/ + fineweb_train_000000.bin + fineweb_train_000001.bin + ... + fineweb_val_000000.bin + fineweb_val_bytes_000000.bin + +Usage: + + python3 prepare_caseops_data.py \\ + --docs ./fineweb10B_raw/docs_selected.jsonl \\ + --out ./data/datasets/fineweb10B_sp8192_caseops/datasets \\ + --sp ./tokenizers/fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model + +Requirements: sentencepiece, numpy. CPU-only. Runs once; reused across seeds. +""" +from __future__ import annotations + +import argparse +import json +import pathlib +import struct +import sys + +import numpy as np +import sentencepiece as spm + +# Local import — lossless_caps.py ships next to this script. +sys.path.insert(0, str(pathlib.Path(__file__).resolve().parent)) +from lossless_caps import ( # noqa: E402 + LOSSLESS_CAPS_CASEOPS_V1, + encode_lossless_caps_v2, + surface_piece_original_byte_counts, +) + + +SHARD_MAGIC = 20240520 +SHARD_VERSION = 1 +SHARD_TOKENS = 10_000_000 # tokens per shard — matches the main pipeline +BOS_ID = 1 # SP model's control token; train_gpt.py:_find_docs requires BOS per doc + + +def _write_shard(out_path: pathlib.Path, arr: np.ndarray) -> None: + """Write a uint16 shard in the standard header-prefixed format.""" + assert arr.dtype == np.uint16 + header = np.zeros(256, dtype=np.int32) + header[0] = SHARD_MAGIC + header[1] = SHARD_VERSION + header[2] = int(arr.size) + with out_path.open("wb") as fh: + fh.write(header.tobytes()) + fh.write(arr.tobytes()) + + +def _iter_docs(docs_path: pathlib.Path): + """Yield doc strings from a jsonl file (one json object per line).""" + with docs_path.open("r", encoding="utf-8") as fh: + for line in fh: + line = line.strip() + if not line: + continue + obj = json.loads(line) + # Support both {"text": ...} and raw strings. + yield obj["text"] if isinstance(obj, dict) else obj + + +def _token_original_byte_counts( + sp: spm.SentencePieceProcessor, + original_text: str, + transformed_text: str, +) -> np.ndarray: + """Per-token canonical (pre-transform) UTF-8 byte counts. + + Delegates to ``surface_piece_original_byte_counts`` in ``lossless_caps.py`` + — the canonical exporter used by the PR #1729 / HF-hosted CaseOps dataset. + Operator pieces (U+E001..U+E004) contribute 0 original bytes; letter pieces + contribute their pre-transform UTF-8 byte count. + """ + proto = sp.encode_as_immutable_proto(transformed_text) + byte_counts = surface_piece_original_byte_counts( + (piece.surface for piece in proto.pieces), + text_transform_name=LOSSLESS_CAPS_CASEOPS_V1, + ) + return np.asarray(list(byte_counts), dtype=np.uint16) + + +def main() -> None: + ap = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter) + ap.add_argument("--docs", required=True, type=pathlib.Path, help="Path to docs_selected.jsonl") + ap.add_argument("--out", required=True, type=pathlib.Path, help="Output datasets dir") + ap.add_argument("--sp", required=True, type=pathlib.Path, help="Path to CaseOps SP model") + ap.add_argument("--val-docs", type=int, default=10_000, help="Validation docs count") + args = ap.parse_args() + + sp = spm.SentencePieceProcessor(model_file=str(args.sp)) + print(f"loaded sp: vocab={sp.vocab_size()}", flush=True) + + train_out = args.out / "datasets" / "fineweb10B_sp8192_lossless_caps_caseops_v1_reserved" + train_out.mkdir(parents=True, exist_ok=True) + + val_buf_tokens: list[int] = [] + val_buf_bytes: list[int] = [] + train_buf: list[int] = [] + val_written = 0 + train_written = 0 + n_docs = 0 + + for text in _iter_docs(args.docs): + transformed = encode_lossless_caps_v2(text) + token_ids = [BOS_ID] + sp.encode(transformed, out_type=int) + if n_docs < args.val_docs: + # Validation doc — also compute byte sidecar + byte_counts = _token_original_byte_counts(sp, text, transformed) + val_buf_tokens.extend(token_ids) + val_buf_bytes.append(0) # BOS contributes 0 original bytes + val_buf_bytes.extend(int(b) for b in byte_counts) + if len(val_buf_tokens) >= SHARD_TOKENS: + _write_shard(train_out / f"fineweb_val_{val_written:06d}.bin", + np.array(val_buf_tokens[:SHARD_TOKENS], dtype=np.uint16)) + _write_shard(train_out / f"fineweb_val_bytes_{val_written:06d}.bin", + np.array(val_buf_bytes[:SHARD_TOKENS], dtype=np.uint16)) + val_buf_tokens = val_buf_tokens[SHARD_TOKENS:] + val_buf_bytes = val_buf_bytes[SHARD_TOKENS:] + val_written += 1 + else: + train_buf.extend(token_ids) + if len(train_buf) >= SHARD_TOKENS: + _write_shard(train_out / f"fineweb_train_{train_written:06d}.bin", + np.array(train_buf[:SHARD_TOKENS], dtype=np.uint16)) + train_buf = train_buf[SHARD_TOKENS:] + train_written += 1 + n_docs += 1 + if n_docs % 10_000 == 0: + print(f" processed {n_docs} docs train_shards={train_written} val_shards={val_written}", flush=True) + + # Flush tail buffers into final (possibly short) shards. + if val_buf_tokens: + _write_shard(train_out / f"fineweb_val_{val_written:06d}.bin", + np.array(val_buf_tokens, dtype=np.uint16)) + _write_shard(train_out / f"fineweb_val_bytes_{val_written:06d}.bin", + np.array(val_buf_bytes, dtype=np.uint16)) + if train_buf: + _write_shard(train_out / f"fineweb_train_{train_written:06d}.bin", + np.array(train_buf, dtype=np.uint16)) + + print(f"done. docs={n_docs} train_shards={train_written + (1 if train_buf else 0)} val_shards={val_written + (1 if val_buf_tokens else 0)}") + + +if __name__ == "__main__": + main() diff --git a/records/submission_2026_04_29_b180_tlr56_SUB106/submission.json b/records/submission_2026_04_29_b180_tlr56_SUB106/submission.json new file mode 100644 index 0000000000..eacfc9851e --- /dev/null +++ b/records/submission_2026_04_29_b180_tlr56_SUB106/submission.json @@ -0,0 +1,29 @@ +{ + "author": "vimeto", + "name": "PR #1797 SparseAttnGate + #1855 hparams + QK_GAIN=6.0 + TTT_LORA_RANK=56 + Per-Group lrzip", + "blurb": "🎉 SUB-1.06 result. Best stack: SparseAttnGate + PR #1855 9-pack + QK_GAIN=6.0 + TTT_LORA_RANK=56 (alpha=112) — sweet spot for low-rank LoRA TTT regularization. Single seed=42 BPB **1.05997** (FIRST sub-1.06000; beats #1855 3-seed mean 1.06108 by -0.00111, beats #1797 mean 1.06157 by -0.00160). Per-group lrzip artifact 15,920,473 bytes + 33KB wrapper = ~15.95MB total. Eval host needs `apt install lrzip`.", + "date": "2026-04-29", + "track": "10min_16mb", + "val_bpb": 1.05997, + "submitted_seed": 42, + "seed_results": { + "42": {"val_bpb": 1.05997, "artifact_bytes_pergroup_lrzip": 15920473} + }, + "ttt_lora_rank_sweep_seed_42": { + "48": 1.06005, + "56": 1.05997, + "64": 1.06014, + "80": 1.06046, + "96": 1.06246 + }, + "key_levers": { + "vs_b178_sparse_gate": "QK_GAIN=6.0 (-0.00015), TTT_LORA_RANK=56 (vs 80, -0.00049)", + "ttt_lora_rank_observation": "rank=56 is regularizer-optimum; sweep shows clear inverted-U around 56." + }, + "artifact_bytes": 15953743, + "code_bytes_compressed_pyminified": 33270, + "quant_bytes_pergroup_lrzip": 15920473, + "cap_margin_bytes": 46257, + "compression": "per-group lrzip (PR #1855 style ZPAQ)", + "deps_runtime": ["lrzip", "brotli", "torch", "numpy", "sentencepiece"] +} diff --git a/records/submission_2026_04_29_b180_tlr56_SUB106/train_gpt.py b/records/submission_2026_04_29_b180_tlr56_SUB106/train_gpt.py new file mode 100644 index 0000000000..c0383eef0d --- /dev/null +++ b/records/submission_2026_04_29_b180_tlr56_SUB106/train_gpt.py @@ -0,0 +1,3405 @@ +import base64, collections, copy, fcntl, glob, io, json, lzma, math, os +import datetime +from pathlib import Path +import random, re, subprocess, sys, time, uuid, numpy as np, sentencepiece as spm, torch, torch.distributed as dist, torch.nn.functional as F +from torch import nn +try: + from flash_attn_interface import ( + flash_attn_func as flash_attn_3_func, + flash_attn_varlen_func, + ) + _HAS_FA3 = True +except ImportError: + from flash_attn import flash_attn_func as flash_attn_3_func, flash_attn_varlen_func + _HAS_FA3 = False +from concurrent.futures import ThreadPoolExecutor +import triton +import triton.language as tl +try: + from triton.tools.tensor_descriptor import TensorDescriptor + _HAS_TENSOR_DESCRIPTOR = True +except ImportError: + TensorDescriptor = None + _HAS_TENSOR_DESCRIPTOR = False + +_IS_ROCM = hasattr(torch.version, "hip") and torch.version.hip is not None +_COMPILER_DISABLE = ( + torch.compiler.disable if hasattr(torch, "compiler") else torch._dynamo.disable +) +_ROCM_CLANG = "/opt/rocm-6.2.4/llvm/bin/clang++" +_ROCM_CLANG_CC = "/opt/rocm-6.2.4/llvm/bin/clang" +_INDUCTOR_CONFIG = None + +if _IS_ROCM: + os.environ.setdefault("CC", _ROCM_CLANG_CC) + os.environ.setdefault("CXX", _ROCM_CLANG) + try: + _INDUCTOR_CONFIG = __import__("torch._inductor.config", fromlist=["config"]) + _INDUCTOR_CONFIG.shape_padding = False + cxx_candidates = [] + for candidate in (os.environ.get("CXX"), _ROCM_CLANG, "clang++", "g++"): + if candidate and candidate not in cxx_candidates: + cxx_candidates.append(candidate) + _INDUCTOR_CONFIG.cpp.cxx = tuple(cxx_candidates) + except Exception: + _INDUCTOR_CONFIG = None + + +class Hyperparameters: + data_dir = os.environ.get("DATA_DIR", "./data/") + seed = int(os.environ.get("SEED", 1337)) + run_id = os.environ.get("RUN_ID", str(uuid.uuid4())) + iterations = int(os.environ.get("ITERATIONS", 20000)) + warmdown_frac = float(os.environ.get("WARMDOWN_FRAC", 0.75)) + warmup_steps = int(os.environ.get("WARMUP_STEPS", 20)) + train_batch_tokens = int(os.environ.get("TRAIN_BATCH_TOKENS", 786432)) + train_seq_len = int(os.environ.get("TRAIN_SEQ_LEN", 2048)) + train_log_every = int(os.environ.get("TRAIN_LOG_EVERY", 500)) + max_wallclock_seconds = float(os.environ.get("MAX_WALLCLOCK_SECONDS", 6e2)) + val_batch_tokens = int(os.environ.get("VAL_BATCH_TOKENS", 524288)) + eval_seq_len = int(os.environ.get("EVAL_SEQ_LEN", 2048)) + val_loss_every = int(os.environ.get("VAL_LOSS_EVERY", 4000)) + sliding_window_enabled = bool(int(os.environ.get("SLIDING_WINDOW_ENABLED", "0"))) + vocab_size = int(os.environ.get("VOCAB_SIZE", 8192)) + num_layers = int(os.environ.get("NUM_LAYERS", 11)) + xsa_last_n = int(os.environ.get("XSA_LAST_N", 11)) + model_dim = int(os.environ.get("MODEL_DIM", 512)) + num_kv_heads = int(os.environ.get("NUM_KV_HEADS", 4)) + num_heads = int(os.environ.get("NUM_HEADS", 8)) + mlp_mult = float(os.environ.get("MLP_MULT", 4.0)) + skip_gates_enabled = bool(int(os.environ.get("SKIP_GATES_ENABLED", "1"))) + tie_embeddings = bool(int(os.environ.get("TIE_EMBEDDINGS", "1"))) + logit_softcap = float(os.environ.get("LOGIT_SOFTCAP", 3e1)) + rope_base = float(os.environ.get("ROPE_BASE", 1e4)) + rope_dims = int(os.environ.get("ROPE_DIMS", 16)) + rope_train_seq_len = int(os.environ.get("ROPE_TRAIN_SEQ_LEN", 2048)) + rope_yarn = bool(int(os.environ.get("ROPE_YARN", "0"))) + ln_scale = bool(int(os.environ.get("LN_SCALE", "1"))) + qk_gain_init = float(os.environ.get("QK_GAIN_INIT", 5.25)) + num_loops = int(os.environ.get("NUM_LOOPS", 2)) + loop_start = int(os.environ.get("LOOP_START", 3)) + loop_end = int(os.environ.get("LOOP_END", 5)) + enable_looping_at = float(os.environ.get("ENABLE_LOOPING_AT", 0.35)) + parallel_start_layer = int(os.environ.get("PARALLEL_START_LAYER", 8)) + parallel_final_lane = os.environ.get("PARALLEL_FINAL_LANE", "mean") + min_lr = float(os.environ.get("MIN_LR", 0.0)) + embed_lr = float(os.environ.get("EMBED_LR", 0.6)) + tied_embed_lr = float(os.environ.get("TIED_EMBED_LR", 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.026)) + scalar_lr = float(os.environ.get("SCALAR_LR", 0.02)) + muon_momentum = float(os.environ.get("MUON_MOMENTUM", 0.97)) + muon_backend_steps = int(os.environ.get("MUON_BACKEND_STEPS", 5)) + muon_momentum_warmup_start = float( + os.environ.get("MUON_MOMENTUM_WARMUP_START", 0.92) + ) + muon_momentum_warmup_steps = int(os.environ.get("MUON_MOMENTUM_WARMUP_STEPS", 1500)) + muon_row_normalize = bool(int(os.environ.get("MUON_ROW_NORMALIZE", "1"))) + 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-08)) + grad_clip_norm = float(os.environ.get("GRAD_CLIP_NORM", 0.3)) + eval_stride = int(os.environ.get("EVAL_STRIDE", 64)) + adam_wd = float(os.environ.get("ADAM_WD", 0.02)) + muon_wd = float(os.environ.get("MUON_WD", 0.095)) + embed_wd = float(os.environ.get("EMBED_WD", 0.085)) + ema_decay = float(os.environ.get("EMA_DECAY", 0.9965)) + ttt_enabled = bool(int(os.environ.get("TTT_ENABLED", "1"))) + ttt_lora_rank = int(os.environ.get("TTT_LORA_RANK", 128)) + ttt_lora_lr = float(os.environ.get("TTT_LORA_LR", 0.0001)) + ttt_chunk_size = int(os.environ.get("TTT_CHUNK_SIZE", 48)) + ttt_eval_seq_len = int(os.environ.get("TTT_EVAL_SEQ_LEN", 2048)) + ttt_batch_size = int(os.environ.get("TTT_BATCH_SIZE", 64)) + ttt_grad_steps = int(os.environ.get("TTT_GRAD_STEPS", 1)) + ttt_weight_decay = float(os.environ.get("TTT_WEIGHT_DECAY", 1.0)) + ttt_beta1 = float(os.environ.get("TTT_BETA1", 0)) + ttt_beta2 = float(os.environ.get("TTT_BETA2", 0.999)) + ttt_k_lora = bool(int(os.environ.get("TTT_K_LORA", "1"))) + ttt_mlp_lora = bool(int(os.environ.get("TTT_MLP_LORA", "1"))) + ttt_o_lora = bool(int(os.environ.get("TTT_O_LORA", "1"))) + ttt_optimizer = os.environ.get("TTT_OPTIMIZER", "adam") + ttt_eval_batches = os.environ.get("TTT_EVAL_BATCHES", "") + val_doc_fraction = float(os.environ.get("VAL_DOC_FRACTION", 1.0)) + compressor = os.environ.get("COMPRESSOR", "brotli") + gptq_calibration_batches = int(os.environ.get("GPTQ_CALIBRATION_BATCHES", 16)) + gptq_reserve_seconds = float(os.environ.get("GPTQ_RESERVE_SECONDS", 4.0)) + phased_ttt_enabled = bool(int(os.environ.get("PHASED_TTT_ENABLED", "1"))) + phased_ttt_prefix_docs = int(os.environ.get("PHASED_TTT_PREFIX_DOCS", 2000)) + phased_ttt_num_phases = int(os.environ.get("PHASED_TTT_NUM_PHASES", 3)) + smear_gate_enabled = bool(int(os.environ.get("SMEAR_GATE", "1"))) + smear_gate_width = int(os.environ.get("SMEAR_GATE_WIDTH", 12)) + gate_attn_out = bool(int(os.environ.get("GATE_ATTN_OUT", "1"))) + gate_attn_width = int(os.environ.get("GATE_ATTN_WIDTH", 24)) + sparse_attn_gate_enabled = bool(int(os.environ.get("SPARSE_ATTN_GATE_ENABLED", "0"))) + sparse_attn_gate_scale = float(os.environ.get("SPARSE_ATTN_GATE_SCALE", 1.0)) + gate_window = int(os.environ.get("GATE_WINDOW", 12)) + if sparse_attn_gate_enabled and gate_attn_out: + raise ValueError("SPARSE_ATTN_GATE_ENABLED=1 and GATE_ATTN_OUT=1 are mutually exclusive") + global_ttt_lr = float(os.environ.get("GLOBAL_TTT_LR", 0.001)) + global_ttt_momentum = float(os.environ.get("GLOBAL_TTT_MOMENTUM", 0.9)) + global_ttt_epochs = int(os.environ.get("GLOBAL_TTT_EPOCHS", 1)) + global_ttt_chunk_tokens = int(os.environ.get("GLOBAL_TTT_CHUNK_TOKENS", 32768)) + global_ttt_batch_seqs = int(os.environ.get("GLOBAL_TTT_BATCH_SEQS", 32)) + global_ttt_warmup_start_lr = float(os.environ.get("GLOBAL_TTT_WARMUP_START_LR", 0.0)) + global_ttt_warmup_chunks = int(os.environ.get("GLOBAL_TTT_WARMUP_CHUNKS", 0)) + global_ttt_grad_clip = float(os.environ.get("GLOBAL_TTT_GRAD_CLIP", 1.0)) + global_ttt_respect_doc_boundaries = bool(int(os.environ.get("GLOBAL_TTT_RESPECT_DOC_BOUNDARIES", "1"))) + matrix_bits = int(os.environ.get("MATRIX_BITS", 6)) + embed_bits = int(os.environ.get("EMBED_BITS", 7)) + matrix_clip_sigmas = float(os.environ.get("MATRIX_CLIP_SIGMAS", 12.85)) + embed_clip_sigmas = float(os.environ.get("EMBED_CLIP_SIGMAS", 15.0)) + mlp_clip_sigmas = float(os.environ.get("MLP_CLIP_SIGMAS", 12.0)) + attn_clip_sigmas = float(os.environ.get("ATTN_CLIP_SIGMAS", 13.0)) + # LQER asymmetric rank-r correction on top-K GPTQ residual tensors (PR #1797). + # Post-GPTQ SVD of residual E = W_fp − W_q; keep rank-r A=U[:,:r]·S[:r], + # B=Vh[:r,:]; dequant reconstructs W_q + A @ B. Asym packing: A in INT2 + # per-matrix scalar, B in INT4 groupwise. ~30 KB / ~0.009 BPB recovery. + lqer_enabled = bool(int(os.environ.get("LQER_ENABLED", "0"))) + lqer_rank = int(os.environ.get("LQER_RANK", 4)) + lqer_top_k = int(os.environ.get("LQER_TOP_K", 3)) + lqer_factor_bits = int(os.environ.get("LQER_FACTOR_BITS", 4)) + lqer_asym_enabled = bool(int(os.environ.get("LQER_ASYM_ENABLED", "1"))) + lqer_asym_group = int(os.environ.get("LQER_ASYM_GROUP", "64")) + # Adaptive LQER allocation: greedy byte-budget allocator over (tensor, rank) + # pairs ranked by SVD energy / byte cost. Per-tensor rank is encoded in + # meta string. Used only when LQER_ADAPTIVE=1; falls back to top_K x rank. + lqer_adaptive = bool(int(os.environ.get("LQER_ADAPTIVE", "0"))) + lqer_budget_bytes = int(os.environ.get("LQER_BUDGET_BYTES", "60000")) + # Accept any non-digit separator so sbatch --export comma parsing doesn't + # mangle this. e.g. LQER_CAND_RANKS=4_6_8_10 works. + lqer_cand_ranks = os.environ.get("LQER_CAND_RANKS", "4_6_8_10") + 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")) + is_main_process = rank == 0 + grad_accum_steps = 8 // world_size + datasets_dir = os.environ.get( + "DATA_PATH", os.path.join(data_dir, "datasets", f"fineweb10B_sp{vocab_size}") + ) + train_files = os.path.join(datasets_dir, "fineweb_train_[0-9]*.bin") + val_files = os.path.join(datasets_dir, "fineweb_val_[0-9]*.bin") + tokenizer_path = os.environ.get( + "TOKENIZER_PATH", + os.path.join(data_dir, "tokenizers", f"fineweb_{vocab_size}_bpe.model"), + ) + artifact_dir = os.environ.get("ARTIFACT_DIR", "") + logfile = ( + os.path.join(artifact_dir, f"{run_id}.txt") + if artifact_dir + else f"logs/{run_id}.txt" + ) + model_path = ( + os.path.join(artifact_dir, "final_model.pt") + if artifact_dir + else "final_model.pt" + ) + quantized_model_path = ( + os.path.join(artifact_dir, "final_model.int6.ptz") + if artifact_dir + else "final_model.int6.ptz" + ) + + +_logger_hparams = None + + +def set_logging_hparams(h): + global _logger_hparams + _logger_hparams = h + + +def log(msg, console=True): + if _logger_hparams is None: + print(msg) + return + if _logger_hparams.is_main_process: + if console: + print(msg) + if _logger_hparams.logfile is not None: + with open(_logger_hparams.logfile, "a", encoding="utf-8") as f: + print(msg, file=f) + + +class ValidationData: + def __init__(self, h, device): + self.sp = spm.SentencePieceProcessor(model_file=h.tokenizer_path) + if int(self.sp.vocab_size()) != h.vocab_size: + raise ValueError( + f"VOCAB_SIZE={h.vocab_size} does not match tokenizer vocab_size={int(self.sp.vocab_size())}" + ) + self.val_tokens = load_validation_tokens(h.val_files, h.eval_seq_len) + ( + self.base_bytes_lut, + self.has_leading_space_lut, + self.is_boundary_token_lut, + ) = build_sentencepiece_luts(self.sp, h.vocab_size, device) + + +def _piece_has_pue_marker(piece): + return any(0xE000 <= ord(c) <= 0xF8FF for c in piece) + + +def build_sentencepiece_luts(sp, vocab_size, device): + sp_vocab_size = int(sp.vocab_size()) + assert ( + sp.piece_to_id("▁") != sp.unk_id() + ), "Tokenizer must have '▁' (space) as its own token for correct BPB byte counting" + 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_) + pue_count = 0 + 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_has_pue_marker(piece): + base_bytes_np[token_id] = 0 + pue_count += 1 + continue + if piece.startswith("▁"): + has_leading_space_np[token_id] = True + piece = piece[1:] + base_bytes_np[token_id] = len(piece.encode("utf-8")) + if pue_count > 0: + print(f"build_sentencepiece_luts: {pue_count} PUE-marker tokens zeroed out") + 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, seq_len): + 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 load_data_shard(file): + header_bytes = 256 * np.dtype(" 0: + pos = start + while pos < end: + seg_starts.append(pos) + pos += max_doc_len + else: + seg_starts.append(start) + boundaries = seg_starts + [total_len] + padded_len = get_next_multiple_of_n(len(boundaries), bucket_size) + cu = torch.full((padded_len,), total_len, dtype=torch.int32, device=device) + cu[: len(boundaries)] = torch.tensor(boundaries, dtype=torch.int32, device=device) + seg_ends = seg_starts[1:] + [total_len] + max_seqlen = max(end - start for start, end in zip(seg_starts, seg_ends)) + return cu, max_seqlen + +class DocumentPackingLoader: + _shard_pool = ThreadPoolExecutor(1) + + def __init__(self, h, device, cu_bucket_size=64): + self.rank = h.rank + self.world_size = h.world_size + self.device = device + self.cu_bucket_size = cu_bucket_size + self.max_seq_len = h.train_seq_len + all_files = [Path(p) for p in sorted(glob.glob(h.train_files))] + if not all_files: + raise FileNotFoundError(f"No files found for pattern: {h.train_files}") + self.files = all_files + self.file_iter = iter(self.files) + self._init_shard(load_data_shard(next(self.file_iter))) + self._next_shard = self._submit_next_shard() + self._batch_pool = ThreadPoolExecutor(1) + self._next_batch = None + + def _init_shard(self, tokens): + global BOS_ID + self.tokens = tokens + self.shard_size = tokens.numel() + if BOS_ID is None: + BOS_ID = 1 + self.bos_idx = ( + (tokens == BOS_ID).nonzero(as_tuple=True)[0].to(torch.int64).cpu().numpy() + ) + if self.bos_idx.size == 0: + self.bos_idx = np.array([0], dtype=np.int64) + self.cursor = int(self.bos_idx[0]) + + def _submit_next_shard(self): + try: + path = next(self.file_iter) + return self._shard_pool.submit(load_data_shard, path) + except StopIteration: + return None + + def _advance_shard(self): + if self._next_shard is None: + self.file_iter = iter(self.files) + self._next_shard = self._shard_pool.submit( + load_data_shard, next(self.file_iter) + ) + self._init_shard(self._next_shard.result()) + self._next_shard = self._submit_next_shard() + + def _local_doc_starts(self, local_start, total_len): + lo = np.searchsorted(self.bos_idx, local_start, side="left") + hi = np.searchsorted(self.bos_idx, local_start + total_len, side="left") + return (self.bos_idx[lo:hi] - local_start).tolist() + + def _prepare_batch(self, num_tokens_local, max_seq_len): + per_rank_span = num_tokens_local + 1 + global_span = per_rank_span * self.world_size + while self.cursor + global_span > self.shard_size: + self._advance_shard() + local_start = self.cursor + self.rank * per_rank_span + buf = self.tokens[local_start : local_start + per_rank_span] + inputs = buf[:-1].to(dtype=torch.int64).pin_memory() + targets = buf[1:].to(dtype=torch.int64).pin_memory() + starts = self._local_doc_starts(local_start, inputs.numel()) + cu_seqlens, max_seqlen = _build_cu_seqlens( + starts, inputs.numel(), inputs.device, max_seq_len, self.cu_bucket_size + ) + cu_seqlens = cu_seqlens.pin_memory() + self.cursor += global_span + return inputs, targets, cu_seqlens, max_seqlen + + def next_batch(self, global_tokens, grad_accum_steps): + num_tokens_local = global_tokens // (self.world_size * grad_accum_steps) + if self._next_batch is not None: + inputs, targets, cu_seqlens, max_seqlen = self._next_batch.result() + else: + inputs, targets, cu_seqlens, max_seqlen = self._prepare_batch( + num_tokens_local, self.max_seq_len + ) + self._next_batch = self._batch_pool.submit( + self._prepare_batch, num_tokens_local, self.max_seq_len + ) + return ( + inputs[None].to(self.device, non_blocking=True), + targets[None].to(self.device, non_blocking=True), + cu_seqlens.to(self.device, non_blocking=True), + max_seqlen, + ) + + +class ShuffledSequenceLoader: + def __init__(self, h, device): + self.world_size = h.world_size + self.seq_len = h.train_seq_len + self.device = device + all_files = [Path(p) for p in sorted(glob.glob(h.train_files))] + if not all_files: + raise FileNotFoundError(f"No files found for pattern: {h.train_files}") + self.files = all_files[h.rank :: h.world_size] + self.rng = np.random.Generator(np.random.PCG64(h.rank)) + self.num_tokens = [_read_num_tokens(f) for f in self.files] + self.start_inds = [[] for _ in self.files] + for si in range(len(self.files)): + self._reset_shard(si) + + def _reset_shard(self, si): + max_phase = min( + self.seq_len - 1, max(0, self.num_tokens[si] - self.seq_len - 1) + ) + phase = int(self.rng.integers(max_phase + 1)) if max_phase > 0 else 0 + num_sequences = (self.num_tokens[si] - 1 - phase) // self.seq_len + sequence_order = self.rng.permutation(num_sequences) + self.start_inds[si] = (phase + sequence_order * self.seq_len).tolist() + + def next_batch(self, global_tokens, grad_accum_steps): + device_tokens = global_tokens // (self.world_size * grad_accum_steps) + device_batch_size = device_tokens // self.seq_len + remaining = np.array([len(s) for s in self.start_inds], dtype=np.float64) + x = torch.empty((device_batch_size, self.seq_len), dtype=torch.int64) + y = torch.empty((device_batch_size, self.seq_len), dtype=torch.int64) + for bi in range(device_batch_size): + total = remaining.sum() + if total <= 0: + for si in range(len(self.files)): + self._reset_shard(si) + remaining = np.array( + [len(s) for s in self.start_inds], dtype=np.float64 + ) + total = remaining.sum() + probs = remaining / total + si = int(self.rng.choice(len(self.files), p=probs)) + start_ind = self.start_inds[si].pop() + remaining[si] -= 1 + mm = _get_shard_memmap(self.files[si]) + window = torch.as_tensor( + np.array(mm[start_ind : start_ind + self.seq_len + 1], dtype=np.int64) + ) + x[bi] = window[:-1] + y[bi] = window[1:] + return x.to(self.device, non_blocking=True), y.to( + self.device, non_blocking=True + ) + + +class RMSNorm(nn.Module): + def __init__(self, eps=None): + super().__init__() + self.eps = eps + + def forward(self, x): + return F.rms_norm(x, (x.size(-1),), eps=self.eps) + + +class CastedLinear(nn.Linear): + def forward(self, x): + w = self.weight.to(x.dtype) + bias = self.bias.to(x.dtype) if self.bias is not None else None + return F.linear(x, w, bias) + + +@triton.jit +def linear_leaky_relu_square_kernel( + a_desc, + b_desc, + c_desc, + aux_desc, + M, + N, + K, + BLOCK_SIZE_M: tl.constexpr, + BLOCK_SIZE_N: tl.constexpr, + BLOCK_SIZE_K: tl.constexpr, + NUM_SMS: tl.constexpr, + FORWARD: tl.constexpr, +): + dtype = tl.bfloat16 + start_pid = tl.program_id(axis=0) + num_pid_m = tl.cdiv(M, BLOCK_SIZE_M) + num_pid_n = tl.cdiv(N, BLOCK_SIZE_N) + k_tiles = tl.cdiv(K, BLOCK_SIZE_K) + num_tiles = num_pid_m * num_pid_n + tile_id_c = start_pid - NUM_SMS + for tile_id in tl.range(start_pid, num_tiles, NUM_SMS, flatten=True): + pid_m = tile_id // num_pid_n + pid_n = tile_id % num_pid_n + offs_am = pid_m * BLOCK_SIZE_M + offs_bn = pid_n * BLOCK_SIZE_N + accumulator = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=tl.float32) + for ki in range(k_tiles): + offs_k = ki * BLOCK_SIZE_K + a = a_desc.load([offs_am, offs_k]) + b = b_desc.load([offs_bn, offs_k]) + accumulator = tl.dot(a, b.T, accumulator) + tile_id_c += NUM_SMS + offs_am_c = offs_am + offs_bn_c = offs_bn + acc = tl.reshape(accumulator, (BLOCK_SIZE_M, 2, BLOCK_SIZE_N // 2)) + acc = tl.permute(acc, (0, 2, 1)) + acc0, acc1 = tl.split(acc) + c0 = acc0.to(dtype) + c1 = acc1.to(dtype) + if not FORWARD: + pre0 = aux_desc.load([offs_am_c, offs_bn_c]) + pre1 = aux_desc.load([offs_am_c, offs_bn_c + BLOCK_SIZE_N // 2]) + c0 = c0 * tl.where(pre0 > 0, 2.0 * pre0, 0.5 * pre0) + c1 = c1 * tl.where(pre1 > 0, 2.0 * pre1, 0.5 * pre1) + c_desc.store([offs_am_c, offs_bn_c], c0) + c_desc.store([offs_am_c, offs_bn_c + BLOCK_SIZE_N // 2], c1) + if FORWARD: + aux0 = tl.where(c0 > 0, c0, 0.5 * c0) + aux1 = tl.where(c1 > 0, c1, 0.5 * c1) + aux_desc.store([offs_am_c, offs_bn_c], aux0 * aux0) + aux_desc.store([offs_am_c, offs_bn_c + BLOCK_SIZE_N // 2], aux1 * aux1) + + +def linear_leaky_relu_square(a, b, aux=None): + M, K = a.shape + N, K2 = b.shape + assert K == K2 + c = torch.empty((M, N), device=a.device, dtype=a.dtype) + forward = aux is None + if aux is None: + aux = torch.empty((M, N), device=a.device, dtype=a.dtype) + num_sms = torch.cuda.get_device_properties(a.device).multi_processor_count + BLOCK_SIZE_M, BLOCK_SIZE_N, BLOCK_SIZE_K = 128, 256, 64 + num_stages = 4 if forward else 3 + a_desc = TensorDescriptor.from_tensor(a, [BLOCK_SIZE_M, BLOCK_SIZE_K]) + b_desc = TensorDescriptor.from_tensor(b, [BLOCK_SIZE_N, BLOCK_SIZE_K]) + c_desc = TensorDescriptor.from_tensor(c, [BLOCK_SIZE_M, BLOCK_SIZE_N // 2]) + aux_desc = TensorDescriptor.from_tensor(aux, [BLOCK_SIZE_M, BLOCK_SIZE_N // 2]) + grid = lambda _meta: ( + min(num_sms, triton.cdiv(M, BLOCK_SIZE_M) * triton.cdiv(N, BLOCK_SIZE_N)), + ) + linear_leaky_relu_square_kernel[grid]( + a_desc, + b_desc, + c_desc, + aux_desc, + M, + N, + K, + BLOCK_SIZE_M=BLOCK_SIZE_M, + BLOCK_SIZE_N=BLOCK_SIZE_N, + BLOCK_SIZE_K=BLOCK_SIZE_K, + NUM_SMS=num_sms, + FORWARD=forward, + num_stages=num_stages, + num_warps=8, + ) + if forward: + return c, aux + return c + + +class FusedLinearLeakyReLUSquareFunction(torch.autograd.Function): + @staticmethod + def forward(ctx, x, w1, w2): + x_flat = x.reshape(-1, x.shape[-1]) + pre, post = linear_leaky_relu_square(x_flat, w1) + out = F.linear(post, w2) + ctx.save_for_backward(x, w1, w2, pre, post) + return out.view(*x.shape[:-1], out.shape[-1]) + + @staticmethod + def backward(ctx, grad_output): + x, w1, w2, pre, post = ctx.saved_tensors + x_flat = x.reshape(-1, x.shape[-1]) + grad_output_flat = grad_output.reshape(-1, grad_output.shape[-1]) + dw2 = grad_output_flat.T @ post + dpre = linear_leaky_relu_square(grad_output_flat, w2.T.contiguous(), aux=pre) + dw1 = dpre.T @ x_flat + dx = dpre @ w1 + return dx.view_as(x), dw1, dw2 + + +FusedLeakyReLUSquareMLP = FusedLinearLeakyReLUSquareFunction.apply + + +class Rotary(nn.Module): + def __init__(self, dim, base=1e4, train_seq_len=1024, rope_dims=0, yarn=True): + super().__init__() + self.dim = dim + self.base = base + self.train_seq_len = train_seq_len + self.yarn = yarn + 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 = None + self._sin_cached = None + + def forward(self, seq_len, device, dtype): + 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 self.yarn and 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.float().to(device) + t = torch.arange(seq_len, device=device, dtype=torch.float32) + 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[:, :seq_len].to(dtype=dtype), self._sin_cached[:, :seq_len].to(dtype=dtype) + + +def apply_rotary_emb(x, cos, sin, rope_dims=0): + 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) + + +def _flash_attn_varlen(q, k, v, cu_seqlens, max_seqlen): + return flash_attn_varlen_func( + q, + k, + v, + cu_seqlens_q=cu_seqlens, + cu_seqlens_k=cu_seqlens, + max_seqlen_q=max_seqlen, + max_seqlen_k=max_seqlen, + causal=True, + window_size=(-1, -1), + ) + + +if _IS_ROCM: + _flash_attn_varlen = _COMPILER_DISABLE(_flash_attn_varlen) + + +class CausalSelfAttention(nn.Module): + def __init__( + self, dim, num_heads, num_kv_heads, rope_base, qk_gain_init, train_seq_len, yarn=True + ): + 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") + 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=train_seq_len, yarn=yarn) + self.use_xsa = False + self.gate_attn_out = False + self.gate_attn_width = 24 + self.sparse_attn_gate = False + self.sparse_attn_gate_scale = 1.0 + self.gate_window = 12 + + 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, q_w, k_w, v_w, out_w, cu_seqlens=None, max_seqlen=0): + 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)).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] + if cu_seqlens is not None: + y = _flash_attn_varlen(q[0], k[0], v[0], cu_seqlens, max_seqlen)[None] + else: + y = flash_attn_3_func(q, k, v, causal=True) + if self.use_xsa: + y = self._xsa_efficient(y, v) + y = y.reshape(bsz, seqlen, self.num_heads, self.head_dim) + if self.gate_attn_out: + g = 2.0 * torch.sigmoid(self.attn_gate_proj(x[:, :, :self.gate_attn_width])) + y = y * g.unsqueeze(-1) + elif self.sparse_attn_gate: + gate_in = x[..., : self.gate_window].contiguous() + g = torch.sigmoid(self.sparse_attn_gate_scale * F.linear(gate_in, self.attn_gate_w.to(x.dtype))) + y = y * g.unsqueeze(-1) + y = y.reshape(bsz, seqlen, dim) + self._last_proj_input = y.detach() if getattr(self, "_calib", False) else None + return F.linear(y, out_w.to(x.dtype)) + + +class MLP(nn.Module): + def __init__(self, dim, mlp_mult): + super().__init__() + self.use_fused = _HAS_TENSOR_DESCRIPTOR + + def forward(self, x, up_w, down_w): + if self.training and self.use_fused: + return FusedLeakyReLUSquareMLP(x, up_w.to(x.dtype), down_w.to(x.dtype)) + hidden = F.leaky_relu(F.linear(x, up_w.to(x.dtype)), negative_slope=0.5).square() + self._last_down_input = hidden.detach() if getattr(self, "_calib", False) else None + return F.linear(hidden, down_w.to(x.dtype)) + + +class Block(nn.Module): + def __init__( + self, + dim, + num_heads, + num_kv_heads, + mlp_mult, + rope_base, + qk_gain_init, + train_seq_len, + layer_idx=0, + ln_scale=False, + yarn=True, + ): + super().__init__() + self.attn_norm = RMSNorm() + self.mlp_norm = RMSNorm() + self.attn = CausalSelfAttention( + dim, num_heads, num_kv_heads, rope_base, qk_gain_init, train_seq_len, yarn=yarn + ) + 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 + + def forward(self, x, x0, q_w, k_w, v_w, out_w, up_w, down_w, cu_seqlens=None, max_seqlen=0): + 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, + q_w, k_w, v_w, out_w, + cu_seqlens=cu_seqlens, + max_seqlen=max_seqlen, + ) + 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) + return x_out + +class GPT(nn.Module): + def __init__(self, h): + super().__init__() + if h.logit_softcap <= 0.0: + raise ValueError(f"logit_softcap must be positive, got {h.logit_softcap}") + self.tie_embeddings = h.tie_embeddings + self.tied_embed_init_std = h.tied_embed_init_std + self.logit_softcap = h.logit_softcap + self.tok_emb = nn.Embedding(h.vocab_size, h.model_dim) + self.num_layers = h.num_layers + head_dim = h.model_dim // h.num_heads + kv_dim = h.num_kv_heads * head_dim + hidden_dim = int(h.mlp_mult * h.model_dim) + self.qo_bank = nn.Parameter(torch.empty(2 * h.num_layers, h.model_dim, h.model_dim)) + self.kv_bank = nn.Parameter(torch.empty(2 * h.num_layers, kv_dim, h.model_dim)) + self.mlp_up_bank = nn.Parameter(torch.empty(h.num_layers, hidden_dim, h.model_dim)) + self.mlp_down_bank = nn.Parameter(torch.empty(h.num_layers, h.model_dim, hidden_dim)) + self.num_encoder_layers = h.num_layers // 2 + self.num_decoder_layers = h.num_layers - self.num_encoder_layers + self.blocks = nn.ModuleList( + [ + Block( + h.model_dim, + h.num_heads, + h.num_kv_heads, + h.mlp_mult, + h.rope_base, + h.qk_gain_init, + h.train_seq_len, + layer_idx=i, + ln_scale=h.ln_scale, + yarn=h.rope_yarn, + ) + for i in range(h.num_layers) + ] + ) + if h.rope_dims > 0: + head_dim = h.model_dim // h.num_heads + for block in self.blocks: + block.attn.rope_dims = h.rope_dims + block.attn.rotary = Rotary( + head_dim, + base=h.rope_base, + train_seq_len=h.train_seq_len, + rope_dims=h.rope_dims, + yarn=h.rope_yarn, + ) + self.final_norm = RMSNorm() + self.lm_head = ( + None + if h.tie_embeddings + else CastedLinear(h.model_dim, h.vocab_size, bias=False) + ) + if self.lm_head is not None: + self.lm_head._zero_init = True + if h.xsa_last_n > 0: + for i in range(max(0, h.num_layers - h.xsa_last_n), h.num_layers): + self.blocks[i].attn.use_xsa = True + if h.gate_attn_out: + for block in self.blocks: + block.attn.gate_attn_out = True + block.attn.gate_attn_width = h.gate_attn_width + block.attn.attn_gate_proj = CastedLinear( + h.gate_attn_width, h.num_heads, bias=False + ) + nn.init.zeros_(block.attn.attn_gate_proj.weight) + if h.sparse_attn_gate_enabled: + for block in self.blocks: + block.attn.sparse_attn_gate = True + block.attn.sparse_attn_gate_scale = h.sparse_attn_gate_scale + block.attn.gate_window = h.gate_window + W = torch.zeros(h.num_heads, h.gate_window, dtype=torch.float32) + block.attn.attn_gate_w = nn.Parameter(W) + self.looping_active = False + if h.num_loops > 0: + loop_seg = list(range(h.loop_start, h.loop_end + 1)) + all_indices = list(range(h.loop_start)) + for _ in range(h.num_loops + 1): + all_indices.extend(loop_seg) + all_indices.extend(range(h.loop_end + 1, h.num_layers)) + num_enc = len(all_indices) // 2 + self.encoder_indices = all_indices[:num_enc] + self.decoder_indices = all_indices[num_enc:] + else: + self.encoder_indices = list(range(self.num_encoder_layers)) + self.decoder_indices = list(range(self.num_encoder_layers, h.num_layers)) + self.num_skip_weights = min( + len(self.encoder_indices), len(self.decoder_indices) + ) + self.skip_weights = nn.Parameter( + torch.ones(self.num_skip_weights, h.model_dim, dtype=torch.float32) + ) + self.skip_gates = ( + nn.Parameter( + torch.zeros(self.num_skip_weights, h.model_dim, dtype=torch.float32) + ) + if h.skip_gates_enabled + else None + ) + self.parallel_start_layer = h.parallel_start_layer + self.parallel_final_lane = h.parallel_final_lane.lower() + self.parallel_post_lambdas = nn.Parameter( + torch.ones(h.num_layers, 2, 2, dtype=torch.float32) + ) + self.parallel_resid_lambdas = nn.Parameter( + torch.full((h.num_layers, 2), 1.1, dtype=torch.float32) + ) + self.smear_gate_enabled = h.smear_gate_enabled + if self.smear_gate_enabled: + self.smear_width = h.smear_gate_width + self.smear_gate = CastedLinear(self.smear_width, 1, bias=False) + nn.init.zeros_(self.smear_gate.weight) + self.smear_lambda = nn.Parameter(torch.zeros(1, dtype=torch.float32)) + self._init_weights() + + def _init_weights(self): + 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) + for i in range(n): + nn.init.orthogonal_(self.qo_bank.data[i], gain=1.0) + nn.init.zeros_(self.qo_bank.data[n + i]) + self.qo_bank.data[n + i].mul_(proj_scale) + nn.init.orthogonal_(self.kv_bank.data[i], gain=1.0) + nn.init.orthogonal_(self.kv_bank.data[n + i], gain=1.0) + for i in range(n): + nn.init.orthogonal_(self.mlp_up_bank.data[i], gain=1.0) + nn.init.zeros_(self.mlp_down_bank.data[i]) + self.mlp_down_bank.data[i].mul_(proj_scale) + 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 _bank_weights(self, i): + n = self.num_layers + return ( + self.qo_bank[i], + self.kv_bank[i], + self.kv_bank[n + i], + self.qo_bank[n + i], + self.mlp_up_bank[i], + self.mlp_down_bank[i], + ) + + def _parallel_block( + self, block_idx, lane0, lane1, x0, + q_w, k_w, v_w, out_w, up_w, down_w, + cu_seqlens=None, max_seqlen=0, + ): + block = self.blocks[block_idx] + mix = block.resid_mix.to(dtype=lane0.dtype) + attn_read = mix[0][None, None, :] * lane0 + mix[1][None, None, :] * x0 + attn_out = block.attn( + block.attn_norm(attn_read) * block.ln_scale_factor, + q_w, k_w, v_w, out_w, + cu_seqlens=cu_seqlens, max_seqlen=max_seqlen, + ) + attn_out = block.attn_scale.to(dtype=attn_out.dtype)[None, None, :] * attn_out + mlp_read = lane1 + mlp_out = block.mlp_scale.to(dtype=lane1.dtype)[None, None, :] * block.mlp( + block.mlp_norm(mlp_read) * block.ln_scale_factor, up_w, down_w + ) + attn_resid = self.parallel_resid_lambdas[block_idx, 0].to(dtype=lane0.dtype) + attn_post = self.parallel_post_lambdas[block_idx, 0].to(dtype=lane0.dtype) + mlp_resid = self.parallel_resid_lambdas[block_idx, 1].to(dtype=lane0.dtype) + mlp_post = self.parallel_post_lambdas[block_idx, 1].to(dtype=lane0.dtype) + lane0 = attn_resid * lane0 + attn_post[0] * attn_out + mlp_post[0] * mlp_out + lane1 = mlp_resid * lane1 + attn_post[1] * attn_out + mlp_post[1] * mlp_out + return lane0, lane1 + + def _final_parallel_hidden(self, lane0, lane1): + if self.parallel_final_lane == "mlp": + return lane1 + if self.parallel_final_lane == "attn": + return lane0 + return 0.5 * (lane0 + lane1) + + def forward_logits(self, input_ids, cu_seqlens=None, max_seqlen=0): + x = self.tok_emb(input_ids) + if self.smear_gate_enabled and x.dim() == 3 and x.size(1) > 1: + sl = self.smear_lambda.to(dtype=x.dtype) + g = sl * torch.sigmoid(self.smear_gate(x[:, 1:, :self.smear_width])) + # BOS mask (PR #1851/#1855): zero smear at positions where the + # CURRENT token is BOS — prevents pulling info into a new doc's + # opening token from the previous doc's last token. + global BOS_ID + if BOS_ID is None: + BOS_ID = 1 + not_bos_cur = (input_ids[:, 1:] != BOS_ID).to(dtype=g.dtype).unsqueeze(-1) + g = g * not_bos_cur + x = torch.cat([x[:, :1], x[:, 1:] + g * x[:, :-1]], dim=1) + x = F.rms_norm(x, (x.size(-1),)) + x0 = x + skips = [] + enc_iter = ( + self.encoder_indices + if self.looping_active + else range(self.num_encoder_layers) + ) + dec_iter = ( + self.decoder_indices + if self.looping_active + else range( + self.num_encoder_layers, + self.num_encoder_layers + self.num_decoder_layers, + ) + ) + for i in enc_iter: + q_w, k_w, v_w, out_w, up_w, down_w = self._bank_weights(i) + x = self.blocks[i](x, x0, q_w, k_w, v_w, out_w, up_w, down_w, cu_seqlens=cu_seqlens, max_seqlen=max_seqlen) + skips.append(x) + psl = self.parallel_start_layer + lane0 = None + lane1 = None + for skip_idx, i in enumerate(dec_iter): + q_w, k_w, v_w, out_w, up_w, down_w = self._bank_weights(i) + if i >= psl and psl > 0: + if lane0 is None: + lane0 = x + lane1 = x + if skip_idx < self.num_skip_weights and skips: + skip = skips.pop() + w = self.skip_weights[skip_idx].to(dtype=lane0.dtype)[None, None, :] + if self.skip_gates is not None: + g = torch.sigmoid(self.skip_gates[skip_idx].to(dtype=lane0.dtype))[None, None, :] + lane0 = torch.lerp(w * skip, lane0, g) + else: + lane0 = lane0 + w * skip + lane0, lane1 = self._parallel_block( + i, lane0, lane1, x0, q_w, k_w, v_w, out_w, up_w, down_w, + cu_seqlens=cu_seqlens, max_seqlen=max_seqlen, + ) + else: + if skip_idx < self.num_skip_weights and skips: + scaled_skip = ( + self.skip_weights[skip_idx].to(dtype=x.dtype)[None, None, :] + * skips.pop() + ) + if self.skip_gates is not None: + g = torch.sigmoid(self.skip_gates[skip_idx].to(dtype=x.dtype))[None, None, :] + x = torch.lerp(scaled_skip, x, g) + else: + x = x + scaled_skip + x = self.blocks[i](x, x0, q_w, k_w, v_w, out_w, up_w, down_w, cu_seqlens=cu_seqlens, max_seqlen=max_seqlen) + if lane0 is not None: + x = self._final_parallel_hidden(lane0, lane1) + x = self.final_norm(x) + if self.tie_embeddings: + logits_proj = F.linear(x, self.tok_emb.weight) + else: + logits_proj = self.lm_head(x) + return self.logit_softcap * torch.tanh(logits_proj / self.logit_softcap) + + def forward(self, input_ids, target_ids, cu_seqlens=None, max_seqlen=0): + logits = self.forward_logits( + input_ids, cu_seqlens=cu_seqlens, max_seqlen=max_seqlen + ) + return F.cross_entropy( + logits.reshape(-1, logits.size(-1)).float(), + target_ids.reshape(-1), + reduction="mean", + ) + + def forward_ttt(self, input_ids, target_ids, lora): + x = self.tok_emb(input_ids) + if self.smear_gate_enabled and x.dim() == 3 and x.size(1) > 1: + sl = self.smear_lambda.to(dtype=x.dtype) + g = sl * torch.sigmoid(self.smear_gate(x[:, 1:, :self.smear_width])) + # BOS mask (PR #1851/#1855): zero smear at positions where the + # CURRENT token is BOS — prevents pulling info into a new doc's + # opening token from the previous doc's last token. + global BOS_ID + if BOS_ID is None: + BOS_ID = 1 + not_bos_cur = (input_ids[:, 1:] != BOS_ID).to(dtype=g.dtype).unsqueeze(-1) + g = g * not_bos_cur + x = torch.cat([x[:, :1], x[:, 1:] + g * x[:, :-1]], dim=1) + x = F.rms_norm(x, (x.size(-1),)) + x0 = x + skips = [] + enc_iter = ( + self.encoder_indices + if self.looping_active + else list(range(self.num_encoder_layers)) + ) + dec_iter = ( + self.decoder_indices + if self.looping_active + else list( + range( + self.num_encoder_layers, + self.num_encoder_layers + self.num_decoder_layers, + ) + ) + ) + slot = 0 + for i in enc_iter: + q_w, k_w, v_w, out_w, up_w, down_w = self._bank_weights(i) + x = self._block_with_lora(self.blocks[i], x, x0, lora, slot, q_w, k_w, v_w, out_w, up_w, down_w) + slot += 1 + skips.append(x) + psl = self.parallel_start_layer + lane0 = None + lane1 = None + for skip_idx, i in enumerate(dec_iter): + q_w, k_w, v_w, out_w, up_w, down_w = self._bank_weights(i) + if i >= psl and psl > 0: + if lane0 is None: + lane0 = x + lane1 = x + if skip_idx < self.num_skip_weights and skips: + skip = skips.pop() + w = self.skip_weights[skip_idx].to(dtype=lane0.dtype)[None, None, :] + if self.skip_gates is not None: + g = torch.sigmoid(self.skip_gates[skip_idx].to(dtype=lane0.dtype))[None, None, :] + lane0 = torch.lerp(w * skip, lane0, g) + else: + lane0 = lane0 + w * skip + lane0, lane1 = self._parallel_block_with_lora( + i, lane0, lane1, x0, lora, slot, + q_w, k_w, v_w, out_w, up_w, down_w, + ) + else: + if skip_idx < self.num_skip_weights and skips: + scaled_skip = ( + self.skip_weights[skip_idx].to(dtype=x.dtype)[None, None, :] + * skips.pop() + ) + if self.skip_gates is not None: + g = torch.sigmoid(self.skip_gates[skip_idx].to(dtype=x.dtype))[None, None, :] + x = torch.lerp(scaled_skip, x, g) + else: + x = x + scaled_skip + x = self._block_with_lora(self.blocks[i], x, x0, lora, slot, q_w, k_w, v_w, out_w, up_w, down_w) + slot += 1 + if lane0 is not None: + x = self._final_parallel_hidden(lane0, lane1) + x = self.final_norm(x) + if self.tie_embeddings: + logits = F.linear(x, self.tok_emb.weight) + else: + logits = self.lm_head(x) + logits = logits + lora.lm_head_lora(x) + logits = self.logit_softcap * torch.tanh(logits / self.logit_softcap) + bsz, sl, V = logits.shape + return F.cross_entropy( + logits.float().reshape(-1, V), target_ids.reshape(-1), reduction="none" + ).reshape(bsz, sl) + + def _block_with_lora(self, block, x, x0, lora, slot, q_w, k_w, v_w, out_w, up_w, down_w): + mix = block.resid_mix.to(dtype=x.dtype) + x_in = mix[0][None, None, :] * x + mix[1][None, None, :] * x0 + n = block.attn_norm(x_in) * block.ln_scale_factor + attn = block.attn + bsz, seqlen, dim = n.shape + q = (F.linear(n, q_w.to(n.dtype)) + lora.q_loras[slot](n)).reshape( + bsz, seqlen, attn.num_heads, attn.head_dim + ) + k = F.linear(n, k_w.to(n.dtype)) + if lora.k_loras is not None: + k = k + lora.k_loras[slot](n) + k = k.reshape(bsz, seqlen, attn.num_kv_heads, attn.head_dim) + v = (F.linear(n, v_w.to(n.dtype)) + lora.v_loras[slot](n)).reshape( + bsz, seqlen, attn.num_kv_heads, attn.head_dim + ) + q = F.rms_norm(q, (q.size(-1),)) + k = F.rms_norm(k, (k.size(-1),)) + cos, sin = attn.rotary(seqlen, n.device, q.dtype) + q = apply_rotary_emb(q, cos, sin, attn.rope_dims) + k = apply_rotary_emb(k, cos, sin, attn.rope_dims) + q = q * attn.q_gain.to(dtype=q.dtype)[None, None, :, None] + y = flash_attn_3_func(q, k, v, causal=True) + if attn.use_xsa: + y = attn._xsa_efficient(y, v) + if attn.gate_attn_out: + g = 2.0 * torch.sigmoid(attn.attn_gate_proj(n[:, :, :attn.gate_attn_width])) + y = y * g.unsqueeze(-1) + elif attn.sparse_attn_gate: + gate_in = n[..., : attn.gate_window].contiguous() + g = torch.sigmoid(attn.sparse_attn_gate_scale * F.linear(gate_in, attn.attn_gate_w.to(n.dtype))) + y = y * g.unsqueeze(-1) + y = y.reshape(bsz, seqlen, dim) + attn_out = F.linear(y, out_w.to(n.dtype)) + if lora.o_loras is not None: + attn_out = attn_out + lora.o_loras[slot](n) + x_out = x_in + block.attn_scale.to(dtype=x_in.dtype)[None, None, :] * attn_out + mlp_n = block.mlp_norm(x_out) * block.ln_scale_factor + mlp_out = block.mlp(mlp_n, up_w, down_w) + if lora.mlp_loras is not None: + mlp_out = mlp_out + lora.mlp_loras[slot](mlp_n) + x_out = x_out + block.mlp_scale.to(dtype=x_out.dtype)[None, None, :] * mlp_out + return x_out + + def _parallel_block_with_lora( + self, block_idx, lane0, lane1, x0, lora, slot, + q_w, k_w, v_w, out_w, up_w, down_w, + ): + block = self.blocks[block_idx] + mix = block.resid_mix.to(dtype=lane0.dtype) + attn_read = mix[0][None, None, :] * lane0 + mix[1][None, None, :] * x0 + n = block.attn_norm(attn_read) * block.ln_scale_factor + attn = block.attn + bsz, seqlen, dim = n.shape + q = (F.linear(n, q_w.to(n.dtype)) + lora.q_loras[slot](n)).reshape( + bsz, seqlen, attn.num_heads, attn.head_dim + ) + k = F.linear(n, k_w.to(n.dtype)) + if lora.k_loras is not None: + k = k + lora.k_loras[slot](n) + k = k.reshape(bsz, seqlen, attn.num_kv_heads, attn.head_dim) + v = (F.linear(n, v_w.to(n.dtype)) + lora.v_loras[slot](n)).reshape( + bsz, seqlen, attn.num_kv_heads, attn.head_dim + ) + q = F.rms_norm(q, (q.size(-1),)) + k = F.rms_norm(k, (k.size(-1),)) + cos, sin = attn.rotary(seqlen, n.device, q.dtype) + q = apply_rotary_emb(q, cos, sin, attn.rope_dims) + k = apply_rotary_emb(k, cos, sin, attn.rope_dims) + q = q * attn.q_gain.to(dtype=q.dtype)[None, None, :, None] + y = flash_attn_3_func(q, k, v, causal=True) + if attn.use_xsa: + y = attn._xsa_efficient(y, v) + if attn.gate_attn_out: + g = 2.0 * torch.sigmoid(attn.attn_gate_proj(n[:, :, :attn.gate_attn_width])) + y = y * g.unsqueeze(-1) + elif attn.sparse_attn_gate: + gate_in = n[..., : attn.gate_window].contiguous() + g = torch.sigmoid(attn.sparse_attn_gate_scale * F.linear(gate_in, attn.attn_gate_w.to(n.dtype))) + y = y * g.unsqueeze(-1) + y = y.reshape(bsz, seqlen, dim) + attn_out = F.linear(y, out_w.to(n.dtype)) + if lora.o_loras is not None: + attn_out = attn_out + lora.o_loras[slot](n) + attn_out = block.attn_scale.to(dtype=attn_out.dtype)[None, None, :] * attn_out + mlp_read = lane1 + mlp_n = block.mlp_norm(mlp_read) * block.ln_scale_factor + mlp_out = block.mlp(mlp_n, up_w, down_w) + if lora.mlp_loras is not None: + mlp_out = mlp_out + lora.mlp_loras[slot](mlp_n) + mlp_out = block.mlp_scale.to(dtype=lane1.dtype)[None, None, :] * mlp_out + attn_resid = self.parallel_resid_lambdas[block_idx, 0].to(dtype=lane0.dtype) + attn_post = self.parallel_post_lambdas[block_idx, 0].to(dtype=lane0.dtype) + mlp_resid = self.parallel_resid_lambdas[block_idx, 1].to(dtype=lane0.dtype) + mlp_post = self.parallel_post_lambdas[block_idx, 1].to(dtype=lane0.dtype) + lane0 = attn_resid * lane0 + attn_post[0] * attn_out + mlp_post[0] * mlp_out + lane1 = mlp_resid * lane1 + attn_post[1] * attn_out + mlp_post[1] * mlp_out + return lane0, lane1 + + +class BatchedLinearLoRA(nn.Module): + _ALPHA = float(os.environ.get("TTT_LORA_ALPHA", "144")) + + def __init__(self, bsz, in_features, out_features, rank): + super().__init__() + self._bound = 1.0 / math.sqrt(in_features) + self._scale = self._ALPHA / rank + self.A = nn.Parameter( + torch.empty(bsz, rank, in_features).uniform_(-self._bound, self._bound) + ) + self.B = nn.Parameter(torch.zeros(bsz, out_features, rank)) + + _WARM_START_A = bool(int(os.environ.get("TTT_WARM_START_A", "1"))) + + def reset(self): + with torch.no_grad(): + if not self._WARM_START_A: + self.A.uniform_(-self._bound, self._bound) + self.B.zero_() + + def forward(self, x): + return ((x @ self.A.transpose(1, 2)) @ self.B.transpose(1, 2)) * self._scale + + +class BatchedTTTLoRA(nn.Module): + def __init__(self, bsz, model, rank, k_lora=True, mlp_lora=True, o_lora=True): + super().__init__() + self.bsz = bsz + dim = model.qo_bank.shape[-1] + vocab = model.tok_emb.num_embeddings + if getattr(model, "looping_active", False): + num_slots = len(model.encoder_indices) + len(model.decoder_indices) + else: + num_slots = len(model.blocks) + kv_dim = model.blocks[0].attn.num_kv_heads * ( + dim // model.blocks[0].attn.num_heads + ) + embed_dim = model.tok_emb.embedding_dim + self.lm_head_lora = BatchedLinearLoRA(bsz, embed_dim, vocab, rank) + self.q_loras = nn.ModuleList( + [BatchedLinearLoRA(bsz, dim, dim, rank) for _ in range(num_slots)] + ) + self.v_loras = nn.ModuleList( + [BatchedLinearLoRA(bsz, dim, kv_dim, rank) for _ in range(num_slots)] + ) + self.k_loras = ( + nn.ModuleList( + [BatchedLinearLoRA(bsz, dim, kv_dim, rank) for _ in range(num_slots)] + ) + if k_lora + else None + ) + self.mlp_loras = ( + nn.ModuleList( + [BatchedLinearLoRA(bsz, dim, dim, rank) for _ in range(num_slots)] + ) + if mlp_lora + else None + ) + self.o_loras = ( + nn.ModuleList( + [BatchedLinearLoRA(bsz, dim, dim, rank) for _ in range(num_slots)] + ) + if o_lora + else None + ) + + def reset(self): + with torch.no_grad(): + self.lm_head_lora.reset() + for loras in [self.q_loras, self.v_loras, self.k_loras, + self.mlp_loras, self.o_loras]: + if loras is not None: + for lora in loras: + lora.reset() + + +# Polar Express per-iteration minimax Newton-Schulz coefficients (PR #1344, +# used in PR #1787's 1.06378 record). Replaces the stock (3.4445, -4.775, +# 2.0315) tuple applied 5x with 5 per-step optimal tuples. Same compute cost; +# tighter polar factor. Set POLAR_EXPRESS_NS=0 to restore stock Muon. +_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), +) +_STOCK_COEFFS = ((3.4445, -4.775, 2.0315),) * 5 +_POLAR_EXPRESS_NS = os.environ.get("POLAR_EXPRESS_NS", "1") == "1" + + +@torch.compile(fullgraph=not _IS_ROCM) +def zeropower_via_newtonschulz5(G, steps=10, eps=1e-07): + 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 + X = X / (X.norm(dim=(-2, -1), keepdim=True) + eps) + coeffs_src = _PE_COEFFS if _POLAR_EXPRESS_NS else _STOCK_COEFFS + coeffs = coeffs_src[:steps] if steps <= len(coeffs_src) else coeffs_src + for a, b, c in coeffs: + 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 + + +class Muon(torch.optim.Optimizer): + def __init__( + self, + params, + lr, + momentum, + backend_steps, + nesterov=True, + weight_decay=0.0, + row_normalize=False, + ): + super().__init__( + params, + dict( + lr=lr, + momentum=momentum, + backend_steps=backend_steps, + nesterov=nesterov, + weight_decay=weight_decay, + row_normalize=row_normalize, + ), + ) + 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, + }) + self._bank_meta.sort(key=lambda m: -m["p"].numel()) + self._built = True + + def launch_reduce_scatters(self): + 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): + 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) + row_normalize = group.get("row_normalize", False) + prev_ag_handle = None + prev_m = None + sharded = self._distributed and hasattr(self, "_rs_futures") + for idx, 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[idx] is not None: + self._rs_futures[idx].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 + if row_normalize: + rn = update.float().norm(dim=-1, keepdim=True).clamp_min(1e-07) + update = update / rn.to(update.dtype) + update = zeropower_via_newtonschulz5(update, steps=backend_steps) + 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 + + +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,skip_gates,parallel_post_lambdas,parallel_resid_lambdas,attn_gate_proj,attn_gate_w", + ).split(",") + if pattern +) + + +PACKED_REPLICATED_GRAD_MAX_NUMEL = 1 << 15 + + +class Optimizers: + def __init__(self, h, base_model): + 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) + if base_model.skip_gates is not None and base_model.skip_gates.numel() > 0: + scalar_params.append(base_model.skip_gates) + if base_model.parallel_post_lambdas is not None: + scalar_params.append(base_model.parallel_post_lambdas) + if base_model.parallel_resid_lambdas is not None: + scalar_params.append(base_model.parallel_resid_lambdas) + if getattr(base_model, "smear_gate_enabled", False): + scalar_params.append(base_model.smear_gate.weight) + scalar_params.append(base_model.smear_lambda) + token_lr = h.tied_embed_lr if h.tie_embeddings else h.embed_lr + tok_params = [ + {"params": [base_model.tok_emb.weight], "lr": token_lr, "base_lr": token_lr} + ] + self.optimizer_tok = torch.optim.AdamW( + tok_params, + betas=(h.beta1, h.beta2), + eps=h.adam_eps, + weight_decay=h.embed_wd, + fused=True, + ) + self.optimizer_muon = Muon( + matrix_params, + lr=h.matrix_lr, + momentum=h.muon_momentum, + backend_steps=h.muon_backend_steps, + weight_decay=h.muon_wd, + row_normalize=h.muon_row_normalize, + ) + for group in self.optimizer_muon.param_groups: + group["base_lr"] = h.matrix_lr + self.optimizer_scalar = torch.optim.AdamW( + [{"params": scalar_params, "lr": h.scalar_lr, "base_lr": h.scalar_lr}], + betas=(h.beta1, h.beta2), + eps=h.adam_eps, + weight_decay=h.adam_wd, + fused=True, + ) + self.optimizers = [ + self.optimizer_tok, + self.optimizer_muon, + self.optimizer_scalar, + ] + self.replicated_params = list(tok_params[0]["params"]) + self.replicated_params.extend(scalar_params) + self.replicated_large_params = [] + self.replicated_packed_params = [] + for p in self.replicated_params: + if p.numel() <= PACKED_REPLICATED_GRAD_MAX_NUMEL: + self.replicated_packed_params.append(p) + else: + self.replicated_large_params.append(p) + + def __iter__(self): + return iter(self.optimizers) + + def zero_grad_all(self): + for opt in self.optimizers: + opt.zero_grad(set_to_none=True) + + def _all_reduce_packed_grads(self): + grads_by_key = collections.defaultdict(list) + for p in self.replicated_packed_params: + if p.grad is not None: + grads_by_key[(p.grad.device, p.grad.dtype)].append(p.grad) + for grads in grads_by_key.values(): + flat = torch.empty( + sum(g.numel() for g in grads), + device=grads[0].device, + dtype=grads[0].dtype, + ) + offset = 0 + for g in grads: + n = g.numel() + flat[offset : offset + n].copy_(g.contiguous().view(-1)) + offset += n + dist.all_reduce(flat, op=dist.ReduceOp.AVG) + offset = 0 + for g in grads: + n = g.numel() + g.copy_(flat[offset : offset + n].view_as(g)) + offset += n + + def step(self, distributed=False): + self.optimizer_muon.launch_reduce_scatters() + if distributed: + reduce_handles = [ + dist.all_reduce(p.grad, op=dist.ReduceOp.AVG, async_op=True) + for p in self.replicated_large_params + if p.grad is not None + ] + self._all_reduce_packed_grads() + for handle in reduce_handles: + handle.wait() + self.optimizer_tok.step() + self.optimizer_scalar.step() + self.optimizer_muon.step() + self.zero_grad_all() + + +def restore_fp32_params(model): + for module in model.modules(): + if isinstance(module, CastedLinear): + module.float() + for name, param in model.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() + if hasattr(model, "qo_bank") and model.qo_bank is not None: + model.qo_bank.data = model.qo_bank.data.float() + model.kv_bank.data = model.kv_bank.data.float() + model.mlp_up_bank.data = model.mlp_up_bank.data.float() + model.mlp_down_bank.data = model.mlp_down_bank.data.float() + + +def collect_hessians(model, train_loader, h, device, n_calibration_batches=64): + hessians = {} + hooks = [] + for i, block in enumerate(model.blocks): + block.attn._calib = True + block.mlp._calib = True + block.mlp.use_fused = False + + def make_attn_hook(layer_idx): + def hook_fn(module, inp, out): + x = inp[0].detach().float() + if x.ndim == 3: + x = x.reshape(-1, x.shape[-1]) + for suffix in ["c_q", "c_k", "c_v"]: + name = f"blocks.{layer_idx}.attn.{suffix}.weight" + if name not in hessians: + hessians[name] = torch.zeros( + x.shape[1], x.shape[1], dtype=torch.float32, device=device + ) + hessians[name].addmm_(x.T, x) + y = module._last_proj_input + if y is not None: + y = y.float() + if y.ndim == 3: + y = y.reshape(-1, y.shape[-1]) + name = f"blocks.{layer_idx}.attn.proj.weight" + if name not in hessians: + hessians[name] = torch.zeros( + y.shape[1], y.shape[1], dtype=torch.float32, device=device + ) + hessians[name].addmm_(y.T, y) + return hook_fn + + def make_mlp_hook(layer_idx): + def hook_fn(module, inp, out): + x = inp[0].detach().float() + if x.ndim == 3: + x = x.reshape(-1, x.shape[-1]) + name = f"blocks.{layer_idx}.mlp.fc.weight" + if name not in hessians: + hessians[name] = torch.zeros( + x.shape[1], x.shape[1], dtype=torch.float32, device=device + ) + hessians[name].addmm_(x.T, x) + h_act = module._last_down_input + if h_act is not None: + h_act = h_act.float() + if h_act.ndim == 3: + h_act = h_act.reshape(-1, h_act.shape[-1]) + name = f"blocks.{layer_idx}.mlp.proj.weight" + if name not in hessians: + hessians[name] = torch.zeros( + h_act.shape[1], h_act.shape[1], dtype=torch.float32, device=device + ) + hessians[name].addmm_(h_act.T, h_act) + return hook_fn + + for i, block in enumerate(model.blocks): + hooks.append(block.attn.register_forward_hook(make_attn_hook(i))) + hooks.append(block.mlp.register_forward_hook(make_mlp_hook(i))) + + # Hessian hooks for embedding factorization projection layers + def make_linear_input_hook(weight_name): + def hook_fn(module, inp, out): + x = inp[0].detach().float() + if x.ndim == 3: + x = x.reshape(-1, x.shape[-1]) + if weight_name not in hessians: + hessians[weight_name] = torch.zeros( + x.shape[1], x.shape[1], dtype=torch.float32, device=device + ) + hessians[weight_name].addmm_(x.T, x) + return hook_fn + + if model.tie_embeddings: + hook_module = model.final_norm + + def make_output_hook(name): + def hook_fn(module, inp, out): + x = out.detach().float() + if x.ndim == 3: + x = x.reshape(-1, x.shape[-1]) + if name not in hessians: + hessians[name] = torch.zeros( + x.shape[1], x.shape[1], dtype=torch.float32, device=device + ) + hessians[name].addmm_(x.T, x) + return hook_fn + + hooks.append( + hook_module.register_forward_hook(make_output_hook("tok_emb.weight")) + ) + model.eval() + with torch.no_grad(): + for _ in range(n_calibration_batches): + x, _ = train_loader.next_batch(h.train_batch_tokens, h.grad_accum_steps) + model.forward_logits(x) + for hook in hooks: + hook.remove() + for i, block in enumerate(model.blocks): + block.attn._calib = False + block.mlp._calib = False + block.mlp.use_fused = True + for name in hessians: + hessians[name] = hessians[name].cpu() / n_calibration_batches + return hessians + + +def gptq_quantize_weight(w, H, clip_sigmas=3.0, clip_range=63, block_size=128): + W_orig = w.float().clone() + rows, cols = W_orig.shape + H = H.float().clone() + dead = torch.diag(H) == 0 + H[dead, dead] = 1 + damp = 0.01 * H.diag().mean() + H.diagonal().add_(damp) + perm = torch.argsort(H.diag(), descending=True) + invperm = torch.argsort(perm) + W_perm = W_orig[:, perm].clone() + W_perm[:, dead[perm]] = 0 + H = H[perm][:, perm] + Hinv = torch.cholesky_inverse(torch.linalg.cholesky(H)) + Hinv = torch.linalg.cholesky(Hinv, upper=True) + row_std = W_orig.std(dim=1) + s = (clip_sigmas * row_std / clip_range).clamp_min(1e-10).to(torch.float16) + sf = s.float() + Q = torch.zeros(rows, cols, dtype=torch.int8) + W_work = W_perm.clone() + for i1 in range(0, cols, block_size): + i2 = min(i1 + block_size, cols) + W_block = W_work[:, i1:i2].clone() + Hinv_block = Hinv[i1:i2, i1:i2] + Err = torch.zeros(rows, i2 - i1) + for j in range(i2 - i1): + w_col = W_block[:, j] + d = Hinv_block[j, j] + q_col = torch.clamp(torch.round(w_col / sf), -clip_range, clip_range) + Q[:, i1 + j] = q_col.to(torch.int8) + err = (w_col - q_col.float() * sf) / d + Err[:, j] = err + W_block[:, j:] -= err.unsqueeze(1) * Hinv_block[j, j:].unsqueeze(0) + if i2 < cols: + W_work[:, i2:] -= Err @ Hinv[i1:i2, i2:] + return Q[:, invperm], s + + +def _lqer_pack(A, B, bits): + rng = 2 ** (bits - 1) - 1 + sA = (A.abs().amax(dim=1).clamp_min(1e-10) / rng).to(torch.float16) + sB = (B.abs().amax(dim=1).clamp_min(1e-10) / rng).to(torch.float16) + qA = torch.clamp(torch.round(A / sA.float().view(-1, 1)), -rng, rng).to(torch.int8) + qB = torch.clamp(torch.round(B / sB.float().view(-1, 1)), -rng, rng).to(torch.int8) + return qA, sA, qB, sB + + +def _lqer_pack_asym(A, B, g=64): + # A: INT2 per-matrix scalar (signed [-2, 1], scale = |A|max/1.5). + sA = (A.abs().amax().clamp_min(1e-10) / 1.5).to(torch.float16) + qA = torch.clamp(torch.round(A / sA.float()), -2, 1).to(torch.int8) + # B: INT4 groupwise over flattened B (signed [-8, 7], per-group scale). + Bf = B.reshape(-1, g) + Bmax = Bf.abs().amax(dim=-1, keepdim=True).clamp_min(1e-10) + sB = (Bmax / 7.5).to(torch.float16).reshape(-1) + qB = torch.clamp(torch.round(Bf / sB.float().reshape(-1, 1)), -8, 7).to( + torch.int8 + ).reshape(B.shape) + return qA, sA, qB, sB + + +def gptq_mixed_quantize(state_dict, hessians, h): + result = {} + meta = {} + lqer_on = bool(getattr(h, "lqer_enabled", False)) + lqer_cands = {} + orig_by_name = {} + for (name, tensor) in state_dict.items(): + t = tensor.detach().cpu().contiguous() + 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 (float16)" + continue + if "tok_emb" in name: + cs = h.embed_clip_sigmas + elif ".mlp." in name: + cs = h.mlp_clip_sigmas + elif ".attn." in name: + cs = h.attn_clip_sigmas + else: + cs = h.matrix_clip_sigmas + bits = h.embed_bits if "tok_emb" in name else h.matrix_bits + clip_range = 2 ** (bits - 1) - 1 + ret = gptq_quantize_weight( + t, hessians[name], clip_sigmas=cs, clip_range=clip_range + ) + q, s = ret + result[name + ".q"] = q + result[name + ".scale"] = s + meta[name] = f"gptq (int{bits})" + if lqer_on: + W_q = q.float() * s.float().view(-1, 1) + E = t.float() - W_q + lqer_cands[name] = (E, float(E.norm())) + orig_by_name[name] = t + if lqer_on and lqer_cands: + asym_on = bool(getattr(h, "lqer_asym_enabled", False)) + asym_g = int(getattr(h, "lqer_asym_group", 64)) + adaptive = bool(getattr(h, "lqer_adaptive", False)) + + def _asym_bytes(rows, cols, r, group): + # A: INT2, scalar fp16 → rows*r*2/8 + 2 + # B: INT4 groupwise → r*cols*4/8 + (r*cols/group)*2 (fp16 scales) + a_bytes = rows * r * 2 // 8 + 2 + b_bytes = r * cols * 4 // 8 + n_groups = (r * cols) // group if group > 0 else 0 + b_bytes += n_groups * 2 + return a_bytes + b_bytes + + if adaptive: + # Split on any non-digit so 4,6,8,10 / 4_6_8_10 / 4-6-8-10 all work. + cand_ranks = sorted( + {int(x) for x in re.findall(r"\d+", str(h.lqer_cand_ranks))} + ) + if not cand_ranks: + cand_ranks = [4, 6, 8, 10] + budget = int(h.lqer_budget_bytes) + # SVD all candidates once. + svd_cache = {} + options = [] + for (name, (E, frob)) in lqer_cands.items(): + U, S, Vh = torch.linalg.svd(E, full_matrices=False) + svd_cache[name] = (U, S, Vh) + rows, cols = E.shape + cum_energy = 0.0 + prev_energy = 0.0 + for r in cand_ranks: + rr = min(r, S.numel()) + energy = float((S[:rr] ** 2).sum().item()) + cost = _asym_bytes(rows, cols, rr, asym_g) + # Marginal gain over previous rank choice for this tensor: + delta_energy = energy - prev_energy + if r > cand_ranks[0]: + prev_cost = _asym_bytes(rows, cols, min(cand_ranks[cand_ranks.index(r) - 1], S.numel()), asym_g) + delta_cost = max(cost - prev_cost, 1) + else: + delta_cost = max(cost, 1) + options.append( + (delta_energy / delta_cost, name, rr, energy, cost) + ) + prev_energy = energy + # Greedy selection: best gain-per-byte, only one rank per tensor (highest r selected). + options.sort(key=lambda x: -x[0]) + best_for = {} # name -> (rank, total_cost) for the highest-rank pick we can afford + remaining = budget + # Two-pass: first include base rank picks, then upgrade ranks if budget allows. + for (gpb, name, r, energy, cost) in options: + if name in best_for: + cur_r, cur_cost = best_for[name] + if r <= cur_r: + continue + extra = cost - cur_cost + if extra > remaining: + continue + best_for[name] = (r, cost) + remaining -= extra + else: + if cost > remaining: + continue + best_for[name] = (r, cost) + remaining -= cost + log( + f"LQER adaptive: budget={budget} bytes, used={budget - remaining} bytes, tensors={len(best_for)}" + ) + for (name, (r, cost)) in sorted(best_for.items()): + log(f" {name}: rank={r} bytes={cost}") + top = [ + (name, (lqer_cands[name][0], lqer_cands[name][1])) + for name in best_for + ] + rank_for = {name: r for (name, (r, _)) in best_for.items()} + else: + top = sorted(lqer_cands.items(), key=lambda kv: -kv[1][1])[: h.lqer_top_k] + rank_for = {name: h.lqer_rank for (name, _) in top} + svd_cache = {} + + for (name, (E, _)) in top: + if name in svd_cache: + U, S, Vh = svd_cache[name] + else: + U, S, Vh = torch.linalg.svd(E, full_matrices=False) + r = min(rank_for[name], S.numel()) + A = (U[:, :r] * S[:r]).contiguous() + B = Vh[:r, :].contiguous() + if asym_on and B.numel() % asym_g == 0: + qA, sA, qB, sB = _lqer_pack_asym(A, B, asym_g) + result[name + ".lqA_a"] = qA + result[name + ".lqAs_a"] = sA + result[name + ".lqB_a"] = qB + result[name + ".lqBs_a"] = sB + meta[name] = meta[name] + f"+lqer_asym_r{r}" + else: + qA, sA, qB, sB = _lqer_pack(A, B, h.lqer_factor_bits) + result[name + ".lqA"] = qA + result[name + ".lqAs"] = sA + result[name + ".lqB"] = qB + result[name + ".lqBs"] = sB + meta[name] = meta[name] + f"+lqer_r{r}" + categories = collections.defaultdict(set) + for (name, cat) in meta.items(): + short = re.sub("\\.\\d+$", "", re.sub("blocks\\.\\d+", "blocks", name)) + categories[cat].add(short) + log("Quantized weights:") + for cat in sorted(categories): + log(f" {cat}: {', '.join(sorted(categories[cat]))}") + return result, meta + + +def dequantize_mixed(result, meta, template_sd): + out = {} + for (name, orig) in template_sd.items(): + info = meta.get(name) + if info is None: + continue + orig_dtype = orig.dtype + if "passthrough" in info: + 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: + W = q.float() * s.float().view(q.shape[0], *[1] * (q.ndim - 1)) + else: + W = q.float() * float(s.item()) + if "lqer_asym" in info: + qA_t = result[name + ".lqA_a"] + sA_t = result[name + ".lqAs_a"] + qB_t = result[name + ".lqB_a"] + sB_t = result[name + ".lqBs_a"] + qA = qA_t.float() * float(sA_t) + g_sz = qB_t.numel() // sB_t.numel() + qB = ( + qB_t.reshape(-1, g_sz).float() * sB_t.float().view(-1, 1) + ).reshape(qB_t.shape) + W = W + qA @ qB + elif "lqer" in info: + qA = result[name + ".lqA"].float() * result[name + ".lqAs"].float().view(-1, 1) + qB = result[name + ".lqB"].float() * result[name + ".lqBs"].float().view(-1, 1) + W = W + qA @ qB + out[name] = W.to(orig_dtype) + return out + + +_BSHF_MAGIC = b"BSHF" + + +def _byte_shuffle(data, stride=2): + if stride <= 1 or len(data) < stride: + return data + src = np.frombuffer(data, dtype=np.uint8) + n = len(src) + out = np.empty(n, dtype=np.uint8) + dest_off = 0 + for pos in range(stride): + chunk = src[pos::stride] + out[dest_off : dest_off + len(chunk)] = chunk + dest_off += len(chunk) + return _BSHF_MAGIC + bytes([stride]) + out.tobytes() + + +def _byte_unshuffle(data): + if len(data) < 5 or data[:4] != _BSHF_MAGIC: + return data + stride = data[4] + if stride < 2: + return data[5:] + payload = np.frombuffer(data, dtype=np.uint8, offset=5) + n = len(payload) + out = np.empty(n, dtype=np.uint8) + src_off = 0 + for pos in range(stride): + chunk_len = n // stride + (1 if pos < n % stride else 0) + out[pos::stride][:chunk_len] = payload[src_off : src_off + chunk_len] + src_off += chunk_len + return out.tobytes() + + +def _compress(data, compressor): + data = _byte_shuffle(data) + if compressor == "lzma": + return lzma.compress(data, preset=6) + elif compressor == "brotli": + import brotli + + return brotli.compress(data, quality=11) + raise ValueError(f"Unknown compressor: {compressor!r}") + + +def _decompress(data, compressor): + if compressor == "lzma": + raw = lzma.decompress(data) + elif compressor == "brotli": + import brotli + + raw = brotli.decompress(data) + else: + raise ValueError(f"Unknown compressor: {compressor!r}") + raw = _byte_unshuffle(raw) + return raw + + +def _unbank_state_dict(state_dict, num_layers): + sd = {} + n = num_layers + for k, v in state_dict.items(): + t = v.detach().cpu() if v is not None else None + if k == "qo_bank": + for i in range(n): + sd[f"blocks.{i}.attn.c_q.weight"] = t[i] + sd[f"blocks.{i}.attn.proj.weight"] = t[n + i] + elif k == "kv_bank": + for i in range(n): + sd[f"blocks.{i}.attn.c_k.weight"] = t[i] + sd[f"blocks.{i}.attn.c_v.weight"] = t[n + i] + elif k == "mlp_up_bank": + for i in range(n): + sd[f"blocks.{i}.mlp.fc.weight"] = t[i] + elif k == "mlp_down_bank": + for i in range(n): + sd[f"blocks.{i}.mlp.proj.weight"] = t[i] + else: + if t is not None: + sd[k] = t + return sd + + +def _rebank_state_dict(flat_sd, num_layers, model_dim, kv_dim, hidden_dim): + sd = {} + n = num_layers + sd["qo_bank"] = torch.zeros(2 * n, model_dim, model_dim) + sd["kv_bank"] = torch.zeros(2 * n, kv_dim, model_dim) + for i in range(n): + sd["qo_bank"][i] = flat_sd[f"blocks.{i}.attn.c_q.weight"] + sd["qo_bank"][n + i] = flat_sd[f"blocks.{i}.attn.proj.weight"] + sd["kv_bank"][i] = flat_sd[f"blocks.{i}.attn.c_k.weight"] + sd["kv_bank"][n + i] = flat_sd[f"blocks.{i}.attn.c_v.weight"] + sd["mlp_up_bank"] = torch.zeros(n, hidden_dim, model_dim) + sd["mlp_down_bank"] = torch.zeros(n, model_dim, hidden_dim) + for i in range(n): + sd["mlp_up_bank"][i] = flat_sd[f"blocks.{i}.mlp.fc.weight"] + sd["mlp_down_bank"][i] = flat_sd[f"blocks.{i}.mlp.proj.weight"] + for k, v in flat_sd.items(): + if not ( + k.startswith("blocks.") + and any( + p in k + for p in [ + ".attn.c_q.", ".attn.c_k.", ".attn.c_v.", + ".attn.proj.", ".mlp.fc.", ".mlp.proj.", + ] + ) + ): + sd[k] = v + return sd + + + +def _compressed_code_size(code): + code_raw = code.encode("utf-8") + try: + result = subprocess.run( + ["pyminify", "--no-rename-locals", "--no-hoist-literals", "--remove-literal-statements", "-"], + input=code_raw, capture_output=True, check=True, + ) + minified = result.stdout + except (FileNotFoundError, subprocess.CalledProcessError): + minified = code_raw + compressed = lzma.compress(minified) + encoded = base64.b85encode(compressed) + wrapper = b'import lzma as L,base64 as B\nexec(L.decompress(B.b85decode("' + encoded + b'")))\n' + return len(code_raw), len(wrapper) + + +def serialize(h, base_model, code): + code_bytes_uncompressed, code_bytes = _compressed_code_size(code) + if h.is_main_process: + torch.save(base_model.state_dict(), h.model_path) + model_bytes = os.path.getsize(h.model_path) + log(f"Serialized model: {model_bytes} bytes") + log(f"Code size (uncompressed): {code_bytes_uncompressed} bytes") + log(f"Code size (compressed): {code_bytes} bytes") + sd_cpu = _unbank_state_dict(base_model.state_dict(), h.num_layers) + device = torch.device("cuda", h.local_rank) + log("GPTQ:collecting Hessians from calibration data...") + t0 = time.perf_counter() + calib_loader = ShuffledSequenceLoader(h, device) + hessians = collect_hessians( + base_model, + calib_loader, + h, + device, + n_calibration_batches=h.gptq_calibration_batches, + ) + log(f"GPTQ:collected {len(hessians)} Hessians in {time.perf_counter()-t0:.1f}s") + quant_result, quant_meta = gptq_mixed_quantize(sd_cpu, hessians, h) + quant_buf = io.BytesIO() + torch.save({"w": quant_result, "m": quant_meta}, quant_buf) + quant_raw = quant_buf.getvalue() + quant_blob = _compress(quant_raw, h.compressor) + quant_file_bytes = len(quant_blob) + bytes_total = quant_file_bytes + code_bytes + if h.is_main_process: + with open(h.quantized_model_path, "wb") as f: + f.write(quant_blob) + log(f"Serialized model quantized+{h.compressor}: {quant_file_bytes} bytes") + log(f"Total submission size quantized+{h.compressor}: {bytes_total} bytes") + return bytes_total, quant_file_bytes + + +def deserialize(h, device): + eval_model = GPT(h).to(device).bfloat16() + restore_fp32_params(eval_model) + flat_template = _unbank_state_dict(eval_model.state_dict(), h.num_layers) + with open(h.quantized_model_path, "rb") as f: + quant_blob_disk = f.read() + quant_state = torch.load( + io.BytesIO(_decompress(quant_blob_disk, h.compressor)), map_location="cpu" + ) + deq_flat = dequantize_mixed(quant_state["w"], quant_state["m"], flat_template) + head_dim = h.model_dim // h.num_heads + kv_dim = h.num_kv_heads * head_dim + hidden_dim = int(h.mlp_mult * h.model_dim) + deq_state = _rebank_state_dict(deq_flat, h.num_layers, h.model_dim, kv_dim, hidden_dim) + eval_model.load_state_dict(deq_state, strict=True) + return eval_model + + +def _loss_bpb(loss_sum, token_count, byte_count): + val_loss = (loss_sum / token_count).item() + val_bpb = val_loss / math.log(2.0) * (token_count.item() / byte_count.item()) + return val_loss, val_bpb + + +def eval_val(h, device, val_data, model, forward_logits_fn=None): + seq_len = h.eval_seq_len + local_batch_tokens = h.val_batch_tokens // (h.world_size * h.grad_accum_steps) + if local_batch_tokens < seq_len: + raise ValueError( + f"VAL_BATCH_SIZE must provide at least one sequence per rank; got VAL_BATCH_SIZE={h.val_batch_tokens}, WORLD_SIZE={h.world_size}, GRAD_ACCUM_STEPS={h.grad_accum_steps}, seq_len={seq_len}" + ) + local_batch_seqs = local_batch_tokens // seq_len + total_seqs = (val_data.val_tokens.numel() - 1) // seq_len + seq_start = total_seqs * h.rank // h.world_size + seq_end = total_seqs * (h.rank + 1) // h.world_size + + # TODO: Don't truncate this. + seq_end = seq_start + ((seq_end - seq_start) // local_batch_seqs) * local_batch_seqs + + 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) + run_forward_logits = ( + (model.module.forward_logits if hasattr(model, "module") else model.forward_logits) + if forward_logits_fn is None + else forward_logits_fn + ) + model.eval() + global BOS_ID + if BOS_ID is None: + BOS_ID = 1 + with torch.no_grad(): + 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_data.val_tokens[raw_start:raw_end].to( + device=device, dtype=torch.int64, non_blocking=True + ) + x = local[:-1] + y = local[1:] + bos_pos = (x == BOS_ID).nonzero(as_tuple=True)[0].tolist() + cu_seqlens, max_seqlen = _build_cu_seqlens( + bos_pos, x.numel(), x.device, h.eval_seq_len, 64 + ) + with torch.autocast(device_type="cuda", dtype=torch.bfloat16, enabled=True): + logits = run_forward_logits( + x[None], cu_seqlens=cu_seqlens, max_seqlen=max_seqlen + ).detach() + per_token_loss = F.cross_entropy( + logits.reshape(-1, logits.size(-1)).float(), + y.reshape(-1), + reduction="none", + ) + val_loss_sum += per_token_loss.to(torch.float64).sum() + val_token_count += float(y.numel()) + prev_ids = x + tgt_ids = y + token_bytes = val_data.base_bytes_lut[tgt_ids].to(dtype=torch.int16) + token_bytes += ( + val_data.has_leading_space_lut[tgt_ids] + & ~val_data.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) + model.train() + return _loss_bpb(val_loss_sum, val_token_count, val_byte_count) + + +def eval_val_sliding(h, device, val_data, base_model, forward_logits_fn=None, batch_seqs=32): + global BOS_ID + if BOS_ID is None: + BOS_ID = 1 + base_model.eval() + run_forward_logits = base_model.forward_logits if forward_logits_fn is None else forward_logits_fn + seq_len = h.eval_seq_len + stride = h.eval_stride + total_tokens = val_data.val_tokens.numel() - 1 + context_size = seq_len - stride + window_starts = [ws for ws in range(0, total_tokens, stride) + if ws + context_size < total_tokens] + total_windows = len(window_starts) + my_s = (total_windows * h.rank) // h.world_size + my_e = (total_windows * (h.rank + 1)) // h.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) + total_batches = (len(my_windows) + batch_seqs - 1) // batch_seqs + is_master = h.rank == 0 + cu_bucket = 64 + t_sw_start = time.perf_counter() + with torch.no_grad(): + for bi in range(0, len(my_windows), batch_seqs): + batch_idx = bi // batch_seqs + if is_master and (batch_idx % 50 == 0 or batch_idx == total_batches - 1): + elapsed = time.perf_counter() - t_sw_start + rl = float(loss_sum.item() / token_count.item()) if token_count.item() > 0 else 0.0 + rb = float((rl / math.log(2.0)) * token_count.item() / byte_count.item()) if byte_count.item() > 0 else 0.0 + log(f"sliding_progress: batch {batch_idx+1}/{total_batches} " + f"tokens:{int(token_count.item())} running_loss:{rl:.4f} running_bpb:{rb:.4f} " + f"elapsed:{elapsed:.1f}s") + batch_ws = my_windows[bi:bi + batch_seqs] + x_parts = [] + y_parts = [] + cu_starts = [] + score_ranges = [] + offset = 0 + for ws in batch_ws: + end = min(ws + seq_len, total_tokens) + wlen = end - ws + chunk_cpu = val_data.val_tokens[ws:end + 1] + bos_pos = (chunk_cpu[:-1] == BOS_ID).nonzero(as_tuple=True)[0].tolist() + if not bos_pos or bos_pos[0] != 0: + bos_pos = [0] + bos_pos + cu_starts.extend(offset + pos for pos in bos_pos) + chunk = chunk_cpu.to(dtype=torch.int64, device=device) + x_parts.append(chunk[:-1]) + y_parts.append(chunk[1:]) + score_ranges.append((offset, wlen, ws)) + offset += wlen + x_cat = torch.cat(x_parts, dim=0)[None] + y_cat = torch.cat(y_parts, dim=0) + boundaries = cu_starts + [offset] + padded_len = get_next_multiple_of_n(len(boundaries), cu_bucket) + cu_seqlens = torch.full((padded_len,), offset, dtype=torch.int32, device=device) + cu_seqlens[:len(boundaries)] = torch.tensor(boundaries, dtype=torch.int32, device=device) + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + logits = run_forward_logits(x_cat, cu_seqlens=cu_seqlens, max_seqlen=seq_len) + flat_nll = F.cross_entropy( + logits.reshape(-1, logits.size(-1)).float(), + y_cat, + reduction="none", + ) + flat_x = x_cat.reshape(-1) + for off, wlen, ws in score_ranges: + s = 0 if ws == 0 else context_size + lo = off + s + hi = off + wlen + scored_nll = flat_nll[lo:hi].to(torch.float64) + loss_sum += scored_nll.sum() + token_count += float(hi - lo) + tgt = y_cat[lo:hi] + prev = flat_x[lo:hi] + tb = val_data.base_bytes_lut[tgt].to(torch.float64) + tb += (val_data.has_leading_space_lut[tgt] & ~val_data.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) + base_model.train() + return _loss_bpb(loss_sum, token_count, byte_count) + + +def _find_docs(all_tokens): + bos_positions = (all_tokens == BOS_ID).nonzero(as_tuple=True)[0].numpy() + docs = [] + for i in range(len(bos_positions)): + start = int(bos_positions[i]) + end = ( + int(bos_positions[i + 1]) + if i + 1 < len(bos_positions) + else all_tokens.numel() + ) + if i + 1 < len(bos_positions): + end += 1 + assert end - start >= 2 + docs.append((start, end - start)) + return docs + + +def _build_ttt_global_batches(doc_entries, h, ascending=False): + batch_size = h.ttt_batch_size + global_doc_entries = sorted(doc_entries, key=lambda x: x[1][1]) + global_batches = [ + global_doc_entries[i : i + batch_size] + for i in range(0, len(global_doc_entries), batch_size) + ] + indexed = list(enumerate(global_batches)) + if not ascending: + indexed.sort(key=lambda ib: -max(dl for _, (_, dl) in ib[1])) + return indexed + + +def _init_batch_counter(path): + with open(path, "wb") as f: + f.write((0).to_bytes(4, "little")) + + +def _claim_next_batch(counter_path, queue_len): + try: + with open(counter_path, "r+b") as f: + fcntl.flock(f, fcntl.LOCK_EX) + idx = int.from_bytes(f.read(4), "little") + f.seek(0) + f.write((idx + 1).to_bytes(4, "little")) + f.flush() + except FileNotFoundError: + return queue_len + return idx + + +def _compute_chunk_window(ci, pred_len, num_chunks, chunk_size, eval_seq_len): + chunk_end = pred_len if ci == num_chunks - 1 else (ci + 1) * chunk_size + win_start = max(0, chunk_end - eval_seq_len) + win_len = chunk_end - win_start + chunk_start = ci * chunk_size + chunk_offset = chunk_start - win_start + chunk_len = chunk_end - chunk_start + return win_start, win_len, chunk_offset, chunk_len + + +def _accumulate_bpb( + ptl, + x, + y, + chunk_offsets, + chunk_lens, + pos_idx, + base_bytes_lut, + has_leading_space_lut, + is_boundary_token_lut, + loss_sum, + byte_sum, + token_count, +): + pos = pos_idx[: x.size(1)].unsqueeze(0) + mask = ( + (chunk_lens.unsqueeze(1) > 0) + & (pos >= chunk_offsets.unsqueeze(1)) + & (pos < (chunk_offsets + chunk_lens).unsqueeze(1)) + ) + mask_f64 = mask.to(torch.float64) + tok_bytes = base_bytes_lut[y].to(torch.float64) + tok_bytes += (has_leading_space_lut[y] & ~is_boundary_token_lut[x]).to( + torch.float64 + ) + loss_sum += (ptl.to(torch.float64) * mask_f64).sum() + byte_sum += (tok_bytes * mask_f64).sum() + token_count += chunk_lens.to(torch.float64).sum() + + +def _loss_bpb_from_sums(loss_sum, token_count, byte_sum): + val_loss = (loss_sum / token_count).item() + val_bpb = val_loss / math.log(2.0) * (token_count.item() / byte_sum.item()) + return val_loss, val_bpb + + +def _split_doc_entries_for_phased(doc_entries, prefix_docs): + prefix_docs = max(0, min(len(doc_entries), int(prefix_docs))) + return doc_entries[:prefix_docs], doc_entries[prefix_docs:] + + +def _add_to_counter(path, delta): + try: + with open(path, "r+b") as f: + fcntl.flock(f, fcntl.LOCK_EX) + cur = int.from_bytes(f.read(8), "little", signed=True) + cur += int(delta) + f.seek(0) + f.write(int(cur).to_bytes(8, "little", signed=True)) + f.flush() + return cur + except FileNotFoundError: + return int(delta) + + +def _init_int64_counter(path): + with open(path, "wb") as f: + f.write((0).to_bytes(8, "little", signed=True)) + + +def _select_ttt_doc_entries(docs, h): + doc_entries = list(enumerate(docs)) + if h.val_doc_fraction < 1.0: + sample_n = max(1, int(round(len(docs) * h.val_doc_fraction))) + sampled_indices = sorted( + random.Random(h.seed).sample(range(len(docs)), sample_n) + ) + return [(i, docs[i]) for i in sampled_indices] + return doc_entries + + +def train_val_ttt_global_sgd_distributed(h, device, val_data, base_model, val_tokens, batch_seqs=None): + global BOS_ID + if BOS_ID is None: + BOS_ID = 1 + base_model.eval() + seq_len = h.eval_seq_len + total_tokens = val_tokens.numel() - 1 + ttt_chunk = h.global_ttt_chunk_tokens + batch_seqs = h.global_ttt_batch_seqs if batch_seqs is None else batch_seqs + num_chunks = (total_tokens + ttt_chunk - 1) // ttt_chunk + ttt_params = [p for p in base_model.parameters()] + for p in ttt_params: + p.requires_grad_(True) + optimizer = torch.optim.SGD( + ttt_params, lr=h.global_ttt_lr, momentum=h.global_ttt_momentum + ) + t_start = time.perf_counter() + for ci in range(num_chunks): + chunk_start = ci * ttt_chunk + chunk_end = min((ci + 1) * ttt_chunk, total_tokens) + is_last_chunk = ci == num_chunks - 1 + if is_last_chunk or h.global_ttt_epochs <= 0: + continue + base_model.train() + chunk_seqs = (chunk_end - chunk_start) // seq_len + if chunk_seqs <= 0: + continue + warmup_chunks = max(0, min(h.global_ttt_warmup_chunks, num_chunks - 1)) + if warmup_chunks > 0 and ci < warmup_chunks: + warmup_denom = max(warmup_chunks - 1, 1) + warmup_t = ci / warmup_denom + lr_now = ( + h.global_ttt_warmup_start_lr + + (h.global_ttt_lr - h.global_ttt_warmup_start_lr) * warmup_t + ) + else: + decay_steps = max(num_chunks - 1 - warmup_chunks, 1) + decay_ci = max(ci - warmup_chunks, 0) + lr_now = h.global_ttt_lr * 0.5 * ( + 1.0 + math.cos(math.pi * decay_ci / decay_steps) + ) + for pg in optimizer.param_groups: + pg["lr"] = lr_now + my_seq_s = chunk_seqs * h.rank // h.world_size + my_seq_e = chunk_seqs * (h.rank + 1) // h.world_size + my_chunk_seqs = my_seq_e - my_seq_s + for _ in range(h.global_ttt_epochs): + for bs in range(0, my_chunk_seqs, batch_seqs): + be = min(bs + batch_seqs, my_chunk_seqs) + actual_bs = my_seq_s + bs + start_tok = chunk_start + actual_bs * seq_len + end_tok = chunk_start + (my_seq_s + be) * seq_len + 1 + if end_tok > val_tokens.numel(): + continue + local = val_tokens[start_tok:end_tok].to(device=device, dtype=torch.int64) + x_flat = local[:-1] + y_flat = local[1:] + optimizer.zero_grad(set_to_none=True) + with torch.enable_grad(): + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + if h.global_ttt_respect_doc_boundaries: + bos_pos = (x_flat == BOS_ID).nonzero(as_tuple=True)[0].tolist() + cu_seqlens, max_seqlen = _build_cu_seqlens( + bos_pos, x_flat.numel(), x_flat.device, h.eval_seq_len, 64 + ) + loss = base_model( + x_flat[None], + y_flat[None], + cu_seqlens=cu_seqlens, + max_seqlen=max_seqlen, + ) + else: + x = x_flat.reshape(-1, seq_len) + y = y_flat.reshape(-1, seq_len) + loss = base_model(x, y) + loss.backward() + if dist.is_available() and dist.is_initialized(): + for p in ttt_params: + if p.grad is not None: + dist.all_reduce(p.grad, op=dist.ReduceOp.SUM) + p.grad.mul_(1.0 / h.world_size) + if h.global_ttt_grad_clip > 0: + torch.nn.utils.clip_grad_norm_(ttt_params, h.global_ttt_grad_clip) + optimizer.step() + base_model.eval() + if h.rank == 0: + elapsed = time.perf_counter() - t_start + log( + f"tttg: c{ci+1}/{num_chunks} lr:{lr_now:.6f} t:{elapsed:.1f}s" + ) + for p in base_model.parameters(): + p.requires_grad_(True) + base_model.eval() + + +def eval_val_ttt_phased(h, base_model, device, val_data, forward_ttt_train): + global BOS_ID + if BOS_ID is None: + BOS_ID = 1 + base_model.eval() + for p in base_model.parameters(): + p.requires_grad_(False) + all_tokens = val_data.val_tokens + all_tokens_idx = all_tokens.to(torch.int32) + docs = _find_docs(all_tokens) + doc_entries = _select_ttt_doc_entries(docs, h) + prefix_doc_limit = max(0, min(len(doc_entries), int(h.phased_ttt_prefix_docs))) + num_phases = max(1, int(h.phased_ttt_num_phases)) + phase_boundaries = [] + for pi in range(num_phases): + boundary = prefix_doc_limit * (pi + 1) // num_phases + phase_boundaries.append(boundary) + current_phase = 0 + current_phase_boundary = phase_boundaries[0] + log( + "ttt_phased:" + f" total_docs:{len(doc_entries)} prefix_docs:{prefix_doc_limit} " + f"suffix_docs:{len(doc_entries) - prefix_doc_limit}" + f" num_phases:{num_phases} boundaries:{phase_boundaries}" + ) + chunk_size, eval_seq_len = h.ttt_chunk_size, h.ttt_eval_seq_len + eval_batch_set = None + if h.ttt_eval_batches: + eval_batch_set = set(int(x) for x in h.ttt_eval_batches.split(",") if x.strip()) + use_ascending = eval_batch_set is not None + global_batches_sorted = _build_ttt_global_batches( + doc_entries, h, ascending=use_ascending + ) + queue_len = len(global_batches_sorted) + counter_path = f"/tmp/ttt_counter_{h.run_id}" + prefix_counter_path = f"/tmp/ttt_prefix_counter_{h.run_id}" + pause_flag_path = f"/tmp/ttt_pause_flag_{h.run_id}" + if h.rank == 0: + _init_batch_counter(counter_path) + _init_int64_counter(prefix_counter_path) + try: + os.remove(pause_flag_path) + except FileNotFoundError: + pass + if dist.is_available() and dist.is_initialized(): + path_list = [counter_path, prefix_counter_path, pause_flag_path] + dist.broadcast_object_list(path_list, src=0) + counter_path, prefix_counter_path, pause_flag_path = path_list + dist.barrier() + loss_sum = torch.zeros((), device=device, dtype=torch.float64) + byte_sum = torch.zeros((), device=device, dtype=torch.float64) + token_count = torch.zeros((), device=device, dtype=torch.float64) + t_start = time.perf_counter() + reusable_lora = BatchedTTTLoRA( + h.ttt_batch_size, base_model, h.ttt_lora_rank, + k_lora=h.ttt_k_lora, mlp_lora=h.ttt_mlp_lora, o_lora=h.ttt_o_lora, + ).to(device) + + def _build_opt(lora): + if h.ttt_optimizer == "sgd": + return torch.optim.SGD( + lora.parameters(), lr=h.ttt_lora_lr, + momentum=h.ttt_beta1, weight_decay=h.ttt_weight_decay, + ) + return torch.optim.AdamW( + lora.parameters(), lr=h.ttt_lora_lr, + betas=(h.ttt_beta1, h.ttt_beta2), + eps=1e-10, weight_decay=h.ttt_weight_decay, fused=True, + ) + + reusable_opt = _build_opt(reusable_lora) + local_scored_docs = [] + global_ttt_done = prefix_doc_limit == 0 + try: + while True: + queue_idx = _claim_next_batch(counter_path, queue_len) + if queue_idx >= queue_len: + break + orig_batch_idx, batch_entries = global_batches_sorted[queue_idx] + batch = [doc for _, doc in batch_entries] + bsz = len(batch) + prev_loss = loss_sum.item() + prev_bytes = byte_sum.item() + prev_tokens = token_count.item() + if bsz == reusable_lora.bsz: + reusable_lora.reset() + for s in reusable_opt.state.values(): + for k, v in s.items(): + if isinstance(v, torch.Tensor): + v.zero_() + elif k == "step": + s[k] = 0 + cur_lora = reusable_lora + cur_opt = reusable_opt + else: + cur_lora = BatchedTTTLoRA( + bsz, base_model, h.ttt_lora_rank, + k_lora=h.ttt_k_lora, mlp_lora=h.ttt_mlp_lora, o_lora=h.ttt_o_lora, + ).to(device) + cur_opt = _build_opt(cur_lora) + pred_lens = [doc_len - 1 for _, doc_len in batch] + num_chunks = [(pl + chunk_size - 1) // chunk_size for pl in pred_lens] + max_nc = max(num_chunks) + num_chunks_t = torch.tensor(num_chunks, dtype=torch.int64, device=device) + for ci in range(max_nc): + active = [ci < nc for nc in num_chunks] + needs_train = any(ci < nc - 1 for nc in num_chunks) + tok_starts = torch.zeros(bsz, dtype=torch.int64) + tok_wls = torch.zeros(bsz, dtype=torch.int64) + chunk_offsets_cpu = torch.zeros(bsz, dtype=torch.int64) + chunk_lens_cpu = torch.zeros(bsz, dtype=torch.int64) + for b in range(bsz): + if not active[b]: + continue + doc_start, doc_len = batch[b] + win_start, win_len, chunk_offset, chunk_len = _compute_chunk_window( + ci, pred_lens[b], num_chunks[b], chunk_size, eval_seq_len + ) + tok_starts[b] = doc_start + win_start + tok_wls[b] = win_len + chunk_offsets_cpu[b] = chunk_offset + chunk_lens_cpu[b] = chunk_len + _, context_size, chunk_offset, _ = _compute_chunk_window( + ci, (ci + 1) * chunk_size, ci + 1, chunk_size, eval_seq_len + ) + col_idx = torch.arange(context_size + 1) + idx = tok_starts.unsqueeze(1) + col_idx.unsqueeze(0) + idx.clamp_(max=all_tokens.numel() - 1) + gathered_gpu = all_tokens_idx[idx].to( + device=device, dtype=torch.int64, non_blocking=True + ) + valid = (col_idx[:context_size].unsqueeze(0) < tok_wls.unsqueeze(1)).to( + device, non_blocking=True + ) + chunk_offsets = chunk_offsets_cpu.to(device, non_blocking=True) + chunk_lens = chunk_lens_cpu.to(device, non_blocking=True) + x = torch.where(valid, gathered_gpu[:, :context_size], 0) + y = torch.where(valid, gathered_gpu[:, 1 : context_size + 1], 0) + ctx_pos = torch.arange(context_size, device=device, dtype=torch.int64) + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + per_tok_loss = forward_ttt_train(x, y, lora=cur_lora) + with torch.no_grad(): + _accumulate_bpb( + per_tok_loss, + x, + y, + chunk_offsets, + chunk_lens, + ctx_pos, + val_data.base_bytes_lut, + val_data.has_leading_space_lut, + val_data.is_boundary_token_lut, + loss_sum, + byte_sum, + token_count, + ) + if needs_train: + activate_chunk_mask = (num_chunks_t - 1 > ci).float() + for gi in range(h.ttt_grad_steps): + if gi > 0: + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + per_tok_loss = forward_ttt_train(x, y, lora=cur_lora) + per_doc = per_tok_loss[ + :, chunk_offset : chunk_offset + chunk_size + ].mean(dim=-1) + cur_opt.zero_grad(set_to_none=True) + (per_doc * activate_chunk_mask).sum().backward() + cur_opt.step() + else: + del per_tok_loss + batch_num = orig_batch_idx + 1 + doc_lens = [dl for _, dl in batch] + should_report = batch_num in eval_batch_set if eval_batch_set is not None else True + if should_report: + cur_tokens = token_count.item() + cur_loss_val = loss_sum.item() + cur_bytes_val = byte_sum.item() + dt = cur_tokens - prev_tokens + db = cur_bytes_val - prev_bytes + if dt > 0 and db > 0: + b_loss = (cur_loss_val - prev_loss) / dt + b_bpb = b_loss / math.log(2.0) * (dt / db) + else: + b_loss = b_bpb = 0.0 + r_loss = cur_loss_val / max(cur_tokens, 1) + r_bpb = r_loss / math.log(2.0) * (cur_tokens / max(cur_bytes_val, 1)) + elapsed = time.perf_counter() - t_start + log( + f"ttp: b{batch_num}/{queue_len} bl:{b_loss:.4f} bb:{b_bpb:.4f} " + f"rl:{r_loss:.4f} rb:{r_bpb:.4f} dl:{min(doc_lens)}-{max(doc_lens)} " + f"gd:{int(global_ttt_done)}" + ) + if not global_ttt_done: + local_scored_docs.extend( + (orig_batch_idx, pos, doc_start, doc_len) + for pos, (doc_start, doc_len) in enumerate(batch) + ) + prefix_done = _add_to_counter(prefix_counter_path, len(batch_entries)) + if prefix_done >= current_phase_boundary: + try: + with open(pause_flag_path, "x"): + pass + except FileExistsError: + pass + should_pause = os.path.exists(pause_flag_path) + if should_pause: + if dist.is_available() and dist.is_initialized(): + dist.barrier() + gathered_scored_docs = [None] * h.world_size + if dist.is_available() and dist.is_initialized(): + dist.all_gather_object(gathered_scored_docs, local_scored_docs) + else: + gathered_scored_docs = [local_scored_docs] + scored_docs_for_global = [] + for rank_docs in gathered_scored_docs: + if rank_docs: + scored_docs_for_global.extend(rank_docs) + scored_docs_for_global.sort(key=lambda x: (x[0], x[1])) + scored_docs_for_global = scored_docs_for_global[:current_phase_boundary] + scored_token_chunks = [ + val_data.val_tokens[doc_start : doc_start + doc_len] + for _, _, doc_start, doc_len in scored_docs_for_global + ] + if scored_token_chunks: + global_ttt_tokens = torch.cat(scored_token_chunks) + else: + global_ttt_tokens = val_data.val_tokens[:0] + if h.rank == 0: + prefix_done = 0 + try: + with open(prefix_counter_path, "rb") as f: + prefix_done = int.from_bytes( + f.read(8), "little", signed=True + ) + except FileNotFoundError: + pass + log( + f"ttpp: phase:{current_phase + 1}/{num_phases} pd:{prefix_done} " + f"gd:{len(scored_docs_for_global)} " + f"t:{time.perf_counter() - t_start:.1f}s" + ) + train_val_ttt_global_sgd_distributed( + h, device, val_data, base_model, global_ttt_tokens + ) + for p in base_model.parameters(): + p.requires_grad_(False) + reusable_lora = BatchedTTTLoRA( + h.ttt_batch_size, base_model, h.ttt_lora_rank, + k_lora=h.ttt_k_lora, mlp_lora=h.ttt_mlp_lora, o_lora=h.ttt_o_lora, + ).to(device) + reusable_opt = _build_opt(reusable_lora) + current_phase += 1 + if current_phase >= num_phases: + global_ttt_done = True + else: + current_phase_boundary = phase_boundaries[current_phase] + if h.rank == 0: + try: + os.remove(pause_flag_path) + except FileNotFoundError: + pass + if dist.is_available() and dist.is_initialized(): + dist.barrier() + if h.rank == 0: + log(f"ttpr: phase:{current_phase}/{num_phases} t:{time.perf_counter() - t_start:.1f}s") + del cur_lora, cur_opt + finally: + pass + if dist.is_available() and dist.is_initialized(): + dist.all_reduce(loss_sum, op=dist.ReduceOp.SUM) + dist.all_reduce(byte_sum, op=dist.ReduceOp.SUM) + dist.all_reduce(token_count, op=dist.ReduceOp.SUM) + for p in base_model.parameters(): + p.requires_grad_(True) + base_model.train() + return _loss_bpb_from_sums(loss_sum, token_count, byte_sum) + + +def timed_eval(label, fn, *args, **kwargs): + torch.cuda.synchronize() + t0 = time.perf_counter() + val_loss, val_bpb = fn(*args, **kwargs) + torch.cuda.synchronize() + elapsed_ms = 1e3 * (time.perf_counter() - t0) + log( + f"{label} val_loss:{val_loss:.8f} val_bpb:{val_bpb:.8f} eval_time:{elapsed_ms:.0f}ms" + ) + return val_loss, val_bpb + + +def train_model(h, device, val_data): + base_model = GPT(h).to(device).bfloat16() + restore_fp32_params(base_model) + _DISABLE_COMPILE = os.environ.get("DISABLE_COMPILE", "0") == "1" + if _DISABLE_COMPILE: + compiled_model = base_model + compiled_forward_logits = base_model.forward_logits + else: + compiled_model = torch.compile(base_model, dynamic=False, fullgraph=False) + compiled_forward_logits = torch.compile( + base_model.forward_logits, dynamic=False, fullgraph=False + ) + model = compiled_model + log(f"model_params:{sum(p.numel()for p in base_model.parameters())}") + optimizers = Optimizers(h, base_model) + train_loader = DocumentPackingLoader(h, device) + max_wallclock_ms = ( + 1e3 * h.max_wallclock_seconds if h.max_wallclock_seconds > 0 else None + ) + if max_wallclock_ms is not None: + max_wallclock_ms -= h.gptq_reserve_seconds * 1e3 + log( + f"gptq:reserving {h.gptq_reserve_seconds:.0f}s, effective={max_wallclock_ms:.0f}ms" + ) + + def training_frac(step, elapsed_ms): + if max_wallclock_ms is None: + return step / max(h.iterations, 1) + return elapsed_ms / max(max_wallclock_ms, 1e-09) + + def lr_mul(frac): + if h.warmdown_frac <= 0: + return 1.0 + if frac >= 1.0 - h.warmdown_frac: + return max((1.0 - frac) / h.warmdown_frac, h.min_lr) + return 1.0 + + def step_fn(step, lr_scale): + optimizers.zero_grad_all() + train_loss = torch.zeros((), device=device) + for micro_step in range(h.grad_accum_steps): + x, y, cu_seqlens, _max_seqlen = train_loader.next_batch( + h.train_batch_tokens, h.grad_accum_steps + ) + with torch.autocast(device_type="cuda", dtype=torch.bfloat16, enabled=True): + loss = model(x, y, cu_seqlens=cu_seqlens, max_seqlen=h.train_seq_len) + train_loss += loss.detach() + (loss / h.grad_accum_steps).backward() + train_loss /= h.grad_accum_steps + frac = ( + min(step / h.muon_momentum_warmup_steps, 1.0) + if h.muon_momentum_warmup_steps > 0 + else 1.0 + ) + muon_momentum = ( + 1 - frac + ) * h.muon_momentum_warmup_start + frac * h.muon_momentum + for group in optimizers.optimizer_muon.param_groups: + group["momentum"] = muon_momentum + for opt in optimizers: + for group in opt.param_groups: + group["lr"] = group["base_lr"] * lr_scale + if h.grad_clip_norm > 0: + torch.nn.utils.clip_grad_norm_(base_model.parameters(), h.grad_clip_norm) + optimizers.step(distributed=h.distributed) + return train_loss + + if h.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() + num_tokens_local = h.train_batch_tokens // h.world_size + for blk in base_model.blocks: + blk.attn.rotary(num_tokens_local, device, torch.bfloat16) + cu_bucket_size = train_loader.cu_bucket_size + warmup_cu_buckets = tuple(cu_bucket_size * i for i in range(1, 5)) + warmup_cu_iters = 3 + x, y, cu_seqlens, _ = train_loader.next_batch( + h.train_batch_tokens, h.grad_accum_steps + ) + log(f"warmup_cu_buckets:{','.join(str(b) for b in warmup_cu_buckets)} iters_each:{warmup_cu_iters}") + def _run_cu_bucket_warmup(): + for bucket_len in warmup_cu_buckets: + boundaries = list(range(0, x.size(1), max(h.train_seq_len, 1))) + if boundaries[-1] != x.size(1): + boundaries.append(x.size(1)) + cu = torch.full((bucket_len,), x.size(1), dtype=torch.int32, device=device) + cu[: len(boundaries)] = torch.tensor(boundaries, dtype=torch.int32, device=device) + for _ in range(warmup_cu_iters): + optimizers.zero_grad_all() + with torch.autocast(device_type="cuda", dtype=torch.bfloat16, enabled=True): + wloss = model(x, y, cu_seqlens=cu, max_seqlen=h.train_seq_len) + (wloss / h.grad_accum_steps).backward() + optimizers.zero_grad_all() + _run_cu_bucket_warmup() + if h.num_loops > 0: + base_model.looping_active = True + _run_cu_bucket_warmup() + base_model.looping_active = False + for warmup_step in range(h.warmup_steps): + step_fn(warmup_step, 1.0) + if ( + warmup_step <= 5 + or (warmup_step + 1) % 10 == 0 + or warmup_step + 1 == h.warmup_steps + ): + log(f"warmup_step: {warmup_step+1}/{h.warmup_steps}") + if h.num_loops > 0: + base_model.looping_active = True + log( + f"loop_warmup:enabled encoder:{base_model.encoder_indices} decoder:{base_model.decoder_indices}" + ) + for warmup_step in range(h.warmup_steps): + step_fn(warmup_step, 1.0) + if ( + warmup_step <= 5 + or (warmup_step + 1) % 10 == 0 + or warmup_step + 1 == h.warmup_steps + ): + log(f"loop_warmup_step: {warmup_step+1}/{h.warmup_steps}") + base_model.looping_active = False + 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) + optimizers.zero_grad_all() + train_loader = DocumentPackingLoader(h, device) + ema_state = { + name: t.detach().float().clone() + for (name, t) in base_model.state_dict().items() + } + ema_decay = h.ema_decay + training_time_ms = 0.0 + stop_after_step = None + torch.cuda.synchronize() + t0 = time.perf_counter() + step = 0 + while True: + last_step = ( + step == h.iterations + or stop_after_step is not None + and step >= stop_after_step + ) + should_validate = ( + last_step or h.val_loss_every > 0 and step % h.val_loss_every == 0 + ) + if should_validate: + torch.cuda.synchronize() + training_time_ms += 1e3 * (time.perf_counter() - t0) + val_loss, val_bpb = eval_val( + h, device, val_data, model, compiled_forward_logits + ) + log( + f"{step}/{h.iterations} val_loss: {val_loss:.4f} val_bpb: {val_bpb:.4f}" + ) + torch.cuda.synchronize() + t0 = time.perf_counter() + if last_step: + if stop_after_step is not None and step < h.iterations: + log( + f"stopping_early: wallclock_cap train_time: {training_time_ms:.0f}ms step: {step}/{h.iterations}" + ) + break + elapsed_ms = training_time_ms + 1e3 * (time.perf_counter() - t0) + frac = training_frac(step, elapsed_ms) + scale = lr_mul(frac) + if ( + h.num_loops > 0 + and not base_model.looping_active + and frac >= h.enable_looping_at + ): + base_model.looping_active = True + log( + f"layer_loop:enabled step:{step} frac:{frac:.3f} encoder:{base_model.encoder_indices} decoder:{base_model.decoder_indices}" + ) + train_loss = step_fn(step, scale) + 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 + approx_training_time_ms = training_time_ms + 1e3 * (time.perf_counter() - t0) + should_log_train = h.train_log_every > 0 and ( + step <= 5 or step % h.train_log_every == 0 or stop_after_step is not None + ) + if should_log_train: + tok_per_sec = step * h.train_batch_tokens / (approx_training_time_ms / 1e3) + log( + f"{step}/{h.iterations} train_loss: {train_loss.item():.4f} train_time: {approx_training_time_ms/60000:.1f}m tok/s: {tok_per_sec:.0f}" + ) + reached_cap = ( + max_wallclock_ms is not None and approx_training_time_ms >= max_wallclock_ms + ) + if h.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 + log( + f"peak memory allocated: {torch.cuda.max_memory_allocated()//1024//1024} MiB reserved: {torch.cuda.max_memory_reserved()//1024//1024} MiB" + ) + log("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) + return base_model, compiled_model, compiled_forward_logits + + +def train_and_eval(h, device): + random.seed(h.seed) + np.random.seed(h.seed) + torch.manual_seed(h.seed) + torch.cuda.manual_seed_all(h.seed) + if h.artifact_dir and h.is_main_process: + os.makedirs(h.artifact_dir, exist_ok=True) + val_data = ValidationData(h, device) + log( + f"train_shards: {len(list(Path(h.datasets_dir).resolve().glob('fineweb_train_*.bin')))}" + ) + log(f"val_tokens: {val_data.val_tokens.numel()-1}") + _TTT_EVAL_ONLY = os.environ.get("TTT_EVAL_ONLY", "0") == "1" + if _TTT_EVAL_ONLY: + artifact_override = os.environ.get("TTT_EVAL_ONLY_ARTIFACT", "").strip() + if artifact_override: + h.quantized_model_path = artifact_override + log( + f"TTT_EVAL_ONLY=1 — skipping train/GPTQ/quant eval; loading " + f"artifact={h.quantized_model_path}" + ) + base_model = None + compiled_model = None + compiled_forward_logits = None + else: + base_model, compiled_model, compiled_forward_logits = train_model( + h, device, val_data + ) + torch._dynamo.reset() + timed_eval( + "diagnostic pre-quantization post-ema", + eval_val, + h, + device, + val_data, + compiled_model, + compiled_forward_logits, + ) + serialize(h, base_model, Path(__file__).read_text(encoding="utf-8")) + if h.distributed: + dist.barrier() + eval_model = deserialize(h, device) + if h.num_loops > 0: + eval_model.looping_active = True + _DISABLE_COMPILE = ( + os.environ.get("DISABLE_COMPILE", "0") == "1" + or os.environ.get("DISABLE_EVAL_COMPILE", "0") == "1" + ) + if _DISABLE_COMPILE: + compiled_model = eval_model + compiled_forward_logits = eval_model.forward_logits + else: + compiled_model = torch.compile(eval_model, dynamic=False, fullgraph=False) + compiled_forward_logits = torch.compile( + eval_model.forward_logits, dynamic=False, fullgraph=False + ) + if not _TTT_EVAL_ONLY: + timed_eval( + "diagnostic quantized", + eval_val, + h, + device, + val_data, + compiled_model, + compiled_forward_logits, + ) + if h.sliding_window_enabled: + timed_eval( + "diagnostic quantized_sliding_window", + eval_val_sliding, + h, + device, + val_data, + eval_model, + forward_logits_fn=compiled_forward_logits, + ) + if h.ttt_enabled: + del eval_model, compiled_model + torch._dynamo.reset() + torch.cuda.empty_cache() + ttt_model = deserialize(h, device) + if h.num_loops > 0: + ttt_model.looping_active = True + for p in ttt_model.parameters(): + p.requires_grad_(False) + + if h.rope_yarn: + _yarn_seqlen = h.train_batch_tokens // h.grad_accum_steps + for block in ttt_model.blocks: + block.attn.rotary(_yarn_seqlen, device, torch.bfloat16) + else: + for block in ttt_model.blocks: + block.attn.rotary._cos_cached = None + block.attn.rotary._sin_cached = None + block.attn.rotary._seq_len_cached = 0 + block.attn.rotary(h.ttt_eval_seq_len, device, torch.bfloat16) + + def _fwd_ttt_inner(input_ids, target_ids, lora): + return ttt_model.forward_ttt(input_ids, target_ids, lora=lora) + + _DISABLE_TTT_COMPILE = ( + os.environ.get("DISABLE_COMPILE", "0") == "1" + or os.environ.get("DISABLE_TTT_COMPILE", "0") == "1" + ) + # dynamic=True caused stride/shape assertions on ROCm inductor; force + # dynamic=False so compile specializes per concrete (bsz, seq_len). + _TTT_COMPILE_DYNAMIC = ( + os.environ.get("TTT_COMPILE_DYNAMIC", "0") == "1" + ) + _TTT_FIXED_SEQ_COMPILE = ( + os.environ.get("TTT_FIXED_SEQ_COMPILE", "0") == "1" + ) + _TTT_COMPILE_BUCKET_SIZE = int( + os.environ.get("TTT_COMPILE_BUCKET_SIZE", "128" if _IS_ROCM else "0") + ) + _fwd_ttt_compiled_inners = {} + + def _get_fwd_ttt_compiled(bsz, seq_len): + key = ("dynamic",) if _TTT_COMPILE_DYNAMIC else (int(bsz), int(seq_len)) + compiled = _fwd_ttt_compiled_inners.get(key) + if compiled is None: + compiled = torch.compile(_fwd_ttt_inner, dynamic=_TTT_COMPILE_DYNAMIC) + _fwd_ttt_compiled_inners[key] = compiled + return compiled + + def _fwd_ttt(input_ids, target_ids, lora): + if _DISABLE_TTT_COMPILE: + return _fwd_ttt_inner(input_ids, target_ids, lora=lora) + orig_len = input_ids.size(1) + target_len = orig_len + if _TTT_FIXED_SEQ_COMPILE: + target_len = h.ttt_eval_seq_len + elif _TTT_COMPILE_BUCKET_SIZE > 0: + target_len = min( + h.ttt_eval_seq_len, + ((orig_len + _TTT_COMPILE_BUCKET_SIZE - 1) // _TTT_COMPILE_BUCKET_SIZE) + * _TTT_COMPILE_BUCKET_SIZE, + ) + if target_len > orig_len: + pad_len = target_len - orig_len + input_ids = F.pad(input_ids, (0, pad_len), value=0) + target_ids = F.pad(target_ids, (0, pad_len), value=0) + compiled = _get_fwd_ttt_compiled(input_ids.size(0), input_ids.size(1)) + out = compiled(input_ids, target_ids, lora=lora) + return out[:, :orig_len] if target_len > orig_len else out + + fwd_ttt_compiled = _fwd_ttt + log(f"ttt_lora:warming up compile (random tokens, no val data)") + global BOS_ID + if BOS_ID is None: + BOS_ID = 1 + t_warmup = time.perf_counter() + warmup_bszes = [h.ttt_batch_size] + for bsz in warmup_bszes: + wl = BatchedTTTLoRA( + bsz, ttt_model, h.ttt_lora_rank, + k_lora=h.ttt_k_lora, mlp_lora=h.ttt_mlp_lora, o_lora=h.ttt_o_lora, + ).to(device) + wo = torch.optim.AdamW( + wl.parameters(), + lr=h.ttt_lora_lr, + betas=(h.ttt_beta1, h.ttt_beta2), + eps=1e-10, + weight_decay=h.ttt_weight_decay, + fused=True, + ) + warmup_ctx_lens = ( + (h.ttt_eval_seq_len,) + if _TTT_FIXED_SEQ_COMPILE and not _DISABLE_TTT_COMPILE + else ( + ( + min( + h.ttt_eval_seq_len, + ((h.ttt_chunk_size + _TTT_COMPILE_BUCKET_SIZE - 1) + // _TTT_COMPILE_BUCKET_SIZE) + * _TTT_COMPILE_BUCKET_SIZE, + ), + h.ttt_eval_seq_len, + ) + if _TTT_COMPILE_BUCKET_SIZE > 0 and not _DISABLE_TTT_COMPILE + else (h.ttt_chunk_size, h.ttt_eval_seq_len) + ) + ) + for ctx_len in sorted(set(warmup_ctx_lens)): + xw = torch.randint(0, h.vocab_size, (bsz, ctx_len), device=device, dtype=torch.int64) + yw = torch.randint(0, h.vocab_size, (bsz, ctx_len), device=device, dtype=torch.int64) + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + ptl = fwd_ttt_compiled(xw, yw, lora=wl) + ptl[:, : min(h.ttt_chunk_size, ctx_len)].mean(dim=-1).sum().backward() + wo.step() + wo.zero_grad(set_to_none=True) + del wl, wo + torch.cuda.empty_cache() + compile_elapsed = time.perf_counter() - t_warmup + log(f"ttt_lora:compile warmup done ({compile_elapsed:.1f}s)") + log("\nbeginning TTT eval timer") + torch.cuda.synchronize() + t_ttt = time.perf_counter() + ttt_val_loss, ttt_val_bpb = eval_val_ttt_phased( + h, ttt_model, device, val_data, forward_ttt_train=fwd_ttt_compiled + ) + torch.cuda.synchronize() + ttt_eval_elapsed = time.perf_counter() - t_ttt + log( + "quantized_ttt_phased " + f"val_loss:{ttt_val_loss:.8f} val_bpb:{ttt_val_bpb:.8f} " + f"eval_time:{1e3*ttt_eval_elapsed:.0f}ms" + ) + log(f"total_eval_time:{ttt_eval_elapsed:.1f}s") + del ttt_model + + +def main(): + world_size = int(os.environ.get("WORLD_SIZE", "1")) + local_rank = int(os.environ.get("LOCAL_RANK", "0")) + distributed = "RANK" in os.environ and "WORLD_SIZE" in os.environ + if not torch.cuda.is_available(): + raise RuntimeError("CUDA is required") + 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" + ) + device = torch.device("cuda", local_rank) + torch.cuda.set_device(device) + if distributed: + dist.init_process_group( + backend="nccl", + device_id=device, + timeout=datetime.timedelta( + hours=float(os.environ.get("DIST_TIMEOUT_HOURS", "12")) + ), + ) + dist.barrier() + torch.backends.cuda.matmul.allow_tf32 = True + torch.backends.cudnn.allow_tf32 = True + torch.set_float32_matmul_precision("high") + 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) + torch._dynamo.config.optimize_ddp = False + torch._dynamo.config.cache_size_limit = int( + os.environ.get("DYNAMO_CACHE_SIZE_LIMIT", "256" if _IS_ROCM else "64") + ) + # Allow disabling inductor shape/stride assertions so the compiled kernel + # runs even when its guessed shape pattern (e.g. ROCm-specific stride + # expressions) doesn't match the actual runtime shape. Set + # TORCHINDUCTOR_SIZE_ASSERTS=0 to silence. + if os.environ.get("TORCHINDUCTOR_SIZE_ASSERTS", "1") == "0": + import torch._inductor.config as _inductor_config + _inductor_config.size_asserts = False + _inductor_config.nan_asserts = False + h = Hyperparameters() + set_logging_hparams(h) + if h.is_main_process: + os.makedirs(h.artifact_dir if h.artifact_dir else "logs", exist_ok=True) + log(100 * "=", console=False) + log("Hyperparameters:", console=True) + for (k, v) in sorted(vars(type(h)).items()): + if not k.startswith("_"): + log(f" {k}: {v}", console=True) + log("=" * 100, console=False) + log("Source code:", console=False) + log("=" * 100, console=False) + with open(__file__, "r", encoding="utf-8") as _src: + log(_src.read(), console=False) + log("=" * 100, console=False) + log(f"Running Python {sys.version}", console=False) + log(f"Running PyTorch {torch.__version__}", console=False) + log("=" * 100, console=False) + train_and_eval(h, device) + if distributed: + dist.destroy_process_group() + + +if __name__ == "__main__": + main()