Files
xtrain/docs/known-issues.md
Gahow Wang a1370446fe docs: T21 — record DDP-dropout wiring gap + fix (known-issues / evolution / dropout doc)
- known-issues.md: new "DDP-dropout wiring" Fixed entry (gap + fix +
  regression test), with the meta-lesson that op/single-GPU unit tests can
  miss launcher-level integration gaps — only the V9-PILOT end-to-end run on
  the real launcher path exposed it.
- 17-dropout.md: annotate the DDP-combination note with the T18 wiring gap
  and its T21 fix.
- evolution.md: T21 row (Infra) recording the fix + meta-lesson.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
2026-06-18 21:22:49 +08:00

25 KiB
Raw Blame History

xtrain — Known Issues & Perf Backlog

已知问题(性能 / 正确性 / 建模)与延后项的活文档:记录现象、复现、根因、拟修复、优先级、状态。 发现即记,修复即标 FIXED(附 commit


Open

(KI-1 fixed in T10. KI-5 fixed in T11. KI-2 fixed in T12. KI-3激活重计算 / gradient checkpointingFIXED in T13——per-block checkpointno-tape forward + recompute-on-backward梯度对非重计算版逐位一致fp32/bf16 max rel 0.00e0dim768 bf16 batch32 峰值显存 31.1→14.6GB53%/ tok/s 20%dim1024 batch32 不开 OOM → 开了 16.6GB 装得下,解锁 v8。见下方 Fixed。)


Fixed

DDP-dropout wiringlauncher 漏接 dropoutFIXED (T21)

  • 背景V9-PILOT 暴露)T18 dropout 在单卡 train.rs 完整接好(--dropout flag → cfg.dropout,每步 model.train()eval 用 model.eval()op 级 + 单卡都测过。但 V9-PILOT 全栈端到端跑DDP 8 卡 + dropout0.1 + flash + GQA + accum + bf16时发现 DDP 路径根本没接 dropouttrain_ddp bin --dropout flag、从不设 cfg.dropout,且 ddp.rs::train_rank 从不调 model.train() → 模型停在默认 eval 模式,ops::dropout 恒等 → DDP 下 dropout 被静默忽略,无论 config 怎么设。模型 + autodiff 完全支持 dropoutT18漏的纯是 DDP launcher / 训练 loop 的 wiring
  • 为何 op/单卡测试没抓到dropout 的测试只覆盖单卡训练循环 + op 级 grad-check,从没在 DDP 路径下跑过 dropout。train_rank 是独立于单卡 train() 的另一条 loop二者共享 model/autodiff 但各自布线 train/eval 纪律 —— 单卡那条对了不代表 DDP 那条对。元教训op 级 + 单 GPU 单元测试能漏掉 launcher 级 integration gap;只有把特性放进真实启动器路径端到端跑pilot 做的事)才暴露。修复随手补了 DDP 路径的回归测试堵这个缺口。
  • 修复(docs/17-dropout.md:① train_ddp.rs--dropout <p> flag默认 0 = 关,对齐旧行为)并设 cfg.dropout;② ddp.rs::train_rank 每步 micro-batch 循环前调 model.train(),镜像单卡 loop 的 train/eval 纪律——关键eval_loss() 内部 model.eval() 翻成 eval 模式且不还原,所以每步重新 assert model.train()dropout 才能跨 eval 边界保持活跃。
  • 正确性(新增 DDP-dropout 回归测试 ddp_dropout_is_live_and_p0_bit_identical,跑真实 train_rank 启动器路径):① GATE A——p=0 下 DDP loss 轨迹 + 末态参数对无 dropout 路径逐位一致ops::dropout(p=0) 是 clone no-op回归保护GATE B——p=0.2 的 loss 轨迹对 p=0 有可观差异>1e-3证 dropout mask 真在训练 forward 应用pre-T21 代码停在 eval 模式 → 二者会逐位相同,此 gate 会 FAILGATE C——run 后 model.is_training()==true(直接证 model.train() 被调用且跨末步 eval 存活);④ p>0 run 无 NaN/Inf。测试故意启用 eval_every < steps 让 eval 中途翻 eval 模式,验证每步 model.train() 的 restore 纪律。默认 --dropout 0 下既有 DDP loss-match + 跨 rank 测试不变(回归保护)。
  • commit:见 T21 提交链(distributed: --dropout flag + model.train() per step in train_rank / test: DDP-dropout regression (live under DDP + p=0 bit-identical) / 文档更新)。

