diff --git a/docs/tutorials/04-port-channel.md b/docs/tutorials/04-port-channel.md index 5dcbed913..096280886 100644 --- a/docs/tutorials/04-port-channel.md +++ b/docs/tutorials/04-port-channel.md @@ -60,12 +60,15 @@ We need to call `proxyService.startProxy()` before running GPU kernels that use The device handle of a `PortChannel` provides the following methods. Since the data transfer is offloaded, each method is supposed to be called by a single GPU thread. - `put()`: Initiates an asynchronous one-way data transfer from the local memory to the remote memory. -- `signal()`: Asynchronously signals the completion of all previous `put()`s to the remote side. +- `accumulate()`: Initiates a 64-bit atomic add on a value in remote memory. +- `signal()`: Asynchronously signals the completion of previous operations to the remote side. - `wait()`: Blocks the calling GPU thread until the corresponding `signal()` is received from the remote side. - `poll()`: Non-blocking version of `wait()`. Returns immediately with a boolean indicating whether the signal has been received. - `flush()`: Synchronizes the local GPU with the `PortChannel`, ensuring that all previous operations are completed. - Fused methods (e.g., `putWithSignal()`): combines multiple sequential operations into a single call for efficiency. +Like `put()`, `accumulate()` is issued asynchronously through the proxy and is ordered with operations issued after it on the same channel. A typical sequence enqueues one or more additions and then uses the application's existing `signal()`/`flush()`/`wait()` protocol before consuming the result. See the `PortChannel::accumulate()` and `Connection::accumulate()` API documentation for transport restrictions, alignment requirements, and concurrency details. + The following diagram illustrates how the `bidirPutKernel()` function in the example code would work when GPU0 is faster than GPU1. The execution order may vary depending on the relative speeds of the GPUs. ```{mermaid} diff --git a/include/mscclpp/atomic_device.hpp b/include/mscclpp/atomic_device.hpp index d00bb50cf..634484e2f 100644 --- a/include/mscclpp/atomic_device.hpp +++ b/include/mscclpp/atomic_device.hpp @@ -46,8 +46,13 @@ constexpr auto memoryOrderRelease = __ATOMIC_RELEASE; constexpr auto memoryOrderAcqRel = __ATOMIC_ACQ_REL; constexpr auto memoryOrderSeqCst = __ATOMIC_SEQ_CST; +#if defined(MSCCLPP_DEVICE_HIP) +constexpr auto scopeSystem = __HIP_MEMORY_SCOPE_SYSTEM; +constexpr auto scopeDevice = __HIP_MEMORY_SCOPE_AGENT; +#else constexpr auto scopeSystem = 0; constexpr auto scopeDevice = 0; +#endif // defined(MSCCLPP_DEVICE_HIP) template MSCCLPP_HOST_DEVICE_INLINE T atomicLoad(const T* ptr, int memoryOrder) { @@ -61,7 +66,11 @@ MSCCLPP_HOST_DEVICE_INLINE void atomicStore(T* ptr, const T& val, int memoryOrde template MSCCLPP_HOST_DEVICE_INLINE T atomicFetchAdd(T* ptr, const T& val, int memoryOrder) { +#if defined(__HIP_DEVICE_COMPILE__) + return __hip_atomic_fetch_add(ptr, val, memoryOrder, scope); +#else return __atomic_fetch_add(ptr, val, memoryOrder); +#endif // defined(__HIP_DEVICE_COMPILE__) } #endif // !defined(MSCCLPP_DEVICE_CUDA) diff --git a/include/mscclpp/core.hpp b/include/mscclpp/core.hpp index a0c9f7494..af5fae48a 100644 --- a/include/mscclpp/core.hpp +++ b/include/mscclpp/core.hpp @@ -649,6 +649,36 @@ class Connection { /// @param newValue The new value to write. void updateAndSync(RegisteredMemory dst, uint64_t dstOffset, uint64_t* src, uint64_t newValue); + /// Add a value to a 64-bit integer in a destination RegisteredMemory. + /// + /// The caller supplies only its own contribution, unlike updateAndSync(), which needs the + /// destination's current value. Addition is modulo 2^64 (the signed operand contributes its + /// two's-complement bit pattern) and commutes, so arrival order does not matter. + /// + /// The addition must be a real read-modify-write at the destination, so how many concurrent + /// writers one address allows depends on the transport: + /// + /// - IB: any number of writers, via NIC atomic fetch-and-add. Throws in no-atomic mode, where + /// the device has no RDMA atomics. + /// - Ethernet: any number of remote writers. The receiving process does the update and + /// serializes its connections. The destination GPU must not write the address concurrently; + /// such a write is lost inside the read-modify-write window. + /// - CudaIpc on ROCm: any number of writers. The proxy runs a kernel, which a caller kernel + /// does not block. + /// - CudaIpc on CUDA: throws. The host cannot read-modify-write device memory, and a + /// proxy-launched kernel cannot run while the caller's kernel waits. Use a device-side atomic + /// on peer memory reached through a MemoryChannel. + /// + /// The 8-byte target word must fit within @p dst, and its final address must be naturally + /// 8-byte aligned. + /// + /// @param dst The destination RegisteredMemory. + /// @param dstOffset The offset in bytes from the start of the destination RegisteredMemory. + /// @param value The 64-bit signed value to add. + /// @throws Error with ErrorCode::InvalidUsage if the target is out of bounds or misaligned, or + /// if the transport cannot accumulate. + void accumulate(RegisteredMemory dst, uint64_t dstOffset, int64_t value); + /// Flush any pending writes to the remote process. /// @param timeoutUsec Timeout in microseconds. Default: -1 (no timeout) void flush(int64_t timeoutUsec = -1); diff --git a/include/mscclpp/env.hpp b/include/mscclpp/env.hpp index 7415119f7..59fb9336a 100644 --- a/include/mscclpp/env.hpp +++ b/include/mscclpp/env.hpp @@ -84,9 +84,10 @@ class Env { /// enabled and will dump traces to this directory. Unset by default. const std::string npkitDumpDir; - /// Env name: `MSCCLPP_CUDAIPC_USE_DEFAULT_STREAM`. If set to true, the CUDA IPC transport will use the default - /// stream for all operations. If set to false, it will use a separate stream for each operation. This is an - /// experimental feature and should be false in most cases. Default is false. + /// Env name: `MSCCLPP_CUDAIPC_USE_DEFAULT_STREAM`. On CUDA, if set to true, the CUDA IPC transport will use the + /// default stream for all operations. If set to false, it will use a separate stream for each operation. ROCm + /// always uses a nonblocking stream so proxy-launched kernels can make progress. This is an experimental feature + /// and should be false in most cases. Default is false. const bool cudaIpcUseDefaultStream; /// Env name: `MSCCLPP_NCCL_LIB_PATH`. The path to the original NCCL/RCCL shared library. If set, it will be used diff --git a/include/mscclpp/fifo_device.hpp b/include/mscclpp/fifo_device.hpp index 4670f47c9..89eae99e8 100644 --- a/include/mscclpp/fifo_device.hpp +++ b/include/mscclpp/fifo_device.hpp @@ -15,10 +15,27 @@ namespace mscclpp { +/// Operation that a trigger asks the proxy to perform. +/// +/// These are opcodes, not flags: compare one by equality, and never combine two. The encoding +/// enumerates the combinations the device API can produce rather than composing them, so a +/// combination nothing emits cannot be expressed, and a trigger whose type field is unset is not +/// a valid operation. using TriggerType = uint64_t; -constexpr TriggerType TriggerData = 0x1; // Trigger a data transfer. -constexpr TriggerType TriggerFlag = 0x2; // Trigger a signaling. -constexpr TriggerType TriggerSync = 0x4; // Trigger a flush. +constexpr TriggerType TriggerNone = 0; // Not an operation; invalid for ProxyService. +constexpr TriggerType TriggerPut = 1; // Transfer data. +constexpr TriggerType TriggerSignal = 2; // Signal the remote semaphore. +constexpr TriggerType TriggerPutWithSignal = 3; // Transfer data, then signal. +constexpr TriggerType TriggerFlush = 4; // Flush the connection. +constexpr TriggerType TriggerAccumulate = 6; // Add a value to remote memory. +constexpr TriggerType TriggerPutWithSignalAndFlush = 7; // Transfer data, signal, then flush. +// 5 is unassigned. + +// Preserve the original flag-combination encodings for existing operations. Triggers are now +// compared as opcodes, but changing these values would break producers and consumers built from +// different revisions. +static_assert(TriggerPut == 1 && TriggerSignal == 2 && TriggerPutWithSignal == 3 && TriggerFlush == 4 && + TriggerPutWithSignalAndFlush == 7); constexpr unsigned int TriggerBitsSize = 32; constexpr unsigned int TriggerBitsOffset = 32; @@ -29,6 +46,8 @@ constexpr unsigned int TriggerBitsSemaphoreId = 10; // there. See FifoDeviceHandle::push(). constexpr unsigned int TriggerBitsFifoReserved = 1; +static_assert(TriggerAccumulate < (1ULL << TriggerBitsType), "trigger opcodes must fit in the type field"); + /// Pair of 64-bit unsigned integers used as a trigger for the proxy. /// Used as a work element in the concurrent FIFO. /// Most significant bit of snd is reserved. diff --git a/include/mscclpp/port_channel.hpp b/include/mscclpp/port_channel.hpp index 18d67524e..b8d5f8501 100644 --- a/include/mscclpp/port_channel.hpp +++ b/include/mscclpp/port_channel.hpp @@ -84,7 +84,7 @@ class ProxyService : public BaseProxyService { std::vector memories_; std::shared_ptr proxy_; std::unordered_map, int> inflightRequests_; - // Latest pending TriggerSync FIFO position per connection. Proxy publishes pos+1 to the + // Latest pending TriggerFlush FIFO position per connection. Proxy publishes pos+1 to the // connection's gpuFlushDonePos_ when the CQ drains, then erases the entry. std::unordered_map, uint64_t> pendingFlushPos_; diff --git a/include/mscclpp/port_channel_device.hpp b/include/mscclpp/port_channel_device.hpp index fd575b4c6..e2306113b 100644 --- a/include/mscclpp/port_channel_device.hpp +++ b/include/mscclpp/port_channel_device.hpp @@ -19,7 +19,7 @@ using MemoryId = uint32_t; namespace detail { #if defined(MSCCLPP_DEVICE_COMPILE) -/// Wait until the proxy has processed and drained the TriggerSync at FIFO position `fifoPos`. +/// Wait until the proxy has processed and drained the TriggerFlush at FIFO position `fifoPos`. /// The proxy publishes `flushDonePos = latestCompletedPos + 1` when the CQ drains, so the /// wait condition `flushDonePos > fifoPos` is satisfied exactly when our own request has /// been completed. Using the FIFO push position as the wait target couples the wait to the @@ -51,7 +51,7 @@ struct BasePortChannelDeviceHandle { : semaphoreId_(semaphoreId), semaphore_(semaphore), fifo_(fifo), flushDonePos_(flushDonePos) {} #if defined(MSCCLPP_DEVICE_COMPILE) - /// Push a TriggerData to the FIFO. + /// Push a TriggerPut to the FIFO. /// @param dstId The ID of destination memory region. /// @param dstOffset The offset into the destination memory region. /// @param srcId The ID of source memory region. @@ -59,10 +59,10 @@ struct BasePortChannelDeviceHandle { /// @param size The size of the transfer. MSCCLPP_DEVICE_INLINE void put(MemoryId dstId, uint64_t dstOffset, MemoryId srcId, uint64_t srcOffset, uint64_t size) { - fifo_.push({TriggerData, dstId, dstOffset, srcId, srcOffset, size, semaphoreId_}); + fifo_.push({TriggerPut, dstId, dstOffset, srcId, srcOffset, size, semaphoreId_}); } - /// Push a TriggerData to the FIFO. + /// Push a TriggerPut to the FIFO. /// @param dstId The ID of destination memory region. /// @param srcId The ID of source memory region. /// @param offset The common offset into the destination and source memory regions. @@ -71,10 +71,10 @@ struct BasePortChannelDeviceHandle { put(dstId, offset, srcId, offset, size); } - /// Push a TriggerFlag to the FIFO. - MSCCLPP_DEVICE_INLINE void signal() { fifo_.push({TriggerFlag, 0, 0, 0, 0, 0, semaphoreId_}); } + /// Push a TriggerSignal to the FIFO. + MSCCLPP_DEVICE_INLINE void signal() { fifo_.push({TriggerSignal, 0, 0, 0, 0, 0, semaphoreId_}); } - /// Push a TriggerData and a TriggerFlag at the same time to the FIFO. + /// Push a TriggerPutWithSignal to the FIFO. /// @param dstId The ID of destination memory region. /// @param dstOffset The offset into the destination memory region. /// @param srcId The ID of source memory region. @@ -82,10 +82,10 @@ struct BasePortChannelDeviceHandle { /// @param size The size of the transfer. MSCCLPP_DEVICE_INLINE void putWithSignal(MemoryId dstId, uint64_t dstOffset, MemoryId srcId, uint64_t srcOffset, uint64_t size) { - fifo_.push({TriggerData | TriggerFlag, dstId, dstOffset, srcId, srcOffset, size, semaphoreId_}); + fifo_.push({TriggerPutWithSignal, dstId, dstOffset, srcId, srcOffset, size, semaphoreId_}); } - /// Push a TriggerData and a TriggerFlag at the same time to the FIFO. + /// Push a TriggerPutWithSignal to the FIFO. /// @param dstId The ID of destination memory region. /// @param srcId The ID of source memory region. /// @param offset The common offset into the destination and source memory regions. @@ -94,7 +94,7 @@ struct BasePortChannelDeviceHandle { putWithSignal(dstId, offset, srcId, offset, size); } - /// Push a TriggerData, a TriggerFlag, and a TriggerSync at the same time to the FIFO. + /// Push a TriggerPutWithSignalAndFlush to the FIFO. /// @param dstId The ID of destination memory region. /// @param dstOffset The offset into the destination memory region. /// @param srcId The ID of source memory region. @@ -103,12 +103,11 @@ struct BasePortChannelDeviceHandle { /// @param maxSpinCount The maximum number of spin counts before asserting. Never assert if negative. MSCCLPP_DEVICE_INLINE void putWithSignalAndFlush(MemoryId dstId, uint64_t dstOffset, MemoryId srcId, uint64_t srcOffset, uint64_t size, int64_t maxSpinCount = 1000000) { - uint64_t pos = - fifo_.push({TriggerData | TriggerFlag | TriggerSync, dstId, dstOffset, srcId, srcOffset, size, semaphoreId_}); + uint64_t pos = fifo_.push({TriggerPutWithSignalAndFlush, dstId, dstOffset, srcId, srcOffset, size, semaphoreId_}); detail::waitFlush(flushDonePos_, pos, maxSpinCount); } - /// Push a TriggerData, a TriggerFlag, and a TriggerSync at the same time to the FIFO. + /// Push a TriggerPutWithSignalAndFlush to the FIFO. /// @param dstId The ID of destination memory region. /// @param srcId The ID of source memory region. /// @param offset The common offset into the destination and source memory regions. @@ -119,13 +118,26 @@ struct BasePortChannelDeviceHandle { putWithSignalAndFlush(dstId, offset, srcId, offset, size, maxSpinCount); } - /// Push a TriggerSync to the FIFO. + /// Push a TriggerFlush to the FIFO. /// @param maxSpinCount The maximum number of spin counts before asserting. Never assert if negative. MSCCLPP_DEVICE_INLINE void flush(int64_t maxSpinCount = 1000000) { - uint64_t pos = fifo_.push({TriggerSync, 0, 0, 0, 0, 0, semaphoreId_}); + uint64_t pos = fifo_.push({TriggerFlush, 0, 0, 0, 0, 0, semaphoreId_}); detail::waitFlush(flushDonePos_, pos, maxSpinCount); } + /// Push an accumulate trigger to the FIFO: add a 64-bit value to remote memory. + /// Connection::accumulate() documents how many concurrent writers each transport allows. + /// @param dstId The ID of destination memory region. + /// @param dstOffset The offset into the destination memory region. + /// @param value The 64-bit signed value to add. + MSCCLPP_DEVICE_INLINE void accumulate(MemoryId dstId, uint64_t dstOffset, int64_t value) { + // The operand occupies fst, spanning the low size and high srcOffset fields. + uint64_t operand = static_cast(value); + ProxyTrigger trigger(TriggerAccumulate, dstId, dstOffset, /*srcId=*/0, operand >> TriggerBitsSize, + static_cast(operand), semaphoreId_); + fifo_.push(trigger); + } + /// Check if the port channel has been signaled. /// @return true if the port channel has been signaled. MSCCLPP_DEVICE_INLINE bool poll() { return semaphore_.poll(); } @@ -149,7 +161,7 @@ struct PortChannelDeviceHandle : public BasePortChannelDeviceHandle { : BasePortChannelDeviceHandle(semaphoreId, semaphore, fifo, flushDonePos), dst_(dst), src_(src) {} #if defined(MSCCLPP_DEVICE_COMPILE) - /// Push a TriggerData to the FIFO. + /// Push a TriggerPut to the FIFO. /// @param dstOffset The offset into the destination memory region. /// @param srcOffset The offset into the source memory region. /// @param size The size of the transfer. @@ -157,12 +169,12 @@ struct PortChannelDeviceHandle : public BasePortChannelDeviceHandle { BasePortChannelDeviceHandle::put(dst_, dstOffset, src_, srcOffset, size); } - /// Push a TriggerData to the FIFO. + /// Push a TriggerPut to the FIFO. /// @param offset The common offset into the destination and source memory regions. /// @param size The size of the transfer. MSCCLPP_DEVICE_INLINE void put(uint64_t offset, uint64_t size) { put(offset, offset, size); } - /// Push a TriggerData and a TriggerFlag at the same time to the FIFO. + /// Push a TriggerPutWithSignal to the FIFO. /// @param dstOffset The offset into the destination memory region. /// @param srcOffset The offset into the source memory region. /// @param size The size of the transfer. @@ -170,12 +182,12 @@ struct PortChannelDeviceHandle : public BasePortChannelDeviceHandle { BasePortChannelDeviceHandle::putWithSignal(dst_, dstOffset, src_, srcOffset, size); } - /// Push a TriggerData and a TriggerFlag at the same time to the FIFO. + /// Push a TriggerPutWithSignal to the FIFO. /// @param offset The common offset into the destination and source memory regions. /// @param size The size of the transfer. MSCCLPP_DEVICE_INLINE void putWithSignal(uint64_t offset, uint64_t size) { putWithSignal(offset, offset, size); } - /// Push a TriggerData, a TriggerFlag, and a TriggerSync at the same time to the FIFO. + /// Push a TriggerPutWithSignalAndFlush to the FIFO. /// @param dstOffset The offset into the destination memory region. /// @param srcOffset The offset into the source memory region. /// @param size The size of the transfer. @@ -185,12 +197,19 @@ struct PortChannelDeviceHandle : public BasePortChannelDeviceHandle { BasePortChannelDeviceHandle::putWithSignalAndFlush(dst_, dstOffset, src_, srcOffset, size, maxSpinCount); } - /// Push a TriggerData, a TriggerFlag, and a TriggerSync at the same time to the FIFO. + /// Push a TriggerPutWithSignalAndFlush to the FIFO. /// @param offset The common offset into the destination and source memory regions. /// @param size The size of the transfer. MSCCLPP_DEVICE_INLINE void putWithSignalAndFlush(uint64_t offset, uint64_t size) { putWithSignalAndFlush(offset, offset, size); } + /// Push an accumulate trigger to the FIFO: add a 64-bit value to the destination memory. + /// See Connection::accumulate() for transport support. + /// @param dstOffset The offset into the destination memory region. + /// @param value The 64-bit signed value to add. + MSCCLPP_DEVICE_INLINE void accumulate(uint64_t dstOffset, int64_t value) { + BasePortChannelDeviceHandle::accumulate(dst_, dstOffset, value); + } #endif // defined(MSCCLPP_DEVICE_COMPILE) }; diff --git a/src/core/accumulate_kernel.cu b/src/core/accumulate_kernel.cu new file mode 100644 index 000000000..0f1169e07 --- /dev/null +++ b/src/core/accumulate_kernel.cu @@ -0,0 +1,33 @@ +// Copyright (c) Microsoft Corporation. +// Licensed under the MIT License. + +#include + +#if defined(MSCCLPP_USE_ROCM) + +#include +#include + +#include "context.hpp" + +namespace mscclpp { + +// System-scope modulo-2^64 atomic add. +__global__ void accumulateU64Kernel(uint64_t* dst, uint64_t value) { + (void)atomicFetchAdd(dst, value, memoryOrderRelaxed); +} + +void CudaIpcStream::accumulate(uint64_t* dst, uint64_t value) { + CudaDeviceGuard deviceGuard(deviceId_); + setStreamIfNeeded(); + // Submit to this connection's stream, which orders the add ahead of any signal or flush that + // follows. On ROCm a kernel runs while the caller's kernel occupies the GPU, so the proxy does + // not wait for the caller. + accumulateU64Kernel<<<1, 1, 0, *stream_>>>(dst, value); + MSCCLPP_CUDATHROW(cudaGetLastError()); + dirty_ = true; +} + +} // namespace mscclpp + +#endif // defined(MSCCLPP_USE_ROCM) diff --git a/src/core/connection.cc b/src/core/connection.cc index e5a516390..cda0dda7d 100644 --- a/src/core/connection.cc +++ b/src/core/connection.cc @@ -9,6 +9,7 @@ #include #include +#include #include #include @@ -38,6 +39,21 @@ static void validateTransport(RegisteredMemory mem, Transport transport, uint64_ } } +static void validateAccumulateBounds(RegisteredMemory mem, uint64_t offset) { + constexpr uint64_t wordSize = sizeof(uint64_t); + if (offset > mem.size() || wordSize > mem.size() - offset) { + THROW(CONN, Error, ErrorCode::InvalidUsage, "RegisteredMemory out of bounds for 64-bit accumulate"); + } +} + +static void validateAccumulateAlignment(uintptr_t base, uint64_t offset) { + constexpr uintptr_t alignment = alignof(uint64_t); + uintptr_t targetAlignment = (base % alignment + offset % alignment) % alignment; + if (targetAlignment != 0) { + THROW(CONN, Error, ErrorCode::InvalidUsage, "accumulate destination must be naturally 8-byte aligned"); + } +} + static bool isSameProcess(const Endpoint& a, const Endpoint& b) { return a.hostHash() == b.hostHash() && a.pidHash() == b.pidHash(); } @@ -76,6 +92,10 @@ MSCCLPP_API_CPP void Connection::updateAndSync(RegisteredMemory dst, uint64_t ds impl_->updateAndSync(dst, dstOffset, src, newValue); } +MSCCLPP_API_CPP void Connection::accumulate(RegisteredMemory dst, uint64_t dstOffset, int64_t value) { + impl_->accumulate(dst, dstOffset, value); +} + MSCCLPP_API_CPP void Connection::flush(int64_t timeoutUsec) { impl_->flush(timeoutUsec); } MSCCLPP_API_CPP Transport Connection::transport() const { return impl_->transport(); } @@ -202,6 +222,32 @@ void CudaIpcConnection::flush(int64_t timeoutUsec) { #endif } +void CudaIpcConnection::accumulate(RegisteredMemory dst, uint64_t dstOffset, [[maybe_unused]] int64_t value) { + validateTransport(dst, remoteTransport()); + validateAccumulateBounds(dst, dstOffset); + validateAccumulateAlignment(reinterpret_cast(dst.data()), dstOffset); +#if defined(MSCCLPP_USE_ROCM) + // A kernel on this connection's stream performs the addition, a real read-modify-write, so + // writers in any number of processes may target one address. The host cannot do it instead: + // GPU memory is host-accessible only to the process that allocated it, and the proxy holds an + // IPC-imported mapping, which is device-only. + uint64_t* dstPtr = reinterpret_cast(reinterpret_cast(dst.data()) + dstOffset); + stream_->accumulate(dstPtr, static_cast(value)); + INFO(CONN, "CudaIpcConnection accumulate: dst ", dstPtr, ", value ", value); +#else + // The host reaches device memory only through the copy engines, which move a value but cannot + // add to one, and a host-side read-modify-write is not atomic. A kernel is atomic but unusable: + // in the caller's context it does not start until the caller's kernel finishes, deadlocking any + // caller that spins on the result; in a separate context it costs 2391 us per operation against + // 19 us for a plain remote store. ROCm has neither limit. + THROW(CONN, Error, ErrorCode::InvalidUsage, + "accumulate is not supported over CudaIpc on CUDA: the host cannot atomically " + "read-modify-write device memory, and a proxy-launched kernel cannot run while the " + "caller's kernel waits. Use a device-side atomic on peer memory reached through a " + "MemoryChannel instead"); +#endif // defined(MSCCLPP_USE_ROCM) +} + // IBConnection void IBConnection::recvThreadFunc() { @@ -500,6 +546,26 @@ void IBConnection::flush(int64_t timeoutUsec) { #endif } +void IBConnection::accumulate(RegisteredMemory dst, uint64_t dstOffset, int64_t value) { + validateTransport(dst, remoteTransport()); + validateAccumulateBounds(dst, dstOffset); + auto dstTransportInfo = getImpl(dst).getTransportInfo(remoteTransport()); + if (dstTransportInfo.ibLocal) { + THROW(CONN, Error, ErrorCode::InvalidUsage, "dst is local, which is not supported"); + } + auto dstMrInfo = dstTransportInfo.ibMrInfo; + validateAccumulateAlignment(static_cast(dstMrInfo.addr), dstOffset); + + if (ibNoAtomic_) { + THROW(CONN, Error, ErrorCode::InvalidUsage, "accumulate is not supported in IB no-atomic mode"); + } + + qp_.lock()->stageSendAtomicAdd(atomicSrcTransportInfo_.ibMr, dstMrInfo, /*wrId=*/0, dstOffset, + static_cast(value), /*signaled=*/true); + qp_.lock()->postSend(); + INFO(CONN, "IBConnection accumulate: dst ", (uint8_t*)dstMrInfo.addr + dstOffset, ", value ", value); +} + void IBConnection::requestFlush() { // No-op: IB sends were already posted by prior conn.write() calls in handleTrigger. // progressFlush() drives completion by polling the send CQ. @@ -733,11 +799,49 @@ bool EthernetConnection::receiveFramePart(void* ptr, int size, bool allowBoundar : "Ethernet peer closed before completing a frame"); } +// Serializes the receive-side read-modify-write of accumulate() across this process's Ethernet +// connections. Ethernet costs ~136 us per operation, so the contention is irrelevant. +static std::mutex& accumulateMutex() { + static std::mutex mtx; + return mtx; +} + +void EthernetConnection::accumulate(RegisteredMemory dst, uint64_t dstOffset, int64_t value) { + validateTransport(dst, remoteTransport()); + validateAccumulateBounds(dst, dstOffset); + validateAccumulateAlignment(reinterpret_cast(dst.originalDataPtr()), dstOffset); + + // Wire format matches write(): [dstPtr(8B)] [size(8B)] [data(size B)]. The MSB of size marks + // the message as an accumulate. + uint64_t* dstPtr = reinterpret_cast(reinterpret_cast(dst.originalDataPtr()) + dstOffset); + constexpr uint64_t accumulateFlag = uint64_t{1} << uint64_t{63}; + uint64_t dataSize = sizeof(uint64_t) | accumulateFlag; + uint64_t messageSize = 0; + + char* dstPtrBytes = reinterpret_cast(&dstPtr); + std::copy(dstPtrBytes, dstPtrBytes + sizeof(dstPtr), sendBuffer_.data() + messageSize); + messageSize += sizeof(dstPtr); + + char* sizeBytes = reinterpret_cast(&dataSize); + std::copy(sizeBytes, sizeBytes + sizeof(dataSize), sendBuffer_.data() + messageSize); + messageSize += sizeof(dataSize); + + uint64_t addValue = static_cast(value); + char* valueBytes = reinterpret_cast(&addValue); + std::copy(valueBytes, valueBytes + sizeof(addValue), sendBuffer_.data() + messageSize); + messageSize += sizeof(addValue); + + sendSocket_->send(sendBuffer_.data(), messageSize); + + INFO(CONN, "EthernetConnection accumulate: dst ", dstPtr, ", value ", value); +} + void EthernetConnection::recvMessages() { - // Declarating Variables + // Declaring Variables char* ptr; uint64_t size; uint64_t recvSize; + constexpr uint64_t accumulateFlag = uint64_t{1} << uint64_t{63}; // Receiving Messages Until Connection is Closed while (!stopping_.load(std::memory_order_acquire)) { @@ -748,9 +852,14 @@ void EthernetConnection::recvMessages() { // Receiving Data Address if (!receiveFramePart(&ptr, sizeof(char*), true)) return; - // Receiving data size + // Receiving data size (MSB may indicate accumulate) if (!receiveFramePart(&size, sizeof(uint64_t), false)) return; + bool isAccumulate = (size & accumulateFlag) != 0; + if (isAccumulate) { + size &= ~accumulateFlag; // Strip the flag to get the data size. + } + #if defined(ENABLE_NPKIT) && defined(ENABLE_NPKIT_EVENT_CONN_ETH_RECV_META_EXIT) NpKit::CollectCpuEvent(NPKIT_EVENT_CONN_ETH_RECV_META_EXIT, uint32_t(size), 0, *NpKit::GetCpuTimestamp(), 1); #endif @@ -759,14 +868,29 @@ void EthernetConnection::recvMessages() { NpKit::CollectCpuEvent(NPKIT_EVENT_CONN_ETH_RECV_DATA_ENTRY, uint32_t(size), 0, *NpKit::GetCpuTimestamp(), 1); #endif - // Receiving Data and Copying Data yo GPU - recvSize = 0; - while (recvSize < size) { - uint64_t messageSize = std::min(recvBufferSize_, (size - recvSize) / sizeof(char)) * sizeof(char); - if (!receiveFramePart(recvBuffer_.data(), messageSize, false)) return; - - mscclpp::gpuMemcpy(ptr + (recvSize / sizeof(char)), recvBuffer_.data(), messageSize, cudaMemcpyHostToDevice); - recvSize += messageSize; + if (isAccumulate && size == sizeof(uint64_t)) { + // Accumulate modulo 2^64: receive the operand, then read, add, and write back. + uint64_t addValue; + if (!receiveFramePart(&addValue, sizeof(uint64_t), false)) return; + + // Every peer terminates its socket in this process, so several recv threads can be here + // at once for one address. The read-modify-write below is not atomic, so serialize it + // against the other recv threads. + const std::lock_guard lock(accumulateMutex()); + uint64_t current; + mscclpp::gpuMemcpy(reinterpret_cast(¤t), ptr, sizeof(uint64_t), cudaMemcpyDeviceToHost); + current += addValue; + mscclpp::gpuMemcpy(ptr, reinterpret_cast(¤t), sizeof(uint64_t), cudaMemcpyHostToDevice); + } else { + // Regular write: receive data and copy to GPU. + recvSize = 0; + while (recvSize < size) { + uint64_t messageSize = std::min(recvBufferSize_, (size - recvSize) / sizeof(char)) * sizeof(char); + if (!receiveFramePart(recvBuffer_.data(), messageSize, false)) return; + + mscclpp::gpuMemcpy(ptr + (recvSize / sizeof(char)), recvBuffer_.data(), messageSize, cudaMemcpyHostToDevice); + recvSize += messageSize; + } } #if defined(ENABLE_NPKIT) && defined(ENABLE_NPKIT_EVENT_CONN_ETH_RECV_DATA_EXIT) diff --git a/src/core/context.cc b/src/core/context.cc index b55939e3a..dd1da0928 100644 --- a/src/core/context.cc +++ b/src/core/context.cc @@ -17,9 +17,14 @@ CudaIpcStream::CudaIpcStream(int deviceId) : stream_(std::make_shared()), deviceId_(deviceId), dirty_(false) {} void CudaIpcStream::setStreamIfNeeded() { - if (!env()->cudaIpcUseDefaultStream && stream_->empty()) { - stream_->set(cudaStreamNonBlocking); - } +#if defined(MSCCLPP_USE_ROCM) + // A proxy-launched ROCm accumulate kernel must not inherit default-stream synchronization: + // the caller may be waiting for that kernel from another stream. Keep all operations on this + // connection ordered on its dedicated nonblocking stream. + if (stream_->empty()) stream_->set(cudaStreamNonBlocking); +#else + if (!env()->cudaIpcUseDefaultStream && stream_->empty()) stream_->set(cudaStreamNonBlocking); +#endif // defined(MSCCLPP_USE_ROCM) } void CudaIpcStream::memcpyD2D(void* dst, const void* src, size_t nbytes) { diff --git a/src/core/include/connection.hpp b/src/core/include/connection.hpp index d02847a60..a11e2d620 100644 --- a/src/core/include/connection.hpp +++ b/src/core/include/connection.hpp @@ -36,6 +36,8 @@ class BaseConnection { virtual void updateAndSync(RegisteredMemory dst, uint64_t dstOffset, uint64_t* src, uint64_t newValue) = 0; + virtual void accumulate(RegisteredMemory dst, uint64_t dstOffset, int64_t value) = 0; + virtual void flush(int64_t timeoutUsec = -1) = 0; /// Start signal forwarding to the given memory address. @@ -95,7 +97,7 @@ class BaseConnection { int maxWriteQueueSize_; // GPU-visible flush-done position (host-pinned memory). ProxyService writes one past the - // highest FIFO position whose TriggerSync request has fully completed on this connection + // highest FIFO position whose TriggerFlush request has fully completed on this connection // (CQ drained for IB, synchronous flush() returned for non-IB). std::shared_ptr gpuFlushDonePos_; }; @@ -114,6 +116,7 @@ class CudaIpcConnection : public BaseConnection { void write(RegisteredMemory dst, uint64_t dstOffset, RegisteredMemory src, uint64_t srcOffset, uint64_t size) override; void updateAndSync(RegisteredMemory dst, uint64_t dstOffset, uint64_t* src, uint64_t newValue) override; + void accumulate(RegisteredMemory dst, uint64_t dstOffset, int64_t value) override; void flush(int64_t timeoutUsec) override; }; @@ -170,6 +173,7 @@ class IBConnection : public BaseConnection { void write(RegisteredMemory dst, uint64_t dstOffset, RegisteredMemory src, uint64_t srcOffset, uint64_t size) override; void updateAndSync(RegisteredMemory dst, uint64_t dstOffset, uint64_t* src, uint64_t newValue) override; + void accumulate(RegisteredMemory dst, uint64_t dstOffset, int64_t value) override; void flush(int64_t timeoutUsec) override; @@ -210,6 +214,7 @@ class EthernetConnection : public BaseConnection { void write(RegisteredMemory dst, uint64_t dstOffset, RegisteredMemory src, uint64_t srcOffset, uint64_t size) override; void updateAndSync(RegisteredMemory dst, uint64_t dstOffset, uint64_t* src, uint64_t newValue) override; + void accumulate(RegisteredMemory dst, uint64_t dstOffset, int64_t value) override; void flush(int64_t timeoutUsec) override; }; diff --git a/src/core/include/context.hpp b/src/core/include/context.hpp index 42d03db15..2c8238c80 100644 --- a/src/core/include/context.hpp +++ b/src/core/include/context.hpp @@ -28,6 +28,12 @@ class CudaIpcStream { void memcpyH2D(void* dst, const void* src, size_t nbytes); +#if defined(MSCCLPP_USE_ROCM) + /// Add a value to a 64-bit integer in peer memory, with a kernel on this stream. ROCm only: + /// on CUDA such a kernel cannot be scheduled while the caller's kernel spins. + void accumulate(uint64_t* dst, uint64_t value); +#endif // defined(MSCCLPP_USE_ROCM) + void sync(); operator cudaStream_t() const { return *stream_; } diff --git a/src/core/port_channel.cc b/src/core/port_channel.cc index 0601ef84b..6cb0dacdc 100644 --- a/src/core/port_channel.cc +++ b/src/core/port_channel.cc @@ -90,7 +90,7 @@ MSCCLPP_API_CPP void ProxyService::startProxy(bool blocking) { proxy_->start(blo MSCCLPP_API_CPP void ProxyService::stopProxy() { proxy_->stop(); - // Drain pending TriggerSync flushes. After a bounded loop, force-unblock any still-pending + // Drain pending TriggerFlush operations. After a bounded loop, force-unblock any still-pending // GPU waiters with a sentinel write (UINT64_MAX > any FIFO position). for (int i = 0; i < 1000 && !pendingFlushPos_.empty(); ++i) { progressFlushes(); @@ -126,22 +126,55 @@ ProxyHandlerResult ProxyService::handleTrigger(ProxyTrigger trigger) { int maxWriteQueueSize = conn.getMaxWriteQueueSize(); auto& numRequests = inflightRequests_[conn.impl_]; - if (trigger.fields.type & TriggerData) { + auto put = [&]() { RegisteredMemory& dst = memories_[trigger.fields.dstMemoryId]; RegisteredMemory& src = memories_[trigger.fields.srcMemoryId]; conn.write(dst, trigger.fields.dstOffset, src, trigger.fields.srcOffset, trigger.fields.size); numRequests++; - } - - if (trigger.fields.type & TriggerFlag) { + }; + auto signal = [&]() { semaphore->signal(); numRequests++; + }; + auto accumulate = [&]() { + RegisteredMemory& dst = memories_[trigger.fields.dstMemoryId]; + // The operand is the full fst word, spanning the size and srcOffset fields. + conn.accumulate(dst, trigger.fields.dstOffset, static_cast(trigger.fst)); + numRequests++; + }; + + bool flushRequested = false; + switch (trigger.fields.type) { + case TriggerPut: + put(); + break; + case TriggerSignal: + signal(); + break; + case TriggerFlush: + flushRequested = true; + break; + case TriggerPutWithSignal: + put(); + signal(); + break; + case TriggerPutWithSignalAndFlush: + put(); + signal(); + flushRequested = true; + break; + case TriggerAccumulate: + accumulate(); + break; + default: + WARN(CONN, "unknown trigger opcode ", uint64_t(trigger.fields.type), ", ignoring the trigger"); + return ProxyHandlerResult::Continue; } - if (trigger.fields.type & TriggerSync) { - // Record this TriggerSync's FIFO position. The GPU caller is spinning on - // flushDonePos_ > pos; progressFlushes() will publish pos+1 once the CQ drains. - // Later TriggerSyncs on the same conn overwrite — CQ drain completes them all at once. + if (flushRequested) { + // Record this flush's FIFO position. The GPU caller is spinning on flushDonePos_ > pos; + // progressFlushes() publishes pos+1 once the CQ drains. A later flush on the same connection + // overwrites this entry, and the CQ drain completes them all at once. conn.impl_->requestFlush(); pendingFlushPos_[conn.impl_] = pos; numRequests = 0; diff --git a/test/mp_unit/mp_unit_tests.hpp b/test/mp_unit/mp_unit_tests.hpp index 8654ccc91..fa0d1a775 100644 --- a/test/mp_unit/mp_unit_tests.hpp +++ b/test/mp_unit/mp_unit_tests.hpp @@ -141,6 +141,22 @@ using DeviceHandle = mscclpp::DeviceHandle; using IbMode = mscclpp::EndpointConfig::Ib::Mode; +// Fan-in: every rank other than 0 accumulates into rank 0's counter, so the destination has many +// concurrent writers. An unserialized read-modify-write loses updates here, which a one-to-one +// test cannot show. +// +// Covers every transport that allows concurrent writers: IB, Ethernet, and CudaIpc on ROCm. See +// Connection::accumulate(). +class PortChannelFanInTest : public CommunicatorTestBase { + protected: + void SetUp() override; + void TearDown() override; + + void testFanIn(bool useIPC, bool useIb, bool useEthernet, IbMode ibMode = IbMode::Default); + + std::shared_ptr proxyService; +}; + class PortChannelOneToOneTest : public CommunicatorTestBase { protected: struct PingPongTestParams { @@ -161,6 +177,9 @@ class PortChannelOneToOneTest : public CommunicatorTestBase { void testPingPongPerf(PingPongTestParams params); void testPacketPingPong(bool useIbOnly, IbMode ibMode = IbMode::Default); void testPacketPingPongPerf(bool useIbOnly, IbMode ibMode = IbMode::Default); + void testAccumulate(bool useIPC, bool useIb, bool useEthernet, IbMode ibMode = IbMode::Default); + void testAccumulateRejected(mscclpp::Transport transport, IbMode ibMode, int tag, const char* backendMessage, + bool checkHugeOffset); void testBandwidth(PingPongTestParams params); void setupMultiQpChannels(int numQps, size_t elemsPerChan, IbMode ibMode, int tagBase, std::vector>& sendBuffs, diff --git a/test/mp_unit/port_channel_tests.cu b/test/mp_unit/port_channel_tests.cu index eec1760cf..d17822597 100644 --- a/test/mp_unit/port_channel_tests.cu +++ b/test/mp_unit/port_channel_tests.cu @@ -1,8 +1,11 @@ // Copyright (c) Microsoft Corporation. // Licensed under the MIT License. +#include #include +#include #include +#include #include "gdr.hpp" #include "mp_unit_tests.hpp" @@ -32,21 +35,75 @@ inline void requireGdrForIbMode(IbMode mode, mscclpp::Transport ibTransport) { } } #define REQUIRE_GDR_FOR_IB_MODE(mode) requireGdrForIbMode((mode), ibTransport) + +inline void requireGdrForHostNoAtomicCollective(int firstRank, int secondRank) { + bool participates = gEnv->rank == firstRank || gEnv->rank == secondRank; + int localReady = (!participates || mscclpp::gdrEnabled()) ? 1 : 0; + int allReady = 0; + MPI_Allreduce(&localReady, &allReady, 1, MPI_INT, MPI_MIN, MPI_COMM_WORLD); + if (!allReady) { + SKIP_TEST() << "HostNoAtomic rejection requires GDRCopy on every participating CUDA rank"; + } +} #else #define REQUIRE_GDR_FOR_IB_MODE(mode) // No extra requirements on non-CUDA platforms. +inline void requireGdrForHostNoAtomicCollective(int, int) {} #endif -// Skip an IPC-only PortChannel test (useIPC=true, useIB=false, useEthernet=false) when CudaIpc -// cannot connect this rank pair. CudaIpc works intra-node always, and cross-node only on MNNVL -// systems (GB200 NVL72 + IMEX). The combined check is "at least 2 ranks per node" OR "fabric -// (MNNVL) handles are usable on this system". -#define REQUIRE_CUDA_IPC_AVAILABLE \ - do { \ - if (gEnv->nRanksPerNode < 2 && !mscclpp::isFabricMemHandleAvailable()) { \ - SKIP_TEST() << "CudaIpc requires intra-node ranks (nRanksPerNode>=2) or MNNVL fabric handles, \ -both unavailable here."; \ - } \ - } while (0) +inline void requireCudaIpcRankZeroPeers(int lastPeer) { + uint64_t localHost = mscclpp::getHostHash(); + std::vector hosts(gEnv->worldSize); + MPI_Allgather(&localHost, sizeof(localHost), MPI_BYTE, hosts.data(), sizeof(localHost), MPI_BYTE, MPI_COMM_WORLD); + + bool needsFabric = false; + for (int peer = 1; peer <= lastPeer; ++peer) needsFabric |= hosts[peer] != hosts[0]; + if (!needsFabric) return; + + int localFabric = 0; + try { + localFabric = mscclpp::isFabricMemHandleAvailable() ? 1 : 0; + } catch (...) { + // A failed capability query is unavailable; every rank makes the same decision below. + } + std::vector fabricAvailable(gEnv->worldSize); + MPI_Allgather(&localFabric, 1, MPI_INT, fabricAvailable.data(), 1, MPI_INT, MPI_COMM_WORLD); + for (int peer = 1; peer <= lastPeer; ++peer) { + if (hosts[peer] != hosts[0] && (!fabricAvailable[0] || !fabricAvailable[peer])) { + SKIP_TEST() << "CudaIpc requires usable fabric memory handles on rank 0 and every cross-host peer"; + } + } +} + +#define REQUIRE_CUDA_IPC_AVAILABLE requireCudaIpcRankZeroPeers(/*lastPeer=*/1) + +inline void requireIbRdmaAtomics(mscclpp::Transport ibTransport) { + int localSupport = 0; + try { + std::string devName = mscclpp::getIBDeviceName(ibTransport); + mscclpp::IbCtx ibCtx(devName); + localSupport = ibCtx.supportsRdmaAtomics() ? 1 : 0; + } catch (...) { + // Treat an unavailable or unusable local IB device as lacking atomic support. The collective + // below then makes every rank skip together. + } + int allSupport = 0; + MPI_Allreduce(&localSupport, &allSupport, 1, MPI_INT, MPI_MIN, MPI_COMM_WORLD); + if (!allSupport) { + SKIP_TEST() << "Positive IB accumulate tests require RDMA atomic support on every rank"; + } +} + +template +bool rejectsInvalidUsage(Func func, const char* message) { + try { + func(); + } catch (const mscclpp::Error& e) { + return e.getErrorCode() == mscclpp::ErrorCode::InvalidUsage && + std::string(e.what()).find(message) != std::string::npos; + } catch (...) { + } + return false; +} void PortChannelOneToOneTest::SetUp() { // Use only two ranks @@ -626,6 +683,193 @@ PERF_TEST(PortChannelOneToOneTest, BandwidthIbHostNoAtomicMode) { .useIPC = false, .useIB = true, .useEthernet = false, .waitWithPoll = false, .ibMode = IbMode::HostNoAtomic}); } +// The high-word coefficients do not cancel: correct net = 2*2^32+4 per block. Reconstructing +// either or both operands from only their low 32 bits produces a different net. +static constexpr int64_t kAccumulatePositive = 3 * (int64_t{1} << 32) + 7; +static constexpr int64_t kAccumulateNegative = -((int64_t{1} << 32) + 3); +static constexpr int64_t kAccumulateNet = kAccumulatePositive + kAccumulateNegative; +static_assert(kAccumulateNet > 0 && kAccumulateNet <= std::numeric_limits::max() / (32 * 20)); + +__global__ void kernelPortChannelAccumulate(int64_t* localBuff, int nTries, mscclpp::DeviceSyncer* syncer, int* ret) { + auto& portChan = gChannelOneToOneTestConstPortChans; + const int numBlocks = gridDim.x; + + for (int iter = 0; iter < nTries; ++iter) { + portChan.accumulate(0, 0); + portChan.accumulate(0, kAccumulatePositive); + portChan.accumulate(0, kAccumulateNegative); + syncer->sync(numBlocks); + + if (blockIdx.x == 0) { + // The signal/flush after every block's additions validates proxy ordering before the peer wait. + portChan.signal(); + portChan.flush(); + portChan.wait(); + + const int64_t perIter = static_cast(numBlocks) * kAccumulateNet; + const int64_t expected = static_cast(iter + 1) * perIter; + const int64_t observed = *(volatile int64_t*)localBuff; + const bool finalIter = iter + 1 == nTries; + if (observed < expected || (finalIter && observed != expected)) { + printf("iter %d (final %d): buff = %lld, expected %lld\n", iter, (int)finalIter, (long long)observed, + (long long)expected); + *ret = 1; + } + } + syncer->sync(numBlocks); + } +} + +__global__ void kernelPortChannelAccumulateWrap(uint64_t* localBuff, int* ret) { + auto& portChan = gChannelOneToOneTestConstPortChans; + portChan.accumulate(0, 1); + portChan.signal(); + portChan.flush(); + portChan.wait(); + if (*(volatile uint64_t*)localBuff != 0) *ret = 1; +} + +void PortChannelOneToOneTest::testAccumulate(bool useIPC, bool useIb, bool useEthernet, IbMode ibMode) { + if (gEnv->rank >= numRanksToUse) return; + + const int nElem = 1; + std::vector portChannels; + auto buff = mscclpp::GpuBuffer(nElem); + MSCCLPP_CUDATHROW(cudaMemset(buff.memory().get(), 0, nElem * sizeof(int64_t))); + + setupMeshConnections(portChannels, useIPC, useIb, useEthernet, buff.memory().get(), nElem * sizeof(int64_t), nullptr, + 0, ibMode); + ASSERT_EQ(portChannels.size(), 1); + + auto handle = portChannels[0].deviceHandle(); + MSCCLPP_CUDATHROW(cudaMemcpyToSymbol(gChannelOneToOneTestConstPortChans, &handle, sizeof(handle))); + auto syncer = mscclpp::detail::gpuCallocShared(); + auto ret = mscclpp::detail::gpuCallocHostShared(); + *ret = 0; + + proxyService->startProxy(); + kernelPortChannelAccumulate<<<32, 1>>>(buff.memory().get(), 20, syncer.get(), ret.get()); + MSCCLPP_CUDATHROW(cudaDeviceSynchronize()); + EXPECT_EQ(*ret, 0); + + // Every transport defines accumulation modulo 2^64. + MSCCLPP_CUDATHROW(cudaMemset(buff.memory().get(), 0xff, sizeof(uint64_t))); + *ret = 0; + communicator->bootstrap()->barrier(); + kernelPortChannelAccumulateWrap<<<1, 1>>>(reinterpret_cast(buff.memory().get()), ret.get()); + MSCCLPP_CUDATHROW(cudaDeviceSynchronize()); + proxyService->stopProxy(); + EXPECT_EQ(*ret, 0); +} + +void PortChannelOneToOneTest::testAccumulateRejected(mscclpp::Transport transport, IbMode ibMode, int tag, + const char* backendMessage, bool checkHugeOffset) { + const bool participates = gEnv->rank < numRanksToUse; + bool rejectsBackend = false; + bool rejectsHuge = !checkHugeOffset; + if (participates) { + const int peer = 1 - gEnv->rank; + auto buff = mscclpp::GpuBuffer(1).memory(); + mscclpp::EndpointConfig cfg; + cfg.transport = transport; + if (transport != mscclpp::Transport::CudaIpc) { + cfg.ib.gidIndex = std::stoi(gEnv->args["ib_gid_index"]); + cfg.ib.mode = ibMode; + } + auto connFuture = communicator->connect(cfg, peer); + auto localMem = communicator->registerMemory(buff.get(), sizeof(int64_t), transport); + communicator->sendMemory(localMem, peer, tag); + auto remoteFuture = communicator->recvMemory(peer, tag); + auto conn = connFuture.get(); + auto remoteMem = remoteFuture.get(); + + rejectsBackend = rejectsInvalidUsage([&] { conn.accumulate(remoteMem, 0, 1); }, backendMessage); + if (checkHugeOffset) { + rejectsHuge = rejectsInvalidUsage([&] { conn.accumulate(remoteMem, std::numeric_limits::max(), 1); }, + "out of bounds"); + } + } + + // Only ranks 0 and 1 own the fixture bootstrap; synchronize the full test world after teardown. + MPI_Barrier(MPI_COMM_WORLD); + if (participates) { + EXPECT_TRUE(rejectsBackend); + EXPECT_TRUE(rejectsHuge); + } +} + +#if defined(__HIP_PLATFORM_AMD__) +TEST(PortChannelOneToOneTest, AccumulateCudaIpc) { + REQUIRE_CUDA_IPC_AVAILABLE; + testAccumulate(true, false, false); +} +#else +TEST(PortChannelOneToOneTest, AccumulateCudaIpcRejected) { + REQUIRE_CUDA_IPC_AVAILABLE; + testAccumulateRejected(mscclpp::Transport::CudaIpc, IbMode::Default, /*tag=*/78, "not supported over CudaIpc on CUDA", + /*checkHugeOffset=*/false); +} +#endif + +TEST(PortChannelOneToOneTest, AccumulateIb) { + REQUIRE_IBVERBS; + requireIbRdmaAtomics(ibTransport); + REQUIRE_GDR_FOR_IB_MODE(IbMode::Host); + testAccumulate(false, true, false, IbMode::Host); +} + +TEST(PortChannelOneToOneTest, AccumulateEthernet) { testAccumulate(false, false, true); } + +TEST(PortChannelOneToOneTest, AccumulateEthernetRejectsInvalidTargets) { + const bool participates = gEnv->rank < numRanksToUse; + bool rejectsUndersized = false; + bool rejectsExactEnd = false; + bool rejectsStraddling = false; + bool rejectsMisalignment = false; + + if (participates) { + const int peer = 1 - gEnv->rank; + // The allocations are deliberately larger than their registrations, so a missing bounds check + // remains within an allocation while the test reports the failure. + auto backing = mscclpp::GpuBuffer(64).memory(); + auto smallBacking = mscclpp::GpuBuffer(64).memory(); + + mscclpp::EndpointConfig cfg; + cfg.transport = mscclpp::Transport::Ethernet; + auto connFuture = communicator->connect(cfg, peer); + auto localMem = communicator->registerMemory(backing.get(), 16, mscclpp::Transport::Ethernet); + auto smallLocalMem = communicator->registerMemory(smallBacking.get(), 7, mscclpp::Transport::Ethernet); + communicator->sendMemory(localMem, peer, /*tag=*/79); + communicator->sendMemory(smallLocalMem, peer, /*tag=*/80); + auto remoteFuture = communicator->recvMemory(peer, /*tag=*/79); + auto smallRemoteFuture = communicator->recvMemory(peer, /*tag=*/80); + + auto conn = connFuture.get(); + auto remoteMem = remoteFuture.get(); + auto smallRemoteMem = smallRemoteFuture.get(); + rejectsUndersized = rejectsInvalidUsage([&] { conn.accumulate(smallRemoteMem, 0, 1); }, "out of bounds"); + rejectsExactEnd = rejectsInvalidUsage([&] { conn.accumulate(remoteMem, 16, 1); }, "out of bounds"); + rejectsStraddling = rejectsInvalidUsage([&] { conn.accumulate(remoteMem, 12, 1); }, "out of bounds"); + rejectsMisalignment = rejectsInvalidUsage([&] { conn.accumulate(remoteMem, 1, 1); }, "aligned"); + } + + // Only ranks 0 and 1 own the fixture bootstrap; synchronize the full test world after teardown. + MPI_Barrier(MPI_COMM_WORLD); + if (participates) { + EXPECT_TRUE(rejectsUndersized); + EXPECT_TRUE(rejectsExactEnd); + EXPECT_TRUE(rejectsStraddling); + EXPECT_TRUE(rejectsMisalignment); + } +} + +TEST(PortChannelOneToOneTest, AccumulateIbHostNoAtomicRejected) { + REQUIRE_IBVERBS; + requireGdrForHostNoAtomicCollective(/*firstRank=*/0, /*secondRank=*/1); + testAccumulateRejected(ibTransport, IbMode::HostNoAtomic, /*tag=*/77, "not supported in IB no-atomic mode", + /*checkHugeOffset=*/true); +} + static constexpr int kMaxQps = 4; __constant__ DeviceHandle gMultiQpPortChans[kMaxQps]; @@ -897,7 +1141,7 @@ PERF_TEST(PortChannelOneToOneTest, MultiQpFlushStressIbHostNoAtomicMode) { // Same-channel concurrent-flush kernel: N GPU threads on the same PortChannel each call // putWithSignalAndFlush in lockstep. Stresses the FIFO-position-based wait target so that -// each caller waits on its own TriggerSync rather than on a globally-incrementing counter +// each caller waits on its own TriggerFlush rather than on a globally-incrementing counter // that could be assigned out-of-order relative to the FIFO push order. __constant__ DeviceHandle gSingleChanForConcurrentFlush; @@ -936,7 +1180,7 @@ void PortChannelOneToOneTest::testSameChanConcurrentFlush(IbMode ibMode) { communicator->bootstrap()->barrier(); // Measure: a successful completion (no deadlock, no CQ error) validates that each - // concurrent-flush caller waited on its own TriggerSync (not someone else's earlier one). + // concurrent-flush caller waited on its own TriggerFlush (not someone else's earlier one). const int nIters = 500; mscclpp::Timer timer; kernelSameChanConcurrentFlush<<<1, nThreads>>>(nIters); @@ -959,3 +1203,138 @@ TEST(PortChannelOneToOneTest, SameChanConcurrentFlushIbHostMode) { REQUIRE_GDR_FOR_IB_MODE(IbMode::Host); testSameChanConcurrentFlush(IbMode::Host); } + +void PortChannelFanInTest::SetUp() { + CommunicatorTestBase::SetUp(); + proxyService = std::make_shared(); +} + +void PortChannelFanInTest::TearDown() { CommunicatorTestBase::TearDown(); } + +// Each rank other than 0 pushes nTries accumulates at rank 0's single counter. +__global__ void kernelFanInAccumulate(int nTries) { + DeviceHandle& portChan = gChannelOneToOneTestConstPortChans; + if (threadIdx.x != 0 || blockIdx.x != 0) return; + for (int i = 0; i < nTries; i++) { + portChan.accumulate(0, kAccumulatePositive); + } + portChan.flush(); +} + +void PortChannelFanInTest::testFanIn(bool useIPC, bool useIb, bool useEthernet, IbMode ibMode) { + const int worldSize = communicator->bootstrap()->getNranks(); + const int rank = communicator->bootstrap()->getRank(); + if (worldSize < 3) { + SKIP_TEST() << "Fan-in test needs at least 3 ranks to have more than one writer."; + return; + } + const int nTries = 200; + + auto buff = mscclpp::GpuBuffer(1); + MSCCLPP_CUDATHROW(cudaMemset(buff.memory().get(), 0, sizeof(int64_t))); + + // Rank 0 is the target; every other rank connects to it. + mscclpp::TransportFlags transport; + if (useIPC) transport |= mscclpp::Transport::CudaIpc; + if (useIb) transport |= ibTransport; + if (useEthernet) transport |= mscclpp::Transport::Ethernet; + + mscclpp::EndpointConfig cfg; + if (useIPC) { + cfg.transport = mscclpp::Transport::CudaIpc; + } else if (useIb) { + cfg.transport = ibTransport; + cfg.ib.gidIndex = std::stoi(gEnv->args["ib_gid_index"]); + cfg.ib.mode = ibMode; + } else { + cfg.transport = mscclpp::Transport::Ethernet; + } + + mscclpp::RegisteredMemory localMem = communicator->registerMemory(buff.memory().get(), sizeof(int64_t), transport); + registeredMemories.push_back(localMem); + + std::vector> connFutures(worldSize); + std::vector> remoteMemFutures(worldSize); + if (rank == 0) { + for (int r = 1; r < worldSize; r++) { + connFutures[r] = communicator->connect(cfg, r); + communicator->sendMemory(localMem, r); + remoteMemFutures[r] = communicator->recvMemory(r); + } + } else { + connFutures[0] = communicator->connect(cfg, 0); + communicator->sendMemory(localMem, 0); + remoteMemFutures[0] = communicator->recvMemory(0); + } + + std::vector portChannels; + if (rank == 0) { + for (int r = 1; r < worldSize; r++) { + auto sema = communicator->buildSemaphore(connFutures[r].get(), r).get(); + mscclpp::SemaphoreId cid = proxyService->addSemaphore(sema); + portChannels.emplace_back(proxyService->portChannel(cid, proxyService->addMemory(remoteMemFutures[r].get()), + proxyService->addMemory(localMem))); + registeredMemories.push_back(remoteMemFutures[r].get()); + } + } else { + auto sema = communicator->buildSemaphore(connFutures[0].get(), 0).get(); + mscclpp::SemaphoreId cid = proxyService->addSemaphore(sema); + portChannels.emplace_back(proxyService->portChannel(cid, proxyService->addMemory(remoteMemFutures[0].get()), + proxyService->addMemory(localMem))); + registeredMemories.push_back(remoteMemFutures[0].get()); + } + + proxyService->startProxy(); + + if (rank != 0) { + std::vector> handles; + handles.push_back(portChannels[0].deviceHandle()); + MSCCLPP_CUDATHROW(cudaMemcpyToSymbol(gChannelOneToOneTestConstPortChans, handles.data(), + sizeof(DeviceHandle))); + kernelFanInAccumulate<<<1, 1>>>(nTries); + MSCCLPP_CUDATHROW(cudaDeviceSynchronize()); + } + + communicator->bootstrap()->barrier(); + + if (rank == 0) { + // EthernetConnection::flush() is a no-op, so wait up to one overall deadline for the receiver + // to apply all updates. Temporary inactivity is not treated as completion. + const int64_t expected = (int64_t)(worldSize - 1) * nTries * kAccumulatePositive; + int64_t observed = 0; + auto deadline = std::chrono::steady_clock::now() + std::chrono::seconds(60); + for (;;) { + mscclpp::gpuMemcpy(reinterpret_cast(&observed), reinterpret_cast(buff.memory().get()), + sizeof(int64_t), cudaMemcpyDeviceToHost); + if (observed == expected || std::chrono::steady_clock::now() >= deadline) break; + std::this_thread::sleep_for(std::chrono::milliseconds(100)); + } + if (observed != expected) { + std::cout << "fan-in lost " << (expected - observed) / kAccumulatePositive << " of " + << (int64_t)(worldSize - 1) * nTries << " accumulates" << std::endl; + } + EXPECT_EQ(observed, expected); + } + + communicator->bootstrap()->barrier(); + proxyService->stopProxy(); + communicator->bootstrap()->barrier(); +} + +#if defined(__HIP_PLATFORM_AMD__) +// CudaIpc supports many writers on ROCm: the kernel is a real read-modify-write, so writers in +// separate processes do not lose updates. +TEST(PortChannelFanInTest, AccumulateCudaIpc) { + requireCudaIpcRankZeroPeers(gEnv->worldSize - 1); + testFanIn(true, false, false); +} +#endif // defined(__HIP_PLATFORM_AMD__) + +TEST(PortChannelFanInTest, AccumulateIb) { + REQUIRE_IBVERBS; + requireIbRdmaAtomics(ibTransport); + REQUIRE_GDR_FOR_IB_MODE(IbMode::Host); + testFanIn(false, true, false, IbMode::Host); +} + +TEST(PortChannelFanInTest, AccumulateEthernet) { testFanIn(false, false, true); }