Files
xtrain/docs/06-performance.md
Gahow Wang 5e8add2a41 docs: Phase T7 — performance
Design doc for the T7 fp32-preserving speedups: cuBLAS matmul fwd/bwd
(row-major⟺col-major layout), GPU AdamW + GPU grad-norm (no per-step
param/grad roundtrip), drop per-op sync + device memset. Includes the
verification table (regression suite green + tok/s 2770→8220 ~3x), the
deferred bf16/recompute follow-up rationale, and the T8 all-reduce note.

Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com>
2026-06-15 17:00:29 +08:00

10 KiB
Raw Blame History

Phase T7: Performance — Design Document

Goal

T6 把真训练打通了TinyStoriesloss 10.83→3.43,采样连贯),但吞吐只有 ~2800 tok/s。 T7 的目标是在不牺牲数值正确性的前提下把训练显著加速——上来先把 fp32 路径里那些纯开销榨干, 再视情况上 bf16 / 激活重计算。

按 xtrain.md T7 note 的优先级,保 fp32 数值不回归的三步是 must-have

  1. matmul fwd/bwd 走 cuBLAS —— 前向 + 两路反向(dA=dC·BᵀdB=Aᵀ·dC)全切 cuBLAS Sgemm。fp32等价于换了求和顺序的同一个 GEMM所以正确性自动保住。
  2. GPU 侧优化器 + grad clip —— 干掉每步把全部参数/梯度 GPU↔host 往返的开销AdamW 的 m/v 状态搬到 device、update 走 kernelglobal grad-norm 用 device reduction只把那一个标量取回 host。
  3. stream / 减 sync —— 不再每个 op 之后都 cudaDeviceSynchronizedefault stream 上 kernel 本就顺序执行host 读数据又都走 stream-ordered 的 cudaMemcpyper-op sync 是纯开销,全删。

不做(本 Phase 范围外):分布式数据并行 / NCCL all-reduceT8、导出回流 xservT9

降级出口bf16 混合精度(④)/ 激活重计算(⑤)改数值、牵动每个 kernel且本 model 太小dim=32属 latency-bound、bf16 tensor-core 收益有限——按 xtrain.md 的 escape hatch①②③ 交付并实测加速后,④⑤ 记为 follow-up不把正确性留在半截状态。详见末节。

Module Layout

csrc/ops/optim.cu                        # 新GPU AdamW step + global grad sumsq reduce + in-place scale
crates/xtrain-cuda/
├── src/cublas.rs                        # 新:持久化 cuBLAS handle + row-major sgemm(含转置位)
├── src/ffi.rs                           # +CUBLAS_OP_T、+optim.cu 的三个 launch_*、+cudaMemset
├── src/memory.rs                        # GpuBuffer::memsetdevice 置零,免 H2D 零拷贝)
├── src/lib.rs                           # pub mod cublasnot(no_cuda)
└── build.rs                             # +optim.cu
crates/xtrain-tensor/
├── src/tensor.rs                        # matmul/matmul_backward 改走 cublas::sgemm删 21 处 per-op sync
└── src/storage.rs                       # device zeros 改用 memset
crates/xtrain-optim/
├── src/lib.rs                           # +GpuAdamWm/v on devicein-place updatehost AdamW 留作参考
├── Cargo.toml                           # xtrain-cuda 升为常规依赖GpuAdamW 要发 kernel
└── tests/adamw_gpu.rs                   # 新GPU AdamW 对 host 参考逐位一致
crates/xtrain-train/
├── src/clip.rs                          # +clip_grad_norm_gpudevice reduce + in-place rescalehost 版留作参考
└── src/train_loop.rs                    # 改用 GpuAdamW + clip_grad_norm_gpu

Key Design Decisions

① cuBLAS matmulrow-major ⟺ col-major

