diff --git a/makefiles/version.mk b/makefiles/version.mk
index 5ba0f0719c..0c0a902ec7 100644
--- a/makefiles/version.mk
+++ b/makefiles/version.mk
@@ -2,5 +2,5 @@
 NCCL_MAJOR   := 2
 NCCL_MINOR   := 21
 NCCL_PATCH   := 5
-NCCL_SUFFIX  := -T3
+NCCL_SUFFIX  := -T4
 PKG_REVISION := 1
diff --git a/src/debug.cc b/src/debug.cc
index 522999b44f..171dd42eb9 100644
--- a/src/debug.cc
+++ b/src/debug.cc
@@ -175,17 +175,18 @@ void ncclDebugLog(ncclDebugLogLevel level, unsigned long flags, const char *file
 
   char buffer[1024];
   size_t len = 0;
+
+  auto delta = std::chrono::steady_clock::now() - ncclEpoch;
+  double timestamp = std::chrono::duration_cast<std::chrono::duration<double>>(delta).count()*1000;
   if (level == NCCL_LOG_WARN) {
-    len = snprintf(buffer, sizeof(buffer), "\n%s:%d:%d [%d] %s:%d NCCL WARN ",
-                   hostname, pid, tid, cudaDev, filefunc, line);
+    len = snprintf(buffer, sizeof(buffer), "\n%s:%d:%d [%d] %f %s:%d NCCL WARN ",
+                   hostname, pid, tid, cudaDev, timestamp, filefunc, line);
     if (ncclParamWarnSetDebugInfo()) ncclDebugLevel = NCCL_LOG_INFO;
   } else if (level == NCCL_LOG_INFO) {
-    len = snprintf(buffer, sizeof(buffer), "%s:%d:%d [%d] NCCL INFO ", hostname, pid, tid, cudaDev);
+    len = snprintf(buffer, sizeof(buffer), "%s:%d:%d [%d] %f NCCL INFO ", hostname, pid, tid, cudaDev, timestamp);
   } else if (level == NCCL_LOG_TRACE && flags == NCCL_CALL) {
-    len = snprintf(buffer, sizeof(buffer), "%s:%d:%d NCCL CALL ", hostname, pid, tid);
+    len = snprintf(buffer, sizeof(buffer), "%s:%d:%d %f NCCL CALL ", hostname, pid, tid, timestamp);
   } else if (level == NCCL_LOG_TRACE) {
-    auto delta = std::chrono::steady_clock::now() - ncclEpoch;
-    double timestamp = std::chrono::duration_cast<std::chrono::duration<double>>(delta).count()*1000;
     len = snprintf(buffer, sizeof(buffer), "%s:%d:%d [%d] %f %s:%d NCCL TRACE ",
                    hostname, pid, tid, cudaDev, timestamp, filefunc, line);
   }