process-per-GPUtorchrun 式独立 CUDA contextCLOSED / 实测负结果 (T17)

  • 背景KI-5T11修掉 per-op cudaMalloc 串行后8 卡 scaling 从 ~1.3× 恢复到 ~5×@8,但残留 ~5×@8 非完美线性。T11 doc / KI-5「残留」推测下一步是 process-per-GPU(每 rank 独立进程 + 独立 CUDA contexttorchrun 式——理由是「N rank 线程共享单 CUDA primary contextkernel-launch/cuBLAS 仍在 context 级串行」。T17 把这条 torchrun 式链路落地并实测,证伪了该推测。

  • 实现(docs/16-process-per-gpu.mdxtrain-distributedproc.rs——launch_processes 每卡 spawn 一个 worker 进程re-exec current_exe + XTRAIN_{RANK,WORLD,LOCAL_RANK,NCCL_ID} envlauncher 一次性铸 ncclUniqueId 后 hex 编码注入子进程 env(无共享 FS/TCP、无轮询、无竞态——id 在子进程出生前就原子就绪worker 读 env → bind device独立 CUDA contextDdpContext::init + build_model + train_rank 全部复用 T8 零改动。新 train_ddp_mp bin + ddp_proc test保留 thread-per-GPU 旧路径(回归 baseline。scope=process-per-GPU onlyZeRO-1 用户 drop

  • 正确性(全绿,无回归)proc vs 单卡 loss 5.67e-7proc vs thread-per-GPU 1.5e-7(两路数值同量级)、跨 rank 1.19e-7<1e-6全回归套 --test-threads=1 全绿 + xserv 闭环 v3 重导 md5 逐位一致 b04fc9f9T17 不碰任何数值路径)。

  • 实测结果关键dash5 8× RTX 5090, dim384 per-rank batch32 seq256, steady-state

    world thread-per-GPU (train_ddp) speedup process-per-GPU (train_ddp_mp) speedup
    1 93257 1.00× 92952 1.00×
    2 149747 1.61× 148809 1.60×
    4 278276 2.98× 273308 2.94×
    8 491360 5.27× 493128 5.31×

    world=8 各重复 2 次thread 493671/493292、proc 491102/494123——两路差异 <1%,落在噪声内。)

  • 诊断(证伪原推测)process-per-GPU world=8 跑时 nvidia-smi 抽样 8 卡全部 9599% util(每卡 ~23GB——GPU 已 compute-bound 喂满、非串行空转KI-5「12/8 在忙」的串行病 T11 allocator 已治好。8 卡满载却仍只 5.3× ⇒ 缺的 ~35% 吞吐去向每步 grad all-reduce + 本机 PCIe-only 拓扑在 8 rank 下的通信开销T11 早点明的「~7% all-reduce + PCIe 余量」那一层8 卡放大),换独立 context 不动这一层。结论本尺度dim3841024、单机 8× PCIe RTX 5090残留非线性是通信/拓扑墙,不是 launch 模型——要再逼近线性须动 all-reduce overlap / NVLink 互联(非本尺度优先)。

  • 方法论一致T11 实测证伪「分桶 all-reduce」、T17 实测证伪「process-per-GPU 解残留串行」——两次都靠 measure 推翻假设而非硬上profile/measure-first净价值:落地 torchrun 式 process-per-GPU 标准链路(项目本职「学训练全栈」)+ 把这个误导性 backlog 项实测钉死关闭默认训练路径不变thread-per-GPUprocess-per-GPU 作并列可选路径留档。

  • commit:见 T17 提交链(distributed: process-per-GPU launcher + worker / distributed: train_ddp_mp bin / test: process-per-GPU DDP correctness / 设计文档 docs: Phase T17 — process-per-GPU DDP design)。


KI-3 · 激活重计算gradient checkpointingFIXED (T13)

  • 触发点v8 surfaced:容量轴放大到 dim1024核心 ~210M+)测是否 capacity-limited。autograd tape 为反向保存所有中间激活,激活显存随 dim 线性增长——dim768 bf16 batch32 已 31.1GBT12 甜点区),dim1024 batch32 再次 OOM(实测撞 32100/32607MiB → OutOfMemory)。

  • 设计per-block gradient checkpointingopt-indocs/12-activation-recompute.md:新增 xtrain_autodiff::checkpoint(segment_fn, input, params) 高阶原语(类比 torch.utils.checkpoint)。前向:把 input/params detach 成局部 leaf 跑 segment_fn,只取输出值,局部 tape 立即 drop → 段内激活释放(不留在外层 tapecheckpoint 节点 parents=[input, ..params]。反向:从保存的 input + 未变的 param 值重跑 segment_fn 重建局部 tape用上游 grad seedVar::backward_seeded,新增——段输出非标量)回传,恢复的 input/param 梯度 push 给真 parents局部 tape drop → 重算激活释放。模型每个 transformer block 前向用它包裹(--recompute flag默认关。切粒度 = 每 block。

  • 正确性exact硬闸门全绿:重计算数学精确(同 segment_fn、同输入、同参数值、确定性 kernel → 重算激活逐位等于原激活)。on-vs-off 梯度对拍 fp32/bf16 双路逐位一致loss rel 0.00e0、logits max rel 0.00e0每个参数梯度 max rel 0.00e0(不是容差内,是逐位)。全套回归开/关重计算全绿autograd 15、structural 5、batched、bf16、overfit 27/27、AdamWGPU bit-exact + host 对 torch、checkpoint roundtrip、DDP loss 对单卡 5.67e-7 + 跨 rank 0.0DDP+recompute 2 卡短训单调降11.079→6.010)。非重计算路径图不变(默认关)→ T10/T11/T12 数值不回归。

  • 显存 + 吞吐payoffdash5 1× RTX 5090 32GB, bf16, batch32 seq256, steady-state

    config per-rank batch 峰值显存 tok/s fits 32GB?
    dim768 (18L/24h ffn2048, core 127M) off 32 31144 MiB 39.7K
    dim768 on 32 14562 MiB53% 31.5K20%
    dim1024 (18L/32h ffn2730, core 226M) off 32 32100 → OOM OOM
    dim1024 on 32 16596 MiB 23.1K 解 OOM

    → dim768重计算砍 ~半激活(31.1→14.6GB53%),代价 tok/s 20%(多一次前向,落在预测 2035%)。dim1024 batch32不开 OOM → 开了 16.6GB 稳稳装下~23K tok/s 正常收敛 → dim1024 解锁v8 可展开

  • commitT13 提交链(autodiff: checkpoint primitive (recompute-on-backward) / model: per-block activation recompute (--recompute) / perf: KI-3 fixed … 本条带 before/after / 文档 docs: Phase T13 — activation recompute)。


