diff --git a/projects/rccl-tests/CMakeLists.txt b/projects/rccl-tests/CMakeLists.txt
index 3f5af2463d98..7b54fced4abd 100644
--- a/projects/rccl-tests/CMakeLists.txt
+++ b/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
diff --git a/projects/rccl-tests/src/CMakeLists.txt b/projects/rccl-tests/src/CMakeLists.txt
index b0d1d3dfc8f0..5f5ca7141d78 100644
--- a/projects/rccl-tests/src/CMakeLists.txt
+++ b/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
diff --git a/projects/rccl-tests/src/alltoall.cu b/projects/rccl-tests/src/alltoall.cu
index ab070f551d14..2d3cf5a892ca 100644
--- a/projects/rccl-tests/src/alltoall.cu
+++ b/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,6 +483,102 @@ 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,
@@ -407,7 +586,9 @@ struct testColl alltoAllTest = {
   AlltoAllGetBw,
   AlltoAllRunColl,
   NULL,
-  NULL
+  NULL,
+  NULL,
+  AlltoAllDeviceTime
 };
 
 void AlltoAllGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, int nranks) {
diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu
index a180ced69c9e..bc9eabb6bfa2 100644
--- a/projects/rccl-tests/src/common.cu
+++ b/projects/rccl-tests/src/common.cu
@@ -28,6 +28,7 @@
 
 #include "verifiable.h"
 #include "util.h"
+#include "gin_sdma_devtime_host.h"
 
 #if defined(NCCL_OS_LINUX)
 extern char **environ;
@@ -156,7 +157,15 @@ std::string rccl_output_file;
 std::string rccl_output_format;
 static int report_cputime = 0;
 static int report_timestamps = 0;
-static int deviceImpl = 0;
+int deviceImpl = 0;  // non-static: per-collective device-timing hooks read this to pick the matching timed kernel
+int deviceTimingMode = 0;
+int devtimeLoop = 10;
+int devtimeSkip = 10;
+int devtimeLoopMid = 0;
+int devtimeLoopLarge = 0;
+int devtimeSkipMid = -1;
+int devtimeSkipLarge = -1;
+int devtimeCheck = 0;
 int unalign = 0;
 int memory_report = 0;
 
@@ -661,6 +670,10 @@ void ReduceProcessMaxIterTimes(double* iterTimes, const double* allProcessTimes,
   }
 }
 
