Add tuning report option

Add -U/--tuning to report the implementation, algorithm, protocol,
kernel variant, synchronization strategy, channel and warp counts,
and NET proxy chunk size in text and JSON output.

Embed the published profiler v5, v6, and v7 ABI headers and export all
three interfaces so NCCL releases from 2.28 onward can resolve a
compatible profiler ABI. Report v5/v6 collectives as GEN and use v7
isSymColl metadata for symmetric detail with NCCL 2.31 and later.

Share handles for concurrent group, task, and proxy events, reset
captured state after every result, and avoid reporting stale P2P and
proxy values. This also corrects proxy chunk sizes across tuning sweeps.

Ignore -U with a notice on pre-2.28 libraries, and warn when it
overrides a user-selected profiler plugin. Export profiler symbols from
Makefile and CMake builds, align the expanded tuning columns, and use
the statement-style PRINT macro so unsupported fields are reported as
N/A. Use JSON schema version 4.

Signed-off-by: David Addison <[email protected]>

README.md

@@ -103,6 +103,7 @@ Collective perf tests support the same set of arguments :
   * `-T,--timeout <time in seconds>` timeout each test after specified number of seconds. Default : disabled.
   * `-M,--memory <0/1>` enable memory usage report. Default : 0.
   * `-u,--unalign <index of first element>` Misalign source and destination buffers. Default : 0.
+  * `-U,--tuning <0/1>` report NCCL tuning information. Requires NCCL 2.28 or newer; symmetric collective identification and kernel variant reporting require NCCL 2.31 or newer. Default : 0.
 
 ### Running multiple operations in parallel
 

src/CMakeLists.txt

@@ -53,6 +53,7 @@ foreach(COLL IN LISTS COLLS)
   add_executable(${COLL}_perf ${COLL}.cu ${PERF_COMMON_SOURCES})
   add_dependencies(${COLL}_perf generate_git_version)
   set_target_properties(${COLL}_perf PROPERTIES
+    ENABLE_EXPORTS TRUE
     RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}
   )
 

src/Makefile

@@ -111,7 +111,7 @@ ${DST_DIR}/%$(NAME_SUFFIX).o: %.cu common.h util.h $(TEST_VERIFIABLE_HDRS)
 	@mkdir -p ${DST_DIR}
 	$(NVCC) -o $@ $(NVCUFLAGS) -c $<
 
-${DST_DIR}/util$(NAME_SUFFIX).o: $(GIT_VERSION_FILE)
+${DST_DIR}/util$(NAME_SUFFIX).o: $(GIT_VERSION_FILE) nccl_profiler.h profiler_v5.h profiler_v6.h profiler_v7.h
 
 ${DST_DIR}/timer.o: timer.cc timer.h
 	@printf "Compiling  %-35s > %s\n" $< $@
@@ -122,12 +122,12 @@ ifeq ($(DSO), 1)
 ${DST_DIR}/%_perf$(NAME_SUFFIX): ${DST_DIR}/%.o ${DST_DIR}/common$(NAME_SUFFIX).o ${DST_DIR}/util$(NAME_SUFFIX).o ${DST_DIR}/timer.o $(TEST_VERIFIABLE_LIBS) $(TEST_OS_LIB)
 	@printf "Linking  %-35s > %s\n" $< $@
 	@mkdir -p ${DST_DIR}
-	$(NVCC) -o $@ $(NVCUFLAGS) $^ -L$(TEST_VERIFIABLE_BUILDDIR) -lverifiable ${NVLDFLAGS} -Xlinker "--enable-new-dtags" -Xlinker "-rpath,\$$ORIGIN:\$$ORIGIN/verifiable"
+	$(NVCC) -o $@ $(NVCUFLAGS) $^ -L$(TEST_VERIFIABLE_BUILDDIR) -lverifiable ${NVLDFLAGS} -Xlinker "--enable-new-dtags" -Xlinker "-rpath,\$$ORIGIN:\$$ORIGIN/verifiable" -Xlinker --export-dynamic
 else
 ${DST_DIR}/%_perf$(NAME_SUFFIX):${DST_DIR}/%.o ${DST_DIR}/common$(NAME_SUFFIX).o ${DST_DIR}/util$(NAME_SUFFIX).o ${DST_DIR}/timer.o $(TEST_VERIFIABLE_OBJS) $(TEST_OS_LIB)
 	@printf "Linking  %-35s > %s\n" $< $@
 	@mkdir -p ${DST_DIR}
-	$(NVCC) -o $@ $(NVCUFLAGS) $^ ${NVLDFLAGS}
+	$(NVCC) -o $@ $(NVCUFLAGS) $^ ${NVLDFLAGS} -Xlinker --export-dynamic
 endif
 
 ifneq ($(MPI), 1)

src/alltoallv.cu

@@ -122,7 +122,10 @@
 #include <stdint.h>
 #include <stdlib.h>
 
-#define PRINT if (is_main_thread) printf
+#define PRINT(...) \
+  do { \
+    if (is_main_thread) printf(__VA_ARGS__); \
+  } while (0)
 
 #define MAX(a,b) ((a) > (b) ? (a) : (b))
 

src/common.cu

@@ -113,6 +113,7 @@ static int deviceImpl = 0;
 static int hostRmaImpl = 0;
 int unalign = 0;
 int memory_report = 0;
+int tuning = 0;
 
 int deviceCtaCount = 16; // Default number of CTAs for device implementation
 int rmaCtxCount = 1;     // Number of RMA contexts to provision for host RMA (-H)
