vh
7b5fd91d3c
Run-01 was killed at step 19 by operator instruction to root-cause before spending a ~13.9h window. Two independent methods now agree on where the step time went, and neither was the hypothesis the consult panel converged on. Scaling fit (3 points, 2 params, residuals <3ms over an 8x range): A = 6.87e-4 s/token, B = 8.85e-8 s/token^2 quadratic share 20.9% @ w=2048 -> 67.8% @ w=16384 No fixed term was needed, which refutes launch-bound outright. Profiler kernel table (device rows only): attention 22,835.8 ms 65.2% fmha_cutlass*_sm80 dense GEMM 2,774.0 ms 7.9% other 5,739.0 ms 16.4% The attention kernels are sm80 — Ampere-generation CUTLASS running on an sm_120 Blackwell card, with the forward on the gmem fallback tier. That is the mechanism behind 100% SM utilisation at 27 of 304 available TFLOPS. Correctness cleared separately: the sliding mask asserts at max 1024 allowed/row, so the 25 windowed layers were genuinely windowed. The same probe found that right-padding is what pins the 5 global layers to an explicit 4D mask and off the is_causal fast path — measured at 9.4% slower for 24% less loss work at fixed width. The largest available win is not the attention kernel. The corpus is 29.9% padding, and bucket-to-pair + shuffle-to-mix takes it to 0.0% for >=35.5% wall clock, no new dependency, unchanged peak memory. Bucket size turned out not to be a diversity knob — roots per accumulation window are flat across a 256x range, so the global micro-batch shuffle does that work alone and the bucket should be tight. Adds docs/pfi/training-throughput-playbook.md as the durable model-agnostic home (sibling to the quantization playbook), the four probes under scripts/training-probes/ with raw output kept for re-derivation, and a §6 to the sizing doc carrying the Gemma-4-specific numbers and round-2 restart parameters. Measured negatives recorded so they are not re-chased: grouped_mm (0.9% slower, and MoE is only 7.9% of the step), CUDA graphs / torch.compile over the expert loop (no fixed cost to amortise), liger fused CE (~1-3% lever), FA4 on sm_120. Round-1 state preserved: 609MB encode cache, order manifest, truncation report, resume script. No checkpoints — it died at step 19 and the first was due at 100, so the lora_B inert-adapter gate never ran and moves to the restart.
127 lines
17 KiB
Plaintext
127 lines
17 KiB
Plaintext
========================================================================
|
|
loading model
|
|
========================================================================
|
|
|
|
Loading weights: 0%| | 0/1013 [00:00<?, ?it/s]
|
|
Loading weights: 0%| | 2/1013 [00:00<01:19, 12.71it/s]
|
|
Loading weights: 0%| | 4/1013 [00:00<01:10, 14.40it/s]
|
|
Loading weights: 2%|▏ | 25/1013 [00:00<00:16, 61.02it/s]
|
|
Loading weights: 5%|▍ | 48/1013 [00:00<00:12, 79.42it/s]
|
|
Loading weights: 7%|▋ | 70/1013 [00:00<00:10, 87.32it/s]
|
|
Loading weights: 9%|▉ | 92/1013 [00:01<00:10, 91.14it/s]
|
|
Loading weights: 11%|█ | 113/1013 [00:01<00:09, 92.58it/s]
|
|
Loading weights: 13%|█▎ | 135/1013 [00:01<00:09, 95.13it/s]
|
|
Loading weights: 15%|█▌ | 157/1013 [00:01<00:08, 97.82it/s]
|
|
Loading weights: 18%|█▊ | 179/1013 [00:02<00:08, 101.21it/s]
|
|
Loading weights: 20%|█▉ | 201/1013 [00:02<00:08, 100.37it/s]
|
|
Loading weights: 22%|██▏ | 223/1013 [00:02<00:07, 102.12it/s]
|
|
Loading weights: 24%|██▍ | 244/1013 [00:02<00:07, 104.64it/s]
|
|
Loading weights: 26%|██▋ | 266/1013 [00:02<00:07, 105.68it/s]
|
|
Loading weights: 28%|██▊ | 288/1013 [00:03<00:06, 106.87it/s]
|
|
Loading weights: 30%|███ | 308/1013 [00:03<00:05, 122.26it/s]
|
|
Loading weights: 32%|███▏ | 322/1013 [00:03<00:05, 117.76it/s]
|
|
Loading weights: 33%|███▎ | 335/1013 [00:03<00:06, 99.26it/s]
|
|
Loading weights: 35%|███▍ | 354/1013 [00:03<00:06, 100.53it/s]
|
|
Loading weights: 37%|███▋ | 375/1013 [00:03<00:06, 104.62it/s]
|
|
Loading weights: 39%|███▉ | 397/1013 [00:04<00:05, 103.34it/s]
|
|
Loading weights: 41%|████▏ | 419/1013 [00:04<00:05, 106.05it/s]
|
|
Loading weights: 44%|████▎ | 441/1013 [00:04<00:05, 109.33it/s]
|
|
Loading weights: 46%|████▌ | 463/1013 [00:04<00:05, 107.04it/s]
|
|
Loading weights: 48%|████▊ | 485/1013 [00:04<00:04, 107.29it/s]
|
|
Loading weights: 50%|█████ | 507/1013 [00:05<00:04, 105.14it/s]
|
|
Loading weights: 52%|█████▏ | 528/1013 [00:05<00:04, 103.18it/s]
|
|
Loading weights: 54%|█████▍ | 550/1013 [00:05<00:04, 102.05it/s]
|
|
Loading weights: 56%|█████▋ | 572/1013 [00:05<00:04, 101.82it/s]
|
|
Loading weights: 59%|█████▊ | 594/1013 [00:05<00:03, 106.80it/s]
|
|
Loading weights: 61%|██████ | 616/1013 [00:06<00:03, 104.33it/s]
|
|
Loading weights: 63%|██████▎ | 638/1013 [00:06<00:03, 103.57it/s]
|
|
Loading weights: 77%|███████▋ | 778/1013 [00:06<00:00, 320.91it/s]
|
|
Loading weights: 91%|█████████ | 920/1013 [00:06<00:00, 532.70it/s]
|
|
Loading weights: 100%|██████████| 1013/1013 [00:06<00:00, 152.35it/s]
|
|
loaded in 9.8s targets=205
|
|
final_logit_softcapping = 30.0
|
|
attn_implementation = sdpa
|
|
|
|
========================================================================
|
|
A. SEQUENCE SCALING (no padding - isolates n)
|
|
========================================================================
|
|
2 x 2,048 1.776 s kept=3227 peak= 53.2 GiB
|
|
2 x 8,192 11.570 s kept=13050 peak= 62.3 GiB
|
|
2 x 16,384 35.017 s kept=25989 peak= 76.6 GiB
|
|
|
|
16384 -> 2048 ratio 19.71x (linear ~8x, launch-bound ~1x, quadratic ~64x)
|
|
16384 -> 8192 ratio 3.03x (linear ~2x, quadratic ~4x)
|
|
|
|
========================================================================
|
|
B. PADDING PENALTY (same real tokens, with vs without pad)
|
|
========================================================================
|
|
2 x 16,384 no padding 35.244 s kept=26048 peak= 76.6 GiB
|
|
2 x 16,384 50% pad on row 1 38.567 s kept=19640 peak= 77.8 GiB
|
|
|
|
========================================================================
|
|
C. ISOLATED CE WALL CLOCK
|
|
========================================================================
|
|
2 x 16,384 (CE timed) 35.329 s kept=26210 peak= 76.6 GiB CE=374 ms (1.1%)
|
|
2 x 4,096 (CE timed) 4.387 s kept=6512 peak= 55.3 GiB CE=93 ms (2.1%)
|
|
|
|
========================================================================
|
|
D. KERNEL TABLE - one fwd+bwd at 2 x 16,384
|
|
========================================================================
|
|
USDT:2026-08-24 22:03:51 574811:574811 SyncActivityProfilerHandler.cpp:52] profiler_start
|
|
USDT:2026-08-24 22:04:27 574811:574811 SyncActivityProfilerHandler.cpp:59] profiler_stop
|
|
------------------------------------------------------- ------------ ------------ ------------ ------------ ------------ ------------ ------------ ------------ ------------ ------------
|
|
Name Self CPU % Self CPU CPU total % CPU total CPU time avg Self CUDA Self CUDA % CUDA total CUDA time avg # of Calls
|
|
------------------------------------------------------- ------------ ------------ ------------ ------------ ------------ ------------ ------------ ------------ ------------ ------------
|
|
aten::_efficient_attention_backward 0.00% 573.671us 0.00% 1.681ms 56.028us 16.145s 45.71% 16.156s 538.546ms 30
|
|
fmha_cutlassB_bf16_aligned_128x64_k65536_sm80(PyTorc... 0.00% 0.000us 0.00% 0.000us 0.000us 16.145s 45.71% 16.145s 538.152ms 30
|
|
aten::_efficient_attention_forward 0.00% 795.660us 0.01% 1.933ms 32.218us 6.691s 18.94% 6.691s 111.519ms 60
|
|
fmha_cutlassF_bf16_aligned_32x128_gmem_sm80(PyTorchM... 0.00% 0.000us 0.00% 0.000us 0.000us 6.691s 18.94% 6.691s 111.519ms 60
|
|
aten::mm 0.51% 180.757ms 0.77% 270.268ms 10.614us 3.745s 10.60% 3.745s 147.063us 25463
|
|
aten::mul 0.13% 45.428ms 0.18% 64.624ms 12.129us 2.721s 7.70% 2.721s 510.606us 5328
|
|
aten::copy_ 0.06% 19.823ms 92.46% 32.574s 6.100ms 1.977s 5.60% 1.977s 370.189us 5340
|
|
void cutlass::Kernel2<cutlass_80_tensorop_bf16_s1681... 0.00% 0.000us 0.00% 0.000us 0.000us 1.542s 4.37% 1.542s 656.804us 2348
|
|
void at::native::elementwise_kernel<128, 2, at::nati... 0.00% 0.000us 0.00% 0.000us 0.000us 1.033s 2.92% 1.033s 545.219us 1894
|
|
void cutlass::Kernel2<cutlass_80_tensorop_bf16_s1681... 0.00% 0.000us 0.00% 0.000us 0.000us 868.642ms 2.46% 868.642ms 583.373us 1489
|
|
void at::native::vectorized_elementwise_kernel<4, at... 0.00% 0.000us 0.00% 0.000us 0.000us 787.028ms 2.23% 787.028ms 395.095us 1992
|
|
void at::native::unrolled_elementwise_kernel<at::nat... 0.00% 0.000us 0.00% 0.000us 0.000us 697.963ms 1.98% 697.963ms 304.521us 2292
|
|
aten::masked_fill_ 0.01% 3.403ms 0.01% 4.927ms 27.373us 576.927ms 1.63% 576.927ms 3.205ms 180
|
|
void at::native::vectorized_elementwise_kernel<4, at... 0.00% 0.000us 0.00% 0.000us 0.000us 561.815ms 1.59% 561.815ms 413.403us 1359
|
|
void at::native::vectorized_elementwise_kernel<4, at... 0.00% 0.000us 0.00% 0.000us 0.000us 469.818ms 1.33% 469.818ms 459.255us 1023
|
|
aten::add_ 0.01% 3.730ms 0.02% 6.640ms 7.209us 448.665ms 1.27% 448.665ms 487.150us 921
|
|
void cutlass::Kernel2<cutlass_80_tensorop_bf16_s1681... 0.00% 0.000us 0.00% 0.000us 0.000us 363.206ms 1.03% 363.206ms 394.789us 920
|
|
aten::add 0.03% 9.745ms 0.04% 14.287ms 10.205us 356.226ms 1.01% 356.226ms 254.447us 1400
|
|
void at::native::elementwise_kernel<128, 2, at::nati... 0.00% 0.000us 0.00% 0.000us 0.000us 352.705ms 1.00% 352.705ms 1.959ms 180
|
|
aten::index 0.01% 4.542ms 0.09% 32.242ms 132.682us 352.404ms 1.00% 352.430ms 1.450ms 243
|
|
void at::native::vectorized_gather_kernel<16, long>(... 0.00% 0.000us 0.00% 0.000us 0.000us 351.768ms 1.00% 351.768ms 1.933ms 182
|
|
Memcpy DtoD (Device -> Device) 0.00% 0.000us 0.00% 0.000us 0.000us 341.588ms 0.97% 341.588ms 634.922us 538
|
|
aten::pow 0.06% 21.418ms 0.11% 37.019ms 18.659us 332.072ms 0.94% 493.271ms 248.624us 1984
|
|
void at::native::vectorized_elementwise_kernel<4, at... 0.00% 0.000us 0.00% 0.000us 0.000us 330.315ms 0.94% 330.315ms 499.720us 661
|
|
void at::native::vectorized_elementwise_kernel<4, at... 0.00% 0.000us 0.00% 0.000us 0.000us 316.108ms 0.89% 316.108ms 383.161us 825
|
|
aten::sum 0.01% 5.258ms 0.02% 7.700ms 13.461us 286.461ms 0.81% 286.464ms 500.811us 572
|
|
aten::native_dropout 0.02% 6.453ms 0.03% 11.139ms 27.168us 258.244ms 0.73% 258.244ms 629.865us 410
|
|
void at::native::(anonymous namespace)::fused_dropou... 0.00% 0.000us 0.00% 0.000us 0.000us 258.244ms 0.73% 258.244ms 629.865us 410
|
|
void at::native::unrolled_elementwise_kernel<at::nat... 0.00% 0.000us 0.00% 0.000us 0.000us 239.055ms 0.68% 239.055ms 583.061us 410
|
|
aten::_index_put_impl_ 0.01% 4.170ms 3.22% 1.134s 7.508ms 235.816ms 0.67% 236.930ms 1.569ms 151
|
|
void at::native::elementwise_kernel<128, 4, at::nati... 0.00% 0.000us 0.00% 0.000us 0.000us 231.434ms 0.66% 231.434ms 385.081us 601
|
|
void at::native::vectorized_elementwise_kernel<4, at... 0.00% 0.000us 0.00% 0.000us 0.000us 226.451ms 0.64% 226.451ms 692.511us 327
|
|
void at::native::elementwise_kernel<128, 4, at::nati... 0.00% 0.000us 0.00% 0.000us 0.000us 224.222ms 0.63% 224.222ms 2.491ms 90
|
|
void cutlass::Kernel2<cutlass_80_simt_sgemm_64x128_8... 0.00% 0.000us 0.00% 0.000us 0.000us 219.030ms 0.62% 219.030ms 534.219us 410
|
|
void at::native::elementwise_kernel<128, 4, at::nati... 0.00% 0.000us 0.00% 0.000us 0.000us 209.409ms 0.59% 209.409ms 1.745ms 120
|
|
aten::div 0.01% 4.911ms 0.02% 6.652ms 13.278us 192.925ms 0.55% 192.925ms 385.080us 501
|
|
void (anonymous namespace)::indexing_backward_kernel... 0.00% 0.000us 0.00% 0.000us 0.000us 183.000ms 0.52% 183.000ms 6.100ms 30
|
|
void at::native::unrolled_elementwise_kernel<at::nat... 0.00% 0.000us 0.00% 0.000us 0.000us 148.335ms 0.42% 148.335ms 988.899us 150
|
|
aten::mean 0.02% 6.214ms 0.02% 8.406ms 12.717us 144.904ms 0.41% 144.904ms 219.220us 661
|
|
void at::native::reduce_kernel<512, 1, at::native::R... 0.00% 0.000us 0.00% 0.000us 0.000us 144.904ms 0.41% 144.904ms 219.220us 661
|
|
aten::_log_softmax 0.00% 527.616us 0.00% 716.309us 13.775us 139.294ms 0.39% 139.294ms 2.679ms 52
|
|
void at::native::(anonymous namespace)::cunn_SoftMax... 0.00% 0.000us 0.00% 0.000us 0.000us 139.294ms 0.39% 139.294ms 2.679ms 52
|
|
void at::native::reduce_kernel<512, 1, at::native::R... 0.00% 0.000us 0.00% 0.000us 0.000us 137.743ms 0.39% 137.743ms 286.368us 481
|
|
void at::native::reduce_kernel<128, 4, at::native::R... 0.00% 0.000us 0.00% 0.000us 0.000us 130.526ms 0.37% 130.526ms 1.088ms 120
|
|
aten::native_dropout_backward 0.00% 1.378ms 0.01% 3.156ms 15.394us 118.690ms 0.34% 118.690ms 578.975us 205
|
|
------------------------------------------------------- ------------ ------------ ------------ ------------ ------------ ------------ ------------ ------------ ------------ ------------
|
|
Self CPU time total: 35.229s
|
|
Self CUDA time total: 35.322s
|
|
|
|
|
|
========================================================================
|
|
E. LAUNCH COUNTS (grouped_mm: 128/layer sequential = no-op, 1 = grouped)
|
|
========================================================================
|
|
kernel count self ms
|
|
aten::_efficient_attention_backward 30 16144.6
|
|
fmha_cutlassB_bf16_aligned_128x64_k65536_sm80(PyTorchMemEf 30 16144.6
|
|
aten::_efficient_attention_forward 60 6691.2
|
|
fmha_cutlassF_bf16_aligned_32x128_gmem_sm80(PyTorchMemEffA 60 6691.2
|
|
aten::mm 25463 3744.6
|
|
aten::mul 5328 2720.5
|
|
aten::copy_ 5340 1976.8
|
|
void cutlass::Kernel2<cutlass_80_tensorop_bf16_s16816gemm_ 2348 1542.2
|
|
void at::native::elementwise_kernel<128, 2, at::native::gp 1894 1032.6
|
|
void cutlass::Kernel2<cutlass_80_tensorop_bf16_s16816gemm_ 1489 868.6
|
|
void at::native::vectorized_elementwise_kernel<4, at::nati 1992 787.0
|
|
void at::native::unrolled_elementwise_kernel<at::native::d 2292 698.0
|
|
aten::masked_fill_ 180 576.9
|
|
void at::native::vectorized_elementwise_kernel<4, at::nati 1359 561.8
|
|
void at::native::vectorized_elementwise_kernel<4, at::nati 1023 469.8
|
|
aten::add_ 921 448.7
|
|
void cutlass::Kernel2<cutlass_80_tensorop_bf16_s16816gemm_ 920 363.2
|
|
aten::add 1400 356.2
|
|
void at::native::elementwise_kernel<128, 2, at::native::gp 180 352.7
|
|
aten::index 243 352.4
|
|
void at::native::vectorized_gather_kernel<16, long>(char*, 182 351.8
|
|
Memcpy DtoD (Device -> Device) 538 341.6
|
|
aten::pow 1984 332.1
|
|
void at::native::vectorized_elementwise_kernel<4, at::nati 661 330.3
|
|
void at::native::vectorized_elementwise_kernel<4, at::nati 825 316.1
|
|
aten::sum 572 286.5
|
|
aten::native_dropout 410 258.2
|
|
void at::native::(anonymous namespace)::fused_dropout_kern 410 258.2
|
|
void at::native::unrolled_elementwise_kernel<at::native::C 410 239.1
|
|
aten::_index_put_impl_ 151 235.8
|
|
|
|
total self CUDA time 70644.4 ms
|
|
GEMM-ish kernels 31761.7 ms (45.0%)
|
|
non-GEMM 38882.7 ms (55.0%)
|