+// Explicit instantiations so other TUs (e.g. gin_sdma_devtime.h via alltoall.cu)
+// can call Allreduce through the common.h declaration.
+template void Allreduce<double>(struct threadArgs* args, double* value, int average);
+template void Allreduce<long long>(struct threadArgs* args, long long* value, int average);
 testResult_t CheckData(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int in_place, int64_t *wrongElts) {
   int nranks = args->nProcs*args->nGpus*args->nThreads;
   size_t count = args->expectedBytes/wordSize(type);
@@ -946,10 +959,21 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
 
   Barrier(args);
 
+  // Device-time-only mode (--device_timing=2): for collectives that provide a
+  // device-timing hook, skip the host/graph timed loop entirely and report the
+  // in-kernel wall_clock64 measurement as the metric. Warmup and datacheck still
+  // run (correctness). Mode 1 keeps the normal timed loop and augments with an
+  // extra line; mode 0 is off.
+  const int devTimeMode =
+      (deviceImpl != 0 && args->collTest->deviceTime != nullptr) ? deviceTimingMode : 0;
+  // Device-time-only applies to the out-of-place pass (the reported metric). The
+  // in-place pass keeps normal timing so its column stays valid.
+  const bool deviceOnly = (devTimeMode == 2) && (in_place == 0);
+
 #if HIP_VERSION >= 50221310
   std::vector<cudaGraph_t> graphs(args->nGpus);
   std::vector<cudaGraphExec_t> graphExec(args->nGpus);
-  if (cudaGraphLaunches >= 1) {
+  if (cudaGraphLaunches >= 1 && !deviceOnly) {
     // Begin cuda graph capture
     for (int i=0; i<args->nGpus; i++) {
       // Thread local mdoe is needed for:
@@ -961,32 +985,34 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
   }
 #endif
 
-  int record = per_iter_timing;
+  int record = per_iter_timing && !deviceOnly;
   if (record) TESTCHECK(initEvents(args, iters));
 
   // Performance Benchmark
   double collOnlySec = 0;
   int pairedEvents = record && blocking_coll == 3;
   timer tim;
-  for (int iter = 0; iter < iters; iter++) {
-    double t0 = tim.elapsed();
-    if (agg_iters>1) NCCLCHECK(ncclGroupStart());
-    if (record) TESTCHECK(recordEvents(args, iters, iter));
-    for (int aiter = 0; aiter < agg_iters; aiter++) {
-      TESTCHECK(startColl(args, type, op, root, in_place, iter*agg_iters+aiter));
-    }
-    if (agg_iters>1) NCCLCHECK(ncclGroupEnd());
-    if (blocking_coll == 3) {
-      TESTCHECK(testStreamSynchronize(args->nGpus, args->streams, args->comms));
-      if (pairedEvents) TESTCHECK(recordEvents(args, iters, iter, 1));
-      collOnlySec += tim.elapsed() - t0;
-      Barrier(args);
+  if (!deviceOnly) {
+    for (int iter = 0; iter < iters; iter++) {
+      double t0 = tim.elapsed();
+      if (agg_iters>1) NCCLCHECK(ncclGroupStart());
+      if (record) TESTCHECK(recordEvents(args, iters, iter));
+      for (int aiter = 0; aiter < agg_iters; aiter++) {
+        TESTCHECK(startColl(args, type, op, root, in_place, iter*agg_iters+aiter));
+      }
+      if (agg_iters>1) NCCLCHECK(ncclGroupEnd());
+      if (blocking_coll == 3) {
+        TESTCHECK(testStreamSynchronize(args->nGpus, args->streams, args->comms));
+        if (pairedEvents) TESTCHECK(recordEvents(args, iters, iter, 1));
+        collOnlySec += tim.elapsed() - t0;
+        Barrier(args);
+      }
     }
   }
   if (record && !pairedEvents) TESTCHECK(recordEvents(args, iters, iters));
 
 #if HIP_VERSION >= 50221310
-  if (cudaGraphLaunches >= 1) {
+  if (cudaGraphLaunches >= 1 && !deviceOnly) {
     // HIP (and CUDA) graph limitation: in single-process multi-GPU mode (-g N),
     // graph nodes are associated with the calling thread's current device at
     // capture time, not the stream's device. RCCL internally calls cudaSetDevice()
@@ -1042,12 +1068,13 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
 #endif
 
   double cputimeSec = blocking_coll == 3 ? collOnlySec : tim.elapsed();
-  cputimeSec /= iters*agg_iters;
+  cputimeSec = rccl_tests_devtime::normalizeElapsedPerIter(cputimeSec, iters, agg_iters);
   TESTCHECK(completeColl(args));
 
   // Exclude per-iteration result processing from the reported time.
-  double deltaSec = (blocking_coll == 3 ? collOnlySec : tim.elapsed())/(iters*agg_iters);
-  if (cudaGraphLaunches >= 1) deltaSec = deltaSec/cudaGraphLaunches;
+  double deltaSec = blocking_coll == 3 ? collOnlySec : tim.elapsed();
+  deltaSec = rccl_tests_devtime::normalizeElapsedPerIter(deltaSec, iters, agg_iters);
+  deltaSec = rccl_tests_devtime::applyCudaGraphLaunchesScale(deltaSec, cudaGraphLaunches);
 
   if (record) {
     // completeColl skips the stream sync in blocking mode with a single aggregated iteration.
@@ -1057,10 +1084,76 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
     TESTCHECK(destroyEvents(args, iters));
   }
 
-  Allreduce(args, &deltaSec, average);
+  if (!deviceOnly) Allreduce(args, &deltaSec, average);
+
+  bool skipMetricRow = false;
+
+  if (deviceOnly) {
+    // The reported metric is the in-kernel wall_clock64 device latency
+    // (per iteration, max across ranks), measured by the collective's hook.
+    // Note: the datacheck loop below re-inits buffers and re-runs the PRODUCTION
+    // kernel via startColl before CheckData, so it validates the production path,
+    // not the timed kernel (whose output is overwritten first). By default the
+    // *TimedKernel bodies call the same __device__ collective functions as the
+    // production kernels, so collective correctness is covered; only the timing-only
+    // loop/skip/stamp wrapper around them is not exercised by that datacheck.
+    // deviceTime() leaves this negative when the timing hook produced no
+    // measurement -- the timed kernel never ran (non-GIN impl, unsupported op/type,
+    // count == 0, or loop < 1). Reporting that as the metric would divide by zero in
+    // getBw (0.00 us / inf GB/s), and the DEVTIME_CHECK below would validate a stale
+    // buffer -- the warmup/production result (false pass), or raw init scratch under
+    // -w 0 (false fail) -- for a kernel that never executed. So WARN once and skip
+    // the metric for this row rather than emit a bogus metric or a misleading check
+    // verdict; datacheck still runs below.
+    // -H 1: zero recvbuff first so the check below observes the timed kernel's
+    // writes rather than the warmup's identical output already sitting there.
+    if (devtimeCheck && datacheck) {
+      for (int i = 0; i < args->nGpus; i++) {
+        CUDACHECK(cudaSetDevice(args->gpus[i]));
+        CUDACHECK(cudaMemset(args->recvbuffs[i], 0, args->expectedBytes));
+      }
+    }
+    double deviceDeltaSec = -1.0;
+    TESTCHECK(args->collTest->deviceTime(args, type, op, root, in_place, &deviceDeltaSec));
+    if (deviceDeltaSec <= 0.0) {
+      if (args->proc == 0 && args->thread == 0) {
+        fprintf(stderr, "# WARN --device_timing=2: no in-kernel device-time "
+                        "measurement for this size (timed kernel did not run: non-GIN impl, "
+                        "unsupported op/type, count==0, or loop<1); skipping row "
+                        "(size %ld bytes)\n", args->expectedBytes);
+      }
+      skipMetricRow = true;
+    } else {
+      deltaSec = deviceDeltaSec;
+      cputimeSec = deviceDeltaSec;
+
+      // Opt-in validation of the TIMED path itself (--devtime_check=1).
+      // Reached only when the timed kernel actually ran (deviceDeltaSec > 0 above), so
+      // recvbuff holds the timed kernel's output -- not a stale warmup/production
+      // buffer. The hook just ran skip+loop back-to-back collectives into recvbuff;
+      // each iteration is a full deterministic collective over the unchanged sendbuff,
+      // so the final recvbuff equals `expected` from initData. Validate it HERE -- the
+      // only point the timed kernel's output is live, before the datacheck loop below
+      // re-inits buffers and overwrites it with the production path. Needs datacheck on
+      // (that is what populates `expected`). Combines wrong-element counts across
+      // threads/ranks so any rank's mismatch fails the run.
+      if (devtimeCheck && datacheck) {
+        int64_t timedWrong = 0;
+        TESTCHECK(CheckData(args, type, op, root, in_place, &timedWrong));
+        long long timedWrong1 = timedWrong;
+        Allreduce(args, &timedWrong1, /*sum*/4);
+        if (timedWrong1) {
+          fprintf(stderr, "\nERROR: --devtime_check: timed-kernel output datacheck "
+                          "failed with %lld wrong elements (size %ld bytes)\n",
+                          timedWrong1, args->expectedBytes);
+          return testInternalError;
+        }
+      }
+    }
+  }
 
 #if HIP_VERSION >= 50221310
-  if (cudaGraphLaunches >= 1) {
+  if (cudaGraphLaunches >= 1 && !deviceOnly) {
     //destroy cuda graph
     for (int i=0; i<args->nGpus; i++) {
 #ifdef __HIP_PLATFORM_AMD__
@@ -1072,8 +1165,24 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
   }
 #endif
 
-  double algBw, busBw;
-  args->collTest->getBw(count, wordSize(type), deltaSec, &algBw, &busBw, args->nProcs*args->nThreads*args->nGpus);
+  // Host-timed path can also yield zero elapsed (timer resolution, fast
+  // collectives, or platform quirks). getBw divides by sec; skip the metric for
+  // this row rather than SIGFPE or emit inf GB/s -- same rationale as the
+  // --device_timing=2 guard. Datacheck still runs below.
+  if (!skipMetricRow && rccl_tests_devtime::shouldSkipBenchTimeRow(deltaSec)) {
+    if (args->proc == 0 && args->thread == 0) {
+      fprintf(stderr, "# WARN: zero or negative elapsed time (deltaSec=%g); "
+                      "skipping row (size %ld bytes)\n", deltaSec, args->expectedBytes);
+    }
+    skipMetricRow = true;
+  }
+
+  double algBw = 0.0, busBw = 0.0;
+  if (!skipMetricRow) {
+    args->collTest->getBw(count, wordSize(type), deltaSec, &algBw, &busBw, args->nProcs*args->nThreads*args->nGpus);
+  } else if (in_place == 0) {
+    writeBenchMarkLineNullBody();
+  }
 
   Barrier(args);
 
@@ -1147,19 +1256,22 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
       if (wrongElts) break;
   }
 
-  double timeUsec = (report_cputime ? cputimeSec : deltaSec)*1.0E6;
-  writeBenchmarkLineBody(timeUsec, algBw, busBw, args->reportErrors, wrongElts, report_cputime, report_timestamps, in_place==0);
+  double timeUsec = 0.0;
+  if (!skipMetricRow) {
+    timeUsec = (report_cputime ? cputimeSec : deltaSec)*1.0E6;
+    writeBenchmarkLineBody(timeUsec, algBw, busBw, args->reportErrors, wrongElts, report_cputime, report_timestamps, in_place==0);
 
-  auto largestMessageSize = std::max(args->sendBytes, args->expectedBytes);
-  if (args->reporter) {
-    if (args->reportErrors) {
-      args->reporter->addResult((args->nThreads * args->nGpus), args->nProcs, args->totalProcs, largestMessageSize, in_place, timeUsec, algBw, busBw, wrongElts);
-    } else {
-      args->reporter->addResult((args->nThreads * args->nGpus), args->nProcs, args->totalProcs, largestMessageSize, in_place, timeUsec, algBw, busBw);
+    auto largestMessageSize = std::max(args->sendBytes, args->expectedBytes);
+    if (args->reporter) {
+      if (args->reportErrors) {
+        args->reporter->addResult((args->nThreads * args->nGpus), args->nProcs, args->totalProcs, largestMessageSize, in_place, timeUsec, algBw, busBw, wrongElts);
+      } else {
+        args->reporter->addResult((args->nThreads * args->nGpus), args->nProcs, args->totalProcs, largestMessageSize, in_place, timeUsec, algBw, busBw);
+      }
     }
   }
 
-  if (record && args->ms) {
+  if (!skipMetricRow && record && args->ms) {
     // Extract max-across-GPUs per iteration (convert ms to seconds)
     double* iterTimes = (double*)calloc(iters, sizeof(double));
     for (int iter = 0; iter < iters; iter++) {
@@ -1181,10 +1293,12 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
     if (allProcessTimes && nProcessRows > 1)
       ReduceProcessMaxIterTimes(iterTimes, allProcessTimes, nProcessRows, iters);
 
-    int summaryIters = iters - per_iter_skip;
-    struct IterStats stats;
-    computeIterStats(iterTimes + per_iter_skip, summaryIters, &stats);
-    writePerIterReport(&stats, iterTimes, iters, per_iter_skip, in_place==0, allProcessTimes, nProcessRows);
+    if (rccl_tests_devtime::shouldComputeIterStats(iters, per_iter_skip)) {
+      const int summaryIters = iters - per_iter_skip;
+      struct IterStats stats;
+      computeIterStats(iterTimes + per_iter_skip, summaryIters, &stats);
+      writePerIterReport(&stats, iterTimes, iters, per_iter_skip, in_place==0, allProcessTimes, nProcessRows);
+    }
 
     if (in_place && report_timestamps) writeTimestamp();
 
@@ -1193,8 +1307,17 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
     free(args->ms);
   }
 
-  args->bw[0] += busBw;
-  args->bw_count[0]++;
+  if (!skipMetricRow) {
+    args->bw[0] += busBw;
+    args->bw_count[0]++;
+  }
+
+  // Mode 1 (augment): print an extra device-only line alongside the normal
+  // graph/hipEvent numbers above. Mode 2 already reported device time as THE
+  // metric before getBw, so no extra line here.
+  if (in_place == 0 && devTimeMode == 1) {
+    TESTCHECK(args->collTest->deviceTime(args, type, op, root, in_place, nullptr));
+  }
   return testSuccess;
 }
 
@@ -1299,6 +1422,9 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char*
     }
     for (size_t size = args->minbytes; size<=args->maxbytes; size = ((args->stepfactor > 1) ? size*args->stepfactor : size+args->stepbytes)) {
       setupArgs(size, type, args);
+#if defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7)
+      args->devtimeAugmentLine[0] = '\0';
+#endif
       writeBenchmarkLinePreamble(std::max(args->sendBytes, args->expectedBytes), args->nbytes / wordSize(type), typeName, opName, root);
       if (enable_out_of_place) {
         TESTCHECK(BenchTime(args, type, op, root, 0));
@@ -1348,6 +1474,12 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char*
         }
       }
       writeBenchmarkLineTerminator(iters, "");
+#if defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7)
+      if (args->devtimeAugmentLine[0] != '\0') {
+        PRINT("%s", args->devtimeAugmentLine);
+        args->devtimeAugmentLine[0] = '\0';
+      }
+#endif
     }
     --repeat;
     ++iter;
@@ -1817,13 +1949,21 @@ int main(int argc, char* argv[], char **envp) {
     {"output_algo_proto_channels", required_argument, 0, 'A'},      //RCCL (changed from M)
     {"per_iter_timing", required_argument, 0, 'I'},
     {"per_iter_skip", required_argument, 0, 'K'},
+    {"device_timing", required_argument, 0, 'B'},
+    {"devtime_loop", required_argument, 0, 'L'},
+    {"devtime_skip", required_argument, 0, 'P'},
+    {"devtime_loop_mid", required_argument, 0, 'W'},
+    {"devtime_loop_large", required_argument, 0, 'Q'},
+    {"devtime_skip_mid", required_argument, 0, 'j'},
+    {"devtime_skip_large", required_argument, 0, 'k'},
+    {"devtime_check", required_argument, 0, 'H'},
     {"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:Y:U:O:q:F:E:Z:X:A:", 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:Y:U:O:q:F:E:Z:X:A:B:L:P:W:Q:H:j:k:", longopts, &longindex);
 
     if (c == -1)
       break;
@@ -2013,6 +2153,35 @@ int main(int argc, char* argv[], char **envp) {
           return -1;
         }
         break;
+      case 'B':
+        deviceTimingMode = (int)strtol(optarg, NULL, 0);
+        if (deviceTimingMode < 0 || deviceTimingMode > 2) {
+          fprintf(stderr, "device_timing (-B) must be 0 (off), 1 (augment), or 2 (device-only), got %d\n",
+                  deviceTimingMode);
+          return -1;
+        }
+        break;
+      case 'L':
+        devtimeLoop = (int)strtol(optarg, NULL, 0);
+        break;
+      case 'P':
+        devtimeSkip = (int)strtol(optarg, NULL, 0);
+        break;
+      case 'W':
+        devtimeLoopMid = (int)strtol(optarg, NULL, 0);
+        break;
+      case 'Q':
+        devtimeLoopLarge = (int)strtol(optarg, NULL, 0);
+        break;
+      case 'j':
+        devtimeSkipMid = (int)strtol(optarg, NULL, 0);
+        break;
+      case 'k':
+        devtimeSkipLarge = (int)strtol(optarg, NULL, 0);
+        break;
+      case 'H':
+        devtimeCheck = (int)strtol(optarg, NULL, 0);
+        break;
       case 'u':
         unalign = (int)strtol(optarg, NULL, 0);
         break;
@@ -2069,6 +2238,14 @@ 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"
+            "[-B,--device_timing <0/1/2> in-kernel device timing: 0=off, 1=augment (extra line), 2=device-only metric (default: 0)] \n\t"
+            "[-L,--devtime_loop <count> timed-kernel loop iterations (default: 10)] \n\t"
+            "[-P,--devtime_skip <count> timed-kernel warmup skip iterations (default: 10)] \n\t"
+            "[-W,--devtime_loop_mid <count> loop at per-peer >= 8 MiB (0=use -L; default: 0)] \n\t"
+            "[-Q,--devtime_loop_large <count> loop at per-peer >= 64 MiB (0=use tier below; default: 0)] \n\t"
+            "[-j,--devtime_skip_mid <count> skip at mid tier (-1=min(-P,2); default: -1)] \n\t"
+            "[-k,--devtime_skip_large <count> skip at large tier (-1=min(-P,1); default: -1)] \n\t"
+            "[-H,--devtime_check <0/1> validate timed-kernel output before datacheck (default: 0)] \n\t"
             "[-h,--help]\n",
           programName);
         return 0;
@@ -2575,7 +2752,9 @@ testResult_t run() {
   }
   envstr = getenv("NCCL_TESTS_MIN_BW");
   const double check_avg_bw = envstr ? atof(envstr) : -1;
-  bw[0] /= bw_count[0];
+  if (bw_count[0] > 0) {
+    bw[0] /= bw_count[0];
+  }
 
   writeResultFooter(errors.data(), bw.data(), check_avg_bw, programName);
   if (memory_report) {
diff --git a/projects/rccl-tests/src/common.h b/projects/rccl-tests/src/common.h
index aa6474dfd7b3..4fa7298517ce 100644
--- a/projects/rccl-tests/src/common.h
+++ b/projects/rccl-tests/src/common.h
@@ -121,6 +121,19 @@ struct testColl {
   // does not wire it or the library lacks the symbol (older librccl).
   testResult_t (*getCollImplInfo)(ncclComm_t comm, size_t count, ncclDataType_t type, ncclRedOp_t op,
       const void* sendbuff, void* recvbuff, int graphCapturing, int* algo, int* proto, int* nchannels);
+  // Optional device-side (in-kernel wall_clock64) timing hook. Non-null only for
+  // collectives that implement it (currently AllToAll). Driven by BenchTime via
+  // --device_timing (0=off, 1=augment, 2=device-time-only):
+  //   - outDeltaSec == nullptr (mode 1): prints an extra device-only
+  //     latency/busbw line alongside the normal graph/hipEvent numbers (report,
+  //     not replace).
+  //   - outDeltaSec != nullptr (mode 2): stores the measured per-iteration
+  //     device latency (seconds, max across ranks) in *outDeltaSec; BenchTime
+  //     then reports that as THE metric, having skipped the graph/hipEvent
+  //     timed loop.
+  // Other collectives leave it nullptr via aggregate initialization, so no other
+  // struct needs to change.
+  testResult_t (*deviceTime)(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int in_place, double* outDeltaSec);
 };
 extern struct testColl allReduceTest;
 extern struct testColl allGatherTest;
@@ -225,6 +238,11 @@ struct threadArgs {
   void** recvRegHandles;
   void** biasRegHandles;
 #endif
+#if defined(ENABLE_DEVICE_API) && NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7)
+  // Mode-1 device-timing augment line (#[a2a-devtime]); buffered here and
+  // flushed after writeBenchmarkLineTerminator so it does not split the row.
+  char devtimeAugmentLine[512];
+#endif
 };
 
 typedef testResult_t (*threadFunc_t)(struct threadArgs* args);
@@ -237,6 +255,11 @@ struct testThread {
 
 // Provided by common.cu
 extern void Barrier(struct threadArgs* args);
+extern testResult_t testStreamSynchronize(int ngpus, cudaStream_t* streams, ncclComm_t* comms);
+// Inter-thread/process reduce: average 0=bcast(r0),1=avg,2=min,3=max,4=sum.
+// Serializes MPI to the last thread (MPI is MPI_THREAD_SINGLE) and writes the
+// combined value back to every thread. Instantiated for double / long long.
+template<typename T> void Allreduce(struct threadArgs* args, T* value, int average);
 extern testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char* typeName, ncclRedOp_t op,  const char* opName, int root);
 extern testResult_t InitDataReduce(void* data, const size_t count, const size_t offset, ncclDataType_t type, ncclRedOp_t op, const uint64_t seed, const int nranks);
 extern testResult_t InitDataApplyBias(void* expected, void* bias, const size_t count, const size_t offset, ncclDataType_t type, ncclRedOp_t op);
@@ -333,6 +356,17 @@ typedef enum { ncclCoarse        = 0,
                nccl_NUM_MTYPES   = 4 } ncclMemoryType_t;
 extern const char *test_memorytypes[nccl_NUM_MTYPES];
 extern int deviceCtaCount; // number of CTAs for device implementation
+extern int deviceImpl;     // selected -D device implementation (0 = host); lets per-collective
+                           // device-timing hooks pick the matching timed kernel.
+// In-kernel device timing (AllToAll -D 3/4): --device_timing 0=off, 1=augment, 2=device-only.
+extern int deviceTimingMode;
+extern int devtimeLoop;       // --devtime_loop (default 10)
+extern int devtimeSkip;       // --devtime_skip (default 10)
+extern int devtimeLoopMid;    // --devtime_loop_mid: loop at per-peer >= 8 MiB (0=disabled)
+extern int devtimeLoopLarge;  // --devtime_loop_large: loop at per-peer >= 64 MiB (0=disabled)
+extern int devtimeSkipMid;    // --devtime_skip_mid (-1: min(base skip, 2))
+extern int devtimeSkipLarge;  // --devtime_skip_large (-1: min(base skip, 1))
+extern int devtimeCheck;      // --devtime_check: validate timed-kernel output
 constexpr int test_opNumMax = (int)ncclNumOps + (NCCL_VERSION_CODE >= NCCL_VERSION(2,11,0) ? 1 : 0);
 extern int test_opnum;
 extern int test_typenum;
diff --git a/projects/rccl-tests/src/gin_sdma_devtime.h b/projects/rccl-tests/src/gin_sdma_devtime.h
new file mode 100644
index 000000000000..d1f5a5b54b57
--- /dev/null
+++ b/projects/rccl-tests/src/gin_sdma_devtime.h
@@ -0,0 +1,135 @@
+/*************************************************************************
+ * Copyright (c) 2024, Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+// Shared device-side timing scaffold for the GIN-SDMA collectives (rocSHMEM
+// methodology). A collective provides (1) a persistent "*TimedKernel" that runs
+// skip+loop back-to-back collective bodies in ONE launch and brackets only the
+// timed region with the GPU fixed-frequency wall clock (wall_clock64()) -- every
+// CTA stamps start_time[blockIdx.x] at i==skip and end_time[blockIdx.x] after the
+// final iteration -- and (2) a thin "*DeviceTime" hook that calls gin_devtime::measure
+// with a lambda that launches that kernel. measure() owns everything mechanical and
+// identical across collectives: sample each GPU's wall-clock rate, allocate the
+// per-CTA start/end stamp buffers, run the launch, reduce the grid busy window
+// (min(start)..max(end) over CTAs) and the slowest rank (MPI MAX), and return the
+// per-iteration device latency in microseconds. This is the pure device-function
+// execution time -- it excludes host launch, stream/graph, and teardown overhead.
+//
+// The reported span is driven by BenchTime (see common.cu) via --device_timing:
+// 1=augment (each hook prints its own extra "#[*-devtime]" line next to the graph
+// numbers), 2=device-time-only (the hook returns per-iteration seconds via
+// outDeltaSec and BenchTime reports it AS the out-of-place metric).
+#ifndef GIN_SDMA_DEVTIME_H_
+#define GIN_SDMA_DEVTIME_H_
+
+#include <vector>
+#include "common.h"
+#include "gin_sdma_devtime_host.h"
+
+namespace gin_devtime {
+
+// GPU fixed-frequency wall-clock rate (kHz) for a specific device ordinal. Sampled
+// per-GPU (rather than once from the current device) so each GPU's cycle delta is
+// converted with its own rate on mixed-clock nodes.
+static inline bool wallClockRateKhz(int dev, int* outKhz) {
+  int khz = 0;
+  if (hipDeviceGetAttribute(&khz, hipDeviceAttributeWallClockRate, dev) != hipSuccess || khz <= 0)
+    return false;
+  *outKhz = khz;
+  return true;
+}
+
+// Run the device-timing measurement. `launch` is a callable
+//   void(int i, long long* d_start, long long* d_end)
+// that launches the collective's persistent timed kernel on GPU i's stream with a
+// <<<gridCtas, ...>>> grid (loop/skip captured by the caller), writing per-CTA
+// start/end wall-clock stamps into d_start/d_end. measure() sets the device before
+// each launch, sizes the stamp buffers to gridCtas, synchronizes, reduces
+// min(start)..max(end) over CTAs into a per-iteration latency (divided by loop),
+// takes the slowest local GPU, then MPI-MAX across ranks. Cross-rank barriers before
+// and after isolate the timing run from the surrounding measurement flow so it can
+// never race the next size's buffer re-init. Returns the per-iteration latency (us)
+// in *outUs.
+template <typename LaunchFn>
+static inline testResult_t measure(struct threadArgs* args, int gridCtas, int loop,
+                                   LaunchFn&& launch, double* outUs) {
+  if (gridCtas < 1) gridCtas = 1;
+  double localMaxUs = 0.0;
+
+  std::vector<long long*> d_start(args->nGpus, nullptr);
+  std::vector<long long*> d_end(args->nGpus, nullptr);
+
+  Barrier(args);
+
+  // Pass 1: allocate stamp buffers on every local GPU, then launch the timed
+  // kernel on every local GPU, before synchronizing any of them. Malloc-first
+  // keeps the error path safe: if GPU i's hipMalloc fails, no kernel is in
+  // flight yet. Launch-all-then-sync avoids the -g N>1 device-barrier deadlock
+  // (GPU 0 would spin waiting for GPUs never launched). Mirrors startColl()/
+  // completeColl().
+  for (int i = 0; i < args->nGpus; i++) {
+    CUDACHECK(hipSetDevice(args->gpus[i]));
+    CUDACHECK(hipMalloc(&d_start[i], (size_t)gridCtas * sizeof(long long)));
+    CUDACHECK(hipMalloc(&d_end[i], (size_t)gridCtas * sizeof(long long)));
+  }
+  for (int i = 0; i < args->nGpus; i++) {
+    CUDACHECK(hipSetDevice(args->gpus[i]));
+    launch(i, d_start[i], d_end[i]);
+  }
+
+  // Pass 2: synchronize each GPU, copy stamps back, reduce the grid busy window.
+  for (int i = 0; i < args->nGpus; i++) {
+    CUDACHECK(hipSetDevice(args->gpus[i]));
+    TESTCHECK(testStreamSynchronize(1, &args->streams[i], &args->comms[i]));
+
+    // Sample THIS GPU's wall-clock rate to convert its own cycle delta; on a
+    // mixed-clock node a single reference rate would skew the non-reference GPUs.
+    int rate = 0;
+    if (!wallClockRateKhz(args->gpus[i], &rate)) {
+      CUDACHECK(hipFree(d_start[i]));
+      CUDACHECK(hipFree(d_end[i]));
+      d_start[i] = d_end[i] = nullptr;
+      for (int j = i + 1; j < args->nGpus; j++) {
+        if (d_start[j]) CUDACHECK(hipFree(d_start[j]));
+        if (d_end[j]) CUDACHECK(hipFree(d_end[j]));
+        d_start[j] = d_end[j] = nullptr;
+      }
+      if (args->proc == 0 && args->thread == 0) {
+        fprintf(stderr, "ERROR: hipDeviceAttributeWallClockRate query failed for GPU %d\n",
+                args->gpus[i]);
+      }
+      return testInternalError;
+    }
+
+    std::vector<long long> h_start(gridCtas), h_end(gridCtas);
+    CUDACHECK(hipMemcpy(h_start.data(), d_start[i], (size_t)gridCtas * sizeof(long long), hipMemcpyDeviceToHost));
+    CUDACHECK(hipMemcpy(h_end.data(), d_end[i], (size_t)gridCtas * sizeof(long long), hipMemcpyDeviceToHost));
+    CUDACHECK(hipFree(d_start[i]));
+    CUDACHECK(hipFree(d_end[i]));
+
+    const long long mn = rccl_tests_devtime::reduceMinStart(h_start.data(), gridCtas);
+    const long long mx = rccl_tests_devtime::reduceMaxEnd(h_end.data(), gridCtas);
+    const double perIterUs =
+        rccl_tests_devtime::gridBusyWindowPerIterUs(mn, mx, rate, loop);
+    if (perIterUs > localMaxUs) localMaxUs = perIterUs;
+  }
+  Barrier(args);
+
+  // Combine across threads AND ranks through the harness helper. A direct
+  // MPI_Allreduce on MPI_COMM_WORLD here would be wrong under -t N>1: MPI is
+  // initialized MPI_THREAD_SINGLE, and all N BenchTime threads would issue
+  // concurrent collectives on the same communicator (can hang). Allreduce()
+  // accumulates across the process's threads, has only the last thread touch MPI,
+  // and broadcasts the reduced value back -- which also fixes localMaxUs (per-GPU
+  // max on this thread) not being combined across threads.
+  double devUs = localMaxUs;
+  Allreduce(args, &devUs, /*max*/3);
+  *outUs = devUs;
+  return testSuccess;
+}
+
+}  // namespace gin_devtime
+
+#endif  // GIN_SDMA_DEVTIME_H_
diff --git a/projects/rccl-tests/src/gin_sdma_devtime_host.h b/projects/rccl-tests/src/gin_sdma_devtime_host.h
new file mode 100644
index 000000000000..d06bf2b452b9
--- /dev/null
+++ b/projects/rccl-tests/src/gin_sdma_devtime_host.h
@@ -0,0 +1,85 @@
+/*************************************************************************
+ * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#ifndef RCCL_TESTS_GIN_SDMA_DEVTIME_HOST_H_
+#define RCCL_TESTS_GIN_SDMA_DEVTIME_HOST_H_
+
+#include <cstdio>
+#include <string>
+
+// Pure host helpers shared by BenchTime (common.cu), gin_sdma_devtime.h, util.cu,
+// and gin_sdma_devtime_host_test.cpp. Keeps zero-time guards, stamp-reduction
+// math, and skipped-row output contract testable without GPU/MPI.
+
+namespace rccl_tests_devtime {
+
+// Out-of-place column filler when BenchTime skipMetricRow is set (matches
+// writeBenchMarkLineNullBody() in util.cu). Preserves row layout so the
+// in-place pass lands in the in-place column, not the out-of-place slot.
+inline constexpr char kSkippedOutOfPlaceColumnFiller[] =
+    "                                ";
+inline constexpr char kSkippedOutOfPlaceJsonKey[] = "out_of_place";
+inline constexpr std::size_t kSkippedOutOfPlaceColumnFillerLen =
+    sizeof(kSkippedOutOfPlaceColumnFiller) - 1;
+
+inline void emitSkippedOutOfPlaceColumnToStdout() {
+  std::fputs(kSkippedOutOfPlaceColumnFiller, stdout);
+}
+
+inline std::string skippedOutOfPlaceJsonNullFragment() {
+  return std::string("\"") + kSkippedOutOfPlaceJsonKey + "\":null";
+}
+
+inline std::string goldenSkippedRowStdoutFragment(const char* afterPreamble,
+                                                  const char* inPlaceBody) {
+  return std::string(afterPreamble) + kSkippedOutOfPlaceColumnFiller + inPlaceBody;
+}
+
+inline double normalizeElapsedPerIter(double elapsedSec, int iters, int aggIters) {
+  const double iterScale = (double)iters * aggIters;
+  return (iterScale > 0.0) ? elapsedSec / iterScale : 0.0;
+}
+
+inline double applyCudaGraphLaunchesScale(double deltaSec, int cudaGraphLaunches) {
+  if (cudaGraphLaunches >= 1) {
+    return deltaSec / (double)cudaGraphLaunches;
+  }
+  return deltaSec;
+}
+
+inline bool shouldSkipBenchTimeRow(double deltaSec) { return deltaSec <= 0.0; }
+
+inline bool shouldComputeIterStats(int iters, int perIterSkip) {
+  return (iters - perIterSkip) > 0;
+}
+
+inline long long reduceMinStart(const long long* starts, int gridCtas) {
+  long long mn = starts[0];
+  for (int c = 1; c < gridCtas; c++)
+    if (starts[c] < mn) mn = starts[c];
+  return mn;
+}
+
+inline long long reduceMaxEnd(const long long* ends, int gridCtas) {
+  long long mx = ends[0];
+  for (int c = 1; c < gridCtas; c++)
+    if (ends[c] > mx) mx = ends[c];
+  return mx;
+}
+
+// Convert a grid busy window (min start .. max end wall-clock cycles) to
+// per-iteration latency in microseconds. Matches gin_devtime::measure().
+inline double gridBusyWindowPerIterUs(long long minStart, long long maxEnd,
+                                      int wallClockRateKhz, int loop) {
+  if (wallClockRateKhz <= 0) return 0.0;
+  const double totalUs =
+      (double)(maxEnd - minStart) / (double)wallClockRateKhz * 1.0e3;
+  return (loop > 0) ? totalUs / (double)loop : totalUs;
+}
+
+}  // namespace rccl_tests_devtime
+
+#endif  // RCCL_TESTS_GIN_SDMA_DEVTIME_HOST_H_
diff --git a/projects/rccl-tests/src/util.cu b/projects/rccl-tests/src/util.cu
index 0f13c26da0a8..a02ab0d9722a 100644
--- a/projects/rccl-tests/src/util.cu
+++ b/projects/rccl-tests/src/util.cu
@@ -18,6 +18,7 @@
 
 #include "rccl/rccl.h"
 #include "util.h"
+#include "gin_sdma_devtime_host.h"
 #include "git_version.h"
 #include "os.h"
 #include <assert.h>
@@ -479,9 +480,10 @@ void writeBenchmarkLineTerminator(int actualIters, const char *name) {
 
 // Handle a cases where we don't write out of place results
 void writeBenchMarkLineNullBody() {
-  PRINT("                                ");  // only do in-place for trace replay
+  rccl_tests_devtime::emitSkippedOutOfPlaceColumnToStdout();
   if(write_json) {
-    jsonKey("out_of_place"); jsonNull();
+    jsonKey(rccl_tests_devtime::kSkippedOutOfPlaceJsonKey);
+    jsonNull();
   }
 }
 
diff --git a/projects/rccl-tests/test/CMakeLists.txt b/projects/rccl-tests/test/CMakeLists.txt
new file mode 100644
index 000000000000..ac07bcec84a6
--- /dev/null
+++ b/projects/rccl-tests/test/CMakeLists.txt
@@ -0,0 +1,40 @@
+# Copyright (c) 2026 Advanced Micro Devices, Inc.
+
+# CPU-only host unit tests for rccl-tests helpers (no GPU/MPI).
+# Parent CMakeLists.txt skips this directory in TheRock super-project builds.
+if(DEFINED THEROCK_STAGE_INSTALL_ROOT)
+  message(STATUS "rccl-tests: host unit tests unavailable in TheRock build")
+  return()
+endif()
+
+if(NOT BUILD_RCCL_TESTS_HOST_UNIT_TESTS)
+  return()
+endif()
+
+find_package(GTest QUIET)
+if(GTest_FOUND)
+  set(_RCCL_TESTS_GTEST_LIB GTest::gtest)
+  set(_RCCL_TESTS_GTEST_MAIN GTest::gtest_main)
+else()
+  include(FetchContent)
+  message(STATUS "rccl-tests: downloading GTest via FetchContent")
+  FetchContent_Declare(googletest
+    GIT_REPOSITORY https://github.com/google/googletest.git
+    GIT_TAG        v1.14.0
+    GIT_SHALLOW    TRUE
+  )
+  FetchContent_MakeAvailable(googletest)
+  set(_RCCL_TESTS_GTEST_LIB gtest)
+  set(_RCCL_TESTS_GTEST_MAIN gtest_main)
+endif()
+
+add_executable(rccl-tests-GinSdmaDevtimeHostUnit gin_sdma_devtime_host_test.cpp)
+target_include_directories(rccl-tests-GinSdmaDevtimeHostUnit PRIVATE
+  ${CMAKE_CURRENT_SOURCE_DIR}/../src)
+target_link_libraries(rccl-tests-GinSdmaDevtimeHostUnit PRIVATE
+  ${_RCCL_TESTS_GTEST_LIB} ${_RCCL_TESTS_GTEST_MAIN})
+set_target_properties(rccl-tests-GinSdmaDevtimeHostUnit PROPERTIES
+  RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/test")
+add_test(NAME rccl-tests-GinSdmaDevtimeHostUnit
+         COMMAND rccl-tests-GinSdmaDevtimeHostUnit)
+set_tests_properties(rccl-tests-GinSdmaDevtimeHostUnit PROPERTIES LABELS "unit")
diff --git a/projects/rccl-tests/test/gin_sdma_devtime_host_test.cpp b/projects/rccl-tests/test/gin_sdma_devtime_host_test.cpp
new file mode 100644
index 000000000000..337b75913d0b
--- /dev/null
+++ b/projects/rccl-tests/test/gin_sdma_devtime_host_test.cpp
@@ -0,0 +1,99 @@
+/*************************************************************************
+ * Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+// Host unit tests for rccl-tests device-timing helpers in gin_sdma_devtime_host.h
+// (BenchTime zero-time guards, gin_sdma_devtime.h stamp-reduction math, and
+// skipped-row stdout/JSON layout contract shared with util.cu).
+// No GPU or MPI required.
+
+#include <gtest/gtest.h>
+
+#include <string>
+
+#include "gin_sdma_devtime_host.h"
+
+using rccl_tests_devtime::applyCudaGraphLaunchesScale;
+using rccl_tests_devtime::emitSkippedOutOfPlaceColumnToStdout;
+using rccl_tests_devtime::goldenSkippedRowStdoutFragment;
+using rccl_tests_devtime::gridBusyWindowPerIterUs;
+using rccl_tests_devtime::kSkippedOutOfPlaceColumnFiller;
+using rccl_tests_devtime::kSkippedOutOfPlaceColumnFillerLen;
+using rccl_tests_devtime::normalizeElapsedPerIter;
+using rccl_tests_devtime::reduceMaxEnd;
+using rccl_tests_devtime::reduceMinStart;
+using rccl_tests_devtime::shouldComputeIterStats;
+using rccl_tests_devtime::shouldSkipBenchTimeRow;
+using rccl_tests_devtime::skippedOutOfPlaceJsonNullFragment;
+
+TEST(GinSdmaDevtimeHost, NormalizeElapsedPerIter) {
+  EXPECT_DOUBLE_EQ(normalizeElapsedPerIter(2.0, 20, 1), 0.1);
+  EXPECT_DOUBLE_EQ(normalizeElapsedPerIter(1.5, 10, 2), 0.075);
+  EXPECT_DOUBLE_EQ(normalizeElapsedPerIter(1.0, 0, 1), 0.0);
+  EXPECT_DOUBLE_EQ(normalizeElapsedPerIter(1.0, 1, 0), 0.0);
+}
+
+TEST(GinSdmaDevtimeHost, ApplyCudaGraphLaunchesScale) {
+  EXPECT_DOUBLE_EQ(applyCudaGraphLaunchesScale(0.4, 0), 0.4);
+  EXPECT_DOUBLE_EQ(applyCudaGraphLaunchesScale(0.4, 4), 0.1);
+  EXPECT_DOUBLE_EQ(applyCudaGraphLaunchesScale(0.4, -1), 0.4);
+}
+
+TEST(GinSdmaDevtimeHost, ShouldSkipBenchTimeRow) {
+  EXPECT_TRUE(shouldSkipBenchTimeRow(0.0));
+  EXPECT_TRUE(shouldSkipBenchTimeRow(-1e-12));
+  EXPECT_FALSE(shouldSkipBenchTimeRow(1e-9));
+}
+
+TEST(GinSdmaDevtimeHost, ShouldComputeIterStats) {
+  EXPECT_TRUE(shouldComputeIterStats(20, 0));
+  EXPECT_TRUE(shouldComputeIterStats(20, 19));
+  EXPECT_FALSE(shouldComputeIterStats(20, 20));
+}
+
+TEST(GinSdmaDevtimeHost, GridBusyWindowReduction) {
+  const long long starts[] = {100, 50, 80};
+  const long long ends[] = {250, 300, 200};
+  EXPECT_EQ(reduceMinStart(starts, 3), 50);
+  EXPECT_EQ(reduceMaxEnd(ends, 3), 300);
+  // (300-50)/100 kHz * 1e3 = 2500 us total; loop=10 -> 250 us/iter
+  EXPECT_DOUBLE_EQ(gridBusyWindowPerIterUs(50, 300, 100, 10), 250.0);
+  EXPECT_DOUBLE_EQ(gridBusyWindowPerIterUs(50, 300, 100, 0), 2500.0);
+}
+
+TEST(GinSdmaDevtimeHost, SkippedOutOfPlaceStdoutGolden) {
+  testing::internal::CaptureStdout();
+  emitSkippedOutOfPlaceColumnToStdout();
+  const std::string out = testing::internal::GetCapturedStdout();
+  EXPECT_EQ(out, kSkippedOutOfPlaceColumnFiller);
+  EXPECT_EQ(out.size(), kSkippedOutOfPlaceColumnFillerLen);
+  EXPECT_EQ(kSkippedOutOfPlaceColumnFillerLen, 32u);
+}
+
+TEST(GinSdmaDevtimeHost, SkippedOutOfPlaceRowLayoutGolden) {
+  // Simulates preamble + skipped OOP metric + in-place body on the same row.
+  // BenchTime must emit the filler (not return early) so in-place metrics stay
+  // in the in-place column and JSON retains out_of_place:null.
+  const char* preamble =
+      "        4096          1024   float     sum       0";
+  const char* inPlaceBody = "   12.34   1.23   2.46    N/A";
+  const std::string row = goldenSkippedRowStdoutFragment(preamble, inPlaceBody);
+
+  const size_t fillerPos = row.find(kSkippedOutOfPlaceColumnFiller);
+  const size_t inPlacePos = row.find(inPlaceBody);
+  ASSERT_NE(fillerPos, std::string::npos);
+  ASSERT_NE(inPlacePos, std::string::npos);
+  EXPECT_LT(fillerPos, inPlacePos);
+  EXPECT_EQ(row.substr(fillerPos, kSkippedOutOfPlaceColumnFillerLen),
+            kSkippedOutOfPlaceColumnFiller);
+
+  const std::string json = std::string("{\"size\":4096,") +
+                           skippedOutOfPlaceJsonNullFragment() + ",\"in_place\":{";
+  const size_t oopJsonPos = json.find(skippedOutOfPlaceJsonNullFragment());
+  const size_t ipJsonPos = json.find("\"in_place\"");
+  ASSERT_NE(oopJsonPos, std::string::npos);
+  ASSERT_NE(ipJsonPos, std::string::npos);
+  EXPECT_LT(oopJsonPos, ipJsonPos);
+}
diff --git a/projects/rccl-tests/test/test_AllToAllDevtime.py b/projects/rccl-tests/test/test_AllToAllDevtime.py
new file mode 100644
index 000000000000..27f803846d4a
--- /dev/null
+++ b/projects/rccl-tests/test/test_AllToAllDevtime.py
@@ -0,0 +1,196 @@
+#################################################################################
+# Copyright (c) 2026 Advanced Micro Devices, Inc. All rights reserved.
+#
+# Permission is hereby granted, free of charge, to any person obtaining a copy
+# of this software and associated documentation files (the "Software"), to deal
+# in the Software without restriction, including without limitation the rights
+# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell cop-
+# ies of the Software, and to permit persons to whom the Software is furnished
+# to do so, subject to the following conditions:
+#
+# The above copyright notice and this permission notice shall be included in all
+# copies or substantial portions of the Software.
+#
+# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR IM-
+# PLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, FITNESS
+# FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR
+# COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER
+# IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNE-
+# CTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+################################################################################
+
+# GPU smoke tests for GIN AllToAll in-kernel device timing (wall_clock64).
+#
+# Exercises the real GinAlltoAllTimedKernel path via alltoall_perf -D 3 with
+# --device_timing modes 1/2 and optional --devtime_check. Requires an MPI
+# launcher, >= 2 GPUs/ranks, and rccl-tests built with ENABLE_DEVICE_API=ON
+# against RCCL >= 2.28.7.
+#
+# Opt-in locally via RCCL_TESTS_GIN_SDMA_DEVTIME=1. The device-api CI job sets
+# that automatically (see projects/rccl/tools/ci/run-device-api-ci.sh).
+#
+# Environment (shared with test_AllToAll.py where noted):
+#   RCCL_TESTS_GIN_SDMA_DEVTIME  enable this module (default: off)
+#   RCCL_TESTS_A2A_NP            MPI ranks (default: detected GPU count)
+#   RCCL_TESTS_MPI_LAUNCHER      launcher binary (default: mpirun)
+#   RCCL_TESTS_MPI_OPTS          extra launcher opts
+#   RCCL_TESTS_A2A_XENV          extra "-x K=V" env beyond the GIN essentials
+#   RCCL_TESTS_A2A_EXE           path to alltoall_perf (default: ../build/...)
+#   RCCL_TESTS_A2A_TIMEOUT_S     per-run timeout seconds (default: 300)
+#   RCCL_TESTS_A2A_CTAS          -V grid CTAs (default: 8)
+#   RCCL_TESTS_A2A_GIN_TYPE      NCCL_GIN_TYPE (default: 2, matches device-api CI)
+
+import os
+import re
+import shlex
+import signal
+import subprocess
+
+import pytest
+
+KiB = 1024
+SMOKE_BYTES = 128 * KiB  # single-size smoke (128 KiB per rank)
+
+path = os.path.dirname(os.path.abspath(__file__))
+executable = os.environ.get(
+    "RCCL_TESTS_A2A_EXE", os.path.join(path, "..", "build", "alltoall_perf"))
+
+_enabled = os.environ.get("RCCL_TESTS_GIN_SDMA_DEVTIME", "") not in (
+    "", "0", "false", "False")
+
+DEVTIME_LINE_RE = re.compile(
+    r"#\[a2a-devtime\].*?\bdevtime\s+([0-9]+(?:\.[0-9]+)?)\s+us")
+
+
+def _detect_ngpus():
+    if os.environ.get("ROCR_VISIBLE_DEVICES") is not None:
+        return len(os.environ["ROCR_VISIBLE_DEVICES"].split(","))
+    if os.environ.get("HIP_VISIBLE_DEVICES") is not None:
+        return len(os.environ["HIP_VISIBLE_DEVICES"].split(","))
+    try:
+        out = subprocess.check_output(
+            'rocminfo | grep "Device Type:.\\s*.GPU" | wc -l', shell=True)
+        return int(out)
+    except Exception:
+        return 0
+
+
+NP = int(os.environ.get("RCCL_TESTS_A2A_NP", "0")) or _detect_ngpus()
+LAUNCHER = os.environ.get("RCCL_TESTS_MPI_LAUNCHER", "mpirun")
+CTAS = os.environ.get("RCCL_TESTS_A2A_CTAS", "8")
+TIMEOUT_S = int(os.environ.get("RCCL_TESTS_A2A_TIMEOUT_S", "300"))
+GIN_TYPE = os.environ.get("RCCL_TESTS_A2A_GIN_TYPE", "2")
+MPI_OPTS = shlex.split(os.environ.get("RCCL_TESTS_MPI_OPTS", ""))
+XENV = shlex.split(os.environ.get("RCCL_TESTS_A2A_XENV", ""))
+
+pytestmark = pytest.mark.skipif(
+    not _enabled,
+    reason="GIN AllToAll devtime smoke tests are opt-in; set "
+           "RCCL_TESTS_GIN_SDMA_DEVTIME=1 on a GIN-capable node to enable.")
+
+
+def _assert_datacheck_clean(out):
+    """Perf binary reports wrong elements via #wrong= or Out of bounds values."""
+    if "#wrong=0" in out:
+        return
+    if "Out of bounds values : 0 OK" in out:
+        return
+    if re.search(r"#wrong=\s*[1-9]", out):
+        pytest.fail("datacheck reported wrong elements:\n{}".format(out[-2000:]))
+    if "Out of bounds values : 0" not in out and "#wrong=" not in out:
+        # Some builds only print the summary line on failure; accept clean exit.
+        return
+
+
+def _run_devtime(request, device_timing_mode, devtime_check=False):
+    """Launch alltoall_perf with GIN (-D 3) and device-timing CLI flags."""
+    if NP < 2:
+        pytest.skip("need >= 2 ranks/GPUs for AllToAll")
+
+    size = str(SMOKE_BYTES)
+    # Match device-api CI gin-d3 essentials; deployment extras via RCCL_TESTS_A2A_XENV.
+    gin_env = []
+    for kv in [
+        "NCCL_CUMEM_ENABLE=1",
+        "HSA_FORCE_FINE_GRAIN_PCIE=1",
+        "NCCL_DMABUF_ENABLE=1",
+        "NCCL_GIN_TYPE={}".format(GIN_TYPE),
+        "HSA_NO_SCRATCH_RECLAIM=1",
+        "NCCL_ENV_PLUGIN=none",
+        "RCCL_ENABLE_INTRANET=1",
+    ] + XENV:
+        gin_env += ["-x", kv]
+
+    hostfile = request.config.getoption("--hostfile")
+    launch = [LAUNCHER, "-np", str(NP)] + MPI_OPTS
+    if hostfile:
+        launch += ["-host", hostfile]
+
+    bench_args = [
+        executable,
+        "-b", size, "-e", size,
+        "-f", "2",
+        "-g", "1",
+        "-R", "2",
+        "-D", "3",
+        "-V", CTAS,
+        "-d", "int32",
+        "-c", "1",
+        "-w", "1",
+        "-n", "3",
+        "-B", str(device_timing_mode),
+        "-L", "5",
+        "-P", "2",
+    ]
+    if devtime_check:
+        bench_args += ["-H", "1"]
+
+    args = launch + gin_env + bench_args
+    cmd = " ".join(shlex.quote(a) for a in args)
+    print(cmd)
+    proc = subprocess.Popen(cmd, shell=True, universal_newlines=True,
+                            stdout=subprocess.PIPE, stderr=subprocess.STDOUT,
+                            start_new_session=True)
+    try:
+        out, _ = proc.communicate(timeout=TIMEOUT_S)
+    except subprocess.TimeoutExpired:
+        try:
+            os.killpg(os.getpgid(proc.pid), signal.SIGKILL)
+        except ProcessLookupError:
+            pass
+        out, _ = proc.communicate()
+        pytest.fail(
+            "AllToAll devtime smoke HANG: no completion within {}s (mode -B {}). "
+            "Output tail:\n{}".format(TIMEOUT_S, device_timing_mode, (out or "")[-2000:]))
+    print(out)
+    return proc.returncode, out
+
+
+def test_AllToAllDevtimeMode1Augment(request):
+    """Mode 1: normal bench plus #[a2a-devtime] from wall_clock64 timed kernel."""
+    rc, out = _run_devtime(request, device_timing_mode=1)
+    assert rc == 0, "alltoall_perf exited {} (mode -B 1)".format(rc)
+    _assert_datacheck_clean(out)
+    match = DEVTIME_LINE_RE.search(out)
+    assert match is not None, (
+        "expected #[a2a-devtime] line in stdout (mode -B 1); tail:\n{}".format(out[-2000:]))
+    devtime_us = float(match.group(1))
+    assert devtime_us > 0.0, "devtime must be positive, got {} us".format(devtime_us)
+
+
+def test_AllToAllDevtimeMode2DeviceOnly(request):
+    """Mode 2: reported metric is in-kernel device latency (no host graph loop)."""
+    rc, out = _run_devtime(request, device_timing_mode=2)
+    assert rc == 0, "alltoall_perf exited {} (mode -B 2)".format(rc)
+    _assert_datacheck_clean(out)
+    assert "WARN --device_timing=2: no in-kernel device-time" not in out, (
+        "device-time-only mode produced no valid measurement:\n{}".format(out[-2000:]))
+
+
+def test_AllToAllDevtimeMode2WithTimedCheck(request):
+    """Mode 2 + --devtime_check: validate timed-kernel output before datacheck."""
+    rc, out = _run_devtime(request, device_timing_mode=2, devtime_check=True)
+    assert rc == 0, "alltoall_perf exited {} (mode -B 2 -H 1)".format(rc)
+    _assert_datacheck_clean(out)
+    assert "ERROR: --devtime_check:" not in out, (
+        "timed-kernel datacheck failed:\n{}".format(out[-2000:]))
diff --git a/projects/rccl/tools/ci/lib/device-api-tests.json b/projects/rccl/tools/ci/lib/device-api-tests.json
index 8644367b1c6f..721fcc14b9d3 100644
--- a/projects/rccl/tools/ci/lib/device-api-tests.json
+++ b/projects/rccl/tools/ci/lib/device-api-tests.json
@@ -3,6 +3,8 @@
     "Device-API benchmark matrix for run-device-api-ci.sh.",
     "Add a mode by appending a suites entry. Each suite runs every bin as:",
     "  mpirun -np $NP <base_env + suite env> -x LD_LIBRARY_PATH ./build/<bin> $bench_args <args>",
+    "After the matrix, run-device-api-ci.sh runs pytest test/test_AllToAllDevtime.py",
+    "(GIN -D 3 device-timing smoke: -B 1/2 and -B 2 -H 1; ENABLE_DEVICE_API build).",
     "bench_args: common bench flags | base_env: -x on every run",
     "debug_env:  -x added only when RCCL_CI_DEBUG=1",
     "  {LOGDIR} in debug_env is replaced by the runner with the per-job debug dir,",
diff --git a/projects/rccl/tools/ci/run-device-api-ci.sh b/projects/rccl/tools/ci/run-device-api-ci.sh
index 71e3d7561d24..aae310f913d8 100755
--- a/projects/rccl/tools/ci/run-device-api-ci.sh
+++ b/projects/rccl/tools/ci/run-device-api-ci.sh
@@ -10,6 +10,8 @@
 #
 # Each bench is wrapped in `timeout` so a hung mpirun/driver can't wedge the job;
 # failures are collected and surfaced at the end (exit non-zero iff any failed).
+# After the JSON matrix, runs pytest test/test_AllToAllDevtime.py (GIN AllToAll
+# in-kernel device-timing smoke: modes -B 1, -B 2, -B 2 -H 1).
 #
 # Test matrix lives in lib/device-api-tests.json; RCCL_CI_DEBUG=1 adds its
 # debug_env to every run.
@@ -168,3 +170,50 @@ if [[ ${#FAILED_RUNS[@]} -ne 0 ]]; then
 fi
 
 echo "All device-api benchmark runs succeeded."
+
+# GIN AllToAll in-kernel device-timing smoke (wall_clock64 / *TimedKernel).
+# Requires ENABLE_DEVICE_API=ON rccl-tests build (alltoall_perf with -B flags).
+run_devtime_smoke() {
+  local pytest_dir="${RCCL_TESTS_DIR}/test"
+  local perf_bin="${RCCL_TESTS_DIR}/${PERF_DIR}/alltoall_perf"
+  if [[ ! -x "${perf_bin}" ]]; then
+    echo "WARN: skip devtime smoke: ${perf_bin} not found"
+    return 0
+  fi
+  if [[ ! -f "${pytest_dir}/test_AllToAllDevtime.py" ]]; then
+    echo "WARN: skip devtime smoke: ${pytest_dir}/test_AllToAllDevtime.py not found"
+    return 0
+  fi
+
+  # Reuse the CI python venv when present; otherwise ensure pytest is importable.
+  # shellcheck source=/dev/null
+  [[ -f "${script_dir}/lib/ensure-python-yaml.sh" ]] && source "${script_dir}/lib/ensure-python-yaml.sh"
+  if ! python3 -c 'import pytest' 2>/dev/null; then
+    echo "==> devtime smoke: pip installing pytest"
+    python3 -m pip install --quiet --disable-pip-version-check pytest
+  fi
+
+  echo "=== devtime-smoke: pytest test_AllToAllDevtime.py (GIN -D 3, -g 1) ==="
+  set +e
+  RCCL_TESTS_GIN_SDMA_DEVTIME=1 \
+  RCCL_TESTS_A2A_EXE="${perf_bin}" \
+  RCCL_TESTS_A2A_NP="${NP}" \
+  RCCL_TESTS_A2A_GIN_TYPE="${RCCL_TESTS_A2A_GIN_TYPE:-2}" \
+  RCCL_TESTS_A2A_TIMEOUT_S="${RCCL_TESTS_A2A_TIMEOUT_S:-300}" \
+    python3 -m pytest "${pytest_dir}/test_AllToAllDevtime.py" -v -p no:cacheprovider
+  local rc=$?
+  set -e
+  if [[ ${rc} -ne 0 ]]; then
+    FAILED_RUNS+=("devtime-smoke:pytest (rc=${rc})")
+  fi
+}
+
+run_devtime_smoke
+
+if [[ ${#FAILED_RUNS[@]} -ne 0 ]]; then
+  echo "=== FAILED RUNS (${#FAILED_RUNS[@]}) ==="
+  printf '  %s\n' "${FAILED_RUNS[@]}"
+  exit 1
+fi
+
+echo "All device-api benchmark and devtime smoke runs succeeded."