diff --git a/.gitignore b/.gitignore
index a0a013e..a37c77a 100644
--- a/.gitignore
+++ b/.gitignore
@@ -2,3 +2,4 @@
 #
 # See LICENCE.txt for license information
 /build
+/build-*
diff --git a/README.md b/README.md
index c0ec663..2fdd37d 100644
--- a/README.md
+++ b/README.md
@@ -32,13 +32,14 @@ NCCL tests can run on multiple processes, multiple threads, and multiple CUDA de
 
 ### Quick examples
 
-Run on single node with 8 GPUs (`-g 8`), scanning from 8 Bytes to 128MBytes :
+Run on single node with 8 GPUs (`-g 8`), scanning from 8 Bytes to 128MiB (Mebibytes), doubling between each test (`-f 2`) :
 
 ```shell
 $ ./build/all_reduce_perf -b 8 -e 128M -f 2 -g 8
 ```
 
-Run 64 MPI processes on nodes with 8 GPUs each, for a total of 64 GPUs spread across 8 nodes :
+Run 64 MPI processes on nodes with 8 GPUs each, for a total of 64 GPUs spread across 8 nodes.
+Scanning from 8 Bytes to 32GiB (Gibibytes), doubling between each test (`-f 2`).
 (NB: The nccl-tests binaries must be compiled with `MPI=1` for this case)
 
 ```shell
@@ -57,10 +58,10 @@ All tests support the same set of arguments :
   * `-t,--nthreads <num threads>` number of threads per process. Default : 1.
   * `-g,--ngpus <GPUs per thread>` number of gpus per thread. Default : 1.
 * Sizes to scan
-  * `-b,--minbytes <min size in bytes>` minimum size to start with. Default : 32M.
-  * `-e,--maxbytes <max size in bytes>` maximum size to end at. Default : 32M.
-  * Increments can be either fixed or a multiplication factor. Only one of those should be used
-    * `-i,--stepbytes <increment size>` fixed increment between sizes. Default : 1M.
+  * `-b,--minbytes <min size in bytes>` minimum size to start with. Default : 32M (Mebibytes).
+  * `-e,--maxbytes <max size in bytes>` maximum size to end at. Default : 32M (Mebibytes).
+  * Increments can be either fixed or a multiplication factor. Only one of those should be used.
+    * `-i,--stepbytes <increment size>` fixed increment between sizes. Default : 1M (Mebibytes).
     * `-f,--stepfactor <increment factor>` multiplication factor between sizes. Default : disabled.
 * NCCL operations arguments
   * `-o,--op <sum/prod/min/max/avg/all>` Specify which reduction operation to perform. Only relevant for reduction operations like Allreduce, Reduce or ReduceScatter. Default : Sum.
@@ -75,12 +76,15 @@ All tests support the same set of arguments :
 * Test operation
   * `-p,--parallel_init <0/1>` use threads to initialize NCCL in parallel. Default : 0.
   * `-c,--check <check iteration count>` perform count iterations, checking correctness of results on each iteration. This can be quite slow on large numbers of GPUs. Default : 1.
-  * `-z,--blocking <0/1>` Make NCCL collective blocking, i.e. have CPUs wait and sync after each collective. Default : 0.
+  * `-z,--blocking <0/1/2>` collective blocking: 1=wait for completion and barrier, 2=wait without barrier. Default : 0.
   * `-G,--cudagraph <num graph launches>` Capture iterations as a CUDA graph and then replay specified number of times. Default : 0.
   * `-C,--report_cputime <0/1>` Report CPU time instead of latency. Default : 0.
   * `-R,--local_register <0/1/2>` enable local (1) or symmetric (2) buffer registration on send/recv buffers. Default : 0.
-  * `-S,--report_timestamps <0/1>` Add timestamp ("%Y-%m-%d %H:%M:%S") to each performance report line. Default : 0.
+  * `-S,--report_timestamps <0/1>` Add timestamp (`"%Y-%m-%d %H:%M:%S"`) to each performance report line. Default : 0.
+  * `-J,--output_file <file>` Write [JSON] output to filepath. Infer type from suffix (only `json` supported presently).
   * `-T,--timeout <time in seconds>` timeout each test after specified number of seconds. Default : disabled.
+  * `-M,--memory_report <0/1>` enable memory usage report. Default : 0.
+  * `-u,--unalign <index of first element>` Misalign source and destination buffers. Default : 0.
 
 ### Running multiple operations in parallel
 
@@ -104,4 +108,4 @@ Note that the reported bandwidth is per group, hence to get the total bandwidth
 
 ## Copyright
 
-NCCL tests are provided under the BSD license. All source code and accompanying documentation is copyright (c) 2016-2025, NVIDIA CORPORATION. All rights reserved.
+NCCL tests are provided under the BSD license. All source code and accompanying documentation is copyright (c) 2016-2026, NVIDIA CORPORATION. All rights reserved.
diff --git a/src/Makefile b/src/Makefile
index c8f787a..2a5ca45 100644
--- a/src/Makefile
+++ b/src/Makefile
@@ -28,7 +28,12 @@ ifeq ($(MPI_IBM),1)
 NVCUFLAGS += -DMPI_SUPPORT
 NVLDFLAGS += -lmpi_ibm
 endif
-LIBRARIES += nccl
+# GPU serials use NVML (libnvidia-ml). Link against the CUDA toolkit's NVML stub so
+# the build succeeds on hosts without the driver's -dev symlink; the real
+# libnvidia-ml is loaded at runtime. libuuid was dropped — the JSON UUID is now
+# generated in-tree (genUuidV4 in util.cu), so no uuid-dev is needed to build.
+NVLDFLAGS += -L$(CUDA_HOME)/lib64/stubs
+LIBRARIES += nccl nvidia-ml
 NVLDFLAGS += $(LIBRARIES:%=-l%)
 
 DST_DIR := $(BUILDDIR)
diff --git a/src/all_reduce.cu b/src/all_reduce.cu
index 755b9e8..260e184 100644
--- a/src/all_reduce.cu
+++ b/src/all_reduce.cu
@@ -64,27 +64,50 @@ void AllReduceGetBw(size_t count, int typesize, double sec, double* algBw, doubl
   *busBw = baseBw * factor;
 }
 
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
-// set devComm reqs for allreduce device kernels
-bool AllReduceGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* reqs) {
-  if (!reqs) return false;
-  memset(reqs, 0, sizeof(*reqs));
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0)
+ // set devComm reqs for allreduce device kernels
+testResult_t AllReduceGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* reqs, ncclCommProperties_t* commProperties) {
+  if (!reqs || !commProperties) return testInternalError;
 
   switch(deviceImpl) {
     case 1: // allReduceLsaKernel
     case 2: // allReduceLsaVectorizedKernel
       reqs->lsaBarrierCount = deviceCtaCount;
-      return true;
+      return testSuccess;
     case 3: // allReduceMultimemKernel
     case 4: // allReduceMultimemVectorizedKernel
+      if (!commProperties->multimemSupport) {
+        fprintf(stderr, "This test requires multimem support, but multimem support is not enabled for this communicator.\n");
+        return testInternalError;
+      }
       reqs->lsaMultimem = true;
       reqs->lsaBarrierCount = deviceCtaCount;
+      return testSuccess;
+    default:
+      return testNotImplemented;
+  }
+}
+#elif NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
+ bool AllReduceGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* reqs) {
+   if (!reqs) return false;
+   memset(reqs, 0, sizeof(*reqs));
+   switch(deviceImpl) {
+    case 1: // allReduceLsaKernel
+    case 2: // allReduceLsaVectorizedKernel
+      reqs->lsaBarrierCount = deviceCtaCount;
+      return true;
+    case 3: // allReduceMultimemKernel
+    case 4: // allReduceMultimemVectorizedKernelMultimem = true;
+    reqs->lsaMultimem = true;
+    reqs->lsaBarrierCount = deviceCtaCount;
       return true;
     default:
       return false;
   }
-}
+ }
+#endif
 
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
 /*
  * Kernel 1: allReduceLsaKernel - Basic LSA-based AllReduce
  *
diff --git a/src/alltoall.cu b/src/alltoall.cu
index 0bc9894..75aaf96 100644
--- a/src/alltoall.cu
+++ b/src/alltoall.cu
@@ -51,7 +51,35 @@ void AlltoAllGetBw(size_t count, int typesize, double sec, double* algBw, double
   *busBw = baseBw * factor;
 }
 
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0)
+// set devComm reqs for alltoall device kernels
+testResult_t AlltoAllGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* reqs, ncclCommProperties_t* commProperties) {
+  if (!reqs || !commProperties) return testInternalError;
+
+  switch(deviceImpl) {
+    case 1: // NvlAlltoAllKernel
+    case 2: // NvlAlltoAllKernelOptimized
+      reqs->lsaBarrierCount = deviceCtaCount;
+      return testSuccess;
+    case 3: // GinAlltoAllKernel
+    case 4: // HybridAlltoAllKernel (LSA+GIN)
+      if (commProperties->ginType == NCCL_GIN_TYPE_NONE) {
+        fprintf(stderr, "This test requires GIN support, but GIN support is not enabled for this communicator.\n");
+        return testInternalError;
+      }
+      reqs->barrierCount = deviceCtaCount;
+      reqs->ginSignalCount = deviceCtaCount;
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2, 29, 7)
+      reqs->ginConnectionType = NCCL_GIN_CONNECTION_FULL;
+#else
+      reqs->ginForceEnable = true;
+#endif
+      return testSuccess;
+    default:
+      return testNotImplemented;
+  }
+}
+#elif NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
 // set devComm reqs for alltoall device kernels
 bool AlltoAllGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* reqs) {
   if (!reqs) return false;
@@ -73,7 +101,9 @@ bool AlltoAllGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* req
       return false;
   }
 }
+#endif
 
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
 // shared scalar AlltoAll implementation used by both kernels
 template <typename T>
 __device__ void AlltoAllScalarImpl(ncclWindow_t sendwin, size_t sendoffset, ncclWindow_t recvwin, size_t recvoffset, size_t count, int rank, int nRanks, int tid, int nthreads) {
diff --git a/src/common.cu b/src/common.cu
index 7a91441..b14c3d5 100644
--- a/src/common.cu
+++ b/src/common.cu
@@ -8,6 +8,7 @@
 #include <pthread.h>
 #include <cstdio>
 #include <type_traits>
+#include <limits>
 #include <getopt.h>
 #include <libgen.h>
 #include <string.h>
@@ -22,6 +23,7 @@
 #pragma weak ncclCommWindowDeregister
 #pragma weak ncclDevCommCreate
 #pragma weak ncclDevCommDestroy
+#pragma weak ncclCommQueryProperties
 
 #define DIVUP(x, y) \
     (((x)+(y)-1)/(y))
@@ -99,6 +101,8 @@ int cudaGraphLaunches = 0;
 static int report_cputime = 0;
 static int report_timestamps = 0;
 static int deviceImpl = 0;
+int unalign = 0;
+int memory_report = 0;
 
 int deviceCtaCount = 16; // Default number of CTAs for device implementation
 
@@ -181,6 +185,35 @@ static void outputFileFinalize(output_file_type_t output_file_type) {
   }
 }
 
+testResult_t initComms(ncclComm_t* comms, int nComms, int firstRank, int nRanks, int* cudaDevs, ncclUniqueId& ncclId) {
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
+  ncclConfig_t config = NCCL_CONFIG_INITIALIZER;
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,27,0)
+  if (ctaPolicy >= 0)
+    config.CTAPolicy = ctaPolicy;
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
+  config.nvlinkCentricSched = 1;
+#endif
+#endif
+  if (ncclTestEngine.initCommConfig)
+    ncclTestEngine.initCommConfig(&config);
+#endif
+
+  NCCLCHECK(ncclGroupStart());
+  for (int i=0; i<nComms; i++) {
+    int rank = firstRank + i;
+    CUDACHECK(cudaSetDevice(cudaDevs[i]));
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
+    NCCLCHECK(ncclCommInitRankConfig(comms+i, nRanks, ncclId, rank, &config));
+#else
+    NCCLCHECK(ncclCommInitRank(comms+i, nRanks, ncclId, rank));
+#endif
+  }
+  NCCLCHECK(ncclGroupEnd());
+  return testSuccess;
+}
+
+// NOTE: We use the binary system, so M=Mebibytes and G=Gibibytes
 static double parsesize(const char *value) {
     long long int units;
     double size;
@@ -509,7 +542,7 @@ testResult_t startColl(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
     // Complete op before returning
     TESTCHECK(testStreamSynchronize(args->nGpus, args->streams, args->comms));
   }
-  if (blocking_coll) Barrier(args);
+  if (blocking_coll == 1) Barrier(args);
   return testSuccess;
 }
 
@@ -686,6 +719,12 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char*
   // Sync to avoid first-call timeout
   Barrier(args);
 
+  // Add forced misalignment
+  for (int i = 0; i < args->nGpus; i++) {
+    args->sendbuffs[i] = (char*)args->sendbuffs[i] + unalign * wordSize(type);
+    args->recvbuffs[i] = (char*)args->recvbuffs[i] + unalign * wordSize(type);
+  }
+
   // Warm-up for all sizes (using a stepfactor of 2)
   for (size_t size = args->minbytes; size <= args->maxbytes; size = size * 2) {
     setupArgs(size, type, args);
@@ -707,21 +746,46 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char*
     }
   } while (--repeat);
 
+  // Revert forced misalignment
+  for (int i = 0; i < args->nGpus; i++) {
+    args->sendbuffs[i] = (char*)args->sendbuffs[i] - unalign * wordSize(type);
+    args->recvbuffs[i] = (char*)args->recvbuffs[i] - unalign * wordSize(type);
+  }
   return testSuccess;
 }
 
+static void getGPUMemoryInfo(int64_t* ptotalGpuMem, int64_t* pfreeGpuMem) {
+  size_t freeGpuMem, totalGpuMem = 0;
+  cudaMemGetInfo(&freeGpuMem, &totalGpuMem);
+  if (ptotalGpuMem != nullptr) *ptotalGpuMem = totalGpuMem;
+  if (pfreeGpuMem != nullptr) *pfreeGpuMem = freeGpuMem;
+}
+
 testResult_t threadRunTests(struct threadArgs* args) {
+  //  capture the free memory before
+  int64_t* totalGpuFreeMem = (int64_t*)calloc(args->nGpus*2, sizeof(int64_t));
+  for (int g = 0; g < args->nGpus; ++g) {
+    CUDACHECK(cudaSetDevice(args->gpus[g]));
+    getGPUMemoryInfo(nullptr, &totalGpuFreeMem[g]);
+  }
+
   // Set device to the first of our GPUs. If we don't do that, some operations
   // will be done on the current GPU (by default : 0) and if the GPUs are in
   // exclusive mode those operations will fail.
   CUDACHECK(cudaSetDevice(args->gpus[0]));
   TESTCHECK(ncclTestEngine.runTest(args, ncclroot, (ncclDataType_t)nccltype, test_typenames[nccltype], (ncclRedOp_t)ncclop, test_opnames[ncclop]));
+
+  // Capture the memory used by the GPUs
+  for (int g = 0; g < args->nGpus; ++g) {
+    CUDACHECK(cudaSetDevice(args->gpus[g]));
+    getGPUMemoryInfo(nullptr, &totalGpuFreeMem[g + args->nGpus]);
+    *args->devMemUsed = std::max(*args->devMemUsed, totalGpuFreeMem[g] - totalGpuFreeMem[g + args->nGpus]);
+  }
+  free(totalGpuFreeMem);
   return testSuccess;
 }
 
 testResult_t threadInit(struct threadArgs* args) {
-  char hostname[1024];
-  getHostName(hostname, 1024);
   int nranks =  args->nProcs*args->nThreads*args->nGpus;
 
   //set main thread again
@@ -729,28 +793,33 @@ testResult_t threadInit(struct threadArgs* args) {
 
   jsonIdentifyWriter(is_main_thread);
 
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
-     ncclConfig_t config = NCCL_CONFIG_INITIALIZER;
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,27,0)
-     if (ctaPolicy >= 0)
-       config.CTAPolicy = ctaPolicy;
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
-     config.nvlinkCentricSched = 1;
-#endif
-#endif
-#endif
+  // Capture GPU memory before initializing the NCCL communicators
+  int64_t* initFreeGpuMem = (int64_t*)calloc(args->nGpus*3, sizeof(int64_t));
+  for (int g = 0; g < args->nGpus; ++g) {
+    CUDACHECK(cudaSetDevice(args->gpus[g]));
+    getGPUMemoryInfo(nullptr, &initFreeGpuMem[g]);
+  }
+
+  int firstRank = args->proc*args->nThreads*args->nGpus + args->thread*args->nGpus;
+  TESTCHECK(initComms(args->comms, args->nGpus, firstRank, nranks, args->gpus, args->ncclId));
 
+  /* Allocate buffers for each GPU (parallel_init: each thread allocates its own) */
+  size_t sendBytes, recvBytes;
+  ncclTestEngine.getBuffSize(&sendBytes, &recvBytes, (size_t)args->maxbytes, (size_t)nranks);
   NCCLCHECK(ncclGroupStart());
-  for (int i=0; i<args->nGpus; i++) {
-    int rank = args->proc*args->nThreads*args->nGpus + args->thread*args->nGpus + i;
+  for (int i = 0; i < args->nGpus; i++) {
     CUDACHECK(cudaSetDevice(args->gpus[i]));
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
-    NCCLCHECK(ncclCommInitRankConfig(args->comms+i, nranks, args->ncclId, rank, &config));
-#else
-    NCCLCHECK(ncclCommInitRank(args->comms+i, nranks, args->ncclId, rank));
-#endif
+    TESTCHECK(AllocateBuffs(args->sendbuffs + i, sendBytes, args->recvbuffs + i, recvBytes, args->expected + i, (size_t)args->maxbytes));
   }
   NCCLCHECK(ncclGroupEnd());
+
+  // Capture the memory used by the GPUs after initializing the NCCL communicators
+  for (int g = 0; g < args->nGpus; ++g) {
+    CUDACHECK(cudaSetDevice(args->gpus[g]));
+    getGPUMemoryInfo(nullptr, &initFreeGpuMem[g + args->nGpus]);
+    *args->initGpuMem = std::max(*args->initGpuMem, initFreeGpuMem[g] - initFreeGpuMem[g + args->nGpus]);
+  }
+
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,19,0)
   NCCLCHECK(ncclGroupStart());
   for (int i=0; i<args->nGpus; i++) {
@@ -767,15 +836,45 @@ testResult_t threadInit(struct threadArgs* args) {
   }
   NCCLCHECK(ncclGroupEnd());
 #endif
+  // Capture memory used by test buffers
+  for (int g = 0; g < args->nGpus; ++g) {
+    CUDACHECK(cudaSetDevice(args->gpus[g]));
+    getGPUMemoryInfo(nullptr, &initFreeGpuMem[g + args->nGpus*2]);
+    args->bufferMemory[args->thread] = std::max(args->bufferMemory[args->thread], initFreeGpuMem[g + args->nGpus] - initFreeGpuMem[g + args->nGpus*2]);
+  }
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
   /* Create device communicators based on test-specific requirements */
   if (deviceImpl) {
-    ncclDevCommRequirements reqs;
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0)
+    if (test_ncclVersion < NCCL_VERSION(2,29,0)) {
+      fprintf(stderr,
+        "Incompatible NCCL versions. nccl-tests was compiled with NCCL %d, but is running with NCCL %d. "
+        "The %d Device API is not compatible with versions before 2.29.\n",
+        NCCL_VERSION_CODE, test_ncclVersion, NCCL_VERSION_CODE);
+      return testInvalidUsage;
+    }
+    ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
+    if (!ncclTestEngine.getDevCommRequirements) {
+      fprintf(stderr, "Device implementation %d is not supported by this test\n", deviceImpl);
+      return testNotImplemented;
+    }
+    ncclCommProperties commProperties = NCCL_COMM_PROPERTIES_INITIALIZER;
+    NCCLCHECK(ncclCommQueryProperties(args->comms[0], &commProperties));
+    TESTCHECK(ncclTestEngine.getDevCommRequirements(deviceImpl, &reqs, &commProperties));
+#else
+    if (test_ncclVersion >= NCCL_VERSION(2,29,0)) {
+      fprintf(stderr, "Incompatible NCCL versions. nccl-tests was compiled with NCCL 2.28, but is running with NCCL %d. "
+        "The 2.28 Device API is not compatible with later.\n",
+        test_ncclVersion);
+      return testInvalidUsage;
+    }
+    ncclDevCommRequirements reqs = {};
     if (!ncclTestEngine.getDevCommRequirements ||
         !ncclTestEngine.getDevCommRequirements(deviceImpl, &reqs)) {
       fprintf(stderr, "Device implementation %d is not supported by this test\n", deviceImpl);
       return testNotImplemented;
     }
+#endif
 
     NCCLCHECK(ncclGroupStart());
     for (int i = 0; i < args->nGpus; i++) {
@@ -783,7 +882,17 @@ testResult_t threadInit(struct threadArgs* args) {
     }
     NCCLCHECK(ncclGroupEnd());
   }
+  // Capture memory used by test buffers
+  int64_t deviceCommMaxMem = 0;
+  for (int g = 0; g < args->nGpus; ++g) {
+    CUDACHECK(cudaSetDevice(args->gpus[g]));
+    int64_t freeGpuMem;
+    getGPUMemoryInfo(nullptr, &freeGpuMem);
+    deviceCommMaxMem = std::max(deviceCommMaxMem, initFreeGpuMem[g + args->nGpus*2] - freeGpuMem);
+  }
+  *args->initGpuMem += deviceCommMaxMem;
 #endif
+  free(initFreeGpuMem);
 
   TESTCHECK(threadRunTests(args));
 
@@ -801,6 +910,7 @@ testResult_t threadLaunch(struct testThread* thread) {
 }
 
 testResult_t AllocateBuffs(void **sendbuff, size_t sendBytes, void **recvbuff, size_t recvBytes, void **expected, size_t nbytes) {
+    nbytes += 8*unalign; // pad with size of max datatype in case all datatypes selected
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,19,0)
     NCCLCHECK(ncclMemAlloc(sendbuff, nbytes));
     NCCLCHECK(ncclMemAlloc(recvbuff, nbytes));
@@ -879,13 +989,15 @@ int main(int argc, char* argv[], char **envp) {
     {"cta_policy", required_argument, 0, 'x'},
     {"device_implementation", required_argument, 0, 'D'},
     {"device_cta_count", required_argument, 0, 'V'},
+    {"memory", required_argument, 0, 'M'},
+    {"unalign", 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:c:o:d:r:z:y:T:hG:C:a:R:x:D:V:J:S:", longopts, &longindex);
+    c = getopt_long(argc, argv, "t:g:b:e:i:f:n:m:w:N:p:c:o:d:r:z:y:T:hG:C:a:R:x:D:V:J:S:M:u:", longopts, &longindex);
 
     if (c == -1)
       break;
@@ -994,6 +1106,9 @@ int main(int argc, char* argv[], char **envp) {
         printf("Option -R (register) is not supported before NCCL 2.19. Ignoring\n");
 #endif
         break;
+      case 'M':
+        memory_report = (int)strtol(optarg, NULL, 0);
+        break;
       case 'x':
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,27,0)
         ctaPolicy = (int)strtol(optarg, NULL, 0);
@@ -1008,8 +1123,7 @@ int main(int argc, char* argv[], char **envp) {
       case 'D':
 	if (test_ncclVersion >= NCCL_VERSION(2,28,0)) {
           deviceImpl = (int)strtol(optarg, NULL, 0);
-	}
-	else {
+        } else {
           fprintf(stderr, "Option -D (device implementation) requires NCCL >= 2.28.0\n");
           return -1;
         }
@@ -1021,13 +1135,15 @@ int main(int argc, char* argv[], char **envp) {
             fprintf(stderr, "device_cta_count (-V) must be positive and less than 128, got %d. "
                     "Using default value 16.\n", deviceCtaCount);
             deviceCtaCount = 16;
-	  }
-        }
-	else {
+          }
+        } else {
           fprintf(stderr, "Option -V (device CTA count) requires NCCL >= 2.28.0\n");
           return -1;
 	}
         break;
+      case 'u':
+        unalign = (int)strtol(optarg, NULL, 0);
+        break;
       case 'h':
       default:
         if (c != 'h') printf("invalid option '%c'\n", c);
@@ -1053,7 +1169,7 @@ int main(int argc, char* argv[], char **envp) {
 #endif
             "[-d,--datatype <nccltype/all>] \n\t"
             "[-r,--root <root>] \n\t"
-            "[-z,--blocking <0/1>] \n\t"
+            "[-z,--blocking <0/1/2> 1=wait for completion and barrier (default behavior), 2=wait without barrier] \n\t"
             "[-y,--stream_null <0/1>] \n\t"
             "[-T,--timeout <time in seconds>] \n\t"
             "[-G,--cudagraph <num graph launches>] \n\t"
@@ -1065,6 +1181,8 @@ int main(int argc, char* argv[], char **envp) {
             "[-x,--cta_policy <0/1/2> set CTA policy (NCCL_CTA_POLICY_DEFAULT (0), NCCL_CTA_POLICY_EFFICIENCY (1), NCCL_CTA_POLICY_ZERO (2)) (default: do not set)] \n\t"
             "[-D,--device_implementation <implementation number> enable device implementation (default: 0, use NCCL implementation; requires -R 2 if > 0)] \n\t"
             "[-V,--device_cta_count <number> set number of CTAs for device implementation (default: 16)] \n\t"
+            "[-M,--memory_report <0/1> enable memory usage report (default: 0)] \n\t"
+            "[-u,--unalign <index of first element> Misalign source and destination buffers (default: 0)] \n\t"
             "[-h,--help]\n",
           basename(argv[0]));
         return 0;
@@ -1213,7 +1331,6 @@ testResult_t run() {
   for (int i=0; i<nGpus*nThreads; i++) {
     gpus[i] = (gpu0 != -1 ? gpu0 : localRank*nThreads*nGpus) + i;
     CUDACHECK(cudaSetDevice(gpus[i]));
-    TESTCHECK(AllocateBuffs(sendbuffs+i, sendBytes, recvbuffs+i, recvBytes, expected+i, (size_t)maxBytes));
     if (streamnull) {
       streams[i] = NULL;
     }
@@ -1256,30 +1373,35 @@ testResult_t run() {
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
   ncclDevComm devComms[nThreads*nGpus];
 #endif
+  int64_t initGpuMem[nThreads];
+  int64_t bufferMemory[nThreads];
+  memset(initGpuMem, 0, sizeof(initGpuMem));
+  memset(bufferMemory, 0, sizeof(bufferMemory));
   if (!parallel_init) {
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
-     ncclConfig_t config = NCCL_CONFIG_INITIALIZER;
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,27,0)
-     if (ctaPolicy >= 0)
-       config.CTAPolicy = ctaPolicy;
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
-     config.nvlinkCentricSched = 1;
-#endif
-#endif
-#endif
-     NCCLCHECK(ncclGroupStart());
-     for (int i=0; i<nGpus*nThreads; i++) {
-       CUDACHECK(cudaSetDevice(gpus[i]));
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
-       NCCLCHECK(ncclCommInitRankConfig(comms+i, ncclProcs*nThreads*nGpus, ncclId, ncclProc*nThreads*nGpus+i, &config));
-#else
-       NCCLCHECK(ncclCommInitRank(comms+i, ncclProcs*nThreads*nGpus, ncclId, ncclProc*nThreads*nGpus+i));
-#endif
+    // Capture the memory used by the GPUs before initializing the NCCL communicators
+    int64_t* initFreeGpuMem = (int64_t*)calloc(nGpus*3, sizeof(int64_t));
+    for (int g = 0; g < nGpus; ++g) {
+      CUDACHECK(cudaSetDevice(gpus[g]));
+      getGPUMemoryInfo(nullptr, &initFreeGpuMem[g]);
+    }
+    //if parallel init is not selected, use main thread to initialize NCCL
+    TESTCHECK(initComms(comms, nGpus*nThreads, ncclProc*nThreads*nGpus, ncclProcs*nThreads*nGpus, gpus, ncclId));
+
+     // Capture the memory used by the GPUs after initializing the NCCL communicators
+     for (int g = 0; g < nGpus; ++g) {
+       CUDACHECK(cudaSetDevice(gpus[g]));
+       getGPUMemoryInfo(nullptr, &initFreeGpuMem[g + nGpus]);
+     }
+     for ( size_t t = 0; t < nThreads; ++t) {
+       for (int g = 0; g < nGpus; ++g) {
+         initGpuMem[t] = std::max(initGpuMem[t], initFreeGpuMem[g] - initFreeGpuMem[g + nGpus]);
+       }
      }
-     NCCLCHECK(ncclGroupEnd());
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,19,0)
      NCCLCHECK(ncclGroupStart());
      for (int i=0; i<nGpus*nThreads; i++) {
+       CUDACHECK(cudaSetDevice(gpus[i]));
+       TESTCHECK(AllocateBuffs(sendbuffs+i, sendBytes, recvbuffs+i, recvBytes, expected+i, (size_t)maxBytes));
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,27,0)
        if (test_ncclVersion >= NCCL_VERSION(2,27,0) && (local_register == SYMMETRIC_REGISTER)) {
          NCCLCHECK(ncclCommWindowRegister(comms[i], sendbuffs[i], maxBytes, (ncclWindow_t*)&sendRegHandles[i], NCCL_WIN_COLL_SYMMETRIC));
@@ -1293,15 +1415,48 @@ testResult_t run() {
      }
      NCCLCHECK(ncclGroupEnd());
 #endif
+     // Capture memory used by after allocating buffers
+     for (int g = 0; g < nGpus; ++g) {
+       CUDACHECK(cudaSetDevice(gpus[g]));
+       getGPUMemoryInfo(nullptr, &initFreeGpuMem[g + nGpus*2]);
+     }
+     for ( size_t t = 0; t < nThreads; ++t) {
+      for (int g = 0; g < nGpus; ++g) {
+        bufferMemory[t] = std::max(bufferMemory[t], initFreeGpuMem[g + nGpus] - initFreeGpuMem[g + nGpus*2]);
+      }
+     }
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
      /* Create device communicators based on test-specific requirements */
      if (deviceImpl) {
-       ncclDevCommRequirements reqs;
-       if (!ncclTestEngine.getDevCommRequirements ||
-           !ncclTestEngine.getDevCommRequirements(deviceImpl, &reqs)) {
-         fprintf(stderr, "Device implementation %d is not supported by this test\n", deviceImpl);
-         return testNotImplemented;
-       }
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0)
+      if (test_ncclVersion < NCCL_VERSION(2,29,0)) {
+        fprintf(stderr,
+          "Incompatible NCCL versions. nccl-tests was compiled with NCCL %d, but is running with NCCL %d. "
+          "The %d Device API is not compatible with versions before 2.29.\n",
+          NCCL_VERSION_CODE, test_ncclVersion, NCCL_VERSION_CODE);
+        return testInvalidUsage;
+      }
+      ncclDevCommRequirements reqs = NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER;
+      if (!ncclTestEngine.getDevCommRequirements) {
+        fprintf(stderr, "Device implementation %d is not supported by this test\n", deviceImpl);
+        return testNotImplemented;
+      }
+      ncclCommProperties commProperties = NCCL_COMM_PROPERTIES_INITIALIZER;
+      NCCLCHECK(ncclCommQueryProperties(comms[0], &commProperties));
+      TESTCHECK(ncclTestEngine.getDevCommRequirements(deviceImpl, &reqs, &commProperties));
+#else
+      if (test_ncclVersion >= NCCL_VERSION(2,29,0)) {
+        fprintf(stderr, "Incompatible NCCL versions. nccl-tests was compiled with NCCL 2.28, but is running with NCCL %d. "
+          "The 2.28 Device API is not compatible with later versions.\n", test_ncclVersion);
+        return testInvalidUsage;
+      }
+      ncclDevCommRequirements reqs = {};
+      if (!ncclTestEngine.getDevCommRequirements ||
+          !ncclTestEngine.getDevCommRequirements(deviceImpl, &reqs)) {
+        fprintf(stderr, "Device implementation %d is not supported by this test\n", deviceImpl);
+        return testNotImplemented;
+      }
+#endif
 
        NCCLCHECK(ncclGroupStart());
        for (int i = 0; i < nGpus * nThreads; i++) {
@@ -1309,15 +1464,28 @@ testResult_t run() {
        }
        NCCLCHECK(ncclGroupEnd());
      }
+     int64_t deviceCommMaxMem = 0;
+     for (int g = 0; g < nGpus; ++g) {
+       CUDACHECK(cudaSetDevice(gpus[g]));
+       int64_t freeGpuMem;
+       getGPUMemoryInfo(nullptr, &freeGpuMem);
+       deviceCommMaxMem = std::max(deviceCommMaxMem, initFreeGpuMem[g + nGpus*2] - freeGpuMem);
+     }
+     for ( size_t t = 0; t < nThreads; ++t) {
+       initGpuMem[t] += deviceCommMaxMem;
+     }
 #endif
+    free(initFreeGpuMem);
   }
 
   int errors[nThreads];
   double bw[nThreads];
+  int64_t devMemUsed[nThreads];
   int bw_count[nThreads];
   for (int t=0; t<nThreads; t++) {
     bw[t] = 0.0;
     errors[t] = bw_count[t] = 0;
+    devMemUsed[t] = std::numeric_limits<int64_t>::min();
   }
 
   fflush(stdout);
@@ -1358,6 +1526,9 @@ testResult_t run() {
     threads[t].args.errors=errors+t;
     threads[t].args.bw=bw+t;
     threads[t].args.bw_count=bw_count+t;
+    threads[t].args.initGpuMem = initGpuMem + t;
+    threads[t].args.bufferMemory = bufferMemory + t;
+    threads[t].args.devMemUsed = devMemUsed + t;
 
     threads[t].args.reportErrors = datacheck;
 
@@ -1376,11 +1547,17 @@ testResult_t run() {
       errors[0] += errors[t];
       bw[0] += bw[t];
       bw_count[0] += bw_count[t];
+      devMemUsed[0] = std::max(devMemUsed[0], devMemUsed[t]);
+      initGpuMem[0] = std::max(initGpuMem[0], initGpuMem[t]);
+      bufferMemory[0] = std::max(bufferMemory[0], bufferMemory[t]);
     }
   }
 
 #ifdef MPI_SUPPORT
   MPI_Allreduce(MPI_IN_PLACE, &errors[0], 1, MPI_INT, MPI_SUM, MPI_COMM_WORLD);
+  MPI_Allreduce(MPI_IN_PLACE, &devMemUsed[0], 1, MPI_INT64_T, MPI_MAX, MPI_COMM_WORLD);
+  MPI_Allreduce(MPI_IN_PLACE, &initGpuMem[0], 1, MPI_INT64_T, MPI_MAX, MPI_COMM_WORLD);
+  MPI_Allreduce(MPI_IN_PLACE, &bufferMemory[0], 1, MPI_INT64_T, MPI_MAX, MPI_COMM_WORLD);
 #endif
 
   if (!parallel_init) {
@@ -1416,10 +1593,18 @@ testResult_t run() {
   }
 
   envstr = getenv("NCCL_TESTS_MIN_BW");
-  double check_avg_bw = envstr ? atof(envstr) : -1;
+  const double check_avg_bw = envstr ? atof(envstr) : -1;
   bw[0] /= bw_count[0];
 
   writeResultFooter(errors, bw, check_avg_bw, program_invocation_short_name);
+  if (memory_report) {
+    memInfo_t memInfos[3];
+    memInfos[0] = { initGpuMem[0], "Initialization" };
+    memInfos[1] = { bufferMemory[0], "User-Allocated" };
+    memInfos[2] = { devMemUsed[0], "Collective" };
+    writeMemInfo(memInfos, 3);
+  }
+  finalizeFooter();
 
 #ifdef MPI_SUPPORT
   MPI_Comm_free(&mpi_comm);
diff --git a/src/common.h b/src/common.h
index b0b250c..990fad6 100644
--- a/src/common.h
+++ b/src/common.h
@@ -6,7 +6,7 @@
 #ifndef __COMMON_H__
 #define __COMMON_H__
 
-#define NCCL_TESTS_VERSION "2.17.6"
+#define NCCL_TESTS_VERSION "2.18.3"
 
 #include "nccl.h"
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
@@ -72,7 +72,8 @@ typedef enum {
   testNcclError = 3,
   testTimeout = 4,
   testNotImplemented = 5,
-  testNumResults = 6
+  testInvalidUsage = 6,
+  testNumResults = 7, // Must be last
 } testResult_t;
 
 // Relay errors up and trace
@@ -111,7 +112,14 @@ struct testEngine {
   void (*getBuffSize)(size_t *sendcount, size_t *recvcount, size_t count, int nranks);
   testResult_t (*runTest)(struct threadArgs* args, int root, ncclDataType_t type,
       const char* typeName, ncclRedOp_t op, const char* opName);
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
+  /* Optional; called from initComms after common fields are set on ncclConfig_t. */
+  void (*initCommConfig)(ncclConfig_t* config);
+#endif
+
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,29,0)
+  testResult_t (*getDevCommRequirements)(int deviceImpl, ncclDevCommRequirements* reqs, ncclCommProperties_t* commProperties);
+#elif NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
   bool (*getDevCommRequirements)(int deviceImpl, ncclDevCommRequirements* reqs);
 #endif
 };
@@ -155,6 +163,10 @@ struct threadArgs {
 
   struct testColl* collTest;
 
+  int64_t* initGpuMem;
+  int64_t* bufferMemory;
+  int64_t* devMemUsed;
+
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,19,0)
   void** sendRegHandles;
   void** recvRegHandles;
@@ -174,7 +186,7 @@ extern void Barrier(struct threadArgs* args);
 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 InitData(void* data, const size_t count, size_t offset, ncclDataType_t type, ncclRedOp_t op, const uint64_t seed, const int nranks, const int rank);
-extern void AllocateBuffs(void **sendbuff, void **recvbuff, void **expected, void **expectedHost, size_t nbytes, int nranks);
+extern testResult_t AllocateBuffs(void **sendbuff, size_t sendBytes, void **recvbuff, size_t recvBytes, void **expected, size_t nbytes);
 
 #include <unistd.h>
 
diff --git a/src/common.mk b/src/common.mk
index 0bd630f..7b4f316 100644
--- a/src/common.mk
+++ b/src/common.mk
@@ -66,7 +66,7 @@ NVCC_GENCODE ?= -gencode=arch=compute_35,code=sm_35 \
                 -gencode=arch=compute_70,code=compute_70
 endif
 
-NVCUFLAGS  := -ccbin $(CXX) $(NVCC_GENCODE) $(CXXSTD)
+NVCUFLAGS  := -ccbin $(CXX) $(NVCC_GENCODE) $(CXXSTD) --extended-lambda
 CXXFLAGS   := $(CXXSTD)
 
 LDFLAGS    := -L${CUDA_LIB} -lcudart -lrt
diff --git a/src/sendrecv.cu b/src/sendrecv.cu
index 2f09667..66e8af0 100644
--- a/src/sendrecv.cu
+++ b/src/sendrecv.cu
@@ -64,6 +64,17 @@ testResult_t SendRecvRunColl(void* sendbuff, size_t sendoffset, void* recvbuff,
   return testSuccess;
 }
 
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
+static void SendRecvInitCommConfig(ncclConfig_t* config) {
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,30,0)
+  // Ring sendrecv uses two distinct peers per rank (send to rank+1, recv from rank-1).
+  config->maxP2pPeers = 2;
+#else
+  (void)config;
+#endif
+}
+#endif
+
 struct testColl sendRecvTest = {
   "SendRecv",
   SendRecvGetCollByteCount,
@@ -114,7 +125,10 @@ testResult_t SendRecvRunTest(struct threadArgs* args, int root, ncclDataType_t t
 
 struct testEngine sendRecvEngine = {
   .getBuffSize = SendRecvGetBuffSize,
-  .runTest = SendRecvRunTest
+  .runTest = SendRecvRunTest,
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
+  .initCommConfig = SendRecvInitCommConfig,
+#endif
 };
 
 #pragma weak ncclTestEngine=sendRecvEngine
diff --git a/src/util.cu b/src/util.cu
index 2bb76f9..40f1935 100644
--- a/src/util.cu
+++ b/src/util.cu
@@ -19,9 +19,18 @@
 #include "util.h"
 #include <assert.h>
 #include <errno.h>
+#include <string>
+#include <iomanip>
+#include <string.h>
+#include <stdio.h>
+#include <math.h>
+#include <nvml.h>
+#include <random>
 
 #define PRINT if (is_main_thread) printf
 
+#define MAX_LINE 2048
+
 extern int nThreads;
 extern int nGpus;
 extern size_t minBytes;
@@ -35,13 +44,14 @@ extern int agg_iters;
 extern int parallel_init;
 extern int blocking_coll;
 extern int cudaGraphLaunches;
+extern int unalign;
 
 static FILE *json_report_fp;
 static thread_local bool write_json;
 
 #define JSON_FILE_VERSION 1
 
-#define TIME_STRING_FORMAT "%Y-%m-%d %H:%M:%S"
+#define TIME_STRING_FORMAT "%Y-%m-%d %H:%M:%SZ"
 
 typedef enum {
   JSON_NONE, // A pseudo-state meaning that the document is empty
@@ -264,6 +274,10 @@ static void jsonDouble(const double val) {
   if(val != val) {
     fprintf(json_report_fp, "\"nan\"");
   }
+  else if(isinf(val)) {
+    // JSON has no infinity literal; emit a string, matching the "nan" handling above.
+    fprintf(json_report_fp, val < 0 ? "\"-inf\"" : "\"inf\"");
+  }
   else {
     fprintf(json_report_fp, "%lf", val);
   }
@@ -274,7 +288,7 @@ static void jsonDouble(const double val) {
 void formatNow(char *buff, int len) {
   time_t now;
   time(&now);
-  struct tm *timeinfo = localtime(&now);
+  struct tm *timeinfo = gmtime(&now);
 
   strftime(buff, len, TIME_STRING_FORMAT, timeinfo);
 }
@@ -342,11 +356,15 @@ void jsonOutputInit(const char *in_path,
   jsonFinishList();
 
   jsonKey("env");
-  jsonStartList();
+  jsonStartObject();
   for(char **e = envp; *e; e++) {
-    jsonStr(*e);
+    char *ptr = strchr(*e, '=');
+    if(!ptr) continue;
+    char key[MAX_LINE];
+    snprintf(key, sizeof(key), "%.*s", (int)(ptr - *e), *e);
+    jsonKey(key); jsonStr(ptr + 1);
   }
-  jsonFinishList();
+  jsonFinishObject();
   jsonKey("nccl_version"); jsonInt(test_ncclVersion);
 }
 
@@ -385,14 +403,23 @@ struct rankInfo_t {
   int device;
   char device_hex[128];
   char devinfo[1024];
+  char gpuSerial[NVML_DEVICE_SERIAL_BUFFER_SIZE];
 };
 
-// Helper function to parse the device info lines passed via MPI to the root rank.
-// This fills 'rank' with the parsed contents of 'instring'.
+// Helper function to parse one device-info line (as gathered to the root rank).
+// This fills 'rank' with the parsed contents at the start of 'instring' and
+// returns the number of bytes consumed, so the caller can advance to the next
+// line and parse again — a single process may report several GPUs, hence
+// several lines, in one buffer. Returns 0 when no complete record is present
+// (end of buffer / trailing garbage), which lets callers loop until exhausted.
+// The sscanf format below hardcodes the gpuSerial field width (%29); this guards
+// against NVML_DEVICE_SERIAL_BUFFER_SIZE drifting out from under that magic number.
+static_assert(sizeof(rankInfo_t::gpuSerial) == 30,
+              "parseRankInfo assumes that rankInfo_t::gpuSerial has a width of 30 bytes.");
 static int parseRankInfo(rankInfo_t *rank, const char *instring) {
-  int end;
-  sscanf(instring,
-         "#  Rank %d Group %d Pid %d on %1024s device %d [%128[^]]] %1024[^\n]\n%n",
+  int end = 0;
+  const int matched = sscanf(instring,
+         "#  Rank %d Group %d Pid %d on %1023s device %d [%127[^]]] %1023[^#]#%29[^\n]\n%n",
          &rank->rank,
          &rank->group,
          &rank->pid,
@@ -400,8 +427,13 @@ static int parseRankInfo(rankInfo_t *rank, const char *instring) {
          &rank->device,
          rank->device_hex,
          rank->devinfo,
+         rank->gpuSerial,
          &end);
-  return end;
+  // All 8 fields must be present for a complete record; %n (end) is only
+  // assigned once the scan reaches the trailing newline, i.e. a whole line was
+  // consumed. Anything short of that (end of buffer, "...\n" truncation marker)
+  // yields 0 so the caller's loop terminates.
+  return (matched == 8) ? end : 0;
 }
 
 static void jsonRankInfo(const rankInfo_t *ri) {
@@ -413,9 +445,72 @@ static void jsonRankInfo(const rankInfo_t *ri) {
   jsonKey("device");      jsonInt(ri->device);
   jsonKey("device_hex");  jsonStr(ri->device_hex);
   jsonKey("device_info"); jsonStr(ri->devinfo);
+  jsonKey("serial");      jsonStr(ri->gpuSerial);
   jsonFinishObject();
 }
 
+
+// Fetch the serial number for one GPU into `serial` (must hold at least
+// NVML_DEVICE_SERIAL_BUFFER_SIZE bytes). NVML must already be initialized by the
+// caller. Returns 0 on success, nonzero if the handle or serial
+// lookup fails (in which case `serial` is left untouched for the caller to handle).
+int getGPUSerial(const char *pciBusId, char *serial) {
+    nvmlDevice_t device;
+
+    // Look up by PCI bus id, not NVML index: NVML's index space is PCI-ordered and
+    // ignores CUDA_VISIBLE_DEVICES, whereas the caller's cudaDev is a CUDA runtime
+    // ordinal that CUDA_VISIBLE_DEVICES / containers / schedulers remap. Keying by
+    // the stable bus id guarantees the serial matches the GPU this rank actually uses.
+    nvmlReturn_t result = nvmlDeviceGetHandleByPciBusId(pciBusId, &device);
+    if (NVML_SUCCESS != result) {
+        fprintf(stderr, "Failed to get device handle for PCI %s: %s\n", pciBusId, nvmlErrorString(result));
+        return 1;
+    }
+
+    result = nvmlDeviceGetSerial(device, serial, NVML_DEVICE_SERIAL_BUFFER_SIZE);
+    if (NVML_SUCCESS != result) {
+        fprintf(stderr, "Failed to get serial number for PCI %s: %s\n", pciBusId, nvmlErrorString(result));
+        return 1;
+    }
+
+    return 0;
+}
+
+// Generate an RFC-4122 v4 UUID string (36 chars + NUL) without an external libuuid,
+// so building needs no uuid-dev. Randomness is fine for a per-run identifier.
+static void genUuidV4(char out[37]) {
+    std::random_device rd;
+    unsigned char b[16];
+    for (int i = 0; i < 16; i += 4) {
+        unsigned int r = rd();
+        b[i+0] =  r        & 0xFF;
+        b[i+1] = (r >> 8)  & 0xFF;
+        b[i+2] = (r >> 16) & 0xFF;
+        b[i+3] = (r >> 24) & 0xFF;
+    }
+    b[6] = (b[6] & 0x0F) | 0x40;  // version 4
+    b[8] = (b[8] & 0x3F) | 0x80;  // variant 10x
+    snprintf(out, 37,
+        "%02x%02x%02x%02x-%02x%02x-%02x%02x-%02x%02x-%02x%02x%02x%02x%02x%02x",
+        b[0],b[1],b[2],b[3],b[4],b[5],b[6],b[7],b[8],b[9],b[10],b[11],b[12],b[13],b[14],b[15]);
+}
+
+// RAII owner for the NVML library handle. Initializes NVML on construction and
+// guarantees a single shutdown when the scope exits, on any path (normal return,
+// early error return, exception). `ready` reflects whether init succeeded; a
+// failed init means serials are unavailable but the rest of the report proceeds.
+struct NvmlSession {
+  const bool ready;
+  NvmlSession() : ready(nvmlInit() == NVML_SUCCESS) {}
+  ~NvmlSession() {
+    if (ready && nvmlShutdown() != NVML_SUCCESS) {
+      fprintf(stderr, "Failed to shut down NVML.\n");
+    }
+  }
+  NvmlSession(const NvmlSession&) = delete;
+  NvmlSession& operator=(const NvmlSession&) = delete;
+};
+
 // Write the start of a benchmark output line containing the bytes &
 // op type, both to stdout and to json if we are writing there.
 void writeBenchmarkLinePreamble(size_t nBytes, size_t nElem, const char typeName[], const char opName[], int root) {
@@ -452,6 +547,12 @@ void writeBenchMarkLineNullBody() {
 }
 
 void getFloatStr(double value, int width, char* str) {
+  if (isinf(value)) {
+    // +inf would loop forever below (value >= val always true as val wraps);
+    // short-circuit both signs.
+    sprintf(str, "%*s", width, value < 0 ? "-inf" : "inf");
+    return;
+  }
   int power = 0;
   for (uint64_t val = 1; value >= val; val *= 10) power++;
 
@@ -512,11 +613,12 @@ void writeBenchmarkLineBody(double timeUsec, double algBw, double busBw, bool re
 testResult_t writeDeviceReport(size_t *maxMem, int localRank, int proc, int totalProcs, int color, const char hostname[], const char *program_name) {
   PRINT("# nccl-tests version %s nccl-headers=%d nccl-library=%d\n", NCCL_TESTS_VERSION, NCCL_VERSION_CODE, test_ncclVersion);
   PRINT("# Collective test starting: %s\n", program_name);
-  PRINT("# nThread %d nGpus %d minBytes %ld maxBytes %ld step: %ld(%s) warmup iters: %d iters: %d agg iters: %d validation: %d graph: %d\n",
+  PRINT("# nThread %d nGpus %d minBytes %ld maxBytes %ld step: %ld(%s) warmup iters: %d iters: %d agg iters: %d validation: %d graph: %d unalign: %d\n",
         nThreads, nGpus, minBytes, maxBytes,
         (stepFactor > 1)?stepFactor:stepBytes, (stepFactor > 1)?"factor":"bytes",
-        warmup_iters, iters, agg_iters, datacheck, cudaGraphLaunches);
-  if (blocking_coll) PRINT("# Blocking Enabled: wait for completion and barrier after each collective \n");
+        warmup_iters, iters, agg_iters, datacheck, cudaGraphLaunches, unalign);
+  if (blocking_coll == 1) PRINT("# Blocking Enabled: wait for completion and barrier after each collective \n");
+  if (blocking_coll > 1)  PRINT("# Blocking Enabled: wait for completion after each collective (no barrier) \n");
   if (parallel_init) PRINT("# Parallel Init Enabled: threads call into NcclInitRank concurrently \n");
   PRINT("#\n");
 
@@ -539,18 +641,26 @@ testResult_t writeDeviceReport(size_t *maxMem, int localRank, int proc, int tota
     jsonKey("aggregated_iterations"); jsonInt(agg_iters);
     jsonKey("validation");            jsonInt(datacheck);
     jsonKey("graph");                 jsonInt(cudaGraphLaunches);
-    jsonKey("blocking_collectives");  jsonBool(blocking_coll);
+    jsonKey("blocking_collectives");  jsonInt(blocking_coll);
     jsonKey("parallel_init");         jsonBool(parallel_init);
   }
 
   PRINT("# Using devices\n");
-#define MAX_LINE 2048
   char line[MAX_LINE];
   int len = 0;
   const char* envstr = getenv("NCCL_TESTS_DEVICE");
   const int gpu0 = envstr ? atoi(envstr) : -1;
   int available_devices;
+  char gpuSerial[NVML_DEVICE_SERIAL_BUFFER_SIZE];
   CUDACHECK(cudaGetDeviceCount(&available_devices));
+
+  // Initialize NVML once for the whole device loop; the guard shuts it down on
+  // every exit path from here on (see NvmlSession).
+  NvmlSession nvml;
+  if (!nvml.ready) {
+    fprintf(stderr, "Failed to initialize NVML; GPU serials will be reported as \"unknown\".\n");
+  }
+
   for (int i=0; i<nThreads*nGpus; i++) {
     const int cudaDev = (gpu0 != -1 ? gpu0 : localRank*nThreads*nGpus) + i;
     const int rank = proc*nThreads*nGpus+i;
@@ -562,12 +672,26 @@ testResult_t writeDeviceReport(size_t *maxMem, int localRank, int proc, int tota
       return testNotImplemented;
     }
     CUDACHECK(cudaGetDeviceProperties(&prop, cudaDev));
+    // Resolve the serial by stable PCI bus id (see getGPUSerial): cudaDev is a CUDA
+    // ordinal that CUDA_VISIBLE_DEVICES can remap, so an NVML-by-index lookup could
+    // return another GPU's serial. cudaDeviceGetPCIBusId yields the "dddd:bb:dd.f"
+    // string NVML expects.
+    char pciBusId[NVML_DEVICE_PCI_BUS_ID_BUFFER_SIZE] = {0};
+    if (!nvml.ready
+        || cudaDeviceGetPCIBusId(pciBusId, sizeof(pciBusId), cudaDev) != cudaSuccess
+        || getGPUSerial(pciBusId, gpuSerial) != 0) {
+      // NVML unavailable or this device's lookup failed — write a defined value
+      // rather than emit uninitialized/stale bytes into the report (and JSON serial).
+      snprintf(gpuSerial, sizeof(gpuSerial), "unknown");
+    }
     if (len < MAX_LINE) {
-      len += snprintf(line+len, MAX_LINE-len, "#  Rank %2d Group %2d Pid %6d on %10s device %2d [%04x:%02x:%02x] %s\n",
-                      rank, color, getpid(), hostname, cudaDev, prop.pciDomainID, prop.pciBusID, prop.pciDeviceID, prop.name);
+      len += snprintf(line+len, MAX_LINE-len, "#  Rank %2d Group %2d Pid %6d on %10s device %2d [%04x:%02x:%02x] %s#%s\n",
+                      rank, color, getpid(), hostname, cudaDev, prop.pciDomainID, prop.pciBusID, prop.pciDeviceID, prop.name, gpuSerial);
     }
     *maxMem = std::min(*maxMem, prop.totalGlobalMem);
   }
+  // NVML is shut down by the NvmlSession guard when this function returns.
+
   if (len >= MAX_LINE) {
     strcpy(line+MAX_LINE-5, "...\n");
   }
@@ -584,9 +708,18 @@ testResult_t writeDeviceReport(size_t *maxMem, int localRank, int proc, int tota
     for (int p = 0; p < totalProcs; p++) {
       PRINT("%s", lines+MAX_LINE*p);
       if(write_json) {
+        // A process that drives multiple GPUs contributes multiple "# Rank"
+        // lines in its MAX_LINE buffer; walk them all so every GPU's serial is
+        // emitted, not just the first. parseRankInfo returns the bytes consumed
+        // (0 at end-of-buffer); the buffer is NUL-terminated within MAX_LINE, so
+        // the walk cannot run past rank p into rank p+1.
+        const char *cursor = lines + MAX_LINE*p;
         rankInfo_t rankinfo;
-        parseRankInfo(&rankinfo, lines + MAX_LINE*p);
-        jsonRankInfo(&rankinfo);
+        int consumed;
+        while ((consumed = parseRankInfo(&rankinfo, cursor)) > 0) {
+          jsonRankInfo(&rankinfo);
+          cursor += consumed;
+        }
       }
     }
     if(write_json) {
@@ -598,11 +731,18 @@ testResult_t writeDeviceReport(size_t *maxMem, int localRank, int proc, int tota
 #else
   PRINT("%s", line);
   if(write_json) {
-    rankInfo_t rankinfo;
-    parseRankInfo(&rankinfo, line);
     jsonKey("devices");
     jsonStartList();
-    jsonRankInfo(&rankinfo);
+    // This process may own several GPUs (nThreads*nGpus > 1), so it wrote one
+    // "# Rank" line per GPU into 'line'. Walk every line so each GPU's serial is
+    // emitted, instead of parsing only the first.
+    const char *cursor = line;
+    rankInfo_t rankinfo;
+    int consumed;
+    while ((consumed = parseRankInfo(&rankinfo, cursor)) > 0) {
+      jsonRankInfo(&rankinfo);
+      cursor += consumed;
+    }
     jsonFinishList();
   }
 #endif
@@ -617,14 +757,15 @@ 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 char* tsFmt = report_timestamps ? TIME_STRING_FORMAT : "";
   const char* timeStr = report_cputime ? "cputime" : "time";
   PRINT("#\n");
   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 %20s\n", "size", "count", "type", "redop", "root",
-        timeStr, "algbw", "busbw", "#wrong", timeStr, "algbw", "busbw", "#wrong", tsLbl);
-  PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %6s  %7s  %6s  %6s  %6s %20s\n", "(B)", "(elements)", "", "", "",
-        "(us)", "(GB/s)", "(GB/s)", "", "(us)", "(GB/s)", "(GB/s)", "", tsFmt);
+  PRINT("# %10s  %12s  %8s  %6s  %6s  %7s  %6s  %6s  %6s  %7s  %6s  %6s  %6s %*s\n", "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)", "", "", "",
+        "(us)", "(GB/s)", "(GB/s)", "", "(us)", "(GB/s)", "(GB/s)", "", tsPad, tsFmt);
 
   if(write_json) {
     jsonKey("results"); jsonStartList();
@@ -640,23 +781,67 @@ void writeResultFooter(const int errors[], const double bw[], double check_avg_b
     jsonFinishList();
   }
 
-  PRINT("# Out of bounds values : %d %s\n", errors[0], errors[0] ? "FAILED" : "OK");
-  PRINT("# Avg bus bandwidth    : %g %s\n", bw[0], check_avg_bw == -1 ? "" : (bw[0] < check_avg_bw*(0.9) ? "FAILED" : "OK"));
+  PRINT("# %-20s : %d %s\n", "Out of bounds values", errors[0], errors[0] ? "FAILED" : "OK");
+  PRINT("# %-20s : %g %s\n", "Avg bus bandwidth", bw[0], check_avg_bw == -1 ? "" : (bw[0] < check_avg_bw*(0.9) ? "FAILED" : "OK"));
   PRINT("#\n");
   PRINT("# Collective test concluded: %s\n", program_name);
 
   if(write_json) {
+    char uuid[37];
+    genUuidV4(uuid);
     jsonKey("out_of_bounds");
     jsonStartObject();
     jsonKey("count");      jsonInt(errors[0]);
     jsonKey("okay");       jsonBool(errors[0] == 0);
     jsonFinishObject();
-    jsonKey("average_bus_bandwidith");
+    jsonKey("average_bus_bandwidth");
     jsonStartObject();
-    jsonKey("bandwidith"); jsonDouble(bw[0]);
+    jsonKey("bandwidth"); jsonDouble(bw[0]);
     jsonKey("okay");       check_avg_bw == -1 ? jsonStr("unchecked") : jsonBool(bw[0] >= check_avg_bw*(0.9));
     jsonFinishObject();
+    jsonKey("collective"); jsonStr(program_name);
+    jsonKey("uuid");       jsonStr(uuid);
+  }
+}
+
+std::string getMemString(double amount) {
+  std::string postfix = " B";
+  if (abs(amount) >= 1024.0*1024.0*1024.0) {
+    postfix = " GB";
+    amount /= 1024.0 * 1024.0 * 1024.0;
+  } else if (abs(amount) >= 1024.0*1024.0) {
+    postfix = " MB";
+    amount /= 1024.0 * 1024.0;
+  } else if (abs(amount) >= 1024.0) {
+    postfix = " KB";
+    amount /= 1024.0;
+  }
+  int precision = 0;
+  if (abs(amount) < 10.0) {
+    precision = 2;
+  } else if (abs(amount) < 100.0) {
+    precision = 1;
+  }
+  std::stringstream ss;
+  ss << std::fixed << std::setprecision(precision) << amount << postfix;
+  return ss.str();
+}
+
+void writeMemInfo(memInfo_t* memInfos, int numMemInfos) {
+
+  std::stringstream ss;
+  uint64_t maxAmount = 0;
+  for (int i = 0; i < numMemInfos; i++) {
+    ss << memInfos[i].name << " "
+      << getMemString(memInfos[i].amount)
+      << " ";
+    if (i < numMemInfos - 1) {
+      ss << "| ";
+    }
+    maxAmount += memInfos[i].amount;
   }
+  ss << "| Total  " << getMemString(maxAmount);
+  PRINT("# %-20s : %s\n", "GPU memory usage", ss.str().c_str());
 }
 
 // Write out remaining errors to stdout/json.
@@ -677,3 +862,6 @@ void writeErrors() {
   }
 }
 
+void finalizeFooter() {
+  PRINT("#\n");
+}
diff --git a/src/util.h b/src/util.h
index bea21ac..059ee7c 100644
--- a/src/util.h
+++ b/src/util.h
@@ -8,6 +8,11 @@
 
 #include "common.h"
 
+struct memInfo_t {
+  int64_t amount;
+  const char* name;
+};
+
 // Try to set up JSON file output. If MPI is used, only rank 0 will proceed.
 // This should be called by only a single thread.
 // If 'in_path' is NULL, we stop.
@@ -32,6 +37,8 @@ void writeBenchmarkLineBody(double timeUsec, double algBw, double busBw, bool re
 testResult_t writeDeviceReport(size_t *maxMem, int localRank, int proc, int totalProcs, int color, const char hostname[], const char *program_name);
 void writeResultHeader(bool report_cputime, bool report_timestamps);
 void writeResultFooter(const int errors[], const double bw[], double check_avg_bw, const char *program_name);
+void finalizeFooter();
+void writeMemInfo(memInfo_t* memInfos, int numMemInfos);
 void writeErrors();
 
 #endif