cuBLAS 是 列主序,我们的张量是 行主序。一个行主序 [r,c]、leading dim = c 的矩阵交给 cuBLAS 被读作它的转置(列主序 [c,r])。要拿到行主序结果 C[m,n] = opA(A)·opB(B),就让 cuBLAS 算它的列主序转置 Cᵀ[n,m] = opB(B)ᵀ·opA(A)ᵀ——Cᵀ 列主序的字节布局正好就是 C 行主序。

cublas::sgemm(trans_a, trans_b, m, n, k, …) 落地为:第一参 = Bop = trans_b ? N : T),第二参 = Aop = trans_a ? N : T),尺寸 (m=n, n=m, k=k)lda/ldb/ldc = 各自存储态行主序的列数

let lda = if trans_a { m } else { k };   // A 存 [m,k] 或 [k,m]
let ldb = if trans_b { k } else { n };   // B 存 [k,n] 或 [n,k]
let ldc = n;                              // Cᵀ 是 [n,m] 列主序 ld=n== 行主序 C[m,n]

trans_a=trans_b=false 这一支与 T3 测试里的 cuBLAS oracle 逐参数一致(同样 OP_N、交换顺序、m=N/n=M/k=K所以前向天然对得上。

反向用 cuBLAS 的转置位省两个 transpose kernelT3 版 matmul_backwarddc.matmul(b.transpose_2d()) + a.transpose_2d().matmul(dc),要起两个 transpose kernel + 两次分配。T7 直接:

dA[M,K] = dC[M,N] · Bᵀ   → sgemm(trans_a=false, trans_b=true,  m=M,n=K,k=N, a=dC, b=B)
dB[K,N] = Aᵀ · dC[M,N]   → sgemm(trans_a=true,  trans_b=false, m=K,n=N,k=M, a=A,  b=dC)

为什么不回归:全程 fp32cuBLAS 与手写 tiled kernel 算的是同一个 GEMM只差求和顺序的 rounding。 所以 T3「fwd 对 cuBLAS / bwd 对 finite-diff」的容差不变下游 autograd grad-check、PyTorch 对拍也不变。

handle 持久化cuBLAS handle 创建很贵T3 oracle 每次调用都 create/destroy。改为 每线程缓存一个 handle,进程生命周期内复用(thread_local! + RefCell<Option<CublasHandle>>)。

② GPU AdamW + GPU grad-norm去掉每步全参往返

T6 的瓶颈之一:AdamW::step 把每个参数的 value + grad 全 D2H 拉回 host、host 上跑 AdamW、再 H2D 写回;clip_grad_norm 同理把全部 grad 拉回 host 算范数。3.26M 参数 × 每步两趟 = 大量 PCIe 往返 + 同步。

GpuAdamWm/v 矩状态以「每参一对 device Tensor」常驻显存update 是一个 in-place kernel——读参数的 .grad()、原地改写参数 buffer参数 leaf 的 storage 是 Arc 共享,原地写对所有 clone 可见leaf 身份跨步稳定,无需 set_value

m ← β1·m + (1β1)·g ;  v ← β2·v + (1β2)·g²
p ← p  lr·( (m/bc1) / (√(v/bc2) + ε) + wd·p )      bc1/bc2 = 1βᵗhost 传入)

数学与 host AdamW::step_host 逐字对应host AdamW 原样保留作 PyTorch 对拍的参考,新增 adamw_gpu 测试拿同一组 params/grads 把 GPU 结果对 host 参考逐位比(实测 max abs err = 0

clip_grad_norm_gpusumsq_accum kernel 对每个 grad 做 block-reduce 后 atomicAdd 到一个 device 标量;只把这一个标量取回 host 求 sqrt、算 clip factor再用 scale_inplace kernel 原地把每个 grad 乘 pre_scale·factor。整步只回传 1 个 float不再拉全部 grad。

③ stream / 减 sync

每个 tensor op 之前 Tensor::zeros 分配输出、之后 cudaDeviceSynchronize——两处都是隐藏开销:

  • per-op sync 全删21 处)default stream 上 kernel 顺序执行;任何 host 读数据都走 to_device(Cpu) → 阻塞且 stream-ordered 的 cudaMemcpy,自然等齐前面的 kernel。所以 op 后那次显式 sync 对正确性纯属多余(只是把异步 kernel 错误提前暴露,可接受地推迟到下一次 sync/memcpy
  • device zeros 改 cudaMemset:原来每个 op 输出都用「host 零 buffer + 阻塞 H2D memcpy」置零那次 H2D 本身就是个 per-op 同步点 + 一次拷贝;换成 device 端 cudaMemsetdefault stream 上异步,不串行化 stream

once-per-step 的 syncclip 取范数前、AdamW step 末尾)保留——量级是每步一次,非每 op。

