test(rccl): clamp isolated-test GPU pool to devices HIP can open

projects/rccl/test/GraphCaptureStreamOrderingTests.cpp

@@ -76,8 +76,16 @@ struct Harness
 
 void initHarness(Harness& h, int nRanks)
 {
-    int devCount = 0;
-    HIP_CHECK(hipGetDeviceCount(&devCount));
+    // A machine with no usable GPU is a reason to skip, not to fail, so
+    // hipErrorNoDevice is folded into a count of zero and falls through to the
+    // guard below. Anything else is a genuine HIP failure and is reported.
+    int              devCount = 0;
+    const hipError_t countErr = hipGetDeviceCount(&devCount);
+    if(countErr == hipErrorNoDevice)
+        devCount = 0;
+    else
+        HIP_CHECK(countErr);
+
     if(devCount < nRanks)
         GTEST_SKIP() << "Requires >= " << nRanks << " GPUs; found " << devCount;
 

projects/rccl/test/common/ProcessIsolatedTestRunner.cpp

@@ -8,7 +8,9 @@
 #include <errno.h>
 #include <fcntl.h>
 #include <gtest/gtest.h>
+#include <hip/hip_runtime_api.h>
 #include <poll.h>
+#include <sys/wait.h>
 #include <unistd.h>
 
 #include <algorithm>
@@ -199,14 +201,69 @@ static std::vector<int> gpuPoolFromKfd()
     return pool;
 }
 
+// Returns the number of devices HIP can actually open, or kProbeFailed if the
+// probe could not produce an answer.
+//
+// The query runs in a short-lived forked child so the parent never initialises
+// the HIP runtime: doing so in the parent would open GPU file descriptors that
+// every subsequently forked test child would inherit, which is the same hazard
+// gpuPoolFromKfd() avoids by reading sysfs instead of a KFD device file.
+//
+// The count travels back as the child's exit status, so it is capped at
+// kProbeMaxCount; a pool larger than that is already far beyond any real node.
+static constexpr int kProbeFailed   = -1;
+static constexpr int kProbeMaxCount = 254;
+static constexpr int kProbeErrorExit = 255;
+
+static int hipDeviceCountViaChild()
+{
+    const pid_t pid = fork();
+    if(pid < 0)
+        return kProbeFailed;
+
+    if(pid == 0)
+    {
+        int count = 0;
+        if(hipGetDeviceCount(&count) != hipSuccess || count < 0)
+            _exit(kProbeErrorExit);
+        _exit(std::min(count, kProbeMaxCount));
+    }
+
+    int status = 0;
+    while(waitpid(pid, &status, 0) < 0 && errno == EINTR)
+    {
+    }
+    if(!WIFEXITED(status))
+        return kProbeFailed;
+
+    const int code = WEXITSTATUS(status);
+    return code == kProbeErrorExit ? kProbeFailed : code;
+}
+
+std::vector<int> clampPoolToDeviceCount(std::vector<int> pool, int hipDeviceCount)
+{
+    if(hipDeviceCount > 0 && static_cast<size_t>(hipDeviceCount) < pool.size())
+    {
+        TEST_INFO(
+            "GPU pool reported %zu device(s) but HIP can open only %d; "
+            "clamping pool to %d",
+            pool.size(), hipDeviceCount, hipDeviceCount
+        );
+        pool.resize(static_cast<size_t>(hipDeviceCount));
+    }
+    return pool;
+}
+
 // Detects available GPU indices: first from HIP_VISIBLE_DEVICES, then from
-// KFD sysfs. Returns empty if neither source yields a pool.
+// KFD sysfs, then clamped to the devices HIP can actually open. Returns empty
+// if neither source yields a pool.
 static std::vector<int> detectGpuPool()
 {
     auto pool = gpuPoolFromEnv();
     if(pool.empty())
         pool = gpuPoolFromKfd();
-    return pool;
+
+    return clampPoolToDeviceCount(std::move(pool), hipDeviceCountViaChild());
 }
 
 // Tracks which physical GPU device indices from a fixed pool are in use.

projects/rccl/test/common/ProcessIsolatedTestRunner.hpp

