From 090e0318d27fea052b29bac35349bbadad928763 Mon Sep 17 00:00:00 2001 From: Shweta Jha Date: Wed, 7 Oct 2026 10:23:52 -0500 Subject: [PATCH 1/5] fix(rccl): return an error when /proc/version cannot be read ncclGetUniqueId() crashed in strstr when a failed open left the kernel version string null. Co-authored-by: Cursor --- projects/rccl/CHANGELOG.md | 1 + projects/rccl/src/init.cc | 10 +++++++-- projects/rccl/test/host/fakes/os_fakes.cc | 3 ++- projects/rccl/test/host/init-test.cc | 27 ++++++++++++++++++++++- 4 files changed, 37 insertions(+), 4 deletions(-) diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index a5dba4d34b43..9cbb8878232a 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -37,6 +37,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * `NCCL_TOPO_SCATTER_START_NET`, which selected how the first network device was scattered across GPUs. Rail and plane assignment replaces it. ### Resolved issues +* Fixed a crash in `ncclGetUniqueId()` when `/proc/version` cannot be opened. Initialization now returns `ncclSystemError` instead of calling `strstr` on a null kernel-version string. * Fixed profiler overhead when no profiler plugin is loaded or only kernel-channel events are enabled. * Fixed profiler API events reporting rank 0 instead of the originating communicator rank. * Fixed P2P IPC registration reuse producing out-of-bounds remote addresses when a registered allocation spans multiple cuMem segments. diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index 992ce1c06f7a..02fc7c3f51a4 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -335,9 +335,15 @@ static ncclResult_t ncclInit() { NCCLCHECK(ncclOsTopoGetStrFromSys("/proc", "version", strValue, sizeof(strValue))); char *verStr, *state; verStr = strtok_r(strValue, " ", &state); - for (int i = 0; i < 2; i++) { + for (int i = 0; i < 2 && verStr != NULL; i++) { verStr = strtok_r(NULL, " ", &state); - if (verStr == NULL) break; + } + // fopen of /proc/version fails (EMFILE, missing /proc) as an empty string, and strtok_r + // then returns NULL. strstr(NULL) is a SIGSEGV; return an error before call_once so a + // later init can retry once descriptors are available. + if (verStr == NULL) { + WARN("Could not read kernel version from /proc/version"); + return ncclSystemError; } INFO(NCCL_INIT, "Kernel version: %s", verStr); if (strstr(verStr, "cray") == NULL) { diff --git a/projects/rccl/test/host/fakes/os_fakes.cc b/projects/rccl/test/host/fakes/os_fakes.cc index 32cef4587386..b3256b18da61 100644 --- a/projects/rccl/test/host/fakes/os_fakes.cc +++ b/projects/rccl/test/host/fakes/os_fakes.cc @@ -60,7 +60,8 @@ ncclResult_t ncclOsSetAffinity(const ncclAffinity& affinity) { return g_ncclOsSetAffinityResult; } -// Must be non-empty and multi-token: ncclInit() strstr()s the strtok_r() of this, and strtok_r("") returns NULL. +// Must be non-empty and multi-token: ncclInit() takes the third strtok_r token of this, and an empty +// string makes that NULL, which ncclInit() rejects with ncclSystemError. // Also not "1" and not the Hyper-V BIOS string, so numa_balancing / bios_version stay on their benign arms. ncclResult_t g_ncclOsTopoGetStrFromSysResult = ncclSuccess; int g_ncclOsTopoGetStrFromSysCalls = 0; diff --git a/projects/rccl/test/host/init-test.cc b/projects/rccl/test/host/init-test.cc index 3b216605bf29..2d739443e8e6 100644 --- a/projects/rccl/test/host/init-test.cc +++ b/projects/rccl/test/host/init-test.cc @@ -4537,11 +4537,36 @@ TEST_F(InitMicrotestIsolated, NcclInit_NumaBalancingOffAndIommuPassthrough_Warns ASSERT_TRUE(LogHas(log, "Kernel version: 6.8.0-microtest")) << "actual log:\n" << log; ASSERT_FALSE(LogHas(log, Dtor_kNumaBalancingWarning)) << "actual log:\n" << log; ASSERT_FALSE(LogHas(log, Dtor_kIommuWarning)) << "actual log:\n" << log; - // Anchors the iommu leg: the probe sits inside the non-cray block (init.cc:310-336), the INFO above does not. + // Anchors the iommu leg: the probe sits inside the non-cray block, the INFO above does not. ASSERT_EQ(1, iommu.calls); }); } +// Empty /proc/version (EMFILE, or any read ncclOsTopoGetStrFromSys reports as "") used to reach +// strstr(NULL, "cray"). The error is returned before call_once, so this stays reachable in a fresh process. +TEST_F(InitMicrotestIsolated, NcclInit_UnreadableKernelVersion_ReturnsSystemError) { + RUN_ISOLATED_TEST( + "Init_NcclInit_UnreadableKernelVersion", + []() { + ScopedHook sysText(g_microTopoGetStrFromSys, + [](const char*, const char* file, char* out, int maxLen) { + if (!out || maxLen <= 0) return ncclSuccess; + if (file && std::strcmp(file, "version") == 0) { + out[0] = '\0'; + return ncclSuccess; + } + std::snprintf(out, maxLen, "0"); + return ncclSuccess; + }); + ScopedHook iommu(g_microIommuPassthroughOk, [](const char*) { return false; }); + ncclResult_t res = ncclSuccess; + const std::string log = RcclUnitTesting::CaptureLog([&] { res = ncclInit(); }); + ASSERT_EQ(ncclSystemError, res); + ASSERT_TRUE(LogHas(log, "Could not read kernel version from /proc/version")) << "actual log:\n" << log; + ASSERT_EQ(0, iommu.calls); + }); +} + // Isolated: ncclInitEnv latches envInitOnceFlag, which no later test could then observe unlatched. TEST_F(InitMicrotestIsolated, InitEnv_PluginFails_LatchesThatErrorAndCallsThePluginOnce) { RUN_ISOLATED_TEST( From 486b2a5b24e999802e2258930465597258f2e748 Mon Sep 17 00:00:00 2001 From: Shweta Jha Date: Wed, 7 Oct 2026 11:26:35 -0500 Subject: [PATCH 2/5] fix(rccl): correct comments on the kernel-version null check The previous notes described a failed fopen as an empty string and treated a later retry as general. Co-authored-by: Cursor --- projects/rccl/src/init.cc | 6 +++--- projects/rccl/test/host/fakes/os_fakes.cc | 4 ++-- projects/rccl/test/host/init-test.cc | 4 ++-- 3 files changed, 7 insertions(+), 7 deletions(-) diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index 02fc7c3f51a4..0b21565d313f 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -338,9 +338,9 @@ static ncclResult_t ncclInit() { for (int i = 0; i < 2 && verStr != NULL; i++) { verStr = strtok_r(NULL, " ", &state); } - // fopen of /proc/version fails (EMFILE, missing /proc) as an empty string, and strtok_r - // then returns NULL. strstr(NULL) is a SIGSEGV; return an error before call_once so a - // later init can retry once descriptors are available. + // A failed read is reported as "" and ncclSuccess. Fewer than three tokens also leaves + // verStr NULL, and strstr(NULL) would SIGSEGV. This runs before call_once, so the failure + // is not latched. if (verStr == NULL) { WARN("Could not read kernel version from /proc/version"); return ncclSystemError; diff --git a/projects/rccl/test/host/fakes/os_fakes.cc b/projects/rccl/test/host/fakes/os_fakes.cc index b3256b18da61..6141ac69ff36 100644 --- a/projects/rccl/test/host/fakes/os_fakes.cc +++ b/projects/rccl/test/host/fakes/os_fakes.cc @@ -60,8 +60,8 @@ ncclResult_t ncclOsSetAffinity(const ncclAffinity& affinity) { return g_ncclOsSetAffinityResult; } -// Must be non-empty and multi-token: ncclInit() takes the third strtok_r token of this, and an empty -// string makes that NULL, which ncclInit() rejects with ncclSystemError. +// Needs at least three tokens: ncclInit() keeps the third strtok_r token. Fewer than three, +// including "", makes verStr NULL and ncclInit() returns ncclSystemError. // Also not "1" and not the Hyper-V BIOS string, so numa_balancing / bios_version stay on their benign arms. ncclResult_t g_ncclOsTopoGetStrFromSysResult = ncclSuccess; int g_ncclOsTopoGetStrFromSysCalls = 0; diff --git a/projects/rccl/test/host/init-test.cc b/projects/rccl/test/host/init-test.cc index 2d739443e8e6..b169991431c1 100644 --- a/projects/rccl/test/host/init-test.cc +++ b/projects/rccl/test/host/init-test.cc @@ -4542,8 +4542,8 @@ TEST_F(InitMicrotestIsolated, NcclInit_NumaBalancingOffAndIommuPassthrough_Warns }); } -// Empty /proc/version (EMFILE, or any read ncclOsTopoGetStrFromSys reports as "") used to reach -// strstr(NULL, "cray"). The error is returned before call_once, so this stays reachable in a fresh process. +// An empty read used to reach strstr(NULL, "cray"). The check runs before call_once, so this arm +// is reachable in-process. TEST_F(InitMicrotestIsolated, NcclInit_UnreadableKernelVersion_ReturnsSystemError) { RUN_ISOLATED_TEST( "Init_NcclInit_UnreadableKernelVersion", From 1357c01ccc720b98522b7a25840b92a44d4e645c Mon Sep 17 00:00:00 2001 From: Shweta Jha Date: Thu, 8 Oct 2026 11:35:41 -0500 Subject: [PATCH 3/5] fix(rccl): retry ncclInit after an unreadable kernel version The empty-/proc/version test only checked the first failure, so a latched initResult would still pass. The changelog now names every public caller of ncclInit. Co-authored-by: Cursor --- projects/rccl/CHANGELOG.md | 2 +- projects/rccl/test/host/init-test.cc | 39 +++++++++++++++++----------- 2 files changed, 25 insertions(+), 16 deletions(-) diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index 9cbb8878232a..7136c30bfec9 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -37,7 +37,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * `NCCL_TOPO_SCATTER_START_NET`, which selected how the first network device was scattered across GPUs. Rail and plane assignment replaces it. ### Resolved issues -* Fixed a crash in `ncclGetUniqueId()` when `/proc/version` cannot be opened. Initialization now returns `ncclSystemError` instead of calling `strstr` on a null kernel-version string. +* Fixed a crash in `ncclInit()` when `/proc/version` cannot be opened. `ncclGetUniqueId()`, `ncclCommInitRank()`, `ncclCommInitAll()`, `ncclCommInitRankConfig()`, `ncclCommInitRankScalable()`, and `ncclCommGrow()` now return `ncclSystemError` instead of calling `strstr` on a null kernel-version string. * Fixed profiler overhead when no profiler plugin is loaded or only kernel-channel events are enabled. * Fixed profiler API events reporting rank 0 instead of the originating communicator rank. * Fixed P2P IPC registration reuse producing out-of-bounds remote addresses when a registered allocation spans multiple cuMem segments. diff --git a/projects/rccl/test/host/init-test.cc b/projects/rccl/test/host/init-test.cc index b169991431c1..05bdb3fd3983 100644 --- a/projects/rccl/test/host/init-test.cc +++ b/projects/rccl/test/host/init-test.cc @@ -4542,28 +4542,37 @@ TEST_F(InitMicrotestIsolated, NcclInit_NumaBalancingOffAndIommuPassthrough_Warns }); } -// An empty read used to reach strstr(NULL, "cray"). The check runs before call_once, so this arm -// is reachable in-process. +// An empty read used to reach strstr(NULL, "cray"). The check runs before call_once, so a later +// ncclInit() in the same process can still succeed once the read works. TEST_F(InitMicrotestIsolated, NcclInit_UnreadableKernelVersion_ReturnsSystemError) { RUN_ISOLATED_TEST( "Init_NcclInit_UnreadableKernelVersion", []() { + ncclResult_t res = ncclSuccess; + std::string log; + { + ScopedHook sysText(g_microTopoGetStrFromSys, + [](const char*, const char* file, char* out, int maxLen) { + if (!out || maxLen <= 0) return ncclSuccess; + if (file && std::strcmp(file, "version") == 0) { + out[0] = '\0'; + return ncclSuccess; + } + return SysFileText(file, "0", out, maxLen); + }); + ScopedHook iommu(g_microIommuPassthroughOk, [](const char*) { return false; }); + log = RcclUnitTesting::CaptureLog([&] { res = ncclInit(); }); + ASSERT_EQ(ncclSystemError, res); + ASSERT_TRUE(LogHas(log, "Could not read kernel version from /proc/version")) << "actual log:\n" << log; + ASSERT_EQ(0, iommu.calls); + } + // The empty-version hook is gone. A regression that stored the error in initResult + // before returning would make this second call fail. ScopedHook sysText(g_microTopoGetStrFromSys, [](const char*, const char* file, char* out, int maxLen) { - if (!out || maxLen <= 0) return ncclSuccess; - if (file && std::strcmp(file, "version") == 0) { - out[0] = '\0'; - return ncclSuccess; - } - std::snprintf(out, maxLen, "0"); - return ncclSuccess; + return SysFileText(file, "0", out, maxLen); }); - ScopedHook iommu(g_microIommuPassthroughOk, [](const char*) { return false; }); - ncclResult_t res = ncclSuccess; - const std::string log = RcclUnitTesting::CaptureLog([&] { res = ncclInit(); }); - ASSERT_EQ(ncclSystemError, res); - ASSERT_TRUE(LogHas(log, "Could not read kernel version from /proc/version")) << "actual log:\n" << log; - ASSERT_EQ(0, iommu.calls); + ASSERT_EQ(ncclSuccess, ncclInit()); }); } From 4bf063c203a4d093e77f2fddd296762387d354ea Mon Sep 17 00:00:00 2001 From: Shweta Jha Date: Thu, 8 Oct 2026 14:46:53 -0500 Subject: [PATCH 4/5] fix(rccl): hook cpuid and iommu around both ncclInit calls The retry call was outside those hooks, so it read the host /proc/cmdline. Both calls now use the same fakes as the neighboring init tests. Co-authored-by: Cursor --- projects/rccl/test/host/init-test.cc | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/projects/rccl/test/host/init-test.cc b/projects/rccl/test/host/init-test.cc index 05bdb3fd3983..51b297b08854 100644 --- a/projects/rccl/test/host/init-test.cc +++ b/projects/rccl/test/host/init-test.cc @@ -4550,6 +4550,8 @@ TEST_F(InitMicrotestIsolated, NcclInit_UnreadableKernelVersion_ReturnsSystemErro []() { ncclResult_t res = ncclSuccess; std::string log; + ScopedHook cpuid(g_microCpuid, CpuidWithoutHypervisorBit); + ScopedHook iommu(g_microIommuPassthroughOk, [](const char*) { return false; }); { ScopedHook sysText(g_microTopoGetStrFromSys, [](const char*, const char* file, char* out, int maxLen) { @@ -4560,19 +4562,20 @@ TEST_F(InitMicrotestIsolated, NcclInit_UnreadableKernelVersion_ReturnsSystemErro } return SysFileText(file, "0", out, maxLen); }); - ScopedHook iommu(g_microIommuPassthroughOk, [](const char*) { return false; }); log = RcclUnitTesting::CaptureLog([&] { res = ncclInit(); }); ASSERT_EQ(ncclSystemError, res); ASSERT_TRUE(LogHas(log, "Could not read kernel version from /proc/version")) << "actual log:\n" << log; ASSERT_EQ(0, iommu.calls); } // The empty-version hook is gone. A regression that stored the error in initResult - // before returning would make this second call fail. + // before returning would make this second call fail. cpuid and iommu stay hooked so + // this call does not depend on the host. ScopedHook sysText(g_microTopoGetStrFromSys, [](const char*, const char* file, char* out, int maxLen) { return SysFileText(file, "0", out, maxLen); }); ASSERT_EQ(ncclSuccess, ncclInit()); + ASSERT_EQ(1, iommu.calls); }); } From c75c16eae704769f5c6e27a223101b7218cffcff Mon Sep 17 00:00:00 2001 From: Shweta Jha Date: Wed, 21 Oct 2026 15:05:17 -0500 Subject: [PATCH 5/5] fix(rccl): correct the unlatched-init test comment The second ncclInit() shows initOnceFlag was left unlatched. NCCLCHECKGOTO overwrites initResult, so that call does not catch a poisoned error. Co-authored-by: Cursor --- projects/rccl/test/host/init-test.cc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/rccl/test/host/init-test.cc b/projects/rccl/test/host/init-test.cc index 51b297b08854..d376d5b8aa88 100644 --- a/projects/rccl/test/host/init-test.cc +++ b/projects/rccl/test/host/init-test.cc @@ -4567,9 +4567,9 @@ TEST_F(InitMicrotestIsolated, NcclInit_UnreadableKernelVersion_ReturnsSystemErro ASSERT_TRUE(LogHas(log, "Could not read kernel version from /proc/version")) << "actual log:\n" << log; ASSERT_EQ(0, iommu.calls); } - // The empty-version hook is gone. A regression that stored the error in initResult - // before returning would make this second call fail. cpuid and iommu stay hooked so - // this call does not depend on the host. + // The empty-version hook is gone. This second call pins that the early return left + // initOnceFlag unlatched, so a later ncclInit() can still initialise. cpuid and iommu + // stay hooked so this call does not depend on the host. ScopedHook sysText(g_microTopoGetStrFromSys, [](const char*, const char* file, char* out, int maxLen) { return SysFileText(file, "0", out, maxLen);