feat(rccl-tests): gin-sdma in-device (wall_clock64) A2A timing mechanism

projects/rccl-tests/CMakeLists.txt

@@ -15,12 +15,16 @@ endif()
 #==================================================================================================
 project(rccl-tests LANGUAGES CXX)
 
+enable_testing()
+
 # Build options
 #==================================================================================================
 option(USE_MPI                     "Build RCCL-tests with MPI support."           OFF)
 option(ENABLE_DEVICE_API           "Build with NCCL Device API support (GIN)."    OFF)
 option(ENABLE_ROCSHMEM_GIN         "Build with rocshmem GIN backend support."     OFF)
 option(BUILD_LOCAL_GPU_TARGET_ONLY "Build only for GPUs detected on this machine" OFF)
+option(BUILD_RCCL_TESTS_HOST_UNIT_TESTS
+       "Build CPU-only host unit tests under test/ (standalone builds only)"   ON)
 
 if (NOT CMAKE_BUILD_TYPE)
   message(WARNING "CMAKE_BUILD_TYPE is not defined. Setting to Release")
@@ -115,7 +119,7 @@ set(CMAKE_CXX_STANDARD   17)   # rocshmem GIN QP device headers require C++17
 set(CMAKE_CXX_EXTENSIONS OFF)  # Without this line, it will add -std=gnu++14 instead, which has some issues.
 set(CPACK_PACKAGING_INSTALL_PREFIX "${ROCM_PATH}" CACHE PATH "Path to install to when packaged.")
 if(ROCM_PATH)
-  #list(APPEND CMAKE_PREFIX_PATH  # Temporary workaround 
+  #list(APPEND CMAKE_PREFIX_PATH  # Temporary workaround
   list(PREPEND CMAKE_PREFIX_PATH  # Add ROCM_PATH to CMake search paths (for finding HIP / HSA
               ${ROCM_PATH}
               ${ROCM_PATH}/hip
@@ -282,6 +286,15 @@ add_subdirectory(verifiable)
 # Add all of the tests
 add_subdirectory(src)
 
+# Host unit tests (GTest) are for standalone/cmake dev builds only. TheRock
+# comm-libs configures rccl-tests as a sub-project; find_package(GTest) there
+# requires super-project BUILD_DEPS wiring we do not have, so skip test/.
+if(BUILD_RCCL_TESTS_HOST_UNIT_TESTS AND NOT DEFINED THEROCK_STAGE_INSTALL_ROOT)
+  add_subdirectory(test)
+elseif(DEFINED THEROCK_STAGE_INSTALL_ROOT)
+  message(STATUS "rccl-tests: skipping host unit tests (TheRock super-project build)")
+endif()
+
 rocm_setup_version(VERSION "2.19.6")
 
 # Create ROCm standard packages

projects/rccl-tests/src/CMakeLists.txt

@@ -73,6 +73,8 @@ set(COMMON_FILES
   collector.cu
   common.h
   common.cu
+  gin_sdma_devtime.h
+  gin_sdma_devtime_host.h
   nccl1_compat.h
   rccl_compat.h
   rccl_float8.h

projects/rccl-tests/src/alltoall.cu

@@ -7,6 +7,7 @@
 
 #include "cuda_runtime.h"
 #include "common.h"
+#include "gin_sdma_devtime.h"  // shared device-side (wall_clock64) timing scaffold
 #if defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
 #include "nccl_device.h"
 #include "rccl_vector_types.h"
@@ -98,7 +99,11 @@ testResult_t AlltoAllGetDevCommRequirements(int deviceImpl, ncclDevCommRequireme
         fprintf(stderr, "This test requires GIN support, but GIN support is not enabled for this communicator.\n");
         return testInvalidUsage;
       }
-      reqs->barrierCount = 1;
+      // barrierCount provisions per-CTA hybrid world barriers (LSA+rail+world) used
+      // by HybridAlltoAllTimedKernel to join CTA 0 (GIN) with LSA CTAs each iter.
+      // The production body still uses barrier index 0 on CTA 0 and lsaBarrier on
+      // CTAs 1..N; ginSignalCount stays 1 because only CTA 0 issues GIN puts.
+      reqs->barrierCount = deviceCtaCount;
       reqs->lsaBarrierCount = deviceCtaCount - 1;
       reqs->ginSignalCount = 1;
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2, 29, 7)
@@ -130,7 +135,7 @@ bool AlltoAllGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* req
       return true;
     case 4: // HybridAlltoAllKernel: CTA 0 = GIN, CTAs 1..N = LSA
       if (deviceCtaCount < 2) return false;
-      reqs->barrierCount = 1;
+      reqs->barrierCount = deviceCtaCount;
       reqs->lsaBarrierCount = deviceCtaCount - 1;
       reqs->ginSignalCount = 1;
       return true;
@@ -253,8 +258,13 @@ __global__ void NvlAlltoAllKernelOptimized(ncclWindow_t sendwin, size_t sendoffs
 }
 
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7) && defined(NCCL_OS_LINUX)
+// Collective body, factored out of the __global__ entry so it can be invoked
+// either once (production kernel) or in a persistent skip+loop (device-timing
+// kernel). It re-derives all per-call sync state at entry (GIN re-reads the
+// accumulated signal), so calling it back-to-back in a loop yields a sequence of
+// complete, correct alltoalls with no external bookkeeping.
 template <typename T>
