Repository navigation
Feat: Context Parallel - #181
Chamberlain0w0 wants to merge 16 commits into
Conversation
81e85ca to
4246e44
Compare
fdf929b to
41c4c95
Compare
39bad6a to
0c6c4d3
Compare
| // COMPUTE_32F alone still permits reduced-precision intermediate reductions | ||
| // for FP16/BF16 outputs. Keep partial sums in FP32 as well. | ||
| CUBLAS_CHECK(cublasSetMathMode( | ||
| blas_handle, CUBLAS_MATH_DISALLOW_REDUCED_PRECISION_REDUCTION)); |
There was a problem hiding this comment.
这里影响默认精度吗,pytorch 脚本是否需要同步修改 torch.backends.cuda.matmul.allow_<dtype>_reduced_precision_reduction
There was a problem hiding this comment.
会影响,需要设置:
torch.backends.cuda.matmul.allow_fp16_reduced_precision_reduction = False
torch.backends.cuda.matmul.allow_bf16_reduced_precision_reduction = FalseThere was a problem hiding this comment.
11.13 更新:不在此 PR 中做任何修改非 CP 路径的数值精度策略,已经恢复
| return x->View({B, H / n_rep, n_rep, T, D})->Sum(2); | ||
| } | ||
|
|
||
| std::shared_ptr<Tensor> ApplyCoreAttention( |
There was a problem hiding this comment.
这个函数最好把 RepeatKVHeads 放在内部执行,因为之后这部分拷贝应该可以通过改 kernel 省下来(可以加个 TODO),移到内部的话可以只改 kernel 调用不改接口。参数可以和 megatron 对齐,传 num_gqa_groups 来推 H_q/H_kv。
https://github.com/NVIDIA/TransformerEngine/blob/c1c2d4d2fa008386f247c2fec56a998e27134e0a/transformer_engine/pytorch/attention/dot_product_attention/dot_product_attention.py#L2781
另外,q@k^T × 1/√D → MaskedFill → Softmax → @v 这套逻辑 causal_self_attention 里也有,这部分能不能直接抽出来改成 ScaledDotProductAttention 共用
There was a problem hiding this comment.
是的,这块 FlashAttn 那个 PR 里是提了一个 ScaledDotProductAttention 的 function 的,但是这个 PR 里主要是想不太动原先的小算子拼接的逻辑。我这块留个 FIXME 吧,我觉得可以放在之后哪个 PR 大改 attn 的时候再顺手改一下?
There was a problem hiding this comment.
我直接在这个 PR 里改了吧。FlashAttn 那个可能还要再往后放,我看了下改动应该不大,主要还是整块整块的结构上的变动
0c6c4d3 to
16f7d18
Compare
16f7d18 to
2712ee2
Compare
|
|
||
| std::shared_ptr<Tensor> ScaledDotProductAttention( | ||
| const std::shared_ptr<Tensor> &q, const std::shared_ptr<Tensor> &k, | ||
| const std::shared_ptr<Tensor> &v, int64_t num_gqa_groups, |
There was a problem hiding this comment.
这里叫 num_gqa_groups, Megatron 同名参数语义是 H_kv,但传的是 n_rep = H_q/H_kv,建议统一一下
Document CP launch options, rank ordering, topology constraints, sequence layouts, communication modes, and test commands. Also update the README parallelism matrix and add a link to the new guide.


概述
本 PR 为 GPT-2、LLaMA 和 Qwen3 引入 Context Parallel(CP),将同一样本的序列分布到多个 rank,降低单卡激活与 attention 中间结果的显存需求,并与现有 TP/SP、PP/VPP 和 ZeRO 训练流程衔接。
主要改动
支持四种 CP 通信模式及其前向、反向传播:
all_gatherp2pa2aa2a+p2pcontiguous布局和可选的zigzag首尾配对布局,后者用于均衡 causal attention 在不同 rank 上的计算量。数值精度
验证
实现边界
当前支持固定长度、等分的 causal self-attention,尚未接入 FlashAttention/fused CP backend,也未覆盖 packed/变长序列和全部并行组合。
FP32 在已验证配置下满足数值容差;BF16 不保证跨并行配置逐位一致,在部分配置下可能会存在累加顺序不一致导致误差较大的情况。