KI-2 · bf16 混合精度fp32 masterFIXED (T12)

  • 触发点v4 surfaceddim768 fp32 在单卡 32GB 里 per-rank batch 32global 256OOM被迫降到 per-rank 16。bf16激活减半找回 batch-32 甜点区,并加速已 compute-bound 的 dim768 GEMM附带xserv 推理 BF16-onlybf16 训练更贴闭环。

  • 设计(标准 AMPopt-indocs/11-bf16-mixed-precision.mdfp32 master weights + AdamW/clip/DDP all-reduce 全程 fp32bf16 compute=linears(q/k/v/o, gate/up/down, lm_head)走 cublasGemmExCUDA_R_16BF in/out, CUBLAS_COMPUTE_32F 累加)+ 激活流 bf16含 attention probs / logitsfp32 稳定=RMSNorm/QK-norm、softmax、RoPE、cross-entropy 内部 upcast→fp32→downcast。无 loss scalingbf16 8-bit 指数=fp32 动态范围)。关键钩子=autodiff cast 算子:前向把 fp32 master leaf 降到 bf16 喂 matmul反向把 bf16 grad 升回 fp32 → fp32 leaf 累加 fp32 grad优化器一行不改。fp32 路径按 dtype 分派、逐字节不变hard gate

  • 正确性(双闸门全绿)

    • fp32 不回归全套在原紧容差绿——autograd 15、structural 5、GEMM 对 cuBLAS 5、batched==looped、overfit 27/27、AdamW GPU bit-exact + host 对 torch、checkpoint 逐位、DDP loss 对单卡 + 跨 rank、xserv 闭环v3 ckpt 用 T12 代码重导 safetensors 与 registry md5 逐位一致 b04fc9f9…)。
    • bf16 looser-tol:同 fp32 master 跑 fp32 vs bf16——loss rel 1.2e-4、logits mean rel 2.0e-3 / p99 6.8e-3、grad worst scaled-mean 1.0e-2,无 NaNgrad 仍 fp32。
    • 收敛dim768 短训 150 步bf16-b16 loss 轨迹对住 fp32-b16step50 4.40 vs 4.40、step149 3.984 vs 3.988),单调降、无发散。
  • 显存 + 吞吐payoffdash5 1× RTX 5090 32GB, dim768/18L/24h×32 ffn2048 seq256, steady-state

    config per-rank batch 峰值显存 tok/s fits 32GB?
    fp32 16 (v4 fallback) 27.2 GB 31.5K
    bf16 16 19.3 GB29% 35.5K+13%
    fp32 32 OOM
    bf16 32甜点区 31.1 GB 40.8K 解 OOM

    同 batch 16bf16 显存 27.2→19.3GB29%、tok/s 31.5K→35.5K+13%bf16 解 fp32-batch32 OOM31.1/32.6GB fitbatch32 达 40.8K tok/s+29% vs fp32-b16。残留norm/softmax/CE 的 fp32 upcast 是 transient但仍占峰值——若 v5 要更大 batch下一杠杆是 KI-3 激活重计算。

  • commit:见 T12 提交链(cuda: bf16 cuBLAS GemmEx / autodiff: bf16 mixed-precision path / train: --bf16 flag / perf: keep bf16 logits / 本条)。


