JhaShweta1 · Oct 7, 2026 · Oct 7, 2026 · Oct 8, 2026 · Oct 8, 2026 · Oct 8, 2026
@@ -66,6 +66,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https:
 
 ### Resolved issues
 * Fixed topology GPU indices that followed the HIP device order instead of the PCI order. This reordered rings and NIC selection, causing an `ncclGather` performance regression on multi-node MI355X systems, and made communicator initialization fail when each rank sees only its own GPU.
+* 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.
 * Anvil SDMA (`NCCL_GIN_TYPE=7`) honours strong GIN signals. A strong signal drains every SDMA queue to that peer before it is sent, including puts that are no longer marked dirty. When the signal is fused with its SDMA put (gfx950), the drain covers only the earlier queue entries and the caller does not wait for the fused put itself. Legacy `ncclGin_SignalInc` is strong by default, so a caller that only needs to order the put it just issued should use `ncclGin_WeakSignalInc`. A barrier on this backend still flushes once and sends weak signals.
 * Fixed VA GIN signals (`ncclGin_VASignalInc`, `ncclGin_VASignalAdd` and their strong and weak forms) on Anvil SDMA (`NCCL_GIN_TYPE=7`). The backend read the VA descriptor as an indexed signal ID and incremented an unrelated signal cell, so the VA signal never arrived. VA signals now update the window cell on the peer, follow the same strong, weak and fused-signal rules as indexed signals, and `ncclGin::resetSignal` on a VA signal clears the local cell.
 * Fixed profiler overhead when no profiler plugin is loaded or only kernel-channel events are enabled.

@@ -357,9 +357,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;
+  }
+  // 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;
   }
   INFO(NCCL_INIT, "Kernel version: %s", verStr);
   if (strstr(verStr, "cray") == NULL) {

@@ -66,7 +66,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.
+// 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;

@@ -4970,7 +4970,44 @@ 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);
+      });
+}
+
+// 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 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) {
+                               if (!out || maxLen <= 0) return ncclSuccess;
+                               if (file && std::strcmp(file, "version") == 0) {
+                                 out[0] = '\0';
+                                 return ncclSuccess;
+                               }
+                               return SysFileText(file, "0", out, maxLen);
+                             });
+          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. 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);
+                           });
+        ASSERT_EQ(ncclSuccess, ncclInit());
         ASSERT_EQ(1, iommu.calls);
       });
 }