Skip to content

PortChannel: add optional GPUNetIO backend with pinned FetchContent dependency - #910

Open
Qinghua Zhou (seagater) wants to merge 72 commits into
mainfrom
qinghuazhou/gpunetio_submodule
Open

Qinghua Zhou (seagater) wants to merge 72 commits into
mainfrom
qinghuazhou/gpunetio_submodule

Conversation

@seagater

@seagater Qinghua Zhou (seagater) commented Sep 22, 2026 •

Copy link
Copy Markdown
Contributor

Summary

  • Add an optional GPU-initiated RDMA backend (EndpointConfig::Ib::Mode::Device, backed by DOCA GPUNetIO / GDAKI) to PortChannel.
  • Fetch the official NVIDIA-DOCA/gpunetio dependency via CMake FetchContent, pinned to commit c5d97fd52346763bcfeaff593dd1d84215156218. Known upstream CQ error-path issues are tracked in NVIDIA-DOCA/gpunetio#23.
  • Template channels and device handles on PortChannelBackend (BasePortChannel<Backend>, PortChannel<Backend>, BasePortChannelDeviceHandleT<Backend, SmemSlots>, and PortChannelDeviceHandleT<Backend, SmemSlots> with Host and Device specializations, with BasePortChannel and PortChannel defaulting to PortChannelBackend::Host, plus non-templated aliases BasePortChannelDeviceHandle and PortChannelDeviceHandle).
  • Construct Device-mode PortChannel directly from std::shared_ptr<PortDevice2DeviceSemaphore> (or Semaphore) and RegisteredMemory (PortChannel channel(semaphore, dst, src)) without ProxyService, and obtain the device handle via channel.deviceHandle() (or channel.deviceHandle<SmemSlots>()).
  • Signal and accumulate with RDMA atomic fetch-and-add: the responder applies an atomic only after placing earlier RDMA writes on the same queue pair, so observing a signal guarantees the preceding data is visible. putWithAccumulate posts a write and an atomic together (putWithSignal uses it on the semaphore).
  • Support queue-pair concurrency scopes via PortChannelDeviceHandle::Concurrency (Device by default, Block, Exclusive). Exclusive operations submit through BlueFlame up to blueFlameLimit (default kDefaultBlueFlameLimit = 512 KiB) using internal 64-byte shared-memory slots (SmemSlots, default 32 per CTA); larger transfers and the shared Block/Device modes ring the regular doorbell (submit_db).

Build and Use

GPUNetIO is built automatically whenever MSCCLPP_USE_CUDA=ON, MSCCLPP_USE_ROCM=OFF, and MSCCLPP_USE_IB=ON (the default), and skipped for ROCm or -DMSCCLPP_USE_IB=OFF builds:

cmake -S . -B build -DCMAKE_BUILD_TYPE=Release -DMSCCLPP_BUILD_TESTS=ON
cmake --build build -j$(nproc)

For an offline GPUNetIO checkout:

cmake -S . -B build -DFETCHCONTENT_SOURCE_DIR_GPUNETIO=/absolute/path/to/gpunetio

Usage Example

mscclpp::EndpointConfig cfg;
cfg.transport = mscclpp::Transport::IB0;
cfg.ib.mode = mscclpp::EndpointConfig::Ib::Mode::Device;
auto connection = communicator->connect(cfg, remoteRank).get();

auto localMem = communicator->registerMemory(buff, bytes, mscclpp::Transport::IB0);
communicator->sendMemory(localMem, remoteRank);
auto remoteMem = communicator->recvMemory(remoteRank).get();

auto semaphore = std::make_shared<mscclpp::PortDevice2DeviceSemaphore>(*communicator, connection);
mscclpp::PortChannel channel(semaphore, remoteMem, localMem);
auto deviceHandle = channel.deviceHandle();

In device code:

__global__ void send(mscclpp::PortChannelDeviceHandleT<mscclpp::PortChannelBackend::Device> channel, size_t bytes) {
  using Concurrency = mscclpp::PortChannelDeviceHandle::Concurrency;
  if (threadIdx.x != 0) return;
  channel.putWithSignal<Concurrency::Exclusive>(0, bytes);
  channel.flush<Concurrency::Exclusive>();
}

See docs/guide/gpunetio.md for full details.

Validation

# Unit tests
./build/bin/unit_tests --exclude-perf-tests

# Multi-process GPUNetIO (IbDeviceMode) unit tests
mpirun -np 2 ./build/bin/mp_unit_tests --filter=IbDeviceMode --exclude-perf-tests

# Multi-process GPUNetIO (IbDeviceMode) performance tests
mpirun -np 2 ./build/bin/mp_unit_tests --filter=IbDeviceMode --only-perf-tests

Performance: GPUNetIO (IbMode::Device) vs. IB CPU Proxy (IbMode::Host / IbMode::HostNoAtomic)

