Added ERROR message class to handle fatal error messages. (#2002)

* Added ERROR message class to handle fatal error messages.

New ERROR message class will print the message in all debug level,
including none.

Change some of the fatal error message to be in ERROR instead of WARN.

Added new error handler function to print out more meaningful error
message in the future.

* Added CHANGELOG entry.

* Update CHANGELOG.md

Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>

* Change to no longer reuse NONE as ERROR. ERROR is now a separated class.

* Update CHANGELOG.md

Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>

---------

Co-authored-by: Jeffrey Novotny <jnovotny@amd.com>

src/debug.cc

@@ -47,6 +47,8 @@ static void ncclDebugInit() {
     ncclDebugFile = stdout;
   }
   if (nccl_debug == NULL) {
+    tempNcclDebugLevel = NCCL_LOG_ERROR;
+  } else if (strcasecmp(nccl_debug, "NONE") == 0) {
     tempNcclDebugLevel = NCCL_LOG_NONE;
   } else if (strcasecmp(nccl_debug, "VERSION") == 0) {
     tempNcclDebugLevel = NCCL_LOG_VERSION;
@@ -372,6 +374,8 @@ void ncclDebugLog(ncclDebugLogLevel level, unsigned long flags, const char *file
     auto delta = std::chrono::steady_clock::now() - ncclEpoch;
     double timestamp = std::chrono::duration_cast<std::chrono::duration<double>>(delta).count()*1000;
     len += snprintf(buffer+len, sizeof(buffer)-len, "[%d] %f %s:%d NCCL TRACE ", cudaDev, timestamp, filefunc, line);
+  } else if (level == NCCL_LOG_ERROR) {
+    len += snprintf(buffer+len, sizeof(buffer)-len, "[%d] [FATAL ERROR]: ", cudaDev);
   }
   len = std::min(len, sizeof(buffer)-1);  // prevent overflows
 

src/include/checks.h

@@ -13,16 +13,14 @@
 #define CUDACHECK(cmd) do {                                 \
     cudaError_t err = cmd;                                  \
     if( err != cudaSuccess ) {                              \
-        WARN("Cuda failure '%s'", cudaGetErrorString(err)); \
-        return ncclUnhandledCudaError;                      \
+        return rcclCudaErrorHandler(err);                   \
     }                                                       \
 } while(false)
 
 #define CUDACHECKGOTO(cmd, RES, label) do {                 \
     cudaError_t err = cmd;                                  \
     if( err != cudaSuccess ) {                              \
-        WARN("Cuda failure '%s'", cudaGetErrorString(err)); \
-        RES = ncclUnhandledCudaError;                       \
+        RES = rcclCudaErrorHandler(err);                    \
         goto label;                                         \
     }                                                       \
 } while(false)

src/include/debug.h

