Add comm split support for XCCL #5
Open
Chao1Han wants to merge 490 commits into
Open
Conversation
…eta-pytorch#577) Summary: Pull Request resolved: meta-pytorch#577 This diff replaces the control message mechanism for NVL memory release with an AsyncSocket-based approach: 1. AsyncSocket-based IPC Release: Changed remReleaseMem() to use notifyRemoteIpcRelease() via AsyncSocket instead of isendCtrlMsg() over CtranIB/CtranSocket backends. 2. Server Address Exchange: Added bootstrap allGather to exchange IPC server addresses (peerIpcServerAddrs_) across all ranks during mapper initialization, enabling direct AsyncSocket communication between any pair of ranks. 3. Request Tracking: Replaced CbCtrlRequest class with IpcReqCb for tracking pending release requests. 4. Removed Callback Registration: Removed releaseMemCb callback and its registration for NVL_RELEASE_MEM control message type. The IpcRegCache now handles incoming release requests directly via its AsyncSocket server. The workflow of `exportMem` and `releaseMem` is as follows. Now `exportMem` and `releaseMem` could arrive at destination out of order. The reason of using AsyncSocket for `releaseMem` is that `IpcRegCache` will eventually be a singleton, and AsyncSocket can ensure all `releaseMem` arrives at peers reliably. {F1985212617} Reviewed By: dsjohns2 Differential Revision: D92077053 fbshipit-source-id: 837c694973946ad89cc4784faad92cf4426608dc
Summary: Pull Request resolved: meta-pytorch#521 Refactor common Scuba file utilities into a shared library that can be used by both NCCLX DataTable and MCCL McclDataTable: - Add ScubaFileUtils.h/.cpp with getScubaFileName(), createScubaFile(), and getTimestamp() - Add scuba_file_utils BUCK target with shared dependencies - Add data_sink BUCK target for DataSink.h header-only library - Update DataTable.cc to use the new ScubaFileUtils functions - Update data_table BUCK target to depend on scuba_file_utils This enables code reuse and eliminates duplicate implementations of file-related utilities between the NCCLX and MCCL logging systems. Reviewed By: arttianezhu Differential Revision: D92017771 fbshipit-source-id: 1288a82f0e88e8255bca2b9ed0ac049a39d8705a
…ontained build (meta-pytorch#520) Summary: Pull Request resolved: meta-pytorch#520 This diff implements a pre-extracted sources snapshot system for rcclx that eliminates ABI mismatch issues by building rcclx from source at build time instead of using pre-compiled bundled dependencies. The key change replaces the previous Manifold-based snapshot approach with a simpler file-based approach where rcclx sources from specific commits are extracted and committed to the repository under comms/rcclx/snapshots/{stable,last-stable}/. The create_snapshot.py script uses sl archive to extract sources from specified commits, then automatically fixes absolute Buck load paths (e.g., fbcode//comms/rcclx/develop:def_build.bzl) to point within the snapshot directory, ensuring each snapshot is self-contained. Additionally, the script adds compiler flags (-Wno-deprecated-this-capture, -Wno-unused-exception-parameter) to suppress warnings that weren't errors when the snapshot code was written but are now treated as errors with current compiler settings. The BUCK file defines rcclx-stable and rcclx-last-stable aliases that point to the snapshot's rcclx-dev target, allowing downstream consumers to build against stable rcclx versions with guaranteed ABI compatibility. Obsolete Manifold-related scripts have been removed as they are no longer needed. Reviewed By: joebos, dmwu Differential Revision: D91904438 fbshipit-source-id: 87baf0ceb2363221c0aee110b3eb8f4de4c438a6
Summary: Pull Request resolved: meta-pytorch#458 As title, removes the dependency on the `ruamel.yaml` Python module for the `extractcvars.py` library. This is ultimately to enable D88748045, as the dependency on `ruamel.yaml` was proving to be problematic: the NCCLX GitHub OSS build was breaking due to the `ruamel.yaml` not being present, and adding the `ruamel.yaml` module to the build environment was unexpectedly not straight-forward. Reviewed By: function47 Differential Revision: D91694592 fbshipit-source-id: f8daaf6329f461f4392c8b030f1b79ca6e4d2505
Summary: Pull Request resolved: meta-pytorch#84 **TODO**: If my cmake-related fixes work, then they should be split into a separate diff before landing. This diff implements a genrule-based solution for generating `nccl_cvars.h` and `nccl_cvars.cc` at build-time. This is intended to eliminate the need to commit these auto-generated files and consequently prevent the frequent rebase conflicts that stemmed from the `SignedSource` updates within these files. **Header of `nccl_cvars.(h|cc)`**: ``` c++ // Copyright (c) Meta Platforms, Inc. and affiliates. // generated SignedSource<<450242072b5f4fa514b07a144a620779>> // Automatically generated by ./comms/utils/cvars/extractcvars.py --- START // DO NOT EDIT!!! ``` The `generated SignedSource...` bit is the culprit of the frequent rebase conflicts. This diff also removes the `nccl_cvars.cc` and `nccl_cvars.h` files now that they'll instead be built/generated/provided by `genrule`. Finally, this diff also updates several build scripts to be compatible with the new `genrule`-based build for the `ncclx-cvars` library: - MCCL build script, `comms/mccl/build/build.sh` - `comms/ncclx/v2_27/maint/oss_build.sh` - `comms/ncclx/v2_28/maint/oss_build.sh` - `comms/github/build_rcclx.sh` - `comms/github/build_ncclx.sh` - `conda/feedstock/nccl/recipe.yaml` **Note**: this diff was originally published, landed, and reverted as D87668052. Reviewed By: tanquer Differential Revision: D88748045 fbshipit-source-id: 13bc81abe24d0bc7ddf1385d003523fa632d5d5d
Summary: Pull Request resolved: meta-pytorch#575 This was backed out (D91603922) when the original Pytorch caffe2 Hipify removal diff was backed out. Turns out we need a little more changes than the diff above in order to correctly run the unit tests. Now that the caffe2 diff is in again (meta-pytorch#544), we need to check this in again. Reviewed By: atalman Differential Revision: D92454446 fbshipit-source-id: a794862662f06519c80108182ff00a599e6e2692
…t mapper destruction (meta-pytorch#579) Summary: Pull Request resolved: meta-pytorch#579 ### Convert IpcRegCache to singleton and add server address exchange **Key Changes:** ``` 1. IpcRegCache Singleton Pattern - Converted IpcRegCache from per-mapper instance to singleton shared by all communicators - Added getInstance() static method using folly::Singleton - Made init() idempotent with CUDA device validation - throws exception if different CUDA devices attempt to use the same singleton 2. Remove Per-Communicator State - Removed logMetaData_ member variable from IpcRegCache - Made logMetaData an optional parameter (default nullptr) for importMem() - Each communicator now passes its own logMetaData when needed 3. Server Address Exchange - Added allGatherIpcServerAddrs() in CtranMapper to exchange AsyncSocket server addresses via bootstrap - Added peerIpcServerAddrs_ storage and getPeerIpcServerAddr() accessor in CtranMapper - Fixed remReleaseMem() to use the gathered peer addresses for async socket communication 4. Remote Memory Release Fix - Removed the atDestruction check in remReleaseMem() - Now always sends release messages even during destruction since IpcRegCache is a singleton that outlives individual mappers 5. Test Updates - Updated all tests to use IpcRegCache::getInstance() instead of ipcRegCachePtr() - Updated importMem() calls to use new parameter order with optional logMetaData ``` **Benefits** ``` - Reduces memory usage by sharing IpcRegCache across communicators - Properly exchanges server addresses needed for async socket communication - Ensures consistent CUDA device usage across all communicators - Fixes potential memory leaks by always releasing remote memory ``` Reviewed By: dsjohns2 Differential Revision: D92090165 fbshipit-source-id: 0623326bbd5944aee611bceb0604ac47e6267be6
Summary: Pull Request resolved: meta-pytorch#574 split_group on N-D mesh with GLOO will fail on key collision on the store. This should fix it. Reviewed By: d4l3k Differential Revision: D92429662 fbshipit-source-id: dd882cf85089fddcf572298c4e629030ce045f39
Summary: Pull Request resolved: meta-pytorch#484 In TBR, we need a way to mock RDMA transport error by raise exception and properly clean up leftover transport resources: ``` SERVER (co_fetchSliceRdma) │ │◄─── FetchSliceRequestRdma │ ├─── RDMA Write 1 ──── error handling ───┐ ┐ ├─── RDMA Write 2 ──── error handling ───┼─ per-write │ timeout(60s) ├─── RDMA Write N ──── error handling ───┘ ┘ (collective) │ │ On ANY error: │ 1. transport->abort() ◄── cancel pending writes │ 2. removeRdmaTransport() ◄── cleanup broken transport │ 3. throw RdmaTransportError (can be captured?) │ │───► FetchSliceResponseRdma (or error) ▼ ``` In this diff - Add optional timeout parameter to write() API for future timeout support (currently only used in mock testing) - Add MockType enum (None, Timeout, Failure) and MockContext struct with hasTimedOut() helper - Add setMockForTest() to inject mock behavior for testing error scenarios - Add abort() to cancel all pending operations with commUserCancel status - Add commTimeout and commUserCancel error codes to commResult_t enum - Mock timeout uses duration from write() parameter, defaults to 60 seconds if not specified - Add unit tests for failure, timeout (immediate, with duration, multiple durations) and abort functionality Reviewed By: tianfengfrank Differential Revision: D91275671 fbshipit-source-id: a5afe41c3ceaeb1f4fe449c1ca5eeda3574d945b
Summary: Pull Request resolved: meta-pytorch#571 Replace FB_COMMCHECKTHROW(ncclToMetaComm(...)) with NCCLCHECK_TEST(...) in ncclx AllReduce tests. FB_COMMCHECKTHROW is defined in comms/utils/checks.h but the test files only include the local "checks.h" which doesn't have this macro. Affected tests: - v2_27: AllReduceTest, AllReduceNumericStableTest, AllReduceUniformTest, AllReduceNumericOffsetTest - v2_28: AllReduceTest, AllReduceNumericStableTest, AllReduceUniformTest, AllReduceNumericOffsetTest TODO(T254162415): AllReduceUniformTest is disabled due to mpirun runtime failure. The test has been broken since 2025-11-06 (codemod c6494ace57c2 removed the checks.h include). Reviewed By: pavanbalaji Differential Revision: D92410157 fbshipit-source-id: 47477173a627f56adc1095d67494ce0f7ce2bdcc
…torch#586) Summary: Pull Request resolved: meta-pytorch#586 The AllReduceARG algorithm has a fundamental design flaw that causes deadlocks when running with multiple ranks. The root cause is contention between the main thread and GPE thread over the epochLock mutex: 1. Main thread calls exchangeUserbuff() which acquires epochLock and performs cross-rank communication (allGatherCtrl) 2. GPE thread needs epochLock to run opIbImpl() but is blocked 3. Kernel waits for GPE to post work via GpeKernelSync 4. This creates a deadlock when Rank 0's main thread waits for Rank 1 in allGatherCtrl, while Rank 1's GPE waits at kernel barrier The algorithm cannot be easily fixed because: - Moving exchangeUserbuff to GPE doesn't work: kernel needs exchange output (intraNodeRemoteRecvBuffs) at launch time, before GPE runs - Even with exchange in GPE, main thread can call exchangeUserbuff for NEXT operation concurrently, causing IB hang - Dynamic argument passing would require major architectural changes that conflict with the kernel elem cleanup design Reviewed By: pavanbalaji Differential Revision: D92428892 fbshipit-source-id: 077625516711bb3a16ecbc01f4377cdab028c85c
…ration (meta-pytorch#576) Summary: Pull Request resolved: meta-pytorch#576 Enable NCCLX GIN device headers to compile with Clang for LLVM bitcode generation. This is needed to link NCCLX GIN templates with Triton-compiled kernels at runtime via cuLink APIs. Key Changes: utility.h: Added NCCL_CHECK_CUDACC macro that correctly detects CUDA compiler mode for both nvcc and Clang. Clang defines the CUDACC macro as empty, so "#if" checks fail but "#ifdef" checks work. Also added: NCCL_IR_EXTERN_C macro for extern "C" linkage when compiling to LLVM bitcode Use attribute((always_inline)) instead of forceinline when compiling with Clang (Clang doesn't support forceinline) All other headers: Replaced the nvcc-style CUDACC checks with NCCL_CHECK_CUDACC throughout: barrier.h, coop.h, core.h, gin.h, gin_barrier.h, ll_a2a.h, lsa_barrier.h, ptr.h gin/gin_device_common.h All impl/ headers (funcs and types files) Why: Triton compiles kernels to LLVM IR using Clang. To call NCCLX GIN operations from Triton kernels, we need to compile the NCCLX device headers to LLVM bitcode and link them at runtime. The existing headers used preprocessor checks that fail under Clang because Clang defines the CUDACC macro as empty rather than 1. Reviewed By: siyengar Differential Revision: D92456185 fbshipit-source-id: 50df7128e4b20fe99a6d865472f3542125963cab
Summary: Pull Request resolved: meta-pytorch#587 When using dist.split_group with NCCL-based backends (nccl, ncclx, rccl, rcclx), force split_share=0 to match the behavior in comms_wrapper.py. This ensures consistent process group creation behavior across both implementations. If pg_options is None and an NCCL backend is used, create a new ProcessGroupNCCL.Options with split_share=0. If pg_options is already a ProcessGroupNCCL.Options instance, set split_share=0 on it. Reviewed By: cenzhaometa Differential Revision: D92472355 fbshipit-source-id: 5761d1578d46faf7c87af134136e7b1b391d8224
Summary: Pull Request resolved: meta-pytorch#588 Update split_group function signature to match the PyTorch torch.distributed.split_group API: - Rename parameter 'group' to 'parent_pg' - Change return type from ProcessGroup to ProcessGroup | None Reviewed By: cenzhaometa Differential Revision: D92472357 fbshipit-source-id: ca357459d0b838ed0475e92b3638d6d32ea060b2
Summary: Pull Request resolved: meta-pytorch#589 Port the _make_nccl_premul_sum functionality from `comms/comms_wrapper.py` to `torchcomms/distwrap/collectives.py`. This function creates a NCCL PREMUL_SUM reduce operation that multiplies inputs by a factor before summing. Changes: - Add _PREMUL_SUM_OPS dict to track dist ops to torchcomms op mappings - Update _convert_reduce_op to check for PREMUL_SUM ops - Add _make_nccl_premul_sum function to create and register PREMUL_SUM ops Reviewed By: cenzhaometa Differential Revision: D92472358 fbshipit-source-id: 8d670b67dc892d9511883f82dc15d7dd8645d38d
Summary: Pull Request resolved: meta-pytorch#590 Port get_mem_allocator and register_mem_pool from `comms/comms_wrapper.py` to `torchcomms/distwrap/collectives.py`. - get_mem_allocator: Returns the memory allocator for a process group and device, using torchcomms instance when enabled or falling back to the PG backend. - register_mem_pool: Registers a memory pool with the backend. This is a no-op when torchcomms is enabled since it manages its own memory pools. Reviewed By: cenzhaometa Differential Revision: D92472356 fbshipit-source-id: 5a45ef777871b1dbfe66028fb4c5f897e51b5c3f
…torch#597) Summary: Pull Request resolved: meta-pytorch#597 In NCCL v2.28, ncclTopoGetVNicParent calls ncclTopoWidenLinks to sum the PCI link widths of constituent physical NICs when placing a bonded vNic in the XML topology. This causes the graph search to see doubled PCI bandwidth (e.g. 96 -> 192) at the bonded NIC's PCI path, which in turn causes the search to route rings through the individual physical NICs (e.g. mlx5_0, mlx5_1) alongside the bonded vNic (e.g. mlx5_0+mlx5_1). Since the bonded vNic and its constituent NICs are the same physical hardware, this just creates more channels without adding real bandwidth. This was observed by comparing NCCL topology logs for two comms-bench jobs with identical placement (32 ranks, 16 nodes, 2 local ranks per node) but different NCCL versions: - comms-bench-test-f3sdvz4w (NCCL 2.27): split communicator uses only bonded vNics (IB/4, IB/5) with 4 channels. PCI bandwidth reported as 96. - comms-bench-test-gbr5jxlj (NCCL 2.28): split communicator uses bonded vNics plus individual NICs (IB/0, IB/1, IB/4 for rank 0) with 8 channels. PCI bandwidth reported as 192 due to ncclTopoWidenLinks summing the link widths. In v2.27, ncclTopoGetVNicParent does not call ncclTopoWidenLinks, so the graph search sees the original PCI bandwidth and only routes through bonded vNics. This change guards the ncclTopoWidenLinks call behind the NCCL_NET_WIDEN_LINKS environment variable. By default (env var unset), link widening remains enabled (preserving current v2.28 behavior). Setting NCCL_NET_WIDEN_LINKS=0 disables it, reverting to v2.27-like behavior where only bonded vNics are used. Reviewed By: SuhitK Differential Revision: D92668667 fbshipit-source-id: 372dcb9a29571e9c4c4675a8f48e782b99a1e6ee
Summary: Pull Request resolved: meta-pytorch#540 1) Add foundational buffer types for IBGDA 2) Add device-side per-peer RDMA transport class that provides GPU-initiated RDMA operations using DOCA GPUNetIO high-level APIs: - put_signal(): RDMA Write with atomic signal for data transfer + notification - put(): RDMA Write without signal - signal(): Atomic signal only (no data transfer) - wait_local(): Wait for local completion (data sent to NIC) - wait_signal(): Spin-wait for remote signal arrival Includes component-level unit tests. Reviewed By: siyengar, GirasoleY, cenzhaometa Differential Revision: D92107374 fbshipit-source-id: 9fc63cf1435885bbd51272d1e65bf2b14bb338ea
…meta-pytorch#596) Summary: Pull Request resolved: meta-pytorch#596 ## Summary D92104317 added `ncclCommGetAsyncError()` polling in the TorchComms timeout watchdog thread to detect asynchronous NCCL errors. This broke the `DeviceApiTest.DeviceWindowCreationFloat` test when GIN (GPU-Initiated Networking) is enabled with GDAKI (`NCCL_GIN_TYPE=3`). ### Root Cause `ncclCommGetAsyncError()` internally calls `ncclGinQueryLastError()`, which iterates over all GIN contexts and queries GDAKI QP error states via `doca_gpu_verbs_query_last_error()`. The problem is that GIN context creation (`ncclGinConnectOnce` → `ncclGinGdakiCreateContext`) happens lazily — triggered inside any collective call (`ncclAllReduce`, etc.) that needs new transport channels. During this initialization, IB QPs are still transitioning through states (INIT → RTR → RTS) and `doca_gpu_verbs_query_last_error()` reports spurious errors that surface as `ncclRemoteError`. Since the watchdog thread runs concurrently and polls `ncclCommGetAsyncError()` on a timer, it races with GIN context creation happening on the main thread inside NCCL collective calls. The watchdog sees the false-positive `ncclRemoteError` and calls `abort()`, crashing all 8 ranks. ### Fix Gate the GIN error query block in `ncclCommGetAsyncError()` on `ginState->connected`, which is only set to `true` at the end of `ncclGinConnectOnce()` after all GIN contexts are fully created and QPs have completed their IB state transitions. This skips the query entirely during the vulnerable initialization window while still checking for real GIN errors once the subsystem is ready. Reviewed By: d4l3k Differential Revision: D92623789 fbshipit-source-id: 7ace969103c3098a6e97254c91ace40cbc8140d0
Summary: Pull Request resolved: meta-pytorch#541 Add device-side multi-peer transport wrapper that manages access to per-peer P2pIbgdaTransportDevice handles This just define the data structure and provide simple utility functions (index to rank mapping etc.) UTs in this diff are simple. API level UTs are in next diff. Reviewed By: siyengar Differential Revision: D92107376 fbshipit-source-id: 40d8b2e8742fc96d614b979f279d3f6b84355930
Summary: Pull Request resolved: meta-pytorch#542 Add host-side multi-peer IBGDA transport manager that handles the full connection lifecycle: - DOCA GPU device initialization and IB device discovery - GPU memory allocation and registration with IB verbs - Queue Pair (QP) creation and configuration for each peer - Out-of-band connection establishment via bootstrap exchange - Device transport construction for use in CUDA kernels The host creates and manages all resources, then exports device handles for GPU kernels to perform RDMA operations directly. Includes comprehensive distributed integration tests (requires 2 MPI ranks) covering: basic construction, put/signal operations, various transfer sizes (1KB-16MB), bidirectional transfers, and stress testing. Reviewed By: siyengar Differential Revision: D92107372 fbshipit-source-id: 885dc2fc876b59907f14810c65daa0c20ca74640
Summary: Pull Request resolved: meta-pytorch#543 Add performance benchmarks for IBGDA transport operations: - Bandwidth benchmarks: Measure put_signal throughput across message sizes (8B to 128MB) - Latency benchmarks: Measure signal-only round-trip latency - Supports configurable iterations and warmup runs - Results reported with min/max/avg statistics Reviewed By: cenzhaometa Differential Revision: D92107373 fbshipit-source-id: a3584ffad06e374615b96a7444634f4dff2f33ab
…pytorch#573) Summary: Pull Request resolved: meta-pytorch#573 $title Reviewed By: cenzhaometa, siyengar Differential Revision: D92364605 fbshipit-source-id: 1f703c5b021d2d652ec5a2b58e5de3dd493e6892
Summary: Pull Request resolved: meta-pytorch#602 Enable Claude AI to perform specialized code reviews for the Pipes library by providing domain-specific context and evaluation criteria. This skill helps ensure diffs touching Pipes code are reviewed against library-specific conventions (naming patterns, memory ordering semantics, ThreadGroup usage, transport patterns) and common pitfalls in GPU communication primitives code. Once committed, the command can be run by any Claude Code instance launched from `comms/pipes` via `/pipes-diff-review <one or more diff IDs>`. **Note**: this is a joint effort with cenzhaometa and dmwu. We've all been discussing how best to maximize productivity with AI tools, including sharing templates, skills, and synchronizing them between developers on the team. Reviewed By: siyengar Differential Revision: D92566303 fbshipit-source-id: c43a0b1699a7e423fbc15bf152564b583ab146c9
Summary: Pull Request resolved: meta-pytorch#604 Adds the packaging library as a declared dependency for the torchcomms wheel to ensure version parsing code has its required module at install time. This prevents build/install issues caused by the missing packaging dependency referenced in setup.py. ``` tristanr@devvm5553 ~/torchcomms (main)> MASTER_ADDR=localhost MASTER_PORT=0 RANK=0 WORLD_SIZE=1 python -c "import torch; import torch.distributed as dist; import torchcomms; comm = torchcomms.new_comm('ncclx', torch.device('cuda'), store=None, name='1234') ; t = torch.zeros(10, device=comm.get_device()); comm.all_reduce(t, torchcomms.ReduceOp.SUM, async_op=False)" Traceback (most recent call last): File "<string>", line 1, in <module> import torch; import torch.distributed as dist; import torchcomms; comm = torchcomms.new_comm('ncclx', torch.device('cuda'), store=None, name='1234'); t = torch.zeros(10, device=comm.get_device()); comm.all_reduce(t, torchcomms.ReduceOp.SUM, async_op=False) ^^^^^^^^^^^^^^^^^ File "/home/tristanr/.conda/envs/torchcomms-nightly-3.13/lib/python3.13/site-packages/torchcomms/__init__.py", line 12, in <module> from torchcomms.functional import is_torch_compile_supported_and_enabled File "/home/tristanr/.conda/envs/torchcomms-nightly-3.13/lib/python3.13/site-packages/torchcomms/functional/__init__.py", line 5, in <module> from packaging import version ModuleNotFoundError: No module named 'packaging' ``` Session: DEV58746194 Reviewed By: amirafzali, dolpm Differential Revision: D92747877 fbshipit-source-id: d512abd5a072cbdf6840f468c78f4eabdfe27b11
…#600) Summary: Pull Request resolved: meta-pytorch#600 As part of T238821628, we are migrating from `std::runtime_error` to `ctran::utils::Exception` throughout CTRAN. This diff adds a custom clang-tidy check cAST that will detect and ban new usages of `std::runtime_error` in the CTRAN library (`comms/ctran/*`). This provides compile-time enforcement of the migration policy. See T247828833 for more details. ## Phabricator Warnings The warnings indicate that the clang-tidy binary used by Phabricator (`clang-tidy-platform010-clang-19`) doesn't recognize the checks `facebook-ban-non-ctran-exceptions` and `facebook-ban-runtime-error-in-ctran`. IIUC, this is basically a timing/deployment problem. This diff adds new clang-tidy checks, but the clang-tidy binary that Phabricator uses hasn't been rebuilt yet to include these new checks. Until that binary is rebuilt and deployed, Phabricator will warn about unrecognized checks. So, once the diff is landed and the clang-tidy toolchain is rebuilt, Phabricator will recognize the checks. ## A Note on Existing Uses of Usages of `std::runtime_error` All existing `std::runtime_error` references within `comms/ctran` are in `catch` clauses (receiving exceptions thrown by external libraries), not in constructor expressions or throw statements. Neither clang-tidy check will fire on them, so these files can continue to catch `std::runtime_error` from other libraries without triggering the ban. Reviewed By: arttianezhu, minsii Differential Revision: D88646705 fbshipit-source-id: edaff458b8a2a12a4bb924dd48c22a96e2093557
…rch#606) Summary: Pull Request resolved: meta-pytorch#606 Refactor CudaGraphsTest.py to create a reusable CUDA graph testing framework and add new test cases for various allreduce patterns. Changes: - Add _run_graph_test() framework method for graph capture/replay/verify - Add try-finally for graph.reset() to ensure proper cleanup before communicator destruction (required by NCCL 2.26+) - Use unique comm names per test to avoid store prefix collision - Handle self.comm1 cleanup in tearDown() for tests needing two comms - Add multi-stream tests with proper branch/rejoin pattern using stream.wait_stream() for dependencies All 8 tests pass: test_single_allreduce_sync, test_single_allreduce_async, test_multiple_allreduce_sync, test_multiple_allreduce_async_wait_at_end, test_multiple_allreduce_mixed, test_multiple_streams_single_comm, test_multiple_streams_multiple_comms, test_two_streams_two_comms_with_dependency Note: cuda event query in timeoutWatchdog needs additional study. Reviewed By: dolpm Differential Revision: D92558623 fbshipit-source-id: ba00cccb8c7128f2502f8690e888c3e1edf4b001
Summary: Pull Request resolved: meta-pytorch#605 Fix integration tests by addressing two issues: 1. Use `comm.get_size()` instead of `torch.cuda.device_count()` to determine world_size in DPTPCommTest and DeviceMeshTest. This ensures the world size reflects the actual communicator group size rather than the local GPU count, which may differ in multi-node or partial-GPU configurations. 2. In MemPoolTest, move the torchcomms.get_mem_allocator() call to after torchcomms.new_comm(), since mempool-related symbols are lazy-loaded at runtime during comm initialization. Reviewed By: d4l3k Differential Revision: D92738739 fbshipit-source-id: 261b571b2b2c1592fbd5ac649236ee3bf52bdf0c
Summary: Pull Request resolved: meta-pytorch#591 This adds a test utility helper class (VerifyAlgoStatsHelper) that enables algorithm verification in NCCL tests via the AlgoStats infrastructure. The utility: - Sets NCCL_ALGOSTATS_ENABLE=1 to enable algorithm statistics collection - Provides verify() method to check if a specific algorithm was used - Supports substring matching for algorithm names (e.g., "PAT", "Ring") This helper is standalone and can be used by any test that needs to verify which algorithm was selected for a collective operation. Reviewed By: YulunW Differential Revision: D92487189 fbshipit-source-id: daa1584d9b8c34b94464b3b042ec1ef8a697425a
…orch#611) Summary: Pull Request resolved: meta-pytorch#611 Reviewed By: dsjohns2 Differential Revision: D92612020 fbshipit-source-id: 6bd9de05c5ae5b468d05f8521623c2a5a792b13d
Summary: Pull Request resolved: meta-pytorch#798 - Wrap long lines in ncclx/CMakeLists.txt - Use ```text fenced code blocks in README.md files Reviewed By: kapilsh Differential Revision: D94061925 fbshipit-source-id: 142236941e6e71b5f5ed6d22d18c2b2daf20a98b
Summary: Pull Request resolved: meta-pytorch#796 The inner loop in the read_sizes lambda used `i` as its loop variable, shadowing the outer loop's `i` from line 429. Since the lambda captures by reference, this could cause incorrect behavior if the outer `i` were accidentally referenced inside the lambda. Rename the inner variable to `j`. Reviewed By: kapilsh Differential Revision: D94061931 fbshipit-source-id: 792ede40f01f0e9f5953b5f7da51b76c76920687
Summary: Pull Request resolved: meta-pytorch#795 Fix CLANGTIDY and INFERMINICPP lint warnings: - TorchCommMCCL.cpp: Use a const empty vector to avoid implicit copies in conditional operator expressions (facebook-conditional-operator-argument-copy). - AllToAllvSingleTest.cpp: Use const auto& instead of copying split_sizes.first into input_split_sizes (PULSE_UNNECESSARY_COPY). Reviewed By: kapilsh Differential Revision: D94061918 fbshipit-source-id: 21bb2df5eae708f28009fe94cf9815fbd523f228
Summary: Pull Request resolved: meta-pytorch#799 Fix PYREMODE lint warnings across torchcomms Python files: - Remove `# pyre-unsafe` headers from 15 files in non-type-checked projects where the header is unnecessary and causes confusing results. - Add missing `# pyre-unsafe` header to FlightRecorderTest.py and test_triton_api.py which are type-checked by Pyre but were missing the required mode header. Reviewed By: kapilsh Differential Revision: D94061915 fbshipit-source-id: 0c3dbc9edae939ee7bc4c3e688ef074681e5f692
…oc only) (meta-pytorch#653) Summary: Pull Request resolved: meta-pytorch#653 In this diff: - Fix the issue of CtranIpc.cc on `cudaMalloc` memType, since `cuMem` API is not yet supported on AMD - `shouldSupportCudaMalloc` should be true on AMD platform ----------------------- Test on AllGatherP: On AMD (use `cudaMalloc`): - Before this diff, - single-node perf: P2162362269 --> AGP pipeline @ 13836.803 GB/s --> exceeding theoretical xGMI peak bandwidth --> why? --> icopy() is skipped --> AGP is not performed at all - After this diff, - single-node perf: P2162378472 --> AGP pipeline @ 60.670 GB/s --> icopy() is performed On NV (use `cuMem`): - Before this diff: - single-node perf: P2195483290 --> AGP pipeline @ 374.535 ==> works fine, no changes needed Reviewed By: dmwu Differential Revision: D91988270 fbshipit-source-id: 0448ce63f41d538c5d707f740aafd182389695bd
…ch#848) Summary: Pull Request resolved: meta-pytorch#848 The AllgatherP benchmark (`allgatherp_bench`) was crashing on all ranks with: ``` *** The MPI_Bcast() function was called before MPI_INIT was invoked. ``` After D872aa6f (ctran test infra decoupling), distributed tests run in TCPStore mode via remote execution, where `MPI_Init()` is never called. However, the benchmark still used raw MPI calls (`MPI_Bcast`, `MPI_Allreduce`, `MPI_Barrier`) which require MPI to be initialized. This diff replaces all direct MPI calls with the bootstrap abstraction (`ctranComm_->bootstrap_->allGather()` and `barrier()`), which works transparently in both MPI and TCPStore modes. It also increases the RE timeout from the default 120s to 300s to accommodate the benchmark's workload. Code history: 1. Dingming Wu (dmwu@meta.com) on 2025-12-14: Created the benchmark (D88081175: [ctran-amd] Add Perf benchmarks for AllgatherP) — the original code used MPI_Bcast/MPI_Allreduce/MPI_Barrier, which worked when tests were launched via mpirun. 2. Min Si (msi@meta.com) on 2026-01-15: D90475557: [ctran][test] Decouple ctran_dist_test_utils — this refactored the test infrastructure to support both MPI and TCPStore modes, with auto-detection via getInitEnvType(). This is what switched the test launch method from MPI to TCPStore for remote execution environments. 3. Zhiyong Wang (zhiyongww@meta.com) on 2026-02-20: D93835157: [testinfra] Fix multinode RE TCPStore connectivity by using container IP — further TCPStore fixes. Reviewed By: dmwu Differential Revision: D94288372 fbshipit-source-id: 2d46ebc569e73cd869ca1b6e9f27a4ccdca0a410
…ASYNC_SOCKET CVar (meta-pytorch#835) Summary: Pull Request resolved: meta-pytorch#835 Add CVar `NCCL_CTRAN_IPC_REGCACHE_ENABLE_ASYNC_SOCKET` (default: true) that can disable the socket-based parts of IpcRegCache: - AsyncSocket server initialization - AllGather of IPC server addresses (peerIpcServerAddrs_) - Remote release notifications via socket (notifyRemoteIpcRelease) This provides a kill-switch to mitigate crash in a recent PAFT job. We will follow up to fix the real issue as well. Note that NVL for windows continues to work because: - Window's NVL IPC descriptor exchange uses the IB/Socket control channel (isendCtrlImpl/irecvCtrlImpl), NOT the IpcRegCache socket server - Window cleanup uses skipRemRelease=true + barrier, so it doesn't need socket-based release notifications - IpcRegCache::regMem(), exportMem(), importMem(), and releaseRemReg() remain fully functional Reviewed By: dboyda, minsii Differential Revision: D94270205 fbshipit-source-id: 7385f05f10b554ab375e6044d8745edbac971d6a
…ytorch#843) Summary: Pull Request resolved: meta-pytorch#843 The DistRegCacheUT tests called notifyRemoteIpcExport and notifyRemoteIpcRelease without checking return values. If the async socket is not initialized (e.g., NCCL_CTRAN_IPC_REGCACHE_ENABLE_ASYNC_SOCKET=0), these functions return commInternalError but the error was silently swallowed, causing sender ranks to proceed while receiver ranks spin forever in std::this_thread::yield() — resulting in a hang instead of a clear error. Wrap all notifyRemoteIpcExport and notifyRemoteIpcRelease calls with COMMCHECK_TEST() so that failures are caught and reported immediately. Reviewed By: elvinlife Differential Revision: D94394698 fbshipit-source-id: b5423e6f252ee495b8227e827194de481eaca48b
…eta-pytorch#837) Summary: Pull Request resolved: meta-pytorch#837 title - Explicitly setting to true when needed like in DistRegCache unit tests. - Update McclCommUT to expect correct num all gathers since no longer calling serveraddr all gather. Left a note to increase when we get rid of cvar. Reviewed By: dboyda Differential Revision: D94369864 fbshipit-source-id: a4270bfd79cb5f593df2589e6270265476e8c684
Summary: Pull Request resolved: meta-pytorch#850 As in title. Reviewed By: t2h-fb Differential Revision: D94414277 fbshipit-source-id: 0417901adc73e73fca08e96e8e2b78b02e5c8867
Summary: Pull Request resolved: meta-pytorch#827 Fix misaligned tmpbuf accesses, the immediate reason is from non-power-of-2 rank counts. Auto-tune computes `chunkSize = partitionMessageBytes / (pipelineDepth * nRanks)`. There are a few opportunities for this to be non-pow2: - When nRanks is non-pow2 - when partitionMessageBytes is e.g. 28, pipelineDepth is 4 Two-layer fix: - Round nRanks to nearest pow2 inside auto-tune for the numChunks calculation, so chunkSize is always pow2 (structurally typeSize-aligned, exact BDP fit). - This should give us more aligned performance (we will followup with measurements here) - Post-hoc alignment in AllReduceRing.cc as safety net for CVAR overrides that bypass auto-tune and can inject arbitrary non-aligned chunkSize values. The invariant test cases should always be valid. Note, We might be removing nRanks factors later, but not waiting on that for the fix. Reviewed By: saifhhasan Differential Revision: D94245412 fbshipit-source-id: 8cf9fe548521b6908d01f84e2be3baaaa59d0099
…ture (meta-pytorch#829) Summary: Pull Request resolved: meta-pytorch#829 Remove the check in enqueueWork() that required abort_process_on_timeout_or_error=true during CUDA graph capture, along with its unit test GraphCaptureRejectsNoAbortMode. - Upcoming diffs will make abort mode the only supported mode, so this guard is no longer needed. Reviewed By: tanquer Differential Revision: D94266658 fbshipit-source-id: f1707e2f76200a16ecf8e8e227f43426c6fa1930
…operly serve as cap (meta-pytorch#828) Summary: Pull Request resolved: meta-pytorch#828 As title. The renames help to prepare for clarity with the next diff. Reviewed By: saifhhasan Differential Revision: D94138983 fbshipit-source-id: f02bad2f143f90c9e349e20973d515a7d83563e2
Summary: Pull Request resolved: meta-pytorch#830 abortNcclComm() should only abort the NCCL communicator, not the process. Process abort decisions belong to the caller (checkAndAbortIfTimedOutOrError, finalize). - Remove the abort() call and its abort_process_on_timeout_or_error guard from abortNcclComm(). Upcoming diffs will make abort mode the only supported mode, handling process abort at the caller level. - Clarify comments and log message: runAbortHooks and detachMemoryHook must both run before commAbort — hooks may need the live comm for debug info, and detachMemoryHook prevents CachingAllocator callbacks from referencing a destroyed comm. - Update PersistentRequest destructor comment: after commAbort the process is no longer guaranteed to abort, so pFree may fail and NCCLX_CHECK_IGNORE handles this gracefully. Reviewed By: tanquer Differential Revision: D94266656 fbshipit-source-id: 1a08af689b14fb40d73529d2a1df647c25e974d6
…torch#831) Summary: Pull Request resolved: meta-pytorch#831 Remove the separate RUN_RMA_TEST env var that gated RMA window tests. RMA capability is already implied by TEST_BACKEND=ncclx + NCCL_CTRAN_ENABLE=true, making the extra env var redundant: - TorchCommTestHelpers.h: simplify `shouldSkipRmaTest()` to check only TEST_BACKEND and NCCL_CTRAN_ENABLE (remove RUN_RMA_TEST check). Boolean parsing matches NCCL's env2bool (y/yes/t/true/1, case-insensitive). - tests.bzl: remove `"RUN_RMA_TEST": "true"` from transport_ctran config - WindowRmaTest.py: add `_should_skip_rma_test()` module-level helper checking TEST_BACKEND and NCCL_CTRAN_ENABLE, replace `skipIf` decorators that checked RUN_RMA_TEST - WindowRmaTest.{cpp,hpp}: replace `checkIfSkip()` (inline RUN_RMA_TEST check) with `shouldSkipRmaTest()` from TorchCommTestHelpers.h, remove the now-unused `checkIfSkip()` method - WindowRmaTimeoutTestMain.cpp: update gating comment to reflect new env vars Reviewed By: tanquer Differential Revision: D94266657 fbshipit-source-id: 49dfb2835b23ce3d1b551068417c6e8dd55cca6b
Summary: Pull Request resolved: meta-pytorch#724 Add Torchcomms support for AllGatherP, similar to PTD's semantics - Initialize, exec, free Reviewed By: JinghanHuang Differential Revision: D92447966 fbshipit-source-id: 7554a876dec234e56e53fb282bc74e4f58dc7bc4
Summary: Pull Request resolved: meta-pytorch#805 test debug server endpoints are hit Reviewed By: d4l3k Differential Revision: D93903496 fbshipit-source-id: d08b50c33da6015276873633094bba229cdd1ee5
Summary: Pull Request resolved: meta-pytorch#833 This diff addresses an issue with the `FlightRecorder` class, specifically with its `size()` API. The fix ensures that the `size()` method returns the correct number of entries in the current epoch, rather than the total entries in the underlying buffer. Reviewed By: kapilsh, d4l3k Differential Revision: D94235338 fbshipit-source-id: 3dff8bfdab47894725c8fb7a44582bcb9e1848a9
Summary: Pull Request resolved: meta-pytorch#700 Add Philox 4x32 counter-based RNG implementation for stochastic rounding in quantized collectives. The implementation is copied from [Jianyu's commit](https://github.com/fairinternal/gb200_moe_sol/commit/50546403adb2e0ef13013dee644ac6d61025dcc6#diff-18537af02b8d0c93856db69d228a907d1f118ae4d857f23b57eb8841c738d2a6) Key components: - philox4x32<N_ROUNDS>(): Core template function implementing the Philox round function with configurable round count - philox_randint4x(): Convenience function that generates 4 random uint32 values from a 64-bit seed and offset Also adds comprehensive unit tests covering: determinism, different seeds/offsets producing different values, uniform distribution (chi-squared test), output channel independence, 64-bit seed/offset utilization, uniqueness of consecutive outputs, different round counts, and multi-threaded consistency. Reviewed By: jspark1105 Differential Revision: D92755791 fbshipit-source-id: 53c3600a58384aa540c2cd6305dceff3149b089b
…orch#821) Summary: Pull Request resolved: meta-pytorch#821 Add FP32-to-BF16 stochastic rounding implementation with hardware and software paths. The implementation was based on the triton implementation. Key components: - Hardware-accelerated path for Blackwell GPUs (SM >= 100) using native PTX `cvt.rs.satfinite.bf16x2.f32` instruction - Software fallback for pre-Blackwell: adds random noise to lower 16 bits before truncation Also adds comprehensive unit tests covering: exact bf16 value preservation, NaN/Inf handling, unbiasedness verification, correct rounding to nearest bf16 neighbors, vectorized bf16x2 operations, determinism with same seed/offset, and round-up probability matching fractional position. Reviewed By: jspark1105 Differential Revision: D92755789 fbshipit-source-id: ab267eba835cbbd8f12dffa6d31c5831f1562d10
…h#849) Summary: Pull Request resolved: meta-pytorch#849 ## Bug Description `P2pNvlTransportDevice::put()` has a critical data corruption bug where `chunkBytes` is accumulated across loop iterations instead of being computed fresh per iteration. This causes `memcpy_vectorized` to copy an ever-growing number of bytes while the source/destination offsets advance per-chunk, resulting in out-of-bounds reads and writes. ## Root Cause The variable `chunkBytes` is declared outside the lambda and uses `+=` (line 853), causing it to accumulate across iterations of `for_each_item_contiguous`. On iteration K, `chunkBytes` contains the sum of all chunk sizes from iterations 0 through K. However, `chunkOffset` correctly advances to chunk K's starting position. This means `memcpy_vectorized` copies (K+1) chunks of data starting from chunk K's offset -- reading and writing well past allocation boundaries. ``` Bug visualization (3 chunks assigned to 1 group, alignedChunkSize=64): Iteration 0: chunkOffset=0, chunkBytes=64 -> copies [0..63] ✓ correct Iteration 1: chunkOffset=64, chunkBytes=128 -> copies [64..191] ✗ OOB! Iteration 2: chunkOffset=128, chunkBytes=192 -> copies [128..319] ✗ OOB! Memory layout: src: [chunk0][chunk1][chunk2][ out of bounds ] 0 64 128 192 320 ^^^^^^^^^^^^^^^^^ Iteration 2 reads past buffer end dst: [chunk0][chunk1][chunk2][ out of bounds ] 0 64 128 192 320 ^^^^^^^^^^^^^^^^^ Iteration 2 writes past buffer end ``` ## Why it hasn't crashed yet The chunk-sizing math typically produces `numChunks <= total_groups`, so `for_each_item_contiguous` assigns at most 1 chunk per group. With a single iteration, the accumulation has no effect. However, this depends on a numerical coincidence: ``` targetChunkSize = nbytes / total_groups alignedChunkSize = ceil(targetChunkSize, 16) // >= targetChunkSize numChunks = ceil(nbytes / alignedChunkSize) // typically <= total_groups ``` The bug CAN trigger when: - Partitioned subgroups have smaller total_groups - Transfer sizes don't divide evenly and produce numChunks > total_groups - Any future change to chunk sizing formula ## Fix Change `chunkBytes` from a mutable outer-scope variable with `+=` to a `const` variable declared inside the lambda (fresh per iteration). This matches the correct pattern used in three other places in the codebase: ``` Before (buggy): After (fixed): std::size_t chunkBytes = 0; std::size_t totalBytesWritten = 0; ... [&](...) { ... [&](...) { chunkBytes += ... const std::size_t chunkBytes = ... memcpy(... chunkBytes ...); memcpy(... chunkBytes ...); } totalBytesWritten += chunkBytes; return chunkBytes; } return totalBytesWritten; ``` Reference implementations (all correct): - `P2pSelfTransportDevice::put()` at line 112 - `P2pNvlTransportDevice::send()` at line 354 - `P2pNvlTransportDevice::recv()` at line 451 ## Impact - **Data corruption**: Chunks after the first copy data from wrong offsets, corrupting peer GPU memory via NVLink silently - **Out-of-bounds reads/writes**: Later iterations read/write past source and destination buffer boundaries - **Incorrect return value**: Returns accumulated total instead of actual bytes written by this group Oncall: ncclx Reviewed By: dmwu Differential Revision: D94437171 fbshipit-source-id: b73e366af42a81c7d81e9f7ac9370ce9cc98c17c
Summary: Pull Request resolved: meta-pytorch#840 This diff introduces `put_signal_with_fence()` and the underlying `signal_with_fence()`, which achieve the same ordering guarantee by setting `IBV_SEND_FENCE` (`DOCA_GPUNETIO_IB_MLX5_WQE_CTRL_FENCE`) on the signal WQE. The fence flag instructs the NIC to complete all prior posted WQEs (the data write) before processing the atomic signal, providing adaptive-routing safety without any GPU-side CQ polling overhead — the fence is handled entirely in NIC hardware. Summary of the three put_signal variants now available: - `put_signal()`: put + wait_local + signal (GPU-side CQ poll, adaptive safe) - `put_signal_with_fence()`: put + signal_with_fence (NIC-level fence, adaptive safe) — **new** - `put_signal_non_adaptive()`: fused put_signal WQE (no fence, NOT adaptive safe) Also extends the IBGDA benchmark suite with: - `signal_with_fence()` as a standalone low-level primitive - Single-shot kernel launchers for all three put_signal variants (correctness verification) - Batched kernel launcher for `put_signal_with_fence` (performance measurement) - A 3-way comparison benchmark (`PutSignalComparison`) with per-size data correctness checks before performance measurement Reviewed By: siyengar Differential Revision: D93763318 fbshipit-source-id: d4fc7f9ebaf8bc01cdfea2ad493316d6d901e22b
…-pytorch#841) Summary: Pull Request resolved: meta-pytorch#841 Adds group-level RDMA APIs to P2pIbgdaTransportDevice in two flavors: **Group-local APIs** (`put_group_local`, `put_signal_group_local`): - Accept a ThreadGroup and a single group's data chunk (already partitioned per group by the caller) - Partition the data across threads within the group - The last thread picks up any remainder bytes when nbytes is not evenly divisible by group_size - All ThreadGroup sizes supported (WARP, MULTIWARP, BLOCK) - group_size == 1: falls back to thread-level put() / put_signal() - group_size > 1: uses put_group_impl() with manual WQE construction **Group-global APIs** (`put_group_global`, `put_signal_group_global`): - Accept a ThreadGroup and a global data buffer shared by all groups - First partition data across groups (last group picks up remainder), then call the group-local API on each group's chunk - For put_signal_group_global, each group issues an atomic fetch-add signal, so the total accumulated signal is (total_groups × signalVal) The private `put_group_impl()` method is the internal building block: the leader reserves WQE slots, broadcasts the base index via `broadcast<uint64_t>()`, each thread prepares its own WQE, then the leader marks all WQEs ready and rings the doorbell. Other improvements in this diff: - Made `ThreadGroup::broadcast` a template function (`broadcast<T>`) that works for both `uint32_t` and `uint64_t`. Added overloaded `shfl()` for both types. Removed the old `broadcast_64()` method; all callsites now use `broadcast<uint64_t>()` directly. - Uses irregular buffer sizes (65000, 65003 bytes) in multi-group tests to exercise remainder handling at both the group-level and thread-level chunking. Reviewed By: snarayankh Differential Revision: D93932634 fbshipit-source-id: b0d49c785ab46c24ecfb99a8c2fcc69310781e55
Summary: Pull Request resolved: meta-pytorch#842 ## Problem `P2pIbgdaTransportDevice::wait_signal()` uses bare infinite spin-loops with no timeout or escape hatch: ``` while (*sig != value) {} // spins forever if peer crashes ``` If a remote peer crashes, a NIC fails, or a network partition occurs, the GPU hangs permanently with no diagnostics and no recovery path. This is a production reliability hazard — every other rank in the communicator will hang forever, requiring a full job kill. The NVLink transport (`SignalState::wait_until()`) already has full `Timeout` support via `Timeout.cuh`, but the IBGDA path was missing it. ``` ┌─────────────┐ RDMA signal ┌─────────────┐ │ Rank 0 │ ───────────────► │ Rank 1 │ │ (sender) │ │ (receiver) │ │ │ │ │ │ put_signal()│ │wait_signal()│ ◄── hangs forever │ │ ╳ NIC failure │ (no timeout│ if signal never │ │ │ today) │ arrives └─────────────┘ └─────────────┘ AFTER THIS DIFF: ┌─────────────┐ RDMA signal ┌─────────────┐ │ Rank 0 │ ───────────────► │ Rank 1 │ │ (sender) │ │ (receiver) │ │ │ │ │ │ put_signal()│ │wait_signal( │ │ │ ╳ NIC failure │ ..., │ │ │ │ timeout) │ ◄── __trap() with │ │ │ │ diagnostic after └─────────────┘ └─────────────┘ timeout expires ``` ## Solution Add an optional `const Timeout& timeout = Timeout()` parameter to `wait_signal()`. The default `Timeout()` constructor creates a "no timeout" config with `timeout_cycles=0`, which makes `checkExpired()` return `false` immediately — one well-predicted branch per spin iteration, effectively zero overhead. When enabled, each spin-wait case calls `TIMEOUT_TRAP_IF_EXPIRED_SINGLE` which checks `clock64() > deadline_cycles` and calls `__trap()` with a diagnostic printf on expiry. The diagnostic message includes: - The comparison operation (EQ, NE, LT, LE, GT, GE) - The signal ID - The expected value - The current signal value This reuses the existing `Timeout.cuh` infrastructure that the NVLink path (`SignalState`, `ChunkState`) already uses, bringing IBGDA to parity. ## API ```cpp // Before (infinite wait, no escape): transport.wait_signal(signalId, IbgdaCmpOp::GE, expectedVal); // After (with timeout - caller must call timeout.start() first): Timeout timeout = makeTimeout(1000, deviceId); // 1 second // ... in kernel: timeout.start(); transport.wait_signal(signalId, IbgdaCmpOp::GE, expectedVal, timeout); // After (without timeout - zero overhead, backward compatible): transport.wait_signal(signalId, IbgdaCmpOp::GE, expectedVal); ``` ## Changes - `P2pIbgdaTransportDevice.cuh`: Added `Timeout` parameter to `wait_signal()` with `TIMEOUT_TRAP_IF_EXPIRED_SINGLE` in all 6 comparison cases. Replaced `#include "ThreadGroup.cuh"` with `#include "Timeout.cuh"` (which transitively includes ThreadGroup). - `comms/pipes/BUCK`: Changed `p2p_ibgda_transport_device` dep from `:thread_group` to `:timeout_utils` (transitively exports thread_group). - Tests: Added `WaitSignalTimeoutTraps` (verifies __trap on expiry) and `WaitSignalNoTimeoutWhenSatisfied` (verifies no interference when signal is already satisfied). Uses separate test fixture with `cudaDeviceReset()` in TearDown to clear trap state. Reviewed By: siyengar Differential Revision: D94320422 fbshipit-source-id: c91ad40085e9c17a3d184243afabc5f6ad31321f
…torch#847) Summary: Pull Request resolved: meta-pytorch#847 CtranMapperUT.cc uses raw `setenv()` calls throughout its tests without proper cleanup. This leads to env var leakage between tests. Update to use RAII. Reviewed By: function47 Differential Revision: D94286006 fbshipit-source-id: cbfa268ac32a1da5af7d7d4b1618abfd5bb1218e
Summary: Pull Request resolved: meta-pytorch#826 Dynamic Connection (DC) is a transport protocol used within RDMA implementations, improving scalability over the widely used Reliably Connected (RC) transport. DC reduces the total number of QPs required system-wide, by having QPs of reliable type dynamically connect and disconnect from any remote node. (see https://docs.nvidia.com/networking/display/rdmacore50/dynamically+connected+(dc)+qps) Currently, we are using RC within RDMA implementations for ctran. In this diff, we experiment with DC using RDMA write tests. This is an initial experiment with the goal of eventually implementing and benchmarking DC as a potential alternative to RC. We have two types of RDMA write tests implemented here: one that does a ring pattern (each rank receives from rank (n-1)%numRanks, sends to rank (n+1)%numRanks), and one that does an all-to-all pattern (receives from every rank and sends to every rank) Reviewed By: mingrany Differential Revision: D91064760 fbshipit-source-id: b8ce82e1b3cc6ccd70d336ca81cc7f9733b2d0d3
Summary: Pull Request resolved: meta-pytorch#823 The AllReduceRing algorithm has two phases: ReduceScatter (RS) with N-1 steps, and AllGather (AG) with N-1 steps. Since AG is pure copy/forward with no reduction, it can be parallelized by sending data in both directions simultaneously. This change adds a reverse AG direction (counter-clockwise / leftward) that runs concurrently with the existing forward AG direction. With bi-directional AG, forward handles ceil((N-1)/2) steps and reverse handles floor((N-1)/2) steps, reducing total AllReduce steps from 2*(N-1) to (N-1) + ceil((N-1)/2), a ~25% reduction for large N. Key design points: - Reverse direction uses separate tmpSendBufRev/tmpRecvBufRev buffers (64MB additional per rank) - Reverse AG starts only after RS completes (rank's reduced shard must be in recvbuff first) - Shard indexing uses shift = (revStep + 1) % N to account for the +1 offset inherent in the RS phase's getStepShardIdx formula (rank r owns shard (r+1)%N after RS) - Host-side dependencies: kRevSendCopy -> kRevSendTrans, kRevRecvTrans -> kRevRecvFlush -> kRevRecvCopy -> kRevSendTrans - Device-side kernel adds _progressRevSend and _progressRevRecv to the main polling loop Files changed: - CtranAlgo.h/cc: Add RING_TMP_SEND_BUF_REV/RING_TMP_RECV_BUF_REV buffer types - AllReduceRingCommon.cuh: Add reverse Op enums, AlgoContext fields, getRevStepShardIdx, reverse round counting and dependency tracking - AllReduceRing.cc: Add host-side reverse progress functions, resource setup, main loop integration - AllReduceRing.cuh: Add device kernel _progressRevSend/_progressRevRecv functions Some tuning parameter - Cap BiDir AllGather to until 4MB by default. Beyond which usual ring based AllGather works. - Use template kernel functions for BiDir for ctring so we can limit the impact to cuda occupancy - P2205706069 (we still see impact for double and int64) - Small/medium messages (64B - 4MB): BiDirAg is 11-20% faster consistently - Large messages (8MB - 16MB): BiDirAg is 14-15% slower - Very large messages (32MB+): Performance is roughly equivalent (within ±2%) - Avg bus bandwidth: Baseline 12.88 GB/s vs BiDirAg 12.07 GB/s ``` Reviewed By: arttianezhu Differential Revision: D93668158 fbshipit-source-id: fd04604a87ecd8a8c2c8ca2fd592a52f68ec7e44
Summary: Pull Request resolved: meta-pytorch#824 As titled. Reviewed By: arttianezhu Differential Revision: D94261337 fbshipit-source-id: b69f42f9c9fb05eb36670b545b51472320a24d2a
…-pytorch#836) Summary: Pull Request resolved: meta-pytorch#836 Also add a clear error message for misconfigured executions. Reviewed By: saifhhasan Differential Revision: D94300251 fbshipit-source-id: 601bc51179b993f76173445e4bcc5f1a0086b78b
…ister (meta-pytorch#859) Summary: Pull Request resolved: meta-pytorch#859 ## Diff Summary Add optional input argument `deviceId` in `globalRegister` and `globalDeregister` in `RegCache`. If `deviceId` is unset, we use auto-detect cudaDev from the buffer pointer; otherwise, we assign `deviceId` to `cudaDev`. Reviewed By: dsjohns2 Differential Revision: D93637327 fbshipit-source-id: b78e90ae7523f8c4ca455762e2d8138e3d0b567d
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
20 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Add TorchCommXCCL::split implement and testcase DeviceMeshTest.py, Another split related cases like DPTPCommTest.py, FSDPCommTest.py, TPCommTest.py need full collective op support.