Measured with mp_unit_tests (PortChannelOneToOneTest) on 2x H100 GPUs over InfiniBand (kDefaultBlueFlameLimit = 512 KiB).

One Thread per QP (Concurrency::Exclusive, BlueFlame)

Benchmark IB Proxy (Host) IB Proxy (HostNoAtomic) GPUNetIO (Exclusive) Improvement vs. Host
PingPong latency (putWithSignal + wait + flush) 22.34 µs/iter 22.88 µs/iter 18.45 µs/iter -17.4% latency
PacketPingPong (LL) latency 6.34 µs/iter 6.51 µs/iter 4.50 µs/iter -29.1% latency
Bidirectional bandwidth (1 KB) 0.0605 GB/s 0.0573 GB/s 0.0698 GB/s +15.3%
1 QP bandwidth (1 KB) 0.0598 GB/s 0.0582 GB/s 0.0715 GB/s +19.6%
4 QPs bandwidth (4 KB total, 1 KB/QP) 0.2123 GB/s 0.1919 GB/s 0.2733 GB/s +28.7%
4 QPs flush stress (1 KB/QP) latency 23.36 µs/iter 24.35 µs/iter 18.66 µs/iter -20.1% latency
4 QPs flush stress (1 KB/QP) aggregate BW 0.1753 GB/s 0.1682 GB/s 0.2196 GB/s +25.3%
Message Size IB Proxy (Host) IB Proxy (HostNoAtomic) GPUNetIO (Exclusive)
1 QP, 64 KB 3.34 GB/s 3.26 GB/s 3.82 GB/s
1 QP, 1 MB 25.85 GB/s 25.49 GB/s 25.87 GB/s
1 QP, 4 MB 40.01 GB/s 39.77 GB/s 39.99 GB/s
1 QP, 128 MB 48.21 GB/s 48.22 GB/s 48.21 GB/s
4 QPs, 256 KB (64 KB/QP) 10.95 GB/s 10.81 GB/s 12.30 GB/s
4 QPs, 4 MB (1 MB/QP) 40.40 GB/s 40.15 GB/s 40.45 GB/s
4 QPs, 512 MB (128 MB/QP) 49.19 GB/s 49.18 GB/s 49.18 GB/s

Shared QPs (Concurrency::Block / Device, regular doorbell)

The proxy runs the same kernels with the same number of threads per QP.

Benchmark IB Proxy (Host) GPUNetIO (Block) GPUNetIO (Device)
PingPong latency (1 thread) 22.34 µs/iter 23.99 µs/iter 24.60 µs/iter
PacketPingPong (LL) latency (1 thread) 6.34 µs/iter 7.26 µs/iter 7.49 µs/iter
Same-channel concurrent flush (4 threads, 1 QP) 29.35 µs/iter 29.09 µs/iter 29.00 µs/iter
4 QPs bandwidth, 2 threads/QP (4 KB total) 0.1686 GB/s 0.1588 GB/s 0.1549 GB/s
4 QPs bandwidth, 2 threads/QP (256 KB total) 9.31 GB/s 7.94 GB/s 7.80 GB/s
4 QPs bandwidth, 2 threads/QP (4 MB total) 39.08 GB/s 37.33 GB/s 37.11 GB/s
4 QPs bandwidth, 2 threads/QP (512 MB total) 49.16 GB/s 49.14 GB/s 49.14 GB/s
4 QPs flush stress, 2 threads/QP (1 KB/QP) latency 28.14 µs/iter 29.54 µs/iter 30.30 µs/iter
4 QPs flush stress, 2 threads/QP (256 KB/QP) latency 46.99 µs/iter 50.97 µs/iter 51.52 µs/iter

In Block and Device modes, each posting thread updates the shared queue pair with GPU atomics and fences, and submits without BlueFlame. For the lowest latency, give each posting thread its own QP in Exclusive mode.

Port the generic integration and safety fixes from qinghuazhou/gpunetio_port_channel_merge_feature_ep onto main, preserving CPU-proxy defaults and accumulate semantics. Exclude all EP kernels and policy.

Pin the unmodified NVIDIA-DOCA/gpunetio v4.0.1 submodule. Add public host channel construction, CPU regressions, two-rank smoke coverage and installation instructions. ON/OFF CUDA builds, six CPU checks, installed-header compilation and lint pass. GPU/NIC runtime remains unverified; documented upstream allocation and CPU-service shutdown issues block merging.
Port GpuNetIoMultiQpBandwidth from the EP integration branch with configurable single-HCA QPs, collective geometry validation, queue-indexed device operations and per-QP atomic scratch. Preserve single-QP defaults and add CPU queue-isolation regression coverage.