@@ -19,6 +19,42 @@
 namespace RcclUnitTesting
 {
 
+/**
+ * @brief Shrinks a detected GPU pool to the number of devices HIP can open.
+ *
+ * Both pool sources over-report. KFD topology sysfs is not namespaced per
+ * container, so a partitioned CI slot still sees the host's GPUs, and an
+ * inherited HIP_VISIBLE_DEVICES may name ordinals this process cannot open.
+ * Handing out an ordinal at or above the real device count makes HIP fail the
+ * child with hipErrorNoDevice instead of running or skipping it.
+ *
+ * The pool holds absolute device ids in the pre-visibility namespace — the
+ * values written verbatim into a child's HIP_VISIBLE_DEVICES — and its order
+ * is significant: KFD sysfs yields ascending 0..n-1, while HIP_VISIBLE_DEVICES
+ * is taken in the order given and may be arbitrary ("3,1,7,5").
+ *
+ * Clamping is therefore by *length*, not by discarding ids >= the count. HIP
+ * makes visible the longest valid prefix of HIP_VISIBLE_DEVICES and reports
+ * its length, so the first @p hipDeviceCount entries are exactly the entries
+ * HIP can open (HIP_VISIBLE_DEVICES=0,5 on a 1-GPU host reports 1, and 5,0
+ * reports none). Dropping ids by value would instead empty a legitimate pool:
+ * HIP_VISIBLE_DEVICES=4,5,6,7 on an 8-GPU host reports 4 devices even though
+ * every id in it is >= 4.
+ *
+ * A non-positive @p hipDeviceCount means the probe failed or found no GPU at
+ * all; the pool is returned unchanged so that "this machine has no GPUs" stays
+ * a per-test skip rather than becoming a pool-detection failure.
+ *
+ * Exposed (rather than kept file-static in the .cpp) so the pool-sizing
+ * decision can be unit-tested without a GPU.
+ *
+ * @param pool            Pool detected from HIP_VISIBLE_DEVICES or KFD sysfs,
+ *                        in the order that source gave them.
+ * @param hipDeviceCount  Devices HIP reports, or <= 0 if unknown.
+ * @return The pool, truncated to @p hipDeviceCount when that is smaller.
+ */
+std::vector<int> clampPoolToDeviceCount(std::vector<int> pool, int hipDeviceCount);
+
 /**
  * @brief Generic thread-safe process isolated test runner
  *

projects/rccl/test/host/CMakeLists.txt

@@ -827,6 +827,7 @@ function(rccl_define_micro_source_lists)
   # rccl-UnitTestsMicroParam -- src/param/{param,c_api}.cc, src/misc/param.cc. Own binary: misc/param.cc
   # defines ncclGetEnv, initEnv and ncclLoadParam, which the shared fakes in rccl-UnitTestsMicro also define.
   set(TEST_MICRO_PARAM_SOURCE_FILES
+    fakes/hip_fakes.cc
     param-test.cc
     ../common/ProcessIsolatedTestRunner.cpp
     ../common/main_altrsmi.cpp
@@ -857,6 +858,7 @@ function(rccl_define_micro_source_lists)
   # rccl-UnitTestsMicroPlugin -- src/plugin/net.cc, net/net_v*.cc, plugin_open.cc. Own binary: gin-plugin-init-test.cc
   # defines ncclPluginLibPaths, ncclClosePluginLib and ncclGetNetPluginLib in rccl-UnitTestsMicro.
   set(TEST_MICRO_PLUGIN_SOURCE_FILES
+    fakes/hip_fakes.cc
     plugin-net-test.cc
     plugin-open-test.cc
     net-v6-test.cc
@@ -902,7 +904,7 @@ function(rccl_define_micro_source_lists)
     fakes/utils_fakes.cc
     ../common/main_altrsmi.cpp
     PARENT_SCOPE)
-    
+
   # rccl-UnitTestsMicroDdaIpc -- src/algorithms/dda/all_reduce/dda_all_reduce_ipc.cu. Own binary:
   # rccl-UnitTestsMicro links fakes/dda_fakes.cc, which defines the same ncclAllReduceDdaIpc* symbols.
   set(TEST_MICRO_DDA_IPC_SOURCE_FILES
@@ -1187,6 +1189,7 @@ function(rccl_define_micro_source_lists)
 
   # src/misc/ IB verbs and direct-verbs loaders
   set(TEST_MICRO_MISC_SOURCE_FILES
+    fakes/hip_fakes.cc
     misc/gdr-peermem-test.cc
     misc/ibvsymbols-test.cc
     misc/dvsymbols-test.cc               # mlx5dvsymbols.cc + ionicdvsymbols.cc
@@ -1496,6 +1499,7 @@ function(rccl_standalone_setup)
     ${RCCL_TEST_DIR}/DdaCollCommonTests.cpp
     ${RCCL_TEST_DIR}/RomeTopoConsensusTests.cpp
     RomeTopoConsensusLogTests.cpp
+    process-isolated-runner-pool-test.cc
     ${RCCL_TEST_DIR}/MiscTests.cpp
     ${RCCL_TEST_DIR}/VersionGateTests.cpp
     ${RCCL_TEST_DIR}/TimeoutTests.cpp

projects/rccl/test/host/process-isolated-runner-pool-test.cc

@@ -0,0 +1,87 @@
+/*************************************************************************
+ * Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+// Pool sizing in test/common/ProcessIsolatedTestRunner.cpp.
+//
+// The runner hands each isolated child a slice of its GPU pool via
+// HIP_VISIBLE_DEVICES. Neither pool source can be trusted to describe the
+// devices this process can open: KFD topology sysfs is not namespaced per
+// container, and an inherited HIP_VISIBLE_DEVICES is just whatever the caller
+// set. clampPoolToDeviceCount() is the one place that reconciles the detected
+// pool with the device count HIP reports, so it is pinned here.
+
+#include <gtest/gtest.h>
+
+#include <vector>
+
+#include "../common/ProcessIsolatedTestRunner.hpp"
+
+namespace
+{
+
+using RcclUnitTesting::clampPoolToDeviceCount;
+
+std::vector<int> poolOf(int n)
+{
+    std::vector<int> pool(static_cast<size_t>(n));
+    for(int i = 0; i < n; ++i) pool[static_cast<size_t>(i)] = i;
+    return pool;
+}
+
+// The CI case: an 8-GPU host topology seen from a container where HIP can open
+// four devices. Ordinals >= 4 must never reach a child.
+TEST(ProcessIsolatedRunnerPool, TruncatesPoolToTheDeviceCountHipReports)
+{
+    EXPECT_EQ(clampPoolToDeviceCount(poolOf(8), 4), (std::vector<int>{0, 1, 2, 3}));
+}
+
+TEST(ProcessIsolatedRunnerPool, KeepsPoolWhenHipReportsAsManyDevicesAsTheSource)
+{
+    EXPECT_EQ(clampPoolToDeviceCount(poolOf(4), 4), poolOf(4));
+}
+
+// A pool smaller than the machine is a deliberate restriction (for example an
+// explicit HIP_VISIBLE_DEVICES); the clamp must not widen it back out.
+TEST(ProcessIsolatedRunnerPool, KeepsPoolWhenHipReportsMoreDevicesThanTheSource)
+{
+    EXPECT_EQ(clampPoolToDeviceCount(poolOf(2), 8), poolOf(2));
+}
+
+// HIP exposes the longest valid prefix of HIP_VISIBLE_DEVICES and reports its
+// length, so the kept entries are the leading ones in source order, whatever
+// ids those happen to be.
+TEST(ProcessIsolatedRunnerPool, KeepsTheLeadingSourceIdsInSourceOrder)
+{
+    EXPECT_EQ(clampPoolToDeviceCount({3, 1, 7, 5}, 2), (std::vector<int>{3, 1}));
+}
+
+// The ids are absolute device numbers, not offsets into the pool, and a count
+// equal to the pool's length means every one of them is openable. Discarding
+// ids >= the count instead of trimming by length would empty this pool, which
+// is what an 8-GPU host restricted to HIP_VISIBLE_DEVICES=4,5,6,7 looks like.
+TEST(ProcessIsolatedRunnerPool, KeepsIdsAboveTheDeviceCountWhenAllOfThemAreVisible)
+{
+    EXPECT_EQ(clampPoolToDeviceCount({4, 5, 6, 7}, 4), (std::vector<int>{4, 5, 6, 7}));
+}
+
+// A failed probe and a genuinely GPU-less machine both leave the pool alone:
+// emptying it would turn a per-test skip into a pool-detection failure.
+TEST(ProcessIsolatedRunnerPool, KeepsPoolWhenTheDeviceCountIsUnknown)
+{
+    EXPECT_EQ(clampPoolToDeviceCount(poolOf(8), -1), poolOf(8));
+}
+
+TEST(ProcessIsolatedRunnerPool, KeepsPoolWhenHipReportsNoDevices)
+{
+    EXPECT_EQ(clampPoolToDeviceCount(poolOf(8), 0), poolOf(8));
+}
+
+TEST(ProcessIsolatedRunnerPool, LeavesAnUndetectedPoolEmpty)
+{
+    EXPECT_TRUE(clampPoolToDeviceCount({}, 8).empty());
+}
+
+} // namespace

projects/rccl/test/host/wrapper_link_stubs.cpp

@@ -101,6 +101,7 @@ int ncclParamWorkArgsBytes() { return 0; }
 // --- HIP runtime stubs (return hipSuccess = 0) ---
 extern "C" {
 int hipHostFree(void*) { return 0; }
+int hipGetDeviceCount(int* count) { if (count) *count = 0; return 0; }
 const char* hipGetErrorString(int) { return "stub"; }
 int hipGetLastError() { return 0; }
 int hipMemAddressFree(void*, size_t) { return 0; }