Upload folder using huggingface_hub
Browse files- 150M_228M_15e4_10B_full.csv +0 -0
- 150M_228M_15e4_10B_full/blumodel_muon_step00015248.pt +3 -0
- 150M_228M_15e4_10B_full/blumodel_muon_step00016201.pt +3 -0
- 150M_228M_15e4_10B_full/blumodel_muon_step00017154.pt +3 -0
- 150M_228M_15e4_10B_full/blumodel_muon_step00018107.pt +3 -0
- 150M_228M_15e4_10B_full/blumodel_muon_step00019060.pt +3 -0
- 150M_228M_15e4_10B_full/blumodel_muon_step00019074.pt +3 -0
- 150M_228M_3e3_6B_incmp.csv +398 -0
- 150M_228M_3e3_6B_incomp/blumodel_muon_step00010706.pt +3 -0
- 150M_228M_75e5_10B_half.csv +0 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00000953.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00001906.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00002859.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00003812.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00004765.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00005718.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00006671.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00007624.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00008577.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00009530.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00010483.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00011436.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00012389.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00013342.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00014295.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00015248.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00016201.pt +3 -0
- 150M_228M_75e5_10B_half/blumodel_muon_step00016332.pt +3 -0
- train_muon.py +1829 -0
150M_228M_15e4_10B_full.csv
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
150M_228M_15e4_10B_full/blumodel_muon_step00015248.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:e2675d12bbea262db9f037d54752b36a1bcda31b8a4c134e750d7c38ad84dc52
|
| 3 |
+
size 2515184199
|
150M_228M_15e4_10B_full/blumodel_muon_step00016201.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:84555bc64bcef2702981e736ac313fedba50cba6bd469c7c64b4d5ca8d3c3add
|
| 3 |
+
size 2515184199
|
150M_228M_15e4_10B_full/blumodel_muon_step00017154.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:80b9c23d170474babe961dc098c1baf1baff746c14cbf76b956c4b0a0d225ef7
|
| 3 |
+
size 2515184199
|
150M_228M_15e4_10B_full/blumodel_muon_step00018107.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:a0090a8afad7b50b640c8f88a73f318437166b319f2cc0bcc0590ff93161dd30
|
| 3 |
+
size 2515184199
|
150M_228M_15e4_10B_full/blumodel_muon_step00019060.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:8cc676ec5c64b3020dabf764d682abbb4820bedd70460f216389cd04d57fa464
|
| 3 |
+
size 2515184199
|
150M_228M_15e4_10B_full/blumodel_muon_step00019074.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:ab976e57eef00ae835c24a03b86a0bbfb7c42a29108cfa7ef47054586f2e5312
|
| 3 |
+
size 2515184199
|
150M_228M_3e3_6B_incmp.csv
ADDED
|
@@ -0,0 +1,398 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
step,train_loss,val_loss,tok_per_sec,grad_norm,dt_ms,lr,elapsed_min,muon_lr
|
| 2 |
+
10309,2.646035,3.057884,241083.302612,0.080317,2174.717180,0.002999,0.036246,0.002999
|
| 3 |
+
10310,2.581861,3.057884,307079.635579,0.083498,1707.335620,0.002999,0.067322,0.002999
|
| 4 |
+
10311,2.633330,3.057884,301997.020155,0.080284,1736.070110,0.002999,0.098820,0.002999
|
| 5 |
+
10312,2.711791,3.057884,317031.979656,0.094343,1653.738530,0.002999,0.126388,0.002999
|
| 6 |
+
10313,2.672399,3.057884,319283.029922,0.100907,1642.079130,0.002999,0.153763,0.002999
|
| 7 |
+
10314,2.651201,3.057884,319127.757069,0.089898,1642.878090,0.002999,0.181149,0.002999
|
| 8 |
+
10315,2.876376,3.057884,319772.240348,0.094825,1639.566960,0.002999,0.208485,0.002999
|
| 9 |
+
10316,2.645422,3.057884,319992.920078,0.098303,1638.436250,0.002999,0.235806,0.002999
|
| 10 |
+
10317,2.569267,3.057884,320027.812964,0.099491,1638.257610,0.002998,0.263117,0.002998
|
| 11 |
+
10318,2.526511,3.057884,319971.904810,0.095968,1638.543860,0.002998,0.290432,0.002998
|
| 12 |
+
10319,2.734799,3.057884,319961.467532,0.085967,1638.597310,0.002998,0.317747,0.002998
|
| 13 |
+
10320,2.614316,3.057884,319905.998325,0.075867,1638.881430,0.002998,0.345069,0.002998
|
| 14 |
+
10321,2.614264,3.057884,319885.991804,0.073269,1638.983930,0.002998,0.372391,0.002998
|
| 15 |
+
10322,2.737565,3.057884,319925.286979,0.072798,1638.782620,0.002997,0.399712,0.002997
|
| 16 |
+
10323,2.688076,3.057884,319928.158712,0.072226,1638.767910,0.002997,0.427032,0.002997
|
| 17 |
+
10324,2.684522,3.057884,319891.013694,0.076770,1638.958200,0.002997,0.454354,0.002997
|
| 18 |
+
10325,2.580991,3.057884,319816.716366,0.078332,1639.338950,0.002997,0.481682,0.002997
|
| 19 |
+
10326,2.612901,3.057884,319771.171562,0.082629,1639.572440,0.002996,0.509014,0.002996
|
| 20 |
+
10327,2.681869,3.057884,319822.838371,0.085332,1639.307570,0.002996,0.536342,0.002996
|
| 21 |
+
10328,2.643904,3.057884,319831.128226,0.080611,1639.265080,0.002996,0.563668,0.002996
|
| 22 |
+
10329,2.624868,3.057884,319825.544380,0.077997,1639.293700,0.002995,0.590995,0.002995
|
| 23 |
+
10330,2.695429,3.057884,319867.236745,0.077656,1639.080030,0.002995,0.618318,0.002995
|
| 24 |
+
10331,2.638742,3.057884,319782.458536,0.077614,1639.514570,0.002995,0.645650,0.002995
|
| 25 |
+
10332,2.680027,3.057884,319788.111099,0.080066,1639.485590,0.002994,0.672980,0.002994
|
| 26 |
+
10333,2.669517,3.057884,319659.210973,0.087468,1640.146700,0.002994,0.700322,0.002994
|
| 27 |
+
10334,2.674735,3.057884,319674.298638,0.086068,1640.069290,0.002994,0.727663,0.002994
|
| 28 |
+
10335,2.683662,3.057884,319687.947178,0.085377,1639.999270,0.002993,0.755002,0.002993
|
| 29 |
+
10336,2.588056,3.057884,319702.214870,0.082420,1639.926080,0.002993,0.782341,0.002993
|
| 30 |
+
10337,2.592314,3.057884,319685.405289,0.078556,1640.012310,0.002993,0.809681,0.002993
|
| 31 |
+
10338,2.659404,3.057884,319722.970505,0.092736,1639.819620,0.002992,0.837018,0.002992
|
| 32 |
+
10339,2.554655,3.057884,319656.926799,0.085486,1640.158420,0.002992,0.864361,0.002992
|
| 33 |
+
10340,2.648113,3.057884,319627.440118,0.088174,1640.309730,0.002991,0.891706,0.002991
|
| 34 |
+
10341,2.685987,3.057884,319721.757769,0.089004,1639.825840,0.002991,0.919042,0.002991
|
| 35 |
+
10342,2.685290,3.057884,319669.213385,0.078556,1640.095380,0.002991,0.946384,0.002991
|
| 36 |
+
10343,2.607982,3.057884,319640.492238,0.085381,1640.242750,0.002990,0.973728,0.002990
|
| 37 |
+
10344,2.616746,3.057884,319631.695873,0.108711,1640.287890,0.002990,1.001073,0.002990
|
| 38 |
+
10345,2.622490,3.057884,319681.003860,0.108881,1640.034890,0.002989,1.028413,0.002989
|
| 39 |
+
10346,2.655280,3.057884,319721.030522,0.097260,1639.829570,0.002989,1.055750,0.002989
|
| 40 |
+
10347,2.706025,3.057884,319694.019417,0.090086,1639.968120,0.002988,1.083090,0.002988
|
| 41 |
+
10348,2.582737,3.057884,319620.970964,0.082688,1640.342930,0.002988,1.110435,0.002988
|
| 42 |
+
10349,2.706316,3.057884,319644.491100,0.087002,1640.222230,0.002987,1.137779,0.002987
|
| 43 |
+
10350,2.585762,3.057884,319714.249535,0.088104,1639.864350,0.002987,1.165116,0.002987
|
| 44 |
+
10351,2.609111,3.057884,319646.500313,0.120095,1640.211920,0.002986,1.192459,0.002986
|
| 45 |
+
10352,2.547157,3.057884,319653.208266,0.107101,1640.177500,0.002986,1.219802,0.002986
|
| 46 |
+
10353,2.709230,3.057884,319682.259170,0.098309,1640.028450,0.002985,1.247142,0.002985
|
| 47 |
+
10354,2.544000,3.057884,319639.428230,0.091700,1640.248210,0.002985,1.274486,0.002985
|
| 48 |
+
10355,2.705685,3.057884,319610.174684,0.088208,1640.398340,0.002984,1.301833,0.002984
|
| 49 |
+
10356,2.560872,3.057884,319621.327540,0.093999,1640.341100,0.002983,1.329178,0.002983
|
| 50 |
+
10357,2.638376,3.057884,319630.614387,0.102006,1640.293440,0.002983,1.356523,0.002983
|
| 51 |
+
10358,2.589150,3.057884,319641.821283,0.092981,1640.235930,0.002982,1.383867,0.002982
|
| 52 |
+
10359,2.681022,3.057884,319590.472041,0.094087,1640.499470,0.002982,1.411215,0.002982
|
| 53 |
+
10360,2.679786,3.057884,319579.416790,0.114467,1640.556220,0.002981,1.438564,0.002981
|
| 54 |
+
10361,2.638465,3.057884,319607.172276,0.106257,1640.413750,0.002980,1.465911,0.002980
|
| 55 |
+
10362,2.634351,3.057884,319598.297868,0.096087,1640.459300,0.002980,1.493259,0.002980
|
| 56 |
+
10363,2.571059,3.057884,319635.369084,0.101446,1640.269040,0.002979,1.520603,0.002979
|
| 57 |
+
10364,2.589033,3.057884,319598.290074,0.089744,1640.459340,0.002979,1.547951,0.002979
|
| 58 |
+
10365,2.617486,3.057884,319631.425014,0.091137,1640.289280,0.002978,1.575295,0.002978
|
| 59 |
+
10366,2.661411,3.057884,319602.293726,0.108345,1640.438790,0.002977,1.602643,0.002977
|
| 60 |
+
10367,2.677341,3.057884,319611.776253,0.109394,1640.390120,0.002977,1.629989,0.002977
|
| 61 |
+
10368,2.596525,3.057884,319608.423110,0.095481,1640.407330,0.002976,1.657335,0.002976
|
| 62 |
+
10369,2.566483,3.057884,319630.871604,0.097378,1640.292120,0.002975,1.684680,0.002975
|
| 63 |
+
10370,2.653575,3.057884,319718.375023,0.107972,1639.843190,0.002974,1.712017,0.002974
|
| 64 |
+
10371,2.579601,3.057884,319624.476355,0.103648,1640.324940,0.002974,1.739362,0.002974
|
| 65 |
+
10372,2.605006,3.057884,319593.943645,0.101006,1640.481650,0.002973,1.766710,0.002973
|
| 66 |
+
10373,2.677480,3.057884,319613.576575,0.109509,1640.380880,0.002972,1.794056,0.002972
|
| 67 |
+
10374,2.627289,3.057884,319651.382163,0.112040,1640.186870,0.002971,1.821399,0.002971
|
| 68 |
+
10375,2.654639,3.057884,319594.356658,0.126767,1640.479530,0.002971,1.848747,0.002971
|
| 69 |
+
10376,2.643901,3.057884,319604.908330,0.103787,1640.425370,0.002970,1.876094,0.002970
|
| 70 |
+
10377,2.653200,3.057884,319623.595617,0.103036,1640.329460,0.002969,1.903439,0.002969
|
| 71 |
+
10378,2.674783,3.057884,319568.973935,0.110426,1640.609830,0.002968,1.930789,0.002968
|
| 72 |
+
10379,2.662562,3.057884,319691.087560,0.086865,1639.983160,0.002968,1.958128,0.002968
|
| 73 |
+
10380,2.763343,3.057884,319639.800438,0.076092,1640.246300,0.002967,1.985472,0.002967
|
| 74 |
+
10381,2.548690,3.057884,319660.791591,0.078707,1640.138590,0.002966,2.012814,0.002966
|
| 75 |
+
10382,2.616013,3.057884,319655.710665,0.079765,1640.164660,0.002965,2.040157,0.002965
|
| 76 |
+
10383,2.728987,3.057884,319608.460130,0.071894,1640.407140,0.002964,2.067504,0.002964
|
| 77 |
+
10384,2.623497,3.057884,319543.809476,0.078327,1640.739030,0.002963,2.094856,0.002963
|
| 78 |
+
10385,2.784320,3.057884,319582.654389,0.093729,1640.539600,0.002963,2.122205,0.002963
|
| 79 |
+
10386,2.754317,3.057884,319643.581019,0.100288,1640.226900,0.002962,2.149548,0.002962
|
| 80 |
+
10387,2.556110,3.057884,319645.560988,0.077960,1640.216740,0.002961,2.176892,0.002961
|
| 81 |
+
10388,2.644885,3.057884,319624.544554,0.075712,1640.324590,0.002960,2.204237,0.002960
|
| 82 |
+
10389,2.512394,3.057884,319626.980254,0.080357,1640.312090,0.002959,2.231583,0.002959
|
| 83 |
+
10390,2.831544,3.057884,319677.107390,0.192073,1640.054880,0.002958,2.258923,0.002958
|
| 84 |
+
10391,2.616250,3.057884,319653.508396,0.093133,1640.175960,0.002957,2.286267,0.002957
|
| 85 |
+
10392,2.559199,3.057884,319623.663813,0.078946,1640.329110,0.002956,2.313612,0.002956
|
| 86 |
+
10393,2.716237,3.057884,319537.655303,0.077445,1640.770630,0.002955,2.340965,0.002955
|
| 87 |
+
10394,2.568694,3.057884,319631.333427,0.066505,1640.289750,0.002954,2.368311,0.002954
|
| 88 |
+
10395,2.644906,3.057884,319559.509529,0.073140,1640.658420,0.002953,2.395661,0.002953
|
| 89 |
+
10396,2.734189,3.057884,319635.448982,0.083121,1640.268630,0.002952,2.423006,0.002952
|
| 90 |
+
10397,2.757720,3.057884,319688.598251,0.095128,1639.995930,0.002951,2.450345,0.002951
|
| 91 |
+
10398,2.773420,3.057884,319636.783831,0.091602,1640.261780,0.002950,2.477690,0.002950
|
| 92 |
+
10399,2.608002,3.057884,319619.895395,0.099007,1640.348450,0.002949,2.505035,0.002949
|
| 93 |
+
10400,2.631424,3.057884,319687.974469,0.095844,1639.999130,0.002948,2.532375,0.002948
|
| 94 |
+
10401,2.836423,3.057884,319619.513491,0.093251,1640.350410,0.002947,2.559721,0.002947
|
| 95 |
+
10402,2.649410,3.057884,319634.525311,0.086392,1640.273370,0.002946,2.587065,0.002946
|
| 96 |
+
10403,2.616033,3.057884,319566.994915,0.094652,1640.619990,0.002945,2.614415,0.002945
|
| 97 |
+
10404,2.648589,3.057884,319668.293416,0.112149,1640.100100,0.002944,2.641757,0.002944
|
| 98 |
+
10405,2.707739,3.057884,319613.888320,0.106043,1640.379280,0.002943,2.669103,0.002943
|
| 99 |
+
10406,2.766337,3.057884,319599.421997,0.096687,1640.453530,0.002942,2.696451,0.002942
|
| 100 |
+
10407,2.616031,3.057884,319633.936815,0.105159,1640.276390,0.002941,2.723796,0.002941
|
| 101 |
+
10408,2.641221,3.057884,319665.843445,0.100178,1640.112670,0.002940,2.751137,0.002940
|
| 102 |
+
10409,2.579833,3.057884,319622.420654,0.101381,1640.335490,0.002939,2.778483,0.002939
|
| 103 |
+
10410,2.570470,3.057884,319693.637337,0.100985,1639.970080,0.002938,2.805822,0.002938
|
| 104 |
+
10411,2.580414,3.057884,319600.620168,0.100138,1640.447380,0.002937,2.833169,0.002937
|
| 105 |
+
10412,2.641308,3.057884,319647.969725,0.116407,1640.204380,0.002936,2.860513,0.002936
|
| 106 |
+
10413,2.606977,3.057884,319660.251723,0.126621,1640.141360,0.002934,2.887855,0.002934
|
| 107 |
+
10414,2.456744,3.057884,319618.100848,0.115935,1640.357660,0.002933,2.915200,0.002933
|
| 108 |
+
10415,2.603387,3.057884,319651.119065,0.078222,1640.188220,0.002932,2.942543,0.002932
|
| 109 |
+
10416,2.625212,3.057884,319620.815084,0.077804,1640.343730,0.002931,2.969889,0.002931
|
| 110 |
+
10417,2.694703,3.057884,319580.899222,0.084717,1640.548610,0.002930,2.997238,0.002930
|
| 111 |
+
10418,2.742951,3.057884,319627.492730,0.077662,1640.309460,0.002929,3.024583,0.002929
|
| 112 |
+
10419,2.844461,3.057884,319640.293468,0.211969,1640.243770,0.002927,3.051927,0.002927
|
| 113 |
+
10420,2.689941,3.057884,319667.112281,0.073109,1640.106160,0.002926,3.079268,0.002926
|
| 114 |
+
10421,2.495415,3.057884,319628.151351,0.067772,1640.306080,0.002925,3.106613,0.002925
|
| 115 |
+
10422,2.551864,3.057884,319608.697828,0.074357,1640.405920,0.002924,3.133960,0.002924
|
| 116 |
+
10423,2.738378,3.057884,319573.931333,0.072473,1640.584380,0.002922,3.161310,0.002922
|
| 117 |
+
10424,2.625021,3.057884,319587.374545,0.074038,1640.515370,0.002921,3.188659,0.002921
|
| 118 |
+
10425,2.863070,3.057884,319577.536987,0.115944,1640.565870,0.002920,3.216008,0.002920
|
| 119 |
+
10426,2.620465,3.057884,319627.687587,0.093652,1640.308460,0.002919,3.243353,0.002919
|
| 120 |
+
10427,2.561589,3.057884,319630.690382,0.080830,1640.293050,0.002917,3.270698,0.002917
|
| 121 |
+
10428,2.683123,3.057884,319650.974848,0.079047,1640.188960,0.002916,3.298040,0.002916
|
| 122 |
+
10429,2.621562,3.057884,319594.691745,0.077472,1640.477810,0.002915,3.325388,0.002915
|
| 123 |
+
10430,2.797495,3.057884,319611.090422,0.074964,1640.393640,0.002914,3.352734,0.002914
|
| 124 |
+
10431,2.756585,3.057884,319625.168089,0.070199,1640.321390,0.002912,3.380080,0.002912
|
| 125 |
+
10432,2.767882,3.057884,319570.760138,0.077297,1640.600660,0.002911,3.407431,0.002911
|
| 126 |
+
10433,2.597656,3.057884,319678.487420,0.076636,1640.047800,0.002910,3.434771,0.002910
|
| 127 |
+
10434,2.596557,3.057884,319609.054376,0.074409,1640.404090,0.002908,3.462118,0.002908
|
| 128 |
+
10435,2.518670,3.057884,319629.762841,0.076837,1640.297810,0.002907,3.489463,0.002907
|
| 129 |
+
10436,2.619940,3.057884,319620.937839,0.087342,1640.343100,0.002906,3.516808,0.002906
|
| 130 |
+
10437,2.614896,3.057884,319592.459145,0.104175,1640.489270,0.002904,3.544156,0.002904
|
| 131 |
+
10438,2.579712,3.057884,319627.859062,0.107138,1640.307580,0.002903,3.571501,0.002903
|
| 132 |
+
10439,2.621733,3.057884,319636.663013,0.101048,1640.262400,0.002901,3.598845,0.002901
|
| 133 |
+
10440,2.503057,3.057884,319646.225531,0.085263,1640.213330,0.002900,3.626188,0.002900
|
| 134 |
+
10441,2.580845,3.057884,319627.319305,0.078154,1640.310350,0.002899,3.653533,0.002899
|
| 135 |
+
10442,2.698635,3.057884,319575.758498,0.069221,1640.575000,0.002897,3.680882,0.002897
|
| 136 |
+
10443,2.571500,3.057884,319636.867624,0.074886,1640.261350,0.002896,3.708227,0.002896
|
| 137 |
+
10444,2.716341,3.057884,319638.179102,0.078948,1640.254620,0.002894,3.735571,0.002894
|
| 138 |
+
10445,2.641122,3.057884,319578.472017,0.085113,1640.561070,0.002893,3.762920,0.002893
|
| 139 |
+
10446,2.715120,3.057884,319607.320350,0.093111,1640.412990,0.002891,3.790266,0.002891
|
| 140 |
+
10447,2.604709,3.057884,319638.744229,0.091830,1640.251720,0.002890,3.817610,0.002890
|
| 141 |
+
10448,2.611566,3.057884,319654.297702,0.088400,1640.171910,0.002889,3.844953,0.002889
|
| 142 |
+
10449,2.691082,3.057884,319607.043686,0.103328,1640.414410,0.002887,3.872300,0.002887
|
| 143 |
+
10450,2.696886,3.057884,319611.068989,0.094205,1640.393750,0.002886,3.899647,0.002886
|
| 144 |
+
10451,2.672131,3.057884,319666.001318,0.080755,1640.111860,0.002884,3.926989,0.002884
|
| 145 |
+
10452,2.647804,3.057884,319612.109427,0.080559,1640.388410,0.002883,3.954335,0.002883
|
| 146 |
+
10453,2.719227,3.057884,319645.071840,0.091364,1640.219250,0.002881,3.981679,0.002881
|
| 147 |
+
10454,2.508566,3.057884,319658.228695,0.095662,1640.151740,0.002880,4.009021,0.002880
|
| 148 |
+
10455,2.690474,3.057884,319650.873508,0.103981,1640.189480,0.002878,4.036363,0.002878
|
| 149 |
+
10456,2.594931,3.057884,319644.282579,0.107731,1640.223300,0.002876,4.063707,0.002876
|
| 150 |
+
10457,2.713377,3.057884,319689.202543,0.093809,1639.992830,0.002875,4.091047,0.002875
|
| 151 |
+
10458,2.631460,3.057884,319642.766432,0.086537,1640.231080,0.002873,4.118391,0.002873
|
| 152 |
+
10459,2.604504,3.057884,319600.727322,0.093811,1640.446830,0.002872,4.145738,0.002872
|
| 153 |
+
10460,2.740415,3.057884,319638.385664,0.089191,1640.253560,0.002870,4.173082,0.002870
|
| 154 |
+
10461,2.793469,3.057884,319638.978076,0.107256,1640.250520,0.002869,4.200426,0.002869
|
| 155 |
+
10462,2.697177,3.057884,319537.458607,0.085313,1640.771640,0.002867,4.227779,0.002867
|
| 156 |
+
10463,2.592596,3.057884,319642.569608,0.079506,1640.232090,0.002865,4.255122,0.002865
|
| 157 |
+
10464,2.562007,3.057884,319654.461409,0.088091,1640.171070,0.002864,4.282465,0.002864
|
| 158 |
+
10465,2.698494,3.057884,319624.503635,0.105216,1640.324800,0.002862,4.309811,0.002862
|
| 159 |
+
10466,2.692487,3.057884,319595.698957,0.108205,1640.472640,0.002861,4.337158,0.002861
|
| 160 |
+
10467,2.622339,3.057884,319611.799633,0.103044,1640.390000,0.002859,4.364505,0.002859
|
| 161 |
+
10468,2.606605,3.057884,319644.964657,0.105802,1640.219800,0.002857,4.391848,0.002857
|
| 162 |
+
10469,2.572092,3.057884,319602.496346,0.099942,1640.437750,0.002856,4.419195,0.002856
|
| 163 |
+
10470,2.632645,3.057884,319623.531314,0.090137,1640.329790,0.002854,4.446540,0.002854
|
| 164 |
+
10471,2.471226,3.057884,319610.624758,0.082590,1640.396030,0.002852,4.473887,0.002852
|
| 165 |
+
10472,2.619698,3.057884,319604.132905,0.094196,1640.429350,0.002850,4.501234,0.002850
|
| 166 |
+
10473,2.572552,3.057884,319644.278683,0.091742,1640.223320,0.002849,4.528577,0.002849
|
| 167 |
+
10474,2.707855,3.057884,319655.570343,0.088665,1640.165380,0.002847,4.555920,0.002847
|
| 168 |
+
10475,2.690413,3.057884,319619.661577,0.091093,1640.349650,0.002845,4.583265,0.002845
|
| 169 |
+
10476,2.569057,3.057884,319588.333009,0.091402,1640.510450,0.002844,4.610614,0.002844
|
| 170 |
+
10477,2.542245,3.057884,319650.351213,0.097850,1640.192160,0.002842,4.637957,0.002842
|
| 171 |
+
10478,2.765682,3.057884,319579.742106,0.119198,1640.554550,0.002840,4.665306,0.002840
|
| 172 |
+
10479,2.575944,3.057884,319606.743644,0.105098,1640.415950,0.002838,4.692653,0.002838
|
| 173 |
+
10480,2.578832,3.057884,319619.934364,0.091078,1640.348250,0.002837,4.719998,0.002837
|
| 174 |
+
10481,2.686681,3.057884,319649.289089,0.101555,1640.197610,0.002835,4.747342,0.002835
|
| 175 |
+
10482,2.625229,3.057884,319674.920419,0.101058,1640.066100,0.002833,4.774682,0.002833
|
| 176 |
+
10483,2.705569,3.057884,319615.123620,0.100343,1640.372940,0.002831,4.802028,0.002831
|
| 177 |
+
10484,2.658226,3.057884,319577.663606,0.108719,1640.565220,0.002829,4.829377,0.002829
|
| 178 |
+
10485,2.493976,3.057884,319637.415209,0.103367,1640.258540,0.002828,4.856722,0.002828
|
| 179 |
+
10486,2.596253,3.057884,319582.358289,0.106996,1640.541120,0.002826,4.884070,0.002826
|
| 180 |
+
10487,2.647026,3.057884,319631.935554,0.101760,1640.286660,0.002824,4.911415,0.002824
|
| 181 |
+
10488,2.614134,3.057884,319580.318715,0.092354,1640.551590,0.002822,4.938764,0.002822
|
| 182 |
+
10489,2.583889,3.057884,319599.538891,0.089481,1640.452930,0.002820,4.966112,0.002820
|
| 183 |
+
10490,2.614840,3.057884,319622.496646,0.097606,1640.335100,0.002819,4.993458,0.002819
|
| 184 |
+
10491,2.675368,3.057884,319612.810850,0.110222,1640.384810,0.002817,5.020804,0.002817
|
| 185 |
+
10492,2.545220,3.057884,319650.583127,0.108203,1640.190970,0.002815,5.048147,0.002815
|
| 186 |
+
10493,2.488529,3.057884,319627.558982,0.104375,1640.309120,0.002813,5.075492,0.002813
|
| 187 |
+
10494,2.621747,3.057884,319635.828974,0.091405,1640.266680,0.002811,5.102836,0.002811
|
| 188 |
+
10495,2.579574,3.057884,319643.072388,0.081315,1640.229510,0.002809,5.130180,0.002809
|
| 189 |
+
10496,2.654875,3.057884,319584.288798,0.099719,1640.531210,0.002807,5.157528,0.002807
|
| 190 |
+
10497,2.692438,3.057884,319601.226075,0.106274,1640.444270,0.002805,5.184876,0.002805
|
| 191 |
+
10498,2.574060,3.057884,319667.199988,0.093198,1640.105710,0.002803,5.212217,0.002803
|
| 192 |
+
10499,2.617799,3.057884,319605.342803,0.087670,1640.423140,0.002801,5.239564,0.002801
|
| 193 |
+
10500,2.553835,3.051036,319630.963189,0.085366,1640.291650,0.002800,5.284262,0.002800
|
| 194 |
+
10501,2.591541,3.051036,319871.241282,0.086118,1639.059510,0.002798,5.311586,0.002798
|
| 195 |
+
10502,2.603822,3.051036,319906.146674,0.086761,1638.880670,0.002796,5.338908,0.002796
|
| 196 |
+
10503,2.656048,3.051036,319874.078866,0.091617,1639.044970,0.002794,5.366231,0.002794
|
| 197 |
+
10504,2.644125,3.051036,319808.150244,0.091598,1639.382860,0.002792,5.393561,0.002792
|
| 198 |
+
10505,2.534880,3.051036,319901.231666,0.090393,1638.905850,0.002790,5.420883,0.002790
|
| 199 |
+
10506,2.605783,3.051036,319836.852753,0.091925,1639.235740,0.002788,5.448210,0.002788
|
| 200 |
+
10507,2.547550,3.051036,319867.751943,0.098600,1639.077390,0.002786,5.475535,0.002786
|
| 201 |
+
10508,2.720389,3.051036,319909.076623,0.093485,1638.865660,0.002784,5.502856,0.002784
|
| 202 |
+
10509,2.639339,3.051036,319870.973920,0.091895,1639.060880,0.002782,5.530180,0.002782
|
| 203 |
+
10510,2.661986,3.051036,319937.992486,0.089784,1638.717540,0.002780,5.557499,0.002780
|
| 204 |
+
10511,2.803216,3.051036,319920.578310,0.090028,1638.806740,0.002778,5.584819,0.002778
|
| 205 |
+
10512,2.688303,3.051036,319907.854667,0.082248,1638.871920,0.002776,5.612140,0.002776
|
| 206 |
+
10513,2.658484,3.051036,319958.120716,0.074884,1638.614450,0.002774,5.639457,0.002774
|
| 207 |
+
10514,2.588421,3.051036,319944.648248,0.070931,1638.683450,0.002771,5.666775,0.002771
|
| 208 |
+
10515,2.576880,3.051036,319787.204101,0.068283,1639.490240,0.002769,5.694106,0.002769
|
| 209 |
+
10516,2.532889,3.051036,319845.071139,0.070746,1639.193620,0.002767,5.721433,0.002767
|
| 210 |
+
10517,2.636373,3.051036,319900.938879,0.065863,1638.907350,0.002765,5.748754,0.002765
|
| 211 |
+
10518,2.630598,3.051036,319871.411069,0.069664,1639.058640,0.002763,5.776078,0.002763
|
| 212 |
+
10519,2.666814,3.051036,319926.960035,0.070672,1638.774050,0.002761,5.803398,0.002761
|
| 213 |
+
10520,2.672949,3.051036,319839.679969,0.103110,1639.221250,0.002759,5.830725,0.002759
|
| 214 |
+
10521,2.664687,3.051036,319914.698526,0.097227,1638.836860,0.002757,5.858045,0.002757
|
| 215 |
+
10522,2.616051,3.051036,319914.079717,0.081522,1638.840030,0.002755,5.885365,0.002755
|
| 216 |
+
10523,2.680714,3.051036,319937.447777,0.083714,1638.720330,0.002753,5.912684,0.002753
|
| 217 |
+
10524,2.623666,3.051036,319858.482534,0.074971,1639.124890,0.002750,5.940009,0.002750
|
| 218 |
+
10525,2.672380,3.051036,319842.239923,0.074018,1639.208130,0.002748,5.967335,0.002748
|
| 219 |
+
10526,2.655396,3.051036,319952.165355,0.076316,1638.644950,0.002746,5.994653,0.002746
|
| 220 |
+
10527,2.548869,3.051036,319883.185236,0.065984,1638.998310,0.002744,6.021976,0.002744
|
| 221 |
+
10528,2.622389,3.051036,319904.192756,0.069845,1638.890680,0.002742,6.049297,0.002742
|
| 222 |
+
10529,2.543951,3.051036,319888.616907,0.075224,1638.970480,0.002740,6.076619,0.002740
|
| 223 |
+
10530,2.603323,3.051036,319905.147264,0.069526,1638.885790,0.002737,6.103940,0.002737
|
| 224 |
+
10531,2.602228,3.051036,319849.322935,0.065717,1639.171830,0.002735,6.131267,0.002735
|
| 225 |
+
10532,2.511177,3.051036,319887.820590,0.070995,1638.974560,0.002733,6.158589,0.002733
|
| 226 |
+
10533,2.572714,3.051036,319889.963632,0.071372,1638.963580,0.002731,6.185912,0.002731
|
| 227 |
+
10534,2.578027,3.051036,319866.934261,0.069808,1639.081580,0.002729,6.213236,0.002729
|
| 228 |
+
10535,2.542672,3.051036,319887.865479,0.069784,1638.974330,0.002726,6.240559,0.002726
|
| 229 |
+
10536,2.642068,3.051036,319859.229922,0.072882,1639.121060,0.002724,6.267884,0.002724
|
| 230 |
+
10537,2.623832,3.051036,319930.247630,0.073846,1638.757210,0.002722,6.295203,0.002722
|
| 231 |
+
10538,2.641641,3.051036,319902.734652,0.077861,1638.898150,0.002720,6.322525,0.002720
|
| 232 |
+
10539,2.536454,3.051036,319882.968598,0.077846,1638.999420,0.002717,6.349848,0.002717
|
| 233 |
+
10540,2.689143,3.051036,319918.895560,0.080883,1638.815360,0.002715,6.377169,0.002715
|
| 234 |
+
10541,2.682338,3.051036,319865.597490,0.089776,1639.088430,0.002713,6.404494,0.002713
|
| 235 |
+
10542,2.559465,3.051036,319876.645226,0.086228,1639.031820,0.002710,6.431817,0.002710
|
| 236 |
+
10543,2.590423,3.051036,319894.064379,0.074885,1638.942570,0.002708,6.459139,0.002708
|
| 237 |
+
10544,2.609137,3.051036,319926.563732,0.075140,1638.776080,0.002706,6.486459,0.002706
|
| 238 |
+
10545,2.634950,3.051036,319873.518762,0.070040,1639.047840,0.002703,6.513783,0.002703
|
| 239 |
+
10546,2.530984,3.051036,319940.895684,0.072124,1638.702670,0.002701,6.541101,0.002701
|
| 240 |
+
10547,2.675320,3.051036,319913.029504,0.074220,1638.845410,0.002699,6.568421,0.002699
|
| 241 |
+
10548,2.523057,3.051036,319861.961915,0.070509,1639.107060,0.002696,6.595746,0.002696
|
| 242 |
+
10549,2.561991,3.051036,319878.637841,0.066621,1639.021610,0.002694,6.623070,0.002694
|
| 243 |
+
10550,2.596606,3.051036,319919.777929,0.070872,1638.810840,0.002692,6.650390,0.002692
|
| 244 |
+
10551,2.530825,3.051036,319887.475128,0.071982,1638.976330,0.002689,6.677712,0.002689
|
| 245 |
+
10552,2.565084,3.051036,319914.995242,0.080476,1638.835340,0.002687,6.705033,0.002687
|
| 246 |
+
10553,2.577397,3.051036,319917.812131,0.080556,1638.820910,0.002685,6.732353,0.002685
|
| 247 |
+
10554,2.597488,3.051036,319952.313749,0.079997,1638.644190,0.002682,6.759670,0.002682
|
| 248 |
+
10555,2.760314,3.051036,319901.339021,0.073902,1638.905300,0.002680,6.786992,0.002680
|
| 249 |
+
10556,2.546053,3.051036,319911.793851,0.070460,1638.851740,0.002678,6.814312,0.002678
|
| 250 |
+
10557,2.691957,3.051036,319878.762748,0.071473,1639.020970,0.002675,6.841637,0.002675
|
| 251 |
+
10558,2.615556,3.051036,319908.085002,0.072449,1638.870740,0.002673,6.868958,0.002673
|
| 252 |
+
10559,2.512166,3.051036,319942.496663,0.067485,1638.694470,0.002670,6.896276,0.002670
|
| 253 |
+
10560,2.664476,3.051036,319919.258658,0.075668,1638.813500,0.002668,6.923596,0.002668
|
| 254 |
+
10561,2.529170,3.051036,319909.919896,0.077503,1638.861340,0.002665,6.950917,0.002665
|
| 255 |
+
10562,2.605489,3.051036,319883.298435,0.072904,1638.997730,0.002663,6.978241,0.002663
|
| 256 |
+
10563,2.479856,3.051036,319893.919943,0.075315,1638.943310,0.002660,7.005563,0.002660
|
| 257 |
+
10564,2.647278,3.051036,319965.052644,0.082513,1638.578950,0.002658,7.032879,0.002658
|
| 258 |
+
10565,2.815350,3.051036,319930.884072,0.339185,1638.753950,0.002656,7.060198,0.002656
|
| 259 |
+
10566,2.538144,3.051036,319846.721891,0.076109,1639.185160,0.002653,7.087524,0.002653
|
| 260 |
+
10567,2.646501,3.051036,319916.016188,0.068290,1638.830110,0.002651,7.114845,0.002651
|
| 261 |
+
10568,2.724876,3.051036,319890.250544,0.071005,1638.962110,0.002648,7.142167,0.002648
|
| 262 |
+
10569,2.591948,3.051036,319878.319724,0.067990,1639.023240,0.002646,7.169491,0.002646
|
| 263 |
+
10570,2.525646,3.051036,319902.555073,0.064275,1638.899070,0.002643,7.196812,0.002643
|
| 264 |
+
10571,2.493958,3.051036,319926.817522,0.071735,1638.774780,0.002641,7.224132,0.002641
|
| 265 |
+
10572,2.664876,3.051036,319940.112772,0.067992,1638.706680,0.002638,7.251450,0.002638
|
| 266 |
+
10573,2.688772,3.051036,319959.958135,0.069894,1638.605040,0.002636,7.278767,0.002636
|
| 267 |
+
10574,2.647939,3.051036,319906.127156,0.080451,1638.880770,0.002633,7.306088,0.002633
|
| 268 |
+
10575,2.795584,3.051036,319946.788147,0.089154,1638.672490,0.002630,7.333406,0.002630
|
| 269 |
+
10576,2.657212,3.051036,319947.565234,0.093406,1638.668510,0.002628,7.360723,0.002628
|
| 270 |
+
10577,2.699704,3.051036,319928.637013,0.084709,1638.765460,0.002625,7.388042,0.002625
|
| 271 |
+
10578,2.545095,3.051036,319905.959285,0.067302,1638.881630,0.002623,7.415364,0.002623
|
| 272 |
+
10579,2.654643,3.051036,319868.755023,0.070425,1639.072250,0.002620,7.442688,0.002620
|
| 273 |
+
10580,2.620667,3.051036,319853.561185,0.078753,1639.150110,0.002618,7.470014,0.002618
|
| 274 |
+
10581,2.590570,3.051036,319912.211590,0.076353,1638.849600,0.002615,7.497335,0.002615
|
| 275 |
+
10582,2.625221,3.051036,319877.213148,0.069685,1639.028910,0.002612,7.524670,0.002612
|
| 276 |
+
10583,2.624538,3.051036,319885.384816,0.070212,1638.987040,0.002610,7.551993,0.002610
|
| 277 |
+
10584,2.696926,3.051036,319922.573427,0.077568,1638.796520,0.002607,7.579313,0.002607
|
| 278 |
+
10585,2.637998,3.051036,319951.495634,0.075602,1638.648380,0.002605,7.606630,0.002605
|
| 279 |
+
10586,2.625880,3.051036,319892.559523,0.072780,1638.950280,0.002602,7.633953,0.002602
|
| 280 |
+
10587,2.543417,3.051036,319930.997304,0.077396,1638.753370,0.002599,7.661272,0.002599
|
| 281 |
+
10588,2.634770,3.051036,319865.931195,0.083200,1639.086720,0.002597,7.688596,0.002597
|
| 282 |
+
10589,2.537612,3.051036,319931.973448,0.086491,1638.748370,0.002594,7.715916,0.002594
|
| 283 |
+
10590,2.513633,3.051036,319879.916174,0.088868,1639.015060,0.002591,7.743240,0.002591
|
| 284 |
+
10591,2.589366,3.051036,319882.705120,0.090307,1639.000770,0.002589,7.770563,0.002589
|
| 285 |
+
10592,2.695950,3.051036,319905.047714,0.090116,1638.886300,0.002586,7.797884,0.002586
|
| 286 |
+
10593,2.564269,3.051036,319994.318460,0.087173,1638.429090,0.002583,7.825198,0.002583
|
| 287 |
+
10594,2.665426,3.051036,319902.598016,0.072874,1638.898850,0.002581,7.852520,0.002581
|
| 288 |
+
10595,2.831981,3.051036,319921.505588,0.074998,1638.801990,0.002578,7.879840,0.002578
|
| 289 |
+
10596,2.514539,3.051036,319895.612190,0.077179,1638.934640,0.002575,7.907162,0.002575
|
| 290 |
+
10597,2.710741,3.051036,319842.196996,0.080509,1639.208350,0.002573,7.934488,0.002573
|
| 291 |
+
10598,2.496753,3.051036,319859.019169,0.077677,1639.122140,0.002570,7.961814,0.002570
|
| 292 |
+
10599,2.580015,3.051036,319896.268013,0.076440,1638.931280,0.002567,7.989136,0.002567
|
| 293 |
+
10600,2.671192,3.051036,319979.979768,0.075587,1638.502510,0.002564,8.016451,0.002564
|
| 294 |
+
10601,2.679725,3.051036,319858.843544,0.081872,1639.123040,0.002562,8.043776,0.002562
|
| 295 |
+
10602,2.659037,3.051036,319902.166638,0.080890,1638.901060,0.002559,8.071098,0.002559
|
| 296 |
+
10603,2.641108,3.051036,319917.566164,0.072960,1638.822170,0.002556,8.098418,0.002556
|
| 297 |
+
10604,2.565979,3.051036,319955.225016,0.071166,1638.629280,0.002554,8.125735,0.002554
|
| 298 |
+
10605,2.522801,3.051036,319917.259680,0.070900,1638.823740,0.002551,8.153056,0.002551
|
| 299 |
+
10606,2.649387,3.051036,319897.413757,0.073748,1638.925410,0.002548,8.180378,0.002548
|
| 300 |
+
10607,2.512681,3.051036,319832.792494,0.077130,1639.256550,0.002545,8.207705,0.002545
|
| 301 |
+
10608,2.555712,3.051036,319773.065346,0.087002,1639.562730,0.002542,8.235040,0.002542
|
| 302 |
+
10609,2.616781,3.051036,319949.033509,0.102688,1638.660990,0.002540,8.262357,0.002540
|
| 303 |
+
10610,2.617699,3.051036,319860.712997,0.105337,1639.113460,0.002537,8.289682,0.002537
|
| 304 |
+
10611,2.610990,3.051036,319924.480717,0.086641,1638.786750,0.002534,8.317002,0.002534
|
| 305 |
+
10612,2.633626,3.051036,319880.372863,0.082724,1639.012720,0.002531,8.344325,0.002531
|
| 306 |
+
10613,2.585816,3.051036,319935.231863,0.086433,1638.731680,0.002528,8.371644,0.002528
|
| 307 |
+
10614,2.740581,3.051036,319885.736127,0.074876,1638.985240,0.002526,8.398967,0.002526
|
| 308 |
+
10615,2.657877,3.051036,319863.640162,0.069786,1639.098460,0.002523,8.426292,0.002523
|
| 309 |
+
10616,2.518871,3.051036,319810.331236,0.074594,1639.371680,0.002520,8.453621,0.002520
|
| 310 |
+
10617,2.637612,3.051036,319851.637177,0.075239,1639.159970,0.002517,8.480947,0.002517
|
| 311 |
+
10618,2.510497,3.051036,319875.222502,0.070262,1639.039110,0.002514,8.508272,0.002514
|
| 312 |
+
10619,2.660985,3.051036,319864.656877,0.067809,1639.093250,0.002511,8.535597,0.002511
|
| 313 |
+
10620,2.566689,3.051036,319933.332251,0.074619,1638.741410,0.002509,8.562915,0.002509
|
| 314 |
+
10621,2.616558,3.051036,319946.296124,0.075489,1638.675010,0.002506,8.590234,0.002506
|
| 315 |
+
10622,2.577765,3.051036,319908.672557,0.083309,1638.867730,0.002503,8.617555,0.002503
|
| 316 |
+
10623,2.569064,3.051036,319906.552688,0.088770,1638.878590,0.002500,8.644876,0.002500
|
| 317 |
+
10624,2.644313,3.051036,319954.742730,0.079174,1638.631750,0.002497,8.672193,0.002497
|
| 318 |
+
10625,2.600619,3.051036,319903.652065,0.071480,1638.893450,0.002494,8.699514,0.002494
|
| 319 |
+
10626,2.528054,3.051036,319873.975433,0.067265,1639.045500,0.002491,8.726838,0.002491
|
| 320 |
+
10627,2.437984,3.051036,319861.850683,0.063434,1639.107630,0.002488,8.754163,0.002488
|
| 321 |
+
10628,2.533826,3.051036,319901.504935,0.066275,1638.904450,0.002485,8.781484,0.002485
|
| 322 |
+
10629,2.619308,3.051036,319866.645440,0.063366,1639.083060,0.002483,8.808809,0.002483
|
| 323 |
+
10630,2.612548,3.051036,319902.262282,0.061399,1638.900570,0.002480,8.836131,0.002480
|
| 324 |
+
10631,2.688943,3.051036,319917.408040,0.064647,1638.822980,0.002477,8.863451,0.002477
|
| 325 |
+
10632,2.664616,3.051036,319912.521967,0.066990,1638.848010,0.002474,8.890772,0.002474
|
| 326 |
+
10633,2.594924,3.051036,319893.080658,0.072295,1638.947610,0.002471,8.918095,0.002471
|
| 327 |
+
10634,2.611122,3.051036,319874.718990,0.071003,1639.041690,0.002468,8.945419,0.002468
|
| 328 |
+
10635,2.570190,3.051036,319931.516612,0.067295,1638.750710,0.002465,8.972738,0.002465
|
| 329 |
+
10636,2.671571,3.051036,319630.390294,0.063783,1640.294590,0.002462,9.000083,0.002462
|
| 330 |
+
10637,2.517806,3.051036,319808.271194,0.067984,1639.382240,0.002459,9.027413,0.002459
|
| 331 |
+
10638,2.718278,3.051036,319768.562037,0.074768,1639.585820,0.002456,9.054746,0.002456
|
| 332 |
+
10639,2.578080,3.051036,319795.406283,0.070484,1639.448190,0.002453,9.082079,0.002453
|
| 333 |
+
10640,2.564280,3.051036,319810.826743,0.063558,1639.369140,0.002450,9.109409,0.002450
|
| 334 |
+
10641,2.535107,3.051036,319922.567569,0.078593,1638.796550,0.002447,9.136729,0.002447
|
| 335 |
+
10642,2.620808,3.051036,319849.067317,0.078063,1639.173140,0.002444,9.164055,0.002444
|
| 336 |
+
10643,2.642947,3.051036,319856.894105,0.068785,1639.133030,0.002441,9.191381,0.002441
|
| 337 |
+
10644,2.641507,3.051036,319888.954563,0.070789,1638.968750,0.002438,9.218704,0.002438
|
| 338 |
+
10645,2.421164,3.051036,319978.624475,0.075252,1638.509450,0.002435,9.246019,0.002435
|
| 339 |
+
10646,2.469529,3.051036,319863.685046,0.068891,1639.098230,0.002432,9.273344,0.002432
|
| 340 |
+
10647,2.558671,3.051036,319909.299152,0.076300,1638.864520,0.002429,9.300665,0.002429
|
| 341 |
+
10648,2.653425,3.051036,319838.476103,0.073404,1639.227420,0.002426,9.327992,0.002426
|
| 342 |
+
10649,2.666234,3.051036,319836.304484,0.065529,1639.238550,0.002423,9.355319,0.002423
|
| 343 |
+
10650,2.587936,3.051036,319775.314124,0.073040,1639.551200,0.002420,9.382651,0.002420
|
| 344 |
+
10651,2.774417,3.051036,319917.593493,0.095976,1638.822030,0.002417,9.409972,0.002417
|
| 345 |
+
10652,2.539855,3.051036,319776.774966,0.061902,1639.543710,0.002414,9.437305,0.002414
|
| 346 |
+
10653,2.685279,3.051036,319834.452876,0.063734,1639.248040,0.002411,9.464632,0.002411
|
| 347 |
+
10654,2.726453,3.051036,319805.436728,0.067314,1639.396770,0.002408,9.491962,0.002408
|
| 348 |
+
10655,2.661300,3.051036,319869.014576,0.073526,1639.070920,0.002405,9.519287,0.002405
|
| 349 |
+
10656,2.544700,3.051036,319911.561559,0.070344,1638.852930,0.002402,9.546607,0.002402
|
| 350 |
+
10657,2.502671,3.051036,319895.118371,0.072678,1638.937170,0.002398,9.573930,0.002398
|
| 351 |
+
10658,2.537918,3.051036,319849.970764,0.071267,1639.168510,0.002395,9.601255,0.002395
|
| 352 |
+
10659,2.595579,3.051036,319904.817382,0.068044,1638.887480,0.002392,9.628576,0.002392
|
| 353 |
+
10660,2.550340,3.051036,319914.538455,0.066611,1638.837680,0.002389,9.655897,0.002389
|
| 354 |
+
10661,2.609774,3.051036,319908.233355,0.066748,1638.869980,0.002386,9.683218,0.002386
|
| 355 |
+
10662,2.601235,3.051036,319858.189825,0.063435,1639.126390,0.002383,9.710543,0.002383
|
| 356 |
+
10663,2.647851,3.051036,319884.924210,0.066107,1638.989400,0.002380,9.737866,0.002380
|
| 357 |
+
10664,2.568462,3.051036,319859.659231,0.062815,1639.118860,0.002377,9.765192,0.002377
|
| 358 |
+
10665,2.614563,3.051036,319962.272027,0.063989,1638.593190,0.002374,9.792508,0.002374
|
| 359 |
+
10666,2.678935,3.051036,319922.524622,0.071495,1638.796770,0.002370,9.819828,0.002370
|
| 360 |
+
10667,2.500076,3.051036,319974.062649,0.066418,1638.532810,0.002367,9.847143,0.002367
|
| 361 |
+
10668,2.561331,3.051036,319864.732984,0.065154,1639.092860,0.002364,9.874468,0.002364
|
| 362 |
+
10669,2.660706,3.051036,319922.389921,0.069331,1638.797460,0.002361,9.901788,0.002361
|
| 363 |
+
10670,2.620308,3.051036,319898.737132,0.066663,1638.918630,0.002358,9.929110,0.002358
|
| 364 |
+
10671,2.481086,3.051036,319933.217065,0.065410,1638.742000,0.002355,9.956428,0.002355
|
| 365 |
+
10672,2.541093,3.051036,319862.168767,0.070190,1639.106000,0.002352,9.983753,0.002352
|
| 366 |
+
10673,2.683227,3.051036,319842.083828,0.072099,1639.208930,0.002348,10.011081,0.002348
|
| 367 |
+
10674,2.566789,3.051036,319925.277218,0.068605,1638.782670,0.002345,10.038400,0.002345
|
| 368 |
+
10675,2.580904,3.051036,319937.801155,0.072783,1638.718520,0.002342,10.065719,0.002342
|
| 369 |
+
10676,2.597879,3.051036,319881.016913,0.074251,1639.009420,0.002339,10.093042,0.002339
|
| 370 |
+
10677,2.633514,3.051036,319858.494243,0.072311,1639.124830,0.002336,10.120367,0.002336
|
| 371 |
+
10678,2.640936,3.051036,319843.508209,0.079450,1639.201630,0.002332,10.147694,0.002332
|
| 372 |
+
10679,2.654224,3.051036,319876.522273,0.079399,1639.032450,0.002329,10.175018,0.002329
|
| 373 |
+
10680,2.556746,3.051036,319905.164833,0.081221,1638.885700,0.002326,10.202339,0.002326
|
| 374 |
+
10681,2.599396,3.051036,319826.303318,0.086565,1639.289810,0.002323,10.229667,0.002323
|
| 375 |
+
10682,2.592875,3.051036,319887.449756,0.094078,1638.976460,0.002320,10.256990,0.002320
|
| 376 |
+
10683,2.569088,3.051036,319857.268771,0.094845,1639.131110,0.002316,10.284316,0.002316
|
| 377 |
+
10684,2.662574,3.051036,319904.946212,0.086158,1638.886820,0.002313,10.311638,0.002313
|
| 378 |
+
10685,2.543761,3.051036,319893.170442,0.080795,1638.947150,0.002310,10.338960,0.002310
|
| 379 |
+
10686,2.545193,3.051036,319846.181396,0.078171,1639.187930,0.002307,10.366286,0.002307
|
| 380 |
+
10687,2.580250,3.051036,319824.252825,0.076685,1639.300320,0.002303,10.393615,0.002303
|
| 381 |
+
10688,2.549900,3.051036,319939.035054,0.074892,1638.712200,0.002300,10.420934,0.002300
|
| 382 |
+
10689,2.707404,3.051036,319883.446765,0.071871,1638.996970,0.002297,10.448257,0.002297
|
| 383 |
+
10690,2.600739,3.051036,319844.883821,0.070457,1639.194580,0.002294,10.475584,0.002294
|
| 384 |
+
10691,2.674275,3.051036,319963.730674,0.066224,1638.585720,0.002290,10.502900,0.002290
|
| 385 |
+
10692,2.572013,3.051036,319887.791312,0.064391,1638.974710,0.002287,10.530223,0.002287
|
| 386 |
+
10693,2.521792,3.051036,319880.415800,0.069601,1639.012500,0.002284,10.557546,0.002284
|
| 387 |
+
10694,2.500074,3.051036,319838.401960,0.073791,1639.227800,0.002281,10.584874,0.002281
|
| 388 |
+
10695,2.579055,3.051036,319931.516612,0.068063,1638.750710,0.002277,10.612194,0.002277
|
| 389 |
+
10696,2.675575,3.051036,319952.065775,0.069365,1638.645460,0.002274,10.639511,0.002274
|
| 390 |
+
10697,2.736763,3.051036,319867.935386,0.064741,1639.076450,0.002271,10.666835,0.002271
|
| 391 |
+
10698,2.598721,3.051036,319905.207777,0.061174,1638.885480,0.002268,10.694157,0.002268
|
| 392 |
+
10699,2.507962,3.051036,319882.685602,0.057756,1639.000870,0.002264,10.721480,0.002264
|
| 393 |
+
10700,2.495724,3.051036,319934.460692,0.080465,1638.735630,0.002261,10.748799,0.002261
|
| 394 |
+
10701,2.673309,3.051036,319869.771770,0.066389,1639.067040,0.002258,10.776123,0.002258
|
| 395 |
+
10702,2.641712,3.051036,319810.916480,0.062941,1639.368680,0.002254,10.803452,0.002254
|
| 396 |
+
10703,2.574183,3.051036,319910.970089,0.063474,1638.855960,0.002251,10.830774,0.002251
|
| 397 |
+
10704,2.662001,3.051036,319849.806855,0.069015,1639.169350,0.002248,10.858100,0.002248
|
| 398 |
+
10705,2.520388,3.051036,319911.481525,0.062682,1638.853340,0.002244,10.885421,0.002244
|
150M_228M_3e3_6B_incomp/blumodel_muon_step00010706.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:086c1617701ba929341a9bd40e6ec9b9d56ac0f3c536be7e03d54d0753a2f855
|
| 3 |
+
size 2515184199
|
150M_228M_75e5_10B_half.csv
ADDED
|
The diff for this file is too large to render.
See raw diff
|
|
|
150M_228M_75e5_10B_half/blumodel_muon_step00000953.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:19c6294243727cd2bc282c9a530ca3c98521a85b9f2c0df4dbe82e8e048a4dcf
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00001906.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:baf6f9d28706c012fdaddc5096e6312240200a176d19d375d968e1786b4a6284
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00002859.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:f968ea83ae6c390c8d6b2668c69ba1dc8646cf7bcc5a9312d801b770f9bcea4c
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00003812.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:a3a0533bba2b0c0f8a38759e7df6745ae41f4fea75f20e544be23784a51201c8
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00004765.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:35c731f2894426e91642b9705ae6e0deb7e0caa7c8c70ff6d065890ec6eb0f7b
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00005718.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:f91dbc49e8ffe592cf3308d2364e16e653a50346c9ec649942bb45fd66d23d61
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00006671.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:2cd83878afac59c02cb76d00b723ec319b95412ad7c2bf5b0cd59aeb58d10c96
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00007624.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:482c93f862bb26681afd5af4d853a36096cc91a22c88c871125069a50c9cc4da
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00008577.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:9c050edb018b294e3a6e84149d3bed6dc8e68fe4444a851c86bd5b41c9703445
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00009530.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:19bd8cdf882c01d35d8d028aa8213bef4b52b85d5421042c76a68fae5ee5919d
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00010483.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:19116842c7aeb68d9ffe03e812646c1095c505d80af2aeb6cf706e6a4b6a44a6
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00011436.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:2427673fceb45df25926a5ddde0d3e763fa235c93160ed0f555edc564be9b89c
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00012389.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:a8e561bf3cdbf3d2674596f98303f90b6715943bb3bcb3d7937187ca98605a83
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00013342.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:f7290fe834c7b667b868eedc829ae77456e0f08bcb92ab008856976548125dd8
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00014295.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:5cb1589be099c8c7ad095906fa3813394c56e992a281c8e65d1d469217153dec
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00015248.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:b336ede3ed3aae5516449c5f5c3f3cb9b15df360210192513f04ed3693476993
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00016201.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:3edf7f3aed504d4f6220f11b736689cbb201f2cf1f757457bba9b560258bf8fc
|
| 3 |
+
size 2515183879
|
150M_228M_75e5_10B_half/blumodel_muon_step00016332.pt
ADDED
|
@@ -0,0 +1,3 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:9ab005c8850f5e064ba08efd426345f3b64d5678bb01f086d6ae9000f4975615
|
| 3 |
+
size 2515183879
|
train_muon.py
ADDED
|
@@ -0,0 +1,1829 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 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 = 768
|
| 438 |
+
n_layers: int = 24
|
| 439 |
+
q_heads: int = 12
|
| 440 |
+
kv_heads: int = 4 # grouped-query: G = q/kv = 4
|
| 441 |
+
head_dim: int = 64 # the fused kernels need 64 or 128
|
| 442 |
+
ffn_hidden: int = 2048 # ~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.003
|
| 462 |
+
min_lr: float = 0.0003 # 10% of max_lr
|
| 463 |
+
|
| 464 |
+
muon_lr: float = 0.003
|
| 465 |
+
muon_min_lr: float = 0.0003
|
| 466 |
+
|
| 467 |
+
max_steps: int = 11444
|
| 468 |
+
warmup_steps: int = 500
|
| 469 |
+
stable_steps: int = int(max_steps * 0.9) - warmup_steps # steps held at max_lr after warmup
|
| 470 |
+
|
| 471 |
+
|
| 472 |
+
# ---- Muon (hidden 2D weights); the AdamW fields above still drive the
|
| 473 |
+
# embedding / lm_head / 1D group, unchanged from the AdamW arm ----
|
| 474 |
+
# Default == max_lr: under adjust_lr_fn=match_rms_adamw the two are the same unit,
|
| 475 |
+
# so out of the box both arms take the same size step and differ only in
|
| 476 |
+
# direction. Same 10% floor, same WSD shape, so the schedules stay aligned.
|
| 477 |
+
|
| 478 |
+
# The rest are torch.optim.Muon's own arguments, passed through unchanged.
|
| 479 |
+
# momentum/nesterov/ns_steps are torch's defaults; adjust_lr_fn is not (see
|
| 480 |
+
# the header — torch defaults to "original").
|
| 481 |
+
muon_momentum: float = 0.95 # beta on the momentum buffer
|
| 482 |
+
muon_ns_steps: int = 5 # Newton-Schulz iterations
|
| 483 |
+
muon_adjust_lr: str = "match_rms_adamw" # "match_rms_adamw" | "original"
|
| 484 |
+
|
| 485 |
+
val_freq: int = 250
|
| 486 |
+
val_loss_steps: int = 20
|
| 487 |
+
|
| 488 |
+
# ---- checkpointing ----
|
| 489 |
+
checkpointing: bool = True
|
| 490 |
+
ckpt_freq: int = int(max_steps / 20) # save every 1800 optimizer steps
|
| 491 |
+
ckpt_keep: int = 40 # retain the most-recent N checkpoints
|
| 492 |
+
# Separate from the AdamW arm's dir: the two checkpoint formats differ (Muon
|
| 493 |
+
# momentum vs Adam m/v) and auto-resume picks the newest file it finds.
|
| 494 |
+
ckpt_dir: str = "ckpt_blu_muon"
|
| 495 |
+
ckpt_prefix: str = "blumodel_muon"
|
| 496 |
+
|
| 497 |
+
# ---- runtime knobs (not in the C++ struct) ----
|
| 498 |
+
fused_rms_rope: bool = True
|
| 499 |
+
activation_checkpoint: bool = False
|
| 500 |
+
loss_chunks: int = 1
|
| 501 |
+
# Depth probe: which layers --diag reports gradient / weight statistics for.
|
| 502 |
+
probe_layers: tuple = (0, 8, 17, 26, 35)
|
| 503 |
+
|
| 504 |
+
|
| 505 |
+
def get_lr(step, cfg, max_lr=None, min_lr=None):
|
| 506 |
+
"""WSD: linear Warmup -> Stable (constant max_lr) -> cosine Decay to min_lr.
|
| 507 |
+
Byte-for-byte the C++ get_lr().
|
| 508 |
+
`max_lr`/`min_lr` override the config pair so the Muon group can ride the
|
| 509 |
+
SAME schedule shape — same warmup length, same stable window, same cosine
|
| 510 |
+
tail — at its own amplitude. Called with no override it is identical to the
|
| 511 |
+
AdamW arm's get_lr(), which is what the embedding/norm group uses.
|
| 512 |
+
"""
|
| 513 |
+
max_lr = cfg.max_lr if max_lr is None else max_lr
|
| 514 |
+
min_lr = cfg.min_lr if min_lr is None else min_lr
|
| 515 |
+
if step < cfg.warmup_steps:
|
| 516 |
+
return max_lr * (step + 1) / cfg.warmup_steps
|
| 517 |
+
if step < cfg.stable_steps + cfg.warmup_steps:
|
| 518 |
+
return max_lr
|
| 519 |
+
tail = cfg.max_steps - (cfg.stable_steps + cfg.warmup_steps) - 1
|
| 520 |
+
decay_ratio = (step - (cfg.stable_steps + cfg.warmup_steps)) / max(tail, 1)
|
| 521 |
+
decay_ratio = min(max(decay_ratio, 0.0), 1.0)
|
| 522 |
+
coeff = 0.5 * (1.0 + math.cos(math.pi * decay_ratio)) # 1 -> 0
|
| 523 |
+
return min_lr + coeff * (max_lr - min_lr)
|
| 524 |
+
|
| 525 |
+
|
| 526 |
+
# =============================================================================
|
| 527 |
+
# 3. Data — port of the C++ DataLoaderLite (Tensor-Implementations/include/
|
| 528 |
+
# DataLoader.h), including its pinned-host / side-stream prefetch ring.
|
| 529 |
+
#
|
| 530 |
+
# Shards are HEADERLESS flat uint16 token files whose basename contains the
|
| 531 |
+
# split ("train" / "val"). For a batch at position p:
|
| 532 |
+
# x = tokens[p : p+B*T] y = tokens[p+1 : p+B*T+1]
|
| 533 |
+
# p strides by B*T (world_size == 1 here). A shard tail that cannot fit a full
|
| 534 |
+
# stride is dropped and the loader wraps to the next shard.
|
| 535 |
+
#
|
| 536 |
+
# The prefetch mirrors the C++ ring: NBUF pinned staging buffers and NBUF device
|
| 537 |
+
# tensors, H2D issued on a side stream, CUDA events ordering producer against
|
| 538 |
+
# consumer. Each slot also carries the (shard, token offset) it was filled from,
|
| 539 |
+
# which is what --diag prints to name the bytes behind a loss spike.
|
| 540 |
+
# =============================================================================
|
| 541 |
+
class DataLoaderLite:
|
| 542 |
+
NBUF = 4
|
| 543 |
+
|
| 544 |
+
def __init__(self, B, T, split, data_root, device, master=True,
|
| 545 |
+
max_tokens_per_shard=100_000_000, rank=0, world_size=1):
|
| 546 |
+
assert split in ("train", "val")
|
| 547 |
+
self.B, self.T = B, T
|
| 548 |
+
self.split = split
|
| 549 |
+
self.max_tokens = max_tokens_per_shard
|
| 550 |
+
self.device = device
|
| 551 |
+
# DDP: every rank reads the SAME shard list but a DISJOINT, interleaved
|
| 552 |
+
# slice of it (offset by rank, strided by world_size), so a training
|
| 553 |
+
# step's world_size*grad_accum micro-batches never repeat a window.
|
| 554 |
+
self.rank = rank
|
| 555 |
+
self.world_size = world_size
|
| 556 |
+
|
| 557 |
+
self.shards = sorted(
|
| 558 |
+
p for p in glob.glob(os.path.join(data_root, "*.bin"))
|
| 559 |
+
if split in os.path.basename(p)
|
| 560 |
+
)
|
| 561 |
+
if not self.shards:
|
| 562 |
+
raise RuntimeError(f"no .bin shards found for split {split} in {data_root}")
|
| 563 |
+
if master:
|
| 564 |
+
print(f"found {len(self.shards)} shards for split {split}")
|
| 565 |
+
|
| 566 |
+
BT = B * T
|
| 567 |
+
self.pin_x = [torch.empty(BT, dtype=torch.int64).pin_memory() for _ in range(self.NBUF)]
|
| 568 |
+
self.pin_y = [torch.empty(BT, dtype=torch.int64).pin_memory() for _ in range(self.NBUF)]
|
| 569 |
+
self.dev_x = [torch.empty(B, T, dtype=torch.int64, device=device) for _ in range(self.NBUF)]
|
| 570 |
+
self.dev_y = [torch.empty(B, T, dtype=torch.int64, device=device) for _ in range(self.NBUF)]
|
| 571 |
+
self.copy_stream = torch.cuda.Stream(device=device)
|
| 572 |
+
self.h2d_done = [torch.cuda.Event() for _ in range(self.NBUF)]
|
| 573 |
+
self.h2d_recorded = [False] * self.NBUF
|
| 574 |
+
self.where = [(0, 0)] * self.NBUF # (shard_index, token_offset)
|
| 575 |
+
self.reset()
|
| 576 |
+
|
| 577 |
+
# -- shard plumbing ------------------------------------------------------
|
| 578 |
+
def _load_shard(self, i):
|
| 579 |
+
self.cur = i
|
| 580 |
+
arr = np.memmap(self.shards[i], dtype=np.uint16, mode="r")
|
| 581 |
+
self.tok = arr
|
| 582 |
+
self.ntok = min(len(arr), self.max_tokens)
|
| 583 |
+
|
| 584 |
+
def _advance_shard(self):
|
| 585 |
+
self._load_shard((self.cur + 1) % len(self.shards))
|
| 586 |
+
self.pos = self.B * self.T * self.rank
|
| 587 |
+
|
| 588 |
+
def reset(self):
|
| 589 |
+
self.copy_stream.synchronize()
|
| 590 |
+
self._load_shard(0)
|
| 591 |
+
self.pos = self.B * self.T * self.rank
|
| 592 |
+
self.primed = False
|
| 593 |
+
self.buf = 0
|
| 594 |
+
|
| 595 |
+
# -- prefetch ------------------------------------------------------------
|
| 596 |
+
def _stage(self, idx):
|
| 597 |
+
BT = self.B * self.T
|
| 598 |
+
stride = BT * self.world_size
|
| 599 |
+
if self.ntok < BT + 1:
|
| 600 |
+
raise RuntimeError(f"shard too small for one batch: {self.shards[self.cur]}")
|
| 601 |
+
if self.pos + BT + 1 > self.ntok:
|
| 602 |
+
self._advance_shard()
|
| 603 |
+
|
| 604 |
+
# Host-side WAR guard: the pinned buffer may still be the source of the
|
| 605 |
+
# H2D issued the last time this slot was filled (that copy is gated on
|
| 606 |
+
# the device, not the host, so a host running ahead of the GPU would
|
| 607 |
+
# otherwise overwrite bytes still in flight and hand the model a batch
|
| 608 |
+
# spliced from two different windows).
|
| 609 |
+
if self.h2d_recorded[idx]:
|
| 610 |
+
self.h2d_done[idx].synchronize()
|
| 611 |
+
|
| 612 |
+
self.where[idx] = (self.cur, self.pos)
|
| 613 |
+
buf = np.asarray(self.tok[self.pos: self.pos + BT + 1]).astype(np.int64)
|
| 614 |
+
self.pin_x[idx].copy_(torch.from_numpy(buf[:-1]))
|
| 615 |
+
self.pin_y[idx].copy_(torch.from_numpy(buf[1:]))
|
| 616 |
+
|
| 617 |
+
with torch.cuda.stream(self.copy_stream):
|
| 618 |
+
self.dev_x[idx].view(-1).copy_(self.pin_x[idx], non_blocking=True)
|
| 619 |
+
self.dev_y[idx].view(-1).copy_(self.pin_y[idx], non_blocking=True)
|
| 620 |
+
self.h2d_done[idx].record(self.copy_stream)
|
| 621 |
+
self.h2d_recorded[idx] = True
|
| 622 |
+
|
| 623 |
+
self.pos += stride
|
| 624 |
+
if self.pos + BT + 1 > self.ntok:
|
| 625 |
+
self._advance_shard()
|
| 626 |
+
|
| 627 |
+
def next_batch(self):
|
| 628 |
+
"""-> (x, y, shard_index, token_offset), x/y int64 [B,T] on device."""
|
| 629 |
+
if not self.primed:
|
| 630 |
+
for i in range(self.NBUF):
|
| 631 |
+
self._stage(i)
|
| 632 |
+
self.primed = True
|
| 633 |
+
self.buf = 0
|
| 634 |
+
else:
|
| 635 |
+
# Refill the slot returned LAST call. Its consumer's kernels are
|
| 636 |
+
# already queued on the compute stream, so the producer must wait
|
| 637 |
+
# for them before overwriting the device tensor.
|
| 638 |
+
self.copy_stream.wait_stream(torch.cuda.current_stream())
|
| 639 |
+
self._stage((self.buf - 1) % self.NBUF)
|
| 640 |
+
|
| 641 |
+
torch.cuda.current_stream().wait_event(self.h2d_done[self.buf])
|
| 642 |
+
i = self.buf
|
| 643 |
+
self.buf = (self.buf + 1) % self.NBUF
|
| 644 |
+
return self.dev_x[i], self.dev_y[i], self.where[i][0], self.where[i][1]
|
| 645 |
+
|
| 646 |
+
def skip_batches(self, n):
|
| 647 |
+
"""Fast-forward `n` batches without touching the GPU (resume path)."""
|
| 648 |
+
self.copy_stream.synchronize()
|
| 649 |
+
self.primed = False
|
| 650 |
+
BT = self.B * self.T
|
| 651 |
+
stride = BT * self.world_size
|
| 652 |
+
for _ in range(n):
|
| 653 |
+
if self.pos + BT + 1 > self.ntok:
|
| 654 |
+
self._advance_shard()
|
| 655 |
+
self.pos += stride
|
| 656 |
+
if self.pos + BT + 1 > self.ntok:
|
| 657 |
+
self._advance_shard()
|
| 658 |
+
|
| 659 |
+
|
| 660 |
+
# =============================================================================
|
| 661 |
+
# 4. Model.
|
| 662 |
+
#
|
| 663 |
+
# Precision (bf16 mixed precision; both --precision modes run under autocast):
|
| 664 |
+
# bf16-params : Linear weights stored bf16, embedding / RMSNorm / QK-norm
|
| 665 |
+
# scales stay fp32, fp32 master weights live in the optimizer
|
| 666 |
+
# — the exact scheme bluscript_single.cpp uses (default).
|
| 667 |
+
# bf16 : all params fp32, autocast casts per-op (standard PyTorch AMP).
|
| 668 |
+
# In both cases: matmuls + attention bf16, RMSNorm / QK-norm reductions, RoPE,
|
| 669 |
+
# softmax and the loss in fp32. No GradScaler (bf16 has fp32's exponent range).
|
| 670 |
+
# =============================================================================
|
| 671 |
+
_FLASH = [SDPBackend.FLASH_ATTENTION, SDPBackend.EFFICIENT_ATTENTION]
|
| 672 |
+
|
| 673 |
+
|
| 674 |
+
class RMSNorm(nn.Module):
|
| 675 |
+
"""RMSNorm over d_model, fp32 weight. F.rms_norm is a fused ATen op and is
|
| 676 |
+
not autocast-eligible, so it runs in the input dtype — fp32 on the residual
|
| 677 |
+
stream, matching the C++ nn::RMSNorm."""
|
| 678 |
+
|
| 679 |
+
def __init__(self, dim, eps):
|
| 680 |
+
super().__init__()
|
| 681 |
+
self.weight = nn.Parameter(torch.ones(dim))
|
| 682 |
+
self.eps = eps
|
| 683 |
+
self.dim = (dim,)
|
| 684 |
+
|
| 685 |
+
def forward(self, x):
|
| 686 |
+
return F.rms_norm(x, self.dim, self.weight, self.eps)
|
| 687 |
+
|
| 688 |
+
|
| 689 |
+
class CausalGQA(nn.Module):
|
| 690 |
+
"""RMSNorm -> fused(QK-norm + RoPE) -> causal grouped-query attention -> Wo -> +x"""
|
| 691 |
+
|
| 692 |
+
def __init__(self, cfg, resid_scale):
|
| 693 |
+
super().__init__()
|
| 694 |
+
self.cfg = cfg
|
| 695 |
+
self.norm = RMSNorm(cfg.d_model, cfg.rms_eps)
|
| 696 |
+
self.wQ = nn.Linear(cfg.d_model, cfg.q_heads * cfg.head_dim, bias=False)
|
| 697 |
+
self.wK = nn.Linear(cfg.d_model, cfg.kv_heads * cfg.head_dim, bias=False)
|
| 698 |
+
self.wV = nn.Linear(cfg.d_model, cfg.kv_heads * cfg.head_dim, bias=False)
|
| 699 |
+
self.wO = nn.Linear(cfg.q_heads * cfg.head_dim, cfg.d_model, bias=False)
|
| 700 |
+
# QK-norm scales, initialised to 1.0 (identity RMS scale), fp32.
|
| 701 |
+
self.q_gamma = nn.Parameter(torch.ones(cfg.head_dim))
|
| 702 |
+
self.k_gamma = nn.Parameter(torch.ones(cfg.head_dim))
|
| 703 |
+
# GPT-2/Llama init: projections std=0.02, residual proj scaled by
|
| 704 |
+
# (2*n_layers)^-0.5 to keep residual-stream variance ~constant.
|
| 705 |
+
for lin in (self.wQ, self.wK, self.wV):
|
| 706 |
+
nn.init.normal_(lin.weight, mean=0.0, std=0.02)
|
| 707 |
+
nn.init.normal_(self.wO.weight, mean=0.0, std=0.02 * resid_scale)
|
| 708 |
+
|
| 709 |
+
def forward(self, x, cos, sin):
|
| 710 |
+
cfg = self.cfg
|
| 711 |
+
B, T, _ = x.shape
|
| 712 |
+
h = self.norm(x)
|
| 713 |
+
|
| 714 |
+
# [B,T,H*D] -> [B,H,T,D]. The transpose stays a view; the fused kernel
|
| 715 |
+
# honours arbitrary (b,h,t) strides, so no .contiguous() is paid here.
|
| 716 |
+
q = self.wQ(h).view(B, T, cfg.q_heads, cfg.head_dim).transpose(1, 2)
|
| 717 |
+
k = self.wK(h).view(B, T, cfg.kv_heads, cfg.head_dim).transpose(1, 2)
|
| 718 |
+
v = self.wV(h).view(B, T, cfg.kv_heads, cfg.head_dim).transpose(1, 2)
|
| 719 |
+
|
| 720 |
+
if cfg.fused_rms_rope:
|
| 721 |
+
q, k = fused_rms_rope(q, k, self.q_gamma, self.k_gamma, cos, sin, cfg.qk_eps)
|
| 722 |
+
else:
|
| 723 |
+
q = rms_rope_eager(q, self.q_gamma, cos, sin, cfg.qk_eps)
|
| 724 |
+
k = rms_rope_eager(k, self.k_gamma, cos, sin, cfg.qk_eps)
|
| 725 |
+
|
| 726 |
+
# Separate fused flash kernel. enable_gqa maps query head i to kv head
|
| 727 |
+
# i // (q_heads/kv_heads), the same grouping as the C++ kernel.
|
| 728 |
+
with sdpa_kernel(_FLASH):
|
| 729 |
+
o = F.scaled_dot_product_attention(
|
| 730 |
+
q.to(v.dtype), k.to(v.dtype), v, is_causal=True, enable_gqa=True)
|
| 731 |
+
|
| 732 |
+
o = o.transpose(1, 2).reshape(B, T, cfg.q_heads * cfg.head_dim)
|
| 733 |
+
return x + self.wO(o)
|
| 734 |
+
|
| 735 |
+
|
| 736 |
+
class SwiGLUMLP(nn.Module):
|
| 737 |
+
"""RMSNorm -> gate_up -> silu(gate)*up -> down -> +x ([gate | up] layout)"""
|
| 738 |
+
|
| 739 |
+
def __init__(self, cfg, resid_scale):
|
| 740 |
+
super().__init__()
|
| 741 |
+
self.norm = RMSNorm(cfg.d_model, cfg.rms_eps)
|
| 742 |
+
self.gate_up = nn.Linear(cfg.d_model, 2 * cfg.ffn_hidden, bias=False)
|
| 743 |
+
self.down = nn.Linear(cfg.ffn_hidden, cfg.d_model, bias=False)
|
| 744 |
+
nn.init.normal_(self.gate_up.weight, mean=0.0, std=0.02)
|
| 745 |
+
nn.init.normal_(self.down.weight, mean=0.0, std=0.02 * resid_scale)
|
| 746 |
+
|
| 747 |
+
def forward(self, x):
|
| 748 |
+
gate, up = self.gate_up(self.norm(x)).chunk(2, dim=-1)
|
| 749 |
+
return x + self.down(F.silu(gate) * up)
|
| 750 |
+
|
| 751 |
+
|
| 752 |
+
class Block(nn.Module):
|
| 753 |
+
def __init__(self, cfg, resid_scale):
|
| 754 |
+
super().__init__()
|
| 755 |
+
self.attn = CausalGQA(cfg, resid_scale)
|
| 756 |
+
self.mlp = SwiGLUMLP(cfg, resid_scale)
|
| 757 |
+
|
| 758 |
+
def forward(self, x, cos, sin):
|
| 759 |
+
return self.mlp(self.attn(x, cos, sin))
|
| 760 |
+
|
| 761 |
+
|
| 762 |
+
class BLU(nn.Module):
|
| 763 |
+
def __init__(self, cfg):
|
| 764 |
+
super().__init__()
|
| 765 |
+
self.cfg = cfg
|
| 766 |
+
resid_scale = 1.0 / math.sqrt(2.0 * cfg.n_layers)
|
| 767 |
+
|
| 768 |
+
self.wte = nn.Embedding(cfg.vocab_size, cfg.d_model)
|
| 769 |
+
nn.init.normal_(self.wte.weight, mean=0.0, std=0.02)
|
| 770 |
+
|
| 771 |
+
self.blocks = nn.ModuleList([Block(cfg, resid_scale) for _ in range(cfg.n_layers)])
|
| 772 |
+
self.norm_f = RMSNorm(cfg.d_model, cfg.rms_eps)
|
| 773 |
+
if not cfg.weight_tying:
|
| 774 |
+
self.lm_head = nn.Linear(cfg.d_model, cfg.vocab_size, bias=False)
|
| 775 |
+
nn.init.normal_(self.lm_head.weight, mean=0.0, std=0.02)
|
| 776 |
+
|
| 777 |
+
cos, sin = build_rope_cache(cfg.context_length, cfg.head_dim, cfg.rope_theta, "cpu")
|
| 778 |
+
self.register_buffer("rope_cos", cos, persistent=False)
|
| 779 |
+
self.register_buffer("rope_sin", sin, persistent=False)
|
| 780 |
+
|
| 781 |
+
def apply_param_dtype(self, mode):
|
| 782 |
+
"""`bf16-params`: store Linear weights in bf16 (embedding, RMSNorm and
|
| 783 |
+
QK-norm scales stay fp32) — the two-dtype-group layout of the C++ model.
|
| 784 |
+
Under autocast this is transparent: the cast to bf16 becomes a no-op."""
|
| 785 |
+
if mode != "bf16-params":
|
| 786 |
+
return
|
| 787 |
+
for m in self.modules():
|
| 788 |
+
if isinstance(m, nn.Linear):
|
| 789 |
+
m.weight.data = m.weight.data.to(torch.bfloat16)
|
| 790 |
+
|
| 791 |
+
def forward(self, idx, targets=None):
|
| 792 |
+
cfg = self.cfg
|
| 793 |
+
T = idx.shape[1]
|
| 794 |
+
cos, sin = self.rope_cos[:T], self.rope_sin[:T]
|
| 795 |
+
|
| 796 |
+
x = self.wte(idx) # fp32 [B,T,d_model]
|
| 797 |
+
for blk in self.blocks:
|
| 798 |
+
if cfg.activation_checkpoint and self.training:
|
| 799 |
+
x = checkpoint(blk, x, cos, sin, use_reentrant=False)
|
| 800 |
+
else:
|
| 801 |
+
x = blk(x, cos, sin)
|
| 802 |
+
x = self.norm_f(x)
|
| 803 |
+
|
| 804 |
+
head_w = self.wte.weight if cfg.weight_tying else self.lm_head.weight
|
| 805 |
+
if targets is None:
|
| 806 |
+
return F.linear(x, head_w), None
|
| 807 |
+
if cfg.loss_chunks > 1:
|
| 808 |
+
return None, _chunked_ce(x, head_w, targets, cfg.loss_chunks)
|
| 809 |
+
# Matches the C++: bf16 matmul, logits promoted to fp32 for the loss
|
| 810 |
+
# (cross_entropy is an autocast-fp32 op, so the promotion is implicit).
|
| 811 |
+
logits = F.linear(x, head_w)
|
| 812 |
+
loss = F.cross_entropy(logits.view(-1, logits.size(-1)), targets.reshape(-1))
|
| 813 |
+
return logits, loss
|
| 814 |
+
|
| 815 |
+
|
| 816 |
+
def _chunked_ce(x, weight, targets, chunks):
|
| 817 |
+
"""lm_head + cross-entropy in row chunks, logits recomputed in backward.
|
| 818 |
+
Peak logit memory drops from [B*T, V] to [B*T/chunks, V] at the cost of one
|
| 819 |
+
extra lm_head matmul. Off by default (`--loss-chunks 1`) so the memory
|
| 820 |
+
profile matches the C++ script, which materialises the full fp32 logits.
|
| 821 |
+
"""
|
| 822 |
+
xf = x.reshape(-1, x.shape[-1])
|
| 823 |
+
tf = targets.reshape(-1)
|
| 824 |
+
n = xf.shape[0]
|
| 825 |
+
size = (n + chunks - 1) // chunks
|
| 826 |
+
total = xf.new_zeros((), dtype=torch.float32)
|
| 827 |
+
for lo in range(0, n, size):
|
| 828 |
+
hi = min(lo + size, n)
|
| 829 |
+
part = checkpoint(_ce_chunk, xf[lo:hi], weight, tf[lo:hi], use_reentrant=False)
|
| 830 |
+
total = total + part * (hi - lo)
|
| 831 |
+
return total / n
|
| 832 |
+
|
| 833 |
+
|
| 834 |
+
def _ce_chunk(x, weight, targets):
|
| 835 |
+
return F.cross_entropy(F.linear(x, weight).float(), targets)
|
| 836 |
+
|
| 837 |
+
|
| 838 |
+
# =============================================================================
|
| 839 |
+
# 5. Optimizer (Muon on the hidden matrices, AdamW on the rest), gradient
|
| 840 |
+
# clipping, checkpointing.
|
| 841 |
+
# =============================================================================
|
| 842 |
+
def _is_muon_param(name, p):
|
| 843 |
+
"""Muon orthogonalizes a matrix, so it needs the parameter to BE a linear
|
| 844 |
+
map between two feature spaces:
|
| 845 |
+
* dim >= 2 — RMSNorm weights and QK-norm gammas are 1D and have no
|
| 846 |
+
singular values to equalise;
|
| 847 |
+
* not the embedding / lm_head — its rows index tokens, not features, so
|
| 848 |
+
the left singular vectors mix unrelated vocabulary entries. Weight
|
| 849 |
+
tying makes this one tensor serving both ends; either way it is AdamW's.
|
| 850 |
+
Everything else — wQ/wK/wV/wO, gate_up, down — is Muon's: 75.5M of the
|
| 851 |
+
114.2M parameters at the default config.
|
| 852 |
+
"""
|
| 853 |
+
return p.dim() >= 2 and not name.startswith(("wte.", "lm_head."))
|
| 854 |
+
|
| 855 |
+
|
| 856 |
+
def _state_bytes(opt):
|
| 857 |
+
"""Bytes held in a torch optimizer's per-parameter state — Adam's m/v, or
|
| 858 |
+
Muon's momentum buffer. Reads 0 before the first step(): both allocate
|
| 859 |
+
lazily."""
|
| 860 |
+
return sum(v.numel() * v.element_size()
|
| 861 |
+
for st in opt.state.values()
|
| 862 |
+
for v in st.values() if torch.is_tensor(v))
|
| 863 |
+
|
| 864 |
+
|
| 865 |
+
class MasterMuon:
|
| 866 |
+
"""torch.optim.Muon on the hidden matrices + torch.optim.AdamW on the rest,
|
| 867 |
+
both driving fp32 master weights.
|
| 868 |
+
Neither optimizer is reimplemented here. What this class owns is the part
|
| 869 |
+
torch cannot know about:
|
| 870 |
+
* ROUTING — torch.optim.Muon accepts whatever you hand it and the docs are
|
| 871 |
+
explicit that only 2D hidden weights belong ("bias, and embedding,
|
| 872 |
+
should be optimized by a standard method such as AdamW"). The split is
|
| 873 |
+
`_is_muon_param`.
|
| 874 |
+
* MASTER WEIGHTS — the parameters are bf16 under --precision bf16-params,
|
| 875 |
+
and a bf16 tensor cannot absorb an lr*update of 1e-4 relative size. So
|
| 876 |
+
both optimizers are constructed over the fp32 MASTERS, fed from the fp32
|
| 877 |
+
landing buffers, and the masters are cast back over the params after the
|
| 878 |
+
step. Unchanged from the AdamW arm — that machinery is what makes the
|
| 879 |
+
two arms comparable, since both optimizers then read a bit-identical
|
| 880 |
+
fp32 gradient and write a bit-identical bf16 parameter.
|
| 881 |
+
Weight decay is decoupled in both (torch.optim.Muon applies it before the
|
| 882 |
+
update, as AdamW does) and is applied to the master for the same reason: a
|
| 883 |
+
decay written to the bf16 param is overwritten by the master write-back
|
| 884 |
+
microseconds later and silently no-ops.
|
| 885 |
+
State cost: Muon holds one momentum buffer per param, against Adam's two.
|
| 886 |
+
"""
|
| 887 |
+
|
| 888 |
+
def __init__(self, named_params, cfg):
|
| 889 |
+
self.cfg = cfg
|
| 890 |
+
self.triples = [] # (param, master, landing_grad), ALL params
|
| 891 |
+
self.muon_masters = [] # masters routed to Muon
|
| 892 |
+
decay, nodecay = [], []
|
| 893 |
+
self.n_muon = self.n_decay = self.n_nodecay = 0
|
| 894 |
+
for name, p in named_params:
|
| 895 |
+
if p.dtype == torch.float32:
|
| 896 |
+
m, g = p, None
|
| 897 |
+
else:
|
| 898 |
+
m = p.detach().float().clone()
|
| 899 |
+
g = torch.zeros_like(m)
|
| 900 |
+
m.grad = g
|
| 901 |
+
self.triples.append((p, m, g))
|
| 902 |
+
if _is_muon_param(name, p):
|
| 903 |
+
self.muon_masters.append(m)
|
| 904 |
+
self.n_muon += m.numel()
|
| 905 |
+
elif p.dim() >= 2:
|
| 906 |
+
decay.append(m)
|
| 907 |
+
self.n_decay += m.numel()
|
| 908 |
+
else:
|
| 909 |
+
nodecay.append(m)
|
| 910 |
+
self.n_nodecay += m.numel()
|
| 911 |
+
self.opt = torch.optim.AdamW(
|
| 912 |
+
[{"params": decay, "weight_decay": cfg.weight_decay},
|
| 913 |
+
{"params": nodecay, "weight_decay": 0.0}],
|
| 914 |
+
lr=cfg.max_lr, betas=(0.9, 0.95), eps=1e-8, fused=True)
|
| 915 |
+
# Everything below is a torch.optim.Muon default except adjust_lr_fn
|
| 916 |
+
# (torch: "original") and weight_decay, which is pinned to the AdamW
|
| 917 |
+
# arm's 0.1 so the decay is identical across both groups and both arms.
|
| 918 |
+
self.muon = torch.optim.Muon(
|
| 919 |
+
self.muon_masters,
|
| 920 |
+
lr=cfg.muon_lr,
|
| 921 |
+
weight_decay=cfg.weight_decay,
|
| 922 |
+
momentum=cfg.muon_momentum,
|
| 923 |
+
nesterov=True,
|
| 924 |
+
ns_steps=cfg.muon_ns_steps,
|
| 925 |
+
adjust_lr_fn=cfg.muon_adjust_lr)
|
| 926 |
+
|
| 927 |
+
@torch.no_grad()
|
| 928 |
+
def step(self, lr, muon_lr):
|
| 929 |
+
for p, m, g in self.triples:
|
| 930 |
+
if g is not None:
|
| 931 |
+
g.copy_(p.grad) # bf16 grad -> fp32 landing buffer
|
| 932 |
+
# AdamW group (embedding / lm_head / norms / gammas) — identical to the
|
| 933 |
+
# baseline's update, on identical gradients.
|
| 934 |
+
for grp in self.opt.param_groups:
|
| 935 |
+
grp["lr"] = lr
|
| 936 |
+
self.opt.step()
|
| 937 |
+
# Muon group. Its nesterov path may consume the gradient in place; that
|
| 938 |
+
# is harmless here because grad_norm and the --diag probes are read
|
| 939 |
+
# before step() and the landing buffers are overwritten wholesale on the
|
| 940 |
+
# next one.
|
| 941 |
+
for grp in self.muon.param_groups:
|
| 942 |
+
grp["lr"] = muon_lr
|
| 943 |
+
self.muon.step()
|
| 944 |
+
for p, m, g in self.triples:
|
| 945 |
+
if g is not None:
|
| 946 |
+
p.data.copy_(m) # fp32 master -> bf16 param
|
| 947 |
+
|
| 948 |
+
def state_breakdown(self):
|
| 949 |
+
"""-> (master_bytes, landing_grad_bytes, adam_mv_bytes, muon_mom_bytes).
|
| 950 |
+
The first two are zero under --precision bf16, where the master IS the
|
| 951 |
+
parameter and the gradient is already fp32. Both optimizers allocate
|
| 952 |
+
their state lazily on the first step(), so the last two read 0 until
|
| 953 |
+
then — the startup report uses the expected_* predictors instead.
|
| 954 |
+
"""
|
| 955 |
+
master = landing = 0
|
| 956 |
+
for p, m, g in self.triples:
|
| 957 |
+
if m is not p:
|
| 958 |
+
master += m.numel() * m.element_size()
|
| 959 |
+
landing += g.numel() * g.element_size()
|
| 960 |
+
mv = _state_bytes(self.opt)
|
| 961 |
+
mom = _state_bytes(self.muon)
|
| 962 |
+
return master, landing, mv, mom
|
| 963 |
+
|
| 964 |
+
def state_bytes(self):
|
| 965 |
+
return sum(self.state_breakdown())
|
| 966 |
+
|
| 967 |
+
def expected_mv_bytes(self):
|
| 968 |
+
"""Adam m/v are lazily allocated, so the startup report has to predict
|
| 969 |
+
them rather than measure them. Muon's params are excluded — they carry a
|
| 970 |
+
momentum buffer instead, reported separately."""
|
| 971 |
+
muon_ids = {id(m) for m in self.muon_masters}
|
| 972 |
+
return sum(m.numel() * m.element_size() * 2
|
| 973 |
+
for _, m, _ in self.triples if id(m) not in muon_ids)
|
| 974 |
+
|
| 975 |
+
def expected_muon_mom_bytes(self):
|
| 976 |
+
"""One fp32 momentum buffer per Muon parameter — the prediction the
|
| 977 |
+
startup report prints before step 0 has allocated it."""
|
| 978 |
+
return sum(m.numel() * m.element_size() for m in self.muon_masters)
|
| 979 |
+
|
| 980 |
+
# -- checkpointing -------------------------------------------------------
|
| 981 |
+
def state_dict(self):
|
| 982 |
+
# `optimizer` tags the arm: resuming a Muon run from an AdamW checkpoint
|
| 983 |
+
# would restore the weights and then silently drop the momentum, so
|
| 984 |
+
# CheckpointManager refuses on a mismatch rather than producing a run
|
| 985 |
+
# that looks fine and is not comparable to either arm.
|
| 986 |
+
return {"optimizer": "muon",
|
| 987 |
+
"adam": self.opt.state_dict(),
|
| 988 |
+
"muon": self.muon.state_dict(),
|
| 989 |
+
"masters": [None if m is p else m for p, m, g in self.triples]}
|
| 990 |
+
|
| 991 |
+
@torch.no_grad()
|
| 992 |
+
def load_state_dict(self, blob):
|
| 993 |
+
"""-> True if the fp32 masters were restored from the checkpoint."""
|
| 994 |
+
if blob.get("optimizer") != "muon":
|
| 995 |
+
raise SystemExit(
|
| 996 |
+
"ERROR: this checkpoint was written by the AdamW arm (train.py). "
|
| 997 |
+
"Point --ckpt-dir somewhere else, or resume it with train.py.")
|
| 998 |
+
self.opt.load_state_dict(blob["adam"])
|
| 999 |
+
if blob.get("muon") is not None:
|
| 1000 |
+
self.muon.load_state_dict(blob["muon"])
|
| 1001 |
+
masters = blob.get("masters")
|
| 1002 |
+
if masters is None or len(masters) != len(self.triples):
|
| 1003 |
+
return False
|
| 1004 |
+
for (p, m, g), saved in zip(self.triples, masters):
|
| 1005 |
+
if saved is not None and m is not p:
|
| 1006 |
+
m.copy_(saved.to(m.device))
|
| 1007 |
+
return True
|
| 1008 |
+
|
| 1009 |
+
@torch.no_grad()
|
| 1010 |
+
def resync_masters(self):
|
| 1011 |
+
"""Re-seed every master from its parameter.
|
| 1012 |
+
Needed after loading a checkpoint that carries no masters: the masters
|
| 1013 |
+
were built in __init__ from the RANDOM INIT, so without this the first
|
| 1014 |
+
step() would write them straight back over the restored weights and
|
| 1015 |
+
silently throw the run away.
|
| 1016 |
+
"""
|
| 1017 |
+
for p, m, g in self.triples:
|
| 1018 |
+
if m is not p:
|
| 1019 |
+
m.copy_(p.detach().float())
|
| 1020 |
+
|
| 1021 |
+
|
| 1022 |
+
def clip_grads_(params, max_norm):
|
| 1023 |
+
"""ONE global L2 norm over every parameter, then one rescale — the textbook
|
| 1024 |
+
definition. (It differs from the ZeRO script, which clips each dtype group
|
| 1025 |
+
against its own norm and can therefore let the true global norm reach
|
| 1026 |
+
sqrt(n_groups) * grad_clip.) The per-tensor norms are promoted to fp32
|
| 1027 |
+
before they are combined, or a reduction over ~250 bf16 norms would lose
|
| 1028 |
+
~3 digits."""
|
| 1029 |
+
grads = [p.grad for p in params if p.grad is not None]
|
| 1030 |
+
norms = torch._foreach_norm(grads)
|
| 1031 |
+
total = torch.linalg.vector_norm(torch.stack([n.float() for n in norms]))
|
| 1032 |
+
if max_norm and max_norm > 0:
|
| 1033 |
+
clip = (max_norm / (total + 1e-6)).clamp(max=1.0)
|
| 1034 |
+
torch._foreach_mul_(grads, clip)
|
| 1035 |
+
return total
|
| 1036 |
+
|
| 1037 |
+
|
| 1038 |
+
class CheckpointManager:
|
| 1039 |
+
"""Synchronous save / keep-N / auto-resume, mirroring the C++
|
| 1040 |
+
CheckpointManager. A checkpoint holds the model parameters, Adam's m/v, the
|
| 1041 |
+
fp32 masters, and the step it should resume at."""
|
| 1042 |
+
|
| 1043 |
+
def __init__(self, ckpt_dir, prefix, keep):
|
| 1044 |
+
self.dir, self.prefix, self.keep = ckpt_dir, prefix, keep
|
| 1045 |
+
os.makedirs(ckpt_dir, exist_ok=True)
|
| 1046 |
+
|
| 1047 |
+
def _path(self, step):
|
| 1048 |
+
return os.path.join(self.dir, f"{self.prefix}_step{step:08d}.pt")
|
| 1049 |
+
|
| 1050 |
+
def _existing(self):
|
| 1051 |
+
out = []
|
| 1052 |
+
for p in glob.glob(os.path.join(self.dir, f"{self.prefix}_step*.pt")):
|
| 1053 |
+
m = re.search(r"_step(\d+)\.pt$", os.path.basename(p))
|
| 1054 |
+
if m:
|
| 1055 |
+
out.append((int(m.group(1)), p))
|
| 1056 |
+
return sorted(out)
|
| 1057 |
+
|
| 1058 |
+
def save(self, step, model, opt, loss, last_val, cfg):
|
| 1059 |
+
tmp = self._path(step) + ".tmp"
|
| 1060 |
+
torch.save({"step": step, "loss": loss, "last_val": last_val,
|
| 1061 |
+
# B/T/global_batch pin the data stride this checkpoint's
|
| 1062 |
+
# step count assumes. world_size is deliberately NOT
|
| 1063 |
+
# recorded: grad_accum = global_batch // (B*T*world_size)
|
| 1064 |
+
# already makes tokens-per-step (and skip_batches'
|
| 1065 |
+
# target offset) independent of how many ranks are
|
| 1066 |
+
# running, so a GPU-count change alone is not a mismatch.
|
| 1067 |
+
"run_cfg": {"B": cfg.B, "T": cfg.T,
|
| 1068 |
+
"global_batch": cfg.global_batch},
|
| 1069 |
+
"model": model.state_dict(),
|
| 1070 |
+
"opt": opt.state_dict()}, tmp)
|
| 1071 |
+
# Rename last: a crash mid-write leaves a .tmp, never a truncated
|
| 1072 |
+
# checkpoint that the next load_latest would pick as the newest.
|
| 1073 |
+
os.replace(tmp, self._path(step))
|
| 1074 |
+
if self.keep > 0:
|
| 1075 |
+
for _, stale in self._existing()[:-self.keep]:
|
| 1076 |
+
os.remove(stale)
|
| 1077 |
+
|
| 1078 |
+
def load_latest(self, model, opt, cfg, master=True):
|
| 1079 |
+
"""-> (start_step, loss, last_val). (0, 0.0, nan) when there is
|
| 1080 |
+
nothing to resume."""
|
| 1081 |
+
existing = self._existing()
|
| 1082 |
+
if not existing:
|
| 1083 |
+
return 0, 0.0, float("nan")
|
| 1084 |
+
_, path = existing[-1]
|
| 1085 |
+
return self.load(path, model, opt, cfg, master)
|
| 1086 |
+
|
| 1087 |
+
def load(self, path, model, opt, cfg, master=True):
|
| 1088 |
+
"""Load a specific checkpoint file (e.g. from --resume-from), rather
|
| 1089 |
+
than auto-discovering the latest one in self.dir.
|
| 1090 |
+
-> (start_step, loss, last_val)."""
|
| 1091 |
+
blob = torch.load(path, map_location="cpu", weights_only=True)
|
| 1092 |
+
model.load_state_dict(blob["model"])
|
| 1093 |
+
if not opt.load_state_dict(blob["opt"]):
|
| 1094 |
+
opt.resync_masters()
|
| 1095 |
+
run_cfg = blob.get("run_cfg")
|
| 1096 |
+
if master and run_cfg is not None:
|
| 1097 |
+
mismatches = [f"{k} checkpoint={v!r} current={getattr(cfg, k)!r}"
|
| 1098 |
+
for k, v in run_cfg.items() if getattr(cfg, k) != v]
|
| 1099 |
+
if mismatches:
|
| 1100 |
+
print("[WARN] resume config differs from the checkpoint's — "
|
| 1101 |
+
"the data offset this step count assumes will be wrong:")
|
| 1102 |
+
for m in mismatches:
|
| 1103 |
+
print(f" {m}")
|
| 1104 |
+
print(f"[resume] loaded {path}")
|
| 1105 |
+
return (int(blob["step"]), float(blob["loss"]),
|
| 1106 |
+
float(blob.get("last_val", float("nan"))))
|
| 1107 |
+
|
| 1108 |
+
|
| 1109 |
+
# =============================================================================
|
| 1110 |
+
# 6. Training.
|
| 1111 |
+
# =============================================================================
|
| 1112 |
+
def _env_int(name, default=0):
|
| 1113 |
+
v = os.environ.get(name)
|
| 1114 |
+
if v in (None, ""):
|
| 1115 |
+
return default
|
| 1116 |
+
try:
|
| 1117 |
+
return int(v)
|
| 1118 |
+
except ValueError:
|
| 1119 |
+
print(f"[WARN] ignoring non-integer {name}={v!r}")
|
| 1120 |
+
return default
|
| 1121 |
+
|
| 1122 |
+
|
| 1123 |
+
def seed_everything(seed):
|
| 1124 |
+
"""Every RNG the run can draw from, off one seed.
|
| 1125 |
+
torch's CPU generator is the one that matters here (the model is built on
|
| 1126 |
+
CPU and moved), but `random` / numpy are seeded too so that anything added
|
| 1127 |
+
later — a shuffled shard order, a sampled probe — is reproducible by
|
| 1128 |
+
default rather than by accident.
|
| 1129 |
+
"""
|
| 1130 |
+
random.seed(seed)
|
| 1131 |
+
np.random.seed(seed)
|
| 1132 |
+
torch.manual_seed(seed)
|
| 1133 |
+
torch.cuda.manual_seed_all(seed)
|
| 1134 |
+
|
| 1135 |
+
|
| 1136 |
+
def apply_determinism(mode):
|
| 1137 |
+
"""Pin the run-to-run drift that seeding does not cover.
|
| 1138 |
+
Three sources, all on the backward pass:
|
| 1139 |
+
* nn.Embedding's dense backward scatters one atomicAdd per token into
|
| 1140 |
+
[vocab, d_model]; with weight tying that is also the lm_head gradient,
|
| 1141 |
+
so it is the largest reduction in the step;
|
| 1142 |
+
* the flash / mem-efficient attention backward splits dq over the key
|
| 1143 |
+
blocks and atomicAdds them together;
|
| 1144 |
+
* cuDNN picks kernels by wall-clock timing when `benchmark` is on, so the
|
| 1145 |
+
chosen algorithm can differ between runs of the same shape.
|
| 1146 |
+
Each is a float reduction in an order the scheduler chooses, so two runs
|
| 1147 |
+
diverge in the last mantissa bits of the first step and, through the
|
| 1148 |
+
optimizer, visibly by a few hundred steps.
|
| 1149 |
+
`strict` raises if some op has no deterministic kernel — better than
|
| 1150 |
+
reporting a reproducible run that is not one. `warn` downgrades that to a
|
| 1151 |
+
warning and keeps training; `off` restores the default fast kernels.
|
| 1152 |
+
"""
|
| 1153 |
+
if mode == "off":
|
| 1154 |
+
return
|
| 1155 |
+
torch.backends.cudnn.benchmark = False
|
| 1156 |
+
torch.backends.cudnn.deterministic = True
|
| 1157 |
+
# Routes embedding-backward and SDPA-backward to their ordered variants, and
|
| 1158 |
+
# enforces the cuBLAS workspace setting made at the top of this file.
|
| 1159 |
+
torch.use_deterministic_algorithms(True, warn_only=(mode == "warn"))
|
| 1160 |
+
|
| 1161 |
+
|
| 1162 |
+
def build_argparser():
|
| 1163 |
+
ap = argparse.ArgumentParser(
|
| 1164 |
+
description="BLU single-GPU training — PyTorch port of bluscript_single.cpp")
|
| 1165 |
+
ap.add_argument("--data-root",
|
| 1166 |
+
default=("/root/.cache/huggingface/hub/datasets--navingv--flux/"
|
| 1167 |
+
"snapshots/5232b3d133d2eb71905725f288ff6b56dc0df773"))
|
| 1168 |
+
ap.add_argument("--device", type=int, default=None,
|
| 1169 |
+
help="CUDA device index (env: BLU_DEVICE)")
|
| 1170 |
+
ap.add_argument("--precision", choices=["bf16-params", "bf16"], default="bf16-params",
|
| 1171 |
+
help="bf16-params (default): Linear weights stored bf16 with "
|
| 1172 |
+
"fp32 master weights in the optimizer — the C++ scheme. "
|
| 1173 |
+
"bf16: fp32 params + autocast (standard PyTorch AMP).")
|
| 1174 |
+
ap.add_argument("--rms-rope", choices=["fused", "eager"], default="fused",
|
| 1175 |
+
help="fused = Triton QK-norm+RoPE kernel (default), "
|
| 1176 |
+
"eager = unfused reference, for A/B on the kernel")
|
| 1177 |
+
ap.add_argument("--compile", action="store_true", help="torch.compile each block")
|
| 1178 |
+
ap.add_argument("--activation-checkpoint", action="store_true",
|
| 1179 |
+
help="recompute each block in backward (trades ~30%% time for memory)")
|
| 1180 |
+
ap.add_argument("--loss-chunks", type=int, default=1,
|
| 1181 |
+
help="chunk lm_head+CE into N pieces with recompute; 1 (default) "
|
| 1182 |
+
"materialises the full logits, as the C++ does")
|
| 1183 |
+
# bluscript_single.cpp defaults unless overridden.
|
| 1184 |
+
ap.add_argument("--B", type=int, default=None, help="micro-batch (env: BLU_B)")
|
| 1185 |
+
ap.add_argument("--T", type=int, default=None)
|
| 1186 |
+
ap.add_argument("--d-model", type=int, default=None, help="smoke-test override")
|
| 1187 |
+
ap.add_argument("--n-layers", type=int, default=None, help="smoke-test override")
|
| 1188 |
+
ap.add_argument("--q-heads", type=int, default=None, help="smoke-test override")
|
| 1189 |
+
ap.add_argument("--kv-heads", type=int, default=None, help="smoke-test override")
|
| 1190 |
+
ap.add_argument("--ffn-hidden", type=int, default=None, help="smoke-test override")
|
| 1191 |
+
ap.add_argument("--vocab-size", type=int, default=None, help="smoke-test override")
|
| 1192 |
+
ap.add_argument("--global-batch", type=int, default=None,
|
| 1193 |
+
help="tokens per optimizer step (env: BLU_GLOBAL_BATCH)")
|
| 1194 |
+
ap.add_argument("--max-steps", type=int, default=None,
|
| 1195 |
+
help="caps the run for benchmarking; warmup/stable shrink "
|
| 1196 |
+
"proportionally so the WSD decay tail stays non-empty "
|
| 1197 |
+
"(env: BLU_MAX_STEPS)")
|
| 1198 |
+
# ---- Muon ----
|
| 1199 |
+
ap.add_argument("--muon-lr", type=float, default=None,
|
| 1200 |
+
help="peak LR for the hidden 2D weights. Under the default "
|
| 1201 |
+
"--muon-adjust-lr match_rms_adamw this is the same unit as "
|
| 1202 |
+
"--max-lr, and defaults to the same value (2e-4) so "
|
| 1203 |
+
"both arms take a matched first run; sweep it upward "
|
| 1204 |
+
"from there (Muon usually takes 2-4x AdamW's LR).")
|
| 1205 |
+
ap.add_argument("--muon-min-lr", type=float, default=None,
|
| 1206 |
+
help="floor of the Muon cosine tail (default 10%% of --muon-lr)")
|
| 1207 |
+
ap.add_argument("--muon-momentum", type=float, default=None,
|
| 1208 |
+
help="torch.optim.Muon momentum (default 0.95, torch's own)")
|
| 1209 |
+
ap.add_argument("--muon-ns-steps", type=int, default=None,
|
| 1210 |
+
help="torch.optim.Muon ns_steps, Newton-Schulz iterations "
|
| 1211 |
+
"(default 5, torch's own); fewer = looser "
|
| 1212 |
+
"orthogonalization, cheaper")
|
| 1213 |
+
ap.add_argument("--muon-adjust-lr", choices=["match_rms_adamw", "original"],
|
| 1214 |
+
default=None,
|
| 1215 |
+
help="torch.optim.Muon adjust_lr_fn. match_rms_adamw "
|
| 1216 |
+
"(default HERE): 0.2*sqrt(max(A,B)), putting --muon-lr "
|
| 1217 |
+
"on the same scale as --max-lr. original (default in "
|
| 1218 |
+
"torch): sqrt(max(1,A/B)), needs a ~5.5x larger "
|
| 1219 |
+
"--muon-lr at d_model=768.")
|
| 1220 |
+
ap.add_argument("--val-freq", type=int, default=None)
|
| 1221 |
+
ap.add_argument("--val-loss-steps", type=int, default=None)
|
| 1222 |
+
ap.add_argument("--diag", action="store_true",
|
| 1223 |
+
help="per-micro-batch loss + depth probes (env: BLU_DIAG=1)")
|
| 1224 |
+
ap.add_argument("--warmup-measure", type=int, default=3,
|
| 1225 |
+
help="steps excluded from the steady-state throughput summary")
|
| 1226 |
+
ap.add_argument("--csv", default="blu_muon_log.csv",
|
| 1227 |
+
help="CSV log path; appended (not truncated) on resume. The "
|
| 1228 |
+
"first 8 columns match the AdamW arm's CSV exactly, so "
|
| 1229 |
+
"the two files plot against each other directly; "
|
| 1230 |
+
"muon_lr is appended as a 9th.")
|
| 1231 |
+
ap.add_argument("--seed", type=int, default=1234)
|
| 1232 |
+
ap.add_argument("--determinism", choices=["strict", "warn", "off"], default="strict",
|
| 1233 |
+
help="strict (default): deterministic CUDA kernels, so a rerun "
|
| 1234 |
+
"of the same launch line reproduces the loss curve bit for "
|
| 1235 |
+
"bit; raises if an op has no deterministic implementation. "
|
| 1236 |
+
"warn: same, but non-deterministic ops warn instead of "
|
| 1237 |
+
"raising. off: fastest kernels, results drift between runs.")
|
| 1238 |
+
# checkpointing
|
| 1239 |
+
ap.add_argument("--no-checkpointing", action="store_true")
|
| 1240 |
+
ap.add_argument("--ckpt-dir", default=None)
|
| 1241 |
+
ap.add_argument("--ckpt-prefix", default=None)
|
| 1242 |
+
ap.add_argument("--ckpt-freq", type=int, default=None)
|
| 1243 |
+
ap.add_argument("--ckpt-keep", type=int, default=None)
|
| 1244 |
+
ap.add_argument("--resume-from", default=None,
|
| 1245 |
+
help="path to a specific *_step*.pt checkpoint to resume "
|
| 1246 |
+
"from, overriding auto-discovery of the latest "
|
| 1247 |
+
"checkpoint in --ckpt-dir")
|
| 1248 |
+
return ap
|
| 1249 |
+
|
| 1250 |
+
|
| 1251 |
+
def config_from_args(args):
|
| 1252 |
+
cfg = ModelConfig()
|
| 1253 |
+
for f in ("B", "T", "d_model", "n_layers", "q_heads", "kv_heads", "ffn_hidden",
|
| 1254 |
+
"vocab_size", "global_batch", "val_freq", "val_loss_steps",
|
| 1255 |
+
"ckpt_dir", "ckpt_prefix", "ckpt_freq", "ckpt_keep",
|
| 1256 |
+
"muon_lr", "muon_min_lr", "muon_momentum", "muon_ns_steps",
|
| 1257 |
+
"muon_adjust_lr"):
|
| 1258 |
+
v = getattr(args, f, None)
|
| 1259 |
+
if v is not None:
|
| 1260 |
+
setattr(cfg, f, v)
|
| 1261 |
+
# A --muon-lr given without a floor keeps the 10% ratio the WSD tail assumes,
|
| 1262 |
+
# instead of silently decaying to the default 2e-5 from wherever it starts.
|
| 1263 |
+
if args.muon_lr is not None and args.muon_min_lr is None:
|
| 1264 |
+
cfg.muon_min_lr = 0.1 * cfg.muon_lr
|
| 1265 |
+
cfg.fused_rms_rope = (args.rms_rope == "fused")
|
| 1266 |
+
cfg.activation_checkpoint = args.activation_checkpoint
|
| 1267 |
+
cfg.loss_chunks = args.loss_chunks
|
| 1268 |
+
cfg.checkpointing = not args.no_checkpointing
|
| 1269 |
+
cfg.context_length = max(cfg.context_length, cfg.T)
|
| 1270 |
+
|
| 1271 |
+
# Mirrors the BLU_MAX_STEPS branch in bluscript_single.cpp: shrink warmup and
|
| 1272 |
+
# stable with the run, or a short benchmark leaves no decay tail (the C++
|
| 1273 |
+
# asserts on it; get_lr here would clamp, which is worse — it would silently
|
| 1274 |
+
# report a schedule the config does not describe).
|
| 1275 |
+
if args.max_steps is not None:
|
| 1276 |
+
ms = args.max_steps
|
| 1277 |
+
cfg.max_steps = ms
|
| 1278 |
+
cfg.warmup_steps = max(1, ms // 10)
|
| 1279 |
+
cfg.stable_steps = max(1, ms // 2)
|
| 1280 |
+
if cfg.warmup_steps + cfg.stable_steps >= cfg.max_steps - 1:
|
| 1281 |
+
cfg.stable_steps = max(1, cfg.max_steps - cfg.warmup_steps - 2)
|
| 1282 |
+
print(f"[bench] max_steps={ms} (warmup={cfg.warmup_steps} "
|
| 1283 |
+
f"stable={cfg.stable_steps})")
|
| 1284 |
+
return cfg
|
| 1285 |
+
|
| 1286 |
+
|
| 1287 |
+
def validate_config(cfg, world_size=1, master=True):
|
| 1288 |
+
"""Cheap checks; each one is a silent-corruption or out-of-bounds failure
|
| 1289 |
+
mode if it does not hold."""
|
| 1290 |
+
bad = []
|
| 1291 |
+
if cfg.T > cfg.context_length:
|
| 1292 |
+
bad.append(f"T ({cfg.T}) > context_length ({cfg.context_length}): "
|
| 1293 |
+
"the RoPE cache is too short")
|
| 1294 |
+
if cfg.q_heads % cfg.kv_heads != 0:
|
| 1295 |
+
bad.append(f"q_heads ({cfg.q_heads}) not divisible by kv_heads ({cfg.kv_heads})")
|
| 1296 |
+
if cfg.head_dim not in (64, 128):
|
| 1297 |
+
bad.append(f"the fused kernels support head_dim 64 or 128, got {cfg.head_dim}")
|
| 1298 |
+
# Under DDP each rank runs B*T tokens per micro-step, so the world's tokens
|
| 1299 |
+
# per micro-step are B*T*world_size; global_batch must land on a whole
|
| 1300 |
+
# number of those or grad_accum below truncates and quietly under-shoots
|
| 1301 |
+
# the configured optimizer-step token count.
|
| 1302 |
+
if cfg.global_batch % (cfg.B * cfg.T * world_size) != 0:
|
| 1303 |
+
bad.append(f"global_batch ({cfg.global_batch}) not divisible by "
|
| 1304 |
+
f"B*T*world_size ({cfg.B * cfg.T * world_size})")
|
| 1305 |
+
if cfg.warmup_steps < 1:
|
| 1306 |
+
bad.append("warmup_steps must be >= 1 (get_lr divides by it)")
|
| 1307 |
+
if cfg.warmup_steps + cfg.stable_steps >= cfg.max_steps - 1:
|
| 1308 |
+
bad.append(f"warmup_steps + stable_steps "
|
| 1309 |
+
f"({cfg.warmup_steps + cfg.stable_steps}) leaves no decay tail "
|
| 1310 |
+
f"before max_steps ({cfg.max_steps})")
|
| 1311 |
+
if not hasattr(torch.optim, "Muon"):
|
| 1312 |
+
bad.append(f"torch {torch.__version__} has no torch.optim.Muon — it first "
|
| 1313 |
+
"appears in the 2.9 docs. Upgrade torch, or run the AdamW "
|
| 1314 |
+
"arm (train.py).")
|
| 1315 |
+
if cfg.muon_ns_steps < 1:
|
| 1316 |
+
bad.append("muon_ns_steps must be >= 1 (0 leaves the update unorthogonalized, "
|
| 1317 |
+
"which is not Muon)")
|
| 1318 |
+
if cfg.muon_momentum <= 0.0 or cfg.muon_momentum >= 1.0:
|
| 1319 |
+
bad.append(f"muon_momentum must be in (0,1), got {cfg.muon_momentum}")
|
| 1320 |
+
if cfg.q_heads * cfg.head_dim != cfg.d_model and master:
|
| 1321 |
+
print(f"WARNING: q_heads*head_dim ({cfg.q_heads * cfg.head_dim}) "
|
| 1322 |
+
f"!= d_model ({cfg.d_model})")
|
| 1323 |
+
if bad:
|
| 1324 |
+
raise SystemExit("ERROR: " + "\nERROR: ".join(bad))
|
| 1325 |
+
|
| 1326 |
+
|
| 1327 |
+
def main():
|
| 1328 |
+
args = build_argparser().parse_args()
|
| 1329 |
+
|
| 1330 |
+
# Env overrides, for parity with the C++ launch line. An explicit CLI flag
|
| 1331 |
+
# wins over the environment.
|
| 1332 |
+
if args.B is None:
|
| 1333 |
+
args.B = _env_int("BLU_B", 0) or None
|
| 1334 |
+
if args.global_batch is None:
|
| 1335 |
+
args.global_batch = _env_int("BLU_GLOBAL_BATCH", 0) or None
|
| 1336 |
+
if args.max_steps is None:
|
| 1337 |
+
args.max_steps = _env_int("BLU_MAX_STEPS", 0) or None
|
| 1338 |
+
diag = args.diag or _env_int("BLU_DIAG", 0) != 0
|
| 1339 |
+
|
| 1340 |
+
if not torch.cuda.is_available():
|
| 1341 |
+
raise SystemExit("ERROR: no CUDA device available")
|
| 1342 |
+
|
| 1343 |
+
# ---- DDP bootstrap ----
|
| 1344 |
+
# torchrun sets RANK/LOCAL_RANK/WORLD_SIZE; their absence means a plain
|
| 1345 |
+
# single-process launch, where --device / BLU_DEVICE pick the GPU instead.
|
| 1346 |
+
# torch.cuda.set_device MUST happen before init_process_group: NCCL creates
|
| 1347 |
+
# its context on whatever device is current at that call, so doing it after
|
| 1348 |
+
# would leave every rank on this node contending for device 0.
|
| 1349 |
+
ddp = int(os.environ.get("RANK", -1)) != -1
|
| 1350 |
+
if ddp:
|
| 1351 |
+
ddp_rank = int(os.environ["RANK"])
|
| 1352 |
+
ddp_local_rank = int(os.environ["LOCAL_RANK"])
|
| 1353 |
+
ddp_world_size = int(os.environ["WORLD_SIZE"])
|
| 1354 |
+
master_process = ddp_rank == 0
|
| 1355 |
+
if ddp_local_rank >= torch.cuda.device_count():
|
| 1356 |
+
raise SystemExit(f"ERROR: LOCAL_RANK={ddp_local_rank} but only "
|
| 1357 |
+
f"{torch.cuda.device_count()} CUDA device(s) visible")
|
| 1358 |
+
torch.cuda.set_device(ddp_local_rank)
|
| 1359 |
+
device = torch.device("cuda", ddp_local_rank)
|
| 1360 |
+
dist.init_process_group(backend="nccl")
|
| 1361 |
+
if master_process and (args.device is not None
|
| 1362 |
+
or os.environ.get("BLU_DEVICE") not in (None, "")):
|
| 1363 |
+
print("[WARN] --device / BLU_DEVICE is ignored under torchrun — "
|
| 1364 |
+
"LOCAL_RANK selects each rank's GPU.")
|
| 1365 |
+
else:
|
| 1366 |
+
ddp_rank, ddp_local_rank, ddp_world_size = 0, 0, 1
|
| 1367 |
+
master_process = True
|
| 1368 |
+
if args.device is None:
|
| 1369 |
+
args.device = _env_int("BLU_DEVICE", 0)
|
| 1370 |
+
if args.device >= torch.cuda.device_count():
|
| 1371 |
+
raise SystemExit(f"ERROR: BLU_DEVICE={args.device} but only "
|
| 1372 |
+
f"{torch.cuda.device_count()} CUDA device(s) visible")
|
| 1373 |
+
torch.cuda.set_device(args.device)
|
| 1374 |
+
device = torch.device("cuda", args.device)
|
| 1375 |
+
|
| 1376 |
+
seed_everything(args.seed)
|
| 1377 |
+
apply_determinism(args.determinism)
|
| 1378 |
+
# TF32 loses mantissa bits but is deterministic, so it is orthogonal to the
|
| 1379 |
+
# setting above and stays on.
|
| 1380 |
+
torch.set_float32_matmul_precision("high")
|
| 1381 |
+
|
| 1382 |
+
cfg = config_from_args(args)
|
| 1383 |
+
if cfg.fused_rms_rope and not _HAS_TRITON:
|
| 1384 |
+
if master_process:
|
| 1385 |
+
print("[WARN] Triton not available — falling back to the eager QK-norm+RoPE")
|
| 1386 |
+
cfg.fused_rms_rope = False
|
| 1387 |
+
validate_config(cfg, ddp_world_size, master_process)
|
| 1388 |
+
|
| 1389 |
+
# Tokens consumed per optimizer step = B * T * grad_accum * world_size.
|
| 1390 |
+
grad_accum = cfg.global_batch // (cfg.B * cfg.T * ddp_world_size)
|
| 1391 |
+
tokens_per_step = cfg.B * cfg.T * grad_accum * ddp_world_size
|
| 1392 |
+
|
| 1393 |
+
if master_process:
|
| 1394 |
+
print("================ Configuration ================")
|
| 1395 |
+
mode_str = (f"DDP x{ddp_world_size} (torchrun, NCCL)" if ddp
|
| 1396 |
+
else "SINGLE GPU (no torchrun / DDP / ZeRO)")
|
| 1397 |
+
print(f" mode : {mode_str}")
|
| 1398 |
+
print(f" cuda device : {ddp_local_rank if ddp else args.device} "
|
| 1399 |
+
f"({torch.cuda.get_device_name(device)})")
|
| 1400 |
+
print(f" d_model : {cfg.d_model}")
|
| 1401 |
+
print(f" n_layers : {cfg.n_layers}")
|
| 1402 |
+
print(f" q_heads/kv_heads: {cfg.q_heads} / {cfg.kv_heads} "
|
| 1403 |
+
f"(G={cfg.q_heads // cfg.kv_heads})")
|
| 1404 |
+
print(f" head_dim : {cfg.head_dim}")
|
| 1405 |
+
print(f" ffn_hidden : {cfg.ffn_hidden} (SwiGLU)")
|
| 1406 |
+
print(f" context_length : {cfg.context_length}")
|
| 1407 |
+
print(f" rope_theta : {cfg.rope_theta:g} (NeoX half-split)")
|
| 1408 |
+
print(f" vocab_size : {cfg.vocab_size}")
|
| 1409 |
+
print(f" weight_tying : {'ON' if cfg.weight_tying else 'OFF'}")
|
| 1410 |
+
print(f" B x T : {cfg.B} x {cfg.T}")
|
| 1411 |
+
print(f" global_batch : {cfg.global_batch} (grad_accum={grad_accum}, "
|
| 1412 |
+
f"world_size={ddp_world_size})")
|
| 1413 |
+
print(" optimizer : MUON (hidden 2D) + AdamW (embed/head/1D)")
|
| 1414 |
+
print(f" max_lr / min_lr : {cfg.max_lr:g} / {cfg.min_lr:g} (AdamW group)")
|
| 1415 |
+
print(f" muon lr / min : {cfg.muon_lr:g} / {cfg.muon_min_lr:g} "
|
| 1416 |
+
f"(adjust_lr_fn={cfg.muon_adjust_lr})")
|
| 1417 |
+
print(f" muon momentum : {cfg.muon_momentum:g} | NS steps {cfg.muon_ns_steps}")
|
| 1418 |
+
print(f" warmup / steps : {cfg.warmup_steps} / {cfg.max_steps}")
|
| 1419 |
+
print(f" WSD stable_steps: {cfg.stable_steps} "
|
| 1420 |
+
f"(decay tail={cfg.max_steps - cfg.warmup_steps - cfg.stable_steps})")
|
| 1421 |
+
print(f" precision : {args.precision}")
|
| 1422 |
+
print(f" seed / determ. : {args.seed} / {args.determinism}")
|
| 1423 |
+
print(f" rms+rope : {'fused' if cfg.fused_rms_rope else 'eager'}"
|
| 1424 |
+
f" | compile {args.compile} | act-ckpt {cfg.activation_checkpoint}"
|
| 1425 |
+
f" | loss_chunks {cfg.loss_chunks}")
|
| 1426 |
+
print("===============================================", flush=True)
|
| 1427 |
+
|
| 1428 |
+
# ---- model ----
|
| 1429 |
+
# Built and seeded identically on every rank (same seed, deterministic CPU
|
| 1430 |
+
# init); DDP's constructor broadcasts rank 0's parameters/buffers anyway,
|
| 1431 |
+
# so this is belt-and-suspenders against any future source of divergence.
|
| 1432 |
+
if master_process:
|
| 1433 |
+
print(f"Initializing model on CUDA device {device.index}...", flush=True)
|
| 1434 |
+
raw_model = BLU(cfg).to(device)
|
| 1435 |
+
raw_model.apply_param_dtype(args.precision)
|
| 1436 |
+
params = list(raw_model.parameters())
|
| 1437 |
+
n_params = sum(p.numel() for p in params)
|
| 1438 |
+
if master_process:
|
| 1439 |
+
print(f"Number of parameters: {n_params}")
|
| 1440 |
+
if args.compile:
|
| 1441 |
+
for blk in raw_model.blocks:
|
| 1442 |
+
blk.compile()
|
| 1443 |
+
|
| 1444 |
+
# DDP wraps the (already dtype-cast, already compiled) raw module. Buffer
|
| 1445 |
+
# broadcasting is off: rope_cos/rope_sin are deterministic functions of
|
| 1446 |
+
# cfg, identical on every rank already, so syncing them every forward
|
| 1447 |
+
# would just be wasted communication.
|
| 1448 |
+
# if ddp:
|
| 1449 |
+
# model = DDP(raw_model, device_ids=[ddp_local_rank], broadcast_buffers=False)
|
| 1450 |
+
# else:
|
| 1451 |
+
# model = raw_model
|
| 1452 |
+
|
| 1453 |
+
if ddp:
|
| 1454 |
+
model = DDP(raw_model, device_ids=[ddp_local_rank], broadcast_buffers=False)
|
| 1455 |
+
torch.cuda.synchronize() # let DDP's constructor-time NCCL broadcast fully land
|
| 1456 |
+
# before DataLoaderLite's independent copy_stream starts
|
| 1457 |
+
# issuing its own async H2D work — see [crash investigation]
|
| 1458 |
+
else:
|
| 1459 |
+
model = raw_model
|
| 1460 |
+
|
| 1461 |
+
|
| 1462 |
+
# ---- optimizer ----
|
| 1463 |
+
# named_parameters(), not parameters(): the routing is by name (the
|
| 1464 |
+
# embedding and lm_head stay on AdamW). Duplicates are dropped by default,
|
| 1465 |
+
# so a weight-tied head is visited once, as `wte.weight`. Built from
|
| 1466 |
+
# raw_model so names never carry DDP's "module." prefix.
|
| 1467 |
+
opt = MasterMuon(raw_model.named_parameters(), cfg)
|
| 1468 |
+
if master_process:
|
| 1469 |
+
print(f"muon params: {opt.n_muon} | adamw decayed: {opt.n_decay}, "
|
| 1470 |
+
f"non-decayed: {opt.n_nodecay}")
|
| 1471 |
+
# Static footprint. Nothing here is sharded — plain DDP replicates the
|
| 1472 |
+
# full optimizer state on every rank (the whole cost of not running
|
| 1473 |
+
# ZeRO-1), so this is the per-GPU charge before a single activation is
|
| 1474 |
+
# allocated.
|
| 1475 |
+
gb = 2 ** 30
|
| 1476 |
+
param_bytes = sum(p.numel() * p.element_size() for p in params)
|
| 1477 |
+
grad_bytes = param_bytes # grads share the param dtype
|
| 1478 |
+
master_bytes, landing_bytes, _, _ = opt.state_breakdown()
|
| 1479 |
+
# Both optimizers allocate their state on the first step(), so at this point
|
| 1480 |
+
# the measured values are 0 and the report has to predict them.
|
| 1481 |
+
mv_bytes = opt.expected_mv_bytes()
|
| 1482 |
+
mom_bytes = opt.expected_muon_mom_bytes()
|
| 1483 |
+
static_bytes = (param_bytes + grad_bytes + master_bytes + landing_bytes
|
| 1484 |
+
+ mv_bytes + mom_bytes)
|
| 1485 |
+
if master_process:
|
| 1486 |
+
print(f"[{'ddp' if ddp else 'single'}] static footprint (per GPU: "
|
| 1487 |
+
f"params + grads + optimizer state)")
|
| 1488 |
+
print(f" parameters : {param_bytes / gb:7.2f} GiB")
|
| 1489 |
+
print(f" gradients : {grad_bytes / gb:7.2f} GiB")
|
| 1490 |
+
print(f" fp32 master weights : {master_bytes / gb:7.2f} GiB")
|
| 1491 |
+
print(f" fp32 grad landing bufs : {landing_bytes / gb:7.2f} GiB")
|
| 1492 |
+
print(f" Muon momentum : {mom_bytes / gb:7.2f} GiB")
|
| 1493 |
+
print(f" Adam exp_avg + exp_avg_sq: {mv_bytes / gb:6.2f} GiB")
|
| 1494 |
+
print(f" static total : {static_bytes / gb:7.2f} GiB", flush=True)
|
| 1495 |
+
|
| 1496 |
+
# ---- data ----
|
| 1497 |
+
# Sharded by rank: each rank reads a disjoint, interleaved slice of the
|
| 1498 |
+
# same shard list (see DataLoaderLite), so DDP's cross-rank gradient
|
| 1499 |
+
# average is over world_size DIFFERENT windows, not world_size copies of
|
| 1500 |
+
# the same one.
|
| 1501 |
+
train_loader = DataLoaderLite(cfg.B, cfg.T, "train", args.data_root, device,
|
| 1502 |
+
master=master_process, rank=ddp_rank,
|
| 1503 |
+
world_size=ddp_world_size)
|
| 1504 |
+
# rank=0, world_size=1 UNCONDITIONALLY, on every rank: val_loss_steps
|
| 1505 |
+
# windows at fixed token offsets 0, BT, 2BT, ... Sharding this by
|
| 1506 |
+
# ddp_rank/ddp_world_size like the train loader would make the reported
|
| 1507 |
+
# val_loss a function of how many GPUs happen to be running (world_size=2
|
| 1508 |
+
# averages 2x as many, differently-offset windows as world_size=1) —
|
| 1509 |
+
# the two numbers would not be comparable across a GPU-count change, which
|
| 1510 |
+
# is what produced the apparent val_loss jump on the 2-GPU -> 1-GPU
|
| 1511 |
+
# resume. Every rank now redundantly evaluates the identical fixed
|
| 1512 |
+
# windows, so the metric means the same thing regardless of world_size.
|
| 1513 |
+
val_loader = DataLoaderLite(cfg.B, cfg.T, "val", args.data_root, device,
|
| 1514 |
+
master=master_process, rank=0, world_size=1)
|
| 1515 |
+
|
| 1516 |
+
# ---- checkpointing: manager + auto-resume ----
|
| 1517 |
+
# Every rank independently loads the same checkpoint file from the same
|
| 1518 |
+
# (assumed shared) filesystem — a read, so no race — which keeps model,
|
| 1519 |
+
# optimizer state, and start_step identical across ranks without an
|
| 1520 |
+
# explicit broadcast. Only rank 0 ever WRITES one (below).
|
| 1521 |
+
start_step, ckpt = 0, None
|
| 1522 |
+
last_val = float("nan")
|
| 1523 |
+
if cfg.checkpointing or args.resume_from:
|
| 1524 |
+
ckpt = CheckpointManager(cfg.ckpt_dir, cfg.ckpt_prefix, cfg.ckpt_keep)
|
| 1525 |
+
if args.resume_from:
|
| 1526 |
+
if master_process:
|
| 1527 |
+
print(f"[resume] loading explicit checkpoint: {args.resume_from}")
|
| 1528 |
+
start_step, resume_loss, last_val = ckpt.load(
|
| 1529 |
+
args.resume_from, raw_model, opt, cfg, master_process)
|
| 1530 |
+
else:
|
| 1531 |
+
start_step, resume_loss, last_val = ckpt.load_latest(
|
| 1532 |
+
raw_model, opt, cfg, master_process)
|
| 1533 |
+
if start_step > 0:
|
| 1534 |
+
if master_process:
|
| 1535 |
+
print(f"[resume] at step {start_step} (loss {resume_loss:.6f}, "
|
| 1536 |
+
f"last val {last_val:.6f})")
|
| 1537 |
+
# Skip the batches already consumed by steps [0, start_step).
|
| 1538 |
+
# global_batch // (B*T*world_size) makes tokens-per-step (and so
|
| 1539 |
+
# this target offset) independent of world_size, so a GPU-count
|
| 1540 |
+
# change alone lands on the correct byte offset as long as B, T,
|
| 1541 |
+
# and global_batch are unchanged (checked above).
|
| 1542 |
+
train_loader.skip_batches(start_step * grad_accum)
|
| 1543 |
+
|
| 1544 |
+
# ---- CSV log ----
|
| 1545 |
+
# Appended (not truncated) when resuming so a restart does not erase the
|
| 1546 |
+
# history of the run it is continuing. Rank 0 only — every rank would
|
| 1547 |
+
# otherwise open and write the same path.
|
| 1548 |
+
log_file = None
|
| 1549 |
+
if args.csv and master_process:
|
| 1550 |
+
resuming = start_step > 0 and os.path.exists(args.csv)
|
| 1551 |
+
log_file = open(args.csv, "a" if resuming else "w")
|
| 1552 |
+
if not resuming:
|
| 1553 |
+
log_file.write("step,train_loss,val_loss,tok_per_sec,grad_norm,dt_ms,lr,"
|
| 1554 |
+
"elapsed_min,muon_lr\n")
|
| 1555 |
+
log_file.flush()
|
| 1556 |
+
|
| 1557 |
+
amp = torch.autocast(device_type="cuda", dtype=torch.bfloat16)
|
| 1558 |
+
probes = [L for L in cfg.probe_layers if L < cfg.n_layers]
|
| 1559 |
+
steady_dts = []
|
| 1560 |
+
t_train_start = time.perf_counter()
|
| 1561 |
+
|
| 1562 |
+
# ---- Ctrl-C: finish the in-flight step, save a checkpoint, exit clean ----
|
| 1563 |
+
# Registered here (not earlier) so a Ctrl-C during setup/data-loading still
|
| 1564 |
+
# gets Python's default immediate-abort behaviour — only the training loop
|
| 1565 |
+
# below gets the graceful path. A second Ctrl-C re-arms the default handler
|
| 1566 |
+
# and re-raises, for when the checkpoint save itself needs interrupting.
|
| 1567 |
+
interrupted = False
|
| 1568 |
+
|
| 1569 |
+
def _handle_sigint(signum, frame):
|
| 1570 |
+
nonlocal interrupted
|
| 1571 |
+
if interrupted:
|
| 1572 |
+
signal.signal(signal.SIGINT, signal.SIG_DFL)
|
| 1573 |
+
raise KeyboardInterrupt
|
| 1574 |
+
interrupted = True
|
| 1575 |
+
if master_process:
|
| 1576 |
+
print("\n[ctrl-c] finishing this step, then checkpointing and exiting "
|
| 1577 |
+
"(press again to abort immediately)...", flush=True)
|
| 1578 |
+
|
| 1579 |
+
# A terminal Ctrl-C is delivered to every process in the foreground
|
| 1580 |
+
# process group, so under torchrun every rank sets `interrupted` at
|
| 1581 |
+
# roughly the same time and the ranks stay in lockstep for the collective
|
| 1582 |
+
# ops (DDP's grad all-reduce, the checkpoint barrier below) still ahead.
|
| 1583 |
+
signal.signal(signal.SIGINT, _handle_sigint)
|
| 1584 |
+
|
| 1585 |
+
if master_process:
|
| 1586 |
+
print("\nStarting training" +
|
| 1587 |
+
(f" (resumed at step {start_step})" if start_step else "") + "...\n", flush=True)
|
| 1588 |
+
|
| 1589 |
+
for step in range(start_step, cfg.max_steps):
|
| 1590 |
+
# ---------- periodic validation ----------
|
| 1591 |
+
if cfg.val_loss_steps > 0 and (step % cfg.val_freq == 0 or step == cfg.max_steps - 1):
|
| 1592 |
+
model.eval()
|
| 1593 |
+
val_loader.reset()
|
| 1594 |
+
acc = 0.0
|
| 1595 |
+
with torch.no_grad(), amp:
|
| 1596 |
+
for _ in range(cfg.val_loss_steps):
|
| 1597 |
+
x, y, _, _ = val_loader.next_batch()
|
| 1598 |
+
acc += model(x, y)[1].item() / cfg.val_loss_steps
|
| 1599 |
+
model.train()
|
| 1600 |
+
# val_loader is rank=0/world_size=1 on every rank (see its
|
| 1601 |
+
# construction above), so every rank just computed the identical
|
| 1602 |
+
# fixed windows redundantly — no cross-rank reduction needed, and
|
| 1603 |
+
# val_loss no longer depends on how many ranks are running.
|
| 1604 |
+
last_val = acc
|
| 1605 |
+
if master_process:
|
| 1606 |
+
print(f" [val] step {step} | val_loss {last_val:.6f}", flush=True)
|
| 1607 |
+
|
| 1608 |
+
# ---------- training step ----------
|
| 1609 |
+
t0 = time.perf_counter()
|
| 1610 |
+
model.zero_grad(set_to_none=True)
|
| 1611 |
+
# Fresh zeros rather than an in-place scale of last step's accumulator:
|
| 1612 |
+
# multiplying by zero does NOT clear a NaN, so a single non-finite
|
| 1613 |
+
# micro-batch would poison the accumulator for the rest of the run.
|
| 1614 |
+
loss_accum = torch.zeros((), device=device)
|
| 1615 |
+
micro_loss, micro_where = [], []
|
| 1616 |
+
|
| 1617 |
+
for micro in range(grad_accum):
|
| 1618 |
+
x, y, shard, offset = train_loader.next_batch()
|
| 1619 |
+
# DDP allreduces gradients on every .backward(); for grad_accum > 1
|
| 1620 |
+
# only the LAST micro-batch of the step should pay for that — the
|
| 1621 |
+
# others accumulate purely locally under no_sync().
|
| 1622 |
+
sync_ctx = (nullcontext() if (not ddp or micro == grad_accum - 1)
|
| 1623 |
+
else model.no_sync())
|
| 1624 |
+
with sync_ctx:
|
| 1625 |
+
with amp:
|
| 1626 |
+
_, loss = model(x, y)
|
| 1627 |
+
loss_accum += loss.detach()
|
| 1628 |
+
# Per-micro-batch loss (--diag). The reported train loss is the mean
|
| 1629 |
+
# over grad_accum micro-steps, so a single pathological micro-batch
|
| 1630 |
+
# is diluted while its gradient dominates the accumulated norm
|
| 1631 |
+
# almost undiluted — which is how a modest loss bump can sit next to
|
| 1632 |
+
# a large norm spike. `micro_where` names the shard and token offset
|
| 1633 |
+
# so the responsible bytes can be decoded offline.
|
| 1634 |
+
#
|
| 1635 |
+
# Costs a device sync per micro-step, so it is off by default.
|
| 1636 |
+
if diag:
|
| 1637 |
+
micro_loss.append(loss.item())
|
| 1638 |
+
micro_where.append((shard, offset))
|
| 1639 |
+
(loss / grad_accum).backward()
|
| 1640 |
+
|
| 1641 |
+
if ddp:
|
| 1642 |
+
# Average each rank's accumulated (summed) loss BEFORE dividing by
|
| 1643 |
+
# grad_accum, so loss_val is the mean over world_size*grad_accum
|
| 1644 |
+
# micro-batches — the true loss behind this optimizer step, not
|
| 1645 |
+
# just rank 0's slice. Also doubles as a sync point: a NaN on any
|
| 1646 |
+
# rank poisons the average for every rank, so the check below
|
| 1647 |
+
# raises everywhere instead of only where it originated (which
|
| 1648 |
+
# would otherwise hang the survivors on the next collective).
|
| 1649 |
+
dist.all_reduce(loss_accum, op=dist.ReduceOp.AVG)
|
| 1650 |
+
loss_val = loss_accum.item() / grad_accum
|
| 1651 |
+
if math.isnan(loss_val) or math.isinf(loss_val):
|
| 1652 |
+
if log_file:
|
| 1653 |
+
log_file.close()
|
| 1654 |
+
# Every rank reaches this identically (the all_reduce above is a
|
| 1655 |
+
# sync point and NaN propagates through AVG), so every rank exits
|
| 1656 |
+
# here together rather than the survivors hanging on the next
|
| 1657 |
+
# collective.
|
| 1658 |
+
if ddp:
|
| 1659 |
+
dist.destroy_process_group()
|
| 1660 |
+
raise SystemExit(f"ERROR: NaN/Inf loss at step {step}")
|
| 1661 |
+
|
| 1662 |
+
# Depth probe (--diag), read BEFORE the update, so these are the
|
| 1663 |
+
# pre-update gradients. Compare the profile on a spiking step against a
|
| 1664 |
+
# quiet one: a blow-up that begins mid-stack and grows toward the loss
|
| 1665 |
+
# points at the forward, while a flat multiplicative lift across all
|
| 1666 |
+
# depths points at the optimizer/update.
|
| 1667 |
+
depth = []
|
| 1668 |
+
if diag:
|
| 1669 |
+
for L in probes:
|
| 1670 |
+
a = raw_model.blocks[L].attn
|
| 1671 |
+
gn = a.wQ.weight.grad
|
| 1672 |
+
depth.append((L,
|
| 1673 |
+
gn.float().norm().item() if gn is not None else -1.0,
|
| 1674 |
+
a.q_gamma.abs().max().item(),
|
| 1675 |
+
a.wQ.weight.abs().max().float().item()))
|
| 1676 |
+
|
| 1677 |
+
grad_norm = clip_grads_(params, cfg.grad_clip).item()
|
| 1678 |
+
lr = get_lr(step, cfg) # AdamW group
|
| 1679 |
+
muon_lr = get_lr(step, cfg, cfg.muon_lr, cfg.muon_min_lr)
|
| 1680 |
+
opt.step(lr, muon_lr)
|
| 1681 |
+
|
| 1682 |
+
torch.cuda.synchronize() # make the per-step timing reflect completed GPU work
|
| 1683 |
+
dt = time.perf_counter() - t0
|
| 1684 |
+
tok_per_sec = tokens_per_step / dt
|
| 1685 |
+
elapsed_min = (time.perf_counter() - t_train_start) / 60.0
|
| 1686 |
+
if step - start_step >= args.warmup_measure:
|
| 1687 |
+
steady_dts.append(dt)
|
| 1688 |
+
|
| 1689 |
+
# ---------- per-step report (rank 0 only — every rank's numbers agree
|
| 1690 |
+
# on loss/grad_norm post-allreduce, so printing from every rank would
|
| 1691 |
+
# only duplicate lines, not add information) ----------
|
| 1692 |
+
if master_process:
|
| 1693 |
+
vs = " - " if last_val != last_val else f"{last_val:.6f}"
|
| 1694 |
+
print(f"step {step:6d} | train {loss_val:.6f} | val {vs} "
|
| 1695 |
+
f"| tok/s {tok_per_sec:10.2f} | norm {grad_norm:.4f} "
|
| 1696 |
+
f"| dt {dt * 1e3:8.2f} ms | lr {lr:.4e} | mlr {muon_lr:.4e} "
|
| 1697 |
+
f"| mem {torch.cuda.max_memory_allocated() / gb:5.2f}/"
|
| 1698 |
+
f"{torch.cuda.max_memory_reserved() / gb:5.2f} GiB", flush=True)
|
| 1699 |
+
|
| 1700 |
+
if diag and micro_loss:
|
| 1701 |
+
# Flag micro-batches that stand out against this step's own median.
|
| 1702 |
+
med = sorted(micro_loss)[len(micro_loss) // 2]
|
| 1703 |
+
marks = " ".join(("*" if m > med + 0.5 else "") + f"{m:.3f}"
|
| 1704 |
+
for m in micro_loss)
|
| 1705 |
+
print(f" [micro] med {med:.4f} | {marks}")
|
| 1706 |
+
for m, (ml, (shard, offset)) in enumerate(zip(micro_loss, micro_where)):
|
| 1707 |
+
if ml > med + 0.5:
|
| 1708 |
+
print(f" [micro] OUTLIER m={m} loss {ml:.4f} (median {med:.4f}) "
|
| 1709 |
+
f"from shard {shard} token offset {offset}")
|
| 1710 |
+
if diag and depth:
|
| 1711 |
+
print(" [depth] ||g(wQ)|| by layer:"
|
| 1712 |
+
+ "".join(f" L{L} {g:.3e}" for L, g, _, _ in depth)
|
| 1713 |
+
+ " | max|q_gamma|:"
|
| 1714 |
+
+ "".join(f" L{L} {gm:.4f}" for L, _, gm, _ in depth)
|
| 1715 |
+
+ " | max|wQ|:"
|
| 1716 |
+
+ "".join(f" L{L} {w:.4f}" for L, _, _, w in depth), flush=True)
|
| 1717 |
+
|
| 1718 |
+
if log_file:
|
| 1719 |
+
log_file.write(f"{step},{loss_val:.6f},{last_val:.6f},{tok_per_sec:.6f},"
|
| 1720 |
+
f"{grad_norm:.6f},{dt * 1e3:.6f},{lr:.6f},{elapsed_min:.6f},"
|
| 1721 |
+
f"{muon_lr:.6f}\n")
|
| 1722 |
+
log_file.flush()
|
| 1723 |
+
|
| 1724 |
+
if ddp:
|
| 1725 |
+
# `interrupted` is set by this process's own SIGINT handler, so in
|
| 1726 |
+
# principle one rank could see it before another (a signal isn't
|
| 1727 |
+
# guaranteed simultaneous delivery, even though in practice a
|
| 1728 |
+
# terminal Ctrl-C hits every rank in the same process group at
|
| 1729 |
+
# once). Left asymmetric, it would desync the checkpoint barrier
|
| 1730 |
+
# and the `break` below against ranks still issuing DDP's regular
|
| 1731 |
+
# per-step collectives — a hang, not just a wrong answer. This
|
| 1732 |
+
# runs unconditionally every step (never gated on a per-rank
|
| 1733 |
+
# value), so it can never itself mismatch across ranks.
|
| 1734 |
+
flag = torch.tensor(1.0 if interrupted else 0.0, device=device)
|
| 1735 |
+
dist.all_reduce(flag, op=dist.ReduceOp.MAX)
|
| 1736 |
+
interrupted = flag.item() > 0
|
| 1737 |
+
|
| 1738 |
+
# ---------- periodic checkpoint (also fires once on Ctrl-C) ----------
|
| 1739 |
+
# Labeled step+1 (the next iteration to run) so load_latest resumes
|
| 1740 |
+
# exactly where this save left off — no duplicated or skipped iteration.
|
| 1741 |
+
# Captures weights + Adam m/v AFTER this step's update. Rank 0 only —
|
| 1742 |
+
# every rank's raw_model/opt state is identical post-allreduce, so a
|
| 1743 |
+
# second writer would only race the first on the same path. The
|
| 1744 |
+
# barrier holds every rank at this point until the save (or its
|
| 1745 |
+
# failure) is done, so a slow write can't let other ranks race ahead
|
| 1746 |
+
# into the next step's collectives before the checkpoint is durable.
|
| 1747 |
+
if cfg.checkpointing and ((step + 1) % cfg.ckpt_freq == 0
|
| 1748 |
+
or step == cfg.max_steps - 1
|
| 1749 |
+
or interrupted):
|
| 1750 |
+
if master_process:
|
| 1751 |
+
try:
|
| 1752 |
+
ckpt.save(step + 1, raw_model, opt, loss_val, last_val, cfg)
|
| 1753 |
+
except Exception as e: # noqa: BLE001 — advisory
|
| 1754 |
+
print(f"[WARN] checkpoint save at step {step + 1} failed: {e}")
|
| 1755 |
+
if ddp:
|
| 1756 |
+
dist.barrier()
|
| 1757 |
+
|
| 1758 |
+
if interrupted:
|
| 1759 |
+
if not cfg.checkpointing and master_process:
|
| 1760 |
+
print(f"[ctrl-c] checkpointing is disabled (--no-checkpointing); "
|
| 1761 |
+
f"stopping at step {step} without saving.", flush=True)
|
| 1762 |
+
break
|
| 1763 |
+
|
| 1764 |
+
# Reclaim the extra allocator high-water mark from the first couple of
|
| 1765 |
+
# steps' one-time warmup allocations (autograd graph caches, etc.).
|
| 1766 |
+
if step in (start_step, start_step + 1):
|
| 1767 |
+
torch.cuda.synchronize()
|
| 1768 |
+
torch.cuda.empty_cache()
|
| 1769 |
+
torch.cuda.reset_peak_memory_stats()
|
| 1770 |
+
|
| 1771 |
+
torch.cuda.synchronize()
|
| 1772 |
+
if log_file:
|
| 1773 |
+
log_file.close()
|
| 1774 |
+
|
| 1775 |
+
# ---- steady-state summary ----
|
| 1776 |
+
# dt/throughput are this rank's own timing; under DDP all ranks run in
|
| 1777 |
+
# lockstep (grad sync + the checkpoint barrier), so rank 0's numbers are
|
| 1778 |
+
# representative and `tokens_per_step` already counts every rank's tokens.
|
| 1779 |
+
if steady_dts and master_process:
|
| 1780 |
+
xs = sorted(steady_dts)
|
| 1781 |
+
med = xs[len(xs) // 2]
|
| 1782 |
+
mean = sum(xs) / len(xs)
|
| 1783 |
+
alloc = torch.cuda.max_memory_allocated() / gb
|
| 1784 |
+
reserved = torch.cuda.max_memory_reserved() / gb
|
| 1785 |
+
# Measured now that Adam's m/v exist, rather than the startup estimate.
|
| 1786 |
+
static = (param_bytes + grad_bytes + opt.state_bytes()) / gb
|
| 1787 |
+
mode_str = f"DDP x{ddp_world_size}" if ddp else "single GPU"
|
| 1788 |
+
print("\n" + "=" * 78)
|
| 1789 |
+
print(f" STEADY-STATE SUMMARY — {mode_str} optimizer=muon+adamw "
|
| 1790 |
+
f"precision={args.precision} "
|
| 1791 |
+
f"rms+rope={'fused' if cfg.fused_rms_rope else 'eager'}")
|
| 1792 |
+
print(f" torch.optim.Muon lr={cfg.muon_lr:g} "
|
| 1793 |
+
f"adjust_lr_fn={cfg.muon_adjust_lr} "
|
| 1794 |
+
f"momentum={cfg.muon_momentum:g} ns_steps={cfg.muon_ns_steps}")
|
| 1795 |
+
print(f" ({len(xs)} steps, first {args.warmup_measure} excluded)")
|
| 1796 |
+
print("=" * 78)
|
| 1797 |
+
print(f" step time : {med * 1e3:10.2f} ms (mean {mean * 1e3:.2f})")
|
| 1798 |
+
print(f" throughput : {tokens_per_step / med:10.1f} tok/s")
|
| 1799 |
+
print(f" tokens / step : {tokens_per_step:10d}")
|
| 1800 |
+
print(f" static (p+g+opt) : {static:10.2f} GiB (per GPU)")
|
| 1801 |
+
print(f" activations (peak) : {max(alloc - static, 0.0):10.2f} GiB")
|
| 1802 |
+
print(f" peak mem allocated : {alloc:10.2f} GiB")
|
| 1803 |
+
print(f" peak mem reserved : {reserved:10.2f} GiB")
|
| 1804 |
+
print("=" * 78)
|
| 1805 |
+
|
| 1806 |
+
if ddp:
|
| 1807 |
+
dist.destroy_process_group()
|
| 1808 |
+
if master_process:
|
| 1809 |
+
print("\n=== Training complete ===")
|
| 1810 |
+
|
| 1811 |
+
|
| 1812 |
+
if __name__ == "__main__":
|
| 1813 |
+
main()
|
| 1814 |
+
|
| 1815 |
+
|
| 1816 |
+
|
| 1817 |
+
|
| 1818 |
+
|
| 1819 |
+
|
| 1820 |
+
|
| 1821 |
+
|
| 1822 |
+
|
| 1823 |
+
|
| 1824 |
+
|
| 1825 |
+
|
| 1826 |
+
|
| 1827 |
+
|
| 1828 |
+
|
| 1829 |
+
|