Repository navigation
π Highlights
-
Dual-CP: inter- + intra-card context parallelism by @Erix025 β FlashQLA now unifies inter-card (cross-GPU) and intra-card CP into a two-level scheme. This extends CP to TP + long-sequence + small-head-count settings and further improves SM utilization.
-
AutoCP v2: a smarter partitioning strategy by @Erix025 β a redesigned auto-CP that decides whether to enable CP and which subsequence length to use, driven by a lightweight per-architecture kernel-latency model and a wave-aware partition search. Adds
infer/trainmodes, SWA gate-decay support, per-arch fitted coefficients (SM90 / SM100), and dedicated documentation undertools/autocp. -
SM120 / SM121 backward pass by @minatoyukinaa β completes Blackwell SM120/121 with a fused backward kernel (previously forward-only), stabilized for auto-CP and variable-length sequences with zero-padding fixes.
π οΈ Additional Improvements
- SM100 (Blackwell) performance β moved the
prepare_hrecurrence totcgen05, improved KKT-solve occupancy and rebalanced backward register caps (@Butterfingrz); async-staged incomplete forward/state tiles and specialized pure full-tile batches (@SakiNana7qi); optimized the fused dQ TMEM path (@byshiue). - New architecture support β SM103 (Blackwell) by @hysmio; SM121 (NVIDIA GB10 / DGX Spark) via SM120 kernels by @zhachar14h-cell.
- Host metadata fast path β reuse of prepared variable-length chunk descriptors to cut CPU overhead for varlen / CP metadata preparation.
- CUDA-graph friendliness β added an sglang-style tensor cache for CUDA graph capture.
- Correctness & stability fixes β improved
dGaccuracy via high-precision matrix transposition; fixed a race condition in cross-partition aggregation; addedfence.proxy.asyncand extra syncs to SM90 / SM100 Blackwell kernels; fixed SM90 backward padding and the autograd interface. - Docs β fixed the low-level API
cp_cacheexample (@eyesofish) and added AutoCP documentation.
π Benchmarks
- Benchmark results on SM90 (H200)
- Benchmark results on SM100 (GB200)
- Benchmark results on SM120 (5090)
- Dual-CP benchmark results on SM90 (H200Γ8)
- Dual-CP benchmark results on SM100 (GB200Γ4)
π Acknowledgements
Special thanks to @Erix025, @minatoyukinaa, @Butterfingrz, @SakiNana7qi, @byshiue, @hysmio, @zhachar14h-cell and @eyesofish for their contributions, and to all community members for their valuable feedback. We welcome your continued participation through issues and pull requests.