diff --git a/projects/rccl-tests/src/CMakeLists.txt b/projects/rccl-tests/src/CMakeLists.txt index 493bf1ba7368..3f671325cce2 100644 --- a/projects/rccl-tests/src/CMakeLists.txt +++ b/projects/rccl-tests/src/CMakeLists.txt @@ -74,6 +74,7 @@ set(COMMON_FILES common.h common.cu gin_sdma_allgather_policy.h + gin_sdma_broadcast_policy.h gin_sdma_devtime.h gin_sdma_devtime_host.h nccl1_compat.h @@ -184,6 +185,16 @@ if (ENABLE_ROCSHMEM_GIN AND TARGET roc::rocshmem AND RCCL_SOURCE_DIR) endif() add_rccl_test(alltoallv) add_rccl_test(broadcast) +if (ENABLE_ROCSHMEM_GIN AND TARGET roc::rocshmem) + # GIN-SDMA hybrid Broadcast (broadcast_perf -D 3): same device-RDC link as + # alltoall_perf/all_gather_perf so the Gin*BroadcastKernel variants resolve + # the GIN device backend (SDMA puts + LSA/LL peer stores). + target_compile_definitions(broadcast_perf PRIVATE ENABLE_ROCSHMEM_GIN) + target_compile_options(broadcast_perf PRIVATE -fgpu-rdc $<$:-ggdb -O0>) + target_link_libraries(broadcast_perf PRIVATE + -Wl,--whole-archive roc::rocshmem -Wl,--no-whole-archive) + target_link_options(broadcast_perf PRIVATE -fgpu-rdc --hip-link -rdynamic) +endif() add_rccl_test(gather) add_rccl_test(hypercube) add_rccl_test(reduce_scatter) diff --git a/projects/rccl-tests/src/broadcast.cu b/projects/rccl-tests/src/broadcast.cu index 074c573b5d44..611fb6bfc131 100644 --- a/projects/rccl-tests/src/broadcast.cu +++ b/projects/rccl-tests/src/broadcast.cu @@ -5,8 +5,154 @@ * See LICENSE.txt for license information ************************************************************************/ +#include + #include "cuda_runtime.h" #include "common.h" +#include "gin_sdma_broadcast_policy.h" // pure host/device Broadcast tier policy (gin_sdma::) +#include "gin_sdma_devtime.h" // shared device-side (wall_clock64) timing scaffold +#if defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0) +#include "nccl_device.h" +// ncclGinAnvilSdmaGPUContext: read by BroadcastGetSdmaThreshold to pick up the +// backend's NCCL_GIN_ANVIL_SDMA_THRESHOLD when no per-collective override is set. +#include "nccl_device/gin/anvil_sdma/gin_anvil_sdma_device_host_common.h" +// v4u / v4u_gptr / RCCL_SYSTEM_SYNCSCOPE / RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS for +// nontemporal system-scope b128 peer stores (the store path LL/host use over xGMI). +#include "nccl_device/rccl_ptr.h" +#include "rccl_vector_types.h" +#endif + +// LL (low-latency, packed data+flag) small-message Broadcast path. Broadcast is +// one-to-all, so it maps directly onto the LL A2A session's bcast() primitive: +// only the root writes the payload (into every peer's epoch-tagged scratch), and +// every rank polls it out of its OWN scratch into its local recvbuff. The +// epoch-tag recv replaces the EXIT barrier, and each rank writes only its own +// recvbuff, so it is immune to the initData recvbuff-memset race. Unlike +// AllGather LL it keeps a single ENTRY barrier for correctness (broadcast lacks +// mutual-recv backpressure; see BroadcastLLImpl for the deadlock rationale). The +// device types come from nccl_device.h, which common.h includes on the device +// API (>=2.28) or any >=2.29 build; gate on that. +#if (defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)) || NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0) +#define BC_HAVE_LL 1 +// The compile ceiling (gin_sdma::kBroadcastLLMaxBytes, 64 KiB) and default +// LL<->LSA cutover (gin_sdma::kBroadcastLLDefaultMaxBytes, 2 KiB) live in +// gin_sdma_broadcast_policy.h so the kernel branch, the requirements sizing and +// the host unit tests all share one source. The actual cutover is runtime-gated +// via env NCCL_GIN_ANVIL_BCAST_LL_MAX_BYTES (0 = disable). LL wins for tiny +// broadcasts by dropping one of the two LSA barriers: on 8x MI355X (50 iters x 3 +// reps, 2026-07-27) it cut small-message latency a robust 13-16% at <=1 KiB and +// ~8% at 2 KiB, crossing over (~+3%) at 4 KiB, hence the 2 KiB default. +// LL scratch handle: bufHandle assigned during ncclDevCommCreate; nSlots set by +// ncclLLA2ACreateRequirement when the LL path is enabled. nSlots==0 => LL not +// configured, so the kernel uses the direct-LSA fan-out. The cutover is tunable +// via NCCL_GIN_ANVIL_BCAST_LL_MAX_BYTES= (0 = disable, unset = default +// BROADCAST_LL_DEFAULT_MAX_BYTES); the value is clamped to BROADCAST_LL_MAX_BYTES +// and the pre-sized slot count makes it the effective cutoff. +// +// thread_local, not plain static: with -t N each host thread runs its own +// requirements/DevCommCreate/runColl sequence, and g_bcastLLReq is linked into +// that thread's ncclDevCommRequirements list (`req.next = reqs->...list`). One +// shared node spliced into N different lists corrupts them; one node per thread +// does not. bufHandle itself is a device-independent uint32 resource index, so a +// thread's single handle stays valid for every device it drives under -g N. +thread_local size_t g_bcastLLMaxBytes = 0; // resolved from env at requirements time +thread_local ncclLLA2AHandle g_bcastLLHandle = {}; +thread_local ncclDevResourceRequirements g_bcastLLReq = {}; +#endif + +// --------------------------------------------------------------------------- +// Self-contained GIN-SDMA launch / threshold / put helpers. +// +// The upstream (stage2j) design carried these in common.h so every collective +// shared them. On this branch the per-collective .cu files stay self-contained +// (see all_gather.cu, which keeps its own threshold resolver and launcher), so +// Broadcast defines exactly the helpers it needs here rather than widening +// common.h. The pure size math (gin_sdma::parseSize/resolveThreshold and the +// put-segmentation kGin* constants) is reused from gin_sdma_broadcast_policy.h +// and the backend put-policy header, so there is no duplicated policy logic. +// --------------------------------------------------------------------------- + +// Sentinel meaning "no per-collective override; use the device/backend value +// (rsCtx->sdmaThreshold, populated from NCCL_GIN_ANVIL_SDMA_THRESHOLD)". +#ifndef TEST_SDMA_THRESHOLD_UNSET +#define TEST_SDMA_THRESHOLD_UNSET (gin_sdma::kThresholdUnset) +#endif + +// Parse a per-collective LSA<->GIN threshold env var (bytes; optional K/M/G +// suffix). Returns TEST_SDMA_THRESHOLD_UNSET when unset/empty/unparseable so the +// kernel falls back to the shared NCCL_GIN_ANVIL_SDMA_THRESHOLD value. Thin +// getenv() wrapper over the pure, unit-tested gin_sdma::parseSize(). +static inline size_t testParseSdmaThresholdEnv(const char* name) { + return gin_sdma::parseSize(getenv(name)); +} + +// Resolve a collective's LSA<->GIN threshold with the fallback chain: +// 1. the collective-specific env var (collVar), if set; +// 2. the shared NCCL_GIN_ANVIL_SDMA_THRESHOLD, if explicitly set; +// 3. the collective's data-driven default (collDefault). +static inline size_t testResolveSdmaThreshold(const char* collVar, size_t collDefault) { + return gin_sdma::resolveThreshold(gin_sdma::parseSize(getenv(collVar)), + gin_sdma::parseSize(getenv("NCCL_GIN_ANVIL_SDMA_THRESHOLD")), + collDefault); +} + +#if defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0) +// NOTE: large-put segmentation is NOT done here anymore. The Anvil-SDMA backend +// ncclGinApi_Put (gin_anvil_sdma.h) internally splits any put into <=128 MiB +// (gin_sdma::kGinPutSegBytes) in-order copies, carrying the remote action on the +// final segment, so a single gin.put() of an arbitrarily large message is both +// overflow-safe (30-bit SDMA copy-count) and hang-safe (MI355X 256 MiB stall) for +// ALL callers. The test therefore issues plain gin.put() calls and exercises the +// same path a real application would; see gin_anvil_sdma_put_policy.h. + +// Like common.h's testLaunchDeviceKernel but threads one extra per-collective +// size_t through to the kernel, and lets the caller pick the grid (CTA count) +// instead of the global deviceCtaCount (-V). +// +// The trailing size_t is deliberately named `kernelArg` rather than after any one +// use: the flat/LL tier passes an LSA<->GIN threshold override, while the ring +// tier passes its pipeline depth (nChunks). The requested grid must be <= the +// lsaBarrier count registered in BroadcastGetDevCommRequirements (sized to +// max(deviceCtaCount, ring CTAs)); the kernels index devComm.lsaBarrier by +// blockIdx.x, so over-launching would corrupt/hang. +template +testResult_t testLaunchDeviceKernelThresholdCtas(F kernel, void* sendbuff, size_t sendoffset, void* recvbuff, size_t recvoffset, size_t count, ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream, size_t kernelArg, int gridCtas) { + if (kernel == nullptr) return testNotImplemented; + ncclDevComm* devComm = (ncclDevComm*)comm; + ncclWindow_t sendwin = (ncclWindow_t)sendbuff; + ncclWindow_t recvwin = (ncclWindow_t)recvbuff; + if (gridCtas < 1) gridCtas = 1; + kernel<<>>(sendwin, sendoffset, recvwin, recvoffset, count, root, *devComm, kernelArg); + return testSuccess; +} + +// Variant that also forwards an LL (low-latency, packed data+flag) handle as the +// last kernel argument. Used by the Broadcast GIN kernel for its tiny-message LL +// fast path; the handle is a small POD ({bufHandle, nSlots}) assigned during +// ncclDevCommCreate. +template +testResult_t testLaunchDeviceKernelThresholdLL(F kernel, void* sendbuff, size_t sendoffset, void* recvbuff, size_t recvoffset, size_t count, ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream, size_t sdmaThresholdOverride, ncclLLA2AHandle llHandle, size_t llMaxBytes, int gridCtas) { + if (kernel == nullptr) return testNotImplemented; + ncclDevComm* devComm = (ncclDevComm*)comm; + ncclWindow_t sendwin = (ncclWindow_t)sendbuff; + ncclWindow_t recvwin = (ncclWindow_t)recvbuff; + if (gridCtas < 1) gridCtas = 1; + kernel<<>>(sendwin, sendoffset, recvwin, recvoffset, count, root, *devComm, sdmaThresholdOverride, llHandle, llMaxBytes); + return testSuccess; +} + +// Like common.h's testLaunchDeviceKernel but the caller picks the grid (CTA count). +template +testResult_t testLaunchDeviceKernelCtas(F kernel, void* sendbuff, size_t sendoffset, void* recvbuff, size_t recvoffset, size_t count, ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream, int gridCtas) { + if (kernel == nullptr) return testNotImplemented; + ncclDevComm* devComm = (ncclDevComm*)comm; + ncclWindow_t sendwin = (ncclWindow_t)sendbuff; + ncclWindow_t recvwin = (ncclWindow_t)recvbuff; + if (gridCtas < 1) gridCtas = 1; + kernel<<>>(sendwin, sendoffset, recvwin, recvoffset, count, root, *devComm); + return testSuccess; +} +#endif // ENABLE_DEVICE_API && >= 2.28.0 // Set RCCL_TESTS_ALLGATHERV_ENABLE=1 to run grouped-broadcast (AllGatherV fusion) mode static int broadcast_grouped = [] { @@ -93,6 +239,867 @@ void BroadcastGetBw(size_t count, size_t typesize, double sec, double* algBw, do } } +// CTA (stripe) count for the GIN ring broadcast kernels. The ring picks its own +// count (its throughput is CTA-bound -- it needs enough CTAs to saturate all xGMI +// links; measured on 8x MI355X @2 GiB: 32 CTAs = 238 GB/s, 64 = 345, 128 = 350; +// SAG does not scale this way). This count is DECOUPLED from -V/deviceCtaCount: +// -V defaults to 16 (tuned for every OTHER collective, several of which regress +// with more CTAs -- e.g. Reduce), so clamping the ring to -V would have starved +// it (16 CTAs ~= 140 GB/s). Instead the ring always self-selects 128 (or the +// NCCL_GIN_ANVIL_BCAST_RING_CTAS override), so the promoted 350 GB/s holds under +// the bare default AND under the gate/board's -V 32. Only power-of-2 counts avoid +// a wave-quantization cliff (non-pow2 like 96 collapse to ~140 GB/s), so the +// result is rounded DOWN to a power of 2 in [1,128]. Read identically here and in +// BroadcastGetDevCommRequirements, which allocates max(deviceCtaCount, this) LSA +// barriers so the launched grid never exceeds the allocated lsaBarrier count +// (kernels index devComm.lsaBarrier by blockIdx.x, so over-launching would +// corrupt/hang). +static inline int bcastRingCtasUncached() { + int pref = 128; + const char* e = getenv("NCCL_GIN_ANVIL_BCAST_RING_CTAS"); + // Base 10 and a full-string check, matching BroadcastParseCtasEnv: the old + // strtol(e, nullptr, 0) accepted "64foo" as 64 and read "010" as octal 8. + if (e != nullptr && e[0] != '\0' && e[0] != '-') { + char* end = nullptr; + long p = strtol(e, &end, 10); + // Clamp in long before narrowing to int: huge values wrap to small positives + // (e.g. 4294967296 -> 0 -> floored to 1) instead of hitting the 128 hard cap. + if (end != e && *end == '\0' && p >= 1) { + if (p > 128) p = 128; + pref = (int)p; + } + } + int v = pref; // ring's own count, independent of -V + if (v < 1) v = 1; + if (v > 128) v = 128; // hard cap (barrier resource / tested range) + int p2 = 1; + while ((p2 << 1) <= v) p2 <<= 1; // largest power of 2 <= v + return p2; +} + +// Cached: BroadcastGetDevCommRequirements uses this to size lsaBarrierCount and +// BroadcastRunColl uses it to pick the launched grid. Re-reading getenv at each +// call site made those two independent reads of a mutable environment, and the +// comment above is explicit that a grid larger than the barrier pool corrupts or +// hangs. One function-local static (thread-safe initialization since C++11) makes +// the launched grid provably the same value the pool was sized from. +static inline int bcastRingCtas() { + static const int v = bcastRingCtasUncached(); + return v; +} + +// Parse NCCL_GIN_ANVIL_BCAST_CTAS (diagnostic pin) into a size_t, returning the +// "unset" sentinel when absent/empty/unparseable so bcastHybridCtas()/bcastSagCtas() +// fall back to the size-adaptive ladder. +static inline size_t BroadcastParseCtasEnv() { + const char* e = getenv("NCCL_GIN_ANVIL_BCAST_CTAS"); + // strtoull() wraps a leading '-' into a huge unsigned value, and stops at the + // first non-digit without complaining, so "-1" and "8foo" would both pin a + // nonsense CTA count. Reject both, plus a literal value colliding with the + // sentinel. Mirrors gin_sdma_allgather::parseAllGatherCtasEnv(). + if (e == nullptr || e[0] == '\0' || e[0] == '-') return gin_sdma::kBroadcastCtasUnset; + char* end = nullptr; + unsigned long long v = strtoull(e, &end, 10); + if (end == e || *end != '\0') return gin_sdma::kBroadcastCtasUnset; + if ((size_t)v == gin_sdma::kBroadcastCtasUnset) return gin_sdma::kBroadcastCtasUnset; + return (size_t)v; +} + +#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0) +testResult_t BroadcastGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* reqs, ncclComm_t comm) { + if (!reqs || !comm) return testInternalError; + + ncclCommProperties_t commProperties = NCCL_COMM_PROPERTIES_INITIALIZER; + if (ncclCommQueryProperties(comm, &commProperties) != ncclSuccess) { + return testNcclError; + } + + switch(deviceImpl) { + case 3: { // GinHybridBroadcastKernel: LSA direct (small) + root GIN puts (large) + if (commProperties.ginType == NCCL_GIN_TYPE_NONE) { + fprintf(stderr, "This test requires GIN support, but GIN support is not enabled for this communicator.\n"); + return testInternalError; + } + // Barriers must cover BOTH the size-adaptive hybrid/SAG grids (pool = + // max(-V, ladder peak)) and the ring's own (larger) CTA count. + { + int barCtas = gin_sdma::bcastPoolCtas(deviceCtaCount, bcastRingCtas()); + reqs->barrierCount = barCtas; + reqs->lsaBarrierCount = barCtas; + } + // >=2 so the scatter+allgather large tier can use two independent + // signal indices (scatter=0, gather=1); the flat/LSA paths use only 0. + reqs->ginSignalCount = gin_sdma::bcastSignalCount(deviceCtaCount); + // LL scratch for the tiny-message fast path (single CTA => nBlocks=1), + // on by default up to BROADCAST_LL_DEFAULT_MAX_BYTES, tunable via + // NCCL_GIN_ANVIL_BCAST_LL_MAX_BYTES (bytes; 0 = disable). Broadcast carries + // a single message (only the root sends), so a receiver needs just cap/8 + // u64 slots -- unlike AllGather/AllToAll which need nRanks*cap/8. + { + g_bcastLLMaxBytes = gin_sdma::resolveLLCap( + testParseSdmaThresholdEnv("NCCL_GIN_ANVIL_BCAST_LL_MAX_BYTES"), + gin_sdma::kBroadcastLLDefaultMaxBytes, gin_sdma::kBroadcastLLMaxBytes); + if (g_bcastLLMaxBytes > 0) { + int llMaxElts = (int)(g_bcastLLMaxBytes / 8); + int nSlots = ncclLLA2ACalcSlots(llMaxElts, /*maxEltSize=*/8); + ncclLLA2ACreateRequirement(/*nBlocks=*/1, nSlots, &g_bcastLLHandle, &g_bcastLLReq); + g_bcastLLReq.next = reqs->resourceRequirementsList; + reqs->resourceRequirementsList = &g_bcastLLReq; + } + } +#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,7) + reqs->ginConnectionType = NCCL_GIN_CONNECTION_FULL; +#else + reqs->ginForceEnable = true; +#endif + return testSuccess; + } + default: + return testNotImplemented; + } +} +#elif defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0) +bool BroadcastGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* reqs) { + if (!reqs) return false; + memset(reqs, 0, sizeof(*reqs)); + + switch(deviceImpl) { +#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7) + case 3: { // GinHybridBroadcastKernel: LSA direct (small) + root GIN puts (large) + // Cover hybrid/SAG size-adaptive grids and the ring's own larger CTA count. + int barCtas = gin_sdma::bcastPoolCtas(deviceCtaCount, bcastRingCtas()); + reqs->barrierCount = barCtas; + reqs->lsaBarrierCount = barCtas; + // >=2 for the scatter+allgather large tier's two signal indices. + reqs->ginSignalCount = gin_sdma::bcastSignalCount(deviceCtaCount); + // No LL resource is provisioned on this older two-argument requirements + // ABI, unlike the >=2.29 arm above. The LL branch still compiles here + // (BC_HAVE_LL follows from ENABLE_DEVICE_API && >=2.28), but + // g_bcastLLHandle.nSlots stays 0, so bcastLLEligible is false and every + // message takes the LSA/GIN tiers. Wiring LL up on 2.28.x needs a build of + // that line to validate against, which this work has not been tested on. + return true; + } +#endif + default: + return false; + } +} +#endif + +#if defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0) +#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7) +__device__ size_t BroadcastGetSdmaThreshold(struct ncclDevComm const& devComm) { + using nccl::utility::loadConst; + if (devComm.ginConnectionCount == 0 || devComm.ginHandles[0] == nullptr) { + return NCCL_GIN_ANVIL_SDMA_THRESHOLD_DEFAULT; + } + ncclGinAnvilSdmaGPUContext* rsCtx = + (ncclGinAnvilSdmaGPUContext*)devComm.ginHandles[0]; + if (rsCtx == nullptr || + loadConst(&rsCtx->layoutMagic) != NCCL_GIN_ANVIL_SDMA_LAYOUT_MAGIC) { + return NCCL_GIN_ANVIL_SDMA_THRESHOLD_DEFAULT; + } + return loadConst(&rsCtx->sdmaThreshold); +} + +// Local send->recv copy on the root (out-of-place). Uses 16-byte vector stores +// when both buffers and the byte count are 16-byte aligned, else scalar. +template +__device__ void BroadcastLocalCopy(T* dst, const T* src, size_t count, int tid, int nthreads) { + const size_t bytes = count * sizeof(T); + const uintptr_t da = (uintptr_t)dst; + const uintptr_t sa = (uintptr_t)src; + if ((bytes % sizeof(uint4)) == 0 && (da % sizeof(uint4)) == 0 && (sa % sizeof(uint4)) == 0) { + uint4* d4 = (uint4*)dst; + const uint4* s4 = (const uint4*)src; + const size_t n4 = bytes / sizeof(uint4); + for (size_t i = tid; i < n4; i += nthreads) d4[i] = s4[i]; + } else { + for (size_t i = tid; i < count; i += nthreads) dst[i] = src[i]; + } +} + +// Streaming peer copy: like BroadcastLocalCopy but the DESTINATION stores use +// nontemporal system-scope 128-bit writes (__builtin_amdgcn_global_store_b128 with +// RCCL_SYSTEM_SYNCSCOPE) instead of cached uint4 stores. This is the store path the +// LL A2A primitive and the host ring use to push xGMI writes at full rate without +// polluting cache; the source read stays a normal (cached) load since it is local. +// Falls back to BroadcastLocalCopy when the b128 builtin or 16B alignment is absent. +template +__device__ void BroadcastPeerCopyStream(T* dst, const T* src, size_t count, int tid, int nthreads) { +#if RCCL_HAVE_GLOBAL_DWORDX4_BUILTINS + const size_t bytes = count * sizeof(T); + const uintptr_t da = (uintptr_t)dst; + const uintptr_t sa = (uintptr_t)src; + if ((bytes % sizeof(uint4)) == 0 && (da % sizeof(uint4)) == 0 && (sa % sizeof(uint4)) == 0) { + const uint4* s4 = (const uint4*)src; + uint4* d4 = (uint4*)dst; + const size_t n4 = bytes / sizeof(uint4); + for (size_t i = tid; i < n4; i += nthreads) { + union { uint4 u; v4u v; } cv; + cv.u = s4[i]; + __builtin_amdgcn_global_store_b128((v4u_gptr)(d4 + i), cv.v, RCCL_SYSTEM_SYNCSCOPE); + } + return; + } +#endif + BroadcastLocalCopy(dst, src, count, tid, nthreads); +} + +// Root reads one send slice and writes the same recv offset on every LSA peer +// (including itself, which also covers the out-of-place self copy). +template +__device__ void BroadcastLsaDirect(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int nRanks, int tid, int nthreads) { + T* src = (T*)ncclGetLocalPointer(sendwin, sendoffset); + for (size_t i = tid; i < count; i += nthreads) { + T value = src[i]; + for (int lp = 0; lp < nRanks; lp++) { + T* dst = (T*)ncclGetLsaPointer(recvwin, recvoffset, lp); + dst[i] = value; + } + } +} + +// LL (low-latency) Broadcast for tiny messages, single CTA. Only the root has +// the payload, so only the root bcast()s its message (as 8-byte units) into the +// epoch-tagged LL scratch of every peer; every rank (root included) then polls +// the message out of its OWN scratch into its local recvbuff. The recv's epoch +// tag replaces the exit barrier (a non-root knows the payload arrived when the +// tag matches), and cross-rank traffic stays in the LL scratch (each rank writes +// only its own recvbuff), so this is immune to the initData recvbuff-memset race. +// +// Why an ENTRY barrier is still required (unlike AllGather LL, which needs none). +// The LL session double-buffers by epoch parity ((epoch&1)*nSlots), so it only +// tolerates a <2-epoch skew between writer and reader. In AllGather every rank +// both sends and receives, so the mutual recv self-throttles all ranks to within +// one epoch. Broadcast has no such backpressure: only the root writes, so across +// the perf loop's back-to-back kernel launches a fast root can run >=2 epochs +// ahead of a lagging non-root and overwrite a slot region that rank is still +// polling with a newer tag -> the reader's tag never matches -> deadlock +// (reproduced on MI355 as a hang). The entry LSA barrier bounds the skew to <1 +// epoch, making the double-buffer safe; the exit barrier stays removed. +__device__ void BroadcastLLImpl(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, + size_t chunkU64, int root, struct ncclDevComm const& devComm, ncclTeam lsa, + ncclLLA2AHandle llHandle) { + const int tid = threadIdx.x; + const int nthreads = blockDim.x; + uint64_t* dst = (uint64_t*)ncclGetLocalPointer(recvwin, recvoffset); + + // Entry barrier: bounds root-vs-laggard epoch skew (see note above). + ncclLsaBarrierSession lsaBar { ncclCoopCta(), devComm, lsa, devComm.lsaBarrier, /*block=*/0 }; + lsaBar.sync(ncclCoopCta(), cuda::memory_order_relaxed); + + ncclLLA2ASession ll { ncclCoopCta(), devComm, lsa, llHandle, /*block=*/0, + /*maxElts=*/(int)chunkU64 }; + + // Root broadcasts its message into every peer's scratch slot [0..chunkU64). + // NB: ll.bcast() already fans out to all N peers per call with an 8-way + // unrolled, register-reusing store loop, so it is more efficient than an + // explicit per-(peer,slot) send flatten across the LL range (measured: the + // flatten helps only <=512 B but regresses 1-2 KiB by 20-40%). Unlike Scatter + // (which had a genuinely serial explicit-send loop), keep bcast() here. + if (devComm.rank == root) { + const uint64_t* src = (const uint64_t*)ncclGetLocalPointer(sendwin, sendoffset); + for (size_t j = tid; j < chunkU64; j += nthreads) { + ll.bcast((int)j, src[j]); + } + } + // Every rank (root included) gathers the message out of its own scratch. + for (size_t j = tid; j < chunkU64; j += nthreads) { + dst[j] = ll.recv((int)j); + } + ll.endEpoch(ncclCoopCta()); +} + +// Single-node hybrid Broadcast (-D 3): flat / star fan-out from root. +// msgBytes <= LL cap (opt-in): LL packed data+flag, single CTA (tiny, no barrier). +// msgBytes <= sdmaThreshold: root writes every peer's recvbuff via LSA. +// msgBytes > sdmaThreshold: root issues one GIN put per non-self peer +// (Anvil picks IPC or SDMA by size). +// Non-root ranks only participate in the entry/exit barriers (LSA path). +// +// sdmaThresholdOverride lets the Broadcast-specific env var +// NCCL_GIN_ANVIL_SDMA_THRESHOLD_BROADCAST tune this LSA<->GIN cutover +// independently. The host always resolves it to a concrete value before launch +// (testResolveSdmaThreshold falls back to NCCL_GIN_ANVIL_SDMA_THRESHOLD and then +// to kBroadcastSdmaThresholdDefault), so the TEST_SDMA_THRESHOLD_UNSET arm below +// is unreachable from this caller; it is kept only so the kernel stays callable +// with the sentinel, matching AllGather. +// +// The LL tier is on by default up to BROADCAST_LL_DEFAULT_MAX_BYTES (disable with +// NCCL_GIN_ANVIL_BCAST_LL_MAX_BYTES=0, which leaves llHandle.nSlots==0). +template +__global__ void GinHybridBroadcastKernel(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm, size_t sdmaThresholdOverride, ncclLLA2AHandle llHandle, size_t llMaxBytes) { + const size_t msgBytes = count * sizeof(T); + const size_t sdmaThreshold = (sdmaThresholdOverride != TEST_SDMA_THRESHOLD_UNSET) + ? sdmaThresholdOverride + : BroadcastGetSdmaThreshold(devComm); + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + const int nthreads = blockDim.x * gridDim.x; + + // Switch on the shared tier function rather than re-deriving the ladder here, + // so the decision the host unit tests assert is literally the one that runs. + // llMaxBytes is the env-RESOLVED cap the host used to size the slot pool and to + // pick this grid, not gin_sdma::kBroadcastLLMaxBytes (the 64 KiB compile + // ceiling). Passing the resolved value keeps both sides on the same bound: the + // two used to agree only because bcastLLEligible's second test (msgBytes/8 <= + // nSlots) re-imposed the cap via the slot count, so any future rounding-up in + // ncclLLA2ACalcSlots would have let the device take LL above a cap the host had + // already sized the grid for LSA against. + const gin_sdma::BcastTier tier = + gin_sdma::bcastKernelTier(msgBytes, sdmaThreshold, llHandle.nSlots, llMaxBytes); + + if (tier != gin_sdma::BcastTier::Flat) { + ncclTeam lsa = ncclTeamLsa(devComm); + + // Tiny messages: LL packed data+flag path (single CTA), barrier-free and + // memset-race-immune. Used when LL is configured (nSlots>0), the message is + // 8-byte aligned, and it fits the pre-sized slot count. + if (tier == gin_sdma::BcastTier::LL) { + if (blockIdx.x != 0) return; // single CTA + const size_t chunkU64 = msgBytes / 8; + BroadcastLLImpl(sendwin, sendoffset, recvwin, recvoffset, chunkU64, root, devComm, lsa, llHandle); + return; + } + + ncclLsaBarrierSession lsaBar { ncclCoopCta(), devComm, lsa, devComm.lsaBarrier, blockIdx.x }; + lsaBar.sync(ncclCoopCta(), cuda::memory_order_relaxed); + if (devComm.rank == root) { + BroadcastLsaDirect(sendwin, sendoffset, recvwin, recvoffset, count, lsa.nRanks, tid, nthreads); + } + lsaBar.sync(ncclCoopCta(), cuda::memory_order_release); + return; + } + + const int ginContext = 0; + const unsigned int signalIndex = 0; + ncclGin gin { devComm, ginContext }; + + // All GIN traffic on this path is confined to block 0, and so are the signal + // read and wait that bracket it. `bar` synchronizes same-index blocks across + // ranks, so it orders block n against block n of every peer and nothing else. + // With the read on every CTA, a non-root's block n could sample the signal + // AFTER the root's block 0 had already landed its put -- block 0 and block n + // share no barrier edge -- and would then wait for base+2, which nobody posts. + // Keeping the read, the puts and the wait all on block 0 puts them on the one + // edge that does order them. Other CTAs still run the grid-strided local copy. + const bool commCta = (blockIdx.x == 0); + + // ncclGin_SignalInc is a *remote* action: each put increments the receiving + // peer's signal (confirmed vs AllGather/AlltoAll, where a rank receives N puts + // and waits base+N). So the root -- which receives no puts -- must not wait on + // its own signal; instead every non-root receives exactly one put and waits + // base+1. flush() alone does not guarantee remote data has settled, so this + // receiver-side waitSignal is what makes the payload visible on non-roots. + const uint64_t signalValue = commCta ? gin.readSignal(signalIndex) : 0; + + ncclBarrierSession bar { ncclCoopCta(), ncclTeamTagWorld(), gin, blockIdx.x }; + // Entry barrier: every rank's recv buffer quiescent before root starts writing. + // Every CTA takes part -- this is a cross-rank barrier keyed by blockIdx.x, so + // dropping a block here would hang the matching block on every other rank. + bar.sync(ncclCoopCta(), cuda::memory_order_relaxed, ncclGinFenceLevel::Relaxed); + + if (devComm.rank == root) { + // Out-of-place self copy (no-op when in-place: local src == local dst). + T* lsrc = (T*)ncclGetLocalPointer(sendwin, sendoffset); + T* ldst = (T*)ncclGetLocalPointer(recvwin, recvoffset); + if (lsrc != ldst) { + BroadcastLocalCopy(ldst, lsrc, count, tid, nthreads); + } + + // Flat fan-out: one put per non-self peer; skip self. The backend segments + // large messages internally (<=128 MiB copies, signal on the final one). + // Strided by blockDim.x rather than grid-wide, so block 0 covers every rank + // even where nRanks exceeds one block's width. + if (commCta) { + for (int r = threadIdx.x; r < devComm.nRanks; r += blockDim.x) { + if (r == root) continue; + gin.put(ncclTeamWorld(devComm), r, + recvwin, recvoffset, + sendwin, sendoffset, + msgBytes, ncclGin_SignalInc{signalIndex}); + } + + // flush(): all source buffers safe to reuse once the puts are pushed. + gin.flush(ncclCoopCta()); + } + } else if (commCta) { + // Each non-root receives exactly one put from root; wait for it to settle. + gin.waitSignal(ncclCoopCta(), signalIndex, signalValue + 1); + } +} + +// Large-message Broadcast via scatter + in-place allgather (van de Geijn). +// +// The flat fan-out (GinHybridBroadcastKernel large path) is latency-optimal +// (1 hop) but its bandwidth is capped by *root egress*: the root alone pushes +// N-1 full copies of the message, so busBw = B_root_egress / N (measured +// ~60 GB/s @128 MiB on 8x MI355X). This kernel breaks that ceiling by making +// every rank forward data, the same way the AllGather GIN path reaches +// ~388 GB/s: +// +// Phase 1 (scatter): the root sends chunk r (= the r-th M/N slice) to rank r, +// so only M total bytes leave the root instead of (N-1)*M. Each non-root +// receives exactly one scatter put (signal 0) and waits base0+1; the root +// copies its own slice locally (its slot is never written by the allgather). +// Phase 2 (allgather): every rank puts its own slice into every peer's +// matching slot, so egress is distributed across all N ranks. Each rank +// receives exactly N-1 puts (signal 1) and waits base1+(N-1). +// +// Two distinct signal indices (scatter=0, gather=1) keep the per-phase completion +// counts from interleaving, so NO inter-phase barrier is needed: a non-root only +// begins forwarding after its waitSignal(0) (a CTA-wide op that also makes the +// scatter data visible), and the root reads its own slice from the stable +// sendwin (not the just-written recvwin), so its local copy needs no ordering +// vs its allgather puts. Only the entry barrier is kept (recvbuff quiescent +// before the root writes -- initData memset race, as in the flat path). +// +// Host-gated to large messages (NCCL_GIN_ANVIL_BCAST_SCATTER_AG_MIN_BYTES, +// default 2 MiB) with count >= nRanks, where the egress ceiling dominates the +// extra round + scatter setup. Requires ginSignalCount >= 2 (set in +// BroadcastGetDevCommRequirements case 3). +template +__global__ void GinScatterAllgatherBroadcastKernel(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm) { + const int N = devComm.nRanks; + const int rank = devComm.rank; + // Even split with the remainder folded into the last rank's slice (shared with + // the host unit tests via gin_sdma::sagChunk). Host guarantees count >= N. + const size_t baseCount = count / (size_t)N; + const gin_sdma::Chunk myChunk = gin_sdma::sagChunk(count, N, rank); + const size_t myCount = myChunk.count; + const size_t myByteOff = myChunk.eltOffset * sizeof(T); + const size_t myBytes = myCount * sizeof(T); + + const int tid = threadIdx.x + blockIdx.x * blockDim.x; + const int nthreads = blockDim.x * gridDim.x; + + ncclGin gin { devComm, /*context=*/0 }; + const unsigned int sigScatter = 0; + const unsigned int sigGather = 1; + + // Both phases put over ranks rather than over elements, so all GIN traffic + // here is block 0's, and both signal reads and both waits stay with it. `bar` + // only orders block n against block n of each peer, so a read on any other CTA + // could sample a signal an incoming put had already advanced and then wait for + // a value nobody posts. See the matching note in GinHybridBroadcastKernel. + const bool commCta = (blockIdx.x == 0); + const uint64_t baseScatter = commCta ? gin.readSignal(sigScatter) : 0; + const uint64_t baseGather = commCta ? gin.readSignal(sigGather) : 0; + + // Entry barrier: every rank's recvbuff quiescent before the root scatters. + // All CTAs participate; it is keyed by blockIdx.x across ranks. + ncclBarrierSession bar { ncclCoopCta(), ncclTeamTagWorld(), gin, blockIdx.x }; + bar.sync(ncclCoopCta(), cuda::memory_order_relaxed, ncclGinFenceLevel::Relaxed); + + // --- Phase 1: scatter chunk r -> rank r (root only) --- + if (rank == root) { + // The root's own slot is never written by the allgather (peers only forward + // their own slice), so fill it locally from the source. Independent of the + // allgather puts below (which read the root slice from sendwin), so no + // intra-CTA ordering is required. No-op when in-place (local src == dst). + T* lsrc = (T*)ncclGetLocalPointer(sendwin, sendoffset); + T* ldst = (T*)ncclGetLocalPointer(recvwin, recvoffset); + if (lsrc != ldst) { + const size_t myElt = (size_t)rank * baseCount; + BroadcastLocalCopy(ldst + myElt, lsrc + myElt, myCount, tid, nthreads); + } + if (commCta) { + for (int r = threadIdx.x; r < N; r += blockDim.x) { + if (r == root) continue; + const gin_sdma::Chunk rChunk = gin_sdma::sagChunk(count, N, r); + const size_t rOff = rChunk.eltOffset * sizeof(T); + // Backend segments large copies internally (<=128 MiB, 30-bit-safe). + gin.put(ncclTeamWorld(devComm), r, + recvwin, recvoffset + rOff, + sendwin, sendoffset + rOff, + rChunk.count * sizeof(T), ncclGin_SignalInc{sigScatter}); + } + } + // No intermediate flush: the scatter and allgather sources are both the + // read-only sendwin, so there is no source-reuse hazard, and completion is + // signal-based. Skipping it lets the scatter and allgather puts pipeline; + // the single flush at the end drains all dirty queues. + } else if (commCta) { + // Wait (CTA-wide) for my scatter slice before I forward it in the allgather. + // Only block 0 forwards, so only block 0 needs the scatter data visible. + gin.waitSignal(ncclCoopCta(), sigScatter, baseScatter + 1); + } + + // --- Phase 2: in-place allgather of the N slices --- + // My slice source: the root reads its stable sendwin (avoids a local-copy vs + // put ordering hazard); non-roots read the scatter result in their recvwin + // (made visible by the waitSignal above). + ncclWindow_t myWin = (rank == root) ? sendwin : recvwin; + const size_t myWinOff = (rank == root) ? (sendoffset + myByteOff) : (recvoffset + myByteOff); + if (commCta) { + for (int r = threadIdx.x; r < N; r += blockDim.x) { + if (r == rank) continue; + // Backend segments large copies internally (<=128 MiB, 30-bit-safe). + gin.put(ncclTeamWorld(devComm), r, + recvwin, recvoffset + myByteOff, + myWin, myWinOff, + myBytes, ncclGin_SignalInc{sigGather}); + } + // Every rank (root included) receives exactly N-1 allgather puts. Block 0 + // issued them all, so it is also the only CTA with queues left to drain. + gin.waitSignal(ncclCoopCta(), sigGather, baseGather + (uint64_t)(N - 1)); + gin.flush(ncclCoopCta()); + } +} + +// Large-message Broadcast via a device-side PIPELINED RING (the van de Geijn ring +// shape RCCL's host path uses). Motivation: the SAG tier above still carries a +// serial, root-only scatter phase -- ~M bytes leaving one GPU on one SDMA channel +// -- that stops being hidden as M grows, so SAG plateaus (~229 GB/s @>=256 MiB) +// while the host ring keeps climbing (~322 GB/s @2 GiB). The ring removes that +// hotspot entirely: the message is split into `nChunks` chunks streamed around a +// Hamiltonian cycle, so at steady state every ring link carries a DIFFERENT chunk +// concurrently and every GPU drives its own egress (unlike SAG's scatter, where +// one root engine drives all N-1 puts). Deep pipelining across chunks hides the +// (N-1)-hop fill latency. +// +// Forwarding uses direct xGMI peer stores issued by the whole CTA +// (ncclGetLsaPointer + vectorized 16-byte stores -- the same SM-copy mechanism +// BroadcastLsaDirect uses). Two alternatives were implemented and measured on +// 8x MI355X, and both lost, so neither ships here: +// * SDMA-put ring (one descriptor per hop issued by tid 0, gated on a +// completion signal): latency-bound, plateaus ~13 GB/s regardless of N. +// * Per-hop point-to-point signaling instead of the per-step LSA barrier +// (2026-08-04, @2 GiB): 191 GB/s vs 232 for the barrier. The barrier bundles +// the cross-GPU release fence with the sync, whereas P2P pays a +// __threadfence_system per chunk/hop and loses lock-step pipelining. +// +// Parallelism / correctness model, shared by both kernels below: +// * Each CTA owns a contiguous stripe [sBase,sEnd) of the buffer and runs an +// INDEPENDENT ring pipeline on it, so multiple CTAs stripe the message for +// aggregate xGMI bandwidth. Block i only ever touches block i's stripe on +// every rank, so a per-block-index LSA barrier is a sufficient team sync (no +// grid-wide sync needed). +// * The stripe is split into C chunks. At pipeline step s the rank at ring +// position p forwards chunk c = s-p to its successor (0<=c {1,3,5,7}, i.e. 4 +// of each GPU's 7 links; N=4 => {1,3}), which is why the table kernel below is +// preferred wherever the full decomposition exists. +template +__global__ void GinRingSmMultiBroadcastKernel(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm, size_t nChunksArg) { + const int N = devComm.nRanks; + const int rank = devComm.rank; + const int tid = threadIdx.x; + const int nthreads = blockDim.x; + + // Strides coprime to N form single Hamiltonian rings; each is one link direction. + int strides[64]; + int nStrides = 0; + for (int s = 1; s < N && nStrides < 64; s++) { + int a = s, b = N; + while (b) { int t = a % b; a = b; b = t; } // gcd(s,N) + if (a == 1) strides[nStrides++] = s; + } + if (nStrides == 0) { strides[0] = 1; nStrides = 1; } // N==1 guard + + const int s = strides[blockIdx.x % nStrides]; + int sinv = 1; // modular inverse of s mod N + for (int k = 1; k < N; k++) { if ((s * k) % N == 1) { sinv = k; break; } } + + const int pos = (int)(((long)(((rank - root) % N + N) % N) * sinv) % N); + const bool isRoot = (pos == 0); + const int nextRank = (rank + s) % N; // this ring's successor + + const gin_sdma::Chunk stripe = gin_sdma::ringStripe(count, (int)blockIdx.x, (int)gridDim.x); + const size_t sBase = stripe.eltOffset; + const size_t sCount = stripe.count; + + int C = (int)((size_t)nChunksArg < sCount ? (size_t)nChunksArg : sCount); + if (C < 1) C = 1; + + T* myRecv = (T*)ncclGetLocalPointer(recvwin, recvoffset); + T* mySend = (T*)ncclGetLocalPointer(sendwin, sendoffset); + T* peerRecv = (T*)ncclGetLsaPointer(recvwin, recvoffset, nextRank); + const bool inPlace = (mySend == myRecv); + T* fwdSrc = isRoot ? (inPlace ? myRecv : mySend) : myRecv; + + ncclTeam lsa = ncclTeamLsa(devComm); + ncclLsaBarrierSession bar { ncclCoopCta(), devComm, lsa, devComm.lsaBarrier, blockIdx.x }; + bar.sync(ncclCoopCta(), cuda::memory_order_relaxed); + + int totalSteps = C + (N - 1) - 1; + if (totalSteps < 1) totalSteps = 1; + + for (int step = 0; step < totalSteps; step++) { + const int c = step - pos; + const bool active = (pos < N - 1) && (c >= 0) && (c < C); + if (active) { + const gin_sdma::Chunk ck = gin_sdma::ringChunk(sBase, sCount, c, C); + const size_t cStart = ck.eltOffset; + const size_t cCount = ck.count; + BroadcastLocalCopy(peerRecv + cStart, fwdSrc + cStart, cCount, tid, nthreads); + if (isRoot && !inPlace) { + BroadcastLocalCopy(myRecv + cStart, mySend + cStart, cCount, tid, nthreads); + } + } + // acq_rel: next-step consumer reads peerRecv written here over xGMI. + bar.sync(ncclCoopCta(), cuda::memory_order_acq_rel); + } +} + +// --------------------------------------------------------------------------- +// FULL edge-disjoint multi-ring SM broadcast (attack (a), maximal): use ALL of +// a GPU's outgoing xGMI links by decomposing the complete symmetric digraph K_N* +// into N-1 arc-disjoint directed Hamiltonian cycles (Tillson: exists for N!=4,6). +// The coprime-stride set (GinRingSmMultiBroadcastKernel) only reaches 4 of 8 +// GPUs' 7 links (offsets 1,3,5,7 -- even offsets don't form single cycles); this +// path uses a host-computed decomposition covering every arc, so each GPU drives +// all N-1 links at once (mirrors how RCCL's 28 rotated-ring channels saturate the +// fabric). The N-1 cycles are found by verified backtracking on the host and +// uploaded to constant memory; CTA b runs ring b%nRings on buffer stripe b. +#define BCAST_RING_MAXR 16 +#define BCAST_RING_MAXN 16 +__constant__ int c_bcastNRings; +__constant__ int c_bcastRingSucc[BCAST_RING_MAXR * BCAST_RING_MAXN]; // [ring*N + rank] -> successor rank +__constant__ int c_bcastRingPos[BCAST_RING_MAXR * BCAST_RING_MAXN]; // [ring*N + rank] -> position in cycle +__constant__ int c_bcastStream; // 1 => peer writes use nontemporal system-scope b128 stores + +// Per-device record of the ring decomposition uploaded to the __constant__ tables +// above. Written by BroadcastRunColl once the decomposition for the current rank +// count is built + uploaded; read by BroadcastDeviceTime to gate whether in-kernel +// timing can run the ring body (needs the tables present on THIS device). +// +// Keyed by device ordinal because __constant__ memory is per device: a -g N run +// drives N devices from one process, and a single process-wide flag would record +// device 0's upload as covering all of them. +struct BcastRingState { + int builtN = -1; + int builtNRings = 0; +}; +#define BCAST_MAX_DEVICES 64 +static BcastRingState g_bcastRingState[BCAST_MAX_DEVICES]; + +// hipify maps cudaMemcpyToSymbolAsync(sym, src, nbytes, offset, stream) to +// hipMemcpyToSymbolAsync without hipMemcpyKind; ROCm 7.x requires the kind arg. +// Template on the symbol so scalar __constant__ ints and array tables both resolve. +template +inline cudaError_t bcastMemcpyConstantAsync(Symbol& symbol, const void* src, + size_t nbytes, cudaStream_t stream) { +#if defined(__HIPCC__) || defined(__HIP_PLATFORM_AMD__) + return hipMemcpyToSymbolAsync(symbol, src, nbytes, 0, hipMemcpyHostToDevice, stream); +#else + return cudaMemcpyToSymbolAsync(symbol, src, nbytes, 0, stream); +#endif +} + +// startColl() calls cudaSetDevice(args->gpus[i]) immediately before runColl on the +// device-API path, so the current device identifies the state slot. Out-of-range +// ordinals fall back to slot 0 rather than scribbling past the array; that only +// costs the (untested, >64 GPU) case the same sharing bug we are fixing here. +static inline BcastRingState& bcastRingState() { + int dev = 0; + if (cudaGetDevice(&dev) != cudaSuccess || dev < 0 || dev >= BCAST_MAX_DEVICES) dev = 0; + return g_bcastRingState[dev]; +} + +// Body extracted so the device-timing kernel below can run it skip+loop times +// under one persistent launch. Each call re-creates its LSA barrier session and +// re-syncs (like the AllGather timed body), so looping is self-contained. +template +__device__ void ginRingSmTableBroadcastBody(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm& devComm, size_t nChunksArg) { + const int N = devComm.nRanks; + const int rank = devComm.rank; + const int tid = threadIdx.x; + const int nthreads = blockDim.x; + + const int nRings = c_bcastNRings; + const int ring = (int)(blockIdx.x % (unsigned)nRings); + const int nextRank = c_bcastRingSucc[ring * N + rank]; + const int posCur = c_bcastRingPos[ring * N + rank]; + const int posRoot = c_bcastRingPos[ring * N + root]; + const int pos = (posCur - posRoot + N) % N; // hops from root along this ring + const bool isRoot = (pos == 0); + + const gin_sdma::Chunk stripe = gin_sdma::ringStripe(count, (int)blockIdx.x, (int)gridDim.x); + const size_t sBase = stripe.eltOffset; + const size_t sCount = stripe.count; + + int C = (int)((size_t)nChunksArg < sCount ? (size_t)nChunksArg : sCount); + if (C < 1) C = 1; + + T* myRecv = (T*)ncclGetLocalPointer(recvwin, recvoffset); + T* mySend = (T*)ncclGetLocalPointer(sendwin, sendoffset); + T* peerRecv = (T*)ncclGetLsaPointer(recvwin, recvoffset, nextRank); + const bool inPlace = (mySend == myRecv); + T* fwdSrc = isRoot ? (inPlace ? myRecv : mySend) : myRecv; + + ncclTeam lsa = ncclTeamLsa(devComm); + ncclLsaBarrierSession bar { ncclCoopCta(), devComm, lsa, devComm.lsaBarrier, blockIdx.x }; + bar.sync(ncclCoopCta(), cuda::memory_order_relaxed); + + int totalSteps = C + (N - 1) - 1; + if (totalSteps < 1) totalSteps = 1; + + for (int step = 0; step < totalSteps; step++) { + const int c = step - pos; + const bool active = (pos < N - 1) && (c >= 0) && (c < C); + if (active) { + const gin_sdma::Chunk ck = gin_sdma::ringChunk(sBase, sCount, c, C); + const size_t cStart = ck.eltOffset; + const size_t cCount = ck.count; + // Streaming (nontemporal system-scope b128) peer stores vs cached uint4 is an + // A/B knob (c_bcastStream); the local root self-copy stays cached. + if (c_bcastStream) { + BroadcastPeerCopyStream(peerRecv + cStart, fwdSrc + cStart, cCount, tid, nthreads); + } else { + BroadcastLocalCopy(peerRecv + cStart, fwdSrc + cStart, cCount, tid, nthreads); + } + if (isRoot && !inPlace) { + BroadcastLocalCopy(myRecv + cStart, mySend + cStart, cCount, tid, nthreads); + } + } + // acq_rel: the consumer at the next step reads peerRecv written here over xGMI; + // release alone does not acquire the predecessor's stores (LSA barrier maps + // release to relaxed acquire on the inbox). + bar.sync(ncclCoopCta(), cuda::memory_order_acq_rel); + } +} + +template +__global__ void GinRingSmTableBroadcastKernel(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm, size_t nChunksArg) { + ginRingSmTableBroadcastBody(sendwin, sendoffset, recvwin, recvoffset, count, root, devComm, nChunksArg); +} + +#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7) +// Device-timing kernel (shared gin_devtime methodology): run skip+loop back-to-back +// ring bodies under ONE persistent launch, bracketing only the timed region with +// wall_clock64() per CTA. min(start)..max(end) over CTAs is the grid busy window. +template +__global__ void GinRingSmTableBroadcastTimedKernel(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm, size_t nChunksArg, int loop, int skip, long long* start_time, long long* end_time) { + for (int i = 0; i < skip + loop; i++) { + if (i == skip) { + __syncthreads(); + if (threadIdx.x == 0) start_time[blockIdx.x] = wall_clock64(); + } + ginRingSmTableBroadcastBody(sendwin, sendoffset, recvwin, recvoffset, count, root, devComm, nChunksArg); + } + __syncthreads(); + if (threadIdx.x == 0) end_time[blockIdx.x] = wall_clock64(); +} +#endif + +// Host-side decomposition of K_N* into (N-1) arc-disjoint directed Hamiltonian +// cycles by verified backtracking (Tillson guarantees existence for N!=4,6; the +// search simply fails and we fall back to the coprime-stride kernel otherwise). +// Fills succ[ring*N+rank] and pos[ring*N+rank]; returns nRings (N-1) or 0. +struct BcastRingDecomp { + int N; + bool used[BCAST_RING_MAXN][BCAST_RING_MAXN]; + int order[BCAST_RING_MAXR][BCAST_RING_MAXN]; + long budget; + bool extend(int r, int* path, bool* vis, int len) { + if (--budget < 0) return false; + const int cur = path[len - 1]; + if (len == N) { + if (used[cur][0]) return false; // must close to start + for (int i = 0; i < N; i++) used[path[i]][path[(i + 1) % N]] = true; + for (int i = 0; i < N; i++) order[r][i] = path[i]; + if (solve(r + 1)) return true; + for (int i = 0; i < N; i++) used[path[i]][path[(i + 1) % N]] = false; + return false; + } + for (int nx = 1; nx < N; nx++) { // start vertex 0 fixed + if (!vis[nx] && !used[cur][nx]) { + vis[nx] = true; path[len] = nx; + if (extend(r, path, vis, len + 1)) return true; + vis[nx] = false; + } + } + return false; + } + bool solve(int r) { + if (r == N - 1) return true; // N-1 arc-disjoint cycles placed + int path[BCAST_RING_MAXN]; bool vis[BCAST_RING_MAXN]; + for (int i = 0; i < N; i++) vis[i] = false; + path[0] = 0; vis[0] = true; + return extend(r, path, vis, 1); + } +}; + +static int buildBcastRingDecomp(int N, int* succ, int* pos) { + if (N < 2 || N > BCAST_RING_MAXN || (N - 1) > BCAST_RING_MAXR) return 0; + // thread_local, not plain static: -t N host threads call this concurrently for + // their own devices and would otherwise share one backtracking scratch object. + thread_local static BcastRingDecomp d; // large; keep off-stack + d.N = N; + for (int i = 0; i < N; i++) for (int j = 0; j < N; j++) d.used[i][j] = false; + d.budget = 20000000L; // ample for N<=8; else fall back + if (!d.solve(0)) return 0; + for (int r = 0; r < N - 1; r++) { + for (int k = 0; k < N; k++) { + const int rk = d.order[r][k]; + succ[r * N + rk] = d.order[r][(k + 1) % N]; + pos[r * N + rk] = k; + } + } + return N - 1; +} +#endif +#endif + +#if defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7) +// Which -D 3 kernel actually launched. A broadcast is correct whichever tier +// serves it, so #wrong stays 0 even when a tier is removed and its messages fall +// through to the next one -- a functional test that only checks the data cannot +// tell the two apart, and the tier it meant to cover ends up covered by nothing. +// Naming the launched kernel gives those tests something that fails. +// +// Opt-in via RCCL_TESTS_BCAST_TIER_LOG so default stdout stays byte-identical +// for the perf harness. Reported once per distinct tier: RunColl is called for +// warmup, every iteration and every root, and one line per launch would flood +// the results table it shares stdout with. +enum class BcastLaunchTier { RingTable, RingMulti, ScatterAllgather, Flat, LsaDirect, LL }; + +static void bcastReportTier(BcastLaunchTier tier) { + static const char* const kTierNames[] = { + "ring-table", "ring-multi", "scatter-allgather", "flat-gin", "lsa-direct", "ll" + }; + // Env read once, like the threshold lookups above: RunColl is on the timed + // path and is called for every iteration and every root. + static const bool enabled = getenv("RCCL_TESTS_BCAST_TIER_LOG") != nullptr; + if (!enabled || !is_main_proc) return; + // -t N host threads reach this concurrently; fetch_or makes "first one wins" + // per tier rather than a read-then-set race that can print twice. + static std::atomic reported{0}; + const unsigned bit = 1u << (unsigned)tier; + if (reported.fetch_or(bit, std::memory_order_relaxed) & bit) return; + // Leading '#' keeps this out of the results table: the pytest row regex is + // anchored on a leading size field, and rccl-tests already prints #-prefixed + // side-channel lines (see #[bcast-devtime]). + printf("#[bcast-tier] %s\n", kTierNames[(unsigned)tier]); + fflush(stdout); +} +#endif + testResult_t BroadcastRunColl(void* sendbuff, size_t sendoffset, void* recvbuff, size_t recvoffset, size_t count, ncclDataType_t type, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream, int deviceImpl, void* bias = nullptr) { if (broadcast_grouped) { // nranks broadcasts with distinct roots inside one group — fuses into AllGatherV ring kernel @@ -126,10 +1133,223 @@ testResult_t BroadcastRunColl(void* sendbuff, size_t sendoffset, void* recvbuff, } #endif } else { - return testNotImplemented; + switch(deviceImpl) { +#if defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7) + case 3: { + // Broadcast-specific LSA<->GIN threshold. Default = 256 KiB (full + // message): on 8x MI355X (NCCL_GIN_TYPE=6) LSA wins <=256K and GIN wins + // >=512K (measured 2026-07-24). Override with + // NCCL_GIN_ANVIL_SDMA_THRESHOLD_BROADCAST, or the shared + // NCCL_GIN_ANVIL_SDMA_THRESHOLD. + static const size_t bcastThr = testResolveSdmaThreshold("NCCL_GIN_ANVIL_SDMA_THRESHOLD_BROADCAST", gin_sdma::kBroadcastSdmaThresholdDefault); + + // Large-message tier: scatter + in-place allgather. Distributes + // egress across all ranks to beat the flat fan-out's root-egress ceiling + // (busBw = B_root_egress / N). Gated to large messages via + // NCCL_GIN_ANVIL_BCAST_SCATTER_AG_MIN_BYTES (bytes, optional K/M/G + // suffix; default 2 MiB; 0 = disable, keeping the proven flat path). + // Crossover measured on 8x MI355X (-V 32, in-place, 2026-07-27): SAG is a + // slight win at 2M (35.9 vs flat 33.4), decisive >=4M (4M 62 vs 42, 128M + // 224 vs 60), and loses <2M (1M 19.4 vs 23.7). So default 2 MiB. + static const size_t bcastSagMin = []() { + size_t v = testParseSdmaThresholdEnv("NCCL_GIN_ANVIL_BCAST_SCATTER_AG_MIN_BYTES"); + return (v == TEST_SDMA_THRESHOLD_UNSET) ? gin_sdma::kBroadcastScatterAgMinDefault : v; + }(); + // Pipelined-ring large tier: distributes forwarding across + // all ranks with no serial root-scatter phase, to beat the SAG plateau at + // very large messages. Enabled by default from 32 MiB (the measured SAG + // crossover) via NCCL_GIN_ANVIL_BCAST_RING_MIN_BYTES, where 0 disables the + // tier and falls back to SAG; the pipeline depth is + // NCCL_GIN_ANVIL_BCAST_RING_CHUNKS (else size-derived). Selection also + // requires the LSA team to cover the world -- see bcastUseRing. + static const size_t bcastRingMin = []() { + size_t v = testParseSdmaThresholdEnv("NCCL_GIN_ANVIL_BCAST_RING_MIN_BYTES"); + return (v == TEST_SDMA_THRESHOLD_UNSET) ? gin_sdma::kBroadcastRingMinDefault : v; + }(); + static const size_t bcastRingChunksEnv = + testParseSdmaThresholdEnv("NCCL_GIN_ANVIL_BCAST_RING_CHUNKS"); + // In the -D 3 path `comm` is the ncclDevComm handle (testLaunchDeviceKernel* + // casts it), NOT a real ncclComm_t -- so read nRanks from the devComm + // struct rather than ncclCommCount (which would fault on a corrupted comm). + const int sagRanks = (int)((struct ncclDevComm*)comm)->nRanks; + // Ring-tier legality, not a performance input: the ring forwards through + // ncclGetLsaPointer with world rank indices, so it is only correct when + // the LSA team is the whole world. See bcastUseRing. + const int lsaSize = (int)((struct ncclDevComm*)comm)->lsaSize; + const size_t msgBytes = count * wordSize(type); + static const size_t bcastCtasEnv = BroadcastParseCtasEnv(); + const int hybridPool = gin_sdma::bcastHybridPoolCtas(deviceCtaCount); + // One tier decision, not two independent predicates: under the shipped + // defaults both the ring (>=32 MiB) and SAG (>=2 MiB) gates are true for + // every large message, so what runs is set by the precedence inside + // bcastLargeTier, where the host tests can assert it. + const gin_sdma::BcastLargeTier largeTier = + gin_sdma::bcastLargeTier(msgBytes, count, sagRanks, lsaSize, bcastRingMin, bcastSagMin); + if (largeTier == gin_sdma::BcastLargeTier::Ring) { + const int ringCtas = bcastRingCtas(); + const size_t nChunks = (size_t)gin_sdma::bcastRingChunks(msgBytes, ringCtas, bcastRingChunksEnv); + // The decomposition is uploaded to __constant__ memory, which exists once + // per DEVICE, so the "already built" gate has to be per device as well. + // With -g N a single process drives N devices through this same code: a + // process-wide flag would let device 0 upload, then let devices 1..N-1 + // skip their upload and run the ring against a zero-filled table + // (nRings==0 => modulo by zero, and every rank reading pos 0 => every + // rank believes it is the root). Distinct host threads own distinct + // devices, so keying on the device also keeps -t N off a shared slot. + BcastRingState& rs = bcastRingState(); + if (rs.builtN != sagRanks) { + // thread_local staging: -t N threads can be here concurrently for + // different devices and must not share the buffers they hand to + // buildBcastRingDecomp / cudaMemcpyToSymbol. + thread_local int h_succ[BCAST_RING_MAXR * BCAST_RING_MAXN]; + thread_local int h_pos[BCAST_RING_MAXR * BCAST_RING_MAXN]; + rs.builtNRings = buildBcastRingDecomp(sagRanks, h_succ, h_pos); + if (rs.builtNRings > 0) { + // Streaming (nontemporal system-scope b128) peer stores default ON; + // NCCL_GIN_ANVIL_BCAST_RING_STREAM=0 reverts to cached uint4 stores. + static const int streamStores = []() { + const char* e = getenv("NCCL_GIN_ANVIL_BCAST_RING_STREAM"); + return (e && e[0] == '0') ? 0 : 1; + }(); + CUDACHECK(bcastMemcpyConstantAsync(c_bcastNRings, &rs.builtNRings, sizeof(int), stream)); + CUDACHECK(bcastMemcpyConstantAsync(c_bcastRingSucc, h_succ, sizeof(int) * sagRanks * rs.builtNRings, stream)); + CUDACHECK(bcastMemcpyConstantAsync(c_bcastRingPos, h_pos, sizeof(int) * sagRanks * rs.builtNRings, stream)); + CUDACHECK(bcastMemcpyConstantAsync(c_bcastStream, &streamStores, sizeof(int), stream)); + } + rs.builtN = sagRanks; + } + bcastReportTier(rs.builtNRings > 0 ? BcastLaunchTier::RingTable + : BcastLaunchTier::RingMulti); + // The SM ring kernels launch their own power-of-2 CTA count (bcastRingCtas, + // default 128) instead of -V/deviceCtaCount: the pipeline is CTA-bound and + // needs ~128 CTAs to saturate all xGMI links (238 -> 350 GB/s @2 GiB). + if (rs.builtNRings > 0) { + TESTCHECK(testLaunchDeviceKernelThresholdCtas(SPECIALIZE_KERNEL(GinRingSmTableBroadcastKernel, type, op), sendbuff, sendoffset, recvbuff, recvoffset, count, type, op, root, comm, stream, nChunks, ringCtas)); + } else { + // No edge-disjoint decomposition exists for N=4 or N=6 (Tillson), and + // the backtracking search can also run out of budget; both land here. + TESTCHECK(testLaunchDeviceKernelThresholdCtas(SPECIALIZE_KERNEL(GinRingSmMultiBroadcastKernel, type, op), sendbuff, sendoffset, recvbuff, recvoffset, count, type, op, root, comm, stream, nChunks, ringCtas)); + } + return testSuccess; + } + if (largeTier == gin_sdma::BcastLargeTier::ScatterAG) { + bcastReportTier(BcastLaunchTier::ScatterAllgather); + const int sagGrid = gin_sdma::bcastSagCtas(msgBytes, sagRanks, bcastThr, bcastCtasEnv, hybridPool); + TESTCHECK(testLaunchDeviceKernelCtas(SPECIALIZE_KERNEL(GinScatterAllgatherBroadcastKernel, type, op), sendbuff, sendoffset, recvbuff, recvoffset, count, type, op, root, comm, stream, sagGrid)); + return testSuccess; + } + { + // BC_HAVE_LL is implied here: its condition is a superset of this + // block's (ENABLE_DEVICE_API && >= 2.28.0), so the LL launcher is the + // only reachable arm and the previous #else was dead. When LL is not + // provisioned -- either NCCL_GIN_ANVIL_BCAST_LL_MAX_BYTES=0, or a + // 2.28.x build whose requirements arm cannot create the LL resource -- + // nSlots stays 0, bcastLLEligible is false, and the kernel takes LSA. + const int llSlots = g_bcastLLHandle.nSlots; + // The hybrid kernel picks LL / direct-LSA / flat-GIN internally, so ask + // the same policy the kernel uses rather than reporting "hybrid" and + // leaving the three indistinguishable to the tests. + switch (gin_sdma::bcastKernelTier(msgBytes, bcastThr, llSlots, g_bcastLLMaxBytes)) { + case gin_sdma::BcastTier::LL: bcastReportTier(BcastLaunchTier::LL); break; + case gin_sdma::BcastTier::LSADirect: bcastReportTier(BcastLaunchTier::LsaDirect); break; + case gin_sdma::BcastTier::Flat: bcastReportTier(BcastLaunchTier::Flat); break; + } + const int hybridGrid = gin_sdma::bcastHybridCtas(msgBytes, bcastThr, llSlots, g_bcastLLMaxBytes, bcastCtasEnv, hybridPool); + TESTCHECK(testLaunchDeviceKernelThresholdLL(SPECIALIZE_KERNEL(GinHybridBroadcastKernel, type, op), sendbuff, sendoffset, recvbuff, recvoffset, count, type, op, root, comm, stream, bcastThr, g_bcastLLHandle, g_bcastLLMaxBytes, hybridGrid)); + } + return testSuccess; + } +#endif + default: + return testNotImplemented; + } + } + return testSuccess; +} + +// Device-side (in-kernel wall_clock64) timing for the GIN-SDMA Broadcast (-D 3), +// via the shared gin_devtime scaffold. Opt-in via --device_timing: 1=augment +// (print an extra #[bcast-devtime] line next to the graph numbers), 2=device-time-only +// (report the in-kernel latency as the out-of-place metric; in-place keeps normal +// timing). loop/skip come from --devtime_loop/--devtime_skip (default 10/10); size- +// tier overrides via --devtime_loop_mid/_large and --devtime_skip_mid/_large. Times +// the default large tier -- the SM edge-disjoint ring -- so it only runs when the +// ring is the active tier (message >= cutover, decomposition built during warmup); +// otherwise it leaves the host-timed metric untouched. +#if defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7) +testResult_t BroadcastDeviceTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int in_place, double* outDeltaSec) { + if (!deviceTimingMode) return testSuccess; + + // Only the GIN hybrid impl (-D 3) provisions the GIN signals/barriers the timed + // body relies on. Other device impls would fault or hang on missing GIN state. + if (deviceImpl != 3) return testSuccess; + + const size_t count = args->nbytes / wordSize(type); // broadcast message count + if (count == 0 || devtimeLoop < 1) return testSuccess; + const size_t msgBytes = count * wordSize(type); + + int loop = devtimeLoop; + int skip = devtimeSkip < 0 ? 0 : devtimeSkip; + if (devtimeLoopLarge > 0 && msgBytes >= (size_t)64 * 1024 * 1024) { + loop = devtimeLoopLarge; + if (devtimeSkipLarge >= 0) skip = devtimeSkipLarge; + else skip = (skip < 1) ? skip : 1; + } else if (devtimeLoopMid > 0 && msgBytes >= (size_t)8 * 1024 * 1024) { + loop = devtimeLoopMid; + if (devtimeSkipMid >= 0) skip = devtimeSkipMid; + else skip = (skip < 2) ? skip : 2; + } + + static const size_t bcastRingMin = []() { + size_t v = testParseSdmaThresholdEnv("NCCL_GIN_ANVIL_BCAST_RING_MIN_BYTES"); + return (v == TEST_SDMA_THRESHOLD_UNSET) ? gin_sdma::kBroadcastRingMinDefault : v; + }(); + const int nRanks = (int)(args->devComms[0].nRanks); + const int lsaSize = (int)(args->devComms[0].lsaSize); + // Only device-time the ring when it is the active tier and its tables are built + // on this device. builtNRings > 0 is exactly the condition under which + // BroadcastRunColl launches the table kernel, so the timed kernel below always + // matches the kernel that actually ran (N=4/6 fall back to the coprime-stride + // kernel, which has no timed variant, and are skipped here). + if (!gin_sdma::bcastUseRing(msgBytes, count, nRanks, lsaSize, bcastRingMin)) return testSuccess; + if (bcastRingState().builtNRings <= 0) return testSuccess; + + auto kernel = SPECIALIZE_KERNEL(GinRingSmTableBroadcastTimedKernel, type, op); + if (kernel == nullptr) return testSuccess; + + static const size_t chunksEnv = testParseSdmaThresholdEnv("NCCL_GIN_ANVIL_BCAST_RING_CHUNKS"); + const int gridCtas = bcastRingCtas(); + const size_t nChunks = (size_t)gin_sdma::bcastRingChunks(msgBytes, gridCtas, chunksEnv); + double devUs = 0.0; + TESTCHECK(gin_devtime::measure(args, gridCtas, loop, + [&](int i, long long* d_start, long long* d_end) { + ncclDevComm* devComm = args->devComms + i; + ncclWindow_t sendwin = (ncclWindow_t)(in_place ? args->recvRegHandles[i] : args->sendRegHandles[i]); + ncclWindow_t recvwin = (ncclWindow_t)args->recvRegHandles[i]; + kernel<<streams[i]>>>(sendwin, 0, recvwin, 0, count, root, *devComm, + nChunks, loop, skip, d_start, d_end); + }, + &devUs)); + + if (outDeltaSec != nullptr) { + *outDeltaSec = (devUs > 0.0) ? devUs * 1.0e-6 : -1.0; + return testSuccess; + } + + if (args->proc == 0 && args->thread == 0 && devUs > 0.0) { + double sec = devUs * 1.0e-6; + double algBw = (double)msgBytes / 1.0e9 / sec; // broadcast busbw factor = 1 + snprintf(args->devtimeAugmentLine, sizeof(args->devtimeAugmentLine), + "#[bcast-devtime] size %12zu B ctas %2d loop %2d skip %2d devtime %10.2f us algbw %8.2f GB/s busbw %8.2f GB/s\n", + msgBytes, gridCtas, loop, skip, devUs, algBw, algBw); } return testSuccess; } +#else +testResult_t BroadcastDeviceTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int in_place, double* outDeltaSec) { + return testSuccess; // device API path not available in this build +} +#endif struct testColl broadcastTest = { "Broadcast", @@ -138,7 +1358,9 @@ struct testColl broadcastTest = { BroadcastGetBw, BroadcastRunColl, BroadcastGetAlgoProtoChannels, - NULL + NULL, + NULL, + BroadcastDeviceTime }; void BroadcastGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) { @@ -186,5 +1408,11 @@ testResult_t BroadcastRunTest(struct threadArgs* args, int root, ncclDataType_t NCCL_WEAK struct testEngine ncclTestEngine = { /* .getBuffSize = */ BroadcastGetBuffSize, - /* .runTest = */ BroadcastRunTest + /* .runTest = */ BroadcastRunTest, +#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0) + /* .initCommConfig = */ nullptr, +#endif +#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0) || (defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)) + /* .getDevCommRequirements = */ BroadcastGetDevCommRequirements, +#endif }; diff --git a/projects/rccl-tests/src/gin_sdma_broadcast_policy.h b/projects/rccl-tests/src/gin_sdma_broadcast_policy.h new file mode 100644 index 000000000000..2ed7f0dd231e --- /dev/null +++ b/projects/rccl-tests/src/gin_sdma_broadcast_policy.h @@ -0,0 +1,361 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Pure (no-GPU, no-RCCL) policy logic for the GIN-SDMA hybrid Broadcast design +// (broadcast.cu, deviceImpl == 3). Single source of truth for: +// - the LSA<->GIN threshold + LL-cap env resolution, +// - the tier-selection predicates the device kernels branch on (LL / LSA / +// flat GIN / scatter+allgather / pipelined ring), +// - the scatter/ring chunk + alignment math, and +// - the devComm signal-count computation. +// +// Every function here is free of `ncclDevComm`, GPU intrinsics and getenv() so +// it can be (a) unit-tested on the host with GoogleTest and (b) called from the +// __global__ kernels so the SAME decision code runs on device. Attributes are +// guarded (GIN_SDMA_HOST_ONLY) so the header also compiles as plain host C++ in +// the test binary. +// +// This header intentionally lives in `namespace gin_sdma` next to the backend +// put-segmentation policy (nccl_device/gin/anvil_sdma/gin_anvil_sdma_put_policy.h, +// also `namespace gin_sdma`). Namespaces are open, so the two coexist as long as +// no symbol is defined twice: the shared kGinPutMaxBytes / kGinSdmaSafeCopyBytes / +// ginPutSegment* live in the put-policy header and are NOT redefined here. +// +// Every default below carries the measurement it came from in the comment above +// it (platform, date, and the numbers on both sides of the crossover), so the +// basis for each threshold is readable here rather than in an external document. + +#ifndef GIN_SDMA_BROADCAST_POLICY_H_ +#define GIN_SDMA_BROADCAST_POLICY_H_ + +#include +#include +#include +#include + +// The device kernels get __host__ __device__; plain host builds (and the host +// unit test, which defines GIN_SDMA_HOST_ONLY) get no attribute so the header +// compiles as ordinary C++ with no device codegen. Guarded with #ifndef so it +// does not clash with the identical definition in gin_anvil_sdma_put_policy.h +// when both headers are pulled into the same device TU (broadcast.cu). +#if (defined(__CUDACC__) || defined(__HIPCC__)) && !defined(GIN_SDMA_HOST_ONLY) +#ifndef GIN_SDMA_HD +#define GIN_SDMA_HD __host__ __device__ +#endif +#else +#ifndef GIN_SDMA_HD +#define GIN_SDMA_HD +#endif +#endif + +namespace gin_sdma { + +// Sentinel meaning "env var unset/empty/unparseable" (mirrors +// TEST_SDMA_THRESHOLD_UNSET in common.h so the two agree). +#ifndef GIN_SDMA_KTHRESHOLDUNSET_DEFINED +#define GIN_SDMA_KTHRESHOLDUNSET_DEFINED 1 +static constexpr size_t kThresholdUnset = (size_t)-1; +#endif + +// ---- Design constants (single source of truth; broadcast.cu aliases these) ---- + +// Broadcast LSA<->GIN default; compared against the full message bytes. +static constexpr size_t kBroadcastSdmaThresholdDefault = 262144; // 256 KiB +// Broadcast scatter+allgather (large tier) default; 0 disables the tier. +static constexpr size_t kBroadcastScatterAgMinDefault = 2ull * 1024 * 1024; // 2 MiB +// Broadcast LL fast path: compile ceiling + default runtime cutover. +static constexpr size_t kBroadcastLLMaxBytes = 65536; // 64 KiB +static constexpr size_t kBroadcastLLDefaultMaxBytes = 2048; // 2 KiB +// Broadcast pipelined-ring (large tier) minimum; 0 disables. Enabled by default +// at the SAG crossover (~32 MiB): below it the (N-1)-hop pipeline over tiny +// per-CTA slices collapses (ring 16M 98 << SAG 145), above it the ring is the +// best GIN option and beats host from 256 MiB up. Checked before SAG; SAG stays +// the opt-out fallback (NCCL_GIN_ANVIL_BCAST_RING_MIN_BYTES=0). +static constexpr size_t kBroadcastRingMinDefault = 32ull * 1024 * 1024; // 32 MiB +// Target bytes per ring pipeline chunk, sized PER CTA (see bcastRingChunks): the +// ring is an (N-1)-hop pipeline whose fill/drain efficiency is C/(C+N-1), so the +// depth must be large. A chunk-depth sweep (8x MI355X, 2G, 128 CTAs) rises +// monotonically past the old 64 cap (32ch 322, 64ch 350), so the cap is raised +// to 256 and the count targets ~64 KiB per chunk per CTA. +static constexpr size_t kBroadcastRingTargetChunk = 64ull * 1024; // 64 KiB per CTA +static constexpr int kBroadcastRingMaxChunks = 256; +static constexpr int kBroadcastRingMinChunks = 16; // fill the pipeline past the N-1 cost + +// ---------------------------- env / threshold ---------------------------- + +// Parse a size string ("123", "4K", "2M", "1G"; decimal only, optional single +// binary suffix). Returns kThresholdUnset for null/empty/negative/non-numeric +// input. +// Pure: takes the raw string so it is testable without touching the process +// environment. (common.h::testParseSdmaThresholdEnv = this on getenv(name).) +inline size_t parseSize(const char* v) { + // strtoull() accepts a leading '-' and wraps it into a huge unsigned value, + // which would read as a ~16 EiB threshold and silently pin every message to + // the LSA tier. Reject the sign before parsing. strtoull also skips leading + // whitespace before the sign, so trim it first (" -4K" must not slip through). + if (v == nullptr) return kThresholdUnset; + while (*v == ' ' || *v == '\t') v++; + if (v[0] == '\0' || v[0] == '-') return kThresholdUnset; + char* end = nullptr; + unsigned long long val = strtoull(v, &end, 10); + if (end == v) return kThresholdUnset; // no digits consumed + // strtoull saturates at ULLONG_MAX on overflow instead of failing, and the + // suffix multiply below can wrap a merely-large value back down to a small + // one. Either way the caller would get a plausible-looking threshold that + // silently mis-tiers every message, so treat both as unparseable. + const unsigned long long kMaxU64 = (unsigned long long)-1; + if (val == kMaxU64) return kThresholdUnset; + if (*end) { // trailing suffix char (end is always set by strtoull) + unsigned long long mult = 1ULL; + switch (*end) { + case 'k': case 'K': mult = 1024ULL; break; + case 'm': case 'M': mult = 1024ULL * 1024ULL; break; + case 'g': case 'G': mult = 1024ULL * 1024ULL * 1024ULL; break; + default: break; // unknown suffix ignored (matches common.h) + } + if (mult > 1ULL && val > kMaxU64 / mult) return kThresholdUnset; + val *= mult; + } + return (size_t)val; +} + +// Resolve a collective LSA<->GIN threshold from already-parsed candidates: +// 1. collective-specific value, if set; +// 2. shared NCCL_GIN_ANVIL_SDMA_THRESHOLD value, if set (global force knob); +// 3. the collective default. +GIN_SDMA_HD inline size_t resolveThreshold(size_t collVal, size_t sharedVal, + size_t collDefault) { + if (collVal != kThresholdUnset) return collVal; + if (sharedVal != kThresholdUnset) return sharedVal; + return collDefault; +} + +// Resolve an LL cutover cap from an already-parsed env value: +// unset -> unsetDefault (Broadcast: 2 KiB), +// then clamp to the compile ceiling, then round down to 8 bytes. +GIN_SDMA_HD inline size_t resolveLLCap(size_t envVal, size_t unsetDefault, + size_t maxCap) { + size_t cap = (envVal == kThresholdUnset) ? unsetDefault : envVal; + if (cap > maxCap) cap = maxCap; + return (cap / 8) * 8; +} + +// ------------------------------- Broadcast ------------------------------- + +// Host gate: use the scatter+allgather large tier instead of the flat/LSA/LL +// hybrid. Requires SAG enabled, >=2 ranks, message at/above the cutover, and at +// least one element per rank (the scatter needs a non-empty slice each). +GIN_SDMA_HD inline bool bcastUseScatterAllgather(size_t msgBytes, size_t count, + int nRanks, size_t sagMin) { + return sagMin != 0 && nRanks >= 2 && msgBytes >= sagMin && + count >= (size_t)nRanks; +} + +// Host gate: use the pipelined-ring large tier (highest priority when enabled). +// Gated by ringMin (NCCL_GIN_ANVIL_BCAST_RING_MIN_BYTES); needs >=2 ranks, +// message at/above the cutover, and count >= nRanks so every ring chunk is +// non-empty on every rank. +// +// lsaSize is devComm.lsaSize, and it must cover the whole world for the ring to +// be legal. Unlike the SAG and flat tiers, which move data with +// gin.put(ncclTeamWorld(...)) and therefore work over the network, the ring +// forwards with direct peer stores through ncclGetLsaPointer(recvwin, off, +// nextRank) and synchronizes with an LSA-team barrier. nextRank comes from the +// world ring order, so it is a WORLD rank index handed to an API that expects an +// LSA peer index. When the LSA team is smaller than the world those two indices +// stop agreeing: the store lands on the wrong peer, or out of the team entirely, +// and the LSA barrier never covers the ranks the ring spans. Requiring +// lsaSize == nRanks makes the two index spaces identical, which is the only +// configuration where the existing code is correct. +GIN_SDMA_HD inline bool bcastUseRing(size_t msgBytes, size_t count, + int nRanks, int lsaSize, size_t ringMin) { + return ringMin != 0 && nRanks >= 2 && lsaSize == nRanks && + msgBytes >= ringMin && count >= (size_t)nRanks; +} + +enum class BcastLargeTier { None, Ring, ScatterAG }; + +// Which large tier the host dispatch selects, as one decision rather than two +// independent predicates. Under the shipped defaults (ring 32 MiB, SAG 2 MiB) +// BOTH predicates are true for every message at or above the ring cutover, so the +// tier that runs is fixed by this precedence -- ring first -- and nothing else. +// Asserting the predicates one at a time is satisfied by "both true" and would +// not catch the order being swapped; asserting this enum across the 1 MiB, 2 MiB +// and 32 MiB boundaries does. +// +// The precedence is what makes lsaSize load-bearing here. Ring is checked first, +// so on a world the LSA team does not cover, an lsaSize-blind gate would hand +// every message at or above the ring cutover to a tier that cannot reach an +// off-team peer, even though SAG sitting right behind it is network-capable. +// Refusing the ring on that topology is what lets the fall-through pick SAG. +GIN_SDMA_HD inline BcastLargeTier bcastLargeTier(size_t msgBytes, size_t count, + int nRanks, int lsaSize, + size_t ringMin, size_t sagMin) { + if (bcastUseRing(msgBytes, count, nRanks, lsaSize, ringMin)) return BcastLargeTier::Ring; + if (bcastUseScatterAllgather(msgBytes, count, nRanks, sagMin)) return BcastLargeTier::ScatterAG; + return BcastLargeTier::None; +} + +// Ring pipeline depth (chunks per CTA stripe). Env override (envChunks) wins when +// set (clamped to <= max); otherwise size-adaptive: target ~64 KiB per chunk of +// THIS CTA's stripe (msgBytes/ctas), clamped to [min,max]. Per-CTA sizing keeps +// the chunk bytes constant across sizes so small messages are not over-chunked. +// The kernel further clamps to <= sCount so none is empty. +GIN_SDMA_HD inline int bcastRingChunks(size_t msgBytes, int ctas, size_t envChunks) { + if (envChunks != kThresholdUnset && envChunks > 0) { + return (envChunks > (size_t)kBroadcastRingMaxChunks) + ? kBroadcastRingMaxChunks : (int)envChunks; + } + if (ctas < 1) ctas = 1; + size_t stripeBytes = msgBytes / (size_t)ctas; + size_t c = stripeBytes / kBroadcastRingTargetChunk; + if (c < (size_t)kBroadcastRingMinChunks) c = (size_t)kBroadcastRingMinChunks; + if (c > (size_t)kBroadcastRingMaxChunks) c = (size_t)kBroadcastRingMaxChunks; + return (int)c; +} + +// In-kernel LL eligibility (inside the msgBytes<=sdmaThreshold branch): +// LL configured, 8-byte aligned, within the compile ceiling, and it fits the +// pre-sized slot count (a broadcast receiver needs msgBytes/8 u64 slots). +GIN_SDMA_HD inline bool bcastLLEligible(size_t msgBytes, int llSlots, + size_t llMaxBytes) { + return llSlots != 0 && (msgBytes % 8 == 0) && msgBytes <= llMaxBytes && + (msgBytes / 8) <= (size_t)llSlots; +} + +enum class BcastTier { LL, LSADirect, Flat }; + +// Tier chosen by GinHybridBroadcastKernel (assumes the host already decided this +// is NOT the scatter+allgather / ring tier). Uses bcastLLEligible so the kernel +// and the tests share the decision. +GIN_SDMA_HD inline BcastTier bcastKernelTier(size_t msgBytes, size_t sdmaThreshold, + int llSlots, size_t llMaxBytes) { + if (msgBytes <= sdmaThreshold) { + if (bcastLLEligible(msgBytes, llSlots, llMaxBytes)) return BcastTier::LL; + return BcastTier::LSADirect; + } + return BcastTier::Flat; +} + +// Sentinel meaning "CTA-count env var unset/empty/unparseable"; bcastHybridCtas() / +// bcastSagCtas() fall back to the size-adaptive ladder when the env value is this. +static constexpr size_t kBroadcastCtasUnset = (size_t)-1; + +// Broadcast -D 3 size-adaptive CTA count (decoupled from -V, like AllGather). Keyed +// off the SAME tier predicates the kernels evaluate: +// * LL tier (tiny, eligible): 1 CTA (kernel returns early for blockIdx.x != 0). +// * LSA-direct tier (msg <= threshold, or SAG slice <= threshold): ~16 CTAs. +// * GIN-put / Anvil-SDMA tier: few (4) CTAs -- only nRanks threads issue puts. +// NCCL_GIN_ANVIL_BCAST_CTAS pins a fixed count (diagnostic), clamped to the pool. +static constexpr int kBroadcastCtasLsa = 16; +static constexpr int kBroadcastCtasSdma = 4; +static constexpr int kBroadcastCtasLL = 1; + +GIN_SDMA_HD inline int bcastHybridMaxCtas() { + return (kBroadcastCtasLsa > kBroadcastCtasSdma) ? kBroadcastCtasLsa : kBroadcastCtasSdma; +} + +// Barrier/lsaBarrier/signal pool for the hybrid + SAG kernels: max(-V, ladder peak). +GIN_SDMA_HD inline int bcastHybridPoolCtas(int deviceCtaCount) { + const int maxTier = bcastHybridMaxCtas(); + return (deviceCtaCount > maxTier) ? deviceCtaCount : maxTier; +} + +// Full pool covering the hybrid/SAG grids AND the ring's independent CTA count. +GIN_SDMA_HD inline int bcastPoolCtas(int deviceCtaCount, int ringCtasPeak) { + int pool = bcastHybridPoolCtas(deviceCtaCount); + return (ringCtasPeak > pool) ? ringCtasPeak : pool; +} + +// Tier predicate for the flat hybrid kernel and SAG slice sizing. +GIN_SDMA_HD inline bool bcastChunkUsesLsaTier(size_t msgBytes, size_t sdmaThreshold) { + return msgBytes <= sdmaThreshold; +} + +// Launched grid for GinHybridBroadcastKernel (-D 3 flat/LL/LSA path). +GIN_SDMA_HD inline int bcastHybridCtas(size_t msgBytes, size_t sdmaThreshold, + int llSlots, size_t llMaxBytes, + size_t envCtas, int poolCtas) { + if (poolCtas < 1) poolCtas = 1; + if (envCtas != kBroadcastCtasUnset && envCtas > 0) + return (envCtas > (size_t)poolCtas) ? poolCtas : (int)envCtas; + if (msgBytes <= sdmaThreshold) { + if (bcastLLEligible(msgBytes, llSlots, llMaxBytes)) + return kBroadcastCtasLL; + const int adaptive = kBroadcastCtasLsa; + return (adaptive > poolCtas) ? poolCtas : adaptive; + } + const int adaptive = kBroadcastCtasSdma; + return (adaptive > poolCtas) ? poolCtas : adaptive; +} + +// Launched grid for GinScatterAllgatherBroadcastKernel. Keys off the per-rank +// allgather slice (msgBytes / nRanks), matching the AllGather adaptive ladder. +GIN_SDMA_HD inline int bcastSagCtas(size_t msgBytes, int nRanks, size_t sdmaThreshold, + size_t envCtas, int poolCtas) { + if (poolCtas < 1) poolCtas = 1; + if (envCtas != kBroadcastCtasUnset && envCtas > 0) + return (envCtas > (size_t)poolCtas) ? poolCtas : (int)envCtas; + size_t sliceBytes = (nRanks > 0) ? (msgBytes / (size_t)nRanks) : msgBytes; + const int adaptive = bcastChunkUsesLsaTier(sliceBytes, sdmaThreshold) + ? kBroadcastCtasLsa : kBroadcastCtasSdma; + return (adaptive > poolCtas) ? poolCtas : adaptive; +} + +// ginSignalCount for Broadcast case 3: >=2 so the SAG two-signal scheme works. +GIN_SDMA_HD inline int bcastSignalCount(int deviceCtaCount) { + return (deviceCtaCount < 2) ? 2 : deviceCtaCount; +} + +// Even split with the remainder folded into the last rank's slice (SAG / ring). +struct Chunk { + size_t count; // elements in this rank's slice + size_t eltOffset; // element offset of this rank's slice +}; +GIN_SDMA_HD inline Chunk sagChunk(size_t totalCount, int nRanks, int rank) { + Chunk c{0, 0}; + if (nRanks <= 0) return c; + size_t base = totalCount / (size_t)nRanks; + size_t tail = totalCount - base * (size_t)(nRanks - 1); + c.count = (rank == nRanks - 1) ? tail : base; + c.eltOffset = (size_t)rank * base; + return c; +} + +// ------------------------------ ring tiling ------------------------------ +// The default large tier is the SM ring, and its two-level split -- a CTA stripe +// of the buffer, then a pipeline chunk of that stripe -- used to be open-coded at +// every kernel that walks it. It lives here so the tiling that actually ships is +// the tiling the host tests assert (sagChunk, which the tests did cover, is used +// only by the scatter+allgather tier). + +// CTA stripe: block b owns [ (count*b)/nBlocks, (count*(b+1))/nBlocks ). Balanced +// to within one element and, by construction, a partition of [0,count). +GIN_SDMA_HD inline Chunk ringStripe(size_t count, int block, int nBlocks) { + Chunk c{0, 0}; + if (nBlocks <= 0 || block < 0 || block >= nBlocks) return c; + const size_t base = (count * (size_t)block) / (size_t)nBlocks; + const size_t end = (count * (size_t)(block + 1)) / (size_t)nBlocks; + c.eltOffset = base; + c.count = end - base; + return c; +} + +// Pipeline chunk c of C within a stripe: the same balanced split applied again, +// offset to the stripe base. Partitions the stripe for c in [0,C). +GIN_SDMA_HD inline Chunk ringChunk(size_t stripeBase, size_t stripeCount, int c, int C) { + Chunk k{0, 0}; + if (C <= 0 || c < 0 || c >= C) return k; + const size_t s = (stripeCount * (size_t)c) / (size_t)C; + const size_t e = (stripeCount * (size_t)(c + 1)) / (size_t)C; + k.eltOffset = stripeBase + s; + k.count = e - s; + return k; +} + +} // namespace gin_sdma + +#endif // GIN_SDMA_BROADCAST_POLICY_H_ diff --git a/projects/rccl-tests/test/test_Broadcast.py b/projects/rccl-tests/test/test_Broadcast.py index 0b66da29f96a..6ed5657663a4 100644 --- a/projects/rccl-tests/test/test_Broadcast.py +++ b/projects/rccl-tests/test/test_Broadcast.py @@ -20,85 +20,671 @@ ################################################################################ import os +import re +import glob +import shlex +import shutil +import signal import subprocess import itertools import math +import tempfile +import time import pytest ngpus = 0 -if os.environ.get('ROCR_VISIBLE_DEVICES') is not None: - ngpus = len(os.environ['ROCR_VISIBLE_DEVICES'].split(",")) -elif os.environ.get('HIP_VISIBLE_DEVICES') is not None: - ngpus = len(os.environ['HIP_VISIBLE_DEVICES'].split(",")) +if os.environ.get("ROCR_VISIBLE_DEVICES") is not None: + ngpus = len(os.environ["ROCR_VISIBLE_DEVICES"].split(",")) +elif os.environ.get("HIP_VISIBLE_DEVICES") is not None: + ngpus = len(os.environ["HIP_VISIBLE_DEVICES"].split(",")) else: - ngpus = int(subprocess.check_output("rocminfo | grep \"Device Type:.\s*.GPU\" | wc -l",shell=True)) + ngpus = int( + subprocess.check_output( + 'rocminfo | grep "Device Type:.\\s*.GPU" | wc -l', shell=True + ) + ) log_ngpus = int(math.log2(ngpus)) nthreads = ["1"] nprocs = ["2"] -ngpus_single = [str(2**x) for x in range(log_ngpus+1)] -ngpus_mpi = ["1","2"] +ngpus_single = [str(2**x) for x in range(log_ngpus + 1)] +ngpus_mpi = ["1", "2"] byte_range = [("4", "128M")] op = ["sum", "prod", "min", "max"] step_factor = ["2"] -datatype = ["int8", "uint8", "int32", "uint32", "int64", "uint64", "half", "float", "double"] -memory_type = ["coarse","fine", "host"] +datatype = [ + "int8", + "uint8", + "int32", + "uint32", + "int64", + "uint64", + "half", + "float", + "double", +] +memory_type = ["coarse", "fine", "host"] path = os.path.dirname(os.path.abspath(__file__)) executable = path + "/../build/broadcast_perf" -@pytest.mark.parametrize("nthreads, ngpus_single, byte_range, op, step_factor, datatype, memory_type", - itertools.product(nthreads, ngpus_single, byte_range, op, step_factor, datatype, memory_type)) -def test_BroadcastSingleProcess(nthreads, ngpus_single, byte_range, op, step_factor, datatype, memory_type): + +@pytest.mark.parametrize( + "nthreads, ngpus_single, byte_range, op, step_factor, datatype, memory_type", + itertools.product( + nthreads, ngpus_single, byte_range, op, step_factor, datatype, memory_type + ), +) +def test_BroadcastSingleProcess( + nthreads, ngpus_single, byte_range, op, step_factor, datatype, memory_type +): try: - args = [executable, - "-t", nthreads, - "-g", ngpus_single, - "-b", byte_range[0], - "-e", byte_range[1], - "-o", op, - "-f", step_factor, - "-d", datatype, - "-Y", memory_type] + args = [ + executable, + "-t", + nthreads, + "-g", + ngpus_single, + "-b", + byte_range[0], + "-e", + byte_range[1], + "-o", + op, + "-f", + step_factor, + "-d", + datatype, + "-Y", + memory_type, + ] if memory_type == "fine": args.insert(0, "HSA_FORCE_FINE_GRAIN_PCIE=1") args_str = " ".join(args) - rccl_test = subprocess.run(args_str, stdout=subprocess.PIPE, universal_newlines=True, shell=True) + rccl_test = subprocess.run( + args_str, stdout=subprocess.PIPE, universal_newlines=True, shell=True + ) except subprocess.CalledProcessError as err: print(rccl_test.stdout) pytest.fail("Broadcast test error(s) detected.") assert rccl_test.returncode == 0 -@pytest.mark.parametrize("nthreads, nprocs, ngpus_mpi, byte_range, op, step_factor, datatype", - itertools.product(nthreads, nprocs, ngpus_mpi, byte_range, op, step_factor, datatype)) -def test_BroadcastMPI(request, nthreads, nprocs, ngpus_mpi, byte_range, op, step_factor, datatype): + +# --------------------------------------------------------------------------- +# GIN-SDMA Broadcast multi-segment regression tests (parity with AllGather / +# AllToAll in test_AllGather.py / test_AllToAll.py). +# +# These drive the real GinHybridBroadcastKernel (deviceImpl 3, NCCL_GIN_TYPE=6) +# at message sizes that cross the 128 MiB SDMA copy clamp in the Anvil-SDMA +# backend Put. For Broadcast the root issues one gin.put() per peer with the +# full message, so -b/-e is the broadcast payload size (not total/NP as in +# AllGather). The segmented suite disables the scatter+allgather and ring large +# tiers so the flat root-fanout GIN path is exercised; the hang guard uses +# default tier selection (ring at 2 GiB). +# +# Opt-in via RCCL_TESTS_GIN_SDMA_BCAST=1 on a GIN-SDMA-capable node. Config: +# RCCL_TESTS_BCAST_NP, RCCL_TESTS_MPI_LAUNCHER, RCCL_TESTS_MPI_OPTS, +# RCCL_TESTS_BCAST_XENV, RCCL_TESTS_BCAST_EXE, RCCL_TESTS_BCAST_CTAS, +# RCCL_TESTS_BCAST_TIMEOUT_S, RCCL_TESTS_BCAST_CONN_RETRIES, +# RCCL_TESTS_BCAST_GIN_TYPE NCCL_GIN_TYPE (default: 6). The ANVIL_SDMA enum +# is 6 on develop but 5 on the NCCL 2.30.7 line, so this must be settable +# rather than baked in -- a wrong value silently loads no GIN plugin. + +MiB = 1024 * 1024 +GiB = 1024 * MiB + +_bcast_enabled = os.environ.get("RCCL_TESTS_GIN_SDMA_BCAST", "") not in ( + "", + "0", + "false", + "False", +) + + +def _env_int(name, default): + """Parse an integer env var; bad values fall back so collection stays usable.""" + raw = os.environ.get(name) + if raw is None or raw == "": + return default + try: + return int(raw) + except ValueError: + return default + + +BCAST_NP = _env_int("RCCL_TESTS_BCAST_NP", 0) or ngpus +BCAST_LAUNCHER = os.environ.get("RCCL_TESTS_MPI_LAUNCHER", "mpirun") +BCAST_CTAS = os.environ.get("RCCL_TESTS_BCAST_CTAS", "8") +BCAST_GIN_TYPE = os.environ.get("RCCL_TESTS_BCAST_GIN_TYPE", "6") +# Manual / SUT defaults. GIN CI (run-gin-ci.sh) overrides both so +# HW_CASES * retries * TIMEOUT_S stays under GIN_PYTEST_TIMEOUT; otherwise +# GNU timeout kills pytest and leaves the mpirun session (start_new_session) +# orphaned because killpg only runs in this TimeoutExpired handler. +BCAST_TIMEOUT_S = _env_int("RCCL_TESTS_BCAST_TIMEOUT_S", 900) +BCAST_CONN_RETRIES = _env_int("RCCL_TESTS_BCAST_CONN_RETRIES", 5) +BCAST_MPI_OPTS = shlex.split(os.environ.get("RCCL_TESTS_MPI_OPTS", "")) +BCAST_XENV = shlex.split(os.environ.get("RCCL_TESTS_BCAST_XENV", "")) +BCAST_EXE = os.environ.get( + "RCCL_TESTS_BCAST_EXE", os.path.join(path, "..", "build", "broadcast_perf") +) + +_bcast_skip = pytest.mark.skipif( + not _bcast_enabled, + reason="GIN-SDMA Broadcast tests are opt-in; set RCCL_TESTS_GIN_SDMA_BCAST=1 on " + "a GIN-SDMA-capable (e.g. 8x MI355X) node to enable.", +) + +_CONN_GATE_RE = re.compile( + r"LSA signal connectivity gate failed|unhandled system error", re.I +) + +# The GIN Anvil-SDMA backend has to actually bind for any of this to mean +# something: a run that falls back to another transport still exits 0 and still +# prints a full results table. NCCL_DEBUG_SUBSYS=INIT,NET below is what puts the +# bind line in the captured output. +_PLUGIN_RE = re.compile(r"gin-anvil-sdma", re.I) + +# One measured rccl-tests results row: +# size count type redop root | time algbw busbw #wrong | time algbw busbw #wrong +# out-of-place columns first, then in-place. Timings come from getFloatStr, which +# falls back to scientific notation when a value will not fit its width, so the +# numeric fields have to admit exponents. The row is not end-anchored: an algo / +# proto / nchannels group and a timestamp may follow the in-place columns, and +# concurrent rank output can splice itself onto the end of the line. +_NUM = r"[-+]?(?:\d+\.?\d*|\.\d+)(?:[eE][-+]?\d+)?" +_WRONG = r"(?:{}|N/A)".format(_NUM) +_ROW_RE = re.compile( + r"^\s*(\d+)\s+(\d+)\s+\S+\s+\S+\s+(\d+)" + r"\s+{n}\s+{n}\s+{n}\s+({w})" + r"\s+{n}\s+{n}\s+{n}\s+({w})".format(n=_NUM, w=_WRONG) +) +_OOB_RE = re.compile(r"Out of bounds values\s*:\s*(\d+)") + +# "#[bcast-tier] ", emitted once per distinct tier by BroadcastRunColl when +# RCCL_TESTS_BCAST_TIER_LOG is set. Named tiers are the contract between the -D 3 +# dispatch and the tier assertions below; see bcastReportTier in broadcast.cu. +_TIER_RE = re.compile(r"^#\[bcast-tier\]\s+(\S+)", re.M) + + +def _bcast_tiers(out): + """Set of -D 3 tier names the run reported.""" + return set(_TIER_RE.findall(out or "")) + + +def _bcast_rows(out): + """Measured rows as (size, root, wrong_out_of_place, wrong_in_place).""" + rows = [] + for line in (out or "").splitlines(): + m = _ROW_RE.match(line) + if m: + rows.append((int(m.group(1)), int(m.group(3)), m.group(4), m.group(5))) + return rows + + +def _wrong_count(field): + """Numeric #wrong, or None when the column is N/A because checks were off.""" + try: + return float(field) + except ValueError: + return None + + +def _data_failed(out): + """True when the run produced demonstrably wrong results. + + Deliberately not a text match on "Wrong": rccl-tests prints a "#wrong" column + header on every run that reaches the results table, so a bare pattern matched + every healthy run and made the connectivity retry below unreachable. N/A is + not a failure here -- it means checking was off, which _assert_bcast_ok + rejects separately rather than burning retries on. + """ + for _, _, oop, ip in _bcast_rows(out): + if any(_wrong_count(f) not in (None, 0.0) for f in (oop, ip)): + return True + m = _OOB_RE.search(out or "") + return bool(m and m.group(1) != "0") + + +def _read_debug_logs(debug_dir): + """Merged per-rank NCCL debug output written under debug_dir.""" + chunks = [] + for log_path in sorted(glob.glob(os.path.join(debug_dir, "nccl-debug.*"))): + try: + with open(log_path, errors="replace") as fh: + chunks.append(fh.read()) + except OSError: + pass + return "\n".join(chunks) + + +def _assert_bcast_ok(rc, out, debug, ctx, expect_tier=None): + """Assert the run bound GIN-SDMA, measured something, and checked clean. + + Exit status alone is not enough: it stays 0 for a run that binds a different + backend, and for one that produces no measured rows at all. + + Nor is the data check enough to say which tier ran. Every -D 3 tier produces + a correct broadcast, so #wrong stays 0 when a tier is deleted and its + messages fall through to the next one -- the test keeps passing while the + kernel it named is no longer reached. expect_tier pins the tier that has to + have run, so removing it fails here instead of silently going uncovered. + """ + tail = (out or "")[-2000:] + assert rc == 0, "Broadcast {} failed (exit {}). Output tail:\n{}".format( + ctx, rc, tail + ) + assert _PLUGIN_RE.search(debug or "") or _PLUGIN_RE.search(out or ""), ( + "Broadcast {} never bound the GIN Anvil-SDMA backend, so the run proves " + "nothing about the GIN path. Output tail:\n{}".format(ctx, tail) + ) + rows = _bcast_rows(out) + assert rows, "Broadcast {} produced no measured rows. Output tail:\n{}".format( + ctx, tail + ) + unchecked = [ + (s, r) + for s, r, oop, ip in rows + if _wrong_count(oop) is None or _wrong_count(ip) is None + ] + assert not unchecked, ( + "Broadcast {} reported #wrong as N/A at (size, root) {}, so the data check " + "never ran and the result is unverified.".format(ctx, unchecked) + ) + bad = [ + (s, r, oop, ip) + for s, r, oop, ip in rows + if _wrong_count(oop) != 0.0 or _wrong_count(ip) != 0.0 + ] + assert ( + not bad + ), "Broadcast {} reported nonzero #wrong (size, root, oop, ip): {}".format(ctx, bad) + m = _OOB_RE.search(out or "") + assert m and m.group(1) == "0", "Broadcast {} out-of-bounds count is {}".format( + ctx, m.group(1) if m else "absent" + ) + if expect_tier is not None: + # A string pins one tier; a collection allows a set of equivalent ones. + # The ring tier needs the latter: it reports ring-table where an + # edge-disjoint decomposition exists and ring-multi where none does + # (N=4 and N=6), and BCAST_NP follows whatever GPUs the host has. + # + # The ring also requires the LSA team to cover the world, so a multi-node + # launch takes scatter+allgather instead and fails here. That is the + # intent: such a run is not exercising the ring the caller asked for. + allowed = ( + {expect_tier} if isinstance(expect_tier, str) else set(expect_tier) + ) + tiers = _bcast_tiers(out) + assert tiers, ( + "Broadcast {} reported no #[bcast-tier] line, so which tier ran is " + "unverified. Is RCCL_TESTS_BCAST_TIER_LOG set for this launch, and is " + "the binary new enough to emit it? Output tail:\n{}".format(ctx, tail) + ) + # Exactly one tier, and an allowed one. These launches pin a single size, + # so a second tier means some roots took a different kernel; merely + # intersecting `allowed` would let that through. + assert len(tiers) == 1 and tiers <= allowed, ( + "Broadcast {} ran tier(s) {} but expected exactly one of {}. The tier " + "gate moved, or the tier under test fell through to another " + "kernel.".format(ctx, sorted(tiers), sorted(allowed)) + ) + + +def _launch_bcast_gin_sdma( + request, msg_bytes, dtype, *, force_flat_gin=False, force_sag_gin=False +): + """Launch broadcast_perf -D 3 at a fixed message size. When force_flat_gin, + disable the SAG/ring large tiers so the root flat gin.put() path is used. + When force_sag_gin, disable only the ring tier so scatter+allgather runs. + + Returns (returncode, stdout, debug_text); debug_text is the merged per-rank + NCCL debug log, kept off stdout so the results table stays parseable.""" + size = str(int(msg_bytes)) + debug_dir = tempfile.mkdtemp(prefix="bcast-gin-dbg-") + gin_env = [ + "NCCL_GIN_ENABLE=1", + "NCCL_GIN_TYPE={}".format(BCAST_GIN_TYPE), + "NCCL_GIN_ANVIL_SDMA_THRESHOLD=0", + "NCCL_GIN_ANVIL_SDMA_THRESHOLD_BROADCAST=0", + # Debug output is what lets _PLUGIN_RE confirm which backend actually + # bound. It goes to per-rank files rather than stdout because ranks + # interleave: on stdout an INFO line splices itself into the middle of a + # results row, which corrupts the very table we need to parse. + "NCCL_DEBUG=INFO", + "NCCL_DEBUG_SUBSYS=INIT,NET", + "NCCL_DEBUG_FILE={}/nccl-debug.%h.%p.log".format(debug_dir), + # Makes BroadcastRunColl name the tier it launched, which is what lets + # _assert_bcast_ok tell the -D 3 tiers apart. One line from the main + # proc, so the results table stays parseable. + "RCCL_TESTS_BCAST_TIER_LOG=1", + ] + if force_flat_gin: + gin_env += [ + "NCCL_GIN_ANVIL_BCAST_SCATTER_AG_MIN_BYTES=0", + "NCCL_GIN_ANVIL_BCAST_RING_MIN_BYTES=0", + ] + elif force_sag_gin: + # Ring is checked before SAG; disable ring only so SAG is exercised. + gin_env += ["NCCL_GIN_ANVIL_BCAST_RING_MIN_BYTES=0"] + gin_env += BCAST_XENV + gin_env = sum([["-x", kv] for kv in gin_env], []) + + hostfile = request.config.getoption("--hostfile") + launch = [BCAST_LAUNCHER, "-np", str(BCAST_NP)] + BCAST_MPI_OPTS + if hostfile: + launch += ["-host", hostfile] + + args = ( + launch + + gin_env + + [ + BCAST_EXE, + "-b", + size, + "-e", + size, + "-f", + "2", + "-g", + "1", + "-R", + "2", + "-D", + "3", # GinHybridBroadcastKernel + "-A", + "1", + "-V", + BCAST_CTAS, + "-d", + dtype, + "-c", + "1", + "-w", + "1", + "-n", + "3", + ] + ) + cmd = " ".join(shlex.quote(a) for a in args) + print(cmd) + try: + proc = subprocess.Popen( + cmd, + shell=True, + universal_newlines=True, + stdout=subprocess.PIPE, + stderr=subprocess.STDOUT, + start_new_session=True, + ) + try: + out, _ = proc.communicate(timeout=BCAST_TIMEOUT_S) + except subprocess.TimeoutExpired: + try: + os.killpg(os.getpgid(proc.pid), signal.SIGKILL) + except ProcessLookupError: + pass + out, _ = proc.communicate() + pytest.fail( + "Broadcast GIN-SDMA HANG: no completion within {}s at msg={} bytes " + "({} MiB), dtype={}. Output tail:\n{}".format( + BCAST_TIMEOUT_S, size, msg_bytes // MiB, dtype, (out or "")[-2000:] + ) + ) + debug = _read_debug_logs(debug_dir) + finally: + shutil.rmtree(debug_dir, ignore_errors=True) + print(out) + return proc.returncode, out, debug + + +def _run_bcast_gin_sdma( + request, msg_bytes, dtype, *, force_flat_gin=False, force_sag_gin=False +): + if BCAST_NP < 2: + pytest.skip("need >= 2 ranks/GPUs for GIN-SDMA Broadcast") + + rc, out, debug = 1, "", "" + for attempt in range(1, max(1, BCAST_CONN_RETRIES) + 1): + rc, out, debug = _launch_bcast_gin_sdma( + request, + msg_bytes, + dtype, + force_flat_gin=force_flat_gin, + force_sag_gin=force_sag_gin, + ) + if rc == 0: + return rc, out, debug + if _data_failed(out): + return rc, out, debug + if _CONN_GATE_RE.search(out or "") and attempt < BCAST_CONN_RETRIES: + print( + "=== connectivity-gate abort (attempt {}/{}); re-launching after settle ===".format( + attempt, BCAST_CONN_RETRIES + ) + ) + time.sleep(3) + continue + return rc, out, debug + return rc, out, debug + + +@_bcast_skip +@pytest.mark.parametrize("msg_mib", [256, 2048]) # 256 MiB (2 seg), 2 GiB (16 seg) +@pytest.mark.parametrize("dtype", ["int32", "int64", "uint8"]) +def test_BroadcastGinSdmaLargeSegmented(request, msg_mib, dtype): + """Flat root-fanout gin.put() at sizes crossing the 128 MiB SDMA segment.""" + rc, out, debug = _run_bcast_gin_sdma( + request, msg_mib * MiB, dtype, force_flat_gin=True + ) + _assert_bcast_ok( + rc, + out, + debug, + "flat-segmented {} MiB dtype={}".format(msg_mib, dtype), + expect_tier="flat-gin", + ) + + +@_bcast_skip +def test_BroadcastGinSdmaScatterAllgather(request): + """256 MiB scatter+allgather tier (ring disabled via RING_MIN_BYTES=0).""" + rc, out, debug = _run_bcast_gin_sdma( + request, 256 * MiB, "int32", force_sag_gin=True + ) + _assert_bcast_ok( + rc, out, debug, "SAG 256 MiB", expect_tier="scatter-allgather" + ) + + +@_bcast_skip +def test_BroadcastGinSdma2GiBHangGuard(request): + """2 GiB completion guard with default tier selection (ring path).""" + rc, out, debug = _run_bcast_gin_sdma( + request, 2 * GiB, "int32", force_flat_gin=False + ) + _assert_bcast_ok( + rc, + out, + debug, + "2 GiB default-tier", + expect_tier=("ring-table", "ring-multi"), + ) + + +@_bcast_skip +def test_BroadcastGinSdma4GiBHangGuard(request): + """4 GiB completion guard with default tier selection (ring path at 4 GiB).""" + rc, out, debug = _run_bcast_gin_sdma( + request, 4 * GiB, "int32", force_flat_gin=False + ) + _assert_bcast_ok( + rc, + out, + debug, + "4 GiB default-tier", + expect_tier=("ring-table", "ring-multi"), + ) + + +# Offline parsing guards: no GIN hardware required. These pin the regression where +# _DATA_FAIL_RE matched the "#wrong" column header and made the connectivity retry +# unreachable, and where scientific-notation timings broke the row regex. +# Named into `-k GinSdma` (rccl-gin-bcast-pytest) because run-gin-ci.sh expands +# pytest args unquoted, so a filter with a space in it is not available. +def test_BroadcastGinSdmaRowRegexAcceptsScientificNotation(): + line = ( + " 134217728 134217728 int32 none 0" + " 1.23e+02 1.00e+02 1.00e+02 0" + " 1.23e+02 1.00e+02 1.00e+02 0" + ) + rows = _bcast_rows(line) + assert rows == [(134217728, 0, "0", "0")] + + +def test_BroadcastGinSdmaDataFailedIgnoresColumnHeader(): + header = ( + "# size count type redop root" + " time algbw busbw #wrong" + ) + assert not _data_failed(header) + + +def test_BroadcastGinSdmaDataFailedDetectsNonzeroWrong(): + line = ( + " 1048576 1048576 int32 none 0" + " 12.34 1.00 1.00 1" + " 12.34 1.00 1.00 0" + ) + assert _data_failed(line) + + +def test_BroadcastGinSdmaDataFailedTreatsNaAsUncheckedNotFailed(): + line = ( + " 1048576 1048576 int32 none 0" + " 12.34 1.00 1.00 N/A" + " 12.34 1.00 1.00 0" + ) + assert not _data_failed(line) + + +def _clean_bcast_output(tier): + """A minimal healthy -D 3 run reporting `tier`, for the guards below.""" + return "\n".join( + [ + "#[bcast-tier] {}".format(tier), + " 268435456 268435456 int32 none 0" + " 12.34 1.00 1.00 0" + " 12.34 1.00 1.00 0", + "# Out of bounds values : 0 OK", + ] + ) + + +# These four need no GPU, but they carry GinSdma in the name on purpose: the CI +# job runs `-k GinSdma` (rccl-gin-bcast-pytest in gin-tests.json), and the runner +# expands its args unquoted, so a filter with a space in it is not available. +def test_BroadcastGinSdmaTierGuardLineIsNotAResultsRow(): + """The tier line shares stdout with the table it must not corrupt.""" + assert _bcast_rows("#[bcast-tier] scatter-allgather") == [] + assert _bcast_tiers(_clean_bcast_output("ring-table")) == {"ring-table"} + assert _bcast_rows(_clean_bcast_output("ring-table")) + + +def test_BroadcastGinSdmaTierGuardAcceptsExpectedTier(): + out = _clean_bcast_output("scatter-allgather") + _assert_bcast_ok(0, out, "gin-anvil-sdma", "guard", expect_tier="scatter-allgather") + # Ring passes against either decomposition outcome. + _assert_bcast_ok( + 0, + _clean_bcast_output("ring-multi"), + "gin-anvil-sdma", + "guard", + expect_tier=("ring-table", "ring-multi"), + ) + + +def test_BroadcastGinSdmaTierGuardRejectsFallThrough(): + """The regression this tier check exists for. + + Deleting the scatter+allgather arm sends its messages to the flat kernel, + which still broadcasts correctly and still reports #wrong 0. Everything + except the tier check passes on that output, so without this the SAG test + would go green while covering nothing. + """ + fell_through = _clean_bcast_output("flat-gin") + _assert_bcast_ok(0, fell_through, "gin-anvil-sdma", "guard", expect_tier="flat-gin") + with pytest.raises(AssertionError, match="expected exactly one of"): + _assert_bcast_ok( + 0, fell_through, "gin-anvil-sdma", "guard", expect_tier="scatter-allgather" + ) + + +def test_BroadcastGinSdmaTierGuardRejectsMissingTierLine(): + """A binary predating the tier log must fail loudly, not pass vacuously.""" + no_tier = "\n".join(_clean_bcast_output("flat-gin").splitlines()[1:]) + with pytest.raises(AssertionError, match="no .*bcast-tier.* line"): + _assert_bcast_ok( + 0, no_tier, "gin-anvil-sdma", "guard", expect_tier="flat-gin" + ) + + +@pytest.mark.parametrize( + "nthreads, nprocs, ngpus_mpi, byte_range, op, step_factor, datatype", + itertools.product( + nthreads, nprocs, ngpus_mpi, byte_range, op, step_factor, datatype + ), +) +def test_BroadcastMPI( + request, nthreads, nprocs, ngpus_mpi, byte_range, op, step_factor, datatype +): try: - mpi_hostfile = request.config.getoption('--hostfile') + mpi_hostfile = request.config.getoption("--hostfile") if not mpi_hostfile: - args = ["mpirun -np", nprocs, - executable, - "-p 1", - "-t", nthreads, - "-g", ngpus_mpi, - "-b", byte_range[0], - "-e", byte_range[1], - "-o", op, - "-f", step_factor, - "-d", datatype] + args = [ + "mpirun -np", + nprocs, + executable, + "-p 1", + "-t", + nthreads, + "-g", + ngpus_mpi, + "-b", + byte_range[0], + "-e", + byte_range[1], + "-o", + op, + "-f", + step_factor, + "-d", + datatype, + ] else: - args = ["mpirun -np", nprocs, - "-host", mpi_hostfile, - executable, - "-p 1", - "-t", nthreads, - "-g", ngpus_mpi, - "-b", byte_range[0], - "-e", byte_range[1], - "-o", op, - "-f", step_factor, - "-d", datatype] + args = [ + "mpirun -np", + nprocs, + "-host", + mpi_hostfile, + executable, + "-p 1", + "-t", + nthreads, + "-g", + ngpus_mpi, + "-b", + byte_range[0], + "-e", + byte_range[1], + "-o", + op, + "-f", + step_factor, + "-d", + datatype, + ] args_str = " ".join(args) print(args_str) rccl_test = subprocess.run(args_str, universal_newlines=True, shell=True) diff --git a/projects/rccl/test/CMakeLists.txt b/projects/rccl/test/CMakeLists.txt index 606930b53a71..aa748ae4fc88 100644 --- a/projects/rccl/test/CMakeLists.txt +++ b/projects/rccl/test/CMakeLists.txt @@ -558,6 +558,26 @@ if(BUILD_TESTS) set_tests_properties(rccl-UnitTestsGinSdmaAllGatherGpu PROPERTIES LABELS "unit;gpu") endif() + # GIN Anvil-SDMA hybrid Broadcast (-D 3) policy host unit test (no GPU). + # Compiles gin_sdma_broadcast_policy.h from the sibling rccl-tests project with + # GIN_SDMA_HOST_ONLY (no librccl / device codegen needed), validating the exact + # LSA<->GIN threshold + LL-cap resolution, tier-selection predicates (LL / LSA / + # flat / scatter+allgather / ring), scatter/ring chunk math and signal-count + # sizing that broadcast.cu relies on. Skip the target if the rccl-tests checkout + # is absent so a standalone rccl build still configures. Runs under `ctest -L unit`. + if (EXISTS "${_gin_sdma_ag_hdr_dir}/gin_sdma_broadcast_policy.h") + add_executable(rccl-UnitTestsGinSdmaBroadcastPolicy gin/gin_sdma_broadcast_policy_test.cpp) + set_source_files_properties(gin/gin_sdma_broadcast_policy_test.cpp PROPERTIES LANGUAGE CXX) + target_include_directories(rccl-UnitTestsGinSdmaBroadcastPolicy PRIVATE + "${_gin_sdma_ag_hdr_dir}") + target_compile_definitions(rccl-UnitTestsGinSdmaBroadcastPolicy PRIVATE GIN_SDMA_HOST_ONLY) + target_link_libraries(rccl-UnitTestsGinSdmaBroadcastPolicy PRIVATE ${GTEST_BOTH_LIBRARIES} Threads::Threads) + set_target_properties(rccl-UnitTestsGinSdmaBroadcastPolicy PROPERTIES + RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/test") + add_test(NAME rccl-UnitTestsGinSdmaBroadcastPolicy COMMAND rccl-UnitTestsGinSdmaBroadcastPolicy) + set_tests_properties(rccl-UnitTestsGinSdmaBroadcastPolicy PROPERTIES LABELS "unit") + endif() + # rccl-UnitTestsFixturesDebug: Tests that access internal symbols compiled into # librccl.so which are only visible in Debug builds (hidden visibility in Release). if (ROCM_VERSION VERSION_GREATER_EQUAL "60400" AND CMAKE_BUILD_TYPE MATCHES "Debug") diff --git a/projects/rccl/test/gin/gin_sdma_broadcast_policy_test.cpp b/projects/rccl/test/gin/gin_sdma_broadcast_policy_test.cpp new file mode 100644 index 000000000000..0eeb6aa9ad52 --- /dev/null +++ b/projects/rccl/test/gin/gin_sdma_broadcast_policy_test.cpp @@ -0,0 +1,537 @@ +/************************************************************************* + * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved. + * + * See LICENSE.txt for license information + ************************************************************************/ + +// Host-only unit tests for the pure GIN Anvil-SDMA hybrid Broadcast policy +// helpers (gin_sdma_broadcast_policy.h). No GPU required: GIN_SDMA_HOST_ONLY +// drops the HIP attributes so the header compiles as plain host C++, and the +// same functions are called by broadcast.cu's __global__ kernels on device, so +// exercising them here exercises the real decision code. These cover the +// threshold/env resolution, LL/LSA/flat/scatter+allgather/ring tier selection, +// scatter/ring chunk + alignment math, and devComm signal-count decisions to +// high line+branch coverage. GPU data-movement correctness (the branch bodies) +// is covered separately by the broadcast_perf functional sweep. + +#include + +#include +#include + +#define GIN_SDMA_HOST_ONLY 1 +#include "gin_sdma_broadcast_policy.h" + +using namespace gin_sdma; + +namespace { + +constexpr size_t kUnset = kThresholdUnset; + +// --------------------------------- parseSize --------------------------------- + +TEST(ParseSize, NullAndEmptyAreUnset) { + EXPECT_EQ(parseSize(nullptr), kUnset); + EXPECT_EQ(parseSize(""), kUnset); +} + +TEST(ParseSize, NonNumericIsUnset) { + EXPECT_EQ(parseSize("abc"), kUnset); // no digits consumed + EXPECT_EQ(parseSize("K"), kUnset); // suffix only, no digits +} + +TEST(ParseSize, NegativeIsUnset) { + // Without an explicit sign check strtoull() wraps these into near-SIZE_MAX + // thresholds, which read as "always LSA" instead of "unset". + EXPECT_EQ(parseSize("-1"), kUnset); + EXPECT_EQ(parseSize("-4K"), kUnset); + EXPECT_EQ(parseSize("-262144"), kUnset); + EXPECT_EQ(parseSize(" -4K"), kUnset); +} + +TEST(ParseSize, PlainDecimal) { + EXPECT_EQ(parseSize("0"), 0u); + EXPECT_EQ(parseSize("1024"), 1024u); + EXPECT_EQ(parseSize("262144"), 262144u); +} + +TEST(ParseSize, BinarySuffixesBothCases) { + EXPECT_EQ(parseSize("4K"), 4u * 1024); + EXPECT_EQ(parseSize("4k"), 4u * 1024); + EXPECT_EQ(parseSize("2M"), 2u * 1024 * 1024); + EXPECT_EQ(parseSize("2m"), 2u * 1024 * 1024); + EXPECT_EQ(parseSize("1G"), 1024u * 1024 * 1024); + EXPECT_EQ(parseSize("1g"), 1024u * 1024 * 1024); +} + +TEST(ParseSize, UnknownSuffixIgnored) { + EXPECT_EQ(parseSize("512B"), 512u); + EXPECT_EQ(parseSize("10X"), 10u); +} + +TEST(ParseSize, OverflowIsUnsetNotASmallThreshold) { + // strtoull saturates at ULLONG_MAX rather than failing, and the suffix multiply + // can then wrap a huge value back down to a plausible-looking small threshold -- + // the same silent mis-tiering shape as a negative value. Both must read as unset + // so the caller falls back to the tested default. + EXPECT_EQ(parseSize("99999999999999999999"), kUnset); + EXPECT_EQ(parseSize("99999999999999999999G"), kUnset); + // Not saturated on its own, but wraps when scaled by the suffix. + EXPECT_EQ(parseSize("17179869184G"), kUnset); // 2^34 GiB = 2^64 bytes + // The largest value that still scales cleanly is accepted. + EXPECT_EQ(parseSize("1024G"), 1024ull * 1024 * 1024 * 1024); +} + +TEST(ParseSize, HexIsNotAcceptedAsHex) { + // Base 10 by contract: "0x10" consumes the leading 0 and stops at 'x', which + // yields 0 -- and 0 is meaningful (it force-enables GIN / disables a tier), so + // this documents that a hex-looking value is NOT silently read as 16. + EXPECT_EQ(parseSize("0x10"), 0u); + EXPECT_NE(parseSize("0x10"), 16u); +} + +// ------------------------------ resolveThreshold ----------------------------- + +TEST(ResolveThreshold, CollectiveValueWins) { + EXPECT_EQ(resolveThreshold(4096, 8192, 262144), 4096u); +} + +TEST(ResolveThreshold, SharedUsedWhenCollUnset) { + EXPECT_EQ(resolveThreshold(kUnset, 8192, 262144), 8192u); +} + +TEST(ResolveThreshold, DefaultWhenBothUnset) { + EXPECT_EQ(resolveThreshold(kUnset, kUnset, 262144), 262144u); +} + +TEST(ResolveThreshold, ExplicitZeroIsHonored) { + // Zero is a valid "force GIN/SDMA always" value, distinct from unset. + EXPECT_EQ(resolveThreshold(0, kUnset, 262144), 0u); + EXPECT_EQ(resolveThreshold(kUnset, 0, 262144), 0u); +} + +// -------------------------------- resolveLLCap ------------------------------- + +TEST(ResolveLLCap, UnsetUsesDefaultThenAligns) { + EXPECT_EQ(resolveLLCap(kUnset, kBroadcastLLDefaultMaxBytes, kBroadcastLLMaxBytes), + 2048u); +} + +TEST(ResolveLLCap, ClampsToMax) { + EXPECT_EQ(resolveLLCap(1u << 20, kBroadcastLLDefaultMaxBytes, kBroadcastLLMaxBytes), + kBroadcastLLMaxBytes); +} + +TEST(ResolveLLCap, RoundsDownToEightBytes) { + EXPECT_EQ(resolveLLCap(1001, 0, 65536), 1000u); // 1001 -> 1000 + EXPECT_EQ(resolveLLCap(1000, 0, 65536), 1000u); // already aligned + EXPECT_EQ(resolveLLCap(7, 0, 65536), 0u); // below one slot +} + +// -------------------------- bcastUseScatterAllgather ------------------------- + +TEST(BcastScatterAllgather, DisabledWhenSagMinZero) { + EXPECT_FALSE(bcastUseScatterAllgather(8u << 20, 1u << 20, 8, 0)); +} + +TEST(BcastScatterAllgather, RequiresAtLeastTwoRanks) { + EXPECT_FALSE(bcastUseScatterAllgather(8u << 20, 1u << 20, 1, 2u << 20)); +} + +TEST(BcastScatterAllgather, RequiresMessageAtOrAboveMin) { + const size_t sagMin = 2u << 20; + EXPECT_FALSE(bcastUseScatterAllgather(sagMin - 1, 1u << 20, 8, sagMin)); + EXPECT_TRUE(bcastUseScatterAllgather(sagMin, 1u << 20, 8, sagMin)); +} + +TEST(BcastScatterAllgather, RequiresCountAtLeastNRanks) { + EXPECT_FALSE(bcastUseScatterAllgather(8u << 20, 7, 8, 1024)); + EXPECT_TRUE(bcastUseScatterAllgather(8u << 20, 8, 8, 1024)); +} + +// -------------------------------- bcastUseRing ------------------------------- + +TEST(BcastUseRing, DisabledWhenRingMinZero) { + EXPECT_FALSE(bcastUseRing(64u << 20, 16u << 20, 8, 8, 0)); +} + +TEST(BcastUseRing, RequiresAtLeastTwoRanks) { + EXPECT_FALSE(bcastUseRing(64u << 20, 16u << 20, 1, 1, kBroadcastRingMinDefault)); +} + +TEST(BcastUseRing, EngagesAtOrAboveMin) { + const size_t ringMin = kBroadcastRingMinDefault; // 32 MiB + EXPECT_FALSE(bcastUseRing(ringMin - 1, 16u << 20, 8, 8, ringMin)); + EXPECT_TRUE(bcastUseRing(ringMin, 16u << 20, 8, 8, ringMin)); +} + +TEST(BcastUseRing, RequiresCountAtLeastNRanks) { + EXPECT_FALSE(bcastUseRing(64u << 20, 7, 8, 8, kBroadcastRingMinDefault)); + EXPECT_TRUE(bcastUseRing(64u << 20, 8, 8, 8, kBroadcastRingMinDefault)); +} + +// The ring forwards with direct peer stores addressed by WORLD rank index +// (ncclGetLsaPointer(recvwin, off, nextRank)) and syncs on an LSA-team barrier, +// so it is only correct when the LSA team is the whole world. Anything smaller +// -- multi-node being the ordinary case -- means the index handed to the LSA API +// is not an LSA peer index at all. +TEST(BcastUseRing, RequiresLsaTeamToCoverTheWorld) { + const size_t ringMin = kBroadcastRingMinDefault; + const size_t bytes = 64u << 20; + const size_t count = 16u << 20; + // Single node: LSA team == world, so world and LSA indices coincide. + EXPECT_TRUE(bcastUseRing(bytes, count, 8, 8, ringMin)); + // Two nodes of 8: every size/shape condition still holds, and the ring is + // still refused, because nextRank would address outside the LSA team. + EXPECT_FALSE(bcastUseRing(bytes, count, 16, 8, ringMin)); + // A single off-team rank is enough to break the index correspondence. + EXPECT_FALSE(bcastUseRing(bytes, count, 9, 8, ringMin)); + // An lsaSize larger than the world is not a topology we know how to reason + // about; require exact agreement rather than lsaSize >= nRanks. + EXPECT_FALSE(bcastUseRing(bytes, count, 8, 16, ringMin)); +} + +// ------------------------------ bcastLargeTier ------------------------------- +// The predicates above are each true in isolation for large messages, so testing +// them one at a time cannot pin down which tier actually runs. These assert the +// resolved decision across the three boundaries that separate the tiers. + +TEST(BcastLargeTier, RingWinsWhereBothPredicatesHold) { + const size_t ringMin = kBroadcastRingMinDefault; // 32 MiB + const size_t sagMin = kBroadcastScatterAgMinDefault; // 2 MiB + const size_t bytes = 64u << 20; + const size_t count = bytes / 4; + // Both gates are open here; precedence -- not the predicates -- picks the ring. + ASSERT_TRUE(bcastUseRing(bytes, count, 8, 8, ringMin)); + ASSERT_TRUE(bcastUseScatterAllgather(bytes, count, 8, sagMin)); + EXPECT_EQ(bcastLargeTier(bytes, count, 8, 8, ringMin, sagMin), + BcastLargeTier::Ring); +} + +TEST(BcastLargeTier, BoundariesSelectEachTier) { + const size_t ringMin = kBroadcastRingMinDefault; // 32 MiB + const size_t sagMin = kBroadcastScatterAgMinDefault; // 2 MiB + auto tier = [&](size_t bytes) { + return bcastLargeTier(bytes, bytes / 4, 8, 8, ringMin, sagMin); + }; + // Below the SAG cutover: neither large tier, so the flat/LSA/LL kernel runs. + EXPECT_EQ(tier(1u << 20), BcastLargeTier::None); + // At the SAG cutover but below the ring cutover. + EXPECT_EQ(tier(sagMin), BcastLargeTier::ScatterAG); + EXPECT_EQ(tier(ringMin - 1), BcastLargeTier::ScatterAG); + // At and above the ring cutover. + EXPECT_EQ(tier(ringMin), BcastLargeTier::Ring); + EXPECT_EQ(tier(4ull << 30), BcastLargeTier::Ring); +} + +TEST(BcastLargeTier, RingOptOutFallsBackToScatterAllgather) { + // NCCL_GIN_ANVIL_BCAST_RING_MIN_BYTES=0 disables the ring; SAG must take over + // rather than the message silently dropping to the flat tier. + const size_t sagMin = kBroadcastScatterAgMinDefault; + EXPECT_EQ(bcastLargeTier(64u << 20, 16u << 20, 8, 8, 0, sagMin), + BcastLargeTier::ScatterAG); + // Both disabled -> flat. + EXPECT_EQ(bcastLargeTier(64u << 20, 16u << 20, 8, 8, 0, 0), + BcastLargeTier::None); +} + +// Ring is checked first, so without the LSA condition every message at or above +// the ring cutover would take a tier that cannot reach an off-team peer. SAG +// moves data with gin.put over the world team, so it is the correct fallback on +// a world the LSA team does not cover, and it sits directly behind the ring. +TEST(BcastLargeTier, MultiNodeFallsBackToScatterAllgatherNotRing) { + const size_t ringMin = kBroadcastRingMinDefault; // 32 MiB + const size_t sagMin = kBroadcastScatterAgMinDefault; // 2 MiB + const size_t count = 16u << 20; + auto tier = [&](size_t bytes, int nRanks, int lsaSize) { + return bcastLargeTier(bytes, count, nRanks, lsaSize, ringMin, sagMin); + }; + // 2x8 ranks, well above the ring cutover: SAG, not Ring. + EXPECT_EQ(tier(64u << 20, 16, 8), BcastLargeTier::ScatterAG); + EXPECT_EQ(tier(4ull << 30, 16, 8), BcastLargeTier::ScatterAG); + // The same sizes on one node still choose the ring, so the gate costs the + // single-node path nothing. + EXPECT_EQ(tier(64u << 20, 8, 8), BcastLargeTier::Ring); + EXPECT_EQ(tier(4ull << 30, 8, 8), BcastLargeTier::Ring); + // Below the ring cutover the LSA team is irrelevant: SAG either way. + EXPECT_EQ(tier(sagMin, 16, 8), BcastLargeTier::ScatterAG); + EXPECT_EQ(tier(sagMin, 8, 8), BcastLargeTier::ScatterAG); +} + +// With the ring disabled by env AND an LSA team that does not cover the world, +// the outcome must still be SAG rather than a silent drop to the flat tier. +TEST(BcastLargeTier, MultiNodeWithRingDisabledStillSelectsScatterAllgather) { + EXPECT_EQ(bcastLargeTier(64u << 20, 16u << 20, 16, 8, 0, + kBroadcastScatterAgMinDefault), + BcastLargeTier::ScatterAG); +} + +// ------------------------------- bcastRingChunks ----------------------------- + +TEST(BcastRingChunks, EnvOverrideWinsClampedToMax) { + // Env pin honored, clamped to the compile ceiling. + EXPECT_EQ(bcastRingChunks(2ull << 30, 128, 32), 32); + EXPECT_EQ(bcastRingChunks(2ull << 30, 128, 100000), + kBroadcastRingMaxChunks); +} + +TEST(BcastRingChunks, SizeAdaptivePerCtaClampedToRange) { + // Small stripe -> floored at the minimum depth. + EXPECT_EQ(bcastRingChunks(1u << 20, 128, kUnset), kBroadcastRingMinChunks); + // Very large message over few CTAs -> capped at the maximum depth. + EXPECT_EQ(bcastRingChunks(2ull << 30, 1, kUnset), kBroadcastRingMaxChunks); + // Interior case: 64 MiB / 8 CTAs = 8 MiB stripe; 8 MiB / 64 KiB = 128 chunks. + EXPECT_EQ(bcastRingChunks(64ull << 20, 8, kUnset), 128); + // Mid message at max depth: 128 MiB / 8 CTAs = 16 MiB stripe -> 256 chunks. + EXPECT_EQ(bcastRingChunks(128u << 20, 8, kUnset), kBroadcastRingMaxChunks); +} + +TEST(BcastRingChunks, GuardsZeroCtas) { + // ctas < 1 is clamped to 1 internally (no divide-by-zero). + EXPECT_GE(bcastRingChunks(2ull << 30, 0, kUnset), kBroadcastRingMinChunks); + EXPECT_LE(bcastRingChunks(2ull << 30, 0, kUnset), kBroadcastRingMaxChunks); +} + +// ------------------------------ bcast LL / tier ------------------------------ + +TEST(BcastLLEligible, AllGates) { + const size_t cap = kBroadcastLLMaxBytes; + EXPECT_FALSE(bcastLLEligible(64, 0, cap)); // no slots + EXPECT_FALSE(bcastLLEligible(60, 1000, cap)); // not 8-aligned + EXPECT_FALSE(bcastLLEligible(cap + 8, 100000, cap)); // above ceiling + EXPECT_FALSE(bcastLLEligible(800, 99, cap)); // 800/8=100 slots > 99 + EXPECT_TRUE(bcastLLEligible(800, 100, cap)); // exactly fits + EXPECT_TRUE(bcastLLEligible(8, 1, cap)); +} + +TEST(BcastKernelTier, LadderSelection) { + const size_t thr = 262144; + const size_t cap = kBroadcastLLMaxBytes; + EXPECT_EQ(bcastKernelTier(thr + 1, thr, 100000, cap), BcastTier::Flat); + EXPECT_EQ(bcastKernelTier(800, thr, 100000, cap), BcastTier::LL); + // <=thr but LL not configured -> LSA direct. + EXPECT_EQ(bcastKernelTier(800, thr, 0, cap), BcastTier::LSADirect); + // <=thr, LL configured but message above LL ceiling -> LSA direct. + EXPECT_EQ(bcastKernelTier(cap + 8, thr, 1u << 20, cap), BcastTier::LSADirect); +} + +TEST(BcastSignalCount, FloorOfTwo) { + EXPECT_EQ(bcastSignalCount(0), 2); + EXPECT_EQ(bcastSignalCount(1), 2); + EXPECT_EQ(bcastSignalCount(2), 2); + EXPECT_EQ(bcastSignalCount(8), 8); +} + +// ---- bcastHybridCtas / bcastSagCtas: size-adaptive ladder, clamped to pool ---- + +TEST(BcastPolicyCtas, HybridLsaTierUsesLsaCtaCount) { + const int pool = bcastHybridPoolCtas(16); + const size_t thr = kBroadcastSdmaThresholdDefault; + EXPECT_EQ(bcastHybridCtas(thr, thr, 100000, kBroadcastLLMaxBytes, kBroadcastCtasUnset, pool), + kBroadcastCtasLsa); + EXPECT_EQ(bcastHybridCtas(1024, thr, 0, kBroadcastLLMaxBytes, kBroadcastCtasUnset, pool), + kBroadcastCtasLsa); +} + +TEST(BcastPolicyCtas, HybridLLTierUsesOneCta) { + const int pool = bcastHybridPoolCtas(16); + const size_t thr = kBroadcastSdmaThresholdDefault; + EXPECT_EQ(bcastHybridCtas(800, thr, 100, kBroadcastLLMaxBytes, kBroadcastCtasUnset, pool), + kBroadcastCtasLL); +} + +TEST(BcastPolicyCtas, HybridSdmaTierUsesSdmaCtaCount) { + const int pool = bcastHybridPoolCtas(16); + const size_t thr = kBroadcastSdmaThresholdDefault; + EXPECT_EQ(bcastHybridCtas(thr + 1, thr, 0, kBroadcastLLMaxBytes, kBroadcastCtasUnset, pool), + kBroadcastCtasSdma); + EXPECT_EQ(bcastHybridCtas(64ull << 20, thr, 0, kBroadcastLLMaxBytes, kBroadcastCtasUnset, pool), + kBroadcastCtasSdma); +} + +TEST(BcastPolicyCtas, SagKeysOffPerRankSlice) { + const int pool = bcastHybridPoolCtas(16); + const size_t thr = kBroadcastSdmaThresholdDefault; + // 128 MiB / 8 ranks = 16 MiB slice -> SDMA tier -> 4 CTAs. + EXPECT_EQ(bcastSagCtas(128ull << 20, 8, thr, kBroadcastCtasUnset, pool), kBroadcastCtasSdma); + // 2 MiB / 8 = 256 KiB slice -> at threshold -> LSA tier -> 16 CTAs. + EXPECT_EQ(bcastSagCtas(2ull << 20, 8, thr, kBroadcastCtasUnset, pool), kBroadcastCtasLsa); +} + +TEST(BcastPolicyCtas, EnvPinHonoredWithinPool) { + const int pool = bcastHybridPoolCtas(64); + const size_t thr = kBroadcastSdmaThresholdDefault; + EXPECT_EQ(bcastHybridCtas(1024, thr, 0, kBroadcastLLMaxBytes, 8, pool), 8); + EXPECT_EQ(bcastHybridCtas(64ull << 20, thr, 0, kBroadcastLLMaxBytes, 8, pool), 8); + EXPECT_EQ(bcastHybridCtas(1024, thr, 0, kBroadcastLLMaxBytes, 0, pool), kBroadcastCtasLsa); +} + +TEST(BcastPolicyCtas, LadderIsMonotonicAcrossTiers) { + // The per-tier assertions above pin which branch was taken, but each compares a + // constant against itself, so the constants' relative ordering -- the property + // the ladder exists for -- is not pinned by any of them. Assert the shape: a + // single-CTA LL tier, a wider LSA tier, and a narrow SDMA tier (only nRanks + // threads issue puts), with every grid inside [1, pool]. + const int pool = bcastHybridPoolCtas(16); + const size_t thr = kBroadcastSdmaThresholdDefault; + const int llGrid = bcastHybridCtas(1024, thr, 4096, kBroadcastLLDefaultMaxBytes, + kBroadcastCtasUnset, pool); + const int lsaGrid = bcastHybridCtas(thr, thr, 0, kBroadcastLLMaxBytes, + kBroadcastCtasUnset, pool); + const int sdmaGrid = bcastHybridCtas(thr + 1, thr, 0, kBroadcastLLMaxBytes, + kBroadcastCtasUnset, pool); + EXPECT_LT(llGrid, lsaGrid) << "LL is a single-CTA tier; LSA must be wider"; + EXPECT_LT(sdmaGrid, lsaGrid) << "GIN-put tier must be narrower than LSA"; + for (int g : {llGrid, lsaGrid, sdmaGrid}) { + EXPECT_GE(g, 1); + EXPECT_LE(g, pool); + } +} + +TEST(BcastPolicyCtas, SagEnvPinHonoredWithinPool) { + const int pool = bcastHybridPoolCtas(64); + const size_t thr = kBroadcastSdmaThresholdDefault; + EXPECT_EQ(bcastSagCtas(128ull << 20, 8, thr, 8, pool), 8); + EXPECT_EQ(bcastSagCtas(2ull << 20, 8, thr, 0, pool), + bcastSagCtas(2ull << 20, 8, thr, kBroadcastCtasUnset, pool)); +} + +TEST(BcastPoolCtas, HybridPoolAndRingPeak) { + EXPECT_EQ(bcastPoolCtas(16, 8), 16); + EXPECT_EQ(bcastPoolCtas(16, 128), 128); + EXPECT_GE(bcastPoolCtas(4, 64), 64); + EXPECT_GE(bcastPoolCtas(16, 128), bcastHybridPoolCtas(16)); +} + +TEST(BcastPolicyCtas, GridNeverExceedsPool) { + const size_t msgs[] = {128, 800, 262144, 262145, 64ull << 20}; + const int deviceVs[] = {1, 4, 16, 32, 128}; + const size_t pins[] = {kBroadcastCtasUnset, 0, 1, 8, 200, 100000}; + for (int v : deviceVs) { + const int pool = bcastHybridPoolCtas(v); + for (size_t m : msgs) { + for (size_t pin : pins) { + const int g = bcastHybridCtas(m, kBroadcastSdmaThresholdDefault, 0, + kBroadcastLLMaxBytes, pin, pool); + EXPECT_GE(g, 1); + EXPECT_LE(g, pool) << "msg=" << m << " V=" << v << " pin=" << pin; + } + } + } +} + +// --------------------------------- sagChunk ---------------------------------- + +TEST(SagChunk, GuardsNonPositiveRanks) { + Chunk c = sagChunk(800, 0, 0); + EXPECT_EQ(c.count, 0u); + EXPECT_EQ(c.eltOffset, 0u); +} + +TEST(SagChunk, EvenSplit) { + Chunk r0 = sagChunk(800, 8, 0); + EXPECT_EQ(r0.count, 100u); + EXPECT_EQ(r0.eltOffset, 0u); + Chunk r3 = sagChunk(800, 8, 3); + EXPECT_EQ(r3.count, 100u); + EXPECT_EQ(r3.eltOffset, 300u); + Chunk r7 = sagChunk(800, 8, 7); + EXPECT_EQ(r7.count, 100u); + EXPECT_EQ(r7.eltOffset, 700u); +} + +TEST(SagChunk, RemainderFoldedIntoLast) { + // 803 / 8 = base 100, last gets 803 - 700 = 103. + Chunk r0 = sagChunk(803, 8, 0); + EXPECT_EQ(r0.count, 100u); + Chunk r7 = sagChunk(803, 8, 7); + EXPECT_EQ(r7.count, 103u); + EXPECT_EQ(r7.eltOffset, 700u); + // Slices must exactly tile the whole message. + size_t total = 0; + for (int r = 0; r < 8; ++r) total += sagChunk(803, 8, r).count; + EXPECT_EQ(total, 803u); +} + +// ---------------------------- ring stripe / chunk ---------------------------- +// The default large tier walks these two splits, so they carry the exactness +// requirement that sagChunk carries for the scatter+allgather tier: every element +// covered exactly once, by exactly one CTA, in exactly one pipeline chunk. + +TEST(RingStripe, GuardsOutOfRangeBlocks) { + EXPECT_EQ(ringStripe(1000, 0, 0).count, 0u); + EXPECT_EQ(ringStripe(1000, -1, 8).count, 0u); + EXPECT_EQ(ringStripe(1000, 8, 8).count, 0u); +} + +TEST(RingStripe, PartitionsBufferExactlyAndContiguously) { + // Deliberately indivisible: 1000 elements over 128 CTAs, and a prime count. + for (size_t count : {size_t(1000), size_t(1048576), size_t(9973)}) { + for (int nBlocks : {1, 7, 8, 128}) { + size_t expectedOffset = 0; + for (int b = 0; b < nBlocks; ++b) { + Chunk s = ringStripe(count, b, nBlocks); + EXPECT_EQ(s.eltOffset, expectedOffset) + << "count=" << count << " nBlocks=" << nBlocks << " b=" << b; + expectedOffset += s.count; + } + // Contiguous from 0 and covering everything => an exact partition. + EXPECT_EQ(expectedOffset, count) + << "count=" << count << " nBlocks=" << nBlocks; + } + } +} + +TEST(RingStripe, BalancedToWithinOneElement) { + const size_t count = 1000; + const int nBlocks = 128; + size_t lo = count, hi = 0; + for (int b = 0; b < nBlocks; ++b) { + size_t c = ringStripe(count, b, nBlocks).count; + lo = (c < lo) ? c : lo; + hi = (c > hi) ? c : hi; + } + EXPECT_LE(hi - lo, 1u); +} + +TEST(RingChunk, GuardsOutOfRangeChunks) { + EXPECT_EQ(ringChunk(0, 1000, 0, 0).count, 0u); + EXPECT_EQ(ringChunk(0, 1000, -1, 16).count, 0u); + EXPECT_EQ(ringChunk(0, 1000, 16, 16).count, 0u); +} + +TEST(RingChunk, PartitionsStripeFromItsBase) { + const size_t stripeBase = 4096; + const size_t stripeCount = 803; // indivisible by C below + const int C = 16; + size_t expectedOffset = stripeBase; + for (int c = 0; c < C; ++c) { + Chunk k = ringChunk(stripeBase, stripeCount, c, C); + EXPECT_EQ(k.eltOffset, expectedOffset) << "c=" << c; + expectedOffset += k.count; + } + EXPECT_EQ(expectedOffset, stripeBase + stripeCount); +} + +TEST(RingChunk, ComposesWithRingStripeToCoverTheBuffer) { + // The two splits together are what the kernel walks: every element of the + // message must land in exactly one (block, chunk) pair. + const size_t count = 9973; // prime + const int nBlocks = 8; + const int C = 7; + size_t total = 0; + size_t expectedOffset = 0; + for (int b = 0; b < nBlocks; ++b) { + Chunk s = ringStripe(count, b, nBlocks); + for (int c = 0; c < C; ++c) { + Chunk k = ringChunk(s.eltOffset, s.count, c, C); + EXPECT_EQ(k.eltOffset, expectedOffset) << "b=" << b << " c=" << c; + expectedOffset += k.count; + total += k.count; + } + } + EXPECT_EQ(total, count); + EXPECT_EQ(expectedOffset, count); +} + +} // namespace diff --git a/projects/rccl/tools/ci/lib/gin-tests.json b/projects/rccl/tools/ci/lib/gin-tests.json index 70f050989969..f662f91ca47b 100644 --- a/projects/rccl/tools/ci/lib/gin-tests.json +++ b/projects/rccl/tools/ci/lib/gin-tests.json @@ -4,7 +4,9 @@ "docker_build_test.bash). Each test runs as:", " mpirun -np $NP <-x env> -x LD_LIBRARY_PATH ", "kind: 'rocshmem' -> $ROCSHMEM_TESTS_BIN_DIR/; 'rccl-tests' -> $RCCL_TESTS_BIN_DIR/;", - "'fixtures' -> $RCCL_FIXTURES_BIN_DIR/ (single-process gtest, no mpirun)", + "'fixtures' -> $RCCL_FIXTURES_BIN_DIR/ (single-process gtest, no mpirun);", + "'pytest' -> python3 -m pytest under $RCCL_TESTS_DIR/test/ (opt-in rccl-tests", + "regressions; exports RCCL_TESTS_GIN_SDMA_BCAST=1 and points BCAST_EXE at broadcast_perf)", "{E} is replaced with NP*MSG_SIZE bytes. Tests are non-gating for now.", "-D 3 is the rocSHMEM-SDMA path the script labels 'Crash' (kept for visibility).", "debug_env: -x added to every test only when RCCL_CI_DEBUG=1.", @@ -38,6 +40,27 @@ "env": ["NCCL_CUMEM_ENABLE=1", "RCCL_ENABLE_INTRANET=1", "NCCL_DMABUF_ENABLE=1", "HSA_NO_SCRATCH_RECLAIM=1", "NCCL_MSCCL_ENABLE=0" ,"NCCL_GIN_ENABLE=1", "NCCL_GIN_TYPE=6"], "args": "-b 128 -e {E} -f 2 -g 1 -R 2 -D 3 -A 1" }, + { + "name": "rccl-gin-rocshmem-sdma-bcast-d3", + "kind": "rccl-tests", + "bin": "broadcast_perf", + "env": ["NCCL_CUMEM_ENABLE=1", "RCCL_ENABLE_INTRANET=1", "NCCL_DMABUF_ENABLE=1", "HSA_NO_SCRATCH_RECLAIM=1", "NCCL_MSCCL_ENABLE=0", "NCCL_GIN_ENABLE=1", "NCCL_GIN_TYPE=6"], + "args": "-b 128 -e {E} -f 2 -g 1 -R 2 -D 3 -A 1" + }, + { + "name": "rccl-gin-bcast-pytest", + "kind": "pytest", + "bin": "test_Broadcast.py", + "env": ["NCCL_CUMEM_ENABLE=1", "RCCL_ENABLE_INTRANET=1", "NCCL_DMABUF_ENABLE=1", "HSA_NO_SCRATCH_RECLAIM=1", "NCCL_MSCCL_ENABLE=0"], + "args": "-k GinSdma -v" + }, + { + "name": "rccl-gin-bcast-policy-ut", + "kind": "fixtures", + "bin": "rccl-UnitTestsGinSdmaBroadcastPolicy", + "env": [], + "args": "" + }, { "name": "rccl-gin-fixtures-anvil-gda", "kind": "fixtures", diff --git a/projects/rccl/tools/ci/lib/gin.sbatch b/projects/rccl/tools/ci/lib/gin.sbatch index 6f918f6c77ef..5e8ace5fc354 100755 --- a/projects/rccl/tools/ci/lib/gin.sbatch +++ b/projects/rccl/tools/ci/lib/gin.sbatch @@ -16,7 +16,8 @@ # 2. lib/build-ompi.sh build OpenMPI (--with-rocm) -> ompi.env # 3. lib/build-rocshmem.sh build rocSHMEM (IPC+SDMA) -> rocshmem.env # 4. projects/rccl install.sh build RCCL (GIN/rocSHMEM), then build only -# rccl-UnitTestsFixtures (GIN_ANVIL_UNIT_TESTS=ON) -> rccl.env +# rccl-UnitTestsFixtures and rccl-UnitTestsGinSdmaBroadcastPolicy +# (GIN_ANVIL_UNIT_TESTS=ON) -> rccl.env # 5. projects/rccl-tests build rccl-tests (DEVICE_API + ROCSHMEM) # 6. run-gin-ci.sh run the GIN/rocSHMEM tests (non-gating) # @@ -116,6 +117,7 @@ rccl_src="${workdir}/projects/rccl" rccl_prefix="${workdir}/.ci-out/rccl-install" rccl_build="${rccl_src}/build/release" rccl_fixtures_bin="${rccl_build}/test/rccl-UnitTestsFixtures" +rccl_bcast_policy_ut_bin="${rccl_build}/test/rccl-UnitTestsGinSdmaBroadcastPolicy" rm -rf "${rccl_prefix}" ( cd "${rccl_src}" @@ -131,21 +133,27 @@ if [[ ! -f "${rccl_prefix}/lib/librccl.so" && ! -f "${rccl_prefix}/lib/librccl.s exit 1 fi -echo "==> [4/6] build rccl-UnitTestsFixtures (GIN_ANVIL_UNIT_TESTS=ON)" +echo "==> [4/6] build rccl-UnitTestsFixtures + GinSdmaBroadcastPolicy (GIN_ANVIL_UNIT_TESTS=ON)" ( cd "${rccl_build}" cmake . \ -DBUILD_TESTS=ON \ -DGIN_ANVIL_UNIT_TESTS=ON \ -DROCSHMEM_INSTALL_DIR="${ROCSHMEM_INSTALL_DIR}" - make -j"${build_jobs}" rccl-UnitTestsFixtures + make -j"${build_jobs}" rccl-UnitTestsFixtures rccl-UnitTestsGinSdmaBroadcastPolicy ) if [[ ! -x "${rccl_fixtures_bin}" ]]; then echo "ERROR: rccl-UnitTestsFixtures not built at ${rccl_fixtures_bin}" >&2 ls -la "$(dirname "${rccl_fixtures_bin}")" 2>/dev/null >&2 || true exit 1 fi +if [[ ! -x "${rccl_bcast_policy_ut_bin}" ]]; then + echo "ERROR: rccl-UnitTestsGinSdmaBroadcastPolicy not built at ${rccl_bcast_policy_ut_bin}" >&2 + ls -la "$(dirname "${rccl_bcast_policy_ut_bin}")" 2>/dev/null >&2 || true + exit 1 +fi echo " rccl-UnitTestsFixtures: ${rccl_fixtures_bin}" +echo " rccl-UnitTestsGinSdmaBroadcastPolicy: ${rccl_bcast_policy_ut_bin}" # Locate the in-tree RCCL CMake package config. rccl-tests' find_package(RCCL) # otherwise resolves roc::rccl to the nightly ROCm (its toolchain prepends @@ -185,6 +193,11 @@ if [[ ! -x "${rccl_tests_build}/alltoall_perf" ]]; then ls -la "${rccl_tests_build}" 2>/dev/null >&2 || true exit 1 fi +if [[ ! -x "${rccl_tests_build}/broadcast_perf" ]]; then + echo "ERROR: broadcast_perf not built under ${rccl_tests_build}" >&2 + ls -la "${rccl_tests_build}" 2>/dev/null >&2 || true + exit 1 +fi # Diag: confirm alltoall_perf exports the rocSHMEM host symbol librccl.so resolves # at runtime (log_flags); a miss here => runtime 'undefined symbol'. @@ -200,6 +213,7 @@ nm -CD "${librccl}" 2>/dev/null | grep -E " U .*${sym}" || echo " not undefined { printf 'export RCCL_INSTALL_PREFIX=%q\n' "${rccl_prefix}" printf 'export RCCL_TESTS_BIN_DIR=%q\n' "${rccl_tests_build}" + printf 'export RCCL_TESTS_DIR=%q\n' "${rccl_tests_src}" printf 'export RCCL_FIXTURES_BIN_DIR=%q\n' "$(dirname "${rccl_fixtures_bin}")" } > "${workdir}/.ci-out/rccl.env" ) > "${build_out}" 2> "${build_err}" || build_rc=$? diff --git a/projects/rccl/tools/ci/lib/parse_gin_config.py b/projects/rccl/tools/ci/lib/parse_gin_config.py index fe7298024231..b31bfd23e409 100755 --- a/projects/rccl/tools/ci/lib/parse_gin_config.py +++ b/projects/rccl/tools/ci/lib/parse_gin_config.py @@ -10,7 +10,11 @@ mca debug_env<-x flags> (appended to every test only when RCCL_CI_DEBUG=1) test<-x env flags> - kind: rocshmem | rccl-tests | fixtures (fixtures: gtest binary, no mpirun) + kind: rocshmem | rccl-tests | fixtures | pytest + rocshmem/rccl-tests: mpirun benchmark; bin is the executable name + fixtures: single-process gtest; bin is the executable name + pytest: rccl-tests opt-in pytest; bin is a file under test/ (e.g. + test_Broadcast.py); args are forwarded to pytest (e.g. -k gin_sdma -v) """ # ASCII unit separator: distinct from whitespace and unlikely in flags/args. diff --git a/projects/rccl/tools/ci/lib/test_parse_gin_config.py b/projects/rccl/tools/ci/lib/test_parse_gin_config.py index 412d0f092db7..125b5db24288 100644 --- a/projects/rccl/tools/ci/lib/test_parse_gin_config.py +++ b/projects/rccl/tools/ci/lib/test_parse_gin_config.py @@ -40,6 +40,34 @@ def test_empty_env_preserves_gtest_filter_in_args(self) -> None: self.assertEqual(parts[4], "") self.assertEqual(parts[5], "--gtest_filter=GinRocshmemGdaTemplateTest.*") + def test_pytest_kind_preserves_args(self) -> None: + config = parse_config( + _write_json( + { + "mca": "", + "debug_env": [], + "tests": [ + { + "name": "bcast-pytest", + "kind": "pytest", + "bin": "test_Broadcast.py", + "env": ["NCCL_CUMEM_ENABLE=1", "HSA_NO_SCRATCH_RECLAIM=1"], + "args": "-k GinSdma -v", + } + ], + } + ) + ) + rows = format_rows(config) + test_row = next(r for r in rows if r.startswith("test" + _FIELD_SEP)) + parts = test_row.split(_FIELD_SEP) + self.assertEqual(parts[2], "pytest") + self.assertEqual(parts[3], "test_Broadcast.py") + self.assertEqual( + parts[4], "-x NCCL_CUMEM_ENABLE=1 -x HSA_NO_SCRATCH_RECLAIM=1" + ) + self.assertEqual(parts[5], "-k GinSdma -v") + def _write_json(data: dict) -> Path: handle = tempfile.NamedTemporaryFile("w", suffix=".json", delete=False) diff --git a/projects/rccl/tools/ci/run-gin-ci.sh b/projects/rccl/tools/ci/run-gin-ci.sh index 7f5d93745774..00a610c3cd07 100755 --- a/projects/rccl/tools/ci/run-gin-ci.sh +++ b/projects/rccl/tools/ci/run-gin-ci.sh @@ -18,6 +18,18 @@ # RCCL_CI_DEBUG Set to 1 to add the config's debug_env to every run # RCCL_CI_DEBUG_DIR Dir for NCCL_DEBUG_FILE output when RCCL_CI_DEBUG=1 # (default: ${SLURM_SUBMIT_DIR:-$PWD}/nccl-debug) +# RCCL_TESTS_DIR rccl-tests source tree (default: $WORKDIR/projects/rccl-tests) +# GIN_PYTEST_TIMEOUT Wall-clock cap for pytest matrix entries (default: 1800s) +# GIN_PYTEST_HW_CASES mpirun cases under -k GinSdma (default: 9). Offline +# parser/tier guards do not launch and are not counted. +# RCCL_TESTS_BCAST_GIN_TYPE NCCL_GIN_TYPE for Broadcast pytest (default: 6) +# RCCL_TESTS_BCAST_TIMEOUT_S / RCCL_TESTS_BCAST_CONN_RETRIES +# Inner pytest launch budget. Unset → derived so +# HW_CASES * retries * TIMEOUT_S is strictly under +# GIN_PYTEST_TIMEOUT minus kill-after and slack. Needed +# because GNU timeout killing pytest does not killpg the +# mpirun session (start_new_session=True); pytest's own +# TimeoutExpired handler is what reaps the group. set -euo pipefail @@ -25,9 +37,14 @@ NP="${NP:-8}" MSG_SIZE="${MSG_SIZE:-33554432}" BENCH_TIMEOUT="${BENCH_TIMEOUT:-600s}" BENCH_KILL_AFTER="${BENCH_KILL_AFTER:-30s}" +GIN_PYTEST_TIMEOUT="${GIN_PYTEST_TIMEOUT:-1800s}" +# Hardware launches selected by rccl-gin-bcast-pytest (-k GinSdma): 6 segmented +# (2 sizes x 3 dtypes) + scatter-allgather + 2 hang guards. +GIN_PYTEST_HW_CASES="${GIN_PYTEST_HW_CASES:-9}" script_dir="$(cd "$(dirname "$0")" && pwd)" WORKDIR="$(cd "${script_dir}/../../../.." && pwd)" +RCCL_TESTS_DIR="${RCCL_TESTS_DIR:-${WORKDIR}/projects/rccl-tests}" CONFIG="${CONFIG:-${script_dir}/lib/gin-tests.json}" PARSER="${script_dir}/lib/parse_gin_config.py" @@ -69,6 +86,7 @@ echo "==> ROCSHMEM_INSTALL_DIR = ${ROCSHMEM_INSTALL_DIR}" echo "==> RCCL_INSTALL_PREFIX = ${RCCL_INSTALL_PREFIX}" echo "==> RCCL_TESTS_BIN_DIR = ${RCCL_TESTS_BIN_DIR}" echo "==> RCCL_FIXTURES_BIN_DIR= ${RCCL_FIXTURES_BIN_DIR}" +echo "==> RCCL_TESTS_DIR = ${RCCL_TESTS_DIR}" echo "==> NP=${NP} MSG_SIZE=${MSG_SIZE} (E=NP*MSG_SIZE=${E})" echo "==> test matrix = ${CONFIG}" @@ -101,21 +119,96 @@ echo "==> ${#TEST_NAMES[@]} tests to run: ${TEST_NAMES[*]}" FAILED_RUNS=() +ensure_pytest() { + # shellcheck source=/dev/null + [[ -f "${script_dir}/lib/ensure-python-yaml.sh" ]] && source "${script_dir}/lib/ensure-python-yaml.sh" + if ! python3 -c 'import pytest' 2>/dev/null; then + echo "==> pip installing pytest" + python3 -m pip install --quiet --disable-pip-version-check pytest + fi +} + +# env_flags is "-x K=V ..."; pytest uses RCCL_TESTS_BCAST_XENV (plain K=V tokens). +gin_env_to_bcast_xenv() { + local raw="$1" + raw="${raw//-x /}" + printf '%s' "${raw}" +} + +# GNU timeout accepts 1800 / 1800s / 30m / 1h. Integer seconds only. +gin_duration_sec() { + local raw="${1// /}" + if [[ "${raw}" =~ ^([0-9]+)([smh]?)$ ]]; then + local n="${BASH_REMATCH[1]}" + case "${BASH_REMATCH[2]}" in + m) echo $((n * 60)) ;; + h) echo $((n * 3600)) ;; + *) echo "${n}" ;; + esac + return 0 + fi + echo "ERROR: cannot parse duration '${1}'" >&2 + return 1 +} + +# pytest's communicate()+killpg is the only path that reaps mpirun (new session). +# Size TIMEOUT_S and CONN_RETRIES so even the worst-case inner budget (every +# hardware case using every retry, each waiting the full per-attempt cap) is +# strictly under GIN_PYTEST_TIMEOUT. Leave already-set env vars alone. +apply_gin_pytest_inner_budget() { + local outer_s kill_s slack usable retries denom timeout_s inner + outer_s="$(gin_duration_sec "${GIN_PYTEST_TIMEOUT}")" || return 1 + kill_s="$(gin_duration_sec "${BENCH_KILL_AFTER}")" || return 1 + slack=60 + usable=$((outer_s - kill_s - slack)) + if ((usable < 30)); then + echo "ERROR: GIN_PYTEST_TIMEOUT=${GIN_PYTEST_TIMEOUT} leaves no room for pytest after kill-after/slack" >&2 + return 1 + fi + if [[ -z "${RCCL_TESTS_BCAST_CONN_RETRIES:-}" ]]; then + export RCCL_TESTS_BCAST_CONN_RETRIES=2 + fi + retries="${RCCL_TESTS_BCAST_CONN_RETRIES}" + if ((retries < 1)); then + retries=1 + export RCCL_TESTS_BCAST_CONN_RETRIES=1 + fi + denom=$((GIN_PYTEST_HW_CASES * retries)) + if ((denom < 1)); then + denom=1 + fi + if [[ -z "${RCCL_TESTS_BCAST_TIMEOUT_S:-}" ]]; then + timeout_s=$(((usable - 1) / denom)) + if ((timeout_s < 30)); then + timeout_s=30 + fi + export RCCL_TESTS_BCAST_TIMEOUT_S="${timeout_s}" + fi + inner=$((GIN_PYTEST_HW_CASES * retries * RCCL_TESTS_BCAST_TIMEOUT_S)) + echo "==> pytest inner budget: RCCL_TESTS_BCAST_TIMEOUT_S=${RCCL_TESTS_BCAST_TIMEOUT_S}s retries=${retries} hw_cases=${GIN_PYTEST_HW_CASES} max=${inner}s < outer ${outer_s}s (kill-after ${kill_s}s, slack ${slack}s)" + if ((inner + kill_s + slack >= outer_s)); then + echo "WARNING: inner budget ${inner}s is not strictly under GIN_PYTEST_TIMEOUT=${outer_s}s; raise the outer cap or lower TIMEOUT_S/retries" >&2 + fi +} + # Word-splitting on flag/arg vars below is intentional. # shellcheck disable=SC2086 run_test() { local name="$1" kind="$2" bin="$3" env_flags="$4" args="$5" - local bin_path + local bin_path bench_timeout="${BENCH_TIMEOUT}" case "${kind}" in rocshmem) bin_path="${ROCSHMEM_TESTS_BIN_DIR}/${bin}" ;; rccl-tests) bin_path="${RCCL_TESTS_BIN_DIR}/${bin}" ;; fixtures) bin_path="${RCCL_FIXTURES_BIN_DIR}/${bin}" ;; + pytest) bench_timeout="${GIN_PYTEST_TIMEOUT}" ;; *) echo " SKIP ${name}: unknown kind '${kind}'"; FAILED_RUNS+=("${name} (unknown kind)"); return ;; esac - if [[ ! -x "${bin_path}" ]]; then - echo " SKIP ${name}: binary not found/executable: ${bin_path}" - FAILED_RUNS+=("${name} (missing ${bin})") - return + if [[ "${kind}" != "pytest" ]]; then + if [[ ! -x "${bin_path}" ]]; then + echo " SKIP ${name}: binary not found/executable: ${bin_path}" + FAILED_RUNS+=("${name} (missing ${bin})") + return + fi fi args="${args//\{E\}/${E}}" if [[ -n "${RCCL_CI_DEBUG:-}" && -n "${DEBUG_ENV}" ]]; then @@ -123,12 +216,47 @@ run_test() { fi echo "=== ${name}: ${bin} ${args} ===" set +e - if [[ "${kind}" == "fixtures" ]]; then - timeout --kill-after="${BENCH_KILL_AFTER}" "${BENCH_TIMEOUT}" \ + if [[ "${kind}" == "pytest" ]]; then + local pytest_dir="${RCCL_TESTS_DIR}/test" + local pytest_file="${pytest_dir}/${bin}" + local bcast_exe="${RCCL_TESTS_BIN_DIR}/broadcast_perf" + local bcast_xenv + if [[ ! -f "${pytest_file}" ]]; then + echo " SKIP ${name}: pytest file not found: ${pytest_file}" + FAILED_RUNS+=("${name} (missing ${bin})") + set -e + return + fi + if [[ ! -x "${bcast_exe}" ]]; then + echo " SKIP ${name}: broadcast_perf not found/executable: ${bcast_exe}" + FAILED_RUNS+=("${name} (missing broadcast_perf)") + set -e + return + fi + ensure_pytest + apply_gin_pytest_inner_budget || { + FAILED_RUNS+=("${name} (inner budget)") + set -e + return + } + bcast_xenv="$(gin_env_to_bcast_xenv "${env_flags}")" + timeout --kill-after="${BENCH_KILL_AFTER}" "${bench_timeout}" \ + env LD_LIBRARY_PATH="${LD_LIBRARY_PATH}" \ + RCCL_TESTS_GIN_SDMA_BCAST=1 \ + RCCL_TESTS_BCAST_EXE="${bcast_exe}" \ + RCCL_TESTS_BCAST_NP="${NP}" \ + RCCL_TESTS_BCAST_GIN_TYPE="${RCCL_TESTS_BCAST_GIN_TYPE:-6}" \ + RCCL_TESTS_BCAST_TIMEOUT_S="${RCCL_TESTS_BCAST_TIMEOUT_S}" \ + RCCL_TESTS_BCAST_CONN_RETRIES="${RCCL_TESTS_BCAST_CONN_RETRIES}" \ + RCCL_TESTS_BCAST_XENV="${bcast_xenv}" \ + RCCL_TESTS_MPI_LAUNCHER="${MPI_HOME}/bin/mpirun" \ + python3 -m pytest "${pytest_file}" ${args} -p no:cacheprovider + elif [[ "${kind}" == "fixtures" ]]; then + timeout --kill-after="${BENCH_KILL_AFTER}" "${bench_timeout}" \ env LD_LIBRARY_PATH="${LD_LIBRARY_PATH}" \ "${bin_path}" ${args} else - timeout --kill-after="${BENCH_KILL_AFTER}" "${BENCH_TIMEOUT}" \ + timeout --kill-after="${BENCH_KILL_AFTER}" "${bench_timeout}" \ mpirun -np "${NP}" ${MCA} ${env_flags} -x LD_LIBRARY_PATH \ "${bin_path}" ${args} fi @@ -136,7 +264,7 @@ run_test() { set -e if [[ ${rc} -ne 0 ]]; then if [[ ${rc} -eq 124 || ${rc} -eq 137 ]]; then - FAILED_RUNS+=("${name} (TIMEOUT >${BENCH_TIMEOUT}, rc=${rc})") + FAILED_RUNS+=("${name} (TIMEOUT >${bench_timeout}, rc=${rc})") else FAILED_RUNS+=("${name} (rc=${rc})") fi