Validated GPUNetIO ON/OFF builds, seven CPU test runs and repository lint. GPU/NIC benchmark execution remains pending; the upstream submodule is unchanged.
@azure-pipelines

Copy link
Copy Markdown
Azure Pipelines:
There may be pipelines that require an authorized user to comment /azp run to run.

Replace the dense-only upstream flatten helper with a checked sparse table adapter that keeps null self-rank slots zero and preserves remote peer/QP indices. Own the uploaded CUDA table in GpuNetIoService without modifying the upstream submodule.

Add actual-upstream-type regression coverage for 28 layouts and 296 rejected inputs. Commit the exact patch verified by the user; ON/OFF builds, CPU regressions and lint were validated separately.

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copilot review overview

🔵 Needs a closer look

Device completion and peer-validation defects remain alongside acknowledged upstream blockers and unverified GPU/NIC runtime behavior.

Review effort: Balanced
Findings: 1 High severity · 3 Medium severity · 3 Low severity

Open (7)
What changed in this PR

Adds an optional GPUNetIO/GDAKI backend to PortChannel while retaining CPU proxy behavior by default.

Changes:

  • Adds collective GPUNetIO service setup and multi-QP device operations.
  • Integrates the upstream GPUNetIO submodule and build/install configuration.
  • Adds CPU-stub, P2P, and bandwidth tests plus usage documentation.
File Description
.gitmodules Registers the GPUNetIO submodule.
CMakeLists.txt Adds and validates the build option.
docs/​quickstart.md Documents optional backend setup.
include/​mscclpp/​gpu_net_io_service.hpp Declares the host service API.
include/​mscclpp/​internal/​port_channel_gpunetio_device_impl.hpp Implements device RDMA operations.
include/​mscclpp/​port_channel.hpp Adds GPUNetIO channel construction.
include/​mscclpp/​port_channel_device.hpp Routes device operations by backend.
include/​mscclpp/​port_channel_gpunetio_device.hpp Defines the device context API.
src/​CMakeLists.txt Includes the optional component.
src/​core/​CMakeLists.txt Links GPUNetIO objects.
src/​core/​include/​ib.hpp Exposes internal verbs handles.
src/​core/​port_channel.cc Constructs GPUNetIO channels.
src/​gpunetio/​CMakeLists.txt Builds and installs the dependency.
src/​gpunetio/​PROVENANCE.md Documents provenance and limitations.
src/​gpunetio/​host/​gpu_net_io_qp_table.hpp Builds sparse device QP tables.
src/​gpunetio/​host/​gpu_net_io_service.cpp Sets up QPs, memory, and device state.
src/​gpunetio/​vendor Pins the upstream GPUNetIO dependency.
test/​mp_unit/​CMakeLists.txt Enables GPUNetIO MP tests.
test/​mp_unit/​port_channel_tests.cu Adds P2P and bandwidth coverage.
test/​unit/​CMakeLists.txt Registers focused unit tests.
test/​unit/​gpunetio_channel_test.cc Tests routing, QPs, and host API behavior.

💡 Configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

Comment thread src/core/port_channel.cc Resolved
Comment thread include/mscclpp/port_channel_device.hpp Resolved
Comment thread src/gpunetio/CMakeLists.txt Resolved
Comment thread include/mscclpp/gpu_net_io_service.hpp Resolved
Comment thread test/unit/gpunetio_channel_test.cc Resolved
Comment thread .gitmodules Resolved
Comment thread include/mscclpp/port_channel_device.hpp Resolved
Comment thread include/mscclpp/port_channel.hpp Resolved
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Refactor flush logic to handle errors from tryFlush.

Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Construct GPUNetIO host channels through the service to validate peers against retained rank/world-size metadata and reject incomplete setup. Restrict QP creation to direct GPU doorbells, valid DBRs and GPU-resident CQs; remove CPU progress fallback and exchange QP failures before connection-info collectives.

Add host peer-validation and direct-only policy regression coverage. Preserve the unmodified upstream submodule. ON/OFF builds, eight CPU test runs and repository lint pass; direct-GPU hardware qualification remains separate.
Replace the Git submodule with ON-only FetchContent pinned to the full upstream v4.0.1 SHA. Share dependency include paths through an interface target and retain offline overrides and installed headers.

Remove unused DO_PRAGMA and NVCC_PRAGMA_UNROLL compatibility definitions. Add CUDA compile probes for clean and caller-defined macro environments. Validate ON/OFF builds, eight CPU tests, disconnected overrides, installed headers, and lint.
Expose explicit connection handles for existing QPs, independently registered local/remote memory, and owned semaphore counters exchanged per connection. Preserve per-channel memory IDs and use the bound QP for writes, signals, atomics and flushes while retaining resource lifetimes.

