diff --git a/records/track_10min_16mb/2026-04-30_TTT_Ensemble/README.md b/records/track_10min_16mb/2026-04-30_TTT_Ensemble/README.md new file mode 100644 index 0000000000..336a4ff351 --- /dev/null +++ b/records/track_10min_16mb/2026-04-30_TTT_Ensemble/README.md @@ -0,0 +1,80 @@ +# Record: TTT Peer-LoRA Ensemble + +**val_bpb = 1.05749** (1 seed) | **~15.99 MB** | 8xH100 SXM | PyTorch 2.10.0+cu130 + +This record introduces peer-LoRA ensembling into the test-time training (TTT) evaluation loop. After each batch's per-doc LoRAs are fully trained, we run k-1 additional forwards using *other* docs' trained LoRAs from the same batch. This is leakage-free: LoRA_p was trained only on doc_p's tokens, so applying it to doc_q reveals no target information. On uncertain tokens (high predictive entropy), we blend own and peer predictions in probability space; confident tokens use only their own prediction. The routing decision is target-free -- it depends only on the model's output distribution, not on validation labels. + +Built on [PR #2014](https://github.com/openai/parameter-golf/pull/2014), descending from @samacqua's work on doc-independent LoRAs. + +## Results + +| Seed | Pre-Quant BPB | Post-Quant BPB | **Post-TTT BPB** | Artifact | +|------|--------------:|---------------:|-----------------:|---------:| +| 42 | 1.05899 | 1.06755 | **1.05749** | 15,986,824 | + +Baseline PR #2014 3-seed mean: val_bpb 1.05855 (as reported by @simonbissonnette). + +Delta: -0.00106 vs PR #2014 baseline (1.05855) + +## Key Changes vs PR #2014 + +### 1. Peer-LoRA ensemble with confidence routing + +After each batch's per-doc LoRAs finish sliding-window training (k docs per batch -> k independent LoRAs), run k-1 peer forwards per doc using other docs' LoRAs: + +- **Stash phase**: during the normal sliding-window eval, stash each doc's per-token NLLs and predictive entropies (entropy of the output distribution -- no target labels used). +- **Peer phase**: for each doc, run k-1 forwards with randomly-selected peer LoRAs from the same batch. `BatchedLinearLoRA.PEER_IDX` routes each batch row to a different doc's LoRA weights. +- **Blend**: on tokens where `predictive_entropy >= threshold` (uncertain), blend: `p = w * p_own + (1 - w) * mean(p_peers)`. Confident tokens use `p_own` only. + +The routing gate is target-free: it uses the model's own entropy, not validation NLLs. This means the ensemble prediction is committed before seeing targets, avoiding post-hoc selection. + +With `threshold = 0.5`, roughly 75% of tokens are routed through the ensemble. + +### 2. TTT hyperparameter tuning + +Per-doc LoRA LR and weight decay were tuned via line search (on a single H100, using `TTT_EVAL_ONLY` to skip retraining): + +| Param | PR #2014 | This submission | +|---|---:|---:| +| `TTT_LORA_LR` | 0.0001 | 0.00015 | +| `TTT_WEIGHT_DECAY` | 0.5 | 0.25 | + +Higher LR lets the per-doc LoRAs fit more aggressively; lower weight decay gives them more freedom. Both changes improve the baseline and the peer ensemble independently. + +## New Env Vars + +| Env var | Default | Description | +|---|---:|---| +| `TTT_PEER_ENSEMBLE_K` | 4 | Peers per batch incl. self (set 1 to disable) | +| `TTT_PEER_CONF_THRESHOLD` | 0.5 | Predictive entropy threshold for routing | +| `TTT_PEER_CONF_BLEND_W` | 0.8 | Weight on own prediction in blend | + +## Reproducing + +Uses the same CaseOps sp8192 dataset/tokenizer as PR #2014, sourced from HuggingFace: + +- Dataset: `romeerp/parameter-golf-caseops-v1` +- Variant: `sp8192_lossless_caps_caseops_v1_reserved` + +```bash +# Install lrzip (artifact compression) +sudo apt-get install -y lrzip + +# Download data +python3 data/cached_challenge_fineweb.py --variant sp8192_lossless_caps_caseops_v1_reserved + +# Run +SEED=42 torchrun --standalone --nproc_per_node=8 train_gpt.py +``` + +All hyperparameters (CASEOPS_ENABLED=1, VOCAB_SIZE=8192, ensemble settings, etc.) are baked into `train_gpt.py`. + +## Hardware / Software + +- 8xH100 80GB SXM +- PyTorch 2.10.0+cu130 +- `lrzip` 0.651 (for `pergroup` compression) + +## Attribution + +See `submission.json`. Built on the PR #2014 stack (@simonbissonnette and earlier contributors). diff --git a/records/track_10min_16mb/2026-04-30_TTT_Ensemble/seed42.log b/records/track_10min_16mb/2026-04-30_TTT_Ensemble/seed42.log new file mode 100644 index 0000000000..ad4b9655fc --- /dev/null +++ b/records/track_10min_16mb/2026-04-30_TTT_Ensemble/seed42.log @@ -0,0 +1,5566 @@ +==================================================================================================== +Hyperparameters: + adam_eps: 1e-08 + adam_wd: 0.02 + artifact_dir: + attn_clip_sigmas: 13.0 + attn_out_gate_enabled: False + attn_out_gate_src: proj + awq_lite_bits: 8 + awq_lite_enabled: True + awq_lite_group_size: 64 + awq_lite_group_top_k: 1 + beta1: 0.9 + beta2: 0.99 + caseops_enabled: True + compile_shape_warmup: True + compile_shape_warmup_iters: 1 + compile_shape_warmup_loop_modes: auto + compressor: pergroup + data_dir: ./data/ + datasets_dir: ./data/datasets/datasets/fineweb10B_sp8192_lossless_caps_caseops_v1_reserved + 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_include_tail: True + eval_seq_len: 3072 + eval_stride: 1536 + fused_ce_enabled: True + gate_window: 12 + gated_attn_enabled: False + gated_attn_init_std: 0.01 + gated_attn_quant_gate: True + 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: logs/8cc79f06-6706-4ab5-ad63-a7a500595515.txt + logit_softcap: 30.0 + loop_end: 5 + loop_start: 3 + lqer_asym_enabled: True + lqer_asym_group: 64 + lqer_enabled: True + lqer_factor_bits: 4 + lqer_gain_select: False + lqer_rank: 4 + lqer_scope: all + lqer_top_k: 3 + matrix_bits: 6 + matrix_clip_sigmas: 12.85 + matrix_lr: 0.026 + max_wallclock_seconds: 600.0 + midrun_cap_log_updates: False + midrun_cap_schedule: + min_lr: 0.1 + mlp_clip_sigmas: 11.5 + mlp_mult: 4.0 + model_dim: 512 + model_path: 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_num_phases: 1 + phased_ttt_prefix_docs: 2500 + qk_gain_init: 5.25 + quantized_model_path: final_model.int6.ptz + rank: 0 + rope_base: 10000.0 + rope_dims: 16 + rope_train_seq_len: 3072 + rope_yarn: False + run_id: 8cc79f06-6706-4ab5-ad63-a7a500595515 + scalar_lr: 0.02 + seed: 42 + seq_change_warmup_steps: 32 + skip_gates_enabled: True + smear_gate_enabled: True + sparse_attn_gate_enabled: True + sparse_attn_gate_init_std: 0.0 + sparse_attn_gate_scale: 0.5 + tie_embeddings: True + tied_embed_init_std: 0.005 + tied_embed_lr: 0.03 + tokenizer_path: ./data/datasets/tokenizers/fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model + train_batch_tokens: 786432 + train_files: ./data/datasets/datasets/fineweb10B_sp8192_lossless_caps_caseops_v1_reserved/fineweb_train_*.bin + train_log_every: 500 + train_seq_len: 3072 + train_seq_schedule: 1024@0.100,2048@0.700,3072@1.000 + train_seq_schedule_mode: wallclock + ttt_batch_size: 24 + ttt_beta1: 0.0 + ttt_beta2: 0.99 + ttt_chunk_size: 48 + ttt_enabled: True + ttt_eval_batches: + ttt_eval_seq_len: 3072 + ttt_grad_steps: 1 + ttt_k_lora: True + ttt_local_lr_mult: 0.75 + ttt_lora_lr: 0.00015 + ttt_lora_rank: 80 + ttt_mask: no_qv + ttt_mlp_lora: True + ttt_o_lora: True + ttt_optimizer: adam + ttt_peer_conf_blend_w: 0.8 + ttt_peer_conf_threshold: 0.5 + ttt_peer_ensemble_k: 4 + ttt_q_lora: False + ttt_short_beta2: 0.99 + ttt_short_chunk_size: 24 + ttt_short_doc_len: 2000 + ttt_short_lora_enabled: False + ttt_short_lora_lr: 0.00015 + ttt_short_lora_rank: 80 + ttt_short_score_first_enabled: True + ttt_short_score_first_steps: 256:8,2000:24 + ttt_short_weight_decay: 0.25 + ttt_train_max_doc_len: 0 + ttt_train_min_doc_len: 0 + ttt_v_lora: False + ttt_warm_start_mean_doc_len: 2000 + ttt_warm_start_mean_enabled: False + ttt_warm_start_mean_momentum: 0.95 + ttt_weight_decay: 0.25 + val_batch_tokens: 524288 + val_bytes_files: ./data/datasets/datasets/fineweb10B_sp8192_lossless_caps_caseops_v1_reserved/fineweb_val_bytes_*.bin + val_doc_fraction: 1.0 + val_files: ./data/datasets/datasets/fineweb10B_sp8192_lossless_caps_caseops_v1_reserved/fineweb_val_*.bin + val_loss_every: 0 + vocab_size: 8192 + warmdown_frac: 0.85 + warmdown_iters: 0 + warmup_steps: 20 + world_size: 8 + xsa_last_n: 11 +==================================================================================================== +Source code: +==================================================================================================== +""" +parent: experiments/parameter-golf/pr2014.py +""" +import base64, collections, copy, fcntl, glob, io, lzma, math, os +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 Tensor, nn +from flash_attn_interface import ( + flash_attn_func as flash_attn_3_func, + flash_attn_varlen_func, +) +from concurrent.futures import ThreadPoolExecutor +import triton +import triton.language as tl +from triton.tools.tensor_descriptor import TensorDescriptor + + +# ===== Fused softcapped cross-entropy (Triton) — training-only path ===== +# Replaces the eager +# logits_softcap = softcap * tanh(logits / softcap) +# F.cross_entropy(logits_softcap.float(), targets, reduction="mean") +# sequence with a single fused kernel that reads logits_proj once, applies +# softcap in-register, and computes (LSE, loss) in one streaming pass. The +# backward kernel mirrors the forward so there's no stored softcapped logits. +# Numerically identical to the eager path up to fp32 accumulation differences. +_FUSED_CE_LIBRARY = "pgsubmission1draft7fusedce" +_FUSED_CE_BLOCK_SIZE = 1024 +_FUSED_CE_NUM_WARPS = 4 + + +@triton.jit +def _softcapped_ce_fwd_kernel( + logits_ptr, losses_ptr, lse_ptr, targets_ptr, + stride_logits_n, stride_logits_v, + n_rows, n_cols, softcap, + block_size: tl.constexpr, +): + row_idx = tl.program_id(0).to(tl.int64) + logits_row_ptr = logits_ptr + row_idx * stride_logits_n + max_val = -float("inf") + sum_exp = 0.0 + A = 2.0 * softcap + inv_C = 2.0 / softcap + for off in range(0, n_cols, block_size): + cols = off + tl.arange(0, block_size) + mask = cols < n_cols + val = tl.load( + logits_row_ptr + cols * stride_logits_v, + mask=mask, other=-float("inf"), + ).to(tl.float32) + z = A * tl.sigmoid(val * inv_C) + z = tl.where(mask, z, -float("inf")) + curr_max = tl.max(z, axis=0) + new_max = tl.maximum(max_val, curr_max) + sum_exp = sum_exp * tl.exp(max_val - new_max) + tl.sum(tl.exp(z - new_max), axis=0) + max_val = new_max + lse = max_val + tl.log(sum_exp) + tl.store(lse_ptr + row_idx, lse) + target = tl.load(targets_ptr + row_idx).to(tl.int32) + target_val = tl.load(logits_row_ptr + target * stride_logits_v).to(tl.float32) + target_z = A * tl.sigmoid(target_val * inv_C) + tl.store(losses_ptr + row_idx, lse - target_z) + + +@triton.jit +def _softcapped_ce_bwd_kernel( + grad_logits_ptr, grad_losses_ptr, lse_ptr, logits_ptr, targets_ptr, + stride_logits_n, stride_logits_v, + stride_grad_n, stride_grad_v, + n_rows, n_cols, softcap, + block_size: tl.constexpr, +): + row_idx = tl.program_id(0).to(tl.int64) + logits_row_ptr = logits_ptr + row_idx * stride_logits_n + grad_row_ptr = grad_logits_ptr + row_idx * stride_grad_n + lse = tl.load(lse_ptr + row_idx) + grad_loss = tl.load(grad_losses_ptr + row_idx).to(tl.float32) + target = tl.load(targets_ptr + row_idx).to(tl.int32) + A = 2.0 * softcap + inv_C = 2.0 / softcap + dz_dx_scale = A * inv_C + for off in range(0, n_cols, block_size): + cols = off + tl.arange(0, block_size) + mask = cols < n_cols + val = tl.load( + logits_row_ptr + cols * stride_logits_v, + mask=mask, other=0.0, + ).to(tl.float32) + sigmoid_u = tl.sigmoid(val * inv_C) + z = A * sigmoid_u + probs = tl.exp(z - lse) + grad_z = grad_loss * (probs - tl.where(cols == target, 1.0, 0.0)) + grad_x = grad_z * (dz_dx_scale * sigmoid_u * (1.0 - sigmoid_u)) + tl.store(grad_row_ptr + cols * stride_grad_v, grad_x, mask=mask) + + +def _validate_softcapped_ce_inputs( + logits: Tensor, targets: Tensor, softcap: float, +) -> tuple[Tensor, Tensor]: + if logits.ndim != 2: + raise ValueError(f"Expected logits.ndim=2, got {logits.ndim}") + if targets.ndim != 1: + raise ValueError(f"Expected targets.ndim=1, got {targets.ndim}") + if logits.shape[0] != targets.shape[0]: + raise ValueError( + f"Expected matching rows, got logits={tuple(logits.shape)} targets={tuple(targets.shape)}" + ) + if not logits.is_cuda or not targets.is_cuda: + raise ValueError("softcapped_cross_entropy requires CUDA tensors") + if softcap <= 0.0: + raise ValueError(f"softcap must be positive, got {softcap}") + if logits.dtype not in (torch.float16, torch.bfloat16, torch.float32): + raise ValueError(f"Unsupported logits dtype: {logits.dtype}") + logits = logits.contiguous() + targets = targets.contiguous() + if targets.dtype != torch.int64: + targets = targets.to(dtype=torch.int64) + return logits, targets + + +@torch.library.custom_op(f"{_FUSED_CE_LIBRARY}::softcapped_ce", mutates_args=()) +def softcapped_ce_op(logits: Tensor, targets: Tensor, softcap: float) -> tuple[Tensor, Tensor]: + logits, targets = _validate_softcapped_ce_inputs(logits, targets, float(softcap)) + n_rows, n_cols = logits.shape + losses = torch.empty((n_rows,), device=logits.device, dtype=torch.float32) + lse = torch.empty((n_rows,), device=logits.device, dtype=torch.float32) + _softcapped_ce_fwd_kernel[(n_rows,)]( + logits, losses, lse, targets, + logits.stride(0), logits.stride(1), + n_rows, n_cols, float(softcap), + block_size=_FUSED_CE_BLOCK_SIZE, num_warps=_FUSED_CE_NUM_WARPS, + ) + return losses, lse + + +@softcapped_ce_op.register_fake +def _(logits: Tensor, targets: Tensor, softcap: float): + if logits.ndim != 2 or targets.ndim != 1: + raise ValueError("softcapped_ce fake impl expects 2D logits and 1D targets") + if logits.shape[0] != targets.shape[0]: + raise ValueError( + f"Expected matching rows, got logits={tuple(logits.shape)} targets={tuple(targets.shape)}" + ) + n_rows = logits.shape[0] + return ( + logits.new_empty((n_rows,), dtype=torch.float32), + logits.new_empty((n_rows,), dtype=torch.float32), + ) + + +@torch.library.custom_op(f"{_FUSED_CE_LIBRARY}::softcapped_ce_backward", mutates_args=()) +def softcapped_ce_backward_op( + logits: Tensor, targets: Tensor, lse: Tensor, grad_losses: Tensor, softcap: float, +) -> Tensor: + logits, targets = _validate_softcapped_ce_inputs(logits, targets, float(softcap)) + lse = lse.contiguous() + grad_losses = grad_losses.contiguous().to(dtype=torch.float32) + if lse.ndim != 1 or grad_losses.ndim != 1: + raise ValueError("Expected 1D lse and grad_losses") + if lse.shape[0] != logits.shape[0] or grad_losses.shape[0] != logits.shape[0]: + raise ValueError( + f"Expected row-aligned lse/grad_losses, got logits={tuple(logits.shape)} " + f"lse={tuple(lse.shape)} grad_losses={tuple(grad_losses.shape)}" + ) + grad_logits = torch.empty_like(logits) + n_rows, n_cols = logits.shape + _softcapped_ce_bwd_kernel[(n_rows,)]( + grad_logits, grad_losses, lse, logits, targets, + logits.stride(0), logits.stride(1), + grad_logits.stride(0), grad_logits.stride(1), + n_rows, n_cols, float(softcap), + block_size=_FUSED_CE_BLOCK_SIZE, num_warps=_FUSED_CE_NUM_WARPS, + ) + return grad_logits + + +@softcapped_ce_backward_op.register_fake +def _(logits: Tensor, targets: Tensor, lse: Tensor, grad_losses: Tensor, softcap: float): + if logits.ndim != 2 or targets.ndim != 1 or lse.ndim != 1 or grad_losses.ndim != 1: + raise ValueError("softcapped_ce_backward fake impl expects 2D logits and 1D row tensors") + if ( + logits.shape[0] != targets.shape[0] + or logits.shape[0] != lse.shape[0] + or logits.shape[0] != grad_losses.shape[0] + ): + raise ValueError("softcapped_ce_backward fake impl expects row-aligned tensors") + return logits.new_empty(logits.shape) + + +def _softcapped_ce_setup_context( + ctx: torch.autograd.function.FunctionCtx, inputs, output, +) -> None: + logits, targets, softcap = inputs + _losses, lse = output + ctx.save_for_backward(logits, targets, lse) + ctx.softcap = float(softcap) + + +def _softcapped_ce_backward( + ctx: torch.autograd.function.FunctionCtx, grad_losses: Tensor, grad_lse: "Tensor | None", +): + del grad_lse + logits, targets, lse = ctx.saved_tensors + grad_logits = torch.ops.pgsubmission1draft7fusedce.softcapped_ce_backward( + logits, targets, lse, grad_losses, ctx.softcap + ) + return grad_logits, None, None + + +softcapped_ce_op.register_autograd( + _softcapped_ce_backward, setup_context=_softcapped_ce_setup_context, +) + + +def softcapped_cross_entropy( + logits: Tensor, targets: Tensor, softcap: float, reduction: str = "mean", +) -> Tensor: + losses, _lse = torch.ops.pgsubmission1draft7fusedce.softcapped_ce( + logits, targets, float(softcap) + ) + if reduction == "none": + return losses + if reduction == "sum": + return losses.sum() + if reduction == "mean": + return losses.mean() + raise ValueError(f"Unsupported reduction={reduction!r}") + + +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.85)) + warmdown_iters = int(os.environ.get("WARMDOWN_ITERS", 0)) + midrun_cap_schedule = os.environ.get("MIDRUN_CAP_SCHEDULE", "").strip() + midrun_cap_log_updates = bool(int(os.environ.get("MIDRUN_CAP_LOG_UPDATES", "0"))) + warmup_steps = int(os.environ.get("WARMUP_STEPS", 20)) + train_batch_tokens = int(os.environ.get("TRAIN_BATCH_TOKENS", 786432)) + # Fused softcapped CE (Triton). Training-only — forward_logits eval path still uses + # eager softcap+F.cross_entropy. Default ON since validated as at-worst neutral. + fused_ce_enabled = bool(int(os.environ.get("FUSED_CE_ENABLED", "1"))) + train_seq_len = int(os.environ.get("TRAIN_SEQ_LEN", 3072)) + train_seq_schedule = os.environ.get("TRAIN_SEQ_SCHEDULE", "1024@0.100,2048@0.700,3072@1.000") + train_seq_schedule_mode = os.environ.get("TRAIN_SEQ_SCHEDULE_MODE", "wallclock").strip().lower() + seq_change_warmup_steps = int(os.environ.get("SEQ_CHANGE_WARMUP_STEPS", 32)) + compile_shape_warmup = bool(int(os.environ.get("COMPILE_SHAPE_WARMUP", "1"))) + compile_shape_warmup_iters = int(os.environ.get("COMPILE_SHAPE_WARMUP_ITERS", "1")) + compile_shape_warmup_loop_modes = os.environ.get("COMPILE_SHAPE_WARMUP_LOOP_MODES", "auto").strip().lower() + 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", 3072)) + val_loss_every = int(os.environ.get("VAL_LOSS_EVERY", 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", 3072)) + 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.1)) + 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.99)) + 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", 1536)) + eval_include_tail = bool(int(os.environ.get("EVAL_INCLUDE_TAIL", "1"))) + 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", 80)) + ttt_lora_lr = float(os.environ.get("TTT_LORA_LR", 0.00015)) + ttt_local_lr_mult = float(os.environ.get("TTT_LOCAL_LR_MULT", 0.75)) + ttt_chunk_size = int(os.environ.get("TTT_CHUNK_SIZE", 48)) + ttt_eval_seq_len = int(os.environ.get("TTT_EVAL_SEQ_LEN", 3072)) + ttt_batch_size = int(os.environ.get("TTT_BATCH_SIZE", 24)) + ttt_grad_steps = int(os.environ.get("TTT_GRAD_STEPS", 1)) + # V19: PR #1886 (renqianluo) + sunnypatneedi research log 2026-04-28 found that + # the Triton fused-CE kernel's fp32-accumulation interacts with warm-start LoRA-A + # to destabilize seeds 314/1337 at TTT_WEIGHT_DECAY=1.0. Raising the default to + # 2.0 prevents seed collapse without measurably moving stable seeds. + ttt_weight_decay = float(os.environ.get("TTT_WEIGHT_DECAY", 0.25)) + ttt_beta1 = float(os.environ.get("TTT_BETA1", 0)) + ttt_beta2 = float(os.environ.get("TTT_BETA2", 0.99)) + ttt_mask = os.environ.get("TTT_MASK", "no_qv").strip().lower() + _ttt_q_default = "1" + _ttt_v_default = "1" + if ttt_mask in ("", "all", "baseline_all"): + pass + elif ttt_mask == "no_q": + _ttt_q_default = "0" + elif ttt_mask == "no_v": + _ttt_v_default = "0" + elif ttt_mask == "no_qv": + _ttt_q_default = "0" + _ttt_v_default = "0" + else: + raise ValueError(f"Unsupported TTT_MASK={ttt_mask!r}") + ttt_q_lora = bool(int(os.environ.get("TTT_Q_LORA", _ttt_q_default))) + ttt_k_lora = bool(int(os.environ.get("TTT_K_LORA", "1"))) + ttt_v_lora = bool(int(os.environ.get("TTT_V_LORA", _ttt_v_default))) + 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", "") + ttt_peer_ensemble_k = int(os.environ.get("TTT_PEER_ENSEMBLE_K", "4")) + ttt_peer_conf_threshold = float(os.environ.get("TTT_PEER_CONF_THRESHOLD", "0.5")) + ttt_peer_conf_blend_w = float(os.environ.get("TTT_PEER_CONF_BLEND_W", "0.8")) + ttt_short_doc_len = int(os.environ.get("TTT_SHORT_DOC_LEN", 2000)) + ttt_short_lora_enabled = bool(int(os.environ.get("TTT_SHORT_LORA_ENABLED", "0"))) + ttt_short_lora_rank = int(os.environ.get("TTT_SHORT_LORA_RANK", ttt_lora_rank)) + ttt_short_lora_lr = float(os.environ.get("TTT_SHORT_LORA_LR", ttt_lora_lr)) + ttt_short_weight_decay = float(os.environ.get("TTT_SHORT_WEIGHT_DECAY", ttt_weight_decay)) + ttt_short_beta2 = float(os.environ.get("TTT_SHORT_BETA2", ttt_beta2)) + ttt_short_score_first_enabled = bool(int(os.environ.get("TTT_SHORT_SCORE_FIRST_ENABLED", "1"))) + ttt_short_chunk_size = int(os.environ.get("TTT_SHORT_CHUNK_SIZE", 24)) + ttt_short_score_first_steps = os.environ.get("TTT_SHORT_SCORE_FIRST_STEPS", "256:8,2000:24") + ttt_train_min_doc_len = int(os.environ.get("TTT_TRAIN_MIN_DOC_LEN", "0")) + ttt_train_max_doc_len = int(os.environ.get("TTT_TRAIN_MAX_DOC_LEN", "0")) + ttt_warm_start_mean_enabled = bool(int(os.environ.get("TTT_WARM_START_MEAN_ENABLED", "0"))) + ttt_warm_start_mean_doc_len = int(os.environ.get("TTT_WARM_START_MEAN_DOC_LEN", ttt_short_doc_len)) + ttt_warm_start_mean_momentum = float(os.environ.get("TTT_WARM_START_MEAN_MOMENTUM", 0.95)) + val_doc_fraction = float(os.environ.get("VAL_DOC_FRACTION", 1.0)) + compressor = os.environ.get("COMPRESSOR", "pergroup") + 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_prefix_docs = int(os.environ.get("PHASED_TTT_PREFIX_DOCS", 2500)) + phased_ttt_num_phases = int(os.environ.get("PHASED_TTT_NUM_PHASES", 1)) + 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", 14.0)) + mlp_clip_sigmas = float(os.environ.get("MLP_CLIP_SIGMAS", 11.5)) + attn_clip_sigmas = float(os.environ.get("ATTN_CLIP_SIGMAS", 13.0)) + # AttnOutGate (per-head multiplicative output gate, PR #1667 MarioPaerle). + # Zero-init weight: 2*sigmoid(0)=1 -> transparent at start. Source defaults to + # block input x ('proj'); 'q' uses raw Q projection output. + attn_out_gate_enabled = bool(int(os.environ.get("ATTN_OUT_GATE_ENABLED", "0"))) + attn_out_gate_src = os.environ.get("ATTN_OUT_GATE_SRC", "proj") + # SmearGate (input-dependent forward-1 token smear, modded-nanogpt @classiclarryd + # via PR #1667). x_t <- x_t + lam * sigmoid(W*x_t[:gate_window]) * x_{t-1}. + # lam=0 + W=0 -> transparent at init. + smear_gate_enabled = bool(int(os.environ.get("SMEAR_GATE_ENABLED", "1"))) + # Window: first GATE_WINDOW dims of the source feed the gate projection. + gate_window = int(os.environ.get("GATE_WINDOW", 12)) + # Gated Attention (Qwen, NeurIPS 2025 Best Paper, arXiv:2505.06708; + # qiuzh20/gated_attention). Per-head sigmoid gate on SDPA output, BEFORE + # out_proj. Gate input = full block input x (paper's headwise G1 variant + # driven from hidden_states). W_g shape (num_heads, dim), plain sigmoid. + # Near-zero init gives g~0.5 at step 0 (half attention output); per-block + # attn_scale (init 1.0) compensates during training. Name contains + # "attn_gate" so CONTROL_TENSOR_NAME_PATTERNS routes it to scalar AdamW. + gated_attn_enabled = bool(int(os.environ.get("GATED_ATTN_ENABLED", "0"))) + gated_attn_init_std = float(os.environ.get("GATED_ATTN_INIT_STD", 0.01)) + # Dedicated int8-per-row quantization for `attn_gate_w` tensors. These are + # small ((num_heads, dim) = (8, 512) = 4096 params) and bypass GPTQ via the + # numel<=65536 passthrough branch -> stored as fp16 (8 KB/layer, ~65 KB total + # compressed). int8-per-row cuts the raw tensor in half with negligible BPB + # impact: scales per head (8 values), symmetric quant over [-127, 127]. + # No Hessian needed (gate weights not in collect_hessians()). + gated_attn_quant_gate = bool(int(os.environ.get("GATED_ATTN_QUANT_GATE", "1"))) + # Sparse Attention Gate (modded-nanogpt-style). Keeps dense SDPA and only + # swaps the output-gate input to the first GATE_WINDOW residual dims. + # W_g: (num_heads, gate_window) = (8, 12) = 96 params/layer (~44K total), + # vs dense GatedAttn's (8, 512) = 4K/layer (~44K diff). Name "attn_gate_w" + # is shared so quant routing and int8 gate passthrough Just Work. Gate + # passthrough int8 still applies via GATED_ATTN_QUANT_GATE=1. + # Mutually exclusive with ATTN_OUT_GATE_ENABLED and GATED_ATTN_ENABLED. + sparse_attn_gate_enabled = bool(int(os.environ.get("SPARSE_ATTN_GATE_ENABLED", "1"))) + sparse_attn_gate_init_std = float(os.environ.get("SPARSE_ATTN_GATE_INIT_STD", 0.0)) + sparse_attn_gate_scale = float(os.environ.get("SPARSE_ATTN_GATE_SCALE", 0.5)) + # LQER asymmetric rank-k correction on top-K quant-error tensors (PR #1530 v2 port). + # Computes SVD of E = W_fp - W_quant, packs top-r A,B as INT2/INT4 (asym) or INTk (sym). + lqer_enabled = bool(int(os.environ.get("LQER_ENABLED", "1"))) + 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")) + lqer_scope = os.environ.get("LQER_SCOPE", "all") + lqer_gain_select = bool(int(os.environ.get("LQER_GAIN_SELECT", "0"))) + awq_lite_enabled = bool(int(os.environ.get("AWQ_LITE_ENABLED", "1"))) + awq_lite_bits = int(os.environ.get("AWQ_LITE_BITS", "8")) + awq_lite_group_top_k = int(os.environ.get("AWQ_LITE_GROUP_TOP_K", "1")) + awq_lite_group_size = int(os.environ.get("AWQ_LITE_GROUP_SIZE", "64")) + 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 + # CaseOps integration: optional override of dataset root + tokenizer path. + # When CASEOPS_ENABLED=1, the wrapper loads a per-token byte sidecar + # (fineweb_val_bytes_*.bin, identical shard layout to val_*.bin) and uses + # it as the canonical raw-byte budget for BPB accounting. The sidecar + # REPLACES the build_sentencepiece_luts byte-counting path entirely. + caseops_enabled = bool(int(os.environ.get("CASEOPS_ENABLED", "1"))) + _default_caseops_data = os.path.join( + data_dir, + "datasets", + "fineweb10B_sp8192_caseops", + "datasets", + "datasets", + "fineweb10B_sp8192_lossless_caps_caseops_v1_reserved", + ) + _default_caseops_tok = os.path.join( + data_dir, + "datasets", + "fineweb10B_sp8192_caseops", + "datasets", + "tokenizers", + "fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model", + ) + if caseops_enabled: + datasets_dir = os.environ.get("DATA_PATH", _default_caseops_data) + tokenizer_path = os.environ.get("TOKENIZER_PATH", _default_caseops_tok) + else: + datasets_dir = os.environ.get( + "DATA_PATH", + os.path.join(data_dir, "datasets", f"fineweb10B_sp{vocab_size}"), + ) + tokenizer_path = os.environ.get( + "TOKENIZER_PATH", + os.path.join(data_dir, "tokenizers", f"fineweb_{vocab_size}_bpe.model"), + ) + train_files = os.path.join(datasets_dir, "fineweb_train_*.bin") + val_files = os.path.join(datasets_dir, "fineweb_val_*.bin") + val_bytes_files = os.path.join(datasets_dir, "fineweb_val_bytes_*.bin") + 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) + + +def parse_train_seq_schedule(schedule, default_seq_len): + if not schedule.strip(): + return [(1.0, int(default_seq_len))] + plan = [] + for raw_stage in schedule.split(","): + raw_stage = raw_stage.strip() + if not raw_stage: + continue + if "@" not in raw_stage: + raise ValueError( + f"Invalid TRAIN_SEQ_SCHEDULE stage `{raw_stage}`; expected format like `1024@0.35`" + ) + seq_raw, progress_raw = raw_stage.split("@", 1) + seq_len = int(seq_raw.strip()) + progress = float(progress_raw.strip()) + if seq_len <= 0: + raise ValueError("TRAIN_SEQ_SCHEDULE sequence lengths must be positive") + if not (0.0 < progress <= 1.0): + raise ValueError("TRAIN_SEQ_SCHEDULE progress fractions must be in (0, 1]") + plan.append((progress, seq_len)) + if not plan: + return [(1.0, int(default_seq_len))] + plan.sort(key=lambda item: item[0]) + if plan[-1][0] < 1.0: + plan.append((1.0, plan[-1][1])) + return plan + + +def parse_scalar_schedule(schedule, default_value): + if not schedule.strip(): + return [(0.0, float(default_value))] + plan = [] + for raw_stage in schedule.split(","): + raw_stage = raw_stage.strip() + if not raw_stage: + continue + if "@" not in raw_stage: + raise ValueError( + f"Invalid scalar schedule stage `{raw_stage}`; expected format like `0.5@0.4`" + ) + value_raw, progress_raw = raw_stage.split("@", 1) + value = float(value_raw.strip()) + progress = float(progress_raw.strip()) + if not (0.0 <= progress <= 1.0): + raise ValueError("Scalar schedule progress fractions must be in [0, 1]") + plan.append((progress, value)) + if not plan: + return [(0.0, float(default_value))] + plan.sort(key=lambda item: item[0]) + if plan[0][0] > 0.0: + plan.insert(0, (0.0, plan[0][1])) + return plan + + +def schedule_value(plan, progress): + value = plan[0][1] + for threshold, candidate in plan: + if progress + 1e-12 >= threshold: + value = candidate + else: + break + return value + + +def max_train_seq_len_from_schedule(plan, default_seq_len): + return max([int(default_seq_len), *[seq_len for _, seq_len in plan]]) + + +def validate_train_seq_plan_compatibility( + plan, + *, + global_tokens, + world_size, + grad_accum_steps, +): + denom = world_size * grad_accum_steps + if denom <= 0: + raise ValueError(f"Invalid world_size * grad_accum_steps={denom}") + if global_tokens % denom != 0: + raise ValueError( + f"TRAIN_BATCH_TOKENS={global_tokens} must be divisible by world_size*grad_accum_steps={denom}" + ) + local_tokens = global_tokens // denom + invalid_seq_lens = sorted( + {seq_len for _, seq_len in plan if local_tokens % seq_len != 0} + ) + if invalid_seq_lens: + raise ValueError( + "TRAIN_SEQ_SCHEDULE contains sequence lengths incompatible with the local micro-batch: " + f"local_tokens={local_tokens}, invalid_seq_lens={invalid_seq_lens}. " + f"Each seq_len must divide {local_tokens} exactly." + ) + return local_tokens + + +def training_progress( + *, + step, + iterations, + elapsed_ms, + max_wallclock_ms, + schedule_mode, +): + if schedule_mode == "step" or max_wallclock_ms is None or max_wallclock_ms <= 0: + return min(max(step / max(iterations, 1), 0.0), 1.0) + if schedule_mode != "wallclock": + raise ValueError( + f"Unsupported TRAIN_SEQ_SCHEDULE_MODE={schedule_mode!r}; expected 'wallclock' or 'step'" + ) + return min(max(elapsed_ms / max(max_wallclock_ms, 1e-9), 0.0), 1.0) + + +def current_train_seq_len( + plan, + *, + step, + iterations, + elapsed_ms, + max_wallclock_ms, + schedule_mode, +): + progress = training_progress( + step=step, + iterations=iterations, + elapsed_ms=elapsed_ms, + max_wallclock_ms=max_wallclock_ms, + schedule_mode=schedule_mode, + ) + for threshold, seq_len in plan: + if progress <= threshold: + return seq_len, progress + return plan[-1][1], progress + + +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, include_tail=h.eval_include_tail + ) + self.caseops_enabled = bool(getattr(h, "caseops_enabled", False)) + if self.caseops_enabled: + self.base_bytes_lut = None + self.has_leading_space_lut = None + self.is_boundary_token_lut = None + else: + ( + self.base_bytes_lut, + self.has_leading_space_lut, + self.is_boundary_token_lut, + ) = build_sentencepiece_luts(self.sp, h.vocab_size, device) + self.val_bytes = None + if self.caseops_enabled: + self.val_bytes = load_validation_byte_sidecar( + h.val_bytes_files, h.eval_seq_len, self.val_tokens.numel() + ) + + +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_) + for token_id in range(sp_vocab_size): + if sp.is_control(token_id) or sp.is_unknown(token_id) or sp.is_unused(token_id): + continue + is_boundary_token_np[token_id] = False + if sp.is_byte(token_id): + base_bytes_np[token_id] = 1 + continue + piece = sp.id_to_piece(token_id) + if piece.startswith("▁"): + has_leading_space_np[token_id] = True + piece = piece[1:] + base_bytes_np[token_id] = len(piece.encode("utf-8")) + return ( + torch.tensor(base_bytes_np, dtype=torch.int16, device=device), + torch.tensor(has_leading_space_np, dtype=torch.bool, device=device), + torch.tensor(is_boundary_token_np, dtype=torch.bool, device=device), + ) + + +def load_validation_tokens(pattern, seq_len, include_tail=True): + # Filter out CaseOps byte sidecar shards which share the val_*.bin glob. + files = [ + Path(p) + for p in sorted(glob.glob(pattern)) + if "_bytes_" not in Path(p).name + ] + if not files: + raise FileNotFoundError(f"No files found for pattern: {pattern}") + tokens = torch.cat([load_data_shard(file) for file in files]).contiguous() + if include_tail: + if tokens.numel() <= 1: + raise ValueError(f"Validation split is too short for TRAIN_SEQ_LEN={seq_len}") + return tokens + 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_validation_byte_sidecar(pattern, seq_len, expected_len): + """Load CaseOps per-token byte sidecar(s). Same shard layout as token shards + (256 int32 header + uint16 array). Each entry = canonical raw-text byte + budget for that token in the corresponding val shard. Returns a CPU + int16 tensor sliced to match expected_len (i.e. val_tokens length).""" + files = [Path(p) for p in sorted(glob.glob(pattern))] + if not files: + raise FileNotFoundError(f"No byte sidecar files for pattern: {pattern}") + shards = [load_data_shard(file) for file in files] + # load_data_shard returns uint16 — that's exactly what the sidecar stores. + bytes_full = torch.cat(shards).contiguous() + if bytes_full.numel() < expected_len: + raise ValueError( + f"Byte sidecar too short: {bytes_full.numel()} < val_tokens {expected_len}" + ) + return bytes_full[:expected_len].to(torch.int32) + + +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._prefetch_queue = [] + + 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() + ) + 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 = torch.empty(per_rank_span - 1, dtype=torch.int64, pin_memory=True) + targets = torch.empty(per_rank_span - 1, dtype=torch.int64, pin_memory=True) + inputs.copy_(buf[:-1]) + targets.copy_(buf[1:]) + 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, max_seq_len=None): + if max_seq_len is None: + max_seq_len = self.max_seq_len + max_seq_len = int(max_seq_len) + if max_seq_len != self.max_seq_len: + self.max_seq_len = max_seq_len + self._prefetch_queue.clear() + num_tokens_local = global_tokens // (self.world_size * grad_accum_steps) + while len(self._prefetch_queue) < 2: + self._prefetch_queue.append( + self._batch_pool.submit(self._prepare_batch, num_tokens_local, self.max_seq_len)) + inputs, targets, cu_seqlens, max_seqlen = self._prefetch_queue.pop(0).result() + self._prefetch_queue.append( + 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 = 256, 128, 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) + + +class CausalSelfAttention(nn.Module): + def __init__( + self, dim, num_heads, num_kv_heads, rope_base, qk_gain_init, train_seq_len, yarn=True, + attn_out_gate=False, attn_out_gate_src="proj", gate_window=12, + gated_attn=False, gated_attn_init_std=0.01, + sparse_attn_gate=False, sparse_attn_gate_init_std=0.0, sparse_attn_gate_scale=1.0, + ): + 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") + if int(attn_out_gate) + int(gated_attn) + int(sparse_attn_gate) > 1: + raise ValueError( + "attn_out_gate, gated_attn, and sparse_attn_gate are mutually exclusive" + ) + 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 + # AttnOutGate (PR #1667 MarioPaerle): per-head multiplicative gate on attention + # output. CastedLinear so restore_fp32_params casts back to fp32 for GPTQ. + # _zero_init -> 2*sigmoid(0)=1 -> transparent at init. + self.attn_out_gate = attn_out_gate + self.attn_out_gate_src = attn_out_gate_src + self.gate_window = gate_window + if attn_out_gate: + self.attn_gate_proj = CastedLinear(gate_window, num_heads, bias=False) + self.attn_gate_proj._zero_init = True + # Gated Attention (arXiv:2505.06708, Qwen, NeurIPS 2025). Per-head sigmoid + # gate on SDPA output, BEFORE out_proj. Gate projection W_g: (num_heads, dim). + # Name "attn_gate_w" contains "attn_gate" substring so it matches + # CONTROL_TENSOR_NAME_PATTERNS and routes to the scalar AdamW group. + # fp32 Parameter -> restore_fp32_params path covers it via the ndim<2 OR + # name-pattern check (name matches "attn_gate"). Cast to x.dtype on use. + self.gated_attn = gated_attn + if gated_attn: + W = torch.empty(num_heads, dim, dtype=torch.float32) + nn.init.normal_(W, mean=0.0, std=gated_attn_init_std) + self.attn_gate_w = nn.Parameter(W) + # Sparse attention head-output gate (modded-nanogpt style). Keeps dense SDPA + # and only narrows the gate input to the first gate_window residual dims. + # W_g: (num_heads, gate_window). y_{t,h} <- sigmoid(scale * W_g_h @ x_t[:gate_window]) * y_{t,h}. + # Shares attn_gate_w name with dense GatedAttn so the quant routing + # (CONTROL_TENSOR_NAME_PATTERNS / attn_gate_w int8 passthrough) is unchanged. + self.sparse_attn_gate = sparse_attn_gate + self.sparse_attn_gate_scale = sparse_attn_gate_scale + if sparse_attn_gate: + W = torch.empty(num_heads, gate_window, dtype=torch.float32) + if sparse_attn_gate_init_std > 0: + nn.init.normal_(W, mean=0.0, std=sparse_attn_gate_init_std) + else: + nn.init.zeros_(W) + self.attn_gate_w = nn.Parameter(W) + + 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_raw kept around as a tap point for attn_out_gate_src='q' (post-projection, + # pre-reshape, pre-RoPE). + q_raw = F.linear(x, q_w.to(x.dtype)) + q = q_raw.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_func( + q[0], + k[0], + v[0], + 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), + )[None] + else: + y = flash_attn_3_func(q, k, v, causal=True) + if self.use_xsa: + y = self._xsa_efficient(y, v) + # AttnOutGate inlined (PR #1667). Inline + .contiguous() barrier so torch.compile + # fullgraph=True is happy (this avoids the @torch.compiler.disable trap that + # crashed gates v3). Per-head gate on (B,T,H,D) tensor: g shape [B,T,H], broadcast + # over D via [..., None]. zero-init weight -> 2*sigmoid(0)=1 -> transparent. + if self.attn_out_gate: + gate_src = q_raw if self.attn_out_gate_src == "q" else x + gate_in = gate_src[..., : self.gate_window].contiguous() + g = 2.0 * torch.sigmoid(self.attn_gate_proj(gate_in)) + y = y * g[..., None] + # Gated Attention (arXiv:2505.06708 G1). Inline + .contiguous() barrier so + # torch.compile fullgraph=True is happy. Per-head gate on (B,T,H,D): g shape + # [B,T,H], broadcast over D via [..., None]. Paper: g = sigmoid(x @ W_g.T) + # where W_g: (H, dim). .to(x.dtype) on fp32 param before broadcast with bf16. + if self.gated_attn: + x_c = x.contiguous() + g = torch.sigmoid(F.linear(x_c, self.attn_gate_w.to(x.dtype))) + y = y * g[..., None] + # Sparse head-output gate: narrower (gate_window) input, same shape g as GatedAttn. + if 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[..., None] + 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 = True + + 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, + attn_out_gate=False, + attn_out_gate_src="proj", + gate_window=12, + gated_attn=False, + gated_attn_init_std=0.01, + sparse_attn_gate=False, + sparse_attn_gate_init_std=0.0, + sparse_attn_gate_scale=1.0, + ): + super().__init__() + self.attn_norm = RMSNorm() + self.mlp_norm = RMSNorm() + self.attn = CausalSelfAttention( + dim, num_heads, num_kv_heads, rope_base, qk_gain_init, train_seq_len, yarn=yarn, + attn_out_gate=attn_out_gate, attn_out_gate_src=attn_out_gate_src, gate_window=gate_window, + gated_attn=gated_attn, gated_attn_init_std=gated_attn_init_std, + sparse_attn_gate=sparse_attn_gate, + sparse_attn_gate_init_std=sparse_attn_gate_init_std, + sparse_attn_gate_scale=sparse_attn_gate_scale, + ) + 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.fused_ce_enabled = bool(h.fused_ce_enabled) + 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, + attn_out_gate=h.attn_out_gate_enabled, + attn_out_gate_src=h.attn_out_gate_src, + gate_window=h.gate_window, + gated_attn=h.gated_attn_enabled, + gated_attn_init_std=h.gated_attn_init_std, + sparse_attn_gate=h.sparse_attn_gate_enabled, + sparse_attn_gate_init_std=h.sparse_attn_gate_init_std, + sparse_attn_gate_scale=h.sparse_attn_gate_scale, + ) + 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 + 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) + ) + # SmearGate (PR #1667 / modded-nanogpt @classiclarryd): + # x_t <- x_t + lam * sigmoid(W * x_t[:gate_window]) * x_{t-1}. + # Per-token forward-1 smear of the embedding lane. W zero-init + lam=0 -> + # transparent at init. Uses CastedLinear so restore_fp32_params handles dtype. + self.smear_gate_enabled = h.smear_gate_enabled + if self.smear_gate_enabled: + self.smear_window = h.gate_window + self.smear_gate = CastedLinear(self.smear_window, 1, bias=False) + self.smear_gate._zero_init = True + self.smear_lambda = nn.Parameter(torch.zeros(1, dtype=torch.float32)) + # V19: Asymmetric Logit Rescale (PR #1923 jorge-asenjo). + # Two learnable softcap scales applied on the EVAL path (forward_logits + + # forward_ttt). Init to logit_softcap so the layer is identity at step 0. + # Train path keeps the single fused softcap to preserve PR #1855 numerics. + self.asym_logit_enabled = bool(int(os.environ.get("ASYM_LOGIT_RESCALE", "0"))) + if self.asym_logit_enabled: + self.softcap_pos = nn.Parameter(torch.tensor(float(h.logit_softcap), dtype=torch.float32)) + self.softcap_neg = nn.Parameter(torch.tensor(float(h.logit_softcap), 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_hidden(self, input_ids, cu_seqlens=None, max_seqlen=0): + """Run the encoder/decoder stack to the final RMSNorm; returns pre-projection hidden. + Shared by eval (softcap+projection via forward_logits) and train (fused CE path).""" + x = self.tok_emb(input_ids) + # SmearGate (PR #1667). lam=0 + W=0 -> identity at init. + # Cross-doc leak fix: zero the prev-token smear at any position whose current token + # is BOS, so the BOS embedding starting doc N+1 in a packed stream is not + # contaminated by doc N's last token (audited issue on PR#1797 base). + if self.smear_gate_enabled: + sl = self.smear_lambda.to(dtype=x.dtype) + gate_in = x[:, 1:, : self.smear_window].contiguous() + g = sl * torch.sigmoid(self.smear_gate(gate_in)) + not_bos = (input_ids[:, 1:] != BOS_ID).to(x.dtype).unsqueeze(-1) + x = torch.cat([x[:, :1], x[:, 1:] + g * x[:, :-1] * not_bos], 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) + return x + + def _project_logits(self, hidden): + if self.tie_embeddings: + return F.linear(hidden, self.tok_emb.weight) + return self.lm_head(hidden) + + def _apply_asym_softcap(self, logits): + # V19: Asymmetric softcap (PR #1923). Splits the logit_softcap scalar into + # learnable positive/negative branches. Score-first preserved: still a + # bounded, normalized post-projection nonlinearity feeding a standard + # softmax over the full vocab. + sp = self.softcap_pos.to(logits.dtype) + sn = self.softcap_neg.to(logits.dtype) + return torch.where(logits > 0, sp * torch.tanh(logits / sp), sn * torch.tanh(logits / sn)) + + def forward_logits(self, input_ids, cu_seqlens=None, max_seqlen=0): + hidden = self._forward_hidden(input_ids, cu_seqlens=cu_seqlens, max_seqlen=max_seqlen) + logits_proj = self._project_logits(hidden) + if self.asym_logit_enabled: + return self._apply_asym_softcap(logits_proj) + return self.logit_softcap * torch.tanh(logits_proj / self.logit_softcap) + + def forward(self, input_ids, target_ids, cu_seqlens=None, max_seqlen=0): + hidden = self._forward_hidden(input_ids, cu_seqlens=cu_seqlens, max_seqlen=max_seqlen) + logits_proj = self._project_logits(hidden) + flat_targets = target_ids.reshape(-1) + # Fused softcapped-CE kernel (training path only). Applies softcap inside the + # Triton kernel; takes pre-softcap logits_proj. Non-fused path matches stock + # PR-1736 numerics exactly (softcap in fp32, then F.cross_entropy on fp32). + if self.fused_ce_enabled: + return softcapped_cross_entropy( + logits_proj.reshape(-1, logits_proj.size(-1)), + flat_targets, + self.logit_softcap, + reduction="mean", + ) + logits = self.logit_softcap * torch.tanh(logits_proj / self.logit_softcap) + return F.cross_entropy( + logits.reshape(-1, logits.size(-1)).float(), + flat_targets, + reduction="mean", + ) + + def forward_ttt(self, input_ids, target_ids, lora, return_entropy: bool = False): + x = self.tok_emb(input_ids) + # SmearGate on the TTT path — same inline compute as forward_logits. + # Cross-doc leak fix: see _forward_hidden comment. + if self.smear_gate_enabled: + sl = self.smear_lambda.to(dtype=x.dtype) + gate_in = x[:, 1:, : self.smear_window].contiguous() + g = sl * torch.sigmoid(self.smear_gate(gate_in)) + not_bos = (input_ids[:, 1:] != BOS_ID).to(x.dtype).unsqueeze(-1) + x = torch.cat([x[:, :1], x[:, 1:] + g * x[:, :-1] * not_bos], 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) + # V19: same asymmetric softcap on the TTT eval path. + if self.asym_logit_enabled: + logits = self._apply_asym_softcap(logits) + else: + logits = self.logit_softcap * torch.tanh(logits / self.logit_softcap) + bsz, sl, V = logits.shape + logits_f = logits.float() + per_tok_loss = F.cross_entropy( + logits_f.reshape(-1, V), target_ids.reshape(-1), reduction="none" + ).reshape(bsz, sl) + if return_entropy: + log_probs = F.log_softmax(logits_f, dim=-1) + entropy = -(log_probs.exp() * log_probs).sum(dim=-1) + return per_tok_loss, entropy + return per_tok_loss + + 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 + # Keep raw Q for AttnOutGate src='q' (matches forward path semantics). + q_raw = F.linear(n, q_w.to(n.dtype)) + if lora.q_loras is not None: + q_raw = q_raw + lora.q_loras[slot](n) + q = q_raw.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)) + if lora.v_loras is not None: + v = v + lora.v_loras[slot](n) + v = v.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) + # AttnOutGate (TTT path) — inline + .contiguous() barrier, same as the eval path. + if attn.attn_out_gate: + gate_src = q_raw if attn.attn_out_gate_src == "q" else n + gate_in = gate_src[..., : attn.gate_window].contiguous() + g = 2.0 * torch.sigmoid(attn.attn_gate_proj(gate_in)) + y = y * g[..., None] + # Gated Attention (TTT path). Gate input is n (post-norm block input), same + # as eval path. .to(n.dtype) on fp32 param before bf16 broadcast. + if attn.gated_attn: + n_c = n.contiguous() + g = torch.sigmoid(F.linear(n_c, attn.attn_gate_w.to(n.dtype))) + y = y * g[..., None] + # Sparse attention head-output gate (TTT path) — must match the eval path in + # forward() exactly, else training (which applied the gate) and TTT eval (which + # skipped it) produce mismatched representations and catastrophic BPB regression. + if 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[..., None] + 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_raw = F.linear(n, q_w.to(n.dtype)) + if lora.q_loras is not None: + q_raw = q_raw + lora.q_loras[slot](n) + q = q_raw.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)) + if lora.v_loras is not None: + v = v + lora.v_loras[slot](n) + v = v.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) + # AttnOutGate (TTT parallel path) — inline + .contiguous() barrier. + if attn.attn_out_gate: + gate_src = q_raw if attn.attn_out_gate_src == "q" else n + gate_in = gate_src[..., : attn.gate_window].contiguous() + g = 2.0 * torch.sigmoid(attn.attn_gate_proj(gate_in)) + y = y * g[..., None] + # Gated Attention (TTT parallel path). Gate input is n (post-norm block input). + if attn.gated_attn: + n_c = n.contiguous() + g = torch.sigmoid(F.linear(n_c, attn.attn_gate_w.to(n.dtype))) + y = y * g[..., None] + # Sparse attention head-output gate (TTT parallel path) — must match the + # eval path in forward() to keep train/eval semantics in sync. + if 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[..., None] + 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")) + _WARM_START_A = bool(int(os.environ.get("TTT_WARM_START_A", "1"))) + PEER_IDX: "torch.Tensor | None" = None + + 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)) + + 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): + peer_idx = BatchedLinearLoRA.PEER_IDX + if peer_idx is None: + A, B = self.A, self.B + else: + A = self.A[peer_idx] + B = self.B[peer_idx] + return ((x @ A.transpose(1, 2)) @ B.transpose(1, 2)) * self._scale + + +class BatchedTTTLoRA(nn.Module): + def __init__( + self, bsz, model, rank, + q_lora=True, k_lora=True, v_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)] + ) + if q_lora + else None + ) + self.v_loras = ( + nn.ModuleList( + [BatchedLinearLoRA(bsz, dim, kv_dim, rank) for _ in range(num_slots)] + ) + if v_lora + else None + ) + 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). +# Replaces the fixed (3.4445, -4.775, 2.0315) coefficients of stock Muon. +# Applied at backend_steps=5 — taking more than 5 iterations from this list +# falls back to the final (converged) tuple via the slice guard below. +_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), +) + + +@torch.compile +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 = _PE_COEFFS[:steps] if steps <= len(_PE_COEFFS) else _PE_COEFFS + 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) + 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, 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, 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, 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,smear_gate,smear_lambda", + ).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) + # SmearGate params live on GPT root (not in .blocks), so add them by hand. + # Both are tiny (gate_window scalars + 1 lambda). Optimized via scalar Adam. + 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) + self._aux_stream = torch.cuda.Stream() + + 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._aux_stream.wait_stream(torch.cuda.current_stream()) + with torch.cuda.stream(self._aux_stream): + self.optimizer_tok.step() + self.optimizer_scalar.step() + self.optimizer_muon.step() + torch.cuda.current_stream().wait_stream(self._aux_stream) + 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 = {} + act_sumsq = {} + act_counts = {} + 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]) + x_sq = x.square().sum(dim=0) + x_count = x.shape[0] + 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) + if name not in act_sumsq: + act_sumsq[name] = torch.zeros( + x.shape[1], dtype=torch.float32, device=device + ) + act_counts[name] = 0 + act_sumsq[name] += x_sq + act_counts[name] += x_count + 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) + if name not in act_sumsq: + act_sumsq[name] = torch.zeros( + y.shape[1], dtype=torch.float32, device=device + ) + act_counts[name] = 0 + act_sumsq[name] += y.square().sum(dim=0) + act_counts[name] += y.shape[0] + 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) + if name not in act_sumsq: + act_sumsq[name] = torch.zeros( + x.shape[1], dtype=torch.float32, device=device + ) + act_counts[name] = 0 + act_sumsq[name] += x.square().sum(dim=0) + act_counts[name] += x.shape[0] + 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) + if name not in act_sumsq: + act_sumsq[name] = torch.zeros( + h_act.shape[1], dtype=torch.float32, device=device + ) + act_counts[name] = 0 + act_sumsq[name] += h_act.square().sum(dim=0) + act_counts[name] += h_act.shape[0] + 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) + if name not in act_sumsq: + act_sumsq[name] = torch.zeros( + x.shape[1], dtype=torch.float32, device=device + ) + act_counts[name] = 0 + act_sumsq[name] += x.square().sum(dim=0) + act_counts[name] += x.shape[0] + 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 + act_stats = {} + for name, sumsq in act_sumsq.items(): + count = max(act_counts.get(name, 0), 1) + act_stats[name] = (sumsq / count).sqrt().cpu() + return hessians, act_stats + + +def gptq_quantize_weight( + w, + H, + clip_sigmas=3.0, + clip_range=63, + block_size=128, + protect_groups=None, + group_size=None, + protect_clip_range=None, +): + 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() + protect_meta = None + protect_mask_perm = None + s_hi = None + sf_hi = None + if ( + protect_groups + and group_size is not None + and protect_clip_range is not None + and protect_clip_range > clip_range + ): + protect_mask = torch.zeros(cols, dtype=torch.bool) + starts = [] + for (start, end) in protect_groups: + if start < 0 or end > cols or end <= start: + continue + protect_mask[start:end] = True + starts.append(start) + if starts: + protect_mask_perm = protect_mask[perm] + s_hi = (clip_sigmas * row_std / protect_clip_range).clamp_min(1e-10).to( + torch.float16 + ) + sf_hi = s_hi.float() + protect_meta = { + "starts": torch.tensor(starts, dtype=torch.int16), + "size": int(group_size), + "s_hi": s_hi, + } + 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] + if protect_mask_perm is not None and bool(protect_mask_perm[i1 + j]): + q_col = torch.clamp( + torch.round(w_col / sf_hi), + -protect_clip_range, + protect_clip_range, + ) + w_recon = q_col.float() * sf_hi + else: + q_col = torch.clamp(torch.round(w_col / sf), -clip_range, clip_range) + w_recon = q_col.float() * sf + Q[:, i1 + j] = q_col.to(torch.int8) + err = (w_col - w_recon) / 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, protect_meta + + +def _quantize_gate_int8_row(w): + # Symmetric int8-per-row quantization for small gate tensors. w shape + # (R, C) -> (R,) scales in fp16, int8 values in [-127, 127]. Single scale + # per row keeps accuracy high while halving storage vs fp16. + W = w.float().contiguous() + row_max = W.abs().amax(dim=1).clamp_min(1e-10) + s = (row_max / 127.0).to(torch.float16) + sf = s.float().view(-1, 1) + q = torch.clamp(torch.round(W / sf), -127, 127).to(torch.int8) + return q, 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 g 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 _lqer_fit_quantized(E, h): + U, S, Vh = torch.linalg.svd(E, full_matrices=False) + r = min(h.lqer_rank, S.numel()) + if r <= 0: + return None + A = (U[:, :r] * S[:r]).contiguous() + B = Vh[:r, :].contiguous() + asym_on = bool(getattr(h, "lqer_asym_enabled", False)) + asym_g = int(getattr(h, "lqer_asym_group", 64)) + if asym_on and B.numel() % asym_g == 0: + qA, sA, qB, sB = _lqer_pack_asym(A, B, asym_g) + A_hat = qA.float() * float(sA) + g_sz = qB.numel() // sB.numel() + B_hat = (qB.reshape(-1, g_sz).float() * sB.float().view(-1, 1)).reshape( + qB.shape + ) + return { + "kind": "asym", + "qA": qA, + "sA": sA, + "qB": qB, + "sB": sB, + "delta": A_hat @ B_hat, + } + qA, sA, qB, sB = _lqer_pack(A, B, h.lqer_factor_bits) + A_hat = qA.float() * sA.float().view(-1, 1) + B_hat = qB.float() * sB.float().view(-1, 1) + return { + "kind": "sym", + "qA": qA, + "sA": sA, + "qB": qB, + "sB": sB, + "delta": A_hat @ B_hat, + } + + +def _awq_lite_group_candidates(w, act_rms, group_size): + cols = w.shape[1] + n_groups = cols // group_size + if n_groups <= 0: + return [] + weight_score = w.float().abs().mean(dim=0) + saliency = act_rms.float() * weight_score + cands = [] + for gi in range(n_groups): + start = gi * group_size + end = start + group_size + score = float(saliency[start:end].sum()) + cands.append((score, start, end)) + return cands + + +def gptq_mixed_quantize(state_dict, hessians, act_stats, h): + result = {} + meta = {} + quant_gate = bool(getattr(h, "gated_attn_quant_gate", False)) + lqer_on = bool(getattr(h, "lqer_enabled", False)) + awq_on = bool(getattr(h, "awq_lite_enabled", False)) + lqer_cands = {} + awq_selected = collections.defaultdict(list) + if awq_on: + awq_cands = [] + for (name, tensor) in state_dict.items(): + t = tensor.detach().cpu().contiguous() + if t.is_floating_point() and t.numel() > 65536 and name in act_stats: + bits = h.embed_bits if "tok_emb" in name else h.matrix_bits + if bits < h.awq_lite_bits: + for score, start, end in _awq_lite_group_candidates( + t, act_stats[name], h.awq_lite_group_size + ): + awq_cands.append((score, name, start, end)) + awq_cands.sort(key=lambda x: -x[0]) + for (_score, name, start, end) in awq_cands[: h.awq_lite_group_top_k]: + awq_selected[name].append((start, end)) + for (name, tensor) in state_dict.items(): + t = tensor.detach().cpu().contiguous() + # Dedicated int8-per-row path for attn_gate_w (bypasses both GPTQ and + # fp16 passthrough). Applied BEFORE the numel<=65536 passthrough check + # so the gate tensor is routed here instead of to fp16. + if ( + quant_gate + and t.is_floating_point() + and t.ndim == 2 + and name.endswith(".attn_gate_w") + # Dense GatedAttn: (num_heads, dim) = (8, 512) = 4096. + # Sparse gate: (num_heads, gate_window) = (8, 12) = 96. + # Both need int8-per-row routing; the 1024 lower bound in stock + # PR-1736 presumed dense-only. Widen to catch both. + and 32 <= t.numel() <= 8192 + ): + gq, gs = _quantize_gate_int8_row(t) + result[name + ".gq"] = gq + result[name + ".gs"] = gs + meta[name] = "gate_int8_row" + continue + 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 + q, s, protect_meta = gptq_quantize_weight( + t, + hessians[name], + clip_sigmas=cs, + clip_range=clip_range, + protect_groups=awq_selected.get(name), + group_size=h.awq_lite_group_size if name in awq_selected else None, + protect_clip_range=(2 ** (h.awq_lite_bits - 1) - 1) + if name in awq_selected + else None, + ) + result[name + ".q"] = q + result[name + ".scale"] = s + meta[name] = f"gptq (int{bits})" + W_q = q.float() * s.float().view(-1, 1) + if protect_meta is not None: + result[name + ".awqg_start"] = protect_meta["starts"] + result[name + ".awqg_s_hi"] = protect_meta["s_hi"] + result[name + ".awqg_size"] = torch.tensor( + protect_meta["size"], dtype=torch.int16 + ) + meta[name] = meta[name] + f"+awqgrpint{h.awq_lite_bits}" + gsz = protect_meta["size"] + for start in protect_meta["starts"].tolist(): + W_q[:, start : start + gsz] = ( + q[:, start : start + gsz].float() + * protect_meta["s_hi"].float().view(-1, 1) + ) + if lqer_on: + # LQER is fit on top of the fully realized GPTQ base, which already + # includes any higher-precision AWQ-protected groups. + scope = str(getattr(h, "lqer_scope", "all")).lower() + scope_ok = ( + scope == "all" + or (scope == "mlp" and ".mlp." in name) + or (scope == "attn" and ".attn." in name) + or (scope == "embed" and "tok_emb" in name) + ) + if scope_ok: + E = t.float() - W_q + err_norm = float(E.norm()) + if err_norm > 0: + lqer_cands[name] = (E, err_norm) + if lqer_on and lqer_cands: + if bool(getattr(h, "lqer_gain_select", False)): + scored = [] + for (name, (E, base_err)) in lqer_cands.items(): + fit = _lqer_fit_quantized(E, h) + if fit is None: + continue + new_err = float((E - fit["delta"]).norm()) + gain = base_err - new_err + if gain > 0: + scored.append((gain, name, fit)) + scored.sort(key=lambda x: -x[0]) + for (_gain, name, fit) in scored[: h.lqer_top_k]: + if fit["kind"] == "asym": + result[name + ".lqA_a"] = fit["qA"] + result[name + ".lqAs_a"] = fit["sA"] + result[name + ".lqB_a"] = fit["qB"] + result[name + ".lqBs_a"] = fit["sB"] + meta[name] = meta[name] + "+lqer_asym" + else: + result[name + ".lqA"] = fit["qA"] + result[name + ".lqAs"] = fit["sA"] + result[name + ".lqB"] = fit["qB"] + result[name + ".lqBs"] = fit["sB"] + meta[name] = meta[name] + "+lqer" + else: + top = sorted(lqer_cands.items(), key=lambda kv: -kv[1][1])[: h.lqer_top_k] + asym_on = bool(getattr(h, "lqer_asym_enabled", False)) + asym_g = int(getattr(h, "lqer_asym_group", 64)) + for (name, (E, _)) in top: + U, S, Vh = torch.linalg.svd(E, full_matrices=False) + r = min(h.lqer_rank, 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] + "+lqer_asym" + 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] + "+lqer" + 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 + if info == "gate_int8_row": + gq = result[name + ".gq"] + gs = result[name + ".gs"] + out[name] = (gq.float() * gs.float().view(-1, 1)).to(orig_dtype) + 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 "awqgrpint" in info: + starts = result[name + ".awqg_start"].tolist() + s_hi = result[name + ".awqg_s_hi"].float() + gsz = int(result[name + ".awqg_size"].item()) + for start in starts: + W[:, start : start + gsz] = ( + q[:, start : start + gsz].float() * s_hi.view(-1, 1) + ) + 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" + + +# ── Per-group lrzip compression (ported from PR#1586 via PR#1667/1729) ──────── + +_GROUP_ORDER = [ + "_tok_emb.weight.q", + "attn.c_k.weight.q", "attn.c_q.weight.q", + "attn.c_v.weight.q", "attn.proj.weight.q", + "mlp.fc.weight.q", "mlp.proj.weight.q", +] +_SIMSORT_KEYS = {"_tok_emb.weight.q", "attn.c_q.weight.q", "mlp.fc.weight.q"} +_PACK_MAGIC = b"PGRP" + + +def _similarity_sort_l1(matrix): + import numpy as _np + n = matrix.shape[0] + used = _np.zeros(n, dtype=bool) + order = [0] + used[0] = True + cur = matrix[0].astype(_np.float32) + for _ in range(n - 1): + dists = _np.sum(_np.abs(matrix[~used].astype(_np.float32) - cur), axis=1) + unused = _np.where(~used)[0] + best = unused[_np.argmin(dists)] + order.append(best) + used[best] = True + cur = matrix[best].astype(_np.float32) + return _np.array(order, dtype=_np.uint16) + + +def _lrzip_compress(data, tmpdir, label): + inp = os.path.join(tmpdir, f"{label}.bin") + out = f"{inp}.lrz" + with open(inp, "wb") as f: + f.write(data) + subprocess.run(["lrzip", "-z", "-L", "9", "-o", out, inp], capture_output=True, check=True) + with open(out, "rb") as f: + result = f.read() + os.remove(inp); os.remove(out) + return result + + +def _lrzip_decompress(data, tmpdir, label): + inp = os.path.join(tmpdir, f"{label}.lrz") + out = os.path.join(tmpdir, f"{label}.bin") + with open(inp, "wb") as f: + f.write(data) + subprocess.run(["lrzip", "-d", "-f", "-o", out, inp], capture_output=True, check=True) + with open(out, "rb") as f: + result = f.read() + os.remove(inp); os.remove(out) + return result + + +def _pack_streams(streams): + import struct + n = len(streams) + hdr = _PACK_MAGIC + struct.pack("= 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, + y_bytes=None, +): + 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) + if y_bytes is not None: + tok_bytes = y_bytes.to(torch.float64) + else: + 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 _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, forward_ttt_score=None): + 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) + target_tokens = sum(doc_len - 1 for _, doc_len in docs) + 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}" + f" target_tokens:{target_tokens}" + f" peer_k:{max(1, int(h.ttt_peer_ensemble_k))}" + f" conf_thresh:{h.ttt_peer_conf_threshold} blend_w:{h.ttt_peer_conf_blend_w}" + ) + chunk_size, eval_seq_len = h.ttt_chunk_size, h.ttt_eval_seq_len + + def _parse_short_score_first_steps(raw): + steps = [] + for item in str(raw).split(","): + item = item.strip() + if not item: + continue + if ":" in item: + doc_raw, chunk_raw = item.split(":", 1) + elif "=" in item: + doc_raw, chunk_raw = item.split("=", 1) + else: + raise ValueError( + "TTT_SHORT_SCORE_FIRST_STEPS must look like '256:16,512:24'" + ) + doc_len = int(doc_raw.strip()) + step_chunk = int(chunk_raw.strip()) + if doc_len <= 0 or step_chunk <= 0: + raise ValueError("TTT short score-first steps must be positive") + steps.append((doc_len, step_chunk)) + steps.sort(key=lambda x: x[0]) + return steps + + short_score_steps = _parse_short_score_first_steps( + h.ttt_short_score_first_steps + ) + + def _score_first_chunk_for_doc(max_doc_len): + if not h.ttt_short_score_first_enabled: + return chunk_size + if short_score_steps: + for doc_limit, step_chunk in short_score_steps: + if max_doc_len <= doc_limit: + return step_chunk + return chunk_size + if max_doc_len <= h.ttt_short_doc_len and h.ttt_short_chunk_size > 0: + return h.ttt_short_chunk_size + return chunk_size + + 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) + ens_loss_sum = torch.zeros((), device=device, dtype=torch.float64) + ens_byte_sum = torch.zeros((), device=device, dtype=torch.float64) + ens_token_count = torch.zeros((), device=device, dtype=torch.float64) + ens_route_token_count = torch.zeros((), device=device, dtype=torch.float64) + ens_batches_covered = 0 + peer_k = max(1, int(h.ttt_peer_ensemble_k)) + t_start = time.perf_counter() + reusable_lora = BatchedTTTLoRA( + h.ttt_batch_size, base_model, h.ttt_lora_rank, + q_lora=h.ttt_q_lora, k_lora=h.ttt_k_lora, v_lora=h.ttt_v_lora, + mlp_lora=h.ttt_mlp_lora, o_lora=h.ttt_o_lora, + ).to(device) + reusable_short_lora = None + reusable_short_opt = None + + def _build_opt(lora, lr=None, weight_decay=None, beta2=None): + lr = h.ttt_lora_lr if lr is None else lr + lr = lr * h.ttt_local_lr_mult + weight_decay = h.ttt_weight_decay if weight_decay is None else weight_decay + beta2 = h.ttt_beta2 if beta2 is None else beta2 + if h.ttt_optimizer == "sgd": + return torch.optim.SGD( + lora.parameters(), lr=lr, + momentum=h.ttt_beta1, weight_decay=weight_decay, + ) + return torch.optim.AdamW( + lora.parameters(), lr=lr, + betas=(h.ttt_beta1, beta2), + eps=1e-10, weight_decay=weight_decay, fused=True, + ) + + def _reset_optimizer_state(opt): + for s in opt.state.values(): + for k, v in s.items(): + if isinstance(v, torch.Tensor): + v.zero_() + elif k == "step": + s[k] = 0 + + def _apply_lora_template(lora, template): + if not template: + return False + with torch.no_grad(): + for name, p in lora.named_parameters(): + t = template.get(name) + if t is None or tuple(t.shape) != tuple(p.shape[1:]): + return False + for name, p in lora.named_parameters(): + t = template[name].to(device=p.device, dtype=p.dtype) + p.copy_(t.unsqueeze(0).expand_as(p)) + return True + + def _update_lora_template(template, lora): + momentum = float(h.ttt_warm_start_mean_momentum) + new_template = {} + with torch.no_grad(): + for name, p in lora.named_parameters(): + mean = p.detach().mean(dim=0).clone() + old = template.get(name) if template else None + if old is not None and tuple(old.shape) == tuple(mean.shape): + mean = old.to(device=mean.device, dtype=mean.dtype).mul(momentum).add( + mean, alpha=1.0 - momentum + ) + new_template[name] = mean + return new_template + + reusable_opt = _build_opt(reusable_lora) + warm_lora_template = None + 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) + doc_lens = [dl for _, dl in batch] + max_doc_len = max(doc_lens) + train_doc_allowed = [ + (h.ttt_train_min_doc_len <= 0 or dl >= h.ttt_train_min_doc_len) + and (h.ttt_train_max_doc_len <= 0 or dl <= h.ttt_train_max_doc_len) + for dl in doc_lens + ] + train_doc_mask_t = torch.tensor( + train_doc_allowed, dtype=torch.float32, device=device + ) + use_short_lora = h.ttt_short_lora_enabled and max_doc_len <= h.ttt_short_doc_len + batch_chunk_size = _score_first_chunk_for_doc(max_doc_len) + use_short_chunks = batch_chunk_size != chunk_size + batch_lora_rank = h.ttt_short_lora_rank if use_short_lora else h.ttt_lora_rank + batch_lora_lr = h.ttt_short_lora_lr if use_short_lora else h.ttt_lora_lr + batch_lora_wd = h.ttt_short_weight_decay if use_short_lora else h.ttt_weight_decay + batch_lora_beta2 = h.ttt_short_beta2 if use_short_lora else h.ttt_beta2 + prev_loss = loss_sum.item() + prev_bytes = byte_sum.item() + prev_tokens = token_count.item() + if use_short_lora and bsz == h.ttt_batch_size: + if reusable_short_lora is None: + reusable_short_lora = BatchedTTTLoRA( + h.ttt_batch_size, base_model, h.ttt_short_lora_rank, + q_lora=h.ttt_q_lora, k_lora=h.ttt_k_lora, v_lora=h.ttt_v_lora, + mlp_lora=h.ttt_mlp_lora, o_lora=h.ttt_o_lora, + ).to(device) + reusable_short_opt = _build_opt( + reusable_short_lora, + lr=h.ttt_short_lora_lr, + weight_decay=h.ttt_short_weight_decay, + beta2=h.ttt_short_beta2, + ) + reusable_short_lora.reset() + _reset_optimizer_state(reusable_short_opt) + cur_lora = reusable_short_lora + cur_opt = reusable_short_opt + elif (not use_short_lora) and bsz == reusable_lora.bsz: + reusable_lora.reset() + _reset_optimizer_state(reusable_opt) + cur_lora = reusable_lora + cur_opt = reusable_opt + else: + cur_lora = BatchedTTTLoRA( + bsz, base_model, batch_lora_rank, + q_lora=h.ttt_q_lora, k_lora=h.ttt_k_lora, v_lora=h.ttt_v_lora, + mlp_lora=h.ttt_mlp_lora, o_lora=h.ttt_o_lora, + ).to(device) + cur_opt = _build_opt( + cur_lora, + lr=batch_lora_lr, + weight_decay=batch_lora_wd, + beta2=batch_lora_beta2, + ) + template_used = False + if ( + h.ttt_warm_start_mean_enabled + and max_doc_len <= h.ttt_warm_start_mean_doc_len + ): + template_used = _apply_lora_template(cur_lora, warm_lora_template) + pred_lens = [doc_len - 1 for _, doc_len in batch] + num_chunks = [(pl + batch_chunk_size - 1) // batch_chunk_size for pl in pred_lens] + max_nc = max(num_chunks) + max_pred_len = max(pred_lens) if pred_lens else 0 + num_chunks_t = torch.tensor(num_chunks, dtype=torch.int64, device=device) + peer_enabled = peer_k > 1 and bsz >= peer_k + if peer_enabled and forward_ttt_score is None: + raise ValueError("peer ensemble requires forward_ttt_score") + peer_nll_stash = own_nll_stash = own_entropy_stash = None + if peer_enabled: + peer_nll_stash = torch.zeros( + bsz, max_pred_len, peer_k - 1, device=device, dtype=torch.float32, + ) + own_nll_stash = torch.zeros(bsz, max_pred_len, device=device, dtype=torch.float32) + own_entropy_stash = torch.zeros(bsz, max_pred_len, device=device, dtype=torch.float32) + for ci in range(max_nc): + active = [ci < nc for nc in num_chunks] + needs_train = any( + train_doc_allowed[b] and ci < nc - 1 + for b, nc in enumerate(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], batch_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) * batch_chunk_size, ci + 1, batch_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) + per_tok_entropy = None + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + if peer_enabled: + per_tok_loss, per_tok_entropy = forward_ttt_score(x, y, lora=cur_lora) + else: + per_tok_loss = forward_ttt_train(x, y, lora=cur_lora) + # CaseOps sidecar-driven byte budget. Mirror the index pattern + # used to build y from all_tokens: y[b, j] corresponds to the + # token at global position tok_starts[b] + 1 + j (when valid). + y_bytes_arg = None + if val_data.caseops_enabled and val_data.val_bytes is not None: + y_idx = ( + tok_starts.unsqueeze(1) + + 1 + + col_idx[:context_size].unsqueeze(0) + ) + y_idx = y_idx.clamp_(max=val_data.val_bytes.numel() - 1) + y_bytes_arg = val_data.val_bytes[y_idx].to( + device=device, dtype=torch.int32, non_blocking=True + ) + # Mirror the `valid` masking used for y so out-of-range tokens + # contribute zero bytes (matches y=0 substitution above). + y_bytes_arg = torch.where( + valid, y_bytes_arg, torch.zeros_like(y_bytes_arg) + ) + 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, + y_bytes=y_bytes_arg, + ) + if own_nll_stash is not None and per_tok_entropy is not None: + with torch.no_grad(): + pred_start = ci * batch_chunk_size + pred_end = min(pred_start + batch_chunk_size, max_pred_len) + stash_range = pred_end - pred_start + col_stash = torch.arange(stash_range, device=device) + gather_idx = chunk_offsets.unsqueeze(1) + col_stash.unsqueeze(0) + gather_idx = gather_idx.clamp_(max=per_tok_loss.size(1) - 1) + scored_slice = torch.gather(per_tok_loss.float(), 1, gather_idx) + valid_stash = (col_stash.unsqueeze(0) < chunk_lens.unsqueeze(1)).float() + own_nll_stash[:, pred_start:pred_end] = scored_slice * valid_stash + entropy_slice = torch.gather(per_tok_entropy.float(), 1, gather_idx) + own_entropy_stash[:, pred_start:pred_end] = entropy_slice * valid_stash + if needs_train: + activate_chunk_mask = (num_chunks_t - 1 > ci).float() * train_doc_mask_t + 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 + batch_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 + if peer_enabled: + doc_starts_cpu = torch.tensor([ds for ds, _ in batch], dtype=torch.int64) + rng = torch.Generator(device="cpu").manual_seed( + (h.seed * 1013 + orig_batch_idx) & 0x7FFFFFFF + ) + peer_ids = torch.empty(bsz, peer_k - 1, dtype=torch.int64) + for b in range(bsz): + others = torch.cat([torch.arange(b), torch.arange(b + 1, bsz)]) + perm = torch.randperm(bsz - 1, generator=rng)[: peer_k - 1] + peer_ids[b] = others[perm] + for peer_slot in range(peer_k - 1): + BatchedLinearLoRA.PEER_IDX = peer_ids[:, peer_slot].to(device) + for c_start in range(0, max_pred_len, eval_seq_len): + c_ctx = min(eval_seq_len, max_pred_len - c_start) + col = torch.arange(c_ctx + 1) + idx_cpu = doc_starts_cpu.unsqueeze(1) + c_start + col.unsqueeze(0) + idx_cpu.clamp_(max=all_tokens.numel() - 1) + gathered = all_tokens_idx[idx_cpu].to(device=device, dtype=torch.int64, non_blocking=True) + valid_len = torch.tensor( + [max(0, min(c_ctx, pl - c_start)) for pl in pred_lens], dtype=torch.int64, + ) + valid_mask = (col[:c_ctx].unsqueeze(0) < valid_len.unsqueeze(1)).to(device, non_blocking=True) + px = torch.where(valid_mask, gathered[:, :c_ctx], 0) + py = torch.where(valid_mask, gathered[:, 1 : c_ctx + 1], 0) + with torch.no_grad(): + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + p_per_tok_loss = forward_ttt_train(px, py, lora=cur_lora) + peer_nll_stash[:, c_start : c_start + c_ctx, peer_slot].copy_( + p_per_tok_loss.float() * valid_mask.float() + ) + BatchedLinearLoRA.PEER_IDX = None + with torch.no_grad(): + own_w = float(h.ttt_peer_conf_blend_w) + peer_mean_prob = torch.exp(-peer_nll_stash).mean(dim=-1) + own_prob = torch.exp(-own_nll_stash) + blended = own_w * own_prob + (1.0 - own_w) * peer_mean_prob + blended_nll = -torch.log(blended.clamp_min(1e-40)) + uncertain = own_entropy_stash >= float(h.ttt_peer_conf_threshold) + ens_nll = torch.where(uncertain, blended_nll, own_nll_stash) + pred_positions = torch.arange(max_pred_len, device=device) + x_idx_cpu = doc_starts_cpu.unsqueeze(1) + pred_positions.unsqueeze(0).cpu() + x_idx_cpu.clamp_(max=all_tokens.numel() - 1) + x_gathered = all_tokens_idx[x_idx_cpu].to(device=device, dtype=torch.int64, non_blocking=True) + y_idx = x_idx_cpu + 1 + y_idx.clamp_(max=all_tokens.numel() - 1) + y_gathered = all_tokens_idx[y_idx].to(device=device, dtype=torch.int64, non_blocking=True) + pred_lens_t = torch.tensor(pred_lens, dtype=torch.int64, device=device) + valid_pred = pred_positions.unsqueeze(0) < pred_lens_t.unsqueeze(1) + mask_f64 = valid_pred.to(torch.float64) + if val_data.caseops_enabled and val_data.val_bytes is not None: + tok_bytes = val_data.val_bytes[y_idx].to(device=device, dtype=torch.int32, non_blocking=True) + tok_bytes = torch.where(valid_pred, tok_bytes, torch.zeros_like(tok_bytes)).to(torch.float64) + else: + tok_bytes = val_data.base_bytes_lut[y_gathered].to(torch.float64) + tok_bytes += ( + val_data.has_leading_space_lut[y_gathered] + & ~val_data.is_boundary_token_lut[x_gathered] + ).to(torch.float64) + ens_loss_sum += (ens_nll.to(torch.float64) * mask_f64).sum() + ens_byte_sum += (tok_bytes * mask_f64).sum() + ens_token_count += mask_f64.sum() + ens_route_token_count += (uncertain.to(torch.float64) * mask_f64).sum() + ens_batches_covered += 1 + if h.ttt_warm_start_mean_enabled: + warm_lora_template = _update_lora_template(warm_lora_template, cur_lora) + batch_num = orig_batch_idx + 1 + 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)} sr:{int(use_short_lora)} " + f"sf:{int(use_short_chunks)} tr:{sum(train_doc_allowed)}/{bsz} " + f"wt:{int(template_used)}" + ) + 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) + if train_doc_allowed[pos] + ) + 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, + q_lora=h.ttt_q_lora, k_lora=h.ttt_k_lora, v_lora=h.ttt_v_lora, + mlp_lora=h.ttt_mlp_lora, o_lora=h.ttt_o_lora, + ).to(device) + reusable_opt = _build_opt(reusable_lora) + reusable_short_lora = None + reusable_short_opt = None + 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) + dist.all_reduce(ens_loss_sum, op=dist.ReduceOp.SUM) + dist.all_reduce(ens_byte_sum, op=dist.ReduceOp.SUM) + dist.all_reduce(ens_token_count, op=dist.ReduceOp.SUM) + dist.all_reduce(ens_route_token_count, op=dist.ReduceOp.SUM) + baseline_loss, baseline_bpb = _loss_bpb_from_sums(loss_sum, token_count, byte_sum) + log( + f"peer_ens:coverage_tokens:{ens_token_count.item():.0f}/{token_count.item():.0f} " + f"batches_covered:{ens_batches_covered}" + ) + if ens_token_count.item() > 0: + cov_loss, cov_bpb = _loss_bpb_from_sums(ens_loss_sum, ens_token_count, ens_byte_sum) + route_frac = (ens_route_token_count / ens_token_count).item() + log( + f"peer_ens:route_frac:{route_frac:.6f} " + f"routed_tokens:{ens_route_token_count.item():.0f}/{ens_token_count.item():.0f}" + ) + log(f"peer_ens:ensemble_covered_only val_loss:{cov_loss:.6f} val_bpb:{cov_bpb:.6f}") + log(f"peer_ens:baseline val_loss:{baseline_loss:.6f} val_bpb:{baseline_bpb:.6f}") + for p in base_model.parameters(): + p.requires_grad_(True) + base_model.train() + if ens_token_count.item() > 0: + return cov_loss, cov_bpb + return baseline_loss, baseline_bpb + + +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) + compiled_model = torch.compile(base_model, dynamic=False, fullgraph=True) + compiled_forward_logits = torch.compile( + base_model.forward_logits, dynamic=False, fullgraph=True + ) + 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) + train_seq_plan = parse_train_seq_schedule(h.train_seq_schedule, h.train_seq_len) + midrun_cap_plan = parse_scalar_schedule(h.midrun_cap_schedule, 1.0) + max_train_seq_len = max_train_seq_len_from_schedule(train_seq_plan, h.train_seq_len) + if max_train_seq_len != h.train_seq_len: + raise ValueError( + f"TRAIN_SEQ_LEN={h.train_seq_len} must match the maximum sequence length in " + f"TRAIN_SEQ_SCHEDULE ({max_train_seq_len})" + ) + local_microbatch_tokens = validate_train_seq_plan_compatibility( + train_seq_plan, + global_tokens=h.train_batch_tokens, + world_size=h.world_size, + grad_accum_steps=h.grad_accum_steps, + ) + log( + "train_seq_schedule:" + + ",".join((f"{seq_len}@{threshold:.3f}" for threshold, seq_len in train_seq_plan)) + ) + if h.midrun_cap_schedule: + log( + "midrun_cap_schedule:" + + ",".join( + (f"{value:.3f}@{threshold:.3f}" for threshold, value in midrun_cap_plan) + ) + ) + log(f"local_microbatch_tokens:{local_microbatch_tokens}") + active_train_seq_len = train_seq_plan[0][1] + seq_change_warmup_start_step = None + midrun_cap_active = False + midrun_cap_prev_scale = schedule_value(midrun_cap_plan, 0.0) + log(f"growth_stage:seq_len:{active_train_seq_len} progress:0.000") + 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(step, elapsed_ms, frac): + if h.warmdown_iters > 0: + if max_wallclock_ms is None: + warmdown_start = max(h.iterations - h.warmdown_iters, 0) + if warmdown_start <= step < h.iterations: + return max( + (h.iterations - step) / max(h.warmdown_iters, 1), + h.min_lr, + ) + return 1.0 + step_ms = elapsed_ms / max(step, 1) + warmdown_ms = h.warmdown_iters * step_ms + remaining_ms = max(max_wallclock_ms - elapsed_ms, 0.0) + if remaining_ms <= warmdown_ms: + return max(remaining_ms / max(warmdown_ms, 1e-9), h.min_lr) + return 1.0 + 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 + + _clip_params = [p for p in base_model.parameters() if p.requires_grad] + def step_fn(step, lr_scale): + 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, active_train_seq_len + ) + with torch.autocast(device_type="cuda", dtype=torch.bfloat16, enabled=True): + loss = model( + x, y, cu_seqlens=cu_seqlens, max_seqlen=active_train_seq_len + ) + train_loss += loss.detach() + (loss / h.grad_accum_steps).backward() + train_loss /= h.grad_accum_steps + if step <= h.muon_momentum_warmup_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_(_clip_params, 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, active_train_seq_len + ) + log(f"warmup_cu_buckets:{','.join(str(b) for b in warmup_cu_buckets)} iters_each:{warmup_cu_iters}") + + def _compile_warmup_work_items(): + if not h.compile_shape_warmup: + items = [(h.train_seq_len, False)] + if h.num_loops > 0: + items.append((h.train_seq_len, True)) + return items + loop_mode = h.compile_shape_warmup_loop_modes + if loop_mode not in {"auto", "inactive", "active", "both"}: + raise ValueError( + "COMPILE_SHAPE_WARMUP_LOOP_MODES must be one of auto,inactive,active,both" + ) + items = [] + stage_start = 0.0 + for stage_end, seq_len in train_seq_plan: + if loop_mode == "inactive" or h.num_loops <= 0: + modes = [False] + elif loop_mode == "active": + modes = [True] + elif loop_mode == "both": + modes = [False, True] + else: + modes = [] + if stage_start < h.enable_looping_at: + modes.append(False) + if stage_end > h.enable_looping_at: + modes.append(True) + if not modes: + modes.append(False) + for loop_active in modes: + item = (seq_len, loop_active) + if item not in items: + items.append(item) + stage_start = stage_end + return items + + def _run_cu_bucket_warmup(seq_len): + for bucket_len in warmup_cu_buckets: + boundaries = list(range(0, x.size(1), max(seq_len, 1))) + if boundaries[-1] != x.size(1): + boundaries.append(x.size(1)) + if bucket_len < len(boundaries): + continue + 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=seq_len) + (wloss / h.grad_accum_steps).backward() + optimizers.zero_grad_all() + + warmup_items = _compile_warmup_work_items() + if h.compile_shape_warmup: + log( + "compile_shape_warmup:start " + + ",".join( + (f"{seq}x{'loop' if loop else 'plain'}" for seq, loop in warmup_items) + ) + ) + for seq_len, loop_active in warmup_items: + base_model.looping_active = bool(loop_active) + if h.compile_shape_warmup: + log( + f"compile_shape_warmup:shape seq_len:{seq_len} loop:{int(loop_active)}" + ) + for _ in range(max(h.compile_shape_warmup_iters if h.compile_shape_warmup else 1, 1)): + _run_cu_bucket_warmup(seq_len) + 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) + _live_state = base_model.state_dict(keep_vars=True) + ema_state = { + name: t.detach().float().clone() + for (name, t) in _live_state.items() + } + _ema_pairs = [(ema_state[name], t) for (name, t) in _live_state.items()] + ema_decay = h.ema_decay + training_time_ms = 0.0 + forced_stop_step = int(os.environ.get("FORCE_STOP_STEP", "0")) + stop_after_step = forced_stop_step if forced_stop_step > 0 else 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) + stage_seq_len, frac = current_train_seq_len( + train_seq_plan, + step=step, + iterations=h.iterations, + elapsed_ms=elapsed_ms, + max_wallclock_ms=max_wallclock_ms, + schedule_mode=h.train_seq_schedule_mode, + ) + if stage_seq_len != active_train_seq_len: + active_train_seq_len = stage_seq_len + log(f"growth_stage:seq_len:{active_train_seq_len} progress:{frac:.3f} step:{step}") + if h.seq_change_warmup_steps > 0 and step > 0: + seq_change_warmup_start_step = step + log( + f"growth_stage_rewarmup:start step:{step} steps:{h.seq_change_warmup_steps} " + f"seq_len:{active_train_seq_len}" + ) + scale = lr_mul(step, elapsed_ms, frac) + cap_scale = schedule_value(midrun_cap_plan, frac) + cap_active = cap_scale < 0.999999 + if cap_active and not midrun_cap_active: + log(f"midrun_cap:start step:{step} progress:{frac:.3f} scale:{cap_scale:.3f}") + elif ( + cap_active + and h.midrun_cap_log_updates + and abs(cap_scale - midrun_cap_prev_scale) > 1e-6 + ): + log(f"midrun_cap:update step:{step} progress:{frac:.3f} scale:{cap_scale:.3f}") + if cap_active: + scale *= cap_scale + midrun_cap_active = cap_active + midrun_cap_prev_scale = cap_scale + if seq_change_warmup_start_step is not None and h.seq_change_warmup_steps > 0: + rewarm_progress = min( + max( + (step - seq_change_warmup_start_step + 1) + / max(h.seq_change_warmup_steps, 1), + 0.0, + ), + 1.0, + ) + scale *= rewarm_progress + if rewarm_progress >= 1.0: + seq_change_warmup_start_step = None + 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 ema_t, t in _ema_pairs: + ema_t.mul_(ema_decay).add_(t.detach(), 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 = ( + forced_stop_step <= 0 + and max_wallclock_ms is not None + and approx_training_time_ms >= max_wallclock_ms + ) + if h.distributed and forced_stop_step <= 0 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" + ) + if h.ema_decay <= 0: + log("averaging:none keeping current weights") + return base_model, compiled_model, compiled_forward_logits + 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): + global BOS_ID + 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: skip training + GPTQ, jump straight to TTT eval on a + # pre-existing quantized artifact. Used to test TTT-only improvements + # (e.g., PR-1767's alpha/warm-start/WD) without retraining. + ttt_eval_only = os.environ.get("TTT_EVAL_ONLY", "0") == "1" + quantize_only = os.environ.get("QUANTIZE_ONLY", "0") == "1" + if ttt_eval_only: + log("TTT_EVAL_ONLY=1 — skipping training + GPTQ, loading saved artifact for TTT eval") + log(f"ttt_lora_alpha: {BatchedLinearLoRA._ALPHA}") + log(f"ttt_warm_start_a: {BatchedLinearLoRA._WARM_START_A}") + log(f"ttt_weight_decay: {h.ttt_weight_decay}") + elif quantize_only: + log("QUANTIZE_ONLY=1 — skipping training, loading saved full-precision checkpoint") + log(f"quantize_only checkpoint: {h.model_path}") + if BOS_ID is None: + BOS_ID = 1 + base_model = GPT(h).to(device).bfloat16() + state = torch.load(h.model_path, map_location="cpu") + template = base_model.state_dict() + for key in ("softcap_pos", "softcap_neg"): + if key not in state and key in template: + state[key] = template[key].detach().cpu().clone() + log(f"quantize_only:added neutral missing {key}") + base_model.load_state_dict(state, strict=True) + del state + serialize(h, base_model, Path(__file__).read_text(encoding="utf-8")) + if h.distributed: + dist.barrier() + 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, + ) + if os.environ.get("PREQUANT_ONLY", "0") == "1": + log("PREQUANT_ONLY=1 — skipping serialize/GPTQ/post-quant eval/TTT") + return + 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 + if not ttt_eval_only: + compiled_model = torch.compile(eval_model, dynamic=False, fullgraph=True) + compiled_forward_logits = torch.compile( + eval_model.forward_logits, dynamic=False, fullgraph=True + ) + timed_eval( + "diagnostic quantized", + eval_val, + h, + device, + val_data, + compiled_model, + compiled_forward_logits, + ) + del eval_model + if h.ttt_enabled: + if not ttt_eval_only: + del compiled_model + if ttt_eval_only: + del eval_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) + + _fwd_ttt_compiled_inner = None + + def _fwd_ttt(input_ids, target_ids, lora): + nonlocal _fwd_ttt_compiled_inner + if _fwd_ttt_compiled_inner is None: + _fwd_ttt_compiled_inner = torch.compile(_fwd_ttt_inner, dynamic=True) + return _fwd_ttt_compiled_inner(input_ids, target_ids, lora=lora) + + fwd_ttt_compiled = _fwd_ttt + + def _fwd_ttt_score_inner(input_ids, target_ids, lora): + return ttt_model.forward_ttt(input_ids, target_ids, lora=lora, return_entropy=True) + + _fwd_ttt_score_compiled_inner = None + + def _fwd_ttt_score(input_ids, target_ids, lora): + nonlocal _fwd_ttt_score_compiled_inner + if _fwd_ttt_score_compiled_inner is None: + _fwd_ttt_score_compiled_inner = torch.compile(_fwd_ttt_score_inner, dynamic=True) + return _fwd_ttt_score_compiled_inner(input_ids, target_ids, lora=lora) + + fwd_ttt_score_compiled = _fwd_ttt_score + log(f"ttt_lora:warming up compile (random tokens, no val data)") + 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, + q_lora=h.ttt_q_lora, k_lora=h.ttt_k_lora, v_lora=h.ttt_v_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 * h.ttt_local_lr_mult, + betas=(h.ttt_beta1, h.ttt_beta2), + eps=1e-10, + weight_decay=h.ttt_weight_decay, + fused=True, + ) + warmup_ctx_lens = [h.ttt_chunk_size, h.ttt_eval_seq_len] + if ( + h.ttt_short_score_first_enabled + and h.ttt_short_chunk_size not in warmup_ctx_lens + ): + warmup_ctx_lens.insert(0, h.ttt_short_chunk_size) + for item in str(h.ttt_short_score_first_steps).split(","): + item = item.strip() + if not item: + continue + sep = ":" if ":" in item else "=" + if sep not in item: + continue + _, chunk_raw = item.split(sep, 1) + step_chunk = int(chunk_raw.strip()) + if step_chunk > 0 and step_chunk not in warmup_ctx_lens: + warmup_ctx_lens.insert(0, step_chunk) + for ctx_len in 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) + if int(h.ttt_peer_ensemble_k) > 1: + for ctx_len in (h.ttt_chunk_size, h.ttt_eval_seq_len): + 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, _entropy = fwd_ttt_score_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) + BatchedLinearLoRA.PEER_IDX = torch.randint(0, bsz, (bsz,), dtype=torch.int64, device=device) + for ctx_len in (h.ttt_eval_seq_len,): + 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.no_grad(): + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + _ = fwd_ttt_compiled(xw, yw, lora=wl) + BatchedLinearLoRA.PEER_IDX = None + 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, + forward_ttt_score=fwd_ttt_score_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) + 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 + dynamo_cache_size_limit = int(os.environ.get("DYNAMO_CACHE_SIZE_LIMIT", "128")) + torch._dynamo.config.cache_size_limit = dynamo_cache_size_limit + if hasattr(torch._dynamo.config, "recompile_limit"): + torch._dynamo.config.recompile_limit = dynamo_cache_size_limit + if hasattr(torch._dynamo.config, "accumulated_cache_size_limit"): + torch._dynamo.config.accumulated_cache_size_limit = max( + int(os.environ.get("DYNAMO_ACCUMULATED_CACHE_SIZE_LIMIT", "1024")), + dynamo_cache_size_limit, + ) + 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() + +==================================================================================================== +Running Python 3.12.13 (main, Mar 10 2026, 18:17:25) [Clang 21.1.4 ] +Running PyTorch 2.10.0+cu130 +==================================================================================================== +train_shards: 80 +val_tokens: 47853343 +model_params:35945671 +train_seq_schedule:1024@0.100,2048@0.700,3072@1.000 +local_microbatch_tokens:98304 +growth_stage:seq_len:1024 progress:0.000 +gptq:reserving 4s, effective=596000ms +warmup_cu_buckets:64,128,192,256 iters_each:3 +compile_shape_warmup:start 1024xplain,2048xplain,2048xloop,3072xloop +compile_shape_warmup:shape seq_len:1024 loop:0 +compile_shape_warmup:shape seq_len:2048 loop:0 +compile_shape_warmup:shape seq_len:2048 loop:1 +compile_shape_warmup:shape seq_len:3072 loop:1 +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 +1/20000 train_loss: 9.0087 train_time: 0.0m tok/s: 17923690 +2/20000 train_loss: 12.8466 train_time: 0.0m tok/s: 12011643 +3/20000 train_loss: 10.2392 train_time: 0.0m tok/s: 10807885 +4/20000 train_loss: 8.6995 train_time: 0.0m tok/s: 10272331 +5/20000 train_loss: 7.9235 train_time: 0.0m tok/s: 9952667 +500/20000 train_loss: 2.6066 train_time: 0.7m tok/s: 8796080 +growth_stage:seq_len:2048 progress:0.100 step:665 +growth_stage_rewarmup:start step:665 steps:32 seq_len:2048 +1000/20000 train_loss: 2.5935 train_time: 1.5m tok/s: 8637671 +1500/20000 train_loss: 2.6242 train_time: 2.3m tok/s: 8542457 +2000/20000 train_loss: 2.6572 train_time: 3.1m tok/s: 8493636 +layer_loop:enabled step:2248 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.5128 train_time: 4.1m tok/s: 8076431 +3000/20000 train_loss: 2.4615 train_time: 5.2m tok/s: 7550238 +3500/20000 train_loss: 2.4710 train_time: 6.4m tok/s: 7214124 +growth_stage:seq_len:3072 progress:0.700 step:3758 +growth_stage_rewarmup:start step:3758 steps:32 seq_len:3072 +4000/20000 train_loss: 2.3917 train_time: 7.5m tok/s: 6962902 +4500/20000 train_loss: 2.3590 train_time: 8.7m tok/s: 6761322 +5000/20000 train_loss: 2.2322 train_time: 9.9m tok/s: 6608662 +5006/20000 val_loss: 2.3422 val_bpb: 1.0703 +stopping_early: wallclock_cap train_time: 595913ms step: 5006/20000 +peak memory allocated: 41710 MiB reserved: 46988 MiB +ema:applying EMA weights +diagnostic pre-quantization post-ema val_loss:2.31746801 val_bpb:1.05898971 eval_time:15789ms +Serialized model: 135417533 bytes +Code size (uncompressed): 201974 bytes +Code size (compressed): 39907 bytes +GPTQ:collecting Hessians from calibration data... +GPTQ:collected 67 Hessians in 3.9s +Quantized weights: + gate_int8_row: blocks.attn.attn_gate_w + gptq (int6): blocks.attn.c_k.weight, blocks.attn.c_q.weight, blocks.attn.c_v.weight, blocks.attn.proj.weight, blocks.mlp.fc.weight, blocks.mlp.proj.weight + gptq (int6)+lqer_asym: blocks.mlp.fc.weight + gptq (int7)+awqgrpint8+lqer_asym: tok_emb.weight + passthrough (float16): 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 +Serialize: per-group lrzip compression... +Serialize: per-group compression done in 109.6s +Serialized model quantized+pergroup: 15946917 bytes +Total submission size quantized+pergroup: 15986824 bytes +Deserialize: per-group lrzip decompression... +Deserialize: decompression done in 20.2s +diagnostic quantized val_loss:2.33619241 val_bpb:1.06754601 eval_time:23419ms +Deserialize: per-group lrzip decompression... +Deserialize: decompression done in 20.2s +ttt_lora:warming up compile (random tokens, no val data) +ttt_lora:compile warmup done (260.5s) + +beginning TTT eval timer +ttt_phased: total_docs:50000 prefix_docs:2500 suffix_docs:47500 num_phases:1 boundaries:[2500] target_tokens:47853343 peer_k:4 conf_thresh:0.5 blend_w:0.8 +ttp: b2084/2084 bl:2.1405 bb:1.0238 rl:2.1405 rb:1.0238 dl:36899-97114 gd:0 sr:0 sf:0 tr:8/8 wt:0 +ttp: b2017/2084 bl:2.2713 bb:1.0413 rl:2.1589 rb:1.0263 dl:3550-3581 gd:0 sr:0 sf:0 tr:24/24 wt:0 +ttpp: phase:1/1 pd:2672 gd:2500 t:291.4s +tttg: c1/344 lr:0.001000 t:0.3s +tttg: c2/344 lr:0.001000 t:0.3s +tttg: c3/344 lr:0.001000 t:0.4s +tttg: c4/344 lr:0.001000 t:0.5s +tttg: c5/344 lr:0.001000 t:0.6s +tttg: c6/344 lr:0.000999 t:0.6s +tttg: c7/344 lr:0.000999 t:0.7s +tttg: c8/344 lr:0.000999 t:0.8s +tttg: c9/344 lr:0.000999 t:0.9s +tttg: c10/344 lr:0.000998 t:0.9s +tttg: c11/344 lr:0.000998 t:1.0s +tttg: c12/344 lr:0.000997 t:1.1s +tttg: c13/344 lr:0.000997 t:1.2s +tttg: c14/344 lr:0.000996 t:1.2s +tttg: c15/344 lr:0.000996 t:1.3s +tttg: c16/344 lr:0.000995 t:1.4s +tttg: c17/344 lr:0.000995 t:1.4s +tttg: c18/344 lr:0.000994 t:1.5s +tttg: c19/344 lr:0.000993 t:1.6s +tttg: c20/344 lr:0.000992 t:1.7s +tttg: c21/344 lr:0.000992 t:1.7s +tttg: c22/344 lr:0.000991 t:1.8s +tttg: c23/344 lr:0.000990 t:1.9s +tttg: c24/344 lr:0.000989 t:2.0s +tttg: c25/344 lr:0.000988 t:2.0s +tttg: c26/344 lr:0.000987 t:2.1s +tttg: c27/344 lr:0.000986 t:2.2s +tttg: c28/344 lr:0.000985 t:2.2s +tttg: c29/344 lr:0.000984 t:2.3s +tttg: c30/344 lr:0.000982 t:2.4s +tttg: c31/344 lr:0.000981 t:2.5s +tttg: c32/344 lr:0.000980 t:2.5s +tttg: c33/344 lr:0.000979 t:2.6s +tttg: c34/344 lr:0.000977 t:2.7s +tttg: c35/344 lr:0.000976 t:2.7s +tttg: c36/344 lr:0.000975 t:2.8s +tttg: c37/344 lr:0.000973 t:2.9s +tttg: c38/344 lr:0.000972 t:3.0s +tttg: c39/344 lr:0.000970 t:3.0s +tttg: c40/344 lr:0.000968 t:3.1s +tttg: c41/344 lr:0.000967 t:3.2s +tttg: c42/344 lr:0.000965 t:3.3s +tttg: c43/344 lr:0.000963 t:3.3s +tttg: c44/344 lr:0.000962 t:3.4s +tttg: c45/344 lr:0.000960 t:3.5s +tttg: c46/344 lr:0.000958 t:3.5s +tttg: c47/344 lr:0.000956 t:3.6s +tttg: c48/344 lr:0.000954 t:3.7s +tttg: c49/344 lr:0.000952 t:3.8s +tttg: c50/344 lr:0.000950 t:3.8s +tttg: c51/344 lr:0.000948 t:3.9s +tttg: c52/344 lr:0.000946 t:4.0s +tttg: c53/344 lr:0.000944 t:4.0s +tttg: c54/344 lr:0.000942 t:4.1s +tttg: c55/344 lr:0.000940 t:4.2s +tttg: c56/344 lr:0.000938 t:4.3s +tttg: c57/344 lr:0.000936 t:4.3s +tttg: c58/344 lr:0.000933 t:4.4s +tttg: c59/344 lr:0.000931 t:4.5s +tttg: c60/344 lr:0.000929 t:4.5s +tttg: c61/344 lr:0.000926 t:4.6s +tttg: c62/344 lr:0.000924 t:4.7s +tttg: c63/344 lr:0.000922 t:4.8s +tttg: c64/344 lr:0.000919 t:4.8s +tttg: c65/344 lr:0.000917 t:4.9s +tttg: c66/344 lr:0.000914 t:5.0s +tttg: c67/344 lr:0.000911 t:5.1s +tttg: c68/344 lr:0.000909 t:5.1s +tttg: c69/344 lr:0.000906 t:5.2s +tttg: c70/344 lr:0.000903 t:5.3s +tttg: c71/344 lr:0.000901 t:5.3s +tttg: c72/344 lr:0.000898 t:5.4s +tttg: c73/344 lr:0.000895 t:5.5s +tttg: c74/344 lr:0.000892 t:5.6s +tttg: c75/344 lr:0.000889 t:5.6s +tttg: c76/344 lr:0.000887 t:5.7s +tttg: c77/344 lr:0.000884 t:5.8s +tttg: c78/344 lr:0.000881 t:5.8s +tttg: c79/344 lr:0.000878 t:5.9s +tttg: c80/344 lr:0.000875 t:6.0s +tttg: c81/344 lr:0.000872 t:6.1s +tttg: c82/344 lr:0.000869 t:6.1s +tttg: c83/344 lr:0.000865 t:6.2s +tttg: c84/344 lr:0.000862 t:6.3s +tttg: c85/344 lr:0.000859 t:6.4s +tttg: c86/344 lr:0.000856 t:6.4s +tttg: c87/344 lr:0.000853 t:6.5s +tttg: c88/344 lr:0.000849 t:6.6s +tttg: c89/344 lr:0.000846 t:6.6s +tttg: c90/344 lr:0.000843 t:6.7s +tttg: c91/344 lr:0.000840 t:6.8s +tttg: c92/344 lr:0.000836 t:6.9s +tttg: c93/344 lr:0.000833 t:6.9s +tttg: c94/344 lr:0.000829 t:7.0s +tttg: c95/344 lr:0.000826 t:7.1s +tttg: c96/344 lr:0.000822 t:7.1s +tttg: c97/344 lr:0.000819 t:7.2s +tttg: c98/344 lr:0.000815 t:7.3s +tttg: c99/344 lr:0.000812 t:7.4s +tttg: c100/344 lr:0.000808 t:7.4s +tttg: c101/344 lr:0.000805 t:7.5s +tttg: c102/344 lr:0.000801 t:7.6s +tttg: c103/344 lr:0.000797 t:7.6s +tttg: c104/344 lr:0.000794 t:7.7s +tttg: c105/344 lr:0.000790 t:7.8s +tttg: c106/344 lr:0.000786 t:7.9s +tttg: c107/344 lr:0.000782 t:7.9s +tttg: c108/344 lr:0.000778 t:8.0s +tttg: c109/344 lr:0.000775 t:8.1s +tttg: c110/344 lr:0.000771 t:8.2s +tttg: c111/344 lr:0.000767 t:8.2s +tttg: c112/344 lr:0.000763 t:8.3s +tttg: c113/344 lr:0.000759 t:8.4s +tttg: c114/344 lr:0.000755 t:8.4s +tttg: c115/344 lr:0.000751 t:8.5s +tttg: c116/344 lr:0.000747 t:8.6s +tttg: c117/344 lr:0.000743 t:8.7s +tttg: c118/344 lr:0.000739 t:8.7s +tttg: c119/344 lr:0.000735 t:8.8s +tttg: c120/344 lr:0.000731 t:8.9s +tttg: c121/344 lr:0.000727 t:8.9s +tttg: c122/344 lr:0.000723 t:9.0s +tttg: c123/344 lr:0.000719 t:9.1s +tttg: c124/344 lr:0.000715 t:9.2s +tttg: c125/344 lr:0.000711 t:9.2s +tttg: c126/344 lr:0.000707 t:9.3s +tttg: c127/344 lr:0.000702 t:9.4s +tttg: c128/344 lr:0.000698 t:9.5s +tttg: c129/344 lr:0.000694 t:9.5s +tttg: c130/344 lr:0.000690 t:9.6s +tttg: c131/344 lr:0.000686 t:9.7s +tttg: c132/344 lr:0.000681 t:9.7s +tttg: c133/344 lr:0.000677 t:9.8s +tttg: c134/344 lr:0.000673 t:9.9s +tttg: c135/344 lr:0.000668 t:10.0s +tttg: c136/344 lr:0.000664 t:10.0s +tttg: c137/344 lr:0.000660 t:10.1s +tttg: c138/344 lr:0.000655 t:10.2s +tttg: c139/344 lr:0.000651 t:10.3s +tttg: c140/344 lr:0.000647 t:10.3s +tttg: c141/344 lr:0.000642 t:10.4s +tttg: c142/344 lr:0.000638 t:10.5s +tttg: c143/344 lr:0.000633 t:10.5s +tttg: c144/344 lr:0.000629 t:10.6s +tttg: c145/344 lr:0.000625 t:10.7s +tttg: c146/344 lr:0.000620 t:10.8s +tttg: c147/344 lr:0.000616 t:10.8s +tttg: c148/344 lr:0.000611 t:10.9s +tttg: c149/344 lr:0.000607 t:11.0s +tttg: c150/344 lr:0.000602 t:11.0s +tttg: c151/344 lr:0.000598 t:11.1s +tttg: c152/344 lr:0.000593 t:11.2s +tttg: c153/344 lr:0.000589 t:11.3s +tttg: c154/344 lr:0.000584 t:11.3s +tttg: c155/344 lr:0.000580 t:11.4s +tttg: c156/344 lr:0.000575 t:11.5s +tttg: c157/344 lr:0.000571 t:11.6s +tttg: c158/344 lr:0.000566 t:11.6s +tttg: c159/344 lr:0.000562 t:11.7s +tttg: c160/344 lr:0.000557 t:11.8s +tttg: c161/344 lr:0.000553 t:11.8s +tttg: c162/344 lr:0.000548 t:11.9s +tttg: c163/344 lr:0.000543 t:12.0s +tttg: c164/344 lr:0.000539 t:12.1s +tttg: c165/344 lr:0.000534 t:12.1s +tttg: c166/344 lr:0.000530 t:12.2s +tttg: c167/344 lr:0.000525 t:12.3s +tttg: c168/344 lr:0.000521 t:12.3s +tttg: c169/344 lr:0.000516 t:12.4s +tttg: c170/344 lr:0.000511 t:12.5s +tttg: c171/344 lr:0.000507 t:12.6s +tttg: c172/344 lr:0.000502 t:12.6s +tttg: c173/344 lr:0.000498 t:12.7s +tttg: c174/344 lr:0.000493 t:12.8s +tttg: c175/344 lr:0.000489 t:12.9s +tttg: c176/344 lr:0.000484 t:12.9s +tttg: c177/344 lr:0.000479 t:13.0s +tttg: c178/344 lr:0.000475 t:13.1s +tttg: c179/344 lr:0.000470 t:13.1s +tttg: c180/344 lr:0.000466 t:13.2s +tttg: c181/344 lr:0.000461 t:13.3s +tttg: c182/344 lr:0.000457 t:13.4s +tttg: c183/344 lr:0.000452 t:13.4s +tttg: c184/344 lr:0.000447 t:13.5s +tttg: c185/344 lr:0.000443 t:13.6s +tttg: c186/344 lr:0.000438 t:13.6s +tttg: c187/344 lr:0.000434 t:13.7s +tttg: c188/344 lr:0.000429 t:13.8s +tttg: c189/344 lr:0.000425 t:13.9s +tttg: c190/344 lr:0.000420 t:13.9s +tttg: c191/344 lr:0.000416 t:14.0s +tttg: c192/344 lr:0.000411 t:14.1s +tttg: c193/344 lr:0.000407 t:14.2s +tttg: c194/344 lr:0.000402 t:14.2s +tttg: c195/344 lr:0.000398 t:14.3s +tttg: c196/344 lr:0.000393 t:14.4s +tttg: c197/344 lr:0.000389 t:14.4s +tttg: c198/344 lr:0.000384 t:14.5s +tttg: c199/344 lr:0.000380 t:14.6s +tttg: c200/344 lr:0.000375 t:14.7s +tttg: c201/344 lr:0.000371 t:14.7s +tttg: c202/344 lr:0.000367 t:14.8s +tttg: c203/344 lr:0.000362 t:14.9s +tttg: c204/344 lr:0.000358 t:14.9s +tttg: c205/344 lr:0.000353 t:15.0s +tttg: c206/344 lr:0.000349 t:15.1s +tttg: c207/344 lr:0.000345 t:15.2s +tttg: c208/344 lr:0.000340 t:15.2s +tttg: c209/344 lr:0.000336 t:15.3s +tttg: c210/344 lr:0.000332 t:15.4s +tttg: c211/344 lr:0.000327 t:15.5s +tttg: c212/344 lr:0.000323 t:15.5s +tttg: c213/344 lr:0.000319 t:15.6s +tttg: c214/344 lr:0.000314 t:15.7s +tttg: c215/344 lr:0.000310 t:15.7s +tttg: c216/344 lr:0.000306 t:15.8s +tttg: c217/344 lr:0.000302 t:15.9s +tttg: c218/344 lr:0.000298 t:16.0s +tttg: c219/344 lr:0.000293 t:16.0s +tttg: c220/344 lr:0.000289 t:16.1s +tttg: c221/344 lr:0.000285 t:16.2s +tttg: c222/344 lr:0.000281 t:16.2s +tttg: c223/344 lr:0.000277 t:16.3s +tttg: c224/344 lr:0.000273 t:16.4s +tttg: c225/344 lr:0.000269 t:16.5s +tttg: c226/344 lr:0.000265 t:16.5s +tttg: c227/344 lr:0.000261 t:16.6s +tttg: c228/344 lr:0.000257 t:16.7s +tttg: c229/344 lr:0.000253 t:16.8s +tttg: c230/344 lr:0.000249 t:16.8s +tttg: c231/344 lr:0.000245 t:16.9s +tttg: c232/344 lr:0.000241 t:17.0s +tttg: c233/344 lr:0.000237 t:17.0s +tttg: c234/344 lr:0.000233 t:17.1s +tttg: c235/344 lr:0.000229 t:17.2s +tttg: c236/344 lr:0.000225 t:17.3s +tttg: c237/344 lr:0.000222 t:17.3s +tttg: c238/344 lr:0.000218 t:17.4s +tttg: c239/344 lr:0.000214 t:17.5s +tttg: c240/344 lr:0.000210 t:17.5s +tttg: c241/344 lr:0.000206 t:17.6s +tttg: c242/344 lr:0.000203 t:17.7s +tttg: c243/344 lr:0.000199 t:17.8s +tttg: c244/344 lr:0.000195 t:17.8s +tttg: c245/344 lr:0.000192 t:17.9s +tttg: c246/344 lr:0.000188 t:18.0s +tttg: c247/344 lr:0.000185 t:18.1s +tttg: c248/344 lr:0.000181 t:18.1s +tttg: c249/344 lr:0.000178 t:18.2s +tttg: c250/344 lr:0.000174 t:18.3s +tttg: c251/344 lr:0.000171 t:18.3s +tttg: c252/344 lr:0.000167 t:18.4s +tttg: c253/344 lr:0.000164 t:18.5s +tttg: c254/344 lr:0.000160 t:18.6s +tttg: c255/344 lr:0.000157 t:18.6s +tttg: c256/344 lr:0.000154 t:18.7s +tttg: c257/344 lr:0.000151 t:18.8s +tttg: c258/344 lr:0.000147 t:18.8s +tttg: c259/344 lr:0.000144 t:18.9s +tttg: c260/344 lr:0.000141 t:19.0s +tttg: c261/344 lr:0.000138 t:19.1s +tttg: c262/344 lr:0.000135 t:19.1s +tttg: c263/344 lr:0.000131 t:19.2s +tttg: c264/344 lr:0.000128 t:19.3s +tttg: c265/344 lr:0.000125 t:19.4s +tttg: c266/344 lr:0.000122 t:19.4s +tttg: c267/344 lr:0.000119 t:19.5s +tttg: c268/344 lr:0.000116 t:19.6s +tttg: c269/344 lr:0.000113 t:19.6s +tttg: c270/344 lr:0.000111 t:19.7s +tttg: c271/344 lr:0.000108 t:19.8s +tttg: c272/344 lr:0.000105 t:19.9s +tttg: c273/344 lr:0.000102 t:19.9s +tttg: c274/344 lr:0.000099 t:20.0s +tttg: c275/344 lr:0.000097 t:20.1s +tttg: c276/344 lr:0.000094 t:20.1s +tttg: c277/344 lr:0.000091 t:20.2s +tttg: c278/344 lr:0.000089 t:20.3s +tttg: c279/344 lr:0.000086 t:20.4s +tttg: c280/344 lr:0.000083 t:20.4s +tttg: c281/344 lr:0.000081 t:20.5s +tttg: c282/344 lr:0.000078 t:20.6s +tttg: c283/344 lr:0.000076 t:20.7s +tttg: c284/344 lr:0.000074 t:20.7s +tttg: c285/344 lr:0.000071 t:20.8s +tttg: c286/344 lr:0.000069 t:20.9s +tttg: c287/344 lr:0.000067 t:20.9s +tttg: c288/344 lr:0.000064 t:21.0s +tttg: c289/344 lr:0.000062 t:21.1s +tttg: c290/344 lr:0.000060 t:21.2s +tttg: c291/344 lr:0.000058 t:21.2s +tttg: c292/344 lr:0.000056 t:21.3s +tttg: c293/344 lr:0.000054 t:21.4s +tttg: c294/344 lr:0.000052 t:21.4s +tttg: c295/344 lr:0.000050 t:21.5s +tttg: c296/344 lr:0.000048 t:21.6s +tttg: c297/344 lr:0.000046 t:21.7s +tttg: c298/344 lr:0.000044 t:21.7s +tttg: c299/344 lr:0.000042 t:21.8s +tttg: c300/344 lr:0.000040 t:21.9s +tttg: c301/344 lr:0.000038 t:22.0s +tttg: c302/344 lr:0.000037 t:22.0s +tttg: c303/344 lr:0.000035 t:22.1s +tttg: c304/344 lr:0.000033 t:22.2s +tttg: c305/344 lr:0.000032 t:22.2s +tttg: c306/344 lr:0.000030 t:22.3s +tttg: c307/344 lr:0.000028 t:22.4s +tttg: c308/344 lr:0.000027 t:22.5s +tttg: c309/344 lr:0.000025 t:22.5s +tttg: c310/344 lr:0.000024 t:22.6s +tttg: c311/344 lr:0.000023 t:22.7s +tttg: c312/344 lr:0.000021 t:22.7s +tttg: c313/344 lr:0.000020 t:22.8s +tttg: c314/344 lr:0.000019 t:22.9s +tttg: c315/344 lr:0.000018 t:23.0s +tttg: c316/344 lr:0.000016 t:23.0s +tttg: c317/344 lr:0.000015 t:23.1s +tttg: c318/344 lr:0.000014 t:23.2s +tttg: c319/344 lr:0.000013 t:23.3s +tttg: c320/344 lr:0.000012 t:23.3s +tttg: c321/344 lr:0.000011 t:23.4s +tttg: c322/344 lr:0.000010 t:23.5s +tttg: c323/344 lr:0.000009 t:23.5s +tttg: c324/344 lr:0.000008 t:23.6s +tttg: c325/344 lr:0.000008 t:23.7s +tttg: c326/344 lr:0.000007 t:23.8s +tttg: c327/344 lr:0.000006 t:23.8s +tttg: c328/344 lr:0.000005 t:23.9s +tttg: c329/344 lr:0.000005 t:24.0s +tttg: c330/344 lr:0.000004 t:24.0s +tttg: c331/344 lr:0.000004 t:24.1s +tttg: c332/344 lr:0.000003 t:24.2s +tttg: c333/344 lr:0.000003 t:24.3s +tttg: c334/344 lr:0.000002 t:24.3s +tttg: c335/344 lr:0.000002 t:24.4s +tttg: c336/344 lr:0.000001 t:24.5s +tttg: c337/344 lr:0.000001 t:24.6s +tttg: c338/344 lr:0.000001 t:24.6s +tttg: c339/344 lr:0.000001 t:24.7s +tttg: c340/344 lr:0.000000 t:24.8s +tttg: c341/344 lr:0.000000 t:24.8s +tttg: c342/344 lr:0.000000 t:24.9s +tttg: c343/344 lr:0.000000 t:25.0s +ttpr: phase:1/1 t:316.8s +ttp: b1972/2084 bl:2.3037 bb:1.0287 rl:2.1727 rb:1.0266 dl:2656-2671 gd:1 sr:0 sf:0 tr:24/24 wt:0 +ttp: b1959/2084 bl:2.2404 bb:1.0306 rl:2.1782 rb:1.0269 dl:2501-2514 gd:1 sr:0 sf:0 tr:24/24 wt:0 +ttp: b1953/2084 bl:2.2920 bb:1.0457 rl:2.1867 rb:1.0284 dl:2441-2454 gd:1 sr:0 sf:0 tr:24/24 wt:0 +ttp: b1947/2084 bl:2.2218 bb:0.9566 rl:2.1890 rb:1.0231 dl:2368-2382 gd:1 sr:0 sf:0 tr:24/24 wt:0 +ttp: b1944/2084 bl:2.4072 bb:1.1086 rl:2.2026 rb:1.0285 dl:2344-2352 gd:1 sr:0 sf:0 tr:24/24 wt:0 +ttp: b1939/2084 bl:2.3912 bb:1.0797 rl:2.2135 rb:1.0316 dl:2301-2307 gd:1 sr:0 sf:0 tr:24/24 wt:0 +ttp: b1929/2084 bl:2.2773 bb:1.0233 rl:2.2168 rb:1.0311 dl:2203-2216 gd:1 sr:0 sf:0 tr:24/24 wt:0 +ttp: b1925/2084 bl:2.2719 bb:1.0648 rl:2.2195 rb:1.0328 dl:2174-2183 gd:1 sr:0 sf:0 tr:24/24 wt:0 +ttp: b1919/2084 bl:2.2953 bb:1.0269 rl:2.2230 rb:1.0325 dl:2129-2135 gd:1 sr:0 sf:0 tr:24/24 wt:0 +ttp: b1913/2084 bl:2.3872 bb:1.0276 rl:2.2301 rb:1.0323 dl:2089-2095 gd:1 sr:0 sf:0 tr:24/24 wt:0 +ttp: b1907/2084 bl:2.3439 bb:1.0490 rl:2.2347 rb:1.0330 dl:2049-2054 gd:1 sr:0 sf:0 tr:24/24 wt:0 +ttp: b1902/2084 bl:2.3008 bb:1.0195 rl:2.2372 rb:1.0324 dl:2015-2022 gd:1 sr:0 sf:0 tr:24/24 wt:0 +ttp: b1893/2084 bl:2.1907 bb:1.0298 rl:2.2355 rb:1.0323 dl:1958-1963 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1891/2084 bl:2.4340 bb:1.1266 rl:2.2424 rb:1.0356 dl:1948-1953 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1883/2084 bl:2.5240 bb:1.0691 rl:2.2516 rb:1.0368 dl:1906-1913 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1879/2084 bl:2.2711 bb:1.0436 rl:2.2522 rb:1.0370 dl:1888-1891 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1874/2084 bl:2.3808 bb:1.0486 rl:2.2560 rb:1.0373 dl:1863-1868 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1864/2084 bl:2.3293 bb:1.0243 rl:2.2581 rb:1.0369 dl:1820-1824 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1862/2084 bl:2.4527 bb:1.0462 rl:2.2635 rb:1.0372 dl:1813-1817 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1847/2084 bl:2.2587 bb:1.0269 rl:2.2633 rb:1.0370 dl:1749-1753 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1844/2084 bl:2.3382 bb:1.0369 rl:2.2652 rb:1.0370 dl:1738-1741 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1837/2084 bl:2.4234 bb:1.0923 rl:2.2690 rb:1.0383 dl:1710-1713 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1829/2084 bl:2.4415 bb:1.0298 rl:2.2730 rb:1.0381 dl:1680-1684 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1822/2084 bl:2.2444 bb:1.0050 rl:2.2724 rb:1.0373 dl:1657-1659 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1817/2084 bl:2.4119 bb:1.0958 rl:2.2754 rb:1.0386 dl:1638-1641 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1811/2084 bl:2.2371 bb:1.0247 rl:2.2746 rb:1.0383 dl:1616-1619 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1805/2084 bl:2.2492 bb:1.0093 rl:2.2741 rb:1.0377 dl:1598-1601 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1799/2084 bl:2.2986 bb:1.0288 rl:2.2746 rb:1.0375 dl:1581-1583 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1794/2084 bl:2.2311 bb:1.0333 rl:2.2737 rb:1.0375 dl:1565-1569 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1788/2084 bl:2.3217 bb:1.0790 rl:2.2746 rb:1.0382 dl:1546-1549 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1782/2084 bl:2.4991 bb:1.0670 rl:2.2786 rb:1.0388 dl:1529-1532 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1774/2084 bl:2.2905 bb:1.0168 rl:2.2788 rb:1.0384 dl:1506-1509 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1766/2084 bl:2.2640 bb:1.0262 rl:2.2786 rb:1.0382 dl:1486-1488 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1758/2084 bl:2.3501 bb:1.0474 rl:2.2798 rb:1.0383 dl:1465-1467 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1749/2084 bl:2.3716 bb:1.1317 rl:2.2812 rb:1.0397 dl:1442-1444 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1741/2084 bl:2.2386 bb:1.0290 rl:2.2805 rb:1.0396 dl:1422-1425 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1733/2084 bl:2.4258 bb:1.0919 rl:2.2827 rb:1.0404 dl:1402-1405 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1725/2084 bl:2.2870 bb:1.0430 rl:2.2828 rb:1.0404 dl:1384-1387 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1717/2084 bl:2.4358 bb:1.0468 rl:2.2849 rb:1.0405 dl:1364-1366 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1709/2084 bl:2.3212 bb:1.0283 rl:2.2854 rb:1.0403 dl:1346-1349 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1701/2084 bl:2.3437 bb:1.0570 rl:2.2862 rb:1.0405 dl:1329-1330 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1693/2084 bl:2.3461 bb:1.0696 rl:2.2870 rb:1.0409 dl:1311-1314 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1686/2084 bl:2.2953 bb:1.0182 rl:2.2871 rb:1.0406 dl:1296-1299 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1679/2084 bl:2.2514 bb:0.9891 rl:2.2866 rb:1.0400 dl:1281-1284 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1671/2084 bl:2.4143 bb:1.1098 rl:2.2882 rb:1.0408 dl:1267-1269 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1663/2084 bl:2.2498 bb:1.0569 rl:2.2877 rb:1.0410 dl:1250-1252 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1655/2084 bl:2.2745 bb:1.0609 rl:2.2876 rb:1.0412 dl:1232-1235 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1647/2084 bl:2.4465 bb:1.1079 rl:2.2894 rb:1.0420 dl:1218-1220 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1639/2084 bl:2.2445 bb:1.0367 rl:2.2889 rb:1.0419 dl:1201-1203 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1631/2084 bl:2.2444 bb:1.0137 rl:2.2884 rb:1.0416 dl:1187-1189 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1623/2084 bl:2.3118 bb:1.0374 rl:2.2886 rb:1.0416 dl:1174-1175 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1615/2084 bl:2.3712 bb:1.0591 rl:2.2895 rb:1.0417 dl:1160-1161 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1607/2084 bl:2.3506 bb:1.0586 rl:2.2901 rb:1.0419 dl:1147-1148 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1599/2084 bl:2.1824 bb:0.9750 rl:2.2890 rb:1.0412 dl:1134-1135 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1591/2084 bl:2.4713 bb:1.0950 rl:2.2908 rb:1.0418 dl:1121-1123 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1583/2084 bl:2.3365 bb:1.0641 rl:2.2912 rb:1.0420 dl:1108-1109 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1575/2084 bl:2.3153 bb:1.0807 rl:2.2914 rb:1.0423 dl:1095-1097 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1567/2084 bl:2.3377 bb:1.0576 rl:2.2919 rb:1.0425 dl:1082-1084 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1559/2084 bl:2.2509 bb:1.0300 rl:2.2915 rb:1.0424 dl:1070-1071 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1551/2084 bl:2.3044 bb:1.0077 rl:2.2916 rb:1.0420 dl:1057-1059 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1543/2084 bl:2.4296 bb:1.0455 rl:2.2928 rb:1.0421 dl:1045-1046 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1535/2084 bl:2.3398 bb:1.1006 rl:2.2932 rb:1.0425 dl:1034-1035 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1528/2084 bl:2.3439 bb:1.1045 rl:2.2936 rb:1.0430 dl:1024-1025 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1522/2084 bl:2.2184 bb:1.0107 rl:2.2930 rb:1.0428 dl:1015-1016 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1515/2084 bl:2.3198 bb:1.0885 rl:2.2932 rb:1.0431 dl:1006-1008 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1508/2084 bl:2.3309 bb:1.0537 rl:2.2935 rb:1.0432 dl:997-998 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1500/2084 bl:2.2974 bb:0.9957 rl:2.2935 rb:1.0428 dl:987-988 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1492/2084 bl:2.2703 bb:0.9893 rl:2.2933 rb:1.0424 dl:976-977 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1485/2084 bl:2.2383 bb:0.9772 rl:2.2929 rb:1.0419 dl:967-967 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1476/2084 bl:2.3780 bb:1.0813 rl:2.2936 rb:1.0422 dl:956-957 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1468/2084 bl:2.2777 bb:1.0155 rl:2.2934 rb:1.0420 dl:947-948 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1459/2084 bl:2.3996 bb:1.0990 rl:2.2942 rb:1.0424 dl:936-937 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1452/2084 bl:2.2943 bb:1.0309 rl:2.2942 rb:1.0423 dl:928-929 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1444/2084 bl:2.2880 bb:1.0580 rl:2.2941 rb:1.0424 dl:919-920 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1436/2084 bl:2.3464 bb:1.0390 rl:2.2945 rb:1.0424 dl:909-910 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1428/2084 bl:2.2463 bb:1.0149 rl:2.2942 rb:1.0422 dl:899-901 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1420/2084 bl:2.4097 bb:1.1075 rl:2.2949 rb:1.0426 dl:890-891 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1412/2084 bl:2.5674 bb:1.1094 rl:2.2966 rb:1.0431 dl:882-883 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1404/2084 bl:2.3726 bb:1.0503 rl:2.2971 rb:1.0431 dl:873-874 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1396/2084 bl:2.4740 bb:1.1000 rl:2.2982 rb:1.0435 dl:863-864 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1389/2084 bl:2.3213 bb:1.0404 rl:2.2983 rb:1.0435 dl:856-857 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1382/2084 bl:2.2262 bb:1.0284 rl:2.2979 rb:1.0434 dl:849-850 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1376/2084 bl:2.5007 bb:1.0708 rl:2.2991 rb:1.0435 dl:842-843 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1367/2084 bl:2.3632 bb:1.0158 rl:2.2995 rb:1.0434 dl:833-834 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1359/2084 bl:2.4214 bb:1.0712 rl:2.3002 rb:1.0435 dl:824-825 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1351/2084 bl:2.1977 bb:0.9716 rl:2.2996 rb:1.0431 dl:815-816 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1343/2084 bl:2.4053 bb:1.0279 rl:2.3002 rb:1.0430 dl:807-808 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1335/2084 bl:2.2828 bb:1.0616 rl:2.3001 rb:1.0431 dl:800-801 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1328/2084 bl:2.2309 bb:1.0060 rl:2.2997 rb:1.0429 dl:792-794 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1323/2084 bl:2.4452 bb:1.1025 rl:2.3005 rb:1.0432 dl:788-788 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1315/2084 bl:2.4370 bb:1.1168 rl:2.3012 rb:1.0436 dl:780-781 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1306/2084 bl:2.3785 bb:1.0360 rl:2.3016 rb:1.0436 dl:772-773 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1299/2084 bl:2.2666 bb:1.0679 rl:2.3014 rb:1.0437 dl:766-767 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1291/2084 bl:2.3539 bb:1.0383 rl:2.3017 rb:1.0437 dl:758-759 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1283/2084 bl:2.3286 bb:1.0459 rl:2.3018 rb:1.0437 dl:751-752 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1275/2084 bl:2.2379 bb:1.0838 rl:2.3015 rb:1.0439 dl:744-745 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1267/2084 bl:2.2753 bb:1.0423 rl:2.3014 rb:1.0439 dl:737-738 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1258/2084 bl:2.4520 bb:1.1345 rl:2.3021 rb:1.0443 dl:729-730 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1251/2084 bl:2.3423 bb:1.0234 rl:2.3022 rb:1.0442 dl:723-724 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1243/2084 bl:2.4253 bb:1.0568 rl:2.3028 rb:1.0442 dl:716-717 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1235/2084 bl:2.0942 bb:0.9875 rl:2.3019 rb:1.0440 dl:709-710 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1227/2084 bl:2.3002 bb:1.0087 rl:2.3019 rb:1.0438 dl:703-704 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1219/2084 bl:2.2421 bb:1.0251 rl:2.3016 rb:1.0437 dl:696-697 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1210/2084 bl:2.3819 bb:1.0851 rl:2.3019 rb:1.0439 dl:690-691 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1203/2084 bl:2.2042 bb:1.0029 rl:2.3015 rb:1.0437 dl:684-685 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1196/2084 bl:2.4135 bb:1.0722 rl:2.3020 rb:1.0439 dl:679-679 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1187/2084 bl:2.2470 bb:1.0371 rl:2.3018 rb:1.0438 dl:671-672 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1179/2084 bl:2.2238 bb:1.0665 rl:2.3015 rb:1.0439 dl:665-666 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1171/2084 bl:2.2535 bb:1.0642 rl:2.3013 rb:1.0440 dl:659-660 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1163/2084 bl:2.3493 bb:1.0470 rl:2.3015 rb:1.0440 dl:653-653 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1155/2084 bl:2.3868 bb:1.0831 rl:2.3018 rb:1.0442 dl:646-647 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1147/2084 bl:2.2641 bb:1.0804 rl:2.3016 rb:1.0443 dl:641-642 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1139/2084 bl:2.2259 bb:1.0180 rl:2.3014 rb:1.0442 dl:636-637 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1130/2084 bl:2.3306 bb:1.0420 rl:2.3015 rb:1.0442 dl:629-630 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1124/2084 bl:2.3133 bb:1.1163 rl:2.3015 rb:1.0445 dl:625-625 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1115/2084 bl:2.2586 bb:1.0113 rl:2.3013 rb:1.0443 dl:618-619 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1108/2084 bl:2.3009 bb:1.0313 rl:2.3013 rb:1.0443 dl:613-614 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1100/2084 bl:2.2823 bb:1.0281 rl:2.3013 rb:1.0442 dl:607-608 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1092/2084 bl:2.2927 bb:1.0303 rl:2.3012 rb:1.0442 dl:601-602 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1083/2084 bl:2.2726 bb:1.0529 rl:2.3011 rb:1.0442 dl:595-596 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1077/2084 bl:2.3327 bb:1.0182 rl:2.3013 rb:1.0441 dl:591-591 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1068/2084 bl:2.3880 bb:1.0722 rl:2.3016 rb:1.0442 dl:584-585 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1059/2084 bl:2.3079 bb:1.0314 rl:2.3016 rb:1.0442 dl:578-579 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1051/2084 bl:2.1635 bb:0.9866 rl:2.3011 rb:1.0440 dl:573-574 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1044/2084 bl:2.3689 bb:1.0373 rl:2.3013 rb:1.0440 dl:568-569 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1035/2084 bl:2.3039 bb:1.0393 rl:2.3013 rb:1.0439 dl:562-563 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1029/2084 bl:2.3679 bb:1.0800 rl:2.3016 rb:1.0441 dl:558-558 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1021/2084 bl:2.2989 bb:1.0483 rl:2.3016 rb:1.0441 dl:553-553 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1013/2084 bl:2.3044 bb:1.0849 rl:2.3016 rb:1.0442 dl:548-548 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b1005/2084 bl:2.2359 bb:0.9956 rl:2.3014 rb:1.0440 dl:543-543 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b995/2084 bl:2.3963 bb:1.0611 rl:2.3016 rb:1.0441 dl:536-537 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b989/2084 bl:2.2311 bb:1.1134 rl:2.3014 rb:1.0443 dl:533-533 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b979/2084 bl:2.4471 bb:1.1362 rl:2.3019 rb:1.0445 dl:526-527 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b973/2084 bl:2.3194 bb:1.0418 rl:2.3019 rb:1.0445 dl:523-523 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b965/2084 bl:2.3347 bb:1.0618 rl:2.3020 rb:1.0446 dl:518-518 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b957/2084 bl:2.1167 bb:1.0327 rl:2.3015 rb:1.0446 dl:513-513 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b947/2084 bl:2.3554 bb:1.0634 rl:2.3016 rb:1.0446 dl:506-507 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b940/2084 bl:2.1923 bb:0.9745 rl:2.3013 rb:1.0444 dl:501-502 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b931/2084 bl:2.4409 bb:1.0904 rl:2.3017 rb:1.0445 dl:496-497 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b924/2084 bl:2.3316 bb:1.0737 rl:2.3018 rb:1.0446 dl:492-493 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b915/2084 bl:2.1701 bb:1.0359 rl:2.3014 rb:1.0446 dl:488-488 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b907/2084 bl:2.4582 bb:1.1055 rl:2.3019 rb:1.0448 dl:483-484 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b901/2084 bl:2.3670 bb:1.1022 rl:2.3020 rb:1.0449 dl:480-480 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b892/2084 bl:2.4730 bb:1.1224 rl:2.3025 rb:1.0451 dl:474-475 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b886/2084 bl:2.4277 bb:1.0650 rl:2.3028 rb:1.0452 dl:471-471 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b878/2084 bl:2.1518 bb:1.0587 rl:2.3024 rb:1.0452 dl:466-466 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b869/2084 bl:2.1972 bb:0.9692 rl:2.3021 rb:1.0450 dl:462-462 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b862/2084 bl:2.4512 bb:1.0875 rl:2.3025 rb:1.0451 dl:458-458 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b854/2084 bl:2.3515 bb:1.1179 rl:2.3026 rb:1.0453 dl:453-453 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b844/2084 bl:2.2896 bb:1.0354 rl:2.3026 rb:1.0452 dl:446-447 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b836/2084 bl:2.3816 bb:1.1326 rl:2.3028 rb:1.0454 dl:442-443 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b829/2084 bl:2.3679 bb:1.1244 rl:2.3029 rb:1.0456 dl:439-439 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b822/2084 bl:2.4277 bb:1.1082 rl:2.3032 rb:1.0458 dl:435-435 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b813/2084 bl:2.3858 bb:1.0558 rl:2.3034 rb:1.0458 dl:429-430 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b808/2084 bl:2.3901 bb:1.1179 rl:2.3036 rb:1.0460 dl:427-427 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b799/2084 bl:2.3577 bb:1.0431 rl:2.3037 rb:1.0459 dl:421-422 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b794/2084 bl:2.4802 bb:1.1045 rl:2.3041 rb:1.0461 dl:419-419 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b787/2084 bl:2.2207 bb:1.0869 rl:2.3040 rb:1.0462 dl:415-415 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b778/2084 bl:2.3105 bb:1.1007 rl:2.3040 rb:1.0463 dl:409-410 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b771/2084 bl:2.3127 bb:1.0550 rl:2.3040 rb:1.0463 dl:406-406 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b761/2084 bl:2.4588 bb:1.0953 rl:2.3043 rb:1.0464 dl:400-401 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b754/2084 bl:2.3146 bb:1.1566 rl:2.3043 rb:1.0466 dl:397-397 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b746/2084 bl:2.3376 bb:1.1265 rl:2.3044 rb:1.0468 dl:393-393 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b738/2084 bl:2.2244 bb:1.0236 rl:2.3042 rb:1.0467 dl:389-389 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b727/2084 bl:2.3962 bb:1.0468 rl:2.3044 rb:1.0467 dl:383-384 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b722/2084 bl:2.4172 bb:1.1396 rl:2.3046 rb:1.0469 dl:381-381 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b714/2084 bl:2.3500 bb:1.0803 rl:2.3047 rb:1.0470 dl:377-377 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b706/2084 bl:2.5198 bb:1.1358 rl:2.3051 rb:1.0471 dl:373-373 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b699/2084 bl:2.4469 bb:1.1258 rl:2.3054 rb:1.0473 dl:370-370 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b691/2084 bl:2.3455 bb:1.0804 rl:2.3055 rb:1.0473 dl:366-366 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b683/2084 bl:2.4618 bb:1.1119 rl:2.3058 rb:1.0475 dl:362-362 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b674/2084 bl:2.3718 bb:1.1595 rl:2.3059 rb:1.0477 dl:358-358 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b667/2084 bl:2.2797 bb:1.0279 rl:2.3059 rb:1.0476 dl:355-355 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b659/2084 bl:2.4349 bb:1.1632 rl:2.3061 rb:1.0478 dl:351-351 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b651/2084 bl:2.4752 bb:1.2208 rl:2.3064 rb:1.0481 dl:347-347 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b641/2084 bl:2.3218 bb:1.0836 rl:2.3064 rb:1.0482 dl:342-343 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b633/2084 bl:2.3735 bb:1.1192 rl:2.3065 rb:1.0483 dl:338-339 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b628/2084 bl:2.4228 bb:1.0969 rl:2.3067 rb:1.0484 dl:336-336 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b617/2084 bl:2.3564 bb:1.1481 rl:2.3068 rb:1.0485 dl:331-332 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b611/2084 bl:2.2365 bb:1.0693 rl:2.3067 rb:1.0485 dl:329-329 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b601/2084 bl:2.2865 bb:1.0411 rl:2.3067 rb:1.0485 dl:324-325 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b595/2084 bl:2.3443 bb:1.1568 rl:2.3067 rb:1.0487 dl:322-322 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b588/2084 bl:2.3594 bb:1.0914 rl:2.3068 rb:1.0488 dl:319-319 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b582/2084 bl:2.3920 bb:1.1356 rl:2.3069 rb:1.0489 dl:316-316 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b571/2084 bl:2.3564 bb:1.1054 rl:2.3070 rb:1.0490 dl:311-312 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b568/2084 bl:2.3223 bb:1.0379 rl:2.3071 rb:1.0490 dl:310-310 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b560/2084 bl:2.2609 bb:1.0571 rl:2.3070 rb:1.0490 dl:307-307 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b552/2084 bl:2.3679 bb:1.1600 rl:2.3071 rb:1.0491 dl:304-304 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b543/2084 bl:2.5027 bb:1.1924 rl:2.3074 rb:1.0493 dl:300-301 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b540/2084 bl:2.5250 bb:1.1230 rl:2.3077 rb:1.0495 dl:299-299 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b532/2084 bl:2.3535 bb:1.0600 rl:2.3078 rb:1.0495 dl:296-296 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b521/2084 bl:2.4173 bb:1.1191 rl:2.3079 rb:1.0496 dl:291-292 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b515/2084 bl:2.2868 bb:1.0289 rl:2.3079 rb:1.0495 dl:288-289 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b508/2084 bl:2.3737 bb:1.0998 rl:2.3080 rb:1.0496 dl:285-286 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b500/2084 bl:2.3216 bb:1.1095 rl:2.3080 rb:1.0497 dl:282-283 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b493/2084 bl:2.4302 bb:1.1185 rl:2.3082 rb:1.0498 dl:280-280 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b484/2084 bl:2.3147 bb:1.1144 rl:2.3082 rb:1.0499 dl:276-277 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b477/2084 bl:2.5798 bb:1.1918 rl:2.3085 rb:1.0501 dl:274-274 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b468/2084 bl:2.2962 bb:1.1252 rl:2.3085 rb:1.0501 dl:270-271 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b460/2084 bl:2.3968 bb:1.1358 rl:2.3086 rb:1.0503 dl:267-268 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b452/2084 bl:2.2644 bb:1.1493 rl:2.3086 rb:1.0504 dl:264-265 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b443/2084 bl:2.4097 bb:1.0593 rl:2.3087 rb:1.0504 dl:261-262 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b440/2084 bl:2.2549 bb:1.1267 rl:2.3086 rb:1.0505 dl:260-260 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b432/2084 bl:2.3005 bb:1.0483 rl:2.3086 rb:1.0505 dl:257-257 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b422/2084 bl:2.5596 bb:1.2266 rl:2.3089 rb:1.0507 dl:254-254 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b414/2084 bl:2.4478 bb:1.1994 rl:2.3091 rb:1.0508 dl:251-251 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b405/2084 bl:2.4833 bb:1.1431 rl:2.3093 rb:1.0510 dl:247-248 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b397/2084 bl:2.3491 bb:1.1491 rl:2.3094 rb:1.0511 dl:244-245 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b388/2084 bl:2.5166 bb:1.2216 rl:2.3096 rb:1.0513 dl:241-242 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b384/2084 bl:2.3755 bb:1.1168 rl:2.3097 rb:1.0513 dl:240-240 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b376/2084 bl:2.4955 bb:1.1637 rl:2.3099 rb:1.0515 dl:237-237 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b367/2084 bl:2.3786 bb:1.1079 rl:2.3100 rb:1.0515 dl:234-234 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b357/2084 bl:2.2753 bb:1.0483 rl:2.3099 rb:1.0515 dl:230-231 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b353/2084 bl:2.4520 bb:1.0987 rl:2.3101 rb:1.0516 dl:229-229 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b343/2084 bl:2.3172 bb:1.0543 rl:2.3101 rb:1.0516 dl:225-226 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b334/2084 bl:2.3180 bb:1.1043 rl:2.3101 rb:1.0516 dl:222-223 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b326/2084 bl:2.2605 bb:1.1196 rl:2.3101 rb:1.0517 dl:219-220 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b320/2084 bl:2.2768 bb:1.0736 rl:2.3100 rb:1.0517 dl:217-218 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b314/2084 bl:2.3615 bb:1.0770 rl:2.3101 rb:1.0517 dl:215-216 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b308/2084 bl:2.4196 bb:1.1211 rl:2.3102 rb:1.0518 dl:213-214 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b303/2084 bl:2.4686 bb:1.0788 rl:2.3104 rb:1.0518 dl:212-212 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b299/2084 bl:2.5307 bb:1.1728 rl:2.3106 rb:1.0520 dl:210-210 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b290/2084 bl:2.6299 bb:1.2325 rl:2.3109 rb:1.0521 dl:207-207 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b280/2084 bl:2.3597 bb:1.1472 rl:2.3109 rb:1.0522 dl:204-204 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b275/2084 bl:2.5073 bb:1.2187 rl:2.3111 rb:1.0524 dl:202-202 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b265/2084 bl:2.4387 bb:1.1746 rl:2.3113 rb:1.0525 dl:199-199 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b256/2084 bl:2.3836 bb:1.1350 rl:2.3113 rb:1.0525 dl:196-196 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b247/2084 bl:2.5788 bb:1.2703 rl:2.3116 rb:1.0527 dl:193-193 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b242/2084 bl:2.4908 bb:1.1427 rl:2.3117 rb:1.0528 dl:191-191 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b233/2084 bl:2.4537 bb:1.1735 rl:2.3119 rb:1.0529 dl:188-188 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b228/2084 bl:2.4509 bb:1.1108 rl:2.3120 rb:1.0530 dl:186-186 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b217/2084 bl:2.4964 bb:1.1316 rl:2.3121 rb:1.0530 dl:183-183 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b209/2084 bl:2.4526 bb:1.1578 rl:2.3123 rb:1.0531 dl:180-180 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b198/2084 bl:2.5217 bb:1.2087 rl:2.3124 rb:1.0532 dl:176-177 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b195/2084 bl:2.4219 bb:1.1550 rl:2.3125 rb:1.0533 dl:175-175 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b183/2084 bl:2.4239 bb:1.1565 rl:2.3126 rb:1.0534 dl:171-172 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b179/2084 bl:2.4478 bb:1.1561 rl:2.3127 rb:1.0535 dl:170-170 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b171/2084 bl:2.3746 bb:1.1650 rl:2.3128 rb:1.0536 dl:167-167 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b160/2084 bl:2.3888 bb:1.1665 rl:2.3128 rb:1.0536 dl:163-164 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b156/2084 bl:2.5248 bb:1.2132 rl:2.3130 rb:1.0538 dl:162-162 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b146/2084 bl:2.5750 bb:1.2189 rl:2.3132 rb:1.0539 dl:158-159 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b142/2084 bl:2.4366 bb:1.2337 rl:2.3133 rb:1.0540 dl:157-157 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b137/2084 bl:2.6357 bb:1.2458 rl:2.3135 rb:1.0541 dl:155-155 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b127/2084 bl:2.3560 bb:1.1564 rl:2.3135 rb:1.0542 dl:151-152 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b124/2084 bl:2.4587 bb:1.1269 rl:2.3136 rb:1.0543 dl:150-150 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b114/2084 bl:2.6527 bb:1.3165 rl:2.3139 rb:1.0544 dl:146-147 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b106/2084 bl:2.6070 bb:1.1517 rl:2.3141 rb:1.0545 dl:144-144 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b99/2084 bl:2.5362 bb:1.1563 rl:2.3142 rb:1.0546 dl:141-142 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b93/2084 bl:2.6379 bb:1.2540 rl:2.3144 rb:1.0547 dl:139-139 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b84/2084 bl:2.5421 bb:1.2155 rl:2.3146 rb:1.0548 dl:135-136 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b76/2084 bl:2.7307 bb:1.2663 rl:2.3148 rb:1.0549 dl:132-133 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b73/2084 bl:2.6035 bb:1.2222 rl:2.3150 rb:1.0550 dl:131-131 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b65/2084 bl:2.5254 bb:1.1873 rl:2.3151 rb:1.0551 dl:128-128 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b54/2084 bl:2.5950 bb:1.2151 rl:2.3153 rb:1.0552 dl:123-124 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b48/2084 bl:2.7354 bb:1.2275 rl:2.3155 rb:1.0553 dl:120-121 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b43/2084 bl:2.5527 bb:1.2568 rl:2.3157 rb:1.0554 dl:118-118 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b35/2084 bl:2.6573 bb:1.1928 rl:2.3158 rb:1.0554 dl:113-114 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b28/2084 bl:2.8248 bb:1.2632 rl:2.3161 rb:1.0555 dl:109-109 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b22/2084 bl:2.8699 bb:1.2832 rl:2.3164 rb:1.0557 dl:105-105 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b13/2084 bl:2.7730 bb:1.2278 rl:2.3166 rb:1.0557 dl:98-99 gd:1 sr:0 sf:1 tr:24/24 wt:0 +ttp: b6/2084 bl:2.6705 bb:1.1681 rl:2.3167 rb:1.0558 dl:89-90 gd:1 sr:0 sf:1 tr:24/24 wt:0 +peer_ens:coverage_tokens:47853343/47853343 batches_covered:263 +peer_ens:route_frac:0.752235 routed_tokens:35996952/47853343 +peer_ens:ensemble_covered_only val_loss:2.314187 val_bpb:1.057491 +peer_ens:baseline val_loss:2.316573 val_bpb:1.058581 +quantized_ttt_phased val_loss:2.31418732 val_bpb:1.05749056 eval_time:588788ms +total_eval_time:588.8s diff --git a/records/track_10min_16mb/2026-04-30_TTT_Ensemble/submission.json b/records/track_10min_16mb/2026-04-30_TTT_Ensemble/submission.json new file mode 100644 index 0000000000..c144376c44 --- /dev/null +++ b/records/track_10min_16mb/2026-04-30_TTT_Ensemble/submission.json @@ -0,0 +1,33 @@ +{ + "author": "varunneal", + "github_id": "varunneal", + "name": "Peer-LoRA Ensemble", + "date": "2026-05-01", + "track": "10min_16mb", + "val_bpb": 1.05749, + "val_bpb_std": null, + "seeds": [42], + "seed_results": { + "42": {"val_bpb": 1.05749, "artifact_bytes": 15986824} + }, + "hardware": "8xH100 80GB SXM", + "pytorch_version": "2.10.0+cu130", + "technique_summary": "PR #2014 stack with peer-LoRA ensemble: after per-doc LoRA training, run K-1 peer forwards and blend on uncertain tokens (entropy >= 0.5) in probability space (K=4, w=0.8). Tuned TTT LR=0.00015, WD=0.25.", + "compliance": { + "train_under_600s": true, + "artifact_under_16mb": true, + "eval_under_600s": true, + "no_slot": true, + "no_pre_quant_ttt": true, + "no_etlb": true, + "no_ngram_cache": true, + "score_first_ttt": true, + "three_seeds": true + }, + "attribution": { + "pr2014_stack": "@simonbissonnette (PR #2014)", + "per_doc_lora_TTT": "@samacqua (PR #1530)", + "mixture_of_models_ensembling": "inspired by Slowrun techniques from @akshayvegesna and others", + "peer_lora_ensemble": "@varunneal" + } +} diff --git a/records/track_10min_16mb/2026-04-30_TTT_Ensemble/submit.sh b/records/track_10min_16mb/2026-04-30_TTT_Ensemble/submit.sh new file mode 100755 index 0000000000..260385e038 --- /dev/null +++ b/records/track_10min_16mb/2026-04-30_TTT_Ensemble/submit.sh @@ -0,0 +1,79 @@ +#!/bin/bash +set -e + +DIR="$(cd "$(dirname "$0")" && pwd)" +REMOTE="ubuntu@91.239.86.233" + +echo "Waiting for seed 314 to finish..." +while true; do + LOGFILE=$(ssh $REMOTE 'for f in ~/parameter-golf/logs/*.txt; do + seed=$(grep "^seed: 314" "$f" 2>/dev/null | head -1) + if [ -n "$seed" ]; then echo "$f"; break; fi + done') + if [ -z "$LOGFILE" ]; then + echo " No seed 314 log found yet, waiting 30s..." + sleep 30 + continue + fi + DONE=$(ssh $REMOTE "grep -c '^quantized_ttt_phased' $LOGFILE 2>/dev/null || echo 0") + if [ "$DONE" -ge 1 ]; then + echo " Found completed log: $LOGFILE" + break + fi + echo " Log exists but not finished, waiting 30s..." + sleep 30 +done + +echo "Pulling log..." +scp "$REMOTE:$LOGFILE" "$DIR/seed314.log" + +echo "Parsing results..." +ENS_BPB=$(grep '^peer_ens:ensemble_covered_only' "$DIR/seed314.log" | grep -oE 'val_bpb:[0-9.]+' | cut -d: -f2) +BASELINE_BPB=$(grep '^peer_ens:baseline' "$DIR/seed314.log" | grep -oE 'val_bpb:[0-9.]+' | cut -d: -f2) +EVAL_TIME=$(grep '^quantized_ttt_phased' "$DIR/seed314.log" | grep -oE 'eval_time:[0-9.]+ms' | cut -d: -f2 | sed 's/ms//') +PRE_QUANT=$(grep '^diagnostic pre-quantization' "$DIR/seed314.log" | grep -oE 'val_bpb:[0-9.]+' | cut -d: -f2) +POST_QUANT=$(grep '^diagnostic quantized' "$DIR/seed314.log" | grep -oE 'val_bpb:[0-9.]+' | cut -d: -f2) +ARTIFACT=$(grep '^Total submission size' "$DIR/seed314.log" | grep -oE '[0-9]+ bytes' | cut -d' ' -f1) + +PRE_Q5=$(printf "%.5f" "$PRE_QUANT") +POST_Q5=$(printf "%.5f" "$POST_QUANT") +ENS_5=$(printf "%.5f" "$ENS_BPB") +DELTA=$(python3 -c "print(f'{$ENS_BPB - 1.05855:.5f}')") + +echo "Results:" +echo " Pre-quant BPB: $PRE_Q5" +echo " Post-quant BPB: $POST_Q5" +echo " Ensemble BPB: $ENS_5" +echo " Baseline BPB: $BASELINE_BPB" +echo " Eval time: ${EVAL_TIME}ms" +echo " Artifact bytes: $ARTIFACT" +echo " Delta vs #2014: $DELTA" + +echo "Updating README.md..." +sed -i '' "s/| 314 | - | - | - | - |/| 314 | $PRE_Q5 | $POST_Q5 | **$ENS_5** | $ARTIFACT |/" "$DIR/README.md" +sed -i '' "s/\*\*val_bpb = TBD\*\* (1 seed)/**val_bpb = $ENS_5** (1 seed)/" "$DIR/README.md" +sed -i '' "s/Delta: TBD/Delta: $DELTA vs PR #2014 baseline (1.05855)/" "$DIR/README.md" + +echo "Updating submission.json..." +python3 -c " +import json +with open('$DIR/submission.json') as f: + d = json.load(f) +d['val_bpb'] = round($ENS_BPB, 5) +d['val_bpb_std'] = None +d['seed_results'] = {'314': {'val_bpb': round($ENS_BPB, 5), 'artifact_bytes': $ARTIFACT}} +with open('$DIR/submission.json', 'w') as f: + json.dump(d, f, indent=2) + f.write('\n') +" + +echo "Committing and pushing..." +cd "$DIR/../../.." +git add records/track_10min_16mb/2026-04-30_TTT_Ensemble/README.md \ + records/track_10min_16mb/2026-04-30_TTT_Ensemble/submission.json \ + records/track_10min_16mb/2026-04-30_TTT_Ensemble/seed314.log +git commit -m "seed 314 results: val_bpb=$ENS_5, K=3 peer-LoRA ensemble" +git push fork ttt-ensemble-clean + +echo "" +echo "DONE! val_bpb=$ENS_5 eval_time=${EVAL_TIME}ms delta=$DELTA" diff --git a/records/track_10min_16mb/2026-04-30_TTT_Ensemble/train_gpt.py b/records/track_10min_16mb/2026-04-30_TTT_Ensemble/train_gpt.py new file mode 100644 index 0000000000..7ef8cafb69 --- /dev/null +++ b/records/track_10min_16mb/2026-04-30_TTT_Ensemble/train_gpt.py @@ -0,0 +1,4701 @@ +""" +parent: experiments/parameter-golf/pr2014.py +""" +import base64, collections, copy, fcntl, glob, io, lzma, math, os +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 Tensor, nn +from flash_attn_interface import ( + flash_attn_func as flash_attn_3_func, + flash_attn_varlen_func, +) +from concurrent.futures import ThreadPoolExecutor +import triton +import triton.language as tl +from triton.tools.tensor_descriptor import TensorDescriptor + + +# ===== Fused softcapped cross-entropy (Triton) — training-only path ===== +# Replaces the eager +# logits_softcap = softcap * tanh(logits / softcap) +# F.cross_entropy(logits_softcap.float(), targets, reduction="mean") +# sequence with a single fused kernel that reads logits_proj once, applies +# softcap in-register, and computes (LSE, loss) in one streaming pass. The +# backward kernel mirrors the forward so there's no stored softcapped logits. +# Numerically identical to the eager path up to fp32 accumulation differences. +_FUSED_CE_LIBRARY = "pgsubmission1draft7fusedce" +_FUSED_CE_BLOCK_SIZE = 1024 +_FUSED_CE_NUM_WARPS = 4 + + +@triton.jit +def _softcapped_ce_fwd_kernel( + logits_ptr, losses_ptr, lse_ptr, targets_ptr, + stride_logits_n, stride_logits_v, + n_rows, n_cols, softcap, + block_size: tl.constexpr, +): + row_idx = tl.program_id(0).to(tl.int64) + logits_row_ptr = logits_ptr + row_idx * stride_logits_n + max_val = -float("inf") + sum_exp = 0.0 + A = 2.0 * softcap + inv_C = 2.0 / softcap + for off in range(0, n_cols, block_size): + cols = off + tl.arange(0, block_size) + mask = cols < n_cols + val = tl.load( + logits_row_ptr + cols * stride_logits_v, + mask=mask, other=-float("inf"), + ).to(tl.float32) + z = A * tl.sigmoid(val * inv_C) + z = tl.where(mask, z, -float("inf")) + curr_max = tl.max(z, axis=0) + new_max = tl.maximum(max_val, curr_max) + sum_exp = sum_exp * tl.exp(max_val - new_max) + tl.sum(tl.exp(z - new_max), axis=0) + max_val = new_max + lse = max_val + tl.log(sum_exp) + tl.store(lse_ptr + row_idx, lse) + target = tl.load(targets_ptr + row_idx).to(tl.int32) + target_val = tl.load(logits_row_ptr + target * stride_logits_v).to(tl.float32) + target_z = A * tl.sigmoid(target_val * inv_C) + tl.store(losses_ptr + row_idx, lse - target_z) + + +@triton.jit +def _softcapped_ce_bwd_kernel( + grad_logits_ptr, grad_losses_ptr, lse_ptr, logits_ptr, targets_ptr, + stride_logits_n, stride_logits_v, + stride_grad_n, stride_grad_v, + n_rows, n_cols, softcap, + block_size: tl.constexpr, +): + row_idx = tl.program_id(0).to(tl.int64) + logits_row_ptr = logits_ptr + row_idx * stride_logits_n + grad_row_ptr = grad_logits_ptr + row_idx * stride_grad_n + lse = tl.load(lse_ptr + row_idx) + grad_loss = tl.load(grad_losses_ptr + row_idx).to(tl.float32) + target = tl.load(targets_ptr + row_idx).to(tl.int32) + A = 2.0 * softcap + inv_C = 2.0 / softcap + dz_dx_scale = A * inv_C + for off in range(0, n_cols, block_size): + cols = off + tl.arange(0, block_size) + mask = cols < n_cols + val = tl.load( + logits_row_ptr + cols * stride_logits_v, + mask=mask, other=0.0, + ).to(tl.float32) + sigmoid_u = tl.sigmoid(val * inv_C) + z = A * sigmoid_u + probs = tl.exp(z - lse) + grad_z = grad_loss * (probs - tl.where(cols == target, 1.0, 0.0)) + grad_x = grad_z * (dz_dx_scale * sigmoid_u * (1.0 - sigmoid_u)) + tl.store(grad_row_ptr + cols * stride_grad_v, grad_x, mask=mask) + + +def _validate_softcapped_ce_inputs( + logits: Tensor, targets: Tensor, softcap: float, +) -> tuple[Tensor, Tensor]: + if logits.ndim != 2: + raise ValueError(f"Expected logits.ndim=2, got {logits.ndim}") + if targets.ndim != 1: + raise ValueError(f"Expected targets.ndim=1, got {targets.ndim}") + if logits.shape[0] != targets.shape[0]: + raise ValueError( + f"Expected matching rows, got logits={tuple(logits.shape)} targets={tuple(targets.shape)}" + ) + if not logits.is_cuda or not targets.is_cuda: + raise ValueError("softcapped_cross_entropy requires CUDA tensors") + if softcap <= 0.0: + raise ValueError(f"softcap must be positive, got {softcap}") + if logits.dtype not in (torch.float16, torch.bfloat16, torch.float32): + raise ValueError(f"Unsupported logits dtype: {logits.dtype}") + logits = logits.contiguous() + targets = targets.contiguous() + if targets.dtype != torch.int64: + targets = targets.to(dtype=torch.int64) + return logits, targets + + +@torch.library.custom_op(f"{_FUSED_CE_LIBRARY}::softcapped_ce", mutates_args=()) +def softcapped_ce_op(logits: Tensor, targets: Tensor, softcap: float) -> tuple[Tensor, Tensor]: + logits, targets = _validate_softcapped_ce_inputs(logits, targets, float(softcap)) + n_rows, n_cols = logits.shape + losses = torch.empty((n_rows,), device=logits.device, dtype=torch.float32) + lse = torch.empty((n_rows,), device=logits.device, dtype=torch.float32) + _softcapped_ce_fwd_kernel[(n_rows,)]( + logits, losses, lse, targets, + logits.stride(0), logits.stride(1), + n_rows, n_cols, float(softcap), + block_size=_FUSED_CE_BLOCK_SIZE, num_warps=_FUSED_CE_NUM_WARPS, + ) + return losses, lse + + +@softcapped_ce_op.register_fake +def _(logits: Tensor, targets: Tensor, softcap: float): + if logits.ndim != 2 or targets.ndim != 1: + raise ValueError("softcapped_ce fake impl expects 2D logits and 1D targets") + if logits.shape[0] != targets.shape[0]: + raise ValueError( + f"Expected matching rows, got logits={tuple(logits.shape)} targets={tuple(targets.shape)}" + ) + n_rows = logits.shape[0] + return ( + logits.new_empty((n_rows,), dtype=torch.float32), + logits.new_empty((n_rows,), dtype=torch.float32), + ) + + +@torch.library.custom_op(f"{_FUSED_CE_LIBRARY}::softcapped_ce_backward", mutates_args=()) +def softcapped_ce_backward_op( + logits: Tensor, targets: Tensor, lse: Tensor, grad_losses: Tensor, softcap: float, +) -> Tensor: + logits, targets = _validate_softcapped_ce_inputs(logits, targets, float(softcap)) + lse = lse.contiguous() + grad_losses = grad_losses.contiguous().to(dtype=torch.float32) + if lse.ndim != 1 or grad_losses.ndim != 1: + raise ValueError("Expected 1D lse and grad_losses") + if lse.shape[0] != logits.shape[0] or grad_losses.shape[0] != logits.shape[0]: + raise ValueError( + f"Expected row-aligned lse/grad_losses, got logits={tuple(logits.shape)} " + f"lse={tuple(lse.shape)} grad_losses={tuple(grad_losses.shape)}" + ) + grad_logits = torch.empty_like(logits) + n_rows, n_cols = logits.shape + _softcapped_ce_bwd_kernel[(n_rows,)]( + grad_logits, grad_losses, lse, logits, targets, + logits.stride(0), logits.stride(1), + grad_logits.stride(0), grad_logits.stride(1), + n_rows, n_cols, float(softcap), + block_size=_FUSED_CE_BLOCK_SIZE, num_warps=_FUSED_CE_NUM_WARPS, + ) + return grad_logits + + +@softcapped_ce_backward_op.register_fake +def _(logits: Tensor, targets: Tensor, lse: Tensor, grad_losses: Tensor, softcap: float): + if logits.ndim != 2 or targets.ndim != 1 or lse.ndim != 1 or grad_losses.ndim != 1: + raise ValueError("softcapped_ce_backward fake impl expects 2D logits and 1D row tensors") + if ( + logits.shape[0] != targets.shape[0] + or logits.shape[0] != lse.shape[0] + or logits.shape[0] != grad_losses.shape[0] + ): + raise ValueError("softcapped_ce_backward fake impl expects row-aligned tensors") + return logits.new_empty(logits.shape) + + +def _softcapped_ce_setup_context( + ctx: torch.autograd.function.FunctionCtx, inputs, output, +) -> None: + logits, targets, softcap = inputs + _losses, lse = output + ctx.save_for_backward(logits, targets, lse) + ctx.softcap = float(softcap) + + +def _softcapped_ce_backward( + ctx: torch.autograd.function.FunctionCtx, grad_losses: Tensor, grad_lse: "Tensor | None", +): + del grad_lse + logits, targets, lse = ctx.saved_tensors + grad_logits = torch.ops.pgsubmission1draft7fusedce.softcapped_ce_backward( + logits, targets, lse, grad_losses, ctx.softcap + ) + return grad_logits, None, None + + +softcapped_ce_op.register_autograd( + _softcapped_ce_backward, setup_context=_softcapped_ce_setup_context, +) + + +def softcapped_cross_entropy( + logits: Tensor, targets: Tensor, softcap: float, reduction: str = "mean", +) -> Tensor: + losses, _lse = torch.ops.pgsubmission1draft7fusedce.softcapped_ce( + logits, targets, float(softcap) + ) + if reduction == "none": + return losses + if reduction == "sum": + return losses.sum() + if reduction == "mean": + return losses.mean() + raise ValueError(f"Unsupported reduction={reduction!r}") + + +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.85)) + warmdown_iters = int(os.environ.get("WARMDOWN_ITERS", 0)) + midrun_cap_schedule = os.environ.get("MIDRUN_CAP_SCHEDULE", "").strip() + midrun_cap_log_updates = bool(int(os.environ.get("MIDRUN_CAP_LOG_UPDATES", "0"))) + warmup_steps = int(os.environ.get("WARMUP_STEPS", 20)) + train_batch_tokens = int(os.environ.get("TRAIN_BATCH_TOKENS", 786432)) + # Fused softcapped CE (Triton). Training-only — forward_logits eval path still uses + # eager softcap+F.cross_entropy. Default ON since validated as at-worst neutral. + fused_ce_enabled = bool(int(os.environ.get("FUSED_CE_ENABLED", "1"))) + train_seq_len = int(os.environ.get("TRAIN_SEQ_LEN", 3072)) + train_seq_schedule = os.environ.get("TRAIN_SEQ_SCHEDULE", "1024@0.100,2048@0.700,3072@1.000") + train_seq_schedule_mode = os.environ.get("TRAIN_SEQ_SCHEDULE_MODE", "wallclock").strip().lower() + seq_change_warmup_steps = int(os.environ.get("SEQ_CHANGE_WARMUP_STEPS", 32)) + compile_shape_warmup = bool(int(os.environ.get("COMPILE_SHAPE_WARMUP", "1"))) + compile_shape_warmup_iters = int(os.environ.get("COMPILE_SHAPE_WARMUP_ITERS", "1")) + compile_shape_warmup_loop_modes = os.environ.get("COMPILE_SHAPE_WARMUP_LOOP_MODES", "auto").strip().lower() + 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", 3072)) + val_loss_every = int(os.environ.get("VAL_LOSS_EVERY", 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", 3072)) + 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.1)) + 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.99)) + 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", 1536)) + eval_include_tail = bool(int(os.environ.get("EVAL_INCLUDE_TAIL", "1"))) + 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", 80)) + ttt_lora_lr = float(os.environ.get("TTT_LORA_LR", 0.00015)) + ttt_local_lr_mult = float(os.environ.get("TTT_LOCAL_LR_MULT", 0.75)) + ttt_chunk_size = int(os.environ.get("TTT_CHUNK_SIZE", 48)) + ttt_eval_seq_len = int(os.environ.get("TTT_EVAL_SEQ_LEN", 3072)) + ttt_batch_size = int(os.environ.get("TTT_BATCH_SIZE", 24)) + ttt_grad_steps = int(os.environ.get("TTT_GRAD_STEPS", 1)) + # V19: PR #1886 (renqianluo) + sunnypatneedi research log 2026-04-28 found that + # the Triton fused-CE kernel's fp32-accumulation interacts with warm-start LoRA-A + # to destabilize seeds 314/1337 at TTT_WEIGHT_DECAY=1.0. Raising the default to + # 2.0 prevents seed collapse without measurably moving stable seeds. + ttt_weight_decay = float(os.environ.get("TTT_WEIGHT_DECAY", 0.25)) + ttt_beta1 = float(os.environ.get("TTT_BETA1", 0)) + ttt_beta2 = float(os.environ.get("TTT_BETA2", 0.99)) + ttt_mask = os.environ.get("TTT_MASK", "no_qv").strip().lower() + _ttt_q_default = "1" + _ttt_v_default = "1" + if ttt_mask in ("", "all", "baseline_all"): + pass + elif ttt_mask == "no_q": + _ttt_q_default = "0" + elif ttt_mask == "no_v": + _ttt_v_default = "0" + elif ttt_mask == "no_qv": + _ttt_q_default = "0" + _ttt_v_default = "0" + else: + raise ValueError(f"Unsupported TTT_MASK={ttt_mask!r}") + ttt_q_lora = bool(int(os.environ.get("TTT_Q_LORA", _ttt_q_default))) + ttt_k_lora = bool(int(os.environ.get("TTT_K_LORA", "1"))) + ttt_v_lora = bool(int(os.environ.get("TTT_V_LORA", _ttt_v_default))) + 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", "") + ttt_peer_ensemble_k = int(os.environ.get("TTT_PEER_ENSEMBLE_K", "4")) + ttt_peer_conf_threshold = float(os.environ.get("TTT_PEER_CONF_THRESHOLD", "0.5")) + ttt_peer_conf_blend_w = float(os.environ.get("TTT_PEER_CONF_BLEND_W", "0.8")) + ttt_short_doc_len = int(os.environ.get("TTT_SHORT_DOC_LEN", 2000)) + ttt_short_lora_enabled = bool(int(os.environ.get("TTT_SHORT_LORA_ENABLED", "0"))) + ttt_short_lora_rank = int(os.environ.get("TTT_SHORT_LORA_RANK", ttt_lora_rank)) + ttt_short_lora_lr = float(os.environ.get("TTT_SHORT_LORA_LR", ttt_lora_lr)) + ttt_short_weight_decay = float(os.environ.get("TTT_SHORT_WEIGHT_DECAY", ttt_weight_decay)) + ttt_short_beta2 = float(os.environ.get("TTT_SHORT_BETA2", ttt_beta2)) + ttt_short_score_first_enabled = bool(int(os.environ.get("TTT_SHORT_SCORE_FIRST_ENABLED", "1"))) + ttt_short_chunk_size = int(os.environ.get("TTT_SHORT_CHUNK_SIZE", 24)) + ttt_short_score_first_steps = os.environ.get("TTT_SHORT_SCORE_FIRST_STEPS", "256:8,2000:24") + ttt_train_min_doc_len = int(os.environ.get("TTT_TRAIN_MIN_DOC_LEN", "0")) + ttt_train_max_doc_len = int(os.environ.get("TTT_TRAIN_MAX_DOC_LEN", "0")) + ttt_warm_start_mean_enabled = bool(int(os.environ.get("TTT_WARM_START_MEAN_ENABLED", "0"))) + ttt_warm_start_mean_doc_len = int(os.environ.get("TTT_WARM_START_MEAN_DOC_LEN", ttt_short_doc_len)) + ttt_warm_start_mean_momentum = float(os.environ.get("TTT_WARM_START_MEAN_MOMENTUM", 0.95)) + val_doc_fraction = float(os.environ.get("VAL_DOC_FRACTION", 1.0)) + compressor = os.environ.get("COMPRESSOR", "pergroup") + 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_prefix_docs = int(os.environ.get("PHASED_TTT_PREFIX_DOCS", 2500)) + phased_ttt_num_phases = int(os.environ.get("PHASED_TTT_NUM_PHASES", 1)) + 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", 14.0)) + mlp_clip_sigmas = float(os.environ.get("MLP_CLIP_SIGMAS", 11.5)) + attn_clip_sigmas = float(os.environ.get("ATTN_CLIP_SIGMAS", 13.0)) + # AttnOutGate (per-head multiplicative output gate, PR #1667 MarioPaerle). + # Zero-init weight: 2*sigmoid(0)=1 -> transparent at start. Source defaults to + # block input x ('proj'); 'q' uses raw Q projection output. + attn_out_gate_enabled = bool(int(os.environ.get("ATTN_OUT_GATE_ENABLED", "0"))) + attn_out_gate_src = os.environ.get("ATTN_OUT_GATE_SRC", "proj") + # SmearGate (input-dependent forward-1 token smear, modded-nanogpt @classiclarryd + # via PR #1667). x_t <- x_t + lam * sigmoid(W*x_t[:gate_window]) * x_{t-1}. + # lam=0 + W=0 -> transparent at init. + smear_gate_enabled = bool(int(os.environ.get("SMEAR_GATE_ENABLED", "1"))) + # Window: first GATE_WINDOW dims of the source feed the gate projection. + gate_window = int(os.environ.get("GATE_WINDOW", 12)) + # Gated Attention (Qwen, NeurIPS 2025 Best Paper, arXiv:2505.06708; + # qiuzh20/gated_attention). Per-head sigmoid gate on SDPA output, BEFORE + # out_proj. Gate input = full block input x (paper's headwise G1 variant + # driven from hidden_states). W_g shape (num_heads, dim), plain sigmoid. + # Near-zero init gives g~0.5 at step 0 (half attention output); per-block + # attn_scale (init 1.0) compensates during training. Name contains + # "attn_gate" so CONTROL_TENSOR_NAME_PATTERNS routes it to scalar AdamW. + gated_attn_enabled = bool(int(os.environ.get("GATED_ATTN_ENABLED", "0"))) + gated_attn_init_std = float(os.environ.get("GATED_ATTN_INIT_STD", 0.01)) + # Dedicated int8-per-row quantization for `attn_gate_w` tensors. These are + # small ((num_heads, dim) = (8, 512) = 4096 params) and bypass GPTQ via the + # numel<=65536 passthrough branch -> stored as fp16 (8 KB/layer, ~65 KB total + # compressed). int8-per-row cuts the raw tensor in half with negligible BPB + # impact: scales per head (8 values), symmetric quant over [-127, 127]. + # No Hessian needed (gate weights not in collect_hessians()). + gated_attn_quant_gate = bool(int(os.environ.get("GATED_ATTN_QUANT_GATE", "1"))) + # Sparse Attention Gate (modded-nanogpt-style). Keeps dense SDPA and only + # swaps the output-gate input to the first GATE_WINDOW residual dims. + # W_g: (num_heads, gate_window) = (8, 12) = 96 params/layer (~44K total), + # vs dense GatedAttn's (8, 512) = 4K/layer (~44K diff). Name "attn_gate_w" + # is shared so quant routing and int8 gate passthrough Just Work. Gate + # passthrough int8 still applies via GATED_ATTN_QUANT_GATE=1. + # Mutually exclusive with ATTN_OUT_GATE_ENABLED and GATED_ATTN_ENABLED. + sparse_attn_gate_enabled = bool(int(os.environ.get("SPARSE_ATTN_GATE_ENABLED", "1"))) + sparse_attn_gate_init_std = float(os.environ.get("SPARSE_ATTN_GATE_INIT_STD", 0.0)) + sparse_attn_gate_scale = float(os.environ.get("SPARSE_ATTN_GATE_SCALE", 0.5)) + # LQER asymmetric rank-k correction on top-K quant-error tensors (PR #1530 v2 port). + # Computes SVD of E = W_fp - W_quant, packs top-r A,B as INT2/INT4 (asym) or INTk (sym). + lqer_enabled = bool(int(os.environ.get("LQER_ENABLED", "1"))) + 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")) + lqer_scope = os.environ.get("LQER_SCOPE", "all") + lqer_gain_select = bool(int(os.environ.get("LQER_GAIN_SELECT", "0"))) + awq_lite_enabled = bool(int(os.environ.get("AWQ_LITE_ENABLED", "1"))) + awq_lite_bits = int(os.environ.get("AWQ_LITE_BITS", "8")) + awq_lite_group_top_k = int(os.environ.get("AWQ_LITE_GROUP_TOP_K", "1")) + awq_lite_group_size = int(os.environ.get("AWQ_LITE_GROUP_SIZE", "64")) + 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 + # CaseOps integration: optional override of dataset root + tokenizer path. + # When CASEOPS_ENABLED=1, the wrapper loads a per-token byte sidecar + # (fineweb_val_bytes_*.bin, identical shard layout to val_*.bin) and uses + # it as the canonical raw-byte budget for BPB accounting. The sidecar + # REPLACES the build_sentencepiece_luts byte-counting path entirely. + caseops_enabled = bool(int(os.environ.get("CASEOPS_ENABLED", "1"))) + _default_caseops_data = os.path.join( + data_dir, + "datasets", + "fineweb10B_sp8192_caseops", + "datasets", + "datasets", + "fineweb10B_sp8192_lossless_caps_caseops_v1_reserved", + ) + _default_caseops_tok = os.path.join( + data_dir, + "datasets", + "fineweb10B_sp8192_caseops", + "datasets", + "tokenizers", + "fineweb_8192_bpe_lossless_caps_caseops_v1_reserved.model", + ) + if caseops_enabled: + datasets_dir = os.environ.get("DATA_PATH", _default_caseops_data) + tokenizer_path = os.environ.get("TOKENIZER_PATH", _default_caseops_tok) + else: + datasets_dir = os.environ.get( + "DATA_PATH", + os.path.join(data_dir, "datasets", f"fineweb10B_sp{vocab_size}"), + ) + tokenizer_path = os.environ.get( + "TOKENIZER_PATH", + os.path.join(data_dir, "tokenizers", f"fineweb_{vocab_size}_bpe.model"), + ) + train_files = os.path.join(datasets_dir, "fineweb_train_*.bin") + val_files = os.path.join(datasets_dir, "fineweb_val_*.bin") + val_bytes_files = os.path.join(datasets_dir, "fineweb_val_bytes_*.bin") + 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) + + +def parse_train_seq_schedule(schedule, default_seq_len): + if not schedule.strip(): + return [(1.0, int(default_seq_len))] + plan = [] + for raw_stage in schedule.split(","): + raw_stage = raw_stage.strip() + if not raw_stage: + continue + if "@" not in raw_stage: + raise ValueError( + f"Invalid TRAIN_SEQ_SCHEDULE stage `{raw_stage}`; expected format like `1024@0.35`" + ) + seq_raw, progress_raw = raw_stage.split("@", 1) + seq_len = int(seq_raw.strip()) + progress = float(progress_raw.strip()) + if seq_len <= 0: + raise ValueError("TRAIN_SEQ_SCHEDULE sequence lengths must be positive") + if not (0.0 < progress <= 1.0): + raise ValueError("TRAIN_SEQ_SCHEDULE progress fractions must be in (0, 1]") + plan.append((progress, seq_len)) + if not plan: + return [(1.0, int(default_seq_len))] + plan.sort(key=lambda item: item[0]) + if plan[-1][0] < 1.0: + plan.append((1.0, plan[-1][1])) + return plan + + +def parse_scalar_schedule(schedule, default_value): + if not schedule.strip(): + return [(0.0, float(default_value))] + plan = [] + for raw_stage in schedule.split(","): + raw_stage = raw_stage.strip() + if not raw_stage: + continue + if "@" not in raw_stage: + raise ValueError( + f"Invalid scalar schedule stage `{raw_stage}`; expected format like `0.5@0.4`" + ) + value_raw, progress_raw = raw_stage.split("@", 1) + value = float(value_raw.strip()) + progress = float(progress_raw.strip()) + if not (0.0 <= progress <= 1.0): + raise ValueError("Scalar schedule progress fractions must be in [0, 1]") + plan.append((progress, value)) + if not plan: + return [(0.0, float(default_value))] + plan.sort(key=lambda item: item[0]) + if plan[0][0] > 0.0: + plan.insert(0, (0.0, plan[0][1])) + return plan + + +def schedule_value(plan, progress): + value = plan[0][1] + for threshold, candidate in plan: + if progress + 1e-12 >= threshold: + value = candidate + else: + break + return value + + +def max_train_seq_len_from_schedule(plan, default_seq_len): + return max([int(default_seq_len), *[seq_len for _, seq_len in plan]]) + + +def validate_train_seq_plan_compatibility( + plan, + *, + global_tokens, + world_size, + grad_accum_steps, +): + denom = world_size * grad_accum_steps + if denom <= 0: + raise ValueError(f"Invalid world_size * grad_accum_steps={denom}") + if global_tokens % denom != 0: + raise ValueError( + f"TRAIN_BATCH_TOKENS={global_tokens} must be divisible by world_size*grad_accum_steps={denom}" + ) + local_tokens = global_tokens // denom + invalid_seq_lens = sorted( + {seq_len for _, seq_len in plan if local_tokens % seq_len != 0} + ) + if invalid_seq_lens: + raise ValueError( + "TRAIN_SEQ_SCHEDULE contains sequence lengths incompatible with the local micro-batch: " + f"local_tokens={local_tokens}, invalid_seq_lens={invalid_seq_lens}. " + f"Each seq_len must divide {local_tokens} exactly." + ) + return local_tokens + + +def training_progress( + *, + step, + iterations, + elapsed_ms, + max_wallclock_ms, + schedule_mode, +): + if schedule_mode == "step" or max_wallclock_ms is None or max_wallclock_ms <= 0: + return min(max(step / max(iterations, 1), 0.0), 1.0) + if schedule_mode != "wallclock": + raise ValueError( + f"Unsupported TRAIN_SEQ_SCHEDULE_MODE={schedule_mode!r}; expected 'wallclock' or 'step'" + ) + return min(max(elapsed_ms / max(max_wallclock_ms, 1e-9), 0.0), 1.0) + + +def current_train_seq_len( + plan, + *, + step, + iterations, + elapsed_ms, + max_wallclock_ms, + schedule_mode, +): + progress = training_progress( + step=step, + iterations=iterations, + elapsed_ms=elapsed_ms, + max_wallclock_ms=max_wallclock_ms, + schedule_mode=schedule_mode, + ) + for threshold, seq_len in plan: + if progress <= threshold: + return seq_len, progress + return plan[-1][1], progress + + +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, include_tail=h.eval_include_tail + ) + self.caseops_enabled = bool(getattr(h, "caseops_enabled", False)) + if self.caseops_enabled: + self.base_bytes_lut = None + self.has_leading_space_lut = None + self.is_boundary_token_lut = None + else: + ( + self.base_bytes_lut, + self.has_leading_space_lut, + self.is_boundary_token_lut, + ) = build_sentencepiece_luts(self.sp, h.vocab_size, device) + self.val_bytes = None + if self.caseops_enabled: + self.val_bytes = load_validation_byte_sidecar( + h.val_bytes_files, h.eval_seq_len, self.val_tokens.numel() + ) + + +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_) + for token_id in range(sp_vocab_size): + if sp.is_control(token_id) or sp.is_unknown(token_id) or sp.is_unused(token_id): + continue + is_boundary_token_np[token_id] = False + if sp.is_byte(token_id): + base_bytes_np[token_id] = 1 + continue + piece = sp.id_to_piece(token_id) + if piece.startswith("▁"): + has_leading_space_np[token_id] = True + piece = piece[1:] + base_bytes_np[token_id] = len(piece.encode("utf-8")) + return ( + torch.tensor(base_bytes_np, dtype=torch.int16, device=device), + torch.tensor(has_leading_space_np, dtype=torch.bool, device=device), + torch.tensor(is_boundary_token_np, dtype=torch.bool, device=device), + ) + + +def load_validation_tokens(pattern, seq_len, include_tail=True): + # Filter out CaseOps byte sidecar shards which share the val_*.bin glob. + files = [ + Path(p) + for p in sorted(glob.glob(pattern)) + if "_bytes_" not in Path(p).name + ] + if not files: + raise FileNotFoundError(f"No files found for pattern: {pattern}") + tokens = torch.cat([load_data_shard(file) for file in files]).contiguous() + if include_tail: + if tokens.numel() <= 1: + raise ValueError(f"Validation split is too short for TRAIN_SEQ_LEN={seq_len}") + return tokens + 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_validation_byte_sidecar(pattern, seq_len, expected_len): + """Load CaseOps per-token byte sidecar(s). Same shard layout as token shards + (256 int32 header + uint16 array). Each entry = canonical raw-text byte + budget for that token in the corresponding val shard. Returns a CPU + int16 tensor sliced to match expected_len (i.e. val_tokens length).""" + files = [Path(p) for p in sorted(glob.glob(pattern))] + if not files: + raise FileNotFoundError(f"No byte sidecar files for pattern: {pattern}") + shards = [load_data_shard(file) for file in files] + # load_data_shard returns uint16 — that's exactly what the sidecar stores. + bytes_full = torch.cat(shards).contiguous() + if bytes_full.numel() < expected_len: + raise ValueError( + f"Byte sidecar too short: {bytes_full.numel()} < val_tokens {expected_len}" + ) + return bytes_full[:expected_len].to(torch.int32) + + +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._prefetch_queue = [] + + 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() + ) + 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 = torch.empty(per_rank_span - 1, dtype=torch.int64, pin_memory=True) + targets = torch.empty(per_rank_span - 1, dtype=torch.int64, pin_memory=True) + inputs.copy_(buf[:-1]) + targets.copy_(buf[1:]) + 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, max_seq_len=None): + if max_seq_len is None: + max_seq_len = self.max_seq_len + max_seq_len = int(max_seq_len) + if max_seq_len != self.max_seq_len: + self.max_seq_len = max_seq_len + self._prefetch_queue.clear() + num_tokens_local = global_tokens // (self.world_size * grad_accum_steps) + while len(self._prefetch_queue) < 2: + self._prefetch_queue.append( + self._batch_pool.submit(self._prepare_batch, num_tokens_local, self.max_seq_len)) + inputs, targets, cu_seqlens, max_seqlen = self._prefetch_queue.pop(0).result() + self._prefetch_queue.append( + 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 = 256, 128, 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) + + +class CausalSelfAttention(nn.Module): + def __init__( + self, dim, num_heads, num_kv_heads, rope_base, qk_gain_init, train_seq_len, yarn=True, + attn_out_gate=False, attn_out_gate_src="proj", gate_window=12, + gated_attn=False, gated_attn_init_std=0.01, + sparse_attn_gate=False, sparse_attn_gate_init_std=0.0, sparse_attn_gate_scale=1.0, + ): + 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") + if int(attn_out_gate) + int(gated_attn) + int(sparse_attn_gate) > 1: + raise ValueError( + "attn_out_gate, gated_attn, and sparse_attn_gate are mutually exclusive" + ) + 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 + # AttnOutGate (PR #1667 MarioPaerle): per-head multiplicative gate on attention + # output. CastedLinear so restore_fp32_params casts back to fp32 for GPTQ. + # _zero_init -> 2*sigmoid(0)=1 -> transparent at init. + self.attn_out_gate = attn_out_gate + self.attn_out_gate_src = attn_out_gate_src + self.gate_window = gate_window + if attn_out_gate: + self.attn_gate_proj = CastedLinear(gate_window, num_heads, bias=False) + self.attn_gate_proj._zero_init = True + # Gated Attention (arXiv:2505.06708, Qwen, NeurIPS 2025). Per-head sigmoid + # gate on SDPA output, BEFORE out_proj. Gate projection W_g: (num_heads, dim). + # Name "attn_gate_w" contains "attn_gate" substring so it matches + # CONTROL_TENSOR_NAME_PATTERNS and routes to the scalar AdamW group. + # fp32 Parameter -> restore_fp32_params path covers it via the ndim<2 OR + # name-pattern check (name matches "attn_gate"). Cast to x.dtype on use. + self.gated_attn = gated_attn + if gated_attn: + W = torch.empty(num_heads, dim, dtype=torch.float32) + nn.init.normal_(W, mean=0.0, std=gated_attn_init_std) + self.attn_gate_w = nn.Parameter(W) + # Sparse attention head-output gate (modded-nanogpt style). Keeps dense SDPA + # and only narrows the gate input to the first gate_window residual dims. + # W_g: (num_heads, gate_window). y_{t,h} <- sigmoid(scale * W_g_h @ x_t[:gate_window]) * y_{t,h}. + # Shares attn_gate_w name with dense GatedAttn so the quant routing + # (CONTROL_TENSOR_NAME_PATTERNS / attn_gate_w int8 passthrough) is unchanged. + self.sparse_attn_gate = sparse_attn_gate + self.sparse_attn_gate_scale = sparse_attn_gate_scale + if sparse_attn_gate: + W = torch.empty(num_heads, gate_window, dtype=torch.float32) + if sparse_attn_gate_init_std > 0: + nn.init.normal_(W, mean=0.0, std=sparse_attn_gate_init_std) + else: + nn.init.zeros_(W) + self.attn_gate_w = nn.Parameter(W) + + 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_raw kept around as a tap point for attn_out_gate_src='q' (post-projection, + # pre-reshape, pre-RoPE). + q_raw = F.linear(x, q_w.to(x.dtype)) + q = q_raw.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_func( + q[0], + k[0], + v[0], + 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), + )[None] + else: + y = flash_attn_3_func(q, k, v, causal=True) + if self.use_xsa: + y = self._xsa_efficient(y, v) + # AttnOutGate inlined (PR #1667). Inline + .contiguous() barrier so torch.compile + # fullgraph=True is happy (this avoids the @torch.compiler.disable trap that + # crashed gates v3). Per-head gate on (B,T,H,D) tensor: g shape [B,T,H], broadcast + # over D via [..., None]. zero-init weight -> 2*sigmoid(0)=1 -> transparent. + if self.attn_out_gate: + gate_src = q_raw if self.attn_out_gate_src == "q" else x + gate_in = gate_src[..., : self.gate_window].contiguous() + g = 2.0 * torch.sigmoid(self.attn_gate_proj(gate_in)) + y = y * g[..., None] + # Gated Attention (arXiv:2505.06708 G1). Inline + .contiguous() barrier so + # torch.compile fullgraph=True is happy. Per-head gate on (B,T,H,D): g shape + # [B,T,H], broadcast over D via [..., None]. Paper: g = sigmoid(x @ W_g.T) + # where W_g: (H, dim). .to(x.dtype) on fp32 param before broadcast with bf16. + if self.gated_attn: + x_c = x.contiguous() + g = torch.sigmoid(F.linear(x_c, self.attn_gate_w.to(x.dtype))) + y = y * g[..., None] + # Sparse head-output gate: narrower (gate_window) input, same shape g as GatedAttn. + if 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[..., None] + 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 = True + + 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, + attn_out_gate=False, + attn_out_gate_src="proj", + gate_window=12, + gated_attn=False, + gated_attn_init_std=0.01, + sparse_attn_gate=False, + sparse_attn_gate_init_std=0.0, + sparse_attn_gate_scale=1.0, + ): + super().__init__() + self.attn_norm = RMSNorm() + self.mlp_norm = RMSNorm() + self.attn = CausalSelfAttention( + dim, num_heads, num_kv_heads, rope_base, qk_gain_init, train_seq_len, yarn=yarn, + attn_out_gate=attn_out_gate, attn_out_gate_src=attn_out_gate_src, gate_window=gate_window, + gated_attn=gated_attn, gated_attn_init_std=gated_attn_init_std, + sparse_attn_gate=sparse_attn_gate, + sparse_attn_gate_init_std=sparse_attn_gate_init_std, + sparse_attn_gate_scale=sparse_attn_gate_scale, + ) + 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.fused_ce_enabled = bool(h.fused_ce_enabled) + 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, + attn_out_gate=h.attn_out_gate_enabled, + attn_out_gate_src=h.attn_out_gate_src, + gate_window=h.gate_window, + gated_attn=h.gated_attn_enabled, + gated_attn_init_std=h.gated_attn_init_std, + sparse_attn_gate=h.sparse_attn_gate_enabled, + sparse_attn_gate_init_std=h.sparse_attn_gate_init_std, + sparse_attn_gate_scale=h.sparse_attn_gate_scale, + ) + 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 + 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) + ) + # SmearGate (PR #1667 / modded-nanogpt @classiclarryd): + # x_t <- x_t + lam * sigmoid(W * x_t[:gate_window]) * x_{t-1}. + # Per-token forward-1 smear of the embedding lane. W zero-init + lam=0 -> + # transparent at init. Uses CastedLinear so restore_fp32_params handles dtype. + self.smear_gate_enabled = h.smear_gate_enabled + if self.smear_gate_enabled: + self.smear_window = h.gate_window + self.smear_gate = CastedLinear(self.smear_window, 1, bias=False) + self.smear_gate._zero_init = True + self.smear_lambda = nn.Parameter(torch.zeros(1, dtype=torch.float32)) + # V19: Asymmetric Logit Rescale (PR #1923 jorge-asenjo). + # Two learnable softcap scales applied on the EVAL path (forward_logits + + # forward_ttt). Init to logit_softcap so the layer is identity at step 0. + # Train path keeps the single fused softcap to preserve PR #1855 numerics. + self.asym_logit_enabled = bool(int(os.environ.get("ASYM_LOGIT_RESCALE", "0"))) + if self.asym_logit_enabled: + self.softcap_pos = nn.Parameter(torch.tensor(float(h.logit_softcap), dtype=torch.float32)) + self.softcap_neg = nn.Parameter(torch.tensor(float(h.logit_softcap), 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_hidden(self, input_ids, cu_seqlens=None, max_seqlen=0): + """Run the encoder/decoder stack to the final RMSNorm; returns pre-projection hidden. + Shared by eval (softcap+projection via forward_logits) and train (fused CE path).""" + x = self.tok_emb(input_ids) + # SmearGate (PR #1667). lam=0 + W=0 -> identity at init. + # Cross-doc leak fix: zero the prev-token smear at any position whose current token + # is BOS, so the BOS embedding starting doc N+1 in a packed stream is not + # contaminated by doc N's last token (audited issue on PR#1797 base). + if self.smear_gate_enabled: + sl = self.smear_lambda.to(dtype=x.dtype) + gate_in = x[:, 1:, : self.smear_window].contiguous() + g = sl * torch.sigmoid(self.smear_gate(gate_in)) + not_bos = (input_ids[:, 1:] != BOS_ID).to(x.dtype).unsqueeze(-1) + x = torch.cat([x[:, :1], x[:, 1:] + g * x[:, :-1] * not_bos], 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) + return x + + def _project_logits(self, hidden): + if self.tie_embeddings: + return F.linear(hidden, self.tok_emb.weight) + return self.lm_head(hidden) + + def _apply_asym_softcap(self, logits): + # V19: Asymmetric softcap (PR #1923). Splits the logit_softcap scalar into + # learnable positive/negative branches. Score-first preserved: still a + # bounded, normalized post-projection nonlinearity feeding a standard + # softmax over the full vocab. + sp = self.softcap_pos.to(logits.dtype) + sn = self.softcap_neg.to(logits.dtype) + return torch.where(logits > 0, sp * torch.tanh(logits / sp), sn * torch.tanh(logits / sn)) + + def forward_logits(self, input_ids, cu_seqlens=None, max_seqlen=0): + hidden = self._forward_hidden(input_ids, cu_seqlens=cu_seqlens, max_seqlen=max_seqlen) + logits_proj = self._project_logits(hidden) + if self.asym_logit_enabled: + return self._apply_asym_softcap(logits_proj) + return self.logit_softcap * torch.tanh(logits_proj / self.logit_softcap) + + def forward(self, input_ids, target_ids, cu_seqlens=None, max_seqlen=0): + hidden = self._forward_hidden(input_ids, cu_seqlens=cu_seqlens, max_seqlen=max_seqlen) + logits_proj = self._project_logits(hidden) + flat_targets = target_ids.reshape(-1) + # Fused softcapped-CE kernel (training path only). Applies softcap inside the + # Triton kernel; takes pre-softcap logits_proj. Non-fused path matches stock + # PR-1736 numerics exactly (softcap in fp32, then F.cross_entropy on fp32). + if self.fused_ce_enabled: + return softcapped_cross_entropy( + logits_proj.reshape(-1, logits_proj.size(-1)), + flat_targets, + self.logit_softcap, + reduction="mean", + ) + logits = self.logit_softcap * torch.tanh(logits_proj / self.logit_softcap) + return F.cross_entropy( + logits.reshape(-1, logits.size(-1)).float(), + flat_targets, + reduction="mean", + ) + + def forward_ttt(self, input_ids, target_ids, lora, return_entropy: bool = False): + x = self.tok_emb(input_ids) + # SmearGate on the TTT path — same inline compute as forward_logits. + # Cross-doc leak fix: see _forward_hidden comment. + if self.smear_gate_enabled: + sl = self.smear_lambda.to(dtype=x.dtype) + gate_in = x[:, 1:, : self.smear_window].contiguous() + g = sl * torch.sigmoid(self.smear_gate(gate_in)) + not_bos = (input_ids[:, 1:] != BOS_ID).to(x.dtype).unsqueeze(-1) + x = torch.cat([x[:, :1], x[:, 1:] + g * x[:, :-1] * not_bos], 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) + # V19: same asymmetric softcap on the TTT eval path. + if self.asym_logit_enabled: + logits = self._apply_asym_softcap(logits) + else: + logits = self.logit_softcap * torch.tanh(logits / self.logit_softcap) + bsz, sl, V = logits.shape + logits_f = logits.float() + per_tok_loss = F.cross_entropy( + logits_f.reshape(-1, V), target_ids.reshape(-1), reduction="none" + ).reshape(bsz, sl) + if return_entropy: + log_probs = F.log_softmax(logits_f, dim=-1) + entropy = -(log_probs.exp() * log_probs).sum(dim=-1) + return per_tok_loss, entropy + return per_tok_loss + + 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 + # Keep raw Q for AttnOutGate src='q' (matches forward path semantics). + q_raw = F.linear(n, q_w.to(n.dtype)) + if lora.q_loras is not None: + q_raw = q_raw + lora.q_loras[slot](n) + q = q_raw.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)) + if lora.v_loras is not None: + v = v + lora.v_loras[slot](n) + v = v.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) + # AttnOutGate (TTT path) — inline + .contiguous() barrier, same as the eval path. + if attn.attn_out_gate: + gate_src = q_raw if attn.attn_out_gate_src == "q" else n + gate_in = gate_src[..., : attn.gate_window].contiguous() + g = 2.0 * torch.sigmoid(attn.attn_gate_proj(gate_in)) + y = y * g[..., None] + # Gated Attention (TTT path). Gate input is n (post-norm block input), same + # as eval path. .to(n.dtype) on fp32 param before bf16 broadcast. + if attn.gated_attn: + n_c = n.contiguous() + g = torch.sigmoid(F.linear(n_c, attn.attn_gate_w.to(n.dtype))) + y = y * g[..., None] + # Sparse attention head-output gate (TTT path) — must match the eval path in + # forward() exactly, else training (which applied the gate) and TTT eval (which + # skipped it) produce mismatched representations and catastrophic BPB regression. + if 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[..., None] + 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_raw = F.linear(n, q_w.to(n.dtype)) + if lora.q_loras is not None: + q_raw = q_raw + lora.q_loras[slot](n) + q = q_raw.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)) + if lora.v_loras is not None: + v = v + lora.v_loras[slot](n) + v = v.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) + # AttnOutGate (TTT parallel path) — inline + .contiguous() barrier. + if attn.attn_out_gate: + gate_src = q_raw if attn.attn_out_gate_src == "q" else n + gate_in = gate_src[..., : attn.gate_window].contiguous() + g = 2.0 * torch.sigmoid(attn.attn_gate_proj(gate_in)) + y = y * g[..., None] + # Gated Attention (TTT parallel path). Gate input is n (post-norm block input). + if attn.gated_attn: + n_c = n.contiguous() + g = torch.sigmoid(F.linear(n_c, attn.attn_gate_w.to(n.dtype))) + y = y * g[..., None] + # Sparse attention head-output gate (TTT parallel path) — must match the + # eval path in forward() to keep train/eval semantics in sync. + if 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[..., None] + 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")) + _WARM_START_A = bool(int(os.environ.get("TTT_WARM_START_A", "1"))) + PEER_IDX: "torch.Tensor | None" = None + + 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)) + + 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): + peer_idx = BatchedLinearLoRA.PEER_IDX + if peer_idx is None: + A, B = self.A, self.B + else: + A = self.A[peer_idx] + B = self.B[peer_idx] + return ((x @ A.transpose(1, 2)) @ B.transpose(1, 2)) * self._scale + + +class BatchedTTTLoRA(nn.Module): + def __init__( + self, bsz, model, rank, + q_lora=True, k_lora=True, v_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)] + ) + if q_lora + else None + ) + self.v_loras = ( + nn.ModuleList( + [BatchedLinearLoRA(bsz, dim, kv_dim, rank) for _ in range(num_slots)] + ) + if v_lora + else None + ) + 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). +# Replaces the fixed (3.4445, -4.775, 2.0315) coefficients of stock Muon. +# Applied at backend_steps=5 — taking more than 5 iterations from this list +# falls back to the final (converged) tuple via the slice guard below. +_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), +) + + +@torch.compile +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 = _PE_COEFFS[:steps] if steps <= len(_PE_COEFFS) else _PE_COEFFS + 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) + 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, 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, 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, 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,smear_gate,smear_lambda", + ).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) + # SmearGate params live on GPT root (not in .blocks), so add them by hand. + # Both are tiny (gate_window scalars + 1 lambda). Optimized via scalar Adam. + 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) + self._aux_stream = torch.cuda.Stream() + + 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._aux_stream.wait_stream(torch.cuda.current_stream()) + with torch.cuda.stream(self._aux_stream): + self.optimizer_tok.step() + self.optimizer_scalar.step() + self.optimizer_muon.step() + torch.cuda.current_stream().wait_stream(self._aux_stream) + 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 = {} + act_sumsq = {} + act_counts = {} + 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]) + x_sq = x.square().sum(dim=0) + x_count = x.shape[0] + 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) + if name not in act_sumsq: + act_sumsq[name] = torch.zeros( + x.shape[1], dtype=torch.float32, device=device + ) + act_counts[name] = 0 + act_sumsq[name] += x_sq + act_counts[name] += x_count + 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) + if name not in act_sumsq: + act_sumsq[name] = torch.zeros( + y.shape[1], dtype=torch.float32, device=device + ) + act_counts[name] = 0 + act_sumsq[name] += y.square().sum(dim=0) + act_counts[name] += y.shape[0] + 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) + if name not in act_sumsq: + act_sumsq[name] = torch.zeros( + x.shape[1], dtype=torch.float32, device=device + ) + act_counts[name] = 0 + act_sumsq[name] += x.square().sum(dim=0) + act_counts[name] += x.shape[0] + 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) + if name not in act_sumsq: + act_sumsq[name] = torch.zeros( + h_act.shape[1], dtype=torch.float32, device=device + ) + act_counts[name] = 0 + act_sumsq[name] += h_act.square().sum(dim=0) + act_counts[name] += h_act.shape[0] + 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) + if name not in act_sumsq: + act_sumsq[name] = torch.zeros( + x.shape[1], dtype=torch.float32, device=device + ) + act_counts[name] = 0 + act_sumsq[name] += x.square().sum(dim=0) + act_counts[name] += x.shape[0] + 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 + act_stats = {} + for name, sumsq in act_sumsq.items(): + count = max(act_counts.get(name, 0), 1) + act_stats[name] = (sumsq / count).sqrt().cpu() + return hessians, act_stats + + +def gptq_quantize_weight( + w, + H, + clip_sigmas=3.0, + clip_range=63, + block_size=128, + protect_groups=None, + group_size=None, + protect_clip_range=None, +): + 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() + protect_meta = None + protect_mask_perm = None + s_hi = None + sf_hi = None + if ( + protect_groups + and group_size is not None + and protect_clip_range is not None + and protect_clip_range > clip_range + ): + protect_mask = torch.zeros(cols, dtype=torch.bool) + starts = [] + for (start, end) in protect_groups: + if start < 0 or end > cols or end <= start: + continue + protect_mask[start:end] = True + starts.append(start) + if starts: + protect_mask_perm = protect_mask[perm] + s_hi = (clip_sigmas * row_std / protect_clip_range).clamp_min(1e-10).to( + torch.float16 + ) + sf_hi = s_hi.float() + protect_meta = { + "starts": torch.tensor(starts, dtype=torch.int16), + "size": int(group_size), + "s_hi": s_hi, + } + 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] + if protect_mask_perm is not None and bool(protect_mask_perm[i1 + j]): + q_col = torch.clamp( + torch.round(w_col / sf_hi), + -protect_clip_range, + protect_clip_range, + ) + w_recon = q_col.float() * sf_hi + else: + q_col = torch.clamp(torch.round(w_col / sf), -clip_range, clip_range) + w_recon = q_col.float() * sf + Q[:, i1 + j] = q_col.to(torch.int8) + err = (w_col - w_recon) / 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, protect_meta + + +def _quantize_gate_int8_row(w): + # Symmetric int8-per-row quantization for small gate tensors. w shape + # (R, C) -> (R,) scales in fp16, int8 values in [-127, 127]. Single scale + # per row keeps accuracy high while halving storage vs fp16. + W = w.float().contiguous() + row_max = W.abs().amax(dim=1).clamp_min(1e-10) + s = (row_max / 127.0).to(torch.float16) + sf = s.float().view(-1, 1) + q = torch.clamp(torch.round(W / sf), -127, 127).to(torch.int8) + return q, 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 g 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 _lqer_fit_quantized(E, h): + U, S, Vh = torch.linalg.svd(E, full_matrices=False) + r = min(h.lqer_rank, S.numel()) + if r <= 0: + return None + A = (U[:, :r] * S[:r]).contiguous() + B = Vh[:r, :].contiguous() + asym_on = bool(getattr(h, "lqer_asym_enabled", False)) + asym_g = int(getattr(h, "lqer_asym_group", 64)) + if asym_on and B.numel() % asym_g == 0: + qA, sA, qB, sB = _lqer_pack_asym(A, B, asym_g) + A_hat = qA.float() * float(sA) + g_sz = qB.numel() // sB.numel() + B_hat = (qB.reshape(-1, g_sz).float() * sB.float().view(-1, 1)).reshape( + qB.shape + ) + return { + "kind": "asym", + "qA": qA, + "sA": sA, + "qB": qB, + "sB": sB, + "delta": A_hat @ B_hat, + } + qA, sA, qB, sB = _lqer_pack(A, B, h.lqer_factor_bits) + A_hat = qA.float() * sA.float().view(-1, 1) + B_hat = qB.float() * sB.float().view(-1, 1) + return { + "kind": "sym", + "qA": qA, + "sA": sA, + "qB": qB, + "sB": sB, + "delta": A_hat @ B_hat, + } + + +def _awq_lite_group_candidates(w, act_rms, group_size): + cols = w.shape[1] + n_groups = cols // group_size + if n_groups <= 0: + return [] + weight_score = w.float().abs().mean(dim=0) + saliency = act_rms.float() * weight_score + cands = [] + for gi in range(n_groups): + start = gi * group_size + end = start + group_size + score = float(saliency[start:end].sum()) + cands.append((score, start, end)) + return cands + + +def gptq_mixed_quantize(state_dict, hessians, act_stats, h): + result = {} + meta = {} + quant_gate = bool(getattr(h, "gated_attn_quant_gate", False)) + lqer_on = bool(getattr(h, "lqer_enabled", False)) + awq_on = bool(getattr(h, "awq_lite_enabled", False)) + lqer_cands = {} + awq_selected = collections.defaultdict(list) + if awq_on: + awq_cands = [] + for (name, tensor) in state_dict.items(): + t = tensor.detach().cpu().contiguous() + if t.is_floating_point() and t.numel() > 65536 and name in act_stats: + bits = h.embed_bits if "tok_emb" in name else h.matrix_bits + if bits < h.awq_lite_bits: + for score, start, end in _awq_lite_group_candidates( + t, act_stats[name], h.awq_lite_group_size + ): + awq_cands.append((score, name, start, end)) + awq_cands.sort(key=lambda x: -x[0]) + for (_score, name, start, end) in awq_cands[: h.awq_lite_group_top_k]: + awq_selected[name].append((start, end)) + for (name, tensor) in state_dict.items(): + t = tensor.detach().cpu().contiguous() + # Dedicated int8-per-row path for attn_gate_w (bypasses both GPTQ and + # fp16 passthrough). Applied BEFORE the numel<=65536 passthrough check + # so the gate tensor is routed here instead of to fp16. + if ( + quant_gate + and t.is_floating_point() + and t.ndim == 2 + and name.endswith(".attn_gate_w") + # Dense GatedAttn: (num_heads, dim) = (8, 512) = 4096. + # Sparse gate: (num_heads, gate_window) = (8, 12) = 96. + # Both need int8-per-row routing; the 1024 lower bound in stock + # PR-1736 presumed dense-only. Widen to catch both. + and 32 <= t.numel() <= 8192 + ): + gq, gs = _quantize_gate_int8_row(t) + result[name + ".gq"] = gq + result[name + ".gs"] = gs + meta[name] = "gate_int8_row" + continue + 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 + q, s, protect_meta = gptq_quantize_weight( + t, + hessians[name], + clip_sigmas=cs, + clip_range=clip_range, + protect_groups=awq_selected.get(name), + group_size=h.awq_lite_group_size if name in awq_selected else None, + protect_clip_range=(2 ** (h.awq_lite_bits - 1) - 1) + if name in awq_selected + else None, + ) + result[name + ".q"] = q + result[name + ".scale"] = s + meta[name] = f"gptq (int{bits})" + W_q = q.float() * s.float().view(-1, 1) + if protect_meta is not None: + result[name + ".awqg_start"] = protect_meta["starts"] + result[name + ".awqg_s_hi"] = protect_meta["s_hi"] + result[name + ".awqg_size"] = torch.tensor( + protect_meta["size"], dtype=torch.int16 + ) + meta[name] = meta[name] + f"+awqgrpint{h.awq_lite_bits}" + gsz = protect_meta["size"] + for start in protect_meta["starts"].tolist(): + W_q[:, start : start + gsz] = ( + q[:, start : start + gsz].float() + * protect_meta["s_hi"].float().view(-1, 1) + ) + if lqer_on: + # LQER is fit on top of the fully realized GPTQ base, which already + # includes any higher-precision AWQ-protected groups. + scope = str(getattr(h, "lqer_scope", "all")).lower() + scope_ok = ( + scope == "all" + or (scope == "mlp" and ".mlp." in name) + or (scope == "attn" and ".attn." in name) + or (scope == "embed" and "tok_emb" in name) + ) + if scope_ok: + E = t.float() - W_q + err_norm = float(E.norm()) + if err_norm > 0: + lqer_cands[name] = (E, err_norm) + if lqer_on and lqer_cands: + if bool(getattr(h, "lqer_gain_select", False)): + scored = [] + for (name, (E, base_err)) in lqer_cands.items(): + fit = _lqer_fit_quantized(E, h) + if fit is None: + continue + new_err = float((E - fit["delta"]).norm()) + gain = base_err - new_err + if gain > 0: + scored.append((gain, name, fit)) + scored.sort(key=lambda x: -x[0]) + for (_gain, name, fit) in scored[: h.lqer_top_k]: + if fit["kind"] == "asym": + result[name + ".lqA_a"] = fit["qA"] + result[name + ".lqAs_a"] = fit["sA"] + result[name + ".lqB_a"] = fit["qB"] + result[name + ".lqBs_a"] = fit["sB"] + meta[name] = meta[name] + "+lqer_asym" + else: + result[name + ".lqA"] = fit["qA"] + result[name + ".lqAs"] = fit["sA"] + result[name + ".lqB"] = fit["qB"] + result[name + ".lqBs"] = fit["sB"] + meta[name] = meta[name] + "+lqer" + else: + top = sorted(lqer_cands.items(), key=lambda kv: -kv[1][1])[: h.lqer_top_k] + asym_on = bool(getattr(h, "lqer_asym_enabled", False)) + asym_g = int(getattr(h, "lqer_asym_group", 64)) + for (name, (E, _)) in top: + U, S, Vh = torch.linalg.svd(E, full_matrices=False) + r = min(h.lqer_rank, 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] + "+lqer_asym" + 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] + "+lqer" + 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 + if info == "gate_int8_row": + gq = result[name + ".gq"] + gs = result[name + ".gs"] + out[name] = (gq.float() * gs.float().view(-1, 1)).to(orig_dtype) + 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 "awqgrpint" in info: + starts = result[name + ".awqg_start"].tolist() + s_hi = result[name + ".awqg_s_hi"].float() + gsz = int(result[name + ".awqg_size"].item()) + for start in starts: + W[:, start : start + gsz] = ( + q[:, start : start + gsz].float() * s_hi.view(-1, 1) + ) + 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" + + +# ── Per-group lrzip compression (ported from PR#1586 via PR#1667/1729) ──────── + +_GROUP_ORDER = [ + "_tok_emb.weight.q", + "attn.c_k.weight.q", "attn.c_q.weight.q", + "attn.c_v.weight.q", "attn.proj.weight.q", + "mlp.fc.weight.q", "mlp.proj.weight.q", +] +_SIMSORT_KEYS = {"_tok_emb.weight.q", "attn.c_q.weight.q", "mlp.fc.weight.q"} +_PACK_MAGIC = b"PGRP" + + +def _similarity_sort_l1(matrix): + import numpy as _np + n = matrix.shape[0] + used = _np.zeros(n, dtype=bool) + order = [0] + used[0] = True + cur = matrix[0].astype(_np.float32) + for _ in range(n - 1): + dists = _np.sum(_np.abs(matrix[~used].astype(_np.float32) - cur), axis=1) + unused = _np.where(~used)[0] + best = unused[_np.argmin(dists)] + order.append(best) + used[best] = True + cur = matrix[best].astype(_np.float32) + return _np.array(order, dtype=_np.uint16) + + +def _lrzip_compress(data, tmpdir, label): + inp = os.path.join(tmpdir, f"{label}.bin") + out = f"{inp}.lrz" + with open(inp, "wb") as f: + f.write(data) + subprocess.run(["lrzip", "-z", "-L", "9", "-o", out, inp], capture_output=True, check=True) + with open(out, "rb") as f: + result = f.read() + os.remove(inp); os.remove(out) + return result + + +def _lrzip_decompress(data, tmpdir, label): + inp = os.path.join(tmpdir, f"{label}.lrz") + out = os.path.join(tmpdir, f"{label}.bin") + with open(inp, "wb") as f: + f.write(data) + subprocess.run(["lrzip", "-d", "-f", "-o", out, inp], capture_output=True, check=True) + with open(out, "rb") as f: + result = f.read() + os.remove(inp); os.remove(out) + return result + + +def _pack_streams(streams): + import struct + n = len(streams) + hdr = _PACK_MAGIC + struct.pack("= 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, + y_bytes=None, +): + 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) + if y_bytes is not None: + tok_bytes = y_bytes.to(torch.float64) + else: + 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 _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, forward_ttt_score=None): + 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) + target_tokens = sum(doc_len - 1 for _, doc_len in docs) + 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}" + f" target_tokens:{target_tokens}" + f" peer_k:{max(1, int(h.ttt_peer_ensemble_k))}" + f" conf_thresh:{h.ttt_peer_conf_threshold} blend_w:{h.ttt_peer_conf_blend_w}" + ) + chunk_size, eval_seq_len = h.ttt_chunk_size, h.ttt_eval_seq_len + + def _parse_short_score_first_steps(raw): + steps = [] + for item in str(raw).split(","): + item = item.strip() + if not item: + continue + if ":" in item: + doc_raw, chunk_raw = item.split(":", 1) + elif "=" in item: + doc_raw, chunk_raw = item.split("=", 1) + else: + raise ValueError( + "TTT_SHORT_SCORE_FIRST_STEPS must look like '256:16,512:24'" + ) + doc_len = int(doc_raw.strip()) + step_chunk = int(chunk_raw.strip()) + if doc_len <= 0 or step_chunk <= 0: + raise ValueError("TTT short score-first steps must be positive") + steps.append((doc_len, step_chunk)) + steps.sort(key=lambda x: x[0]) + return steps + + short_score_steps = _parse_short_score_first_steps( + h.ttt_short_score_first_steps + ) + + def _score_first_chunk_for_doc(max_doc_len): + if not h.ttt_short_score_first_enabled: + return chunk_size + if short_score_steps: + for doc_limit, step_chunk in short_score_steps: + if max_doc_len <= doc_limit: + return step_chunk + return chunk_size + if max_doc_len <= h.ttt_short_doc_len and h.ttt_short_chunk_size > 0: + return h.ttt_short_chunk_size + return chunk_size + + 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) + ens_loss_sum = torch.zeros((), device=device, dtype=torch.float64) + ens_byte_sum = torch.zeros((), device=device, dtype=torch.float64) + ens_token_count = torch.zeros((), device=device, dtype=torch.float64) + ens_route_token_count = torch.zeros((), device=device, dtype=torch.float64) + ens_batches_covered = 0 + peer_k = max(1, int(h.ttt_peer_ensemble_k)) + t_start = time.perf_counter() + reusable_lora = BatchedTTTLoRA( + h.ttt_batch_size, base_model, h.ttt_lora_rank, + q_lora=h.ttt_q_lora, k_lora=h.ttt_k_lora, v_lora=h.ttt_v_lora, + mlp_lora=h.ttt_mlp_lora, o_lora=h.ttt_o_lora, + ).to(device) + reusable_short_lora = None + reusable_short_opt = None + + def _build_opt(lora, lr=None, weight_decay=None, beta2=None): + lr = h.ttt_lora_lr if lr is None else lr + lr = lr * h.ttt_local_lr_mult + weight_decay = h.ttt_weight_decay if weight_decay is None else weight_decay + beta2 = h.ttt_beta2 if beta2 is None else beta2 + if h.ttt_optimizer == "sgd": + return torch.optim.SGD( + lora.parameters(), lr=lr, + momentum=h.ttt_beta1, weight_decay=weight_decay, + ) + return torch.optim.AdamW( + lora.parameters(), lr=lr, + betas=(h.ttt_beta1, beta2), + eps=1e-10, weight_decay=weight_decay, fused=True, + ) + + def _reset_optimizer_state(opt): + for s in opt.state.values(): + for k, v in s.items(): + if isinstance(v, torch.Tensor): + v.zero_() + elif k == "step": + s[k] = 0 + + def _apply_lora_template(lora, template): + if not template: + return False + with torch.no_grad(): + for name, p in lora.named_parameters(): + t = template.get(name) + if t is None or tuple(t.shape) != tuple(p.shape[1:]): + return False + for name, p in lora.named_parameters(): + t = template[name].to(device=p.device, dtype=p.dtype) + p.copy_(t.unsqueeze(0).expand_as(p)) + return True + + def _update_lora_template(template, lora): + momentum = float(h.ttt_warm_start_mean_momentum) + new_template = {} + with torch.no_grad(): + for name, p in lora.named_parameters(): + mean = p.detach().mean(dim=0).clone() + old = template.get(name) if template else None + if old is not None and tuple(old.shape) == tuple(mean.shape): + mean = old.to(device=mean.device, dtype=mean.dtype).mul(momentum).add( + mean, alpha=1.0 - momentum + ) + new_template[name] = mean + return new_template + + reusable_opt = _build_opt(reusable_lora) + warm_lora_template = None + 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) + doc_lens = [dl for _, dl in batch] + max_doc_len = max(doc_lens) + train_doc_allowed = [ + (h.ttt_train_min_doc_len <= 0 or dl >= h.ttt_train_min_doc_len) + and (h.ttt_train_max_doc_len <= 0 or dl <= h.ttt_train_max_doc_len) + for dl in doc_lens + ] + train_doc_mask_t = torch.tensor( + train_doc_allowed, dtype=torch.float32, device=device + ) + use_short_lora = h.ttt_short_lora_enabled and max_doc_len <= h.ttt_short_doc_len + batch_chunk_size = _score_first_chunk_for_doc(max_doc_len) + use_short_chunks = batch_chunk_size != chunk_size + batch_lora_rank = h.ttt_short_lora_rank if use_short_lora else h.ttt_lora_rank + batch_lora_lr = h.ttt_short_lora_lr if use_short_lora else h.ttt_lora_lr + batch_lora_wd = h.ttt_short_weight_decay if use_short_lora else h.ttt_weight_decay + batch_lora_beta2 = h.ttt_short_beta2 if use_short_lora else h.ttt_beta2 + prev_loss = loss_sum.item() + prev_bytes = byte_sum.item() + prev_tokens = token_count.item() + if use_short_lora and bsz == h.ttt_batch_size: + if reusable_short_lora is None: + reusable_short_lora = BatchedTTTLoRA( + h.ttt_batch_size, base_model, h.ttt_short_lora_rank, + q_lora=h.ttt_q_lora, k_lora=h.ttt_k_lora, v_lora=h.ttt_v_lora, + mlp_lora=h.ttt_mlp_lora, o_lora=h.ttt_o_lora, + ).to(device) + reusable_short_opt = _build_opt( + reusable_short_lora, + lr=h.ttt_short_lora_lr, + weight_decay=h.ttt_short_weight_decay, + beta2=h.ttt_short_beta2, + ) + reusable_short_lora.reset() + _reset_optimizer_state(reusable_short_opt) + cur_lora = reusable_short_lora + cur_opt = reusable_short_opt + elif (not use_short_lora) and bsz == reusable_lora.bsz: + reusable_lora.reset() + _reset_optimizer_state(reusable_opt) + cur_lora = reusable_lora + cur_opt = reusable_opt + else: + cur_lora = BatchedTTTLoRA( + bsz, base_model, batch_lora_rank, + q_lora=h.ttt_q_lora, k_lora=h.ttt_k_lora, v_lora=h.ttt_v_lora, + mlp_lora=h.ttt_mlp_lora, o_lora=h.ttt_o_lora, + ).to(device) + cur_opt = _build_opt( + cur_lora, + lr=batch_lora_lr, + weight_decay=batch_lora_wd, + beta2=batch_lora_beta2, + ) + template_used = False + if ( + h.ttt_warm_start_mean_enabled + and max_doc_len <= h.ttt_warm_start_mean_doc_len + ): + template_used = _apply_lora_template(cur_lora, warm_lora_template) + pred_lens = [doc_len - 1 for _, doc_len in batch] + num_chunks = [(pl + batch_chunk_size - 1) // batch_chunk_size for pl in pred_lens] + max_nc = max(num_chunks) + max_pred_len = max(pred_lens) if pred_lens else 0 + num_chunks_t = torch.tensor(num_chunks, dtype=torch.int64, device=device) + peer_enabled = peer_k > 1 and bsz >= peer_k + if peer_enabled and forward_ttt_score is None: + raise ValueError("peer ensemble requires forward_ttt_score") + peer_nll_stash = own_nll_stash = own_entropy_stash = None + if peer_enabled: + peer_nll_stash = torch.zeros( + bsz, max_pred_len, peer_k - 1, device=device, dtype=torch.float32, + ) + own_nll_stash = torch.zeros(bsz, max_pred_len, device=device, dtype=torch.float32) + own_entropy_stash = torch.zeros(bsz, max_pred_len, device=device, dtype=torch.float32) + for ci in range(max_nc): + active = [ci < nc for nc in num_chunks] + needs_train = any( + train_doc_allowed[b] and ci < nc - 1 + for b, nc in enumerate(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], batch_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) * batch_chunk_size, ci + 1, batch_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) + per_tok_entropy = None + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + if peer_enabled: + per_tok_loss, per_tok_entropy = forward_ttt_score(x, y, lora=cur_lora) + else: + per_tok_loss = forward_ttt_train(x, y, lora=cur_lora) + # CaseOps sidecar-driven byte budget. Mirror the index pattern + # used to build y from all_tokens: y[b, j] corresponds to the + # token at global position tok_starts[b] + 1 + j (when valid). + y_bytes_arg = None + if val_data.caseops_enabled and val_data.val_bytes is not None: + y_idx = ( + tok_starts.unsqueeze(1) + + 1 + + col_idx[:context_size].unsqueeze(0) + ) + y_idx = y_idx.clamp_(max=val_data.val_bytes.numel() - 1) + y_bytes_arg = val_data.val_bytes[y_idx].to( + device=device, dtype=torch.int32, non_blocking=True + ) + # Mirror the `valid` masking used for y so out-of-range tokens + # contribute zero bytes (matches y=0 substitution above). + y_bytes_arg = torch.where( + valid, y_bytes_arg, torch.zeros_like(y_bytes_arg) + ) + 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, + y_bytes=y_bytes_arg, + ) + if own_nll_stash is not None and per_tok_entropy is not None: + with torch.no_grad(): + pred_start = ci * batch_chunk_size + pred_end = min(pred_start + batch_chunk_size, max_pred_len) + stash_range = pred_end - pred_start + col_stash = torch.arange(stash_range, device=device) + gather_idx = chunk_offsets.unsqueeze(1) + col_stash.unsqueeze(0) + gather_idx = gather_idx.clamp_(max=per_tok_loss.size(1) - 1) + scored_slice = torch.gather(per_tok_loss.float(), 1, gather_idx) + valid_stash = (col_stash.unsqueeze(0) < chunk_lens.unsqueeze(1)).float() + own_nll_stash[:, pred_start:pred_end] = scored_slice * valid_stash + entropy_slice = torch.gather(per_tok_entropy.float(), 1, gather_idx) + own_entropy_stash[:, pred_start:pred_end] = entropy_slice * valid_stash + if needs_train: + activate_chunk_mask = (num_chunks_t - 1 > ci).float() * train_doc_mask_t + 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 + batch_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 + if peer_enabled: + doc_starts_cpu = torch.tensor([ds for ds, _ in batch], dtype=torch.int64) + rng = torch.Generator(device="cpu").manual_seed( + (h.seed * 1013 + orig_batch_idx) & 0x7FFFFFFF + ) + peer_ids = torch.empty(bsz, peer_k - 1, dtype=torch.int64) + for b in range(bsz): + others = torch.cat([torch.arange(b), torch.arange(b + 1, bsz)]) + perm = torch.randperm(bsz - 1, generator=rng)[: peer_k - 1] + peer_ids[b] = others[perm] + for peer_slot in range(peer_k - 1): + BatchedLinearLoRA.PEER_IDX = peer_ids[:, peer_slot].to(device) + for c_start in range(0, max_pred_len, eval_seq_len): + c_ctx = min(eval_seq_len, max_pred_len - c_start) + col = torch.arange(c_ctx + 1) + idx_cpu = doc_starts_cpu.unsqueeze(1) + c_start + col.unsqueeze(0) + idx_cpu.clamp_(max=all_tokens.numel() - 1) + gathered = all_tokens_idx[idx_cpu].to(device=device, dtype=torch.int64, non_blocking=True) + valid_len = torch.tensor( + [max(0, min(c_ctx, pl - c_start)) for pl in pred_lens], dtype=torch.int64, + ) + valid_mask = (col[:c_ctx].unsqueeze(0) < valid_len.unsqueeze(1)).to(device, non_blocking=True) + px = torch.where(valid_mask, gathered[:, :c_ctx], 0) + py = torch.where(valid_mask, gathered[:, 1 : c_ctx + 1], 0) + with torch.no_grad(): + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + p_per_tok_loss = forward_ttt_train(px, py, lora=cur_lora) + peer_nll_stash[:, c_start : c_start + c_ctx, peer_slot].copy_( + p_per_tok_loss.float() * valid_mask.float() + ) + BatchedLinearLoRA.PEER_IDX = None + with torch.no_grad(): + own_w = float(h.ttt_peer_conf_blend_w) + peer_mean_prob = torch.exp(-peer_nll_stash).mean(dim=-1) + own_prob = torch.exp(-own_nll_stash) + blended = own_w * own_prob + (1.0 - own_w) * peer_mean_prob + blended_nll = -torch.log(blended.clamp_min(1e-40)) + uncertain = own_entropy_stash >= float(h.ttt_peer_conf_threshold) + ens_nll = torch.where(uncertain, blended_nll, own_nll_stash) + pred_positions = torch.arange(max_pred_len, device=device) + x_idx_cpu = doc_starts_cpu.unsqueeze(1) + pred_positions.unsqueeze(0).cpu() + x_idx_cpu.clamp_(max=all_tokens.numel() - 1) + x_gathered = all_tokens_idx[x_idx_cpu].to(device=device, dtype=torch.int64, non_blocking=True) + y_idx = x_idx_cpu + 1 + y_idx.clamp_(max=all_tokens.numel() - 1) + y_gathered = all_tokens_idx[y_idx].to(device=device, dtype=torch.int64, non_blocking=True) + pred_lens_t = torch.tensor(pred_lens, dtype=torch.int64, device=device) + valid_pred = pred_positions.unsqueeze(0) < pred_lens_t.unsqueeze(1) + mask_f64 = valid_pred.to(torch.float64) + if val_data.caseops_enabled and val_data.val_bytes is not None: + tok_bytes = val_data.val_bytes[y_idx].to(device=device, dtype=torch.int32, non_blocking=True) + tok_bytes = torch.where(valid_pred, tok_bytes, torch.zeros_like(tok_bytes)).to(torch.float64) + else: + tok_bytes = val_data.base_bytes_lut[y_gathered].to(torch.float64) + tok_bytes += ( + val_data.has_leading_space_lut[y_gathered] + & ~val_data.is_boundary_token_lut[x_gathered] + ).to(torch.float64) + ens_loss_sum += (ens_nll.to(torch.float64) * mask_f64).sum() + ens_byte_sum += (tok_bytes * mask_f64).sum() + ens_token_count += mask_f64.sum() + ens_route_token_count += (uncertain.to(torch.float64) * mask_f64).sum() + ens_batches_covered += 1 + if h.ttt_warm_start_mean_enabled: + warm_lora_template = _update_lora_template(warm_lora_template, cur_lora) + batch_num = orig_batch_idx + 1 + 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)} sr:{int(use_short_lora)} " + f"sf:{int(use_short_chunks)} tr:{sum(train_doc_allowed)}/{bsz} " + f"wt:{int(template_used)}" + ) + 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) + if train_doc_allowed[pos] + ) + 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, + q_lora=h.ttt_q_lora, k_lora=h.ttt_k_lora, v_lora=h.ttt_v_lora, + mlp_lora=h.ttt_mlp_lora, o_lora=h.ttt_o_lora, + ).to(device) + reusable_opt = _build_opt(reusable_lora) + reusable_short_lora = None + reusable_short_opt = None + 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) + dist.all_reduce(ens_loss_sum, op=dist.ReduceOp.SUM) + dist.all_reduce(ens_byte_sum, op=dist.ReduceOp.SUM) + dist.all_reduce(ens_token_count, op=dist.ReduceOp.SUM) + dist.all_reduce(ens_route_token_count, op=dist.ReduceOp.SUM) + baseline_loss, baseline_bpb = _loss_bpb_from_sums(loss_sum, token_count, byte_sum) + log( + f"peer_ens:coverage_tokens:{ens_token_count.item():.0f}/{token_count.item():.0f} " + f"batches_covered:{ens_batches_covered}" + ) + if ens_token_count.item() > 0: + cov_loss, cov_bpb = _loss_bpb_from_sums(ens_loss_sum, ens_token_count, ens_byte_sum) + route_frac = (ens_route_token_count / ens_token_count).item() + log( + f"peer_ens:route_frac:{route_frac:.6f} " + f"routed_tokens:{ens_route_token_count.item():.0f}/{ens_token_count.item():.0f}" + ) + log(f"peer_ens:ensemble_covered_only val_loss:{cov_loss:.6f} val_bpb:{cov_bpb:.6f}") + log(f"peer_ens:baseline val_loss:{baseline_loss:.6f} val_bpb:{baseline_bpb:.6f}") + for p in base_model.parameters(): + p.requires_grad_(True) + base_model.train() + if ens_token_count.item() > 0: + return cov_loss, cov_bpb + return baseline_loss, baseline_bpb + + +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) + compiled_model = torch.compile(base_model, dynamic=False, fullgraph=True) + compiled_forward_logits = torch.compile( + base_model.forward_logits, dynamic=False, fullgraph=True + ) + 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) + train_seq_plan = parse_train_seq_schedule(h.train_seq_schedule, h.train_seq_len) + midrun_cap_plan = parse_scalar_schedule(h.midrun_cap_schedule, 1.0) + max_train_seq_len = max_train_seq_len_from_schedule(train_seq_plan, h.train_seq_len) + if max_train_seq_len != h.train_seq_len: + raise ValueError( + f"TRAIN_SEQ_LEN={h.train_seq_len} must match the maximum sequence length in " + f"TRAIN_SEQ_SCHEDULE ({max_train_seq_len})" + ) + local_microbatch_tokens = validate_train_seq_plan_compatibility( + train_seq_plan, + global_tokens=h.train_batch_tokens, + world_size=h.world_size, + grad_accum_steps=h.grad_accum_steps, + ) + log( + "train_seq_schedule:" + + ",".join((f"{seq_len}@{threshold:.3f}" for threshold, seq_len in train_seq_plan)) + ) + if h.midrun_cap_schedule: + log( + "midrun_cap_schedule:" + + ",".join( + (f"{value:.3f}@{threshold:.3f}" for threshold, value in midrun_cap_plan) + ) + ) + log(f"local_microbatch_tokens:{local_microbatch_tokens}") + active_train_seq_len = train_seq_plan[0][1] + seq_change_warmup_start_step = None + midrun_cap_active = False + midrun_cap_prev_scale = schedule_value(midrun_cap_plan, 0.0) + log(f"growth_stage:seq_len:{active_train_seq_len} progress:0.000") + 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(step, elapsed_ms, frac): + if h.warmdown_iters > 0: + if max_wallclock_ms is None: + warmdown_start = max(h.iterations - h.warmdown_iters, 0) + if warmdown_start <= step < h.iterations: + return max( + (h.iterations - step) / max(h.warmdown_iters, 1), + h.min_lr, + ) + return 1.0 + step_ms = elapsed_ms / max(step, 1) + warmdown_ms = h.warmdown_iters * step_ms + remaining_ms = max(max_wallclock_ms - elapsed_ms, 0.0) + if remaining_ms <= warmdown_ms: + return max(remaining_ms / max(warmdown_ms, 1e-9), h.min_lr) + return 1.0 + 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 + + _clip_params = [p for p in base_model.parameters() if p.requires_grad] + def step_fn(step, lr_scale): + 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, active_train_seq_len + ) + with torch.autocast(device_type="cuda", dtype=torch.bfloat16, enabled=True): + loss = model( + x, y, cu_seqlens=cu_seqlens, max_seqlen=active_train_seq_len + ) + train_loss += loss.detach() + (loss / h.grad_accum_steps).backward() + train_loss /= h.grad_accum_steps + if step <= h.muon_momentum_warmup_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_(_clip_params, 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, active_train_seq_len + ) + log(f"warmup_cu_buckets:{','.join(str(b) for b in warmup_cu_buckets)} iters_each:{warmup_cu_iters}") + + def _compile_warmup_work_items(): + if not h.compile_shape_warmup: + items = [(h.train_seq_len, False)] + if h.num_loops > 0: + items.append((h.train_seq_len, True)) + return items + loop_mode = h.compile_shape_warmup_loop_modes + if loop_mode not in {"auto", "inactive", "active", "both"}: + raise ValueError( + "COMPILE_SHAPE_WARMUP_LOOP_MODES must be one of auto,inactive,active,both" + ) + items = [] + stage_start = 0.0 + for stage_end, seq_len in train_seq_plan: + if loop_mode == "inactive" or h.num_loops <= 0: + modes = [False] + elif loop_mode == "active": + modes = [True] + elif loop_mode == "both": + modes = [False, True] + else: + modes = [] + if stage_start < h.enable_looping_at: + modes.append(False) + if stage_end > h.enable_looping_at: + modes.append(True) + if not modes: + modes.append(False) + for loop_active in modes: + item = (seq_len, loop_active) + if item not in items: + items.append(item) + stage_start = stage_end + return items + + def _run_cu_bucket_warmup(seq_len): + for bucket_len in warmup_cu_buckets: + boundaries = list(range(0, x.size(1), max(seq_len, 1))) + if boundaries[-1] != x.size(1): + boundaries.append(x.size(1)) + if bucket_len < len(boundaries): + continue + 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=seq_len) + (wloss / h.grad_accum_steps).backward() + optimizers.zero_grad_all() + + warmup_items = _compile_warmup_work_items() + if h.compile_shape_warmup: + log( + "compile_shape_warmup:start " + + ",".join( + (f"{seq}x{'loop' if loop else 'plain'}" for seq, loop in warmup_items) + ) + ) + for seq_len, loop_active in warmup_items: + base_model.looping_active = bool(loop_active) + if h.compile_shape_warmup: + log( + f"compile_shape_warmup:shape seq_len:{seq_len} loop:{int(loop_active)}" + ) + for _ in range(max(h.compile_shape_warmup_iters if h.compile_shape_warmup else 1, 1)): + _run_cu_bucket_warmup(seq_len) + 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) + _live_state = base_model.state_dict(keep_vars=True) + ema_state = { + name: t.detach().float().clone() + for (name, t) in _live_state.items() + } + _ema_pairs = [(ema_state[name], t) for (name, t) in _live_state.items()] + ema_decay = h.ema_decay + training_time_ms = 0.0 + forced_stop_step = int(os.environ.get("FORCE_STOP_STEP", "0")) + stop_after_step = forced_stop_step if forced_stop_step > 0 else 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) + stage_seq_len, frac = current_train_seq_len( + train_seq_plan, + step=step, + iterations=h.iterations, + elapsed_ms=elapsed_ms, + max_wallclock_ms=max_wallclock_ms, + schedule_mode=h.train_seq_schedule_mode, + ) + if stage_seq_len != active_train_seq_len: + active_train_seq_len = stage_seq_len + log(f"growth_stage:seq_len:{active_train_seq_len} progress:{frac:.3f} step:{step}") + if h.seq_change_warmup_steps > 0 and step > 0: + seq_change_warmup_start_step = step + log( + f"growth_stage_rewarmup:start step:{step} steps:{h.seq_change_warmup_steps} " + f"seq_len:{active_train_seq_len}" + ) + scale = lr_mul(step, elapsed_ms, frac) + cap_scale = schedule_value(midrun_cap_plan, frac) + cap_active = cap_scale < 0.999999 + if cap_active and not midrun_cap_active: + log(f"midrun_cap:start step:{step} progress:{frac:.3f} scale:{cap_scale:.3f}") + elif ( + cap_active + and h.midrun_cap_log_updates + and abs(cap_scale - midrun_cap_prev_scale) > 1e-6 + ): + log(f"midrun_cap:update step:{step} progress:{frac:.3f} scale:{cap_scale:.3f}") + if cap_active: + scale *= cap_scale + midrun_cap_active = cap_active + midrun_cap_prev_scale = cap_scale + if seq_change_warmup_start_step is not None and h.seq_change_warmup_steps > 0: + rewarm_progress = min( + max( + (step - seq_change_warmup_start_step + 1) + / max(h.seq_change_warmup_steps, 1), + 0.0, + ), + 1.0, + ) + scale *= rewarm_progress + if rewarm_progress >= 1.0: + seq_change_warmup_start_step = None + 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 ema_t, t in _ema_pairs: + ema_t.mul_(ema_decay).add_(t.detach(), 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 = ( + forced_stop_step <= 0 + and max_wallclock_ms is not None + and approx_training_time_ms >= max_wallclock_ms + ) + if h.distributed and forced_stop_step <= 0 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" + ) + if h.ema_decay <= 0: + log("averaging:none keeping current weights") + return base_model, compiled_model, compiled_forward_logits + 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): + global BOS_ID + 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: skip training + GPTQ, jump straight to TTT eval on a + # pre-existing quantized artifact. Used to test TTT-only improvements + # (e.g., PR-1767's alpha/warm-start/WD) without retraining. + ttt_eval_only = os.environ.get("TTT_EVAL_ONLY", "0") == "1" + quantize_only = os.environ.get("QUANTIZE_ONLY", "0") == "1" + if ttt_eval_only: + log("TTT_EVAL_ONLY=1 — skipping training + GPTQ, loading saved artifact for TTT eval") + log(f"ttt_lora_alpha: {BatchedLinearLoRA._ALPHA}") + log(f"ttt_warm_start_a: {BatchedLinearLoRA._WARM_START_A}") + log(f"ttt_weight_decay: {h.ttt_weight_decay}") + elif quantize_only: + log("QUANTIZE_ONLY=1 — skipping training, loading saved full-precision checkpoint") + log(f"quantize_only checkpoint: {h.model_path}") + if BOS_ID is None: + BOS_ID = 1 + base_model = GPT(h).to(device).bfloat16() + state = torch.load(h.model_path, map_location="cpu") + template = base_model.state_dict() + for key in ("softcap_pos", "softcap_neg"): + if key not in state and key in template: + state[key] = template[key].detach().cpu().clone() + log(f"quantize_only:added neutral missing {key}") + base_model.load_state_dict(state, strict=True) + del state + serialize(h, base_model, Path(__file__).read_text(encoding="utf-8")) + if h.distributed: + dist.barrier() + 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, + ) + if os.environ.get("PREQUANT_ONLY", "0") == "1": + log("PREQUANT_ONLY=1 — skipping serialize/GPTQ/post-quant eval/TTT") + return + 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 + if not ttt_eval_only: + compiled_model = torch.compile(eval_model, dynamic=False, fullgraph=True) + compiled_forward_logits = torch.compile( + eval_model.forward_logits, dynamic=False, fullgraph=True + ) + timed_eval( + "diagnostic quantized", + eval_val, + h, + device, + val_data, + compiled_model, + compiled_forward_logits, + ) + del eval_model + if h.ttt_enabled: + if not ttt_eval_only: + del compiled_model + if ttt_eval_only: + del eval_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) + + _fwd_ttt_compiled_inner = None + + def _fwd_ttt(input_ids, target_ids, lora): + nonlocal _fwd_ttt_compiled_inner + if _fwd_ttt_compiled_inner is None: + _fwd_ttt_compiled_inner = torch.compile(_fwd_ttt_inner, dynamic=True) + return _fwd_ttt_compiled_inner(input_ids, target_ids, lora=lora) + + fwd_ttt_compiled = _fwd_ttt + + def _fwd_ttt_score_inner(input_ids, target_ids, lora): + return ttt_model.forward_ttt(input_ids, target_ids, lora=lora, return_entropy=True) + + _fwd_ttt_score_compiled_inner = None + + def _fwd_ttt_score(input_ids, target_ids, lora): + nonlocal _fwd_ttt_score_compiled_inner + if _fwd_ttt_score_compiled_inner is None: + _fwd_ttt_score_compiled_inner = torch.compile(_fwd_ttt_score_inner, dynamic=True) + return _fwd_ttt_score_compiled_inner(input_ids, target_ids, lora=lora) + + fwd_ttt_score_compiled = _fwd_ttt_score + log(f"ttt_lora:warming up compile (random tokens, no val data)") + 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, + q_lora=h.ttt_q_lora, k_lora=h.ttt_k_lora, v_lora=h.ttt_v_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 * h.ttt_local_lr_mult, + betas=(h.ttt_beta1, h.ttt_beta2), + eps=1e-10, + weight_decay=h.ttt_weight_decay, + fused=True, + ) + warmup_ctx_lens = [h.ttt_chunk_size, h.ttt_eval_seq_len] + if ( + h.ttt_short_score_first_enabled + and h.ttt_short_chunk_size not in warmup_ctx_lens + ): + warmup_ctx_lens.insert(0, h.ttt_short_chunk_size) + for item in str(h.ttt_short_score_first_steps).split(","): + item = item.strip() + if not item: + continue + sep = ":" if ":" in item else "=" + if sep not in item: + continue + _, chunk_raw = item.split(sep, 1) + step_chunk = int(chunk_raw.strip()) + if step_chunk > 0 and step_chunk not in warmup_ctx_lens: + warmup_ctx_lens.insert(0, step_chunk) + for ctx_len in 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) + if int(h.ttt_peer_ensemble_k) > 1: + for ctx_len in (h.ttt_chunk_size, h.ttt_eval_seq_len): + 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, _entropy = fwd_ttt_score_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) + BatchedLinearLoRA.PEER_IDX = torch.randint(0, bsz, (bsz,), dtype=torch.int64, device=device) + for ctx_len in (h.ttt_eval_seq_len,): + 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.no_grad(): + with torch.autocast(device_type="cuda", dtype=torch.bfloat16): + _ = fwd_ttt_compiled(xw, yw, lora=wl) + BatchedLinearLoRA.PEER_IDX = None + 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, + forward_ttt_score=fwd_ttt_score_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) + 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 + dynamo_cache_size_limit = int(os.environ.get("DYNAMO_CACHE_SIZE_LIMIT", "128")) + torch._dynamo.config.cache_size_limit = dynamo_cache_size_limit + if hasattr(torch._dynamo.config, "recompile_limit"): + torch._dynamo.config.recompile_limit = dynamo_cache_size_limit + if hasattr(torch._dynamo.config, "accumulated_cache_size_limit"): + torch._dynamo.config.accumulated_cache_size_limit = max( + int(os.environ.get("DYNAMO_ACCUMULATED_CACHE_SIZE_LIMIT", "1024")), + dynamo_cache_size_limit, + ) + 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()