[diffusion] feat: capture-safe pynccl all-to-all - #33775
Conversation
|
/tag-and-rerun-ci |
| """ | ||
| if enable is None: | ||
| enable = self.available | ||
| old_disabled = self.disabled |
There was a problem hiding this comment.
self.disabled is False after successful initialization, so this context restores it to False on exit. The final disabled-communicator check in the new test will therefore execute the all-to-all and fail. Should the communicator be disabled before entering the capture window, or should the test set that state explicitly?
There was a problem hiding this comment.
Right on both counts — and this caught a real verification gap: the transport was probe-verified on GPUs, but this test file itself had not run there (the devbox was reclaimed before it could). This communicator initializes enabled, so the check now sets the state explicitly with change_state(enable=False), and the whole file has now actually run on 2xH200: PYNCCL_A2A_CAPTURE PASS.
| if stream is None: | ||
| stream = current_stream() | ||
| chunk = input_.numel() // self.world_size | ||
| send_counts = input_split_sizes or [chunk] * self.world_size |
There was a problem hiding this comment.
dist.all_to_all_single defines split sizes along dim 0, but this treats them as flattened element counts and validates them against numel(). For an input shaped (4, 3) with splits [1, 3], PyTorch sends 3 and 9 elements, while this code asserts 4 == 12. Could we either restrict this API to 1-D flattened tensors or convert dim-0 splits to element counts and add an uneven 2-D parity test?
There was a problem hiding this comment.
Fixed — splits now convert dim-0 rows to element counts (sum validated against size(0)), so n-D tensors split identically to dist.all_to_all_single, and the test gained an uneven 2-D parity case compared bitwise against the dist result. Passing on 2xH200.
181cd9f to
ea243a1
Compare
Review caught both: split sizes were treated as flattened element counts where dist.all_to_all_single defines them along dim 0 -- rows now convert to element counts so n-D tensors split identically to torch, with an uneven 2-D parity case in the test. And this communicator initializes enabled, so the raise-when-disabled check now sets its state explicitly instead of relying on change_state's exit value. Verified on 2xH200: the full test file passes (bitwise parity, capture+replay, uneven 2-D splits, disabled refusal). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
ea243a1 to
e0680ca
Compare
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Motivation
A CUDA graph cannot hold a
torch.distributedProcessGroupNCCL collective: its host-side per-op bookkeeping advances once at capture and never again, so on replay the ranks stop agreeing about which collective is in flight and hang. This is the reason whole-forward DiT capture (#31852) must refuse sharded runs today, and it is not a NCCL limitation — srt already runs collectives inside captured regions by calling raw nccl through itspyncclwrapper, which carries no per-call host state. The diffusion tree's vendoredpynccllacked both an all-to-all and the capture-window switch.Modifications
pynccl_wrapper.py: bindncclGroupStart/ncclGroupEnd.pynccl.py:all_to_all_singlebuilt from groupedncclSend/ncclRecv(equal and uneven splits, matchingdist.all_to_all_singlesemantics), pluschange_stateso the communicator is live only inside a capture window rather than becoming a mode. A disabled communicator raises instead of silently no-oping, since a silent skip would return whatever the output buffer already held.test_pynccl_a2a_capture_2_gpu.py, registered in the 2-gpu suite: bitwise parity withdist.all_to_all_singleeagerly, then capture once and replay against three fresh inputs, then assert the disabled state refuses.Nothing consumes this in the serving path yet: routing the sequence-parallel all-to-all through it during whole-forward capture is the follow-up on top of #31852, which is what lets
--dit-cuda-graph fullrun sharded (verified end-to-end there: capture succeeds at ulysses 2 and 4, outputs bitwise-identical to eager).Test
4×H200,
world=2andworld=4: eager result bitwise-equal todist.all_to_all_single; captured graph replays correctly for three fresh inputs (PYNCCL_A2A_CAPTURE world=2 PASS,world=4 PASS). The three files here are byte-identical to the branch those runs used; the 2-GPU case is the CI guard.Checklist
🤖 Generated with Claude Code
CI States
Latest PR Test (Base): ✅ Run #31076290463
Latest PR Test (Extra): ❌ Run #31076289925