-__global__ void GinAlltoAllKernel(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm) {
+__device__ void ginAlltoAllBody(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm) {
   int ginContext = 0;
   unsigned int signalIndex = blockIdx.x;
   ncclGin gin { devComm, ginContext };
@@ -285,8 +295,20 @@ __global__ void GinAlltoAllKernel(ncclWindow_t sendwin, size_t sendoffset, ncclW
   bar.sync(ncclCoopCta(), cuda::memory_order_release, ncclGinFenceLevel::Relaxed);
 }
 
+// Production entry: one collective call. Thin wrapper over the body so the
+// device-timed kernel and the production path share identical code.
 template <typename T>
-__global__ void HybridAlltoAllKernel(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm) {
+__global__ void GinAlltoAllKernel(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm) {
+  ginAlltoAllBody<T>(sendwin, sendoffset, recvwin, recvoffset, count, root, devComm);
+}
+
+// Hybrid LSA+GIN alltoall: CTA 0 handles remote peers via GIN,
+// CTAs 1..N handle intra-node peers via LSA.
+// Production body: CTA 0 uses hybrid world barrier index 0; LSA CTAs use
+// lsaBarrier (blockIdx.x - 1). DevComm barrierCount == deviceCtaCount also
+// provisions per-CTA hybrid world barriers for HybridAlltoAllTimedKernel joins.
+template <typename T>
+__device__ void hybridAlltoAllBody(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm) {
   ncclTeam world = ncclTeamWorld(devComm);
   ncclTeam lsa = ncclTeamLsa(devComm);
   const int startLsa = world.rank - lsa.rank;
@@ -343,6 +365,67 @@ __global__ void HybridAlltoAllKernel(ncclWindow_t sendwin, size_t sendoffset, nc
     lsaBar.sync(ncclCoopCta(), cuda::memory_order_release);
   }
 }
+
+// Production entry: one collective call. Thin wrapper over the body so the
+// device-timed kernel and the production path share identical code.
+template <typename T>
+__global__ void HybridAlltoAllKernel(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm) {
+  hybridAlltoAllBody<T>(sendwin, sendoffset, recvwin, recvoffset, count, root, devComm);
+}
+
+// Join all CTAs (GIN CTA 0 + LSA CTAs 1..N) in the timed-kernel loop. The
+// production body uses split sync domains; ncclBarrierSession(world) spans LSA,
+// rail, and world GIN barriers so every CTA waits together (see NCCL hybrid
+// alltoall example). Requires barrierCount == gridDim.x in devComm setup.
+__device__ inline void hybridTimedGridJoin(struct ncclDevComm devComm) {
+  int ginContext = 0;
+  ncclGin gin { devComm, ginContext };
+  ncclBarrierSession<ncclCoopCta> bar { ncclCoopCta(), ncclTeamTagWorld(), gin, (uint32_t)blockIdx.x };
+  bar.sync(ncclCoopCta(), cuda::memory_order_acquire, ncclGinFenceLevel::Relaxed);
+}
+
+// Device-side timing kernels (rocSHMEM AllToAll methodology, AICOMRCCL-1459).
+// A single persistent launch runs (skip + loop) back-to-back collective bodies
+// and brackets only the timed region with the GPU fixed-frequency wall clock
+// (wall_clock64()), so the reported span excludes host launch, teardown, and
+// stream/graph overhead -- it is the pure device-function execution time.
+//
+// skip warmup iterations run first (steady-state caches/queues, discarded).
+// At i == skip every CTA records its start stamp; after the final iteration
+// every CTA records its end stamp. The host reduces min(start)/max(end) across
+// CTAs (the grid's true busy window) and MAX across ranks (the slowest rank
+// closes the collective), then divides by loop for per-iteration latency.
+//
+// Because the body re-derives all per-call sync state at entry (GIN re-reads the
+// accumulated signal; LSA rebuilds its barrier session), looping is correct with
+// no extra signal bookkeeping.
+template <typename T>
+__global__ void GinAlltoAllTimedKernel(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm, int loop, int skip, long long* start_time, long long* end_time) {
+  for (int i = 0; i < skip + loop; i++) {
+    if (i == skip) {
+      __syncthreads();
+      if (threadIdx.x == 0) start_time[blockIdx.x] = wall_clock64();
+    }
+    ginAlltoAllBody<T>(sendwin, sendoffset, recvwin, recvoffset, count, root, devComm);
+  }
+  __syncthreads();
+  if (threadIdx.x == 0) end_time[blockIdx.x] = wall_clock64();
+}
+
+template <typename T>
+__global__ void HybridAlltoAllTimedKernel(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int root, struct ncclDevComm devComm, int loop, int skip, long long* start_time, long long* end_time) {
+  for (int i = 0; i < skip + loop; i++) {
+    if (i == skip) {
+      hybridTimedGridJoin(devComm);
+      __syncthreads();
+      if (threadIdx.x == 0) start_time[blockIdx.x] = wall_clock64();
+    }
+    hybridAlltoAllBody<T>(sendwin, sendoffset, recvwin, recvoffset, count, root, devComm);
+  }
+  hybridTimedGridJoin(devComm);
+  __syncthreads();
+  if (threadIdx.x == 0) end_time[blockIdx.x] = wall_clock64();
+}
 #endif
 #endif
 
@@ -400,14 +483,112 @@ testResult_t AlltoAllRunColl(void* sendbuff, size_t sendoffset, void* recvbuff,
   return testSuccess;
 }
 
+// Device-side (in-kernel wall_clock64) timing for AllToAll (GIN and Hybrid
+// tiers). Opt-in via --device_timing: launches the persistent timed kernel once
+// for the current size, brackets only the (skip+loop) steady-state collectives
+// with the GPU wall clock, reduces the grid busy window (min start .. max end
+// over CTAs) and the slowest rank (MPI MAX), and reports the per-iteration
+// device latency. loop/skip come from --devtime_loop/--devtime_skip (default
+// 10/10); size-tier overrides via --devtime_loop_mid/_large and
+// --devtime_skip_mid/_large. The timed kernel launches with the same
+// <<<deviceCtaCount, 512>>> grid the production path uses, and selects the
+// GIN (deviceImpl 3) or Hybrid (deviceImpl 4) body to match RunColl.
+#if defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7)
+testResult_t AlltoAllDeviceTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int in_place, double* outDeltaSec) {
+  if (!deviceTimingMode) return testSuccess;
+
+  // Only the GIN (deviceImpl 3) and Hybrid (deviceImpl 4) tiers provision the GIN
+  // signals/barriers the timed bodies rely on. The Nvl LSA impls (-D1/-D2) only
+  // provision the LSA barrier, so running the GIN timed kernel there would fault
+  // on an out-of-range signal index or wait forever -- skip them.
+  if (deviceImpl != 3 && deviceImpl != 4) return testSuccess;
+
+  const size_t count = args->nbytes / wordSize(type);
+  if (count == 0 || devtimeLoop < 1) return testSuccess;
+  const size_t perPeerBytes = count * wordSize(type);
+
+  // By default use the exact skip/loop counts at every size so the device-timing
+  // window matches the host tests' -w/-n. Optional mid/large tier overrides cap
+  // iteration counts for very large chunks (bounded runtime; per-call copy dominates).
+  // Clamp skip to >= 0: the timed kernel only stamps start_time[] at i == skip, so
+  // a negative skip would leave start_time[] as uninitialized cudaMalloc memory and
+  // measure() would reduce mx - mn over garbage (reported as the mode-2 metric).
+  int loop = devtimeLoop;
+  int skip = devtimeSkip < 0 ? 0 : devtimeSkip;
+  if (devtimeLoopLarge > 0 && perPeerBytes >= (size_t)64 * 1024 * 1024) {
+    loop = devtimeLoopLarge;
+    if (devtimeSkipLarge >= 0) skip = devtimeSkipLarge;
+    else skip = (skip < 1) ? skip : 1;
+  } else if (devtimeLoopMid > 0 && perPeerBytes >= (size_t)8 * 1024 * 1024) {
+    loop = devtimeLoopMid;
+    if (devtimeSkipMid >= 0) skip = devtimeSkipMid;
+    else skip = (skip < 2) ? skip : 2;
+  }
+
+  int gridCtas = deviceCtaCount;
+  if (gridCtas < 1) gridCtas = 1;
+
+  // Match the kernel RunColl would launch for the selected -D device impl.
+  const bool hybrid = (deviceImpl == 4);
+  auto kernel = hybrid ? SPECIALIZE_KERNEL(HybridAlltoAllTimedKernel, type, op)
+                       : SPECIALIZE_KERNEL(GinAlltoAllTimedKernel, type, op);
+  const char* tierName = hybrid ? "HYB" : "GIN";
+  if (kernel == nullptr) return testSuccess;
+
+  // Shared scaffold: allocates per-CTA start/end stamps, launches the timed kernel,
+  // reduces min(start)..max(end) over CTAs and MPI-MAX across ranks -> per-iter us.
+  double devUs = 0.0;
+  TESTCHECK(gin_devtime::measure(args, gridCtas, loop,
+      [&](int i, long long* d_start, long long* d_end) {
+        ncclDevComm* devComm = args->devComms + i;
+        ncclWindow_t sendwin = (ncclWindow_t)args->sendRegHandles[i];
+        ncclWindow_t recvwin = (ncclWindow_t)args->recvRegHandles[i];
+        kernel<<<gridCtas, 512, 0, args->streams[i]>>>(
+            sendwin, 0, recvwin, 0, count, root, *devComm, loop, skip, d_start, d_end);
+      },
+      &devUs));
+
+  int nRanksGlobal = args->nProcs * args->nThreads * args->nGpus;
+
+  // Mode 2 (device-time-only): hand the per-iteration latency (seconds) back to
+  // BenchTime, which reports it as THE time/busbw metric on the normal result
+  // line. Stay silent here so that line is not split.
+  if (outDeltaSec != nullptr) {
+    // Negative sentinel = "no valid measurement" (the timed grid produced no
+    // positive busy window). The caller (BenchTime mode 2) then WARNs and skips the
+    // row instead of reporting a bogus 0.00 us / inf GB/s.
+    *outDeltaSec = (devUs > 0.0) ? devUs * 1.0e-6 : -1.0;
+    return testSuccess;
+  }
+
+  // Mode 1 (augment): buffer the extra device-only line; TimeTest flushes it after
+  // writeBenchmarkLineTerminator so the row is not split between OOP/IP columns.
+  if (args->proc == 0 && args->thread == 0 && devUs > 0.0) {
+    double sec = devUs * 1.0e-6;
+    double algBw = (double)(perPeerBytes * (size_t)nRanksGlobal) / 1.0e9 / sec;
+    double busBw = algBw * ((double)(nRanksGlobal - 1) / (double)nRanksGlobal);
+    snprintf(args->devtimeAugmentLine, sizeof(args->devtimeAugmentLine),
+             "#[a2a-devtime] size %12zu B  tier %-3s  ctas %2d  loop %2d skip %2d  devtime %10.2f us  algbw %8.2f GB/s  busbw %8.2f GB/s\n",
+             perPeerBytes * (size_t)nRanksGlobal, tierName, gridCtas, loop, skip, devUs, algBw, busBw);
+  }
+  return testSuccess;
+}
+#else
+testResult_t AlltoAllDeviceTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int in_place, double* outDeltaSec) {
+  return testSuccess;  // device API path not available in this build
+}
+#endif
+
 struct testColl alltoAllTest = {
   "AlltoAll",
   AlltoAllGetCollByteCount,
   AlltoAllInitData,
   AlltoAllGetBw,
   AlltoAllRunColl,
   NULL,
-  NULL
+  NULL,
+  NULL,
+  AlltoAllDeviceTime
 };
 
 void AlltoAllGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) {