@@ -1221,13 +1222,14 @@ int main(int argc, char* argv[], char **envp) {
     {"per_iter_timing", required_argument, 0, 'I'},
     {"per_iter_skip", required_argument, 0, 'K'},
     {"host_rma_implementation", required_argument, 0, 'H'},
+    {"tuning", required_argument, 0, 'U'},
     {"help", no_argument, 0, 'h'},
     {}
   };
 
   while(1) {
     int c;
-    c = getopt_long(argc, argv, "t:g:b:e:i:f:n:m:w:N:p:I:K:c:o:d:r:z:y:T:hG:C:a:R:x:D:V:J:S:M:u:H:", longopts, &longindex);
+    c = getopt_long(argc, argv, "t:g:b:e:i:f:n:m:w:N:p:I:K:c:o:d:r:z:y:T:hG:C:a:R:x:D:V:J:S:M:u:H:U:", longopts, &longindex);
 
     if (c == -1)
       break;
@@ -1345,6 +1347,13 @@ int main(int argc, char* argv[], char **envp) {
       case 'K':
         per_iter_skip = (int)strtol(optarg, NULL, 0);
         break;
+      case 'U':
+        tuning = (int)strtol(optarg, NULL, 0);
+        if (tuning && test_ncclVersion < NCCL_VERSION(2,28,0)) {
+          printf("Option -U (tuning) is not supported before NCCL 2.28. Ignoring\n");
+          tuning = 0;
+        }
+        break;
       case 'x':
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,27,0)
         ctaPolicy = (int)strtol(optarg, NULL, 0);
@@ -1445,6 +1454,7 @@ int main(int argc, char* argv[], char **envp) {
             "    i_p99 uses nearest-rank percentile and may equal i_max\n\t"
             "    with <100 samples) (default: 0)] \n\t"
             "[-K,--per_iter_skip <count> exclude leading samples from -I summary stats (default: 0)] \n\t"
+            "[-U,--tuning <0/1> report NCCL tuning info (NCCL >= 2.28; full detail >= 2.31) (default: 0)] \n\t"
             "[-h,--help]\n",
           programName);
         return 0;
@@ -1498,6 +1508,18 @@ int main(int argc, char* argv[], char **envp) {
     per_iter_skip = 0;
   }
 
+  if (tuning) {
+    const char* profilerPlugin = getenv("NCCL_PROFILER_PLUGIN");
+    if (profilerPlugin && strcmp(profilerPlugin, "STATIC_PLUGIN") != 0) {
+      fprintf(stderr, "Option -U overrides NCCL_PROFILER_PLUGIN=%s with STATIC_PLUGIN\n", profilerPlugin);
+    }
+#if defined(NCCL_OS_WINDOWS)
+    _putenv_s("NCCL_PROFILER_PLUGIN", "STATIC_PLUGIN");
+#else
+    setenv("NCCL_PROFILER_PLUGIN", "STATIC_PLUGIN", 1);
+#endif
+  }
+
 #ifdef MPI_SUPPORT
   MPI_Init(&argc, &argv);
 #endif

src/nccl_profiler.h

@@ -0,0 +1,86 @@
+/*************************************************************************
+ * Copyright (c) 2024-2026, NVIDIA CORPORATION. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#ifndef NCCL_PROFILER_H_
+#define NCCL_PROFILER_H_
+
+#include "nccl.h"
+#include <stdint.h>
+#include <stdlib.h>
+
+#if defined(NCCL_OS_WINDOWS)
+typedef unsigned long pid_t;
+#else
+#include <sys/types.h>
+#endif
+
+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} ncclDebugLogLevel;
+typedef void (*ncclDebugLogger_t)(ncclDebugLogLevel level, unsigned long flags, const char *file, int line, const char *fmt, ...);
+
+enum {
+  ncclProfileGroup          = (1 << 0),  // group event type
+  ncclProfileColl           = (1 << 1),  // host collective call event type
+  ncclProfileP2p            = (1 << 2),  // host point-to-point call event type
+  ncclProfileProxyOp        = (1 << 3),  // proxy operation event type
+  ncclProfileProxyStep      = (1 << 4),  // proxy step event type
+  ncclProfileProxyCtrl      = (1 << 5),  // proxy control event type
+  ncclProfileKernelCh       = (1 << 6),  // kernel channel event type
+  ncclProfileNetPlugin      = (1 << 7),  // network plugin-defined events
+  ncclProfileGroupApi       = (1 << 8),  // group API events
+  ncclProfileCollApi        = (1 << 9),  // collective API events
+  ncclProfileP2pApi         = (1 << 10), // point-to-point API events
+  ncclProfileKernelLaunch   = (1 << 11), // kernel launch events
+  ncclProfileCeColl         = (1 << 12), // CE collective operation
+  ncclProfileCeSync         = (1 << 13), // CE synchronization operation
+  ncclProfileCeBatch        = (1 << 14), // CE batch operation
+  ncclProfileKernelPhase    = (1 << 15), // kernel barrier phase sub-event
+};
+
+typedef enum {
+  ncclProfilerProxyOpSendPosted        = 0,
+  ncclProfilerProxyOpSendRemFifoWait   = 1,
+  ncclProfilerProxyOpSendTransmitted   = 2,
+  ncclProfilerProxyOpSendDone          = 3,
+  ncclProfilerProxyOpRecvPosted        = 4,
+  ncclProfilerProxyOpRecvReceived      = 5,
+  ncclProfilerProxyOpRecvTransmitted   = 6,
+  ncclProfilerProxyOpRecvDone          = 7,
+  ncclProfilerProxyStepSendGPUWait     = 8,
+  ncclProfilerProxyStepSendWait        = 9,
+  ncclProfilerProxyStepRecvWait        = 10,
+  ncclProfilerProxyStepRecvFlushWait   = 11,
+  ncclProfilerProxyStepRecvGPUWait     = 12,
+  ncclProfilerProxyCtrlIdle            = 13,
+  ncclProfilerProxyCtrlActive          = 14,
+  ncclProfilerProxyCtrlSleep           = 15,
+  ncclProfilerProxyCtrlWakeup          = 16,
+  ncclProfilerProxyCtrlAppend          = 17,
+  ncclProfilerProxyCtrlAppendEnd       = 18,
+  ncclProfilerProxyOpInProgress_v4     = 19,
+  ncclProfilerProxyStepSendPeerWait_v4 = 20,
+  ncclProfilerNetPluginUpdate          = 21,
+  ncclProfilerKernelChStop             = 22,
+  ncclProfilerGroupStartApiStop        = 23,
+  ncclProfilerGroupEndApiStart         = 24,
+  ncclProfilerCeCollStart               = 25,
+  ncclProfilerCeCollComplete            = 26,
+  ncclProfilerCeSyncStart               = 27,
+  ncclProfilerCeSyncComplete            = 28,
+  ncclProfilerCeBatchStart              = 29,
+  ncclProfilerCeBatchComplete           = 30,
+  ncclProfilerKernelPhaseStop           = 31,
+} ncclProfilerEventState_t;
+
+typedef ncclProfilerEventState_t ncclProfilerEventState_v5_t;
+typedef ncclProfilerEventState_t ncclProfilerEventState_v6_t;
+typedef ncclProfilerEventState_t ncclProfilerEventState_v7_t;
+
+// Copied verbatim from the published NCCL profiler example headers.
+#include "profiler_v5.h"
+#include "profiler_v6.h"
+#include "profiler_v7.h"
+
+#endif

src/profiler_v5.h

@@ -0,0 +1,153 @@
+/*************************************************************************
+ * SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * See LICENSE.txt for more license information
+ *************************************************************************/
+
+#ifndef PROFILER_V5_H_
+#define PROFILER_V5_H_
+#include <stdbool.h>
+
+typedef struct {
+  uint64_t type;                // event type descriptor: ncclProfileGroupApi, ...
+  void* parentObj;              // pointer to the profiler parent object
+  int rank;                     // originating rank
+  union {
+    struct {
+      bool graphCaptured;
+      int groupDepth;
+    } groupApi;
+
+    struct {
+      const char* func;
+      size_t count;
+      const char* datatype;
+      int root;
+      void* stream;
+      bool graphCaptured;
+    } collApi;
+
+    struct {
+      const char* func;
+      size_t count;
+      const char* datatype;
+      void* stream;
+      bool graphCaptured;
+    } p2pApi;
+
+    struct {
+      void* stream;
+    } kernelLaunch;
+
+    struct {
+      uint64_t seqNumber;
+      const char* func;
+      void const* sendBuff;
+      void* recvBuff;
+      size_t count;
+      int root;
+      const char* datatype;
+      uint8_t nChannels;
+      uint8_t nWarps;
+      const char* algo;
+      const char* proto;
+      void* parentGroup; // for backward compatibility with v4
+    } coll;
+
+    struct {
+      const char* func;
+      void* buff;
+      const char* datatype;
+      size_t count;
+      int peer;
+      uint8_t nChannels;
+      void* parentGroup; // for backward compatibility with v4
+    } p2p;
+
+    struct {
+      pid_t pid;                // pid of the originating process
+      uint8_t channelId;        // channel id for this proxy operation
+      int peer;                 // remote rank for send/recv
+      int nSteps;               // number of steps for this proxy operation
+      int chunkSize;            // amount of data transferred by this proxy operation
+      int isSend;
+    } proxyOp;
+
+    struct {
+      int step;
+    } proxyStep;
+
+    struct {
+      uint8_t channelId;
+      uint64_t pTimer;          // start timestamp from GPU globaltimer
+    } kernelCh;
+
+    struct {
+      int64_t id;
+      void* data;
+    } netPlugin;
+  };
+} ncclProfilerEventDescr_v5_t;
+
+typedef union {
+  struct {
+    size_t transSize;
+  } proxyStep;
+
+  struct {
+    int appendedProxyOps;
+  } proxyCtrl;
+
+  struct {
+    void* data;
+  } netPlugin;
+
+  struct {
+    uint64_t pTimer;
+  } kernelCh;
+} ncclProfilerEventStateArgs_v5_t;
+
+typedef struct {
+  const char* name;
+
+  // init - initialize the profiler plugin
+  // Input
+  //  - context        : opaque profiler context object for separating profiler behavior across comms
+  //  - commId         : communicator id
+  //  - commName       : user assigned communicator name
+  //  - nNodes         : number of nodes in communicator
+  //  - nranks         : number of ranks in communicator
+  //  - rank           : rank identifier in communicator
+  //  - logfn          : logger function
+  // Output
+  //  - eActivationMask: bitmask of active events set by the plugin
+  ncclResult_t (*init)(void** context, uint64_t commId, int* eActivationMask, const char* commName, int nNodes, int nranks, int rank, ncclDebugLogger_t logfn);
+
+  // startEvent - initialize and start a new event for the supplied event descriptor inside the eventset
+  // Input
+  //  - context: opaque profiler context object
+  //  - eDescr : pointer to ncclProfilerEventDescr_t object
+  // Output
+  //  - eHandle: return event handle for supplied event descriptor object
+  ncclResult_t (*startEvent)(void* context, void** eHandle, ncclProfilerEventDescr_v5_t* eDescr);
+
+  // stopEvent - stop/finalize an event inside and event set
+  // Input
+  //  - eHandle: handle to event object
+  ncclResult_t (*stopEvent)(void* eHandle);
+
+  // recordEventState - record event state transitions and event attribute updates
+  // Input
+  //  - eHandle   : handle to event object created through startEvent
+  //  - eStateArgs: optional argument used to capture event attribute updates associated with the state transition
+  //  - eState    : event state transition
+  ncclResult_t (*recordEventState)(void* eHandle, ncclProfilerEventState_v5_t eState, ncclProfilerEventStateArgs_v5_t* eStateArgs);
+
+  // finalize - finalize the profiler plugin
+  // Input
+  //  - context: opaque profiler context object
+  ncclResult_t (*finalize)(void* context);
+} ncclProfiler_v5_t;
+
+#endif

src/profiler_v6.h

@@ -0,0 +1,147 @@
+/*************************************************************************
+ * SPDX-FileCopyrightText: Copyright (c) 2025-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * See LICENSE.txt for more license information
+ *************************************************************************/
+
+#ifndef PROFILER_V6_H_
+#define PROFILER_V6_H_
+
+#include "profiler_v5.h"
+
+// Extend v5 descriptor with CE-specific fields
+typedef struct {
+  uint64_t type;                // event type descriptor
+  void* parentObj;              // pointer to the profiler parent object
+  int rank;                     // originating rank
+  union {
+    // All v5 descriptors (groupApi, collApi, p2pApi, kernelLaunch, coll, p2p, proxyOp, proxyStep, kernelCh, netPlugin)
+    struct {
+      bool graphCaptured;
+      int groupDepth;
+    } groupApi;
+
+    struct {
+      const char* func;
+      size_t count;
+      const char* datatype;
+      int root;
+      void* stream;
+      bool graphCaptured;
+    } collApi;
+
+    struct {
+      const char* func;
+      size_t count;
+      const char* datatype;
+      void* stream;
+      bool graphCaptured;
+    } p2pApi;
+
+    struct {
+      void* stream;
+    } kernelLaunch;
+
+    struct {
+      uint64_t seqNumber;
+      const char* func;
+      void const* sendBuff;
+      void* recvBuff;
+      size_t count;
+      int root;
+      const char* datatype;
+      uint8_t nChannels;
+      uint8_t nWarps;
+      const char* algo;
+      const char* proto;
+      void* parentGroup;
+    } coll;
+
+    struct {
+      const char* func;
+      void* buff;
+      const char* datatype;
+      size_t count;
+      int peer;
+      uint8_t nChannels;
+      void* parentGroup;
+    } p2p;
+
+    struct {
+      pid_t pid;
+      uint8_t channelId;
+      int peer;
+      int nSteps;
+      int chunkSize;
+      int isSend;
+    } proxyOp;
+
+    struct {
+      int step;
+    } proxyStep;
+
+    struct {
+      uint8_t channelId;
+      uint64_t pTimer;
+    } kernelCh;
+
+    struct {
+      int64_t id;
+      void* data;
+    } netPlugin;
+
+    // v6 CE-specific descriptors
+    struct {
+      uint64_t seqNumber;
+      const char* func;
+      void const* sendBuff;
+      void* recvBuff;
+      size_t count;
+      int root;
+      const char* datatype;
+      const char* syncStrategy;
+      bool intraBatchSync;
+      uint32_t batchSize;
+      uint32_t numBatches;
+      uint32_t ceSeqNum;
+      void* stream;
+    } ceColl;
+
+    struct {
+      bool isComplete;
+      int nRanks;
+    } ceCollSync;
+
+    struct {
+      int numOps;
+      size_t totalBytes;
+      bool useIntraSync;
+    } ceCollBatch;
+  };
+} ncclProfilerEventDescr_v6_t;
+
+// v6 uses same state args as v5 (no CE-specific state args needed)
+// CE events don't use recordEventState - plugin manages all timing internally
+typedef ncclProfilerEventStateArgs_v5_t ncclProfilerEventStateArgs_v6_t;
+
+typedef struct {
+  const char* name;
+
+  // init - initialize the profiler plugin
+  ncclResult_t (*init)(void** context, uint64_t commId, int* eActivationMask, const char* commName, int nNodes, int nranks, int rank, ncclDebugLogger_t logfn);
+
+  // startEvent - initialize and start a new event
+  ncclResult_t (*startEvent)(void* context, void** eHandle, ncclProfilerEventDescr_v6_t* eDescr);
+
+  // stopEvent - stop/finalize an event
+  ncclResult_t (*stopEvent)(void* eHandle);
+
+  // recordEventState - record event state transitions and updates
+  ncclResult_t (*recordEventState)(void* eHandle, ncclProfilerEventState_v6_t eState, ncclProfilerEventStateArgs_v6_t* eStateArgs);
+
+  // finalize - finalize the profiler plugin
+  ncclResult_t (*finalize)(void* context);
+} ncclProfiler_v6_t;
+
+#endif // PROFILER_V6_H_

src/profiler_v7.h

@@ -0,0 +1,159 @@
+/*************************************************************************
+ * SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+ * SPDX-License-Identifier: Apache-2.0
+ *
+ * See LICENSE.txt for more license information
+ *************************************************************************/
+
+#ifndef PROFILER_V7_H_
+#define PROFILER_V7_H_
+
+#include "profiler_v6.h"
+
+// Extend v6 descriptors with a per-call user profiler tag
+typedef struct {
+  uint64_t type;                // event type descriptor
+  void* parentObj;              // pointer to the profiler parent object
+  int rank;                     // originating rank
+  union {
+    struct {
+      bool graphCaptured;
+      int groupDepth;
+    } groupApi;
+
+    struct {
+      const char* func;
+      size_t count;
+      const char* datatype;
+      int root;
+      void* stream;
+      bool graphCaptured;
+      uint64_t userTag;         // v7: per-call user profiler tag (0 == untagged)
+    } collApi;
+
+    struct {
+      const char* func;
+      size_t count;
+      const char* datatype;
+      void* stream;
+      bool graphCaptured;
+      uint64_t userTag;         // v7: per-call user profiler tag (0 == untagged)
+    } p2pApi;
+
+    struct {
+      void* stream;
+    } kernelLaunch;
+
+    struct {
+      uint64_t seqNumber;
+      const char* func;
+      void const* sendBuff;
+      void* recvBuff;
+      size_t count;
+      int root;
+      const char* datatype;
+      uint8_t nChannels;
+      uint8_t nWarps;
+      const char* algo;
+      const char* proto;
+      void* parentGroup;
+      const char* kernelVariant;
+      bool isSymColl;
+      uint64_t userTag;         // v7: per-call user profiler tag (0 == untagged)
+    } coll;
+
+    struct {
+      const char* func;
+      void* buff;
+      const char* datatype;
+      size_t count;
+      int peer;
+      uint8_t nChannels;
+      void* parentGroup;
+      uint64_t userTag;         // v7: per-call user profiler tag (0 == untagged)
+    } p2p;
+
+    struct {
+      pid_t pid;
+      uint8_t channelId;
+      int peer;
+      int nSteps;
+      int chunkSize;
+      int isSend;
+    } proxyOp;
+
+    struct {
+      int step;
+    } proxyStep;
+
+    struct {
+      uint8_t channelId;
+      uint64_t pTimer;
+    } kernelCh;
+
+    struct {
+      int64_t id;
+      void* data;
+    } netPlugin;
+
+    struct {
+      uint64_t seqNumber;
+      const char* func;
+      void const* sendBuff;
+      void* recvBuff;
+      size_t count;
+      int root;
+      const char* datatype;
+      const char* syncStrategy;
+      bool intraBatchSync;
+      uint32_t batchSize;
+      uint32_t numBatches;
+      uint32_t ceSeqNum;
+      void* stream;
+      uint64_t userTag;         // v7: per-call user profiler tag (0 == untagged)
+    } ceColl;
+
+    struct {
+      bool isComplete;
+      int nRanks;
+    } ceCollSync;
+
+    struct {
+      int numOps;
+      size_t totalBytes;
+      bool useIntraSync;
+    } ceCollBatch;
+
+    // v7: kernel barrier phase sub-event
+    struct {
+      uint8_t channelId;
+      uint8_t phaseId;          // 0=initial_sync, 1=compute, 2=final_sync
+      const char* phaseName;
+      uint64_t pTimer;          // start timestamp (GPU globaltimer)
+    } kernelPhase;
+  };
+} ncclProfilerEventDescr_v7_t;
+
+// v7 uses the same state args as v6/v5
+typedef ncclProfilerEventStateArgs_v6_t ncclProfilerEventStateArgs_v7_t;
+
+typedef struct {
+  const char* name;
+
+  // init - initialize the profiler plugin
+  ncclResult_t (*init)(void** context, uint64_t commId, int* eActivationMask, const char* commName, int nNodes, int nranks, int rank, ncclDebugLogger_t logfn);
+
+  // startEvent - initialize and start a new event
+  ncclResult_t (*startEvent)(void* context, void** eHandle, ncclProfilerEventDescr_v7_t* eDescr);
+
+  // stopEvent - stop/finalize an event
+  ncclResult_t (*stopEvent)(void* eHandle);
+
+  // recordEventState - record event state transitions and updates
+  ncclResult_t (*recordEventState)(void* eHandle, ncclProfilerEventState_v7_t eState, ncclProfilerEventStateArgs_v7_t* eStateArgs);
+
+  // finalize - finalize the profiler plugin
+  ncclResult_t (*finalize)(void* context);
+} ncclProfiler_v7_t;
+
+#endif // PROFILER_V7_H_

src/util.cu

@@ -16,6 +16,7 @@
 // except some of the usual ones.
 
 #include "nccl.h"
+#include "nccl_profiler.h"
 #include "util.h"
 #include "nccl_tests_git_version.h"
 #include "os.h"
@@ -27,8 +28,12 @@
 #include <string>
 #include <sstream>
 #include <iomanip>
+#include <atomic>
 
-#define PRINT if (is_main_thread) printf
+#define PRINT(...) \
+  do { \
+    if (is_main_thread) printf(__VA_ARGS__); \
+  } while (0)
 
 extern int nThreads;
 extern int nGpus;
@@ -46,11 +51,52 @@ extern int per_iter_timing;
 extern int per_iter_skip;
 extern int cudaGraphLaunches;
 extern int unalign;
+extern int tuning;
 
 static FILE *json_report_fp;
 static thread_local bool write_json;
 
-#define JSON_FILE_VERSION 3
+struct groupEvent {
+  uint64_t type;
+  std::atomic<int> count;
+};
+
+struct taskEvent {
+  uint64_t type;
+  std::atomic<int> count;
+  std::atomic<bool> ready;
+  union {
+    struct {
+      const char* implementation;
+      const char* algo;
+      const char* proto;
+      const char* kernelVariant;
+      int nChannels;
+      uint8_t nWarps;
+    } coll;
+
+    struct {
+      uint8_t nChannels;
+    } p2p;
+
+    struct {
+      const char* syncStrategy;
+    } ce;
+  };
+};
+
+struct proxyEvent {
+  uint64_t type;
+  std::atomic<int> count;
+  std::atomic<bool> ready;
+  int chunkSize;
+};
+
+static struct groupEvent group;
+static struct taskEvent task;
+static struct proxyEvent proxy;
+
+#define JSON_FILE_VERSION 4
 
 #define TIME_STRING_FORMAT "%Y-%m-%d %H:%M:%S"
 
@@ -469,12 +515,62 @@ void writeBenchmarkLinePreamble(size_t nBytes, size_t nElem, const char typeName
 
 // Finish a result record we were writing to stdout/json
 void writeBenchmarkLineTerminator(int actualIters, const char *name) {
+  const bool taskReady = task.ready.load(std::memory_order_acquire);
+  const bool proxyReady = proxy.ready.load(std::memory_order_acquire);
+  const char* implementation = "N/A";
+  const char* algo = "N/A";
+  const char* proto = "N/A";
+  const char* kernelVariant = "N/A";
+  const char* syncStrategy = "N/A";
+  int nChannels = -1;
+  int nWarps = -1;
+
+  if (taskReady) {
+    if (task.type == ncclProfileColl) {
+      implementation = task.coll.implementation;
+      algo = task.coll.algo;
+      proto = task.coll.proto;
+      kernelVariant = task.coll.kernelVariant;
+      nChannels = task.coll.nChannels;
+      nWarps = task.coll.nWarps;
+    } else if (task.type == ncclProfileP2p) {
+      implementation = "P2P";
+      nChannels = task.p2p.nChannels;
+    } else if (task.type == ncclProfileCeColl) {
+      implementation = "CE";
+      syncStrategy = task.ce.syncStrategy;
+    }
+  }
+
+  if (tuning) {
+    PRINT("  %-4s  %-14s  %-8s  %-18s  %-8s", implementation, algo, proto, kernelVariant, syncStrategy);
+    if (nChannels >= 0) PRINT("  %-9d", nChannels);
+    else PRINT("  %-9s", "N/A");
+    if (nWarps >= 0) PRINT("  %-6d", nWarps);
+    else PRINT("  %-6s", "N/A");
+    if (proxyReady) PRINT("  %-12d", proxy.chunkSize);
+    else PRINT("  %-12s", "N/A");
+  }
   PRINT("\n");
   if(write_json) {
+    if (tuning) {
+      jsonKey("tuning"); jsonStartObject();
+      jsonKey("implementation"); jsonStr(implementation);
+      jsonKey("algo"); jsonStr(algo);
+      jsonKey("proto"); jsonStr(proto);
+      jsonKey("kernelVariant"); jsonStr(kernelVariant);
+      jsonKey("syncStrategy"); jsonStr(syncStrategy);
+      jsonKey("#channels"); nChannels >= 0 ? jsonInt(nChannels) : jsonStr("N/A");
+      jsonKey("#warps"); nWarps >= 0 ? jsonInt(nWarps) : jsonStr("N/A");
+      jsonKey("netChunkSize"); proxyReady ? jsonInt(proxy.chunkSize) : jsonStr("N/A");
+      jsonFinishObject();
+    }
     jsonKey("actual_iterations"); jsonInt(actualIters);
     jsonKey("experiment_name");   jsonStr(name);
     jsonFinishObject();
   }
+  task.ready.store(false, std::memory_order_relaxed);
+  proxy.ready.store(false, std::memory_order_relaxed);
 }
 
 // Handle a cases where we don't write out of place results
@@ -740,26 +836,52 @@ testResult_t writeDeviceReport(size_t *maxMem, int localRank, int proc, int tota
 // Json results object and contained table list are left open
 void writeResultHeader(bool report_cputime, bool report_timestamps) {
   const char* tsLbl  = report_timestamps ? "timestamp" : "";
-  const int tsPad = report_timestamps ? 19 : 0;
+  const int tsPad = report_timestamps ? 21 : 0;
   const char* tsFmt = report_timestamps ? TIME_STRING_FORMAT : "";
   const char* timeStr = report_cputime ? "cputime" : "time";
   PRINT("#\n");
   if (per_iter_timing) {
-    PRINT("# %10s  %12s  %8s  %6s  %6s              out-of-place (+ per-iteration)                                in-place (+ per-iteration)\n", "", "", "", "", "");
-    PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %6s  %7s  %7s  %7s  %7s  %7s  %6s  %6s  %6s  %7s  %7s  %7s  %7s %*s\n",
+    PRINT("# %10s  %12s  %8s  %6s  %6s              out-of-place (+ per-iteration)                                in-place (+ per-iteration)%*s",
+          "", "", "", "", "", tsPad, "");
+    if (tuning) {
+      PRINT("%87s%44s", "tuning", "");
+    }
+    PRINT("\n");
+    PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %6s  %7s  %7s  %7s  %7s  %7s  %6s  %6s  %6s  %7s  %7s  %7s  %7s%*s",
           "size", "count", "type", "redop", "root",
           timeStr, "algbw", "busbw", "#wrong", "i_min", "i_max", "i_p99", "i_cv%",
           timeStr, "algbw", "busbw", "#wrong", "i_min", "i_max", "i_p99", "i_cv%", tsPad, tsLbl);
-    PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %6s  %7s  %7s  %7s  %7s  %7s  %6s  %6s  %6s  %7s  %7s  %7s  %7s %*s\n",
+    if (tuning) {
+      PRINT("  %-4s  %-14s  %-8s  %-18s  %-8s  %-9s  %-6s  %-12s", "impl", "algo", "proto", "kernelVariant", "sync", "#channels", "#warps", "netChunkSize");
+    }
+    PRINT("\n");
+    PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %6s  %7s  %7s  %7s  %7s  %7s  %6s  %6s  %6s  %7s  %7s  %7s  %7s%*s",
           "(B)", "(elements)", "", "", "",
           "(us)", "(GB/s)", "(GB/s)", "", "(us)", "(us)", "(us)", "(%)",
           "(us)", "(GB/s)", "(GB/s)", "", "(us)", "(us)", "(us)", "(%)", tsPad, tsFmt);
+    if (tuning) {
+      PRINT("  %-4s  %-14s  %-8s  %-18s  %-8s  %-9s  %-6s  %-12s", "", "", "", "", "", "", "", "");
+    }
+    PRINT("\n");
   } else {
-    PRINT("# %10s  %12s  %8s  %6s  %6s           out-of-place                       in-place          \n", "", "", "", "", "");
-    PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %6s  %7s  %6s  %6s  %6s %*s\n", "size", "count", "type", "redop", "root",
+    PRINT("# %10s  %12s  %8s  %6s  %6s           out-of-place                       in-place%*s",
+          "", "", "", "", "", tsPad, "");
+    if (tuning) {
+      PRINT("%63s%44s", "tuning", "");
+    }
+    PRINT("\n");
+    PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %6s  %7s  %6s  %6s  %6s%*s", "size", "count", "type", "redop", "root",
           timeStr, "algbw", "busbw", "#wrong", timeStr, "algbw", "busbw", "#wrong", tsPad, tsLbl);
-    PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %6s  %7s  %6s  %6s  %6s %*s\n", "(B)", "(elements)", "", "", "",
+    if (tuning) {
+      PRINT("  %-4s  %-14s  %-8s  %-18s  %-8s  %-9s  %-6s  %-12s", "impl", "algo", "proto", "kernelVariant", "sync", "#channels", "#warps", "netChunkSize");
+    }
+    PRINT("\n");
+    PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %6s  %7s  %6s  %6s  %6s%*s", "(B)", "(elements)", "", "", "",
           "(us)", "(GB/s)", "(GB/s)", "", "(us)", "(GB/s)", "(GB/s)", "", tsPad, tsFmt);
+    if (tuning) {
+      PRINT("  %-4s  %-14s  %-8s  %-18s  %-8s  %-9s  %-6s  %-12s", "", "", "", "", "", "", "", "");
+    }
+    PRINT("\n");
   }
 
   if(write_json) {
@@ -857,3 +979,179 @@ void writeErrors() {
 void finalizeFooter() {
   PRINT("#\n");
 }
+
+static int profilerContext;
+
+static ncclResult_t ncclProfilerInit(void** ctx, uint64_t id, int* eMask, const char* name, int nodes, int ranks, int rank, ncclDebugLogger_t logfn) {
+  *ctx = &profilerContext;
+  *eMask = (ncclProfileColl | ncclProfileP2p | ncclProfileProxyOp | ncclProfileCeColl);
+  return ncclSuccess;
+}
+
+static void ncclProfilerStartCollEvent(void** eHandle, uint64_t type, const char* implementation, const char* algo,
+                                       const char* proto, const char* kernelVariant, uint8_t nChannels, uint8_t nWarps) {
+  if (task.count.fetch_add(1, std::memory_order_relaxed) == 0) {
+    task.type = type;
+    task.coll.implementation = implementation;
+    task.coll.algo = algo ? algo : "N/A";
+    task.coll.proto = proto ? proto : "N/A";
+    task.coll.kernelVariant = kernelVariant ? kernelVariant : "N/A";
+    task.coll.nChannels = nChannels == 0 ? -1 : nChannels;
+    task.coll.nWarps = nWarps;
+  }
+  *eHandle = &task;
+}
+
+static void ncclProfilerStartP2pEvent(void** eHandle, uint64_t type, uint8_t nChannels) {
+  if (task.count.fetch_add(1, std::memory_order_relaxed) == 0) {
+    task.type = type;
+    task.p2p.nChannels = nChannels;
+  }
+  *eHandle = &task;
+}
+
+static void ncclProfilerStartCeEvent(void** eHandle, uint64_t type, const char* syncStrategy) {
+  if (task.count.fetch_add(1, std::memory_order_relaxed) == 0) {
+    task.type = type;
+    task.ce.syncStrategy = syncStrategy ? syncStrategy : "N/A";
+  }
+  *eHandle = &task;
+}
+
+static void ncclProfilerStartProxyEvent(void** eHandle, uint64_t type, int chunkSize) {
+  // The chunk size is the same for all channels, so capture it only once.
+  if (proxy.count.fetch_add(1, std::memory_order_relaxed) == 0) {
+    proxy.type = type;
+    proxy.chunkSize = chunkSize;
+  }
+  *eHandle = &proxy;
+}
+
+template <typename EventDescr>
+static ncclResult_t ncclProfilerStartEventCommon(void** eHandle, EventDescr* eDescr) {
+  switch (eDescr->type) {
+    case ncclProfileGroup:
+      group.type = eDescr->type;
+      group.count.fetch_add(1, std::memory_order_relaxed);
+      *eHandle = &group;
+      break;
+    case ncclProfileP2p:
+      ncclProfilerStartP2pEvent(eHandle, eDescr->type, eDescr->p2p.nChannels);
+      break;
+    case ncclProfileProxyOp:
+      ncclProfilerStartProxyEvent(eHandle, eDescr->type, eDescr->proxyOp.chunkSize);
+      break;
+    default:;
+  }
+  return ncclSuccess;
+}
+
+static ncclResult_t ncclProfilerStartEvent_v5(void*, void** eHandle, ncclProfilerEventDescr_v5_t* eDescr) {
+  if (eDescr->type == ncclProfileColl) {
+    ncclProfilerStartCollEvent(eHandle, eDescr->type, "GEN", eDescr->coll.algo, eDescr->coll.proto, nullptr,
+                               eDescr->coll.nChannels, eDescr->coll.nWarps);
+    return ncclSuccess;
+  }
+  return ncclProfilerStartEventCommon(eHandle, eDescr);
+}
+
+static ncclResult_t ncclProfilerStartEvent_v6(void*, void** eHandle, ncclProfilerEventDescr_v6_t* eDescr) {
+  if (eDescr->type == ncclProfileCeColl) {
+    ncclProfilerStartCeEvent(eHandle, eDescr->type, eDescr->ceColl.syncStrategy);
+    return ncclSuccess;
+  }
+  if (eDescr->type == ncclProfileColl) {
+    ncclProfilerStartCollEvent(eHandle, eDescr->type, "GEN", eDescr->coll.algo, eDescr->coll.proto, nullptr,
+                               eDescr->coll.nChannels, eDescr->coll.nWarps);
+    return ncclSuccess;
+  }
+  return ncclProfilerStartEventCommon(eHandle, eDescr);
+}
+
+static ncclResult_t ncclProfilerStartEvent_v7(void*, void** eHandle, ncclProfilerEventDescr_v7_t* eDescr) {
+  if (eDescr->type == ncclProfileColl) {
+    const char* kernelVariant = eDescr->coll.isSymColl ? eDescr->coll.kernelVariant : nullptr;
+    if (kernelVariant && eDescr->coll.func) {
+      size_t len = strlen(eDescr->coll.func);
+      if (strncmp(kernelVariant, eDescr->coll.func, len) == 0 && kernelVariant[len] == '_') kernelVariant += len + 1;
+    }
+    ncclProfilerStartCollEvent(eHandle, eDescr->type, eDescr->coll.isSymColl ? "SYM" : "GEN",
+                               eDescr->coll.isSymColl ? "N/A" : eDescr->coll.algo,
+                               eDescr->coll.isSymColl ? "N/A" : eDescr->coll.proto, kernelVariant,
+                               eDescr->coll.nChannels, eDescr->coll.nWarps);
+    return ncclSuccess;
+  }
+  if (eDescr->type == ncclProfileCeColl) {
+    ncclProfilerStartCeEvent(eHandle, eDescr->type, eDescr->ceColl.syncStrategy);
+    return ncclSuccess;
+  }
+  return ncclProfilerStartEventCommon(eHandle, eDescr);
+}
+
+static ncclResult_t ncclProfilerStopEvent(void* eHandle) {
+  uint64_t type = *(uint64_t *)eHandle;
+  switch (type) {
+    case ncclProfileGroup: {
+      struct groupEvent* e = (struct groupEvent *)eHandle;
+      e->count.fetch_sub(1, std::memory_order_relaxed);
+    } break;
+    case ncclProfileColl:
+    case ncclProfileP2p:
+    case ncclProfileCeColl: {
+      struct taskEvent* e = (struct taskEvent *)eHandle;
+      if (e->count.fetch_sub(1, std::memory_order_relaxed) == 1)
+        e->ready.store(true, std::memory_order_release);
+    } break;
+    case ncclProfileProxyOp: {
+      struct proxyEvent* e = (struct proxyEvent *)eHandle;
+      if (e->count.fetch_sub(1, std::memory_order_relaxed) == 1)
+        e->ready.store(true, std::memory_order_release);
+    } break;
+    default:;
+  }
+  return ncclSuccess;
+}
+
+static ncclResult_t ncclProfilerRecordEventState_v5(void* eHandle, ncclProfilerEventState_v5_t eState, ncclProfilerEventStateArgs_v5_t* eStateArgs) {
+  return ncclSuccess;
+}
+
+static ncclResult_t ncclProfilerRecordEventState_v6(void* eHandle, ncclProfilerEventState_v6_t eState, ncclProfilerEventStateArgs_v6_t* eStateArgs) {
+  return ncclSuccess;
+}
+
+static ncclResult_t ncclProfilerRecordEventState_v7(void* eHandle, ncclProfilerEventState_v7_t eState, ncclProfilerEventStateArgs_v7_t* eStateArgs) {
+  return ncclSuccess;
+}
+
+static ncclResult_t ncclProfilerFinalize(void* ctx) {
+  return ncclSuccess;
+}
+
+// perftest exposes profiler interface to nccl
+ncclProfiler_v5_t ncclProfiler_v5 = {
+  /* .name = */ "perftest",
+  /* .init = */ ncclProfilerInit,
+  /* .startEvent = */ ncclProfilerStartEvent_v5,
+  /* .stopEvent = */ ncclProfilerStopEvent,
+  /* .recordEventState = */ ncclProfilerRecordEventState_v5,
+  /* .finalize = */ ncclProfilerFinalize,
+};
+
+ncclProfiler_v6_t ncclProfiler_v6 = {
+  /* .name = */ "perftest",
+  /* .init = */ ncclProfilerInit,
+  /* .startEvent = */ ncclProfilerStartEvent_v6,
+  /* .stopEvent = */ ncclProfilerStopEvent,
+  /* .recordEventState = */ ncclProfilerRecordEventState_v6,
+  /* .finalize = */ ncclProfilerFinalize,
+};
+
+ncclProfiler_v7_t ncclProfiler_v7 = {
+  /* .name = */ "perftest",
+  /* .init = */ ncclProfilerInit,
+  /* .startEvent = */ ncclProfilerStartEvent_v7,
+  /* .stopEvent = */ ncclProfilerStopEvent,
+  /* .recordEventState = */ ncclProfilerRecordEventState_v7,
+  /* .finalize = */ ncclProfilerFinalize,
+};