CUDA-graph capture 是 optional bonus本 Phase 未做。

验证方法

两道闸都要过

A. 数值不回归fp32 容差不变,全绿)——dash5 实跑:

测试 结果
T3 GEMMfwd vs cuBLAS / bwd vs finite-diff rel-err 容差不变 5/5 ok
T4 autograd grad-check每 op finite-diff ≤2e-2 不变 12/12 ok
T5 结构 grad-check + overfit + PyTorch 对拍 overfit 27/27、logits relerr、21 参梯度 rtol 不变 overfit 2.821→0.004 (27/27)parity logits relerr 1.5e-4、21 grads OK
T6 AdamW vs torch + checkpoint round-trip 轨迹/终参 rtol 不变、逐位一致 AdamW relerr 4.6e-6ckpt logit diff 0.0
T7 GPU AdamW vs host 参考 逐位一致 max abs err 0.0

B. 吞吐提升——同 model/configdim 32、4 层、vocab 50257、seq 64、batch 8、~3.26M 参60 步计时取稳态:

步骤 tok/s 备注
baseline (T6) ~2770 起点
① cuBLAS matmul ~3310 matmul 非主瓶颈model 小、latency-bound
② GPU AdamW + grad-norm ~4070 去掉每步全参 GPU↔host 往返
③ drop per-op sync + memset ~8220 删 21 处 per-op sync 是大头

端到端real_training 800 步,新快路):~8500 tok/s 稳态loss 10.81→3.90(avg10),采样 Once upon a time, there was a little girl named Lily. She was very happy to play with her mom.—— 收敛与 T6 fp32 同轨。

净加速 ~3×零数值回归。

④⑤ Follow-up本 Phase 未做,记给后续)

  • ④ bf16 混合精度fp32 mastermatmul/激活走 bf16、optimizer 持 fp32 master 拷贝。本 model dim=32 太小、属 launch/latency-boundbf16 tensor-core 算力收益有限,唯一够大的 lm_head [64,32]@[32,50257] 主要吃带宽;且 bf16 改数值、要单独加宽容差 + 重验收敛,风险/收益此规模下不划算。等模型放大或上 T8 多卡再做更值。
  • ⑤ 激活重计算:反向重算 block 激活省显存。当前单序列、显存不紧,优先级低。

两者按 escape hatch 推迟,①②③ 的 fp32 加速已完整交付且全测绿。

T8 衔接(数据并行 all-reduce

T7 之后梯度常驻 device.grad() 是 device tensor优化器 update 也全在 device——这正好对接 T8 的 NCCL 数据并行:

  • 各 rank 本地 backward 后,梯度已在显存里,直接对 params.grad() 张量 all-reduce(无需先拉回 host
  • all-reduce 取 均值后,每 rank 各自跑 GpuAdamW.step——因为各 rank 梯度一致、优化器状态从相同 init 同步演化,参数自然保持一致(无需再同步参数)。
  • grad clip 的 global-norm 在 all-reduce 之后算:clip_grad_norm_gpusumsq_accum 已是 device reduction多卡只需把那个标量再 all-reduce 一次(或对已 all-reduce 的梯度本地算,因梯度已一致,结果天然相同)。