KI-5 · DDP 弱扩展性 — FIXED (T11, device caching/pool allocator)

  • 根因T11 重诊断all-reduce 不是瓶颈):每个 tape op 输出走 Tensor::zerosGpuBuffer::alloccudaMalloc(同步、进程级串行的 driver 调用)。单进程 thread-per-GPU 下 N rank 每步几百次 alloc 在单 CUDA context 排队串行(NOCOMM=1 完全不通信时 fwd+bwd 仍 136→780ms 膨胀 ~6×nvidia-smi 抽样 8 卡只 12 张在忙轮流跑);单卡也吃这笔 per-op alloc。

  • 原拟修复「分桶 all-reduce」经 T11 实测证伪并 revertgrad all-reduce 每步只占 ~67%,融成一发对 1/2/4/8 卡几乎无差(详见下方历史诊断)。

  • 修复xtrain-cudadevice caching/pool allocatordocs/10-caching-allocator.md)——GpuBuffer::alloc 从 per-device、size-classed free-list 取miss 才 cudaMallocDrop 归还 free-listcudaFree。训练定形状→命中率极高warm-up 后每步 cudaMalloc≈0。线程安全global registry 按 device id 分桶,每 device 的 free-list 各自 Mutexregistry 锁只在 clone 出 Arc<Mutex<_>> 时极短持有→跨 device 真并发buffer 记 alloc 时的 deviceDrop 归还对应 pool。透明:物理 cap 可向上取整但 len()/memset/copy bounds 都用请求 len尾部字节永不读到→数值逐位不变。memset 保留(复用 buffer 有陈旧字节skip-memset uninit 本次不做malloc 已是瓶颈memset async 开销小,逐 op 证明全覆盖风险大)。

  • before → afterdash5, 8× RTX 5090, dim384/12L per-rank batch 32 seq 256, steady-state tok/s; before=d422c68 after=pooled

    world before tok/s before speedup after tok/s after speedup
    1 39801 1.00× 92385 1.00×
    2 47229 1.19× 146821 1.59×
    4 52854 1.33× 269867 2.92×
    8 48996 1.23× 461270 4.99×

    单卡 40226→92638 tok/s (~2.3×)8 卡 49K→461K tok/s (9.4×)scaling 从 ~1.3× 封顶恢复到 ~5×@88 卡 nvidia-smi 抽样 全 8 卡 9599% utilKI-5 时只 12/8 忙。loss 轨迹逐位对住(单卡 10.9026→4.8453 before/after 一致)。

  • 正确性(全绿,无回归)15 算子 grad-check、5 结构、GEMM 对 cuBLAS、batched==looped、overfit 27/27、AdamW GPU bit-exact + host 对 torch、checkpoint 逐位、DDP loss 对单卡 5.67e-7 + 跨 rank diff 0.0loosened <1e-6)、xserv 闭环v3 ckpt 重导 safetensors 与 registry md5 逐位一致 + xserv 加载服务贪心 "Once upon a time," 对住)。

  • 顺手DDP ddp_correctness 的 cross-rank ==0.0<1e-6(本机 PCIe-only NCCL run-to-run 跨 rank 非逐位可复现diff≤1.2e-7 几 ULP 无害,承重闸门是 loss-match 5.67e-7ddp_throughput_scaling 扩到 world=8。

  • 残留~5×@8 非完美线性grad all-reduce ~7% + 8 卡 PCIe/launch 余量),但弱扩展悬崖已消。曾以为下一步是 process-per-GPU(每 rank 独立 CUDA contexttorchrun 式)——T17 实测证伪该方向见下方「process-per-GPUT17残留是通信/PCIe 墙,不是单 CUDA context 的 launch/cuBLAS 串行。

  • commit:见 T11 提交链(cuda: device caching allocator / perf: KI-5 … 那条带 before/after

  • 历史诊断保留如下(证伪「分桶 all-reduce」的过程


KI-5 历史诊断 · DDP 弱扩展性 — T10 暴露T11 重诊断all-reduce 不是瓶颈)

  • 现象batched forward 修掉单卡 launch-bound 后dim384/per-rank batch 321 卡 40.3K → 4 卡 47.2K tok/sglobal仅 ~1.17×。

  • T11 实测dash5, 8× RTX 5090, dim384/12L, per-rank batch 32, seq 256, 原 ungrouped all-reduce, 50 步均, ms/step

    world fwd+bwd grad all-reduce clip+opt+zero TOTAL tok/s(global) speedup
    1 136 0 8.6 145 36582 1.00×
    2 202 21 15 238 47267 1.29×
    4 342 29 21 392 51466 1.41×
    8 780 54 47 882 47719 1.30×

    → grad all-reduce 每步只占 ~67%;真正爆炸的是逐 rank 的 fwd+bwd 时间随 world 线性膨胀(同一 per-rank workload136→780ms~6×

  • 「分桶 all-reduce」拟修复经 T11 实测证伪(无收益):把 ~150 个 per-tensor ncclAllReducencclGroupStart/End 融成一发 → 1/2/4/8 卡 = 1.00/1.30/1.42/1.34×与不分桶几乎无差all-reduce 本就只占 7%。flat-buffer 分桶同理。故回退revert b8b5821保留原 ungrouped 路径。

  • 附带发现T8 correctness 测试的 max|p0p1| == 0.0 在本机 flaky(与 T11 无关)。原 ungrouped 代码同一 GPU 重跑 6 次 cross-rank diff = {0.0, 0.0, 5.96e-8, 5.96e-8, 1.19e-7, 1.19e-7},只 ~1/3 命中 0.0。即本机/本版 NCCL 的 all-reduce run-to-run 跨 rank 不是逐位可复现PCIe-only 拓扑下 algorithm/chunk 选择不稳。diff 都 ≤1.19e-7几 ULP数值无害loss-match 仍 ~6e-7== 0.0 断言过严 → 建议改为 < 1e-6 紧容差(留作 follow-up本次未改测试)。

  • 重新定位的根因单进程 thread-per-GPU 模型下N 个 rank 线程各自跑独立训练却互相串行——NOCOMM=1(完全不做任何跨 rank 通信/barrier时 fwd+bwd 仍 136→378→800ms 膨胀;nvidia-smi 抽样显示 8 卡同一时刻只有 12 张在忙、轮流跑。排除项CPU 不缺187 核, load 2.5nvcc --default-stream per-thread 不解决。剩余怀疑:每个 op 输出走 Tensor::zeroscudaMalloc+cudaMemset,而 cudaMalloc 是同步、进程级串行的 driver 调用;单 CUDA context 下 N rank 每步几百次 alloc 互相排队——即 DDP 真瓶颈是 per-op 显存分配 / driver 调用在单进程内串行,不是梯度通信。

  • 真正的修复方向(待定,非 T11 范围):① caching/pool allocatorop 输出复用显存,消掉每步几百次 cudaMalloc,单卡也受益);或 ② process-per-GPU(每 rank 独立 CUDA contexttorchrun 式,彻底解串行,但要改 launcher + 跨进程 UniqueId 分发)。先做 ① 再实测是否解 DDP 串行。

  • 重启条件:多卡 v4 需要扩展性时做。单卡 batched 已 40K tok/sv3 即单卡训完),多卡当前只有 ~1.4× 上限v4 若要多卡须先修上面的真瓶颈。


KI-1 · 单序列 launch-bound"DDP 弱扩展性"的根因)— FIXED (T10, batched forward)

  • 修复T10 给 model + autograd 加 batch 维——linears 摊平成 [B*S, dim] 一个大 GEMM 填满 GPUattention 走 fused 批量 SDPAcublasSgemmStridedBatched ×2 + 一个 causal-softmax kernelRoPE 位置 per-sequence 复位(row % S);训练 loop 用真 batch 一次 forward/backward 替代 "loop B 次 + SUM"。详见 docs/09-batched-forward.md

  • before → afterdim384/12L/12h, batch 16, seq 256, 1 卡, back-to-back A/B

    tok/s GPU util 显存
    before单序列 launch-bound ~1653 015% ~3 GB
    afterbatched 25627batch16/ 40263batch32 37% 均值 / 54% 峰 ~10 GB

    → 单卡 ~15.5×batch16/ ~24×batch32util 015% → 3754%。

  • 正确性(全绿,无回归)15 算子 grad-check新增 batched-rope / transpose_4d12 / batched-attention dQ/dK/dV、batched==looped 单序列logits 0.0、grad 6.4e-4PyTorch 对拍 B>1loss 5e-8 / logits 6.9e-6 / 全参数 grad 在 rtol 2e-2、overfit 27/27、checkpoint 逐位、AdamW 对 torch、DDP loss 对单卡 5.7e-7 + 跨 rank 参数 bit-identical(0.0)、xserv 加载导出权重对 xtrain 贪心仍逐 token 一致top token 同序、BF16 漂移 ~0.03)。

  • commit:见 T10 提交链(perf: KI-1 fixed — GPU util / tok/s 那条带 before/after

  • DDP 残留弱扩展性 → KI-5(这是 batching 后新暴露的 all-reduce 瓶颈,不是 KI-1 的单序列根因)。

  • 历史诊断保留如下v2 暴露 → v3 重诊断的过程,证明根因不是 all-reduce


KI-1 历史诊断 · DDP 弱扩展性(吞吐受单序列 launch-bound 限制)— v2 暴露v3 重新诊断

  • 现象4 卡 DDP 仅 ~3.2K tok/s几乎不快于单卡≈2× over 单卡远低于近线性T8 在 tiny micro-bench 为 3.0×@4
  • 复现dim384/12L, world=4, seq 256
  • v3 实测dash5, 4× RTX 5090, dim384, 隔离 back-to-back A/B
    global_batch 每卡 tok/s4卡 GPU util 显存
    32 8 3163 569%spiky ~23 GB / 32 GB
    256 64 3200 015% ~23 GB / 32 GB
    加大 8× batch 仅 +1.2% 吞吐(噪声内)。1 卡 dim384 ≈ 1653 tok/s4 卡 3163 ≈ 2.1×。
  • 原"拟修复"(加大 global batch经 v3 实测 falsifiedgbatch256 时每 token 的 all-reduce 次数只有 gbatch32 的 1/8若瓶颈是 all-reduce 应大幅提速——实际没有 → all-reduce / 通信不是瓶颈
  • 重新诊断的根因:瓶颈是单序列模型设计T5每个 sequence 各跑一次独立 forward/backward逐 op kernel-launch 开销,见 docs/06 延迟瓶颈。GPU util 仅 015%、显存仅占 ~8% → 严重 launch-bound / under-utilizedGEMM 太小喂不饱 GPU。加大 batch 只是按比例增加串行 launch 次数无法摊薄。4 卡相对单卡 ~2× 的固定天花板来自跨 rank 同步税,但不是靠调 batch 能修的。
  • 真正的修复(需实作,非调参)
    1. batched多序列forward——把一个 step 的多条序列在 batch 维一次性过模型,让 GEMM 大到能填满 GPU这是 launch-bound 的根本解,但要改 T4/T5 的 single-sequence autograd/model工作量大、有正确性风险
    2. 在 (1) 之后,梯度 all-reduce 分桶 + 与 backward 重叠bucketed / overlapped all-reduce才会有意义当前 all-reduce 已非瓶颈,做了也无收益)。
  • 参考docs/07-distributed.mddocs/06-performance.md

Deferred来自 T7放大后重启

(KI-3 已在 T13 FIXED见上方 Fixed。本节暂无待启项。)


Modeling notes

KI-4 · 大词表 embedding 占比过高 — 接受的建模权衡用户拍板T19 DROPPED

  • gpt2 vocab=50257 在 dim 小时让 embed+lm_head 主导参数v1 25.7M/34M、v2 38.6M/66.9Mcore transformer 才是学习主体。
  • 决定(不是 open是一个被接受的权衡:曾计划 T19 训一个更贴语料的小 vocab 来压 embed 占比,用户拍板 DROP——保住 xserv gpt2-tokenizer 闭环(项目皇冠:导出权重回流 xserv 逐 token 一致)比清理 embedding 占比更值。小 dim 下 embed 占比高=复用 gpt2 大词表的已知、接受的代价
  • 缓解路径仍在:更大 dim 时 core 自然成为主体(继续 scaling 即可摊薄占比v8 dim1024 core 226M 已主导);若日后愿意放弃该版闭环再重启小词表(见 ~/toc/projects/xtrain.md T19

集成 / 测试注记pre-existing非回归记账

DDP 三测并行会争卡 deadlock → --test-threads=1

  • xtrain-distributed 的三个 DDP 测试thread-per-GPU correctness / scaling、process-per-GPU ddp_proc)若被 cargo 并行跑,会在共享的 2 卡上互相争 GPU/NCCL 资源 deadlock(不是数值 bug是测试调度
  • 跑法cargo test ... -- --test-threads=1(或把 DDP 测试标 serial串行跑即全绿。Phase-2 全回归 capture 均在 --test-threads=1 下取得。

fresh-train md5 run-to-run 不定 → 有效确定性闸门是「导出重确定性」而非「fresh-train 复现」

  • 现象从随机初始化全新训练fresh-train产出的 ckpt其 md5 run-to-run 不逐位复现
  • 根因:反向里多处 atomicAdd 归约(如跨行 dK/dV、扇出累加的浮点加法序非确定GPU 原子操作完成序不固定)→ 末位 ULP 抖动逐步累积 → ckpt 字节不同。属本机/本版的已知浮点非确定性,不是正确性回归loss 轨迹仍同量级收敛)。
  • 因此项目的确定性硬闸门定义为「导出export重确定性」:拿同一个已固定的 ckpt 重新导出 HF-safetensorsmd5 与 registry 逐位一致b04fc9f9,两阶段全程)——这条是确定性的、承重的;不要求 fresh-train 字节复现。