From 51f87fbb438b8eb3a5c6a4cdfaad0d70f724578b Mon Sep 17 00:00:00 2001
From: Dingming Wu <dmwu0506@gmail.com>
Date: Wed, 14 May 2025 08:12:45 -0700
Subject: [PATCH] Detect if HSA_NO_SCRATCH_RECLAIM is set after initEnv()
 (#1683)

* Detect if HSA_NO_SCRATCH_RECLAIM is set after initEnv()

 For rocm older than 6.4, we need to set HSA_NO_SCRATCH_RECLAIM=1 to use LL128 protocol.
This Env is set outside of RCCL, add the logging to detect whether its set during runtime.

* check hip runtime ver via hipRuntimeGetVersion

* move the detection to ncclinit func

* correct rocm version integer

* update warning message

* avoid unnecessary info msg on hsa_no_scratch_reclaim detection
---
 src/init.cc | 7 +++++++
 1 file changed, 7 insertions(+)

diff --git a/src/init.cc b/src/init.cc
index c1e1403f0e..9feaeadac8 100644
--- a/src/init.cc
+++ b/src/init.cc
@@ -186,6 +186,13 @@ static ncclResult_t ncclInit() {
         WARN("Missing \"HSA_FORCE_FINE_GRAIN_PCIE=1\" from environment which can lead to low RCCL performance, system instablity or hang!");
 #endif
     }
+  const char* hsaScratchEnv = getenv("HSA_NO_SCRATCH_RECLAIM");
+  int hipRuntimeVersion = 0;
+  // hipVer is an integer e.g., 6.2.41133 -> 60241133
+  CUDACHECK(hipRuntimeGetVersion(&hipRuntimeVersion));
+  if ((!hsaScratchEnv || strcmp(hsaScratchEnv,"1") != 0) && hipRuntimeVersion < 60400000){
+    WARN("HSA_NO_SCRATCH_RECLAIM=1 must be set to avoid RCCL perf hit for rocm older than 6.4,, rocm ver:%d", hipRuntimeVersion);
+  }
   pthread_once(&initOnceControl, initOnceFunc);
   return initResult;
 }