jayasriii commited on
Commit
331a7d7
·
verified ·
1 Parent(s): 1d4e5ea

Upload folder using huggingface_hub

Browse files
390M_7e4_mbs8.csv ADDED
The diff for this file is too large to render. See raw diff
 
390M_7e4_mbs8/blumodel_muon_step00011000.pt ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:b17f6889751272ce554c19cb3e0dfb296ad07629de723b6a86938a2e0d6051e5
3
+ size 4205800344
390M_7e4_mbs8/blumodel_muon_step00012000.pt ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:1a5027dde4964dcb1d255592e48fa1200440eae94cee94076ab6077cf26996ca
3
+ size 4205800344
390M_7e4_mbs8/blumodel_muon_step00013000.pt ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:e889866bd69bfe1fc68bad5fab505460cc33082ab06ffc68303da4b5f38430f4
3
+ size 4205800344
390M_7e4_mbs8/blumodel_muon_step00014000.pt ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:5749d1dcb13260c74f9ea2c7cd4f291be4ae642efc3b20df3a0ef704842c86b7
3
+ size 4205800344
390M_7e4_mbs8/blumodel_muon_step00015000.pt ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:818831300460e23e9fe8f11b34861a1bc8bf33708641152c7bd4e0c040712d80
3
+ size 4205800344
390M_7e4_mbs8/blumodel_muon_step00016000.pt ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:3ef932ecc9a0e47ca248de1ef0f6ea62f4fd3472d83e49abaacd91d5aa48eea1
3
+ size 4205800344
390M_7e4_mbs8/blumodel_muon_step00017000.pt ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:8c0bc42dc631e17940961774bb5723b1497a1ffcefbae61c63501406aba985b7
3
+ size 4205800344
390M_7e4_mbs8/blumodel_muon_step00017535.pt ADDED
@@ -0,0 +1,3 @@
 
 
 
 
1
+ version https://git-lfs.github.com/spec/v1
2
+ oid sha256:f30de3281f040dd7ccf5d6459720483f68c73618dcd1343fd15d089bbb70c9fb
3
+ size 4205800344
train_390M.py ADDED
@@ -0,0 +1,1849 @@
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
 