@@ -25,6 +25,7 @@ void ncclDebugLog(ncclDebugLogLevel level, unsigned long flags, const char *file
 extern thread_local int ncclDebugNoWarn;
 extern char ncclLastError[];
 
+#define ERROR(...) ncclDebugLog(NCCL_LOG_ERROR, NCCL_ALL, __FILE__, __LINE__, __VA_ARGS__)
 #define VERSION(...) ncclDebugLog(NCCL_LOG_VERSION, NCCL_ALL, __FILE__, __LINE__, __VA_ARGS__)
 #define WARN(...) ncclDebugLog(NCCL_LOG_WARN, NCCL_ALL, __FILE__, __LINE__, __VA_ARGS__)
 #define INFO(FLAGS, ...) ncclDebugLog(NCCL_LOG_INFO, (FLAGS), __func__, __LINE__, __VA_ARGS__)
@@ -40,4 +41,22 @@ void ncclSetThreadName(pthread_t thread, const char *fmt, ...);
 
 void ncclResetDebugInit();
 
+// RCCL custom error message handling.
+static inline ncclResult_t rcclCudaErrorHandler(cudaError_t err) {
+
+    // Print the cuda error
+    ERROR("HIP failure: '%s'", cudaGetErrorString(err));
+
+    // Special error message here:
+    switch (err) {
+    case cudaErrorStreamCaptureInvalidated:
+	    ERROR("Application is trying to use an invalidated stream to launch RCCL kernel. "
+	          "This operation is invalid. RCCL is exiting.");
+	    break;
+    default:
+	    break;
+    }
+    return ncclUnhandledCudaError;
+}
+
 #endif

src/include/nccl_common.h

@@ -12,11 +12,12 @@
 
 typedef enum {
   NCCL_LOG_NONE = 0,
-  NCCL_LOG_VERSION = 1,
-  NCCL_LOG_WARN = 2,
-  NCCL_LOG_INFO = 3,
-  NCCL_LOG_ABORT = 4,
-  NCCL_LOG_TRACE = 5
+  NCCL_LOG_ERROR = 1,
+  NCCL_LOG_VERSION = 2,
+  NCCL_LOG_WARN = 3,
+  NCCL_LOG_INFO = 4,
+  NCCL_LOG_ABORT = 5,
+  NCCL_LOG_TRACE = 6
 } ncclDebugLogLevel;
 
 typedef enum {

src/init.cc

@@ -149,7 +149,8 @@ ncclResult_t checkHsaEnvSetting() {
   INFO(NCCL_INIT, "Hipruntime version: %d, firmware version: %d", hipRuntimeVersion, firmwareVersion);
   if (!validHsaScratchEnvSetting(hsaScratchEnv, hipRuntimeVersion, firmwareVersion, devProp.gcnArchName)) {
     // Always print out this warning message
-    printf("Fatal Error: HSA_NO_SCRATCH_RECLAIM=1 must be set to avoid performance degradation with HIP Runtime version:%d, GPU Firmware version:%d\n", hipRuntimeVersion, firmwareVersion);
+    ERROR("HSA_NO_SCRATCH_RECLAIM=1 must be set to avoid performance degradation with the current HIP configuration. (Runtime version:%d, GPU Firmware version:%d)", hipRuntimeVersion, firmwareVersion);
+    ERROR("Please set HSA_NO_SCRATCH_RECLAIM=1 and rerun.");
     return ncclSystemError;
   }
   return ncclSuccess;
@@ -2433,7 +2434,7 @@ static ncclResult_t ncclCommInitRankDev(ncclComm_t* newcomm, int nranks, int nId
   // first call ncclInit, this will setup the environment
   NCCLCHECKGOTO(ncclInit(), res, fail);
 
-  if (ncclDebugLevel > NCCL_LOG_WARN || (ncclDebugLevel != NCCL_LOG_NONE && myrank == 0)) {
+  if (ncclDebugLevel > NCCL_LOG_WARN || (ncclDebugLevel >= NCCL_LOG_VERSION && myrank == 0)) {
     static pthread_once_t once = PTHREAD_ONCE_INIT;
     pthread_once(&once, showVersion);
   }

src/misc/strongstream.cc

@@ -259,7 +259,7 @@ ncclResult_t ncclStrongStreamRelease(
           CUDACHECK(cudaEventRecord(ss->serialEvent, ss->liveStream));
         }
         if (ss->liveAcquiredBy != localThreadId() && ncclParamLaunchRaceFatal()) {
-          WARN("%s", launchRaceFatalMsg);
+          ERROR("%s", launchRaceFatalMsg);
           return ncclInvalidUsage;
         }
       } else {
@@ -321,7 +321,7 @@ ncclResult_t ncclStrongStreamRelease(
         #endif
 
         if (cap->acquiredBy != localThreadId() && ncclParamLaunchRaceFatal()) {
-          WARN("%s", launchRaceFatalMsg);
+          ERROR("%s", launchRaceFatalMsg);
           return ncclInvalidUsage;
         }
       }

src/transport/net_ib.cc

@@ -139,7 +139,8 @@ static void ncclIbStatsFatalError(struct ncclIbStats* stat){
 }
 static ncclResult_t ncclIbStatsCheckFatalCount(struct ncclIbStats* stat, const char* funcName) {
   if (ncclParamIbAsyncEvents() && __atomic_load_n(&stat->fatalErrorCount, __ATOMIC_RELAXED)) {
-    WARN("communicator encountered a fatal error (detected in %s)\n", funcName);
+    ERROR("RCCL encountered a communication fatal error (detected in %s)\n", funcName);
+    ERROR("RCCL cannot recover from this network failure and now exiting. Please check the network health.");
     return ncclSystemError;
   }
   return ncclSuccess;