Migrate P2P and multi-QP bandwidth tests to the public binding API and add two-channel isolation coverage. ON/OFF builds, eight CPU tests, installed-header checks and lint pass. GPU/NIC execution remains pending; uniform collective QP setup is retained.
Rename gin-prefixed device members, constructor parameters and helpers to gpuNetIo and update their test references. Remove NCCL-specific comparisons from public backend comments. No behavior or layout changes.

Validated naming-only equivalence, ON/OFF builds, eight CPU tests and repository lint.
Copilot AI review requested due to automatic review settings September 25, 2026 21:23
Comment thread include/mscclpp/gpunetio_device.hpp Resolved

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Comment thread src/gpunetio/CMakeLists.txt Resolved
Comment thread src/gpunetio/CMakeLists.txt Resolved
Comment thread src/gpunetio/host/gpu_net_io_policy.hpp Resolved
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Copilot AI review requested due to automatic review settings September 25, 2026 23:04
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>
Co-authored-by: Copilot Autofix powered by AI <175728472+Copilot@users.noreply.github.com>

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Comment thread cmake/GpuNetIo.cmake Resolved
Comment thread include/mscclpp/gpunetio_device.hpp Resolved
Comment thread src/core/gpunetio.cc Resolved
Comment thread src/core/port_channel.cc Resolved
Comment thread src/core/semaphore.cc Resolved
GPUNetIO queue pairs are created on an IB device's protection domain, so
the DOCA device and per-GPU handles opened on it now belong to that
IbCtx (IbCtx::gpuNetIoDevice(), created on first use and released before
the protection domain). Context no longer keeps a separate map of
GPUNetIO contexts.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Copilot-Session: 95bbe000-7818-4c60-8184-773c93680b01
Copilot AI balanced review requested due to automatic review settings November 4, 2026 19:04

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 Needs a closer look

The change breaks CMake generation when optional MLX5 headers are missing and allocates semaphore counters on the wrong GPU when the current device differs from the connection’s device, risking illegal memory access.

5 open issues

Review effort: Balanced

Comment thread cmake/GpuNetIo.cmake Resolved
Comment thread include/mscclpp/port_channel.hpp Resolved
Comment thread include/mscclpp/port_channel.hpp Resolved
Comment thread include/mscclpp/port_channel_device.hpp Resolved
Comment thread include/mscclpp/device.hpp Resolved
Comment thread include/mscclpp/port_channel_device.hpp Resolved
Configure per-CTA WQE capacity through device-handle templates and accept runtime slot indices with doorbell fallback. Remove caller-owned shared-memory binding and update existing callers and documentation.

Drop the local upstream CQ cleanup patch and track its error-handling issues upstream.

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Copilot-Session: 55ef1b7e-0ccb-4b59-ad58-48a168608302
Copilot AI balanced review requested due to automatic review settings November 7, 2026 00:09

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Comment thread include/mscclpp/gpunetio_device.hpp Resolved
Comment thread include/mscclpp/semaphore_device.hpp Resolved
Comment thread src/core/CMakeLists.txt Resolved
Comment thread src/core/gpunetio.cc Resolved
Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Copilot-Session: 55ef1b7e-0ccb-4b59-ad58-48a168608302
Copilot AI balanced review requested due to automatic review settings November 7, 2026 01:23

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Comment thread cmake/GpuNetIo.cmake Resolved
…elDeviceHandle

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Copilot-Session: 55ef1b7e-0ccb-4b59-ad58-48a168608302
Copilot AI balanced review requested due to automatic review settings November 7, 2026 02:54

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

…eHandle

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Copilot-Session: 55ef1b7e-0ccb-4b59-ad58-48a168608302
Copilot AI balanced review requested due to automatic review settings November 7, 2026 06:31

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Comment thread include/mscclpp/port_channel.hpp Resolved
…_IBVERBS

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Copilot-Session: 55ef1b7e-0ccb-4b59-ad58-48a168608302
Copilot AI balanced review requested due to automatic review settings November 9, 2026 02:00

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

#include "atomic_device.hpp"
#endif

#if defined(MSCCLPP_DEVICE_CUDA) && defined(USE_IBVERBS)
Comment thread src/core/gpunetio.cc
MSCCLPP_DOCA_THROW(doca_verbs_qp_attr_set_ack_timeout(attr, 18));
MSCCLPP_DOCA_THROW(doca_verbs_qp_attr_set_retry_cnt(attr, 7));
MSCCLPP_DOCA_THROW(doca_verbs_qp_attr_set_rnr_retry(attr, 7));
MSCCLPP_DOCA_THROW(doca_verbs_qp_attr_set_max_rd_atomic(attr, 1));
…hannel

Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com>
Copilot-Session: 55ef1b7e-0ccb-4b59-ad58-48a168608302
Copilot AI balanced review requested due to automatic review settings November 9, 2026 02:29

This branch has not been deployed

No deployments
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

6 participants