1
+
2
+ #!/usr/bin/env python3
3
+ # =============================================================================
4
+ # train_muon.py — the MUON arm of train.py. Same model, same data order, same
5
+ # schedule shape, same seed: the ONLY thing that changes is how a gradient
6
+ # becomes a weight update.
7
+ #
8
+ # Fork discipline — everything outside section 5 is byte-identical to train.py,
9
+ # so `diff train.py train_muon.py` is exactly the experiment being run. If you
10
+ # edit a kernel or the model in one file, port it to the other or the comparison
11
+ # stops meaning anything.
12
+ #
13
+ # The optimizer itself is torch.optim.Muon — stock, not reimplemented. This file
14
+ # supplies only the two things it does not do for you: the ROUTING (which
15
+ # parameters it may have) and the fp32 MASTER WEIGHTS (it updates the master,
16
+ # not the bf16 param). Requires a torch new enough to ship it; main() checks and
17
+ # says so.
18
+ #
19
+ # What Muon does: AdamW rescales each weight ENTRY by its own gradient history.
20
+ # Muon rescales each weight MATRIX by its singular values — it takes the
21
+ # momentum buffer M and replaces it with the nearest semi-orthogonal matrix
22
+ # (all singular values driven to ~1) via a 5-step Newton-Schulz iteration, then
23
+ # steps along that. Directions the gradient has barely explored get the same
24
+ # step size as the dominant ones, which is where the reported speedup comes
25
+ # from. Only 2D hidden weights can be orthogonalized, so this is a HYBRID:
26
+ #
27
+ # Muon : every 2D Linear weight inside the blocks (75.5M params, 66%)
28
+ # AdamW : the token embedding / lm_head (38.6M params)
29
+ # + RMSNorm weights + QK-norm gammas (1D) (22.3k params)
30
+ #
31
+ # The embedding is deliberately NOT given to Muon: its rows are per-token, not a
32
+ # linear map between two feature spaces, and orthogonalizing across the vocab
33
+ # axis mixes unrelated tokens. That split is the standard Muon recipe, and it
34
+ # keeps the AdamW arm's exact update on 34% of the parameters — so a difference
35
+ # in the loss curve is attributable to the hidden matrices alone.
36
+ #
37
+ # Everything train_ddp.py / train_zero1.py import from blu_common / blu_data /
38
+ # blu_kernels / blu_model is inlined here, in that order, so this file runs on
39
+ # its own:
40
+ # section 1 fused QK-norm + RoPE (Triton) <- blu_kernels.py
41
+ # section 2 config + WSD schedule <- blu_common.py
42
+ # section 3 data loader <- blu_data.py
43
+ # section 4 model <- blu_model.py
44
+ # section 5 optimizer / clipping / checkpointing
45
+ # section 6 training
46
+ #
47
+ # Architecture:
48
+ # token embedding
49
+ # -> N x { RMSNorm -> [fused QK-norm + RoPE] -> causal GQA (flash) -> Wo -> +x
50
+ # RMSNorm -> SwiGLU FFN (gate_up -> silu(gate)*up -> down) -> +x }
51
+ # -> final RMSNorm
52
+ # -> lm_head (weight-tied with the token embedding)
53
+ # -> cross-entropy
54
+ # -> backward -> AdamW (decoupled weight decay) on a WSD (warmup / stable /
55
+ # cosine-decay) schedule
56
+ #
57
+ # Kept from the AdamW arm, unchanged, because they are confounds if they move:
58
+ # * fp32 master weights for every non-fp32 param — Muon updates the MASTER,
59
+ # reading the same fp32 landing buffer AdamW reads, so both optimizers see
60
+ # bit-identical gradients;
61
+ # * decoupled weight decay 0.1 on dim>=2 params, applied to the master;
62
+ # * ONE global-norm gradient clip over every parameter, BEFORE the optimizer.
63
+ # (Muon's orthogonalization makes the update scale-invariant, so clipping is
64
+ # nearly a no-op for its group — but it still gates the AdamW group, and
65
+ # removing it would change two things at once.)
66
+ #
67
+ # Memory, at the 114M config this repo actually runs (printed at startup):
68
+ #
69
+ # parameters (bf16 Linear + fp32 embed/norms) 0.28 GiB
70
+ # gradients (same dtypes as the params) 0.28 GiB
71
+ # fp32 master weights (the bf16 Linears) 0.28 GiB
72
+ # fp32 gradient landing buffers 0.28 GiB
73
+ # Muon momentum (75.5M params, fp32 x1) 0.28 GiB
74
+ # AdamW m + v (38.7M params, fp32 x2) 0.29 GiB
75
+ # -------------------------------------------------------
76
+ # static total 1.70 GiB (AdamW arm: 1.98)
77
+ #
78
+ # Muon carries ONE state tensor per param where Adam carries two, so the saving
79
+ # scales with how much of the model is hidden matrices: 0.28 GiB here (14% of
80
+ # static), because this config spends a third of its parameters on a 50304-row
81
+ # embedding that stays on AdamW. At the 1.73B config in the original header the
82
+ # same split saves ~5.5 GiB.
83
+ #
84
+ # Activations dominate either way (63.96 GiB measured at B=16 x T=4096), so the
85
+ # OOM escape hatches are unchanged: --B 8, --loss-chunks 4, --activation-checkpoint.
86
+ #
87
+ # Fidelity notes (why a run here is statistically but not bit-wise identical):
88
+ # * init draws come from torch's Philox, not the C++ single advancing
89
+ # std::mt19937, so the two runs start from different samples of the same
90
+ # distributions;
91
+ # * QK-norm reads the projection output in bf16 here (autocast) where the C++
92
+ # casts it up to fp32 first; the reduction and the rotation are fp32 in both.
93
+ #
94
+ # Determinism: two runs of the same launch line on the same GPU produce
95
+ # bit-identical losses. Seeding is only half of that — the default CUDA kernels
96
+ # for the embedding backward and the flash/mem-efficient attention backward
97
+ # accumulate partial sums with atomicAdd, so the order they land in (and with it
98
+ # the fp32 rounding) follows the scheduler, not the program. --determinism
99
+ # strict (the default) forces the ordered variants of those kernels; it costs a
100
+ # few percent of step time, which --determinism off buys back.
101
+ #
102
+ # Env overrides (the bluscript_single.cpp names, so an existing launch line keeps
103
+ # working; the equivalent CLI flag wins if both are given):
104
+ # BLU_DEVICE=<idx> CUDA device to train on (default 0)
105
+ # BLU_MAX_STEPS=<n> short benchmark run (rescales WSD) (default off)
106
+ # BLU_B=<n> micro-batch size (default 4)
107
+ # BLU_GLOBAL_BATCH=<tokens> tokens per optimizer step (default 524288)
108
+ # BLU_DIAG=1 per-micro-batch loss + depth probes
109
+ #
110
+ # Multi-GPU: DDP, one process per GPU, via torchrun. RANK / LOCAL_RANK /
111
+ # WORLD_SIZE (set by torchrun, not this file) switch the bootstrap in main();
112
+ # their absence is the single-process path above, unchanged. --device /
113
+ # BLU_DEVICE are ignored under torchrun — LOCAL_RANK picks each rank's GPU.
114
+ # --global-batch is still the TOTAL tokens per optimizer step across every
115
+ # rank (grad_accum = global_batch / (B*T*world_size)), so the same
116
+ # --global-batch reproduces the same schedule at any world size. Data is
117
+ # rank-sharded (DataLoaderLite reads a disjoint, rank-offset slice of the
118
+ # shards); gradients and the reported train/val loss are averaged across
119
+ # ranks; only rank 0 writes the CSV and checkpoints. Same launch line as
120
+ # single-GPU otherwise:
121
+ # torchrun --standalone --nproc_per_node=8 train_muon.py
122
+ #
123
+ # Learning rate — the one knob that decides whether the comparison is fair.
124
+ #
125
+ # Muon's update is spectrally normalized, so its natural scale has nothing to do
126
+ # with AdamW's. torch.optim.Muon ships both published conventions behind its
127
+ # `adjust_lr_fn`; --muon-adjust-lr passes the string straight through:
128
+ #
129
+ # match_rms_adamw Moonshot's: gamma <- 0.2 * gamma * sqrt(max(A,B)). The
130
+ # (DEFAULT HERE) Newton-Schulz output has singular values ~1, so its
131
+ # entries have RMS 1/sqrt(max(A,B)); this lifts them to
132
+ # RMS 0.2, which is what a converged AdamW update looks
133
+ # like. The consequence: --muon-lr is on the SAME SCALE
134
+ # as --max-lr, and the default (2e-4, equal to the AdamW
135
+ # arm's max_lr) gives both arms the same step size on the
136
+ # first run — only the update DIRECTION differs. That is
137
+ # the controlled experiment; run it before tuning.
138
+ # original Keller's modded-nanogpt scaling, sqrt(max(1, A/B)).
139
+ # Entries stay at RMS 1/sqrt(max(A,B)), so an equivalent
140
+ # LR is ~5.5x larger at d_model=768 (0.2*sqrt(768)).
141
+ # Use it to reproduce published Muon numbers.
142
+ #
143
+ # NOTE: torch's own default for adjust_lr_fn is `original`. This file defaults
144
+ # to `match_rms_adamw` instead, because an LR-matched first run is the point of
145
+ # having two arms. If you cite a number from here, cite the flag with it.
146
+ #
147
+ # Muon usually tolerates 2-4x more LR than AdamW once it is working, so the
148
+ # honest protocol is: baseline run, matched run, then a short --max-steps 200
149
+ # sweep over --muon-lr before claiming a win.
150
+ #
151
+ # Launch (the two arms; separate CSV + checkpoint dirs so they never collide):
152
+ # python3 train.py # AdamW baseline
153
+ # python3 train_muon.py # Muon, LR-matched
154
+ # python3 train_muon.py --muon-lr 6e-4 # after a sweep
155
+ # python3 train_muon.py --max-steps 200 --muon-lr 4e-4 # short sweep point
156
+ # =============================================================================
157
+
158
+ import argparse
159
+ import glob
160
+ import math
161
+ import os
162
+ import random
163
+ import re
164
+ import signal
165
+ import time
166
+ from contextlib import nullcontext
167
+ from dataclasses import dataclass
168
+
169
+ # cuBLAS reads CUBLAS_WORKSPACE_CONFIG once, when it creates its handle on the
170
+ # first GEMM, and torch refuses to run fp32 cuBLAS GEMMs under
171
+ # use_deterministic_algorithms() unless it is one of the two allowed values. It
172
+ # has to be in the environment before torch touches CUDA, hence: before the
173
+ # import, not inside main(). An explicit setting from the launch line wins.
174
+ os.environ.setdefault("CUBLAS_WORKSPACE_CONFIG", ":4096:8")
175
+
176
+ import numpy as np
177
+ import torch
178
+ import torch.distributed as dist
179
+ import torch.nn as nn
180
+ import torch.nn.functional as F
181
+ from torch.nn.attention import SDPBackend, sdpa_kernel
182
+ from torch.nn.parallel import DistributedDataParallel as DDP
183
+ from torch.utils.checkpoint import checkpoint
184
+
185
+ try:
186
+ import triton
187
+ import triton.language as tl
188
+ _HAS_TRITON = True
189
+ except ImportError: # eager fallback below
190
+ _HAS_TRITON = False
191
+
192
+
193
+ # =============================================================================
194
+ # 1. Fused QK-norm (RMSNorm over head_dim) + RoPE, in Triton.
195
+ #
196
+ # PyTorch counterpart of the C++ `autograd::fused_rms_rope`. One kernel does,
197
+ # per (batch, head, position) row:
198
+ #
199
+ # s = rsqrt(mean(x^2) + eps) # RMS over head_dim
200
+ # y = x * s * gamma # learnable per-head_dim scale
201
+ # o = rope_neox(y, cos[t], sin[t]) # NeoX / Llama half-split rotation
202
+ #
203
+ # Attention is deliberately NOT fused in here — it stays a separate op
204
+ # (F.scaled_dot_product_attention on the flash backend), mirroring the C++ split
205
+ # between `fused_rms_rope` and `gqa_fused_flash_attn_forward/backward`.
206
+ #
207
+ # Why fuse: the eager formulation (`rms_rope_eager`) materialises an fp32 copy of
208
+ # x, the squared tensor, the normalised tensor, the gamma-scaled tensor and a
209
+ # `cat` of the two rotated halves — roughly 6 extra [B,H,T,D] buffers held for
210
+ # backward. The fused version saves only x (which the projection's backward
211
+ # already holds) plus one fp32 scalar per row.
212
+ #
213
+ # Layout notes
214
+ # - x is [B,H,T,D] and may be a non-contiguous transpose of a [B,T,H,D] view;
215
+ # arbitrary (b,h,t) strides are honoured, D must be stride-1.
216
+ # - D must be a power of two (the C++ kernel requires 64 or 128).
217
+ # - cos/sin are [T_ctx, D//2] fp32, as produced by `build_rope_cache`.
218
+ # - the row index is r = (b*H + h)*T + t, so the output is contiguous
219
+ # [B,H,T,D] in that same order.
220
+ # =============================================================================
221
+ if _HAS_TRITON:
222
+
223
+ @triton.jit
224
+ def _fused_rms_rope_fwd(
225
+ X, GAMMA, COS, SIN, OUT, RSTD,
226
+ sx_b, sx_h, sx_t, # input strides, in elements
227
+ T, H, N, # seq len, n_heads, total rows
228
+ eps,
229
+ D: tl.constexpr, HALF: tl.constexpr, BLOCK_R: tl.constexpr,
230
+ ):
231
+ pid = tl.program_id(0)
232
+ rows = pid * BLOCK_R + tl.arange(0, BLOCK_R) # [R]
233
+ mask_r = rows < N
234
+
235
+ # r = (b*H + h)*T + t -> recover t (for RoPE) and the (b, h) base.
236
+ t = rows % T
237
+ bh = rows // T
238
+ b = bh // H
239
+ h = bh % H
240
+ base = b * sx_b + h * sx_h + t * sx_t # [R]
241
+
242
+ dh = tl.arange(0, HALF) # [HALF]
243
+ p1 = X + base[:, None] + dh[None, :]
244
+ x1 = tl.load(p1, mask=mask_r[:, None], other=0.0).to(tl.float32)
245
+ x2 = tl.load(p1 + HALF, mask=mask_r[:, None], other=0.0).to(tl.float32)
246
+
247
+ # RMS over the FULL head_dim (both halves), fp32 accumulation.
248
+ ss = tl.sum(x1 * x1, 1) + tl.sum(x2 * x2, 1) # [R]
249
+ rstd = 1.0 / tl.sqrt(ss / D + eps) # [R]
250
+
251
+ g1 = tl.load(GAMMA + dh).to(tl.float32)
252
+ g2 = tl.load(GAMMA + HALF + dh).to(tl.float32)
253
+ y1 = x1 * rstd[:, None] * g1[None, :]
254
+ y2 = x2 * rstd[:, None] * g2[None, :]
255
+
256
+ cp = t[:, None] * HALF + dh[None, :]
257
+ c = tl.load(COS + cp, mask=mask_r[:, None], other=0.0)
258
+ s = tl.load(SIN + cp, mask=mask_r[:, None], other=0.0)
259
+
260
+ # NeoX half-split rotation.
261
+ o1 = y1 * c - y2 * s
262
+ o2 = y2 * c + y1 * s
263
+
264
+ op = OUT + rows[:, None] * D + dh[None, :]
265
+ tl.store(op, o1.to(OUT.dtype.element_ty), mask=mask_r[:, None])
266
+ tl.store(op + HALF, o2.to(OUT.dtype.element_ty), mask=mask_r[:, None])
267
+ tl.store(RSTD + rows, rstd, mask=mask_r)
268
+
269
+ # -------------------------------------------------------------------------
270
+ # backward
271
+ #
272
+ # RoPE is an orthogonal rotation, so its vjp is the inverse rotation:
273
+ # dy1 = do1*cos + do2*sin ; dy2 = do2*cos - do1*sin
274
+ # then, with n = x * rstd (the normalised row),
275
+ # dgamma += dy * n
276
+ # dn = dy * gamma
277
+ # dx = rstd * (dn - n * mean(dn * n))
278
+ # -------------------------------------------------------------------------
279
+ @triton.jit
280
+ def _fused_rms_rope_bwd(
281
+ X, GAMMA, COS, SIN, DOUT, RSTD, DX, DGP,
282
+ sx_b, sx_h, sx_t, # input strides
283
+ sd_b, sd_h, sd_t, # dx strides
284
+ T, H, N,
285
+ D: tl.constexpr, HALF: tl.constexpr, BLOCK_R: tl.constexpr,
286
+ ):
287
+ pid = tl.program_id(0)
288
+ rows = pid * BLOCK_R + tl.arange(0, BLOCK_R)
289
+ mask_r = rows < N
290
+
291
+ t = rows % T
292
+ bh = rows // T
293
+ b = bh // H
294
+ h = bh % H
295
+ base = b * sx_b + h * sx_h + t * sx_t
296
+
297
+ dh = tl.arange(0, HALF)
298
+ p1 = X + base[:, None] + dh[None, :]
299
+ x1 = tl.load(p1, mask=mask_r[:, None], other=0.0).to(tl.float32)
300
+ x2 = tl.load(p1 + HALF, mask=mask_r[:, None], other=0.0).to(tl.float32)
301
+
302
+ rstd = tl.load(RSTD + rows, mask=mask_r, other=1.0)
303
+ n1 = x1 * rstd[:, None]
304
+ n2 = x2 * rstd[:, None]
305
+
306
+ dp = DOUT + rows[:, None] * D + dh[None, :]
307
+ do1 = tl.load(dp, mask=mask_r[:, None], other=0.0).to(tl.float32)
308
+ do2 = tl.load(dp + HALF, mask=mask_r[:, None], other=0.0).to(tl.float32)
309
+
310
+ cp = t[:, None] * HALF + dh[None, :]
311
+ c = tl.load(COS + cp, mask=mask_r[:, None], other=0.0)
312
+ s = tl.load(SIN + cp, mask=mask_r[:, None], other=0.0)
313
+
314
+ # inverse rotation
315
+ dy1 = do1 * c + do2 * s
316
+ dy2 = do2 * c - do1 * s
317
+
318
+ # per-program partial dgamma; summed by the caller (deterministic, no atomics)
319
+ tl.store(DGP + pid * D + dh, tl.sum(dy1 * n1, 0))
320
+ tl.store(DGP + pid * D + HALF + dh, tl.sum(dy2 * n2, 0))
321
+
322
+ g1 = tl.load(GAMMA + dh).to(tl.float32)
323
+ g2 = tl.load(GAMMA + HALF + dh).to(tl.float32)
324
+ dn1 = dy1 * g1[None, :]
325
+ dn2 = dy2 * g2[None, :]
326
+
327
+ dot = (tl.sum(dn1 * n1, 1) + tl.sum(dn2 * n2, 1)) / D # [R]
328
+ dx1 = rstd[:, None] * (dn1 - n1 * dot[:, None])
329
+ dx2 = rstd[:, None] * (dn2 - n2 * dot[:, None])
330
+
331
+ dbase = b * sd_b + h * sd_h + t * sd_t
332
+ q1 = DX + dbase[:, None] + dh[None, :]
333
+ tl.store(q1, dx1.to(DX.dtype.element_ty), mask=mask_r[:, None])
334
+ tl.store(q1 + HALF, dx2.to(DX.dtype.element_ty), mask=mask_r[:, None])
335
+
336
+
337
+ _BLOCK_R = 16
338
+ _NUM_WARPS = 4
339
+
340
+
341
+ def _launch_fwd(x, gamma, cos, sin, eps):
342
+ B, H, T, D = x.shape
343
+ N = B * H * T
344
+ out = torch.empty((B, H, T, D), dtype=x.dtype, device=x.device)
345
+ rstd = torch.empty(N, dtype=torch.float32, device=x.device)
346
+ grid = (triton.cdiv(N, _BLOCK_R),)
347
+ _fused_rms_rope_fwd[grid](
348
+ x, gamma, cos, sin, out, rstd,
349
+ x.stride(0), x.stride(1), x.stride(2),
350
+ T, H, N, eps,
351
+ D=D, HALF=D // 2, BLOCK_R=_BLOCK_R, num_warps=_NUM_WARPS,
352
+ )
353
+ return out, rstd
354
+
355
+
356
+ def _launch_bwd(x, gamma, cos, sin, dout, rstd):
357
+ B, H, T, D = x.shape
358
+ N = B * H * T
359
+ dout = dout.contiguous()
360
+ dx = torch.empty_like(x)
361
+ nprog = triton.cdiv(N, _BLOCK_R)
362
+ dgp = torch.empty((nprog, D), dtype=torch.float32, device=x.device)
363
+ _fused_rms_rope_bwd[(nprog,)](
364
+ x, gamma, cos, sin, dout, rstd, dx, dgp,
365
+ x.stride(0), x.stride(1), x.stride(2),
366
+ dx.stride(0), dx.stride(1), dx.stride(2),
367
+ T, H, N,
368
+ D=D, HALF=D // 2, BLOCK_R=_BLOCK_R, num_warps=_NUM_WARPS,
369
+ )
370
+ return dx, dgp.sum(0).to(gamma.dtype)
371
+
372
+
373
+ class _FusedRMSRoPE(torch.autograd.Function):
374
+ """QK-norm + RoPE for q and k in one autograd node (two kernel launches).
375
+ Mirrors the C++ `fused_rms_rope(q, k, q_gamma, k_gamma, cache)`: q/k keep
376
+ their incoming dtype (bf16 under autocast), the reduction and the rotation
377
+ run in fp32 internally, gamma and the cache stay fp32.
378
+ """
379
+
380
+ @staticmethod
381
+ def forward(ctx, q, k, q_gamma, k_gamma, cos, sin, eps):
382
+ qo, q_rstd = _launch_fwd(q, q_gamma, cos, sin, eps)
383
+ ko, k_rstd = _launch_fwd(k, k_gamma, cos, sin, eps)
384
+ ctx.save_for_backward(q, k, q_gamma, k_gamma, cos, sin, q_rstd, k_rstd)
385
+ return qo, ko
386
+
387
+ @staticmethod
388
+ def backward(ctx, dqo, dko):
389
+ q, k, q_gamma, k_gamma, cos, sin, q_rstd, k_rstd = ctx.saved_tensors
390
+ dq, dqg = _launch_bwd(q, q_gamma, cos, sin, dqo, q_rstd)
391
+ dk, dkg = _launch_bwd(k, k_gamma, cos, sin, dko, k_rstd)
392
+ return dq, dk, dqg, dkg, None, None, None
393
+
394
+
395
+ def fused_rms_rope(q, k, q_gamma, k_gamma, cos, sin, eps):
396
+ """q: [B,Hq,T,D], k: [B,Hkv,T,D] -> same shapes, contiguous.
397
+ `cos`/`sin` are [>=T, D//2] fp32; only the first T rows are read.
398
+ """
399
+ assert q.shape[-1] == k.shape[-1], "q and k must share head_dim"
400
+ D = q.shape[-1]
401
+ assert D % 2 == 0 and (D & (D - 1)) == 0, f"head_dim must be a power of two, got {D}"
402
+ assert q.stride(-1) == 1 and k.stride(-1) == 1, "head_dim must be stride-1"
403
+ T = q.shape[2]
404
+ return _FusedRMSRoPE.apply(q, k, q_gamma, k_gamma,
405
+ cos[:T].contiguous(), sin[:T].contiguous(), eps)
406
+
407
+
408
+ def rms_rope_eager(x, gamma, cos, sin, eps):
409
+ """Unfused QK-norm then NeoX RoPE, fp32 throughout — the reference the fused
410
+ kernel is measured against (`--rms-rope eager`), and the fallback when
411
+ Triton is unavailable."""
412
+ T, D = x.shape[2], x.shape[-1]
413
+ xf = x.float()
414
+ xf = xf * torch.rsqrt(xf.pow(2).mean(-1, keepdim=True) + eps) * gamma
415
+ x1, x2 = xf[..., : D // 2], xf[..., D // 2:]
416
+ c = cos[:T][None, None, :, :]
417
+ s = sin[:T][None, None, :, :]
418
+ return torch.cat([x1 * c - x2 * s, x2 * c + x1 * s], dim=-1).to(x.dtype)
419
+
420
+
421
+ def build_rope_cache(seq_len, head_dim, theta, device):
422
+ """NeoX / Llama RoPE cache: cos, sin of shape [seq_len, head_dim//2], fp32."""
423
+ inv_freq = 1.0 / (theta ** (torch.arange(0, head_dim, 2, dtype=torch.float32,
424
+ device=device) / head_dim))
425
+ pos = torch.arange(seq_len, dtype=torch.float32, device=device)
426
+ freqs = torch.outer(pos, inv_freq) # [T, D/2]
427
+ return torch.cos(freqs).contiguous(), torch.sin(freqs).contiguous()
428
+
429
+
430
+ # =============================================================================
431
+ # 2. Configuration — 1:1 with ModelConfig in bluscript_single.cpp — and the WSD
432
+ # learning-rate schedule.
433
+ # =============================================================================
434
+ @dataclass
435
+ class ModelConfig:
436
+ # ---- model ----
437
+ d_model: int = 960
438
+ n_layers: int = 30
439
+ q_heads: int = 15
440
+ kv_heads: int = 5 # grouped-query: G = q/kv = 4
441
+ head_dim: int = 64 # the fused kernels need 64 or 128
442
+ ffn_hidden: int = 2560 # ~8/3 * d_model, rounded to /256
443
+ context_length: int = 4096 # RoPE cache length; must be >= T
444
+ rope_theta: float = 500000.0
445
+ vocab_size: int = 50304 # 50257 (tiktoken) padded for the GPU
446
+ weight_tying: bool = False
447
+
448
+ rms_eps: float = 1e-5 # block RMSNorm epsilon
449
+ qk_eps: float = 1e-6 # QK-norm (fused) epsilon
450
+
451
+ # ---- optimization / data ----
452
+ B: int = 4 # micro-batch
453
+ T: int = 4096 # sequence length (<= context_length)
454
+ global_batch: int = 524288 # tokens per optimizer step
455
+ grad_clip: float = 1.0
456
+ weight_decay: float = 0.1
457
+
458
+ # WSD: linear warmup -> constant at max_lr -> cosine decay to min_lr.
459
+ # Requires warmup_steps + stable_steps < max_steps - 1 (the remainder is
460
+ # the decay tail).
461
+ max_lr: float = 0.00071
462
+ min_lr: float = 0.000071 # 10% of max_lr
463
+
464
+ muon_lr: float = 0.00071
465
+ muon_min_lr: float = 0.000071
466
+
467
+ max_steps: int = 26703
468
+
469
+ warmup_steps: int = 500
470
+ stable_steps: int = int(max_steps * 0.9) - warmup_steps # steps held at max_lr after warmup
471
+
472
+
473
+ # ---- Muon (hidden 2D weights); the AdamW fields above still drive the
474
+ # embedding / lm_head / 1D group, unchanged from the AdamW arm ----
475
+ # Default == max_lr: under adjust_lr_fn=match_rms_adamw the two are the same unit,
476
+ # so out of the box both arms take the same size step and differ only in
477
+ # direction. Same 10% floor, same WSD shape, so the schedules stay aligned.
478
+
479
+ # The rest are torch.optim.Muon's own arguments, passed through unchanged.
480
+ # momentum/nesterov/ns_steps are torch's defaults; adjust_lr_fn is not (see
481
+ # the header — torch defaults to "original").
482
+ muon_momentum: float = 0.95 # beta on the momentum buffer
483
+ muon_ns_steps: int = 5 # Newton-Schulz iterations
484
+ muon_adjust_lr: str = "match_rms_adamw" # "match_rms_adamw" | "original"
485
+
486
+ val_freq: int = 250
487
+ val_loss_steps: int = 20
488
+
489
+ # ---- checkpointing ----
490
+ checkpointing: bool = True
491
+ ckpt_freq: int = 1000 # save every 1800 optimizer steps
492
+ ckpt_keep: int = 40 # retain the most-recent N checkpoints
493
+ # Separate from the AdamW arm's dir: the two checkpoint formats differ (Muon
494
+ # momentum vs Adam m/v) and auto-resume picks the newest file it finds.
495
+ ckpt_dir: str = "ckpt_blu_muon"
496
+ ckpt_prefix: str = "blumodel_muon"
497
+
498
+ # ---- runtime knobs (not in the C++ struct) ----
499
+ fused_rms_rope: bool = True
500
+ activation_checkpoint: bool = False
501
+ loss_chunks: int = 1
502
+ # Depth probe: which layers --diag reports gradient / weight statistics for.
503
+ probe_layers: tuple = (0, 8, 17, 26, 35)
504
+
505
+
506
+ def get_lr(step, cfg, max_lr=None, min_lr=None):
507
+ """WSD: linear Warmup -> Stable (constant max_lr) -> cosine Decay to min_lr.
508
+ Byte-for-byte the C++ get_lr().
509
+ `max_lr`/`min_lr` override the config pair so the Muon group can ride the
510
+ SAME schedule shape — same warmup length, same stable window, same cosine
511
+ tail — at its own amplitude. Called with no override it is identical to the
512
+ AdamW arm's get_lr(), which is what the embedding/norm group uses.
513
+ """
514
+ max_lr = cfg.max_lr if max_lr is None else max_lr
515
+ min_lr = cfg.min_lr if min_lr is None else min_lr
516
+ if step < cfg.warmup_steps:
517
+ return max_lr * (step + 1) / cfg.warmup_steps
518
+ if step < cfg.stable_steps + cfg.warmup_steps:
519
+ return max_lr
520
+ tail = cfg.max_steps - (cfg.stable_steps + cfg.warmup_steps) - 1
521
+ decay_ratio = (step - (cfg.stable_steps + cfg.warmup_steps)) / max(tail, 1)
522
+ decay_ratio = min(max(decay_ratio, 0.0), 1.0)
523
+ coeff = 0.5 * (1.0 + math.cos(math.pi * decay_ratio)) # 1 -> 0
524
+ return min_lr + coeff * (max_lr - min_lr)
525
+
526
+
527
+ # =============================================================================
528
+ # 3. Data — port of the C++ DataLoaderLite (Tensor-Implementations/include/
529
+ # DataLoader.h), including its pinned-host / side-stream prefetch ring.
530
+ #
531
+ # Shards are HEADERLESS flat uint16 token files. A shard is "val" if "val"
532
+ # appears anywhere in its basename, "train" otherwise — this covers both an
533
+ # explicit flux_train_*/flux_val_* naming and an unmarked-train naming like
534
+ # flux_20b_shard_*.bin alongside a separately named *val* shard. For a batch
535
+ # at position p:
536
+ # x = tokens[p : p+B*T] y = tokens[p+1 : p+B*T+1]
537
+ # p strides by B*T (world_size == 1 here). A shard tail that cannot fit a full
538
+ # stride is dropped and the loader wraps to the next shard.
539
+ #
540
+ # The prefetch mirrors the C++ ring: NBUF pinned staging buffers and NBUF device
541
+ # tensors, H2D issued on a side stream, CUDA events ordering producer against
542
+ # consumer. Each slot also carries the (shard, token offset) it was filled from,
543
+ # which is what --diag prints to name the bytes behind a loss spike.
544
+ # =============================================================================
545
+ class DataLoaderLite:
546
+ NBUF = 4
547
+
548
+ def __init__(self, B, T, split, data_root, device, master=True,
549
+ max_tokens_per_shard=100_000_000, rank=0, world_size=1):
550
+ assert split in ("train", "val")
551
+ self.B, self.T = B, T
552
+ self.split = split
553
+ self.max_tokens = max_tokens_per_shard
554
+ self.device = device
555
+ # DDP: every rank reads the SAME shard list but a DISJOINT, interleaved
556
+ # slice of it (offset by rank, strided by world_size), so a training
557
+ # step's world_size*grad_accum micro-batches never repeat a window.
558
+ self.rank = rank
559
+ self.world_size = world_size
560
+
561
+ # Shard naming isn't consistent across datasets: some mark the split
562
+ # explicitly in both names (flux_train_*.bin / flux_val_*.bin), others
563
+ # only mark val and leave train shards unmarked (flux_20b_shard_*.bin
564
+ # + a separate *val* file). Treat "contains val" as the val split and
565
+ # everything else as train so both conventions resolve correctly.
566
+ self.shards = sorted(
567
+ p for p in glob.glob(os.path.join(data_root, "*.bin"))
568
+ if ("val" in os.path.basename(p).lower()) == (split == "val")
569
+ )
570
+ if not self.shards:
571
+ raise RuntimeError(f"no .bin shards found for split {split} in {data_root}")
572
+ if master:
573
+ print(f"found {len(self.shards)} shards for split {split}")
574
+
575
+ BT = B * T
576
+ self.pin_x = [torch.empty(BT, dtype=torch.int64).pin_memory() for _ in range(self.NBUF)]
577
+ self.pin_y = [torch.empty(BT, dtype=torch.int64).pin_memory() for _ in range(self.NBUF)]
578
+ self.dev_x = [torch.empty(B, T, dtype=torch.int64, device=device) for _ in range(self.NBUF)]
579
+ self.dev_y = [torch.empty(B, T, dtype=torch.int64, device=device) for _ in range(self.NBUF)]
580
+ self.copy_stream = torch.cuda.Stream(device=device)
581
+ self.h2d_done = [torch.cuda.Event() for _ in range(self.NBUF)]
582
+ self.h2d_recorded = [False] * self.NBUF
583
+ self.where = [(0, 0)] * self.NBUF # (shard_index, token_offset)
584
+ self.reset()
585
+
586
+ # -- shard plumbing ------------------------------------------------------
587
+ def _load_shard(self, i):
588
+ self.cur = i
589
+ arr = np.memmap(self.shards[i], dtype=np.uint16, mode="r")
590
+ self.tok = arr
591
+ self.ntok = min(len(arr), self.max_tokens)
592
+
593
+ def _advance_shard(self):
594
+ self._load_shard((self.cur + 1) % len(self.shards))
595
+ self.pos = self.B * self.T * self.rank
596
+
597
+ def reset(self):
598
+ self.copy_stream.synchronize()
599
+ self._load_shard(0)
600
+ self.pos = self.B * self.T * self.rank
601
+ self.primed = False
602
+ self.buf = 0
603
+
604
+ # -- prefetch ------------------------------------------------------------
605
+ def _stage(self, idx):
606
+ BT = self.B * self.T
607
+ stride = BT * self.world_size
608
+ if self.ntok < BT + 1:
609
+ raise RuntimeError(f"shard too small for one batch: {self.shards[self.cur]}")
610
+ if self.pos + BT + 1 > self.ntok:
611
+ self._advance_shard()
612
+
613
+ # Host-side WAR guard: the pinned buffer may still be the source of the
614
+ # H2D issued the last time this slot was filled (that copy is gated on
615
+ # the device, not the host, so a host running ahead of the GPU would
616
+ # otherwise overwrite bytes still in flight and hand the model a batch
617
+ # spliced from two different windows).
618
+ if self.h2d_recorded[idx]:
619
+ self.h2d_done[idx].synchronize()
620
+
621
+ self.where[idx] = (self.cur, self.pos)
622
+ buf = np.asarray(self.tok[self.pos: self.pos + BT + 1]).astype(np.int64)
623
+ self.pin_x[idx].copy_(torch.from_numpy(buf[:-1]))
624
+ self.pin_y[idx].copy_(torch.from_numpy(buf[1:]))
625
+
626
+ with torch.cuda.stream(self.copy_stream):
627
+ self.dev_x[idx].view(-1).copy_(self.pin_x[idx], non_blocking=True)
628
+ self.dev_y[idx].view(-1).copy_(self.pin_y[idx], non_blocking=True)
629
+ self.h2d_done[idx].record(self.copy_stream)
630
+ self.h2d_recorded[idx] = True
631
+
632
+ self.pos += stride
633
+ if self.pos + BT + 1 > self.ntok:
634
+ self._advance_shard()
635
+
636
+ def next_batch(self):
637
+ """-> (x, y, shard_index, token_offset), x/y int64 [B,T] on device."""
638
+ if not self.primed:
639
+ for i in range(self.NBUF):
640
+ self._stage(i)
641
+ self.primed = True
642
+ self.buf = 0
643
+ else:
644
+ # Refill the slot returned LAST call. Its consumer's kernels are
645
+ # already queued on the compute stream, so the producer must wait
646
+ # for them before overwriting the device tensor.
647
+ self.copy_stream.wait_stream(torch.cuda.current_stream())
648
+ self._stage((self.buf - 1) % self.NBUF)
649
+
650
+ torch.cuda.current_stream().wait_event(self.h2d_done[self.buf])
651
+ i = self.buf
652
+ self.buf = (self.buf + 1) % self.NBUF
653
+ return self.dev_x[i], self.dev_y[i], self.where[i][0], self.where[i][1]
654
+
655
+ def skip_batches(self, n):
656
+ """Fast-forward `n` batches without touching the GPU (resume path)."""
657
+ self.copy_stream.synchronize()
658
+ self.primed = False
659
+ BT = self.B * self.T
660
+ stride = BT * self.world_size
661
+ for _ in range(n):
662
+ if self.pos + BT + 1 > self.ntok:
663
+ self._advance_shard()
664
+ self.pos += stride
665
+ if self.pos + BT + 1 > self.ntok:
666
+ self._advance_shard()
667
+
668
+
669
+ # =============================================================================
670
+ # 4. Model.
671
+ #
672
+ # Precision (bf16 mixed precision; both --precision modes run under autocast):
673
+ # bf16-params : Linear weights stored bf16, embedding / RMSNorm / QK-norm
674
+ # scales stay fp32, fp32 master weights live in the optimizer
675
+ # — the exact scheme bluscript_single.cpp uses (default).
676
+ # bf16 : all params fp32, autocast casts per-op (standard PyTorch AMP).
677
+ # In both cases: matmuls + attention bf16, RMSNorm / QK-norm reductions, RoPE,
678
+ # softmax and the loss in fp32. No GradScaler (bf16 has fp32's exponent range).
679
+ # =============================================================================
680
+ _FLASH = [SDPBackend.FLASH_ATTENTION, SDPBackend.EFFICIENT_ATTENTION]
681
+
682
+
683
+ class RMSNorm(nn.Module):
684
+ """RMSNorm over d_model, fp32 weight. F.rms_norm is a fused ATen op and is
685
+ not autocast-eligible, so it runs in the input dtype — fp32 on the residual
686
+ stream, matching the C++ nn::RMSNorm."""
687
+
688
+ def __init__(self, dim, eps):
689
+ super().__init__()
690
+ self.weight = nn.Parameter(torch.ones(dim))
691
+ self.eps = eps
692
+ self.dim = (dim,)
693
+
694
+ def forward(self, x):
695
+ return F.rms_norm(x, self.dim, self.weight, self.eps)
696
+
697
+
698
+ class CausalGQA(nn.Module):
699
+ """RMSNorm -> fused(QK-norm + RoPE) -> causal grouped-query attention -> Wo -> +x"""
700
+
701
+ def __init__(self, cfg, resid_scale):
702
+ super().__init__()
703
+ self.cfg = cfg
704
+ self.norm = RMSNorm(cfg.d_model, cfg.rms_eps)
705
+ self.wQ = nn.Linear(cfg.d_model, cfg.q_heads * cfg.head_dim, bias=False)
706
+ self.wK = nn.Linear(cfg.d_model, cfg.kv_heads * cfg.head_dim, bias=False)
707
+ self.wV = nn.Linear(cfg.d_model, cfg.kv_heads * cfg.head_dim, bias=False)
708
+ self.wO = nn.Linear(cfg.q_heads * cfg.head_dim, cfg.d_model, bias=False)
709
+ # QK-norm scales, initialised to 1.0 (identity RMS scale), fp32.
710
+ self.q_gamma = nn.Parameter(torch.ones(cfg.head_dim))
711
+ self.k_gamma = nn.Parameter(torch.ones(cfg.head_dim))
712
+ # GPT-2/Llama init: projections std=0.02, residual proj scaled by
713
+ # (2*n_layers)^-0.5 to keep residual-stream variance ~constant.
714
+ for lin in (self.wQ, self.wK, self.wV):
715
+ nn.init.normal_(lin.weight, mean=0.0, std=0.02)
716
+ nn.init.normal_(self.wO.weight, mean=0.0, std=0.02 * resid_scale)
717
+
718
+ def forward(self, x, cos, sin):
719
+ cfg = self.cfg
720
+ B, T, _ = x.shape
721
+ h = self.norm(x)
722
+
723
+ # [B,T,H*D] -> [B,H,T,D]. The transpose stays a view; the fused kernel
724
+ # honours arbitrary (b,h,t) strides, so no .contiguous() is paid here.
725
+ q = self.wQ(h).view(B, T, cfg.q_heads, cfg.head_dim).transpose(1, 2)
726
+ k = self.wK(h).view(B, T, cfg.kv_heads, cfg.head_dim).transpose(1, 2)
727
+ v = self.wV(h).view(B, T, cfg.kv_heads, cfg.head_dim).transpose(1, 2)
728
+
729
+ if cfg.fused_rms_rope:
730
+ q, k = fused_rms_rope(q, k, self.q_gamma, self.k_gamma, cos, sin, cfg.qk_eps)
731
+ else:
732
+ q = rms_rope_eager(q, self.q_gamma, cos, sin, cfg.qk_eps)
733
+ k = rms_rope_eager(k, self.k_gamma, cos, sin, cfg.qk_eps)
734
+
735
+ # Separate fused flash kernel. enable_gqa maps query head i to kv head
736
+ # i // (q_heads/kv_heads), the same grouping as the C++ kernel.
737
+ with sdpa_kernel(_FLASH):
738
+ o = F.scaled_dot_product_attention(
739
+ q.to(v.dtype), k.to(v.dtype), v, is_causal=True, enable_gqa=True)
740
+
741
+ o = o.transpose(1, 2).reshape(B, T, cfg.q_heads * cfg.head_dim)
742
+ return x + self.wO(o)
743
+
744
+
745
+ class SwiGLUMLP(nn.Module):
746
+ """RMSNorm -> gate_up -> silu(gate)*up -> down -> +x ([gate | up] layout)"""
747
+
748
+ def __init__(self, cfg, resid_scale):
749
+ super().__init__()
750
+ self.norm = RMSNorm(cfg.d_model, cfg.rms_eps)
751
+ self.gate_up = nn.Linear(cfg.d_model, 2 * cfg.ffn_hidden, bias=False)
752
+ self.down = nn.Linear(cfg.ffn_hidden, cfg.d_model, bias=False)
753
+ nn.init.normal_(self.gate_up.weight, mean=0.0, std=0.02)
754
+ nn.init.normal_(self.down.weight, mean=0.0, std=0.02 * resid_scale)
755
+
756
+ def forward(self, x):
757
+ gate, up = self.gate_up(self.norm(x)).chunk(2, dim=-1)
758
+ return x + self.down(F.silu(gate) * up)
759
+
760
+
761
+ class Block(nn.Module):
762
+ def __init__(self, cfg, resid_scale):
763
+ super().__init__()
764
+ self.attn = CausalGQA(cfg, resid_scale)
765
+ self.mlp = SwiGLUMLP(cfg, resid_scale)
766
+
767
+ def forward(self, x, cos, sin):
768
+ return self.mlp(self.attn(x, cos, sin))
769
+
770
+
771
+ class BLU(nn.Module):
772
+ def __init__(self, cfg):
773
+ super().__init__()
774
+ self.cfg = cfg
775
+ resid_scale = 1.0 / math.sqrt(2.0 * cfg.n_layers)
776
+
777
+ self.wte = nn.Embedding(cfg.vocab_size, cfg.d_model)
778
+ nn.init.normal_(self.wte.weight, mean=0.0, std=0.02)
779
+
780
+ self.blocks = nn.ModuleList([Block(cfg, resid_scale) for _ in range(cfg.n_layers)])
781
+ self.norm_f = RMSNorm(cfg.d_model, cfg.rms_eps)
782
+ if not cfg.weight_tying:
783
+ self.lm_head = nn.Linear(cfg.d_model, cfg.vocab_size, bias=False)
784
+ nn.init.normal_(self.lm_head.weight, mean=0.0, std=0.02)
785
+
786
+ cos, sin = build_rope_cache(cfg.context_length, cfg.head_dim, cfg.rope_theta, "cpu")
787
+ self.register_buffer("rope_cos", cos, persistent=False)
788
+ self.register_buffer("rope_sin", sin, persistent=False)
789
+
790
+ def apply_param_dtype(self, mode):
791
+ """`bf16-params`: store Linear weights in bf16 (embedding, RMSNorm and
792
+ QK-norm scales stay fp32) — the two-dtype-group layout of the C++ model.
793
+ Under autocast this is transparent: the cast to bf16 becomes a no-op."""
794
+ if mode != "bf16-params":
795
+ return
796
+ for m in self.modules():
797
+ if isinstance(m, nn.Linear):
798
+ m.weight.data = m.weight.data.to(torch.bfloat16)
799
+
800
+ def forward(self, idx, targets=None):
801
+ cfg = self.cfg
802
+ T = idx.shape[1]
803
+ cos, sin = self.rope_cos[:T], self.rope_sin[:T]
804
+
805
+ x = self.wte(idx) # fp32 [B,T,d_model]
806
+ for blk in self.blocks:
807
+ if cfg.activation_checkpoint and self.training:
808
+ x = checkpoint(blk, x, cos, sin, use_reentrant=False)
809
+ else:
810
+ x = blk(x, cos, sin)
811
+ x = self.norm_f(x)
812
+
813
+ head_w = self.wte.weight if cfg.weight_tying else self.lm_head.weight
814
+ if targets is None:
815
+ return F.linear(x, head_w), None
816
+ if cfg.loss_chunks > 1:
817
+ return None, _chunked_ce(x, head_w, targets, cfg.loss_chunks)
818
+ # Matches the C++: bf16 matmul, logits promoted to fp32 for the loss
819
+ # (cross_entropy is an autocast-fp32 op, so the promotion is implicit).
820
+ logits = F.linear(x, head_w)
821
+ loss = F.cross_entropy(logits.view(-1, logits.size(-1)), targets.reshape(-1))
822
+ return logits, loss
823
+
824
+
825
+ def _chunked_ce(x, weight, targets, chunks):
826
+ """lm_head + cross-entropy in row chunks, logits recomputed in backward.
827
+ Peak logit memory drops from [B*T, V] to [B*T/chunks, V] at the cost of one
828
+ extra lm_head matmul. Off by default (`--loss-chunks 1`) so the memory
829
+ profile matches the C++ script, which materialises the full fp32 logits.
830
+ """
831
+ xf = x.reshape(-1, x.shape[-1])
832
+ tf = targets.reshape(-1)
833
+ n = xf.shape[0]
834
+ size = (n + chunks - 1) // chunks
835
+ total = xf.new_zeros((), dtype=torch.float32)
836
+ for lo in range(0, n, size):
837
+ hi = min(lo + size, n)
838
+ part = checkpoint(_ce_chunk, xf[lo:hi], weight, tf[lo:hi], use_reentrant=False)
839
+ total = total + part * (hi - lo)
840
+ return total / n
841
+
842
+
843
+ def _ce_chunk(x, weight, targets):
844
+ return F.cross_entropy(F.linear(x, weight).float(), targets)
845
+
846
+
847
+ # =============================================================================
848
+ # 5. Optimizer (Muon on the hidden matrices, AdamW on the rest), gradient
849
+ # clipping, checkpointing.
850
+ # =============================================================================
851
+ def _is_muon_param(name, p):
852
+ """Muon orthogonalizes a matrix, so it needs the parameter to BE a linear
853
+ map between two feature spaces:
854
+ * dim >= 2 — RMSNorm weights and QK-norm gammas are 1D and have no
855
+ singular values to equalise;
856
+ * not the embedding / lm_head — its rows index tokens, not features, so
857
+ the left singular vectors mix unrelated vocabulary entries. Weight
858
+ tying makes this one tensor serving both ends; either way it is AdamW's.
859
+ Everything else — wQ/wK/wV/wO, gate_up, down — is Muon's: 75.5M of the
860
+ 114.2M parameters at the default config.
861
+ """
862
+ return p.dim() >= 2 and not name.startswith(("wte.", "lm_head."))
863
+
864
+
865
+ def _state_bytes(opt):
866
+ """Bytes held in a torch optimizer's per-parameter state — Adam's m/v, or
867
+ Muon's momentum buffer. Reads 0 before the first step(): both allocate
868
+ lazily."""
869
+ return sum(v.numel() * v.element_size()
870
+ for st in opt.state.values()
871
+ for v in st.values() if torch.is_tensor(v))
872
+
873
+
874
+ class MasterMuon:
875
+ """torch.optim.Muon on the hidden matrices + torch.optim.AdamW on the rest,
876
+ both driving fp32 master weights.
877
+ Neither optimizer is reimplemented here. What this class owns is the part
878
+ torch cannot know about:
879
+ * ROUTING — torch.optim.Muon accepts whatever you hand it and the docs are
880
+ explicit that only 2D hidden weights belong ("bias, and embedding,
881
+ should be optimized by a standard method such as AdamW"). The split is
882
+ `_is_muon_param`.
883
+ * MASTER WEIGHTS — the parameters are bf16 under --precision bf16-params,
884
+ and a bf16 tensor cannot absorb an lr*update of 1e-4 relative size. So
885
+ both optimizers are constructed over the fp32 MASTERS, fed from the fp32
886
+ landing buffers, and the masters are cast back over the params after the
887
+ step. Unchanged from the AdamW arm — that machinery is what makes the
888
+ two arms comparable, since both optimizers then read a bit-identical
889
+ fp32 gradient and write a bit-identical bf16 parameter.
890
+ Weight decay is decoupled in both (torch.optim.Muon applies it before the
891
+ update, as AdamW does) and is applied to the master for the same reason: a
892
+ decay written to the bf16 param is overwritten by the master write-back
893
+ microseconds later and silently no-ops.
894
+ State cost: Muon holds one momentum buffer per param, against Adam's two.
895
+ """
896
+
897
+ def __init__(self, named_params, cfg):
898
+ self.cfg = cfg
899
+ self.triples = [] # (param, master, landing_grad), ALL params
900
+ self.muon_masters = [] # masters routed to Muon
901
+ decay, nodecay = [], []
902
+ self.n_muon = self.n_decay = self.n_nodecay = 0
903
+ for name, p in named_params:
904
+ if p.dtype == torch.float32:
905
+ m, g = p, None
906
+ else:
907
+ m = p.detach().float().clone()
908
+ g = torch.zeros_like(m)
909
+ m.grad = g
910
+ self.triples.append((p, m, g))
911
+ if _is_muon_param(name, p):
912
+ self.muon_masters.append(m)
913
+ self.n_muon += m.numel()
914
+ elif p.dim() >= 2:
915
+ decay.append(m)
916
+ self.n_decay += m.numel()
917
+ else:
918
+ nodecay.append(m)
919
+ self.n_nodecay += m.numel()
920
+ self.opt = torch.optim.AdamW(
921
+ [{"params": decay, "weight_decay": cfg.weight_decay},
922
+ {"params": nodecay, "weight_decay": 0.0}],
923
+ lr=cfg.max_lr, betas=(0.9, 0.95), eps=1e-8, fused=True)
924
+ # Everything below is a torch.optim.Muon default except adjust_lr_fn
925
+ # (torch: "original") and weight_decay, which is pinned to the AdamW
926
+ # arm's 0.1 so the decay is identical across both groups and both arms.
927
+ self.muon = torch.optim.Muon(
928
+ self.muon_masters,
929
+ lr=cfg.muon_lr,
930
+ weight_decay=cfg.weight_decay,
931
+ momentum=cfg.muon_momentum,
932
+ nesterov=True,
933
+ ns_steps=cfg.muon_ns_steps,
934
+ adjust_lr_fn=cfg.muon_adjust_lr)
935
+
936
+ @torch.no_grad()
937
+ def step(self, lr, muon_lr):
938
+ for p, m, g in self.triples:
939
+ if g is not None:
940
+ g.copy_(p.grad) # bf16 grad -> fp32 landing buffer
941
+ # AdamW group (embedding / lm_head / norms / gammas) — identical to the
942
+ # baseline's update, on identical gradients.
943
+ for grp in self.opt.param_groups:
944
+ grp["lr"] = lr
945
+ self.opt.step()
946
+ # Muon group. Its nesterov path may consume the gradient in place; that
947
+ # is harmless here because grad_norm and the --diag probes are read
948
+ # before step() and the landing buffers are overwritten wholesale on the
949
+ # next one.
950
+ for grp in self.muon.param_groups:
951
+ grp["lr"] = muon_lr
952
+ self.muon.step()
953
+ for p, m, g in self.triples:
954
+ if g is not None:
955
+ p.data.copy_(m) # fp32 master -> bf16 param
956
+
957
+ def state_breakdown(self):
958
+ """-> (master_bytes, landing_grad_bytes, adam_mv_bytes, muon_mom_bytes).
959
+ The first two are zero under --precision bf16, where the master IS the
960
+ parameter and the gradient is already fp32. Both optimizers allocate
961
+ their state lazily on the first step(), so the last two read 0 until
962
+ then — the startup report uses the expected_* predictors instead.
963
+ """
964
+ master = landing = 0
965
+ for p, m, g in self.triples:
966
+ if m is not p:
967
+ master += m.numel() * m.element_size()
968
+ landing += g.numel() * g.element_size()
969
+ mv = _state_bytes(self.opt)
970
+ mom = _state_bytes(self.muon)
971
+ return master, landing, mv, mom
972
+
973
+ def state_bytes(self):
974
+ return sum(self.state_breakdown())
975
+
976
+ def expected_mv_bytes(self):
977
+ """Adam m/v are lazily allocated, so the startup report has to predict
978
+ them rather than measure them. Muon's params are excluded — they carry a
979
+ momentum buffer instead, reported separately."""
980
+ muon_ids = {id(m) for m in self.muon_masters}
981
+ return sum(m.numel() * m.element_size() * 2
982
+ for _, m, _ in self.triples if id(m) not in muon_ids)
983
+
984
+ def expected_muon_mom_bytes(self):
985
+ """One fp32 momentum buffer per Muon parameter — the prediction the
986
+ startup report prints before step 0 has allocated it."""
987
+ return sum(m.numel() * m.element_size() for m in self.muon_masters)
988
+
989
+ # -- checkpointing -------------------------------------------------------
990
+ def state_dict(self):
991
+ # `optimizer` tags the arm: resuming a Muon run from an AdamW checkpoint
992
+ # would restore the weights and then silently drop the momentum, so
993
+ # CheckpointManager refuses on a mismatch rather than producing a run
994
+ # that looks fine and is not comparable to either arm.
995
+ return {"optimizer": "muon",
996
+ "adam": self.opt.state_dict(),
997
+ "muon": self.muon.state_dict(),
998
+ "masters": [None if m is p else m for p, m, g in self.triples]}
999
+
1000
+ @torch.no_grad()
1001
+ def load_state_dict(self, blob):
1002
+ """-> True if the fp32 masters were restored from the checkpoint."""
1003
+ if blob.get("optimizer") != "muon":
1004
+ raise SystemExit(
1005
+ "ERROR: this checkpoint was written by the AdamW arm (train.py). "
1006
+ "Point --ckpt-dir somewhere else, or resume it with train.py.")
1007
+ self.opt.load_state_dict(blob["adam"])
1008
+ if blob.get("muon") is not None:
1009
+ self.muon.load_state_dict(blob["muon"])
1010
+ masters = blob.get("masters")
1011
+ if masters is None or len(masters) != len(self.triples):
1012
+ return False
1013
+ for (p, m, g), saved in zip(self.triples, masters):
1014
+ if saved is not None and m is not p:
1015
+ m.copy_(saved.to(m.device))
1016
+ return True
1017
+
1018
+ @torch.no_grad()
1019
+ def resync_masters(self):
1020
+ """Re-seed every master from its parameter.
1021
+ Needed after loading a checkpoint that carries no masters: the masters
1022
+ were built in __init__ from the RANDOM INIT, so without this the first
1023
+ step() would write them straight back over the restored weights and
1024
+ silently throw the run away.
1025
+ """
1026
+ for p, m, g in self.triples:
1027
+ if m is not p:
1028
+ m.copy_(p.detach().float())
1029
+
1030
+
1031
+ def clip_grads_(params, max_norm):
1032
+ """ONE global L2 norm over every parameter, then one rescale — the textbook
1033
+ definition. (It differs from the ZeRO script, which clips each dtype group
1034
+ against its own norm and can therefore let the true global norm reach
1035
+ sqrt(n_groups) * grad_clip.) The per-tensor norms are promoted to fp32
1036
+ before they are combined, or a reduction over ~250 bf16 norms would lose
1037
+ ~3 digits."""
1038
+ grads = [p.grad for p in params if p.grad is not None]
1039
+ norms = torch._foreach_norm(grads)
1040
+ total = torch.linalg.vector_norm(torch.stack([n.float() for n in norms]))
1041
+ if max_norm and max_norm > 0:
1042
+ clip = (max_norm / (total + 1e-6)).clamp(max=1.0)
1043
+ torch._foreach_mul_(grads, clip)
1044
+ return total
1045
+
1046
+
1047
+ class CheckpointManager:
1048
+ """Synchronous save / keep-N / auto-resume, mirroring the C++
1049
+ CheckpointManager. A checkpoint holds the model parameters, Adam's m/v, the
1050
+ fp32 masters, and the step it should resume at."""
1051
+
1052
+ def __init__(self, ckpt_dir, prefix, keep):
1053
+ self.dir, self.prefix, self.keep = ckpt_dir, prefix, keep
1054
+ os.makedirs(ckpt_dir, exist_ok=True)
1055
+
1056
+ def _path(self, step):
1057
+ return os.path.join(self.dir, f"{self.prefix}_step{step:08d}.pt")
1058
+
1059
+ def _existing(self):
1060
+ out = []
1061
+ for p in glob.glob(os.path.join(self.dir, f"{self.prefix}_step*.pt")):
1062
+ m = re.search(r"_step(\d+)\.pt$", os.path.basename(p))
1063
+ if m:
1064
+ out.append((int(m.group(1)), p))
1065
+ return sorted(out)
1066
+
1067
+ def save(self, step, model, opt, loss, last_val, cfg):
1068
+ tmp = self._path(step) + ".tmp"
1069
+ torch.save({"step": step, "loss": loss, "last_val": last_val,
1070
+ # B/T/global_batch pin the data stride this checkpoint's
1071
+ # step count assumes. world_size is deliberately NOT
1072
+ # recorded: grad_accum = global_batch // (B*T*world_size)
1073
+ # already makes tokens-per-step (and skip_batches'
1074
+ # target offset) independent of how many ranks are
1075
+ # running, so a GPU-count change alone is not a mismatch.
1076
+ "run_cfg": {"B": cfg.B, "T": cfg.T,
1077
+ "global_batch": cfg.global_batch},
1078
+ "model": model.state_dict(),
1079
+ "opt": opt.state_dict()}, tmp)
1080
+ # Rename last: a crash mid-write leaves a .tmp, never a truncated
1081
+ # checkpoint that the next load_latest would pick as the newest.
1082
+ os.replace(tmp, self._path(step))
1083
+ if self.keep > 0:
1084
+ for _, stale in self._existing()[:-self.keep]:
1085
+ os.remove(stale)
1086
+
1087
+ def load_latest(self, model, opt, cfg, master=True):
1088
+ """-> (start_step, loss, last_val). (0, 0.0, nan) when there is
1089
+ nothing to resume."""
1090
+ existing = self._existing()
1091
+ if not existing:
1092
+ return 0, 0.0, float("nan")
1093
+ _, path = existing[-1]
1094
+ return self.load(path, model, opt, cfg, master)
1095
+
1096
+ def load(self, path, model, opt, cfg, master=True):
1097
+ """Load a specific checkpoint file (e.g. from --resume-from), rather
1098
+ than auto-discovering the latest one in self.dir.
1099
+ -> (start_step, loss, last_val)."""
1100
+ blob = torch.load(path, map_location="cpu", weights_only=True)
1101
+ model.load_state_dict(blob["model"])
1102
+ if not opt.load_state_dict(blob["opt"]):
1103
+ opt.resync_masters()
1104
+ run_cfg = blob.get("run_cfg")
1105
+ if master and run_cfg is not None:
1106
+ mismatches = [f"{k} checkpoint={v!r} current={getattr(cfg, k)!r}"
1107
+ for k, v in run_cfg.items() if getattr(cfg, k) != v]
1108
+ if mismatches:
1109
+ print("[WARN] resume config differs from the checkpoint's — "
1110
+ "the data offset this step count assumes will be wrong:")
1111
+ for m in mismatches:
1112
+ print(f" {m}")
1113
+ print(f"[resume] loaded {path}")
1114
+ return (int(blob["step"]), float(blob["loss"]),
1115
+ float(blob.get("last_val", float("nan"))))
1116
+
1117
+
1118
+ # =============================================================================
1119
+ # 6. Training.
1120
+ # =============================================================================
1121
+ def _env_int(name, default=0):
1122
+ v = os.environ.get(name)
1123
+ if v in (None, ""):
1124
+ return default
1125
+ try:
1126
+ return int(v)
1127
+ except ValueError:
1128
+ print(f"[WARN] ignoring non-integer {name}={v!r}")
1129
+ return default
1130
+
1131
+
1132
+ def seed_everything(seed):
1133
+ """Every RNG the run can draw from, off one seed.
1134
+ torch's CPU generator is the one that matters here (the model is built on
1135
+ CPU and moved), but `random` / numpy are seeded too so that anything added
1136
+ later — a shuffled shard order, a sampled probe — is reproducible by
1137
+ default rather than by accident.
1138
+ """
1139
+ random.seed(seed)
1140
+ np.random.seed(seed)
1141
+ torch.manual_seed(seed)
1142
+ torch.cuda.manual_seed_all(seed)
1143
+
1144
+
1145
+ def apply_determinism(mode):
1146
+ """Pin the run-to-run drift that seeding does not cover.
1147
+ Three sources, all on the backward pass:
1148
+ * nn.Embedding's dense backward scatters one atomicAdd per token into
1149
+ [vocab, d_model]; with weight tying that is also the lm_head gradient,
1150
+ so it is the largest reduction in the step;
1151
+ * the flash / mem-efficient attention backward splits dq over the key
1152
+ blocks and atomicAdds them together;
1153
+ * cuDNN picks kernels by wall-clock timing when `benchmark` is on, so the
1154
+ chosen algorithm can differ between runs of the same shape.
1155
+ Each is a float reduction in an order the scheduler chooses, so two runs
1156
+ diverge in the last mantissa bits of the first step and, through the
1157
+ optimizer, visibly by a few hundred steps.
1158
+ `strict` raises if some op has no deterministic kernel — better than
1159
+ reporting a reproducible run that is not one. `warn` downgrades that to a
1160
+ warning and keeps training; `off` restores the default fast kernels.
1161
+ """
1162
+ if mode == "off":
1163
+ return
1164
+ torch.backends.cudnn.benchmark = False
1165
+ torch.backends.cudnn.deterministic = True
1166
+ # Routes embedding-backward and SDPA-backward to their ordered variants, and
1167
+ # enforces the cuBLAS workspace setting made at the top of this file.
1168
+ torch.use_deterministic_algorithms(True, warn_only=(mode == "warn"))
1169
+
1170
+
1171
+ def build_argparser():
1172
+ ap = argparse.ArgumentParser(
1173
+ description="BLU single-GPU training — PyTorch port of bluscript_single.cpp")
1174
+ ap.add_argument("--data-root",
1175
+ default=("/root/.cache/huggingface/hub/datasets--jayasriii--flux_20b/"
1176
+ "/snapshots/ddc6cc7e43365a6b9096c30053df0fb001c94f4c"))
1177
+ ap.add_argument("--device", type=int, default=None,
1178
+ help="CUDA device index (env: BLU_DEVICE)")
1179
+ ap.add_argument("--precision", choices=["bf16-params", "bf16"], default="bf16-params",
1180
+ help="bf16-params (default): Linear weights stored bf16 with "
1181
+ "fp32 master weights in the optimizer — the C++ scheme. "
1182
+ "bf16: fp32 params + autocast (standard PyTorch AMP).")
1183
+ ap.add_argument("--rms-rope", choices=["fused", "eager"], default="fused",
1184
+ help="fused = Triton QK-norm+RoPE kernel (default), "
1185
+ "eager = unfused reference, for A/B on the kernel")
1186
+ ap.add_argument("--compile", action="store_true", help="torch.compile each block")
1187
+ ap.add_argument("--activation-checkpoint", action="store_true",
1188
+ help="recompute each block in backward (trades ~30%% time for memory)")
1189
+ ap.add_argument("--loss-chunks", type=int, default=1,
1190
+ help="chunk lm_head+CE into N pieces with recompute; 1 (default) "
1191
+ "materialises the full logits, as the C++ does")
1192
+ # bluscript_single.cpp defaults unless overridden.
1193
+ ap.add_argument("--B", type=int, default=None, help="micro-batch (env: BLU_B)")
1194
+ ap.add_argument("--T", type=int, default=None)
1195
+ ap.add_argument("--d-model", type=int, default=None, help="smoke-test override")
1196
+ ap.add_argument("--n-layers", type=int, default=None, help="smoke-test override")
1197
+ ap.add_argument("--q-heads", type=int, default=None, help="smoke-test override")
1198
+ ap.add_argument("--kv-heads", type=int, default=None, help="smoke-test override")
1199
+ ap.add_argument("--ffn-hidden", type=int, default=None, help="smoke-test override")
1200
+ ap.add_argument("--vocab-size", type=int, default=None, help="smoke-test override")
1201
+ ap.add_argument("--global-batch", type=int, default=None,
1202
+ help="tokens per optimizer step (env: BLU_GLOBAL_BATCH)")
1203
+ ap.add_argument("--max-steps", type=int, default=None,
1204
+ help="caps the run for benchmarking; warmup/stable shrink "
1205
+ "proportionally so the WSD decay tail stays non-empty "
1206
+ "(env: BLU_MAX_STEPS)")
1207
+ # ---- Muon ----
1208
+ ap.add_argument("--muon-lr", type=float, default=None,
1209
+ help="peak LR for the hidden 2D weights. Under the default "
1210
+ "--muon-adjust-lr match_rms_adamw this is the same unit as "
1211
+ "--max-lr, and defaults to the same value (2e-4) so "
1212
+ "both arms take a matched first run; sweep it upward "
1213
+ "from there (Muon usually takes 2-4x AdamW's LR).")
1214
+ ap.add_argument("--muon-min-lr", type=float, default=None,
1215
+ help="floor of the Muon cosine tail (default 10%% of --muon-lr)")
1216
+ ap.add_argument("--muon-momentum", type=float, default=None,
1217
+ help="torch.optim.Muon momentum (default 0.95, torch's own)")
1218
+ ap.add_argument("--muon-ns-steps", type=int, default=None,
1219
+ help="torch.optim.Muon ns_steps, Newton-Schulz iterations "
1220
+ "(default 5, torch's own); fewer = looser "
1221
+ "orthogonalization, cheaper")
1222
+ ap.add_argument("--muon-adjust-lr", choices=["match_rms_adamw", "original"],
1223
+ default=None,
1224
+ help="torch.optim.Muon adjust_lr_fn. match_rms_adamw "
1225
+ "(default HERE): 0.2*sqrt(max(A,B)), putting --muon-lr "
1226
+ "on the same scale as --max-lr. original (default in "
1227
+ "torch): sqrt(max(1,A/B)), needs a ~5.5x larger "
1228
+ "--muon-lr at d_model=768.")
1229
+ ap.add_argument("--val-freq", type=int, default=None)
1230
+ ap.add_argument("--val-loss-steps", type=int, default=None)
1231
+ ap.add_argument("--diag", action="store_true",
1232
+ help="per-micro-batch loss + depth probes (env: BLU_DIAG=1)")
1233
+ ap.add_argument("--warmup-measure", type=int, default=3,
1234
+ help="steps excluded from the steady-state throughput summary")
1235
+ ap.add_argument("--csv", default="blu_muon_log.csv",
1236
+ help="CSV log path; appended (not truncated) on resume. The "
1237
+ "first 8 columns match the AdamW arm's CSV exactly, so "
1238
+ "the two files plot against each other directly; "
1239
+ "muon_lr is appended as a 9th.")
1240
+ ap.add_argument("--seed", type=int, default=1234)
1241
+ ap.add_argument("--determinism", choices=["strict", "warn", "off"], default="strict",
1242
+ help="strict (default): deterministic CUDA kernels, so a rerun "
1243
+ "of the same launch line reproduces the loss curve bit for "
1244
+ "bit; raises if an op has no deterministic implementation. "
1245
+ "warn: same, but non-deterministic ops warn instead of "
1246
+ "raising. off: fastest kernels, results drift between runs.")
1247
+ # checkpointing
1248
+ ap.add_argument("--no-checkpointing", action="store_true")
1249
+ ap.add_argument("--ckpt-dir", default=None)
1250
+ ap.add_argument("--ckpt-prefix", default=None)
1251
+ ap.add_argument("--ckpt-freq", type=int, default=None)
1252
+ ap.add_argument("--ckpt-keep", type=int, default=None)
1253
+ ap.add_argument("--resume-from", default=None,
1254
+ help="path to a specific *_step*.pt checkpoint to resume "
1255
+ "from, overriding auto-discovery of the latest "
1256
+ "checkpoint in --ckpt-dir")
1257
+ return ap
1258
+
1259
+
1260
+ def config_from_args(args):
1261
+ cfg = ModelConfig()
1262
+ for f in ("B", "T", "d_model", "n_layers", "q_heads", "kv_heads", "ffn_hidden",
1263
+ "vocab_size", "global_batch", "val_freq", "val_loss_steps",
1264
+ "ckpt_dir", "ckpt_prefix", "ckpt_freq", "ckpt_keep",
1265
+ "muon_lr", "muon_min_lr", "muon_momentum", "muon_ns_steps",
1266
+ "muon_adjust_lr"):
1267
+ v = getattr(args, f, None)
1268
+ if v is not None:
1269
+ setattr(cfg, f, v)
1270
+ # A --muon-lr given without a floor keeps the 10% ratio the WSD tail assumes,
1271
+ # instead of silently decaying to the default 2e-5 from wherever it starts.
1272
+ if args.muon_lr is not None and args.muon_min_lr is None:
1273
+ cfg.muon_min_lr = 0.1 * cfg.muon_lr
1274
+ cfg.fused_rms_rope = (args.rms_rope == "fused")
1275
+ cfg.activation_checkpoint = args.activation_checkpoint
1276
+ cfg.loss_chunks = args.loss_chunks
1277
+ cfg.checkpointing = not args.no_checkpointing
1278
+ cfg.context_length = max(cfg.context_length, cfg.T)
1279
+
1280
+ # Mirrors the BLU_MAX_STEPS branch in bluscript_single.cpp: shrink warmup and
1281
+ # stable with the run, or a short benchmark leaves no decay tail (the C++
1282
+ # asserts on it; get_lr here would clamp, which is worse — it would silently
1283
+ # report a schedule the config does not describe).
1284
+ if args.max_steps is not None:
1285
+ ms = args.max_steps
1286
+ cfg.max_steps = ms
1287
+ cfg.warmup_steps = max(1, ms // 10)
1288
+ cfg.stable_steps = max(1, ms // 2)
1289
+ if cfg.warmup_steps + cfg.stable_steps >= cfg.max_steps - 1:
1290
+ cfg.stable_steps = max(1, cfg.max_steps - cfg.warmup_steps - 2)
1291
+ print(f"[bench] max_steps={ms} (warmup={cfg.warmup_steps} "
1292
+ f"stable={cfg.stable_steps})")
1293
+ return cfg
1294
+
1295
+
1296
+ def validate_config(cfg, world_size=1, master=True):
1297
+ """Cheap checks; each one is a silent-corruption or out-of-bounds failure
1298
+ mode if it does not hold."""
1299
+ bad = []
1300
+ if cfg.T > cfg.context_length:
1301
+ bad.append(f"T ({cfg.T}) > context_length ({cfg.context_length}): "
1302
+ "the RoPE cache is too short")
1303
+ if cfg.q_heads % cfg.kv_heads != 0:
1304
+ bad.append(f"q_heads ({cfg.q_heads}) not divisible by kv_heads ({cfg.kv_heads})")
1305
+ if cfg.head_dim not in (64, 128):
1306
+ bad.append(f"the fused kernels support head_dim 64 or 128, got {cfg.head_dim}")
1307
+ # Under DDP each rank runs B*T tokens per micro-step, so the world's tokens
1308
+ # per micro-step are B*T*world_size; global_batch must land on a whole
1309
+ # number of those or grad_accum below truncates and quietly under-shoots
1310
+ # the configured optimizer-step token count.
1311
+ if cfg.global_batch % (cfg.B * cfg.T * world_size) != 0:
1312
+ bad.append(f"global_batch ({cfg.global_batch}) not divisible by "
1313
+ f"B*T*world_size ({cfg.B * cfg.T * world_size})")
1314
+ if cfg.warmup_steps < 1:
1315
+ bad.append("warmup_steps must be >= 1 (get_lr divides by it)")
1316
+ if cfg.warmup_steps + cfg.stable_steps >= cfg.max_steps - 1:
1317
+ bad.append(f"warmup_steps + stable_steps "
1318
+ f"({cfg.warmup_steps + cfg.stable_steps}) leaves no decay tail "
1319
+ f"before max_steps ({cfg.max_steps})")
1320
+ if not hasattr(torch.optim, "Muon"):
1321
+ bad.append(f"torch {torch.__version__} has no torch.optim.Muon — it first "
1322
+ "appears in the 2.9 docs. Upgrade torch, or run the AdamW "
1323
+ "arm (train.py).")
1324
+ if cfg.muon_ns_steps < 1:
1325
+ bad.append("muon_ns_steps must be >= 1 (0 leaves the update unorthogonalized, "
1326
+ "which is not Muon)")
1327
+ if cfg.muon_momentum <= 0.0 or cfg.muon_momentum >= 1.0:
1328
+ bad.append(f"muon_momentum must be in (0,1), got {cfg.muon_momentum}")
1329
+ if cfg.q_heads * cfg.head_dim != cfg.d_model and master:
1330
+ print(f"WARNING: q_heads*head_dim ({cfg.q_heads * cfg.head_dim}) "
1331
+ f"!= d_model ({cfg.d_model})")
1332
+ if bad:
1333
+ raise SystemExit("ERROR: " + "\nERROR: ".join(bad))
1334
+
1335
+
1336
+ def main():
1337
+ args = build_argparser().parse_args()
1338
+
1339
+ # Env overrides, for parity with the C++ launch line. An explicit CLI flag
1340
+ # wins over the environment.
1341
+ if args.B is None:
1342
+ args.B = _env_int("BLU_B", 0) or None
1343
+ if args.global_batch is None:
1344
+ args.global_batch = _env_int("BLU_GLOBAL_BATCH", 0) or None
1345
+ if args.max_steps is None:
1346
+ args.max_steps = _env_int("BLU_MAX_STEPS", 0) or None
1347
+ diag = args.diag or _env_int("BLU_DIAG", 0) != 0
1348
+
1349
+ if not torch.cuda.is_available():
1350
+ raise SystemExit("ERROR: no CUDA device available")
1351
+
1352
+ # ---- DDP bootstrap ----
1353
+ # torchrun sets RANK/LOCAL_RANK/WORLD_SIZE; their absence means a plain
1354
+ # single-process launch, where --device / BLU_DEVICE pick the GPU instead.
1355
+ # torch.cuda.set_device MUST happen before init_process_group: NCCL creates
1356
+ # its context on whatever device is current at that call, so doing it after
1357
+ # would leave every rank on this node contending for device 0.
1358
+ ddp = int(os.environ.get("RANK", -1)) != -1
1359
+ if ddp:
1360
+ ddp_rank = int(os.environ["RANK"])
1361
+ ddp_local_rank = int(os.environ["LOCAL_RANK"])
1362
+ ddp_world_size = int(os.environ["WORLD_SIZE"])
1363
+ master_process = ddp_rank == 0
1364
+ if ddp_local_rank >= torch.cuda.device_count():
1365
+ raise SystemExit(f"ERROR: LOCAL_RANK={ddp_local_rank} but only "
1366
+ f"{torch.cuda.device_count()} CUDA device(s) visible")
1367
+ torch.cuda.set_device(ddp_local_rank)
1368
+ device = torch.device("cuda", ddp_local_rank)
1369
+ dist.init_process_group(backend="nccl")
1370
+ if master_process and (args.device is not None
1371
+ or os.environ.get("BLU_DEVICE") not in (None, "")):
1372
+ print("[WARN] --device / BLU_DEVICE is ignored under torchrun — "
1373
+ "LOCAL_RANK selects each rank's GPU.")
1374
+ else:
1375
+ ddp_rank, ddp_local_rank, ddp_world_size = 0, 0, 1
1376
+ master_process = True
1377
+ if args.device is None:
1378
+ args.device = _env_int("BLU_DEVICE", 0)
1379
+ if args.device >= torch.cuda.device_count():
1380
+ raise SystemExit(f"ERROR: BLU_DEVICE={args.device} but only "
1381
+ f"{torch.cuda.device_count()} CUDA device(s) visible")
1382
+ torch.cuda.set_device(args.device)
1383
+ device = torch.device("cuda", args.device)
1384
+
1385
+ seed_everything(args.seed)
1386
+ apply_determinism(args.determinism)
1387
+ # TF32 loses mantissa bits but is deterministic, so it is orthogonal to the
1388
+ # setting above and stays on.
1389
+ torch.set_float32_matmul_precision("high")
1390
+
1391
+ cfg = config_from_args(args)
1392
+ if cfg.fused_rms_rope and not _HAS_TRITON:
1393
+ if master_process:
1394
+ print("[WARN] Triton not available — falling back to the eager QK-norm+RoPE")
1395
+ cfg.fused_rms_rope = False
1396
+ validate_config(cfg, ddp_world_size, master_process)
1397
+
1398
+ # Tokens consumed per optimizer step = B * T * grad_accum * world_size.
1399
+ grad_accum = cfg.global_batch // (cfg.B * cfg.T * ddp_world_size)
1400
+ tokens_per_step = cfg.B * cfg.T * grad_accum * ddp_world_size
1401
+
1402
+ if master_process:
1403
+ print("================ Configuration ================")
1404
+ mode_str = (f"DDP x{ddp_world_size} (torchrun, NCCL)" if ddp
1405
+ else "SINGLE GPU (no torchrun / DDP / ZeRO)")
1406
+ print(f" mode : {mode_str}")
1407
+ print(f" cuda device : {ddp_local_rank if ddp else args.device} "
1408
+ f"({torch.cuda.get_device_name(device)})")
1409
+ print(f" d_model : {cfg.d_model}")
1410
+ print(f" n_layers : {cfg.n_layers}")
1411
+ print(f" q_heads/kv_heads: {cfg.q_heads} / {cfg.kv_heads} "
1412
+ f"(G={cfg.q_heads // cfg.kv_heads})")
1413
+ print(f" head_dim : {cfg.head_dim}")
1414
+ print(f" ffn_hidden : {cfg.ffn_hidden} (SwiGLU)")
1415
+ print(f" context_length : {cfg.context_length}")
1416
+ print(f" rope_theta : {cfg.rope_theta:g} (NeoX half-split)")
1417
+ print(f" vocab_size : {cfg.vocab_size}")
1418
+ print(f" weight_tying : {'ON' if cfg.weight_tying else 'OFF'}")
1419
+ print(f" B x T : {cfg.B} x {cfg.T}")
1420
+ print(f" global_batch : {cfg.global_batch} (grad_accum={grad_accum}, "
1421
+ f"world_size={ddp_world_size})")
1422
+ print(" optimizer : MUON (hidden 2D) + AdamW (embed/head/1D)")
1423
+ print(f" max_lr / min_lr : {cfg.max_lr:g} / {cfg.min_lr:g} (AdamW group)")
1424
+ print(f" muon lr / min : {cfg.muon_lr:g} / {cfg.muon_min_lr:g} "
1425
+ f"(adjust_lr_fn={cfg.muon_adjust_lr})")
1426
+ print(f" muon momentum : {cfg.muon_momentum:g} | NS steps {cfg.muon_ns_steps}")
1427
+ print(f" warmup / steps : {cfg.warmup_steps} / {cfg.max_steps}")
1428
+ print(f" WSD stable_steps: {cfg.stable_steps} "
1429
+ f"(decay tail={cfg.max_steps - cfg.warmup_steps - cfg.stable_steps})")
1430
+ print(f" precision : {args.precision}")
1431
+ print(f" seed / determ. : {args.seed} / {args.determinism}")
1432
+ print(f" rms+rope : {'fused' if cfg.fused_rms_rope else 'eager'}"
1433
+ f" | compile {args.compile} | act-ckpt {cfg.activation_checkpoint}"
1434
+ f" | loss_chunks {cfg.loss_chunks}")
1435
+ print("===============================================", flush=True)
1436
+
1437
+ # ---- model ----
1438
+ # Built and seeded identically on every rank (same seed, deterministic CPU
1439
+ # init); DDP's constructor broadcasts rank 0's parameters/buffers anyway,
1440
+ # so this is belt-and-suspenders against any future source of divergence.
1441
+ if master_process:
1442
+ print(f"Initializing model on CUDA device {device.index}...", flush=True)
1443
+ raw_model = BLU(cfg).to(device)
1444
+ raw_model.apply_param_dtype(args.precision)
1445
+ params = list(raw_model.parameters())
1446
+ n_params = sum(p.numel() for p in params)
1447
+ if master_process:
1448
+ print(f"Number of parameters: {n_params}")
1449
+ if args.compile:
1450
+ for blk in raw_model.blocks:
1451
+ blk.compile()
1452
+
1453
+ # DDP wraps the (already dtype-cast, already compiled) raw module. Buffer
1454
+ # broadcasting is off: rope_cos/rope_sin are deterministic functions of
1455
+ # cfg, identical on every rank already, so syncing them every forward
1456
+ # would just be wasted communication.
1457
+ # if ddp:
1458
+ # model = DDP(raw_model, device_ids=[ddp_local_rank], broadcast_buffers=False)
1459
+ # else:
1460
+ # model = raw_model
1461
+
1462
+ if ddp:
1463
+ model = DDP(raw_model, device_ids=[ddp_local_rank], broadcast_buffers=False)
1464
+ torch.cuda.synchronize() # let DDP's constructor-time NCCL broadcast fully land
1465
+ # before DataLoaderLite's independent copy_stream starts
1466
+ # issuing its own async H2D work — see [crash investigation]
1467
+ else:
1468
+ model = raw_model
1469
+
1470
+
1471
+ # ---- optimizer ----
1472
+ # named_parameters(), not parameters(): the routing is by name (the
1473
+ # embedding and lm_head stay on AdamW). Duplicates are dropped by default,
1474
+ # so a weight-tied head is visited once, as `wte.weight`. Built from
1475
+ # raw_model so names never carry DDP's "module." prefix.
1476
+ opt = MasterMuon(raw_model.named_parameters(), cfg)
1477
+ if master_process:
1478
+ print(f"muon params: {opt.n_muon} | adamw decayed: {opt.n_decay}, "
1479
+ f"non-decayed: {opt.n_nodecay}")
1480
+ # Static footprint. Nothing here is sharded — plain DDP replicates the
1481
+ # full optimizer state on every rank (the whole cost of not running
1482
+ # ZeRO-1), so this is the per-GPU charge before a single activation is
1483
+ # allocated.
1484
+ gb = 2 ** 30
1485
+ param_bytes = sum(p.numel() * p.element_size() for p in params)
1486
+ grad_bytes = param_bytes # grads share the param dtype
1487
+ master_bytes, landing_bytes, _, _ = opt.state_breakdown()
1488
+ # Both optimizers allocate their state on the first step(), so at this point
1489
+ # the measured values are 0 and the report has to predict them.
1490
+ mv_bytes = opt.expected_mv_bytes()
1491
+ mom_bytes = opt.expected_muon_mom_bytes()
1492
+ static_bytes = (param_bytes + grad_bytes + master_bytes + landing_bytes
1493
+ + mv_bytes + mom_bytes)
1494
+ if master_process:
1495
+ print(f"[{'ddp' if ddp else 'single'}] static footprint (per GPU: "
1496
+ f"params + grads + optimizer state)")
1497
+ print(f" parameters : {param_bytes / gb:7.2f} GiB")
1498
+ print(f" gradients : {grad_bytes / gb:7.2f} GiB")
1499
+ print(f" fp32 master weights : {master_bytes / gb:7.2f} GiB")
1500
+ print(f" fp32 grad landing bufs : {landing_bytes / gb:7.2f} GiB")
1501
+ print(f" Muon momentum : {mom_bytes / gb:7.2f} GiB")
1502
+ print(f" Adam exp_avg + exp_avg_sq: {mv_bytes / gb:6.2f} GiB")
1503
+ print(f" static total : {static_bytes / gb:7.2f} GiB", flush=True)
1504
+
1505
+ # ---- data ----
1506
+ # Sharded by rank: each rank reads a disjoint, interleaved slice of the
1507
+ # same shard list (see DataLoaderLite), so DDP's cross-rank gradient
1508
+ # average is over world_size DIFFERENT windows, not world_size copies of
1509
+ # the same one.
1510
+ train_loader = DataLoaderLite(cfg.B, cfg.T, "train", args.data_root, device,
1511
+ master=master_process, rank=ddp_rank,
1512
+ world_size=ddp_world_size)
1513
+ # rank=0, world_size=1 UNCONDITIONALLY, on every rank: val_loss_steps
1514
+ # windows at fixed token offsets 0, BT, 2BT, ... Sharding this by
1515
+ # ddp_rank/ddp_world_size like the train loader would make the reported
1516
+ # val_loss a function of how many GPUs happen to be running (world_size=2
1517
+ # averages 2x as many, differently-offset windows as world_size=1) —
1518
+ # the two numbers would not be comparable across a GPU-count change, which
1519
+ # is what produced the apparent val_loss jump on the 2-GPU -> 1-GPU
1520
+ # resume. Every rank now redundantly evaluates the identical fixed
1521
+ # windows, so the metric means the same thing regardless of world_size.
1522
+ val_loader = DataLoaderLite(cfg.B, cfg.T, "val", args.data_root, device,
1523
+ master=master_process, rank=0, world_size=1)
1524
+
1525
+ # ---- checkpointing: manager + auto-resume ----
1526
+ # Every rank independently loads the same checkpoint file from the same
1527
+ # (assumed shared) filesystem — a read, so no race — which keeps model,
1528
+ # optimizer state, and start_step identical across ranks without an
1529
+ # explicit broadcast. Only rank 0 ever WRITES one (below).
1530
+ start_step, ckpt = 0, None
1531
+ last_val = float("nan")
1532
+ if cfg.checkpointing or args.resume_from:
1533
+ ckpt = CheckpointManager(cfg.ckpt_dir, cfg.ckpt_prefix, cfg.ckpt_keep)
1534
+ if args.resume_from:
1535
+ if master_process:
1536
+ print(f"[resume] loading explicit checkpoint: {args.resume_from}")
1537
+ start_step, resume_loss, last_val = ckpt.load(
1538
+ args.resume_from, raw_model, opt, cfg, master_process)
1539
+ else:
1540
+ start_step, resume_loss, last_val = ckpt.load_latest(
1541
+ raw_model, opt, cfg, master_process)
1542
+ if start_step > 0:
1543
+ if master_process:
1544
+ print(f"[resume] at step {start_step} (loss {resume_loss:.6f}, "
1545
+ f"last val {last_val:.6f})")
1546
+ # Skip the batches already consumed by steps [0, start_step).
1547
+ # global_batch // (B*T*world_size) makes tokens-per-step (and so
1548
+ # this target offset) independent of world_size, so a GPU-count
1549
+ # change alone lands on the correct byte offset as long as B, T,
1550
+ # and global_batch are unchanged (checked above).
1551
+ train_loader.skip_batches(start_step * grad_accum)
1552
+
1553
+ # ---- CSV log ----
1554
+ # Appended (not truncated) when resuming so a restart does not erase the
1555
+ # history of the run it is continuing. Rank 0 only — every rank would
1556
+ # otherwise open and write the same path.
1557
+ log_file = None
1558
+ if args.csv and master_process:
1559
+ resuming = start_step > 0 and os.path.exists(args.csv)
1560
+ log_file = open(args.csv, "a" if resuming else "w")
1561
+ if not resuming:
1562
+ log_file.write("step,train_loss,val_loss,tok_per_sec,grad_norm,dt_ms,lr,"
1563
+ "elapsed_min,muon_lr\n")
1564
+ log_file.flush()
1565
+
1566
+ amp = torch.autocast(device_type="cuda", dtype=torch.bfloat16)
1567
+ probes = [L for L in cfg.probe_layers if L < cfg.n_layers]
1568
+ steady_dts = []
1569
+ t_train_start = time.perf_counter()
1570
+
1571
+ # ---- Ctrl-C: finish the in-flight step, save a checkpoint, exit clean ----
1572
+ # Registered here (not earlier) so a Ctrl-C during setup/data-loading still
1573
+ # gets Python's default immediate-abort behaviour — only the training loop
1574
+ # below gets the graceful path. A second Ctrl-C re-arms the default handler
1575
+ # and re-raises, for when the checkpoint save itself needs interrupting.
1576
+ interrupted = False
1577
+
1578
+ def _handle_sigint(signum, frame):
1579
+ nonlocal interrupted
1580
+ if interrupted:
1581
+ signal.signal(signal.SIGINT, signal.SIG_DFL)
1582
+ raise KeyboardInterrupt
1583
+ interrupted = True
1584
+ if master_process:
1585
+ print("\n[ctrl-c] finishing this step, then checkpointing and exiting "
1586
+ "(press again to abort immediately)...", flush=True)
1587
+
1588
+ # A terminal Ctrl-C is delivered to every process in the foreground
1589
+ # process group, so under torchrun every rank sets `interrupted` at
1590
+ # roughly the same time and the ranks stay in lockstep for the collective
1591
+ # ops (DDP's grad all-reduce, the checkpoint barrier below) still ahead.
1592
+ signal.signal(signal.SIGINT, _handle_sigint)
1593
+
1594
+ if master_process:
1595
+ print("\nStarting training" +
1596
+ (f" (resumed at step {start_step})" if start_step else "") + "...\n", flush=True)
1597
+
1598
+ for step in range(start_step, cfg.max_steps):
1599
+ # ---------- periodic validation ----------
1600
+ if cfg.val_loss_steps > 0 and (step % cfg.val_freq == 0 or step == cfg.max_steps - 1):
1601
+ model.eval()
1602
+ val_loader.reset()
1603
+ acc = 0.0
1604
+ with torch.no_grad(), amp:
1605
+ for _ in range(cfg.val_loss_steps):
1606
+ x, y, _, _ = val_loader.next_batch()
1607
+ acc += model(x, y)[1].item() / cfg.val_loss_steps
1608
+ model.train()
1609
+ # val_loader is rank=0/world_size=1 on every rank (see its
1610
+ # construction above), so every rank just computed the identical
1611
+ # fixed windows redundantly — no cross-rank reduction needed, and
1612
+ # val_loss no longer depends on how many ranks are running.
1613
+ last_val = acc
1614
+ if master_process:
1615
+ print(f" [val] step {step} | val_loss {last_val:.6f}", flush=True)
1616
+
1617
+ # ---------- training step ----------
1618
+ t0 = time.perf_counter()
1619
+ model.zero_grad(set_to_none=True)
1620
+ # Fresh zeros rather than an in-place scale of last step's accumulator:
1621
+ # multiplying by zero does NOT clear a NaN, so a single non-finite
1622
+ # micro-batch would poison the accumulator for the rest of the run.
1623
+ loss_accum = torch.zeros((), device=device)
1624
+ micro_loss, micro_where = [], []
1625
+
1626
+ for micro in range(grad_accum):
1627
+ x, y, shard, offset = train_loader.next_batch()
1628
+ # DDP allreduces gradients on every .backward(); for grad_accum > 1
1629
+ # only the LAST micro-batch of the step should pay for that — the
1630
+ # others accumulate purely locally under no_sync().
1631
+ sync_ctx = (nullcontext() if (not ddp or micro == grad_accum - 1)
1632
+ else model.no_sync())
1633
+ with sync_ctx:
1634
+ with amp:
1635
+ _, loss = model(x, y)
1636
+ loss_accum += loss.detach()
1637
+ # Per-micro-batch loss (--diag). The reported train loss is the mean
1638
+ # over grad_accum micro-steps, so a single pathological micro-batch
1639
+ # is diluted while its gradient dominates the accumulated norm
1640
+ # almost undiluted — which is how a modest loss bump can sit next to
1641
+ # a large norm spike. `micro_where` names the shard and token offset
1642
+ # so the responsible bytes can be decoded offline.
1643
+ #
1644
+ # Costs a device sync per micro-step, so it is off by default.
1645
+ if diag:
1646
+ micro_loss.append(loss.item())
1647
+ micro_where.append((shard, offset))
1648
+ (loss / grad_accum).backward()
1649
+
1650
+ if ddp:
1651
+ # Average each rank's accumulated (summed) loss BEFORE dividing by
1652
+ # grad_accum, so loss_val is the mean over world_size*grad_accum
1653
+ # micro-batches — the true loss behind this optimizer step, not
1654
+ # just rank 0's slice. Also doubles as a sync point: a NaN on any
1655
+ # rank poisons the average for every rank, so the check below
1656
+ # raises everywhere instead of only where it originated (which
1657
+ # would otherwise hang the survivors on the next collective).
1658
+ dist.all_reduce(loss_accum, op=dist.ReduceOp.AVG)
1659
+ loss_val = loss_accum.item() / grad_accum
1660
+ if math.isnan(loss_val) or math.isinf(loss_val):
1661
+ if log_file:
1662
+ log_file.close()
1663
+ # Every rank reaches this identically (the all_reduce above is a
1664
+ # sync point and NaN propagates through AVG), so every rank exits
1665
+ # here together rather than the survivors hanging on the next
1666
+ # collective.
1667
+ if ddp:
1668
+ dist.destroy_process_group()
1669
+ raise SystemExit(f"ERROR: NaN/Inf loss at step {step}")
1670
+
1671
+ # Depth probe (--diag), read BEFORE the update, so these are the
1672
+ # pre-update gradients. Compare the profile on a spiking step against a
1673
+ # quiet one: a blow-up that begins mid-stack and grows toward the loss
1674
+ # points at the forward, while a flat multiplicative lift across all
1675
+ # depths points at the optimizer/update.
1676
+ depth = []
1677
+ if diag:
1678
+ for L in probes:
1679
+ a = raw_model.blocks[L].attn
1680
+ gn = a.wQ.weight.grad
1681
+ depth.append((L,
1682
+ gn.float().norm().item() if gn is not None else -1.0,
1683
+ a.q_gamma.abs().max().item(),
1684
+ a.wQ.weight.abs().max().float().item()))
1685
+
1686
+ grad_norm = clip_grads_(params, cfg.grad_clip).item()
1687
+ lr = get_lr(step, cfg) # AdamW group
1688
+ muon_lr = get_lr(step, cfg, cfg.muon_lr, cfg.muon_min_lr)
1689
+ opt.step(lr, muon_lr)
1690
+
1691
+ torch.cuda.synchronize() # make the per-step timing reflect completed GPU work
1692
+ dt = time.perf_counter() - t0
1693
+ tok_per_sec = tokens_per_step / dt
1694
+ elapsed_min = (time.perf_counter() - t_train_start) / 60.0
1695
+ if step - start_step >= args.warmup_measure:
1696
+ steady_dts.append(dt)
1697
+
1698
+ # ---------- per-step report (rank 0 only — every rank's numbers agree
1699
+ # on loss/grad_norm post-allreduce, so printing from every rank would
1700
+ # only duplicate lines, not add information) ----------
1701
+ if master_process:
1702
+ vs = " - " if last_val != last_val else f"{last_val:.6f}"
1703
+ print(f"step {step:6d} | train {loss_val:.6f} | val {vs} "
1704
+ f"| tok/s {tok_per_sec:10.2f} | norm {grad_norm:.4f} "
1705
+ f"| dt {dt * 1e3:8.2f} ms | lr {lr:.4e} | mlr {muon_lr:.4e} "
1706
+ f"| mem {torch.cuda.max_memory_allocated() / gb:5.2f}/"
1707
+ f"{torch.cuda.max_memory_reserved() / gb:5.2f} GiB", flush=True)
1708
+
1709
+ if diag and micro_loss:
1710
+ # Flag micro-batches that stand out against this step's own median.
1711
+ med = sorted(micro_loss)[len(micro_loss) // 2]
1712
+ marks = " ".join(("*" if m > med + 0.5 else "") + f"{m:.3f}"
1713
+ for m in micro_loss)
1714
+ print(f" [micro] med {med:.4f} | {marks}")
1715
+ for m, (ml, (shard, offset)) in enumerate(zip(micro_loss, micro_where)):
1716
+ if ml > med + 0.5:
1717
+ print(f" [micro] OUTLIER m={m} loss {ml:.4f} (median {med:.4f}) "
1718
+ f"from shard {shard} token offset {offset}")
1719
+ if diag and depth:
1720
+ print(" [depth] ||g(wQ)|| by layer:"
1721
+ + "".join(f" L{L} {g:.3e}" for L, g, _, _ in depth)
1722
+ + " | max|q_gamma|:"
1723
+ + "".join(f" L{L} {gm:.4f}" for L, _, gm, _ in depth)
1724
+ + " | max|wQ|:"
1725
+ + "".join(f" L{L} {w:.4f}" for L, _, _, w in depth), flush=True)
1726
+
1727
+ if log_file:
1728
+ log_file.write(f"{step},{loss_val:.6f},{last_val:.6f},{tok_per_sec:.6f},"
1729
+ f"{grad_norm:.6f},{dt * 1e3:.6f},{lr:.6f},{elapsed_min:.6f},"
1730
+ f"{muon_lr:.6f}\n")
1731
+ log_file.flush()
1732
+
1733
+ if ddp:
1734
+ # `interrupted` is set by this process's own SIGINT handler, so in
1735
+ # principle one rank could see it before another (a signal isn't
1736
+ # guaranteed simultaneous delivery, even though in practice a
1737
+ # terminal Ctrl-C hits every rank in the same process group at
1738
+ # once). Left asymmetric, it would desync the checkpoint barrier
1739
+ # and the `break` below against ranks still issuing DDP's regular
1740
+ # per-step collectives — a hang, not just a wrong answer. This
1741
+ # runs unconditionally every step (never gated on a per-rank
1742
+ # value), so it can never itself mismatch across ranks.
1743
+ flag = torch.tensor(1.0 if interrupted else 0.0, device=device)
1744
+ dist.all_reduce(flag, op=dist.ReduceOp.MAX)
1745
+ interrupted = flag.item() > 0
1746
+
1747
+ # ---------- periodic checkpoint (also fires once on Ctrl-C) ----------
1748
+ # Labeled step+1 (the next iteration to run) so load_latest resumes
1749
+ # exactly where this save left off — no duplicated or skipped iteration.
1750
+ # Captures weights + Adam m/v AFTER this step's update. Rank 0 only —
1751
+ # every rank's raw_model/opt state is identical post-allreduce, so a
1752
+ # second writer would only race the first on the same path. The
1753
+ # barrier holds every rank at this point until the save (or its
1754
+ # failure) is done, so a slow write can't let other ranks race ahead
1755
+ # into the next step's collectives before the checkpoint is durable.
1756
+ if cfg.checkpointing and ((step + 1) % cfg.ckpt_freq == 0
1757
+ or step == cfg.max_steps - 1
1758
+ or interrupted):
1759
+ if master_process:
1760
+ try:
1761
+ ckpt.save(step + 1, raw_model, opt, loss_val, last_val, cfg)
1762
+ except Exception as e: # noqa: BLE001 — advisory
1763
+ print(f"[WARN] checkpoint save at step {step + 1} failed: {e}")
1764
+ if ddp:
1765
+ dist.barrier()
1766
+
1767
+ if interrupted:
1768
+ if not cfg.checkpointing and master_process:
1769
+ print(f"[ctrl-c] checkpointing is disabled (--no-checkpointing); "
1770
+ f"stopping at step {step} without saving.", flush=True)
1771
+ break
1772
+
1773
+ # Reclaim the extra allocator high-water mark from the first couple of
1774
+ # steps' one-time warmup allocations (autograd graph caches, etc.).
1775
+ if step in (start_step, start_step + 1):
1776
+ torch.cuda.synchronize()
1777
+ torch.cuda.empty_cache()
1778
+ torch.cuda.reset_peak_memory_stats()
1779
+
1780
+ torch.cuda.synchronize()
1781
+ if log_file:
1782
+ log_file.close()
1783
+
1784
+ # ---- steady-state summary ----
1785
+ # dt/throughput are this rank's own timing; under DDP all ranks run in
1786
+ # lockstep (grad sync + the checkpoint barrier), so rank 0's numbers are
1787
+ # representative and `tokens_per_step` already counts every rank's tokens.
1788
+ if steady_dts and master_process:
1789
+ xs = sorted(steady_dts)
1790
+ med = xs[len(xs) // 2]
1791
+ mean = sum(xs) / len(xs)
1792
+ alloc = torch.cuda.max_memory_allocated() / gb
1793
+ reserved = torch.cuda.max_memory_reserved() / gb
1794
+ # Measured now that Adam's m/v exist, rather than the startup estimate.
1795
+ static = (param_bytes + grad_bytes + opt.state_bytes()) / gb
1796
+ mode_str = f"DDP x{ddp_world_size}" if ddp else "single GPU"
1797
+ print("\n" + "=" * 78)
1798
+ print(f" STEADY-STATE SUMMARY — {mode_str} optimizer=muon+adamw "
1799
+ f"precision={args.precision} "
1800
+ f"rms+rope={'fused' if cfg.fused_rms_rope else 'eager'}")
1801
+ print(f" torch.optim.Muon lr={cfg.muon_lr:g} "
1802
+ f"adjust_lr_fn={cfg.muon_adjust_lr} "
1803
+ f"momentum={cfg.muon_momentum:g} ns_steps={cfg.muon_ns_steps}")
1804
+ print(f" ({len(xs)} steps, first {args.warmup_measure} excluded)")
1805
+ print("=" * 78)
1806
+ print(f" step time : {med * 1e3:10.2f} ms (mean {mean * 1e3:.2f})")
1807
+ print(f" throughput : {tokens_per_step / med:10.1f} tok/s")
1808
+ print(f" tokens / step : {tokens_per_step:10d}")
1809
+ print(f" static (p+g+opt) : {static:10.2f} GiB (per GPU)")
1810
+ print(f" activations (peak) : {max(alloc - static, 0.0):10.2f} GiB")
1811
+ print(f" peak mem allocated : {alloc:10.2f} GiB")
1812
+ print(f" peak mem reserved : {reserved:10.2f} GiB")
1813
+ print("=" * 78)
1814
+
1815
+ if ddp:
1816
+ dist.destroy_process_group()
1817
+ if master_process:
1818
+ print("\n=== Training complete ===")
1819
+
1820
+
1821
+ if __name__ == "__main__":
1822
+ main()
1823
+
1824
+
1825
+
1826
+
1827
+
1828
+
1829
+
1830
+
1831
+
1832
+
1833
+
1834
+
1835
+
1836
+
1837
+
1838
+
1839
+
1840
+
1841
+
1842
+
1843
+
1844
+
1845
+
1846
+
1847
+
1848
+
1849
+