Add Windows support for nccl tests

Add a portable OS abstraction and CMake build entry points so the tests can build on Windows with MSVC and CUDA. Update the perf and verifiable sources to avoid POSIX-only APIs and compiler extensions where needed. Signed-off-by: Unmesh Deodhar <[email protected]>
2026-09-26 03:57:38 +00:00 · 2026-06-25 10:17:21 -05:00
parent 036a7eda0d
commit 2c52b4e910
25 changed files with 865 additions and 215 deletions
@@ -0,0 +1,136 @@
+#
+# Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved.
+#
+# See LICENSE.txt for license information
+#
+
+cmake_minimum_required(VERSION 3.25)
+
+project(nccl_tests LANGUAGES CXX CUDA)
+
+if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES)
+  set(CMAKE_BUILD_TYPE Release CACHE STRING "Build type" FORCE)
+endif()
+
+if(CMAKE_SYSTEM_NAME STREQUAL "Linux")
+  set(NCCL_OS_LINUX ON)
+elseif(CMAKE_SYSTEM_NAME STREQUAL "Windows")
+  set(NCCL_OS_WINDOWS ON)
+else()
+  message(FATAL_ERROR "Unsupported OS: ${CMAKE_SYSTEM_NAME}")
+endif()
+
+find_package(CUDAToolkit REQUIRED)
+find_package(MPI QUIET)
+
+if(CUDAToolkit_VERSION VERSION_GREATER_EQUAL 13.0)
+  set(CMAKE_CXX_STANDARD 17)
+  set(CMAKE_CUDA_STANDARD 17)
+else()
+  set(CMAKE_CXX_STANDARD 14)
+  set(CMAKE_CUDA_STANDARD 14)
+endif()
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+set(CMAKE_CUDA_STANDARD_REQUIRED ON)
+set(CMAKE_CXX_EXTENSIONS OFF)
+
+if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES OR CMAKE_CUDA_ARCHITECTURES STREQUAL "")
+  if(CUDAToolkit_VERSION VERSION_GREATER_EQUAL 13.0)
+    set(CMAKE_CUDA_ARCHITECTURES "75;80;90;100;110;120")
+  elseif(CUDAToolkit_VERSION VERSION_GREATER_EQUAL 12.8)
+    set(CMAKE_CUDA_ARCHITECTURES "60;70;80;90;100;120")
+  elseif(CUDAToolkit_VERSION VERSION_GREATER_EQUAL 11.8)
+    set(CMAKE_CUDA_ARCHITECTURES "60;70;80;90")
+  else()
+    set(CMAKE_CUDA_ARCHITECTURES "60;70;80")
+  endif()
+endif()
+
+set(NCCL_HOME "" CACHE PATH "Path to NCCL install or build directory")
+
+find_path(NCCL_INCLUDE_DIR
+  NAMES nccl.h
+  HINTS
+    ${NCCL_HOME}/include
+    ${NCCL_HOME}/build/include
+    $ENV{NCCL_HOME}/include
+    $ENV{NCCL_HOME}/build/include
+    /usr/include
+    /usr/local/include
+)
+
+find_library(NCCL_LIBRARY
+  NAMES nccl libnccl nccl_static
+  HINTS
+    ${NCCL_HOME}/lib
+    ${NCCL_HOME}/lib/x64
+    ${NCCL_HOME}/build/lib
+    ${NCCL_HOME}/build/lib/Release
+    ${NCCL_HOME}/build/lib/Debug
+    ${NCCL_HOME}/build/src
+    ${NCCL_HOME}/build/src/Release
+    ${NCCL_HOME}/build/src/Debug
+    $ENV{NCCL_HOME}/lib
+    $ENV{NCCL_HOME}/lib/x64
+    $ENV{NCCL_HOME}/build/lib
+    $ENV{NCCL_HOME}/build/lib/Release
+    $ENV{NCCL_HOME}/build/lib/Debug
+    $ENV{NCCL_HOME}/build/src
+    $ENV{NCCL_HOME}/build/src/Release
+    $ENV{NCCL_HOME}/build/src/Debug
+    /usr/lib
+    /usr/local/lib
+    /usr/lib/x86_64-linux-gnu
+)
+
+if(NOT NCCL_INCLUDE_DIR OR NOT NCCL_LIBRARY)
+  message(FATAL_ERROR "Could not find NCCL. Set -DNCCL_HOME=/path/to/nccl or provide NCCL_INCLUDE_DIR and NCCL_LIBRARY.")
+endif()
+
+add_library(nccl UNKNOWN IMPORTED)
+set_target_properties(nccl PROPERTIES
+  IMPORTED_LOCATION "${NCCL_LIBRARY}"
+  INTERFACE_INCLUDE_DIRECTORIES "${NCCL_INCLUDE_DIR}"
+)
+
+add_library(nccl_tests_options INTERFACE)
+target_compile_definitions(nccl_tests_options INTERFACE
+  $<$<BOOL:${NCCL_OS_LINUX}>:NCCL_OS_LINUX>
+  $<$<BOOL:${NCCL_OS_LINUX}>:NCCL_GIN_PROXY_ENABLE=1>
+  $<$<BOOL:${NCCL_OS_WINDOWS}>:NCCL_OS_WINDOWS>
+  $<$<BOOL:${NCCL_OS_WINDOWS}>:WIN32_LEAN_AND_MEAN>
+  $<$<BOOL:${NCCL_OS_WINDOWS}>:NOMINMAX>
+  $<$<BOOL:${NCCL_OS_WINDOWS}>:NCCL_GIN_PROXY_ENABLE=0>
+)
+
+if(MSVC)
+  target_compile_options(nccl_tests_options INTERFACE
+    $<$<COMPILE_LANGUAGE:CXX>:/W3>
+    $<$<COMPILE_LANGUAGE:CXX>:/wd4267>
+    $<$<COMPILE_LANGUAGE:CXX>:/wd4244>
+    $<$<COMPILE_LANGUAGE:CXX>:/wd4996>
+    $<$<COMPILE_LANGUAGE:CXX>:/wd4146>
+    $<$<COMPILE_LANGUAGE:CXX>:/wd4197>
+    $<$<COMPILE_LANGUAGE:CXX>:/wd5105>
+    $<$<COMPILE_LANGUAGE:CXX>:/wd4805>
+    $<$<COMPILE_LANGUAGE:CXX>:/wd4018>
+    $<$<COMPILE_LANGUAGE:CXX>:/FS>
+    $<$<COMPILE_LANGUAGE:CXX>:/Zc:preprocessor>
+    $<$<COMPILE_LANGUAGE:CUDA>:--expt-extended-lambda>
+    $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>
+    $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/Zc:preprocessor>
+    $<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=/FS>
+  )
+else()
+  target_compile_options(nccl_tests_options INTERFACE
+    $<$<COMPILE_LANGUAGE:CXX>:-Wall>
+    $<$<COMPILE_LANGUAGE:CXX>:-Wno-unused-function>
+    $<$<COMPILE_LANGUAGE:CXX>:-Wno-sign-compare>
+    $<$<COMPILE_LANGUAGE:CUDA>:--expt-extended-lambda>
+    $<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>
+  )
+endif()
+
+add_subdirectory(os)
+add_subdirectory(verifiable)
+add_subdirectory(src)
@@ -0,0 +1,24 @@
+#
+# Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved.
+#
+# See LICENSE.txt for license information
+#
+
+# OS abstraction layer for NCCL tests.
+
+if(NCCL_OS_LINUX)
+  set(OS_SOURCES linux.cc)
+elseif(NCCL_OS_WINDOWS)
+  set(OS_SOURCES windows.cc)
+else()
+  message(FATAL_ERROR "Unsupported OS for nccl-tests: ${CMAKE_SYSTEM_NAME}")
+endif()
+
+add_library(nccl_test_os STATIC ${OS_SOURCES})
+
+target_include_directories(nccl_test_os
+  PUBLIC
+    ${CMAKE_CURRENT_SOURCE_DIR}
+)
+
+target_link_libraries(nccl_test_os PRIVATE nccl_tests_options)
@@ -0,0 +1,35 @@
+#
+# Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved.
+#
+# See LICENSE.txt for license information
+#
+
+include ../src/common.mk
+
+.PHONY: build clean
+
+BUILDDIR ?= ../build
+DST_DIR := $(BUILDDIR)/os
+
+ifneq ($(OS),Windows_NT)
+  OS_SRC := linux.cc
+else
+  OS_SRC := windows.cc
+endif
+
+OBJ_FILES := $(OS_SRC:%.cc=${DST_DIR}/%.o)
+
+build: ${DST_DIR}/libnccl_test_os.a
+
+clean:
+	rm -rf ${DST_DIR}
+
+${DST_DIR}/%.o: %.cc os.h
+	@printf "Compiling  %-35s > %s\n" $< $@
+	@mkdir -p ${DST_DIR}
+	$(CXX) $(CXXFLAGS) -I../src -o $@ -c $<
+
+${DST_DIR}/libnccl_test_os.a: $(OBJ_FILES)
+	@printf "Creating archive %-35s > %s\n" $^ $@
+	@mkdir -p ${DST_DIR}
+	ar rcs $@ $^
@@ -0,0 +1,71 @@
+/*************************************************************************
+ * Copyright (c) 2016-2025, NVIDIA CORPORATION. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#include "os.h"
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <strings.h>
+
+#define HOSTID_FILE "/proc/sys/kernel/random/boot_id"
+
+static uint64_t getHash(const char* string, size_t n) {
+  uint64_t result = 5381;
+  for (size_t c = 0; c < n; c++) {
+    result = ((result << 5) + result) ^ string[c];
+  }
+  return result;
+}
+
+uint64_t ncclTestGetHostHash(const char* hostname) {
+  char hostHash[1024];
+
+  snprintf(hostHash, sizeof(hostHash), "%s", hostname);
+  int offset = strlen(hostHash);
+
+  FILE *file = fopen(HOSTID_FILE, "r");
+  if (file != NULL) {
+    char *p;
+    if (fscanf(file, "%ms", &p) == 1) {
+      strncpy(hostHash+offset, p, sizeof(hostHash)-offset-1);
+      free(p);
+    }
+    fclose(file);
+  }
+
+  hostHash[sizeof(hostHash)-1] = '\0';
+
+  return getHash(hostHash, strlen(hostHash));
+}
+
+int ncclTestGetHostname(char* name, size_t len) {
+  return gethostname(name, len);
+}
+
+int ncclTestGetPid() {
+  return (int)getpid();
+}
+
+int ncclTestStrcasecmp(const char* s1, const char* s2) {
+  return strcasecmp(s1, s2);
+}
+
+int ncclTestStrncasecmp(const char* s1, const char* s2, size_t n) {
+  return strncasecmp(s1, s2, n);
+}
+
+int ncclTestAsprintf(char** strp, const char* fmt, ...) {
+  va_list args;
+  va_start(args, fmt);
+  int result = vasprintf(strp, fmt, args);
+  va_end(args);
+  return result;
+}
+
+void ncclTestSetlinebuf(FILE* stream) {
+  setlinebuf(stream);
+}
@@ -0,0 +1,71 @@
+/*************************************************************************
+ * Copyright (c) 2016-2025, NVIDIA CORPORATION. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#ifndef __NCCL_TEST_OS_H__
+#define __NCCL_TEST_OS_H__
+
+#include <stdint.h>
+#include <stddef.h>
+#include <stdio.h>
+#include <stdarg.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+uint64_t ncclTestGetHostHash(const char* hostname);
+int ncclTestGetHostname(char* name, size_t len);
+int ncclTestGetPid();
+int ncclTestStrcasecmp(const char* s1, const char* s2);
+int ncclTestStrncasecmp(const char* s1, const char* s2, size_t n);
+int ncclTestAsprintf(char** strp, const char* fmt, ...);
+void ncclTestSetlinebuf(FILE* stream);
+
+#ifdef __cplusplus
+}
+#endif
+
+#if defined(NCCL_OS_LINUX)
+
+#include <getopt.h>
+#include <unistd.h>
+
+#define NCCL_WEAK __attribute__((weak))
+
+#elif defined(NCCL_OS_WINDOWS)
+
+#include <windows.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+struct option {
+  const char *name;
+  int has_arg;
+  int *flag;
+  int val;
+};
+
+#define no_argument         0
+#define required_argument   1
+#define optional_argument   2
+
+extern char *optarg;
+extern int optind, opterr, optopt;
+
+int getopt_long(int argc, char * const argv[], const char *optstring,
+                const struct option *longopts, int *longindex);
+
+#ifdef __cplusplus
+}
+#endif
+
+#define NCCL_WEAK
+
+#endif
+
+#endif
@@ -0,0 +1,225 @@
+/*************************************************************************
+ * Copyright (c) 2016-2025, NVIDIA CORPORATION. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#include "os.h"
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <windows.h>
+
+static uint64_t getHash(const char* string, size_t n) {
+  uint64_t result = 5381;
+  for (size_t c = 0; c < n; c++) {
+    result = ((result << 5) + result) ^ string[c];
+  }
+  return result;
+}
+
+static bool getWindowsMachineGuid(char* guid, size_t len) {
+  HKEY hKey;
+  LONG result = RegOpenKeyExA(HKEY_LOCAL_MACHINE,
+                              "SOFTWARE\\Microsoft\\Cryptography",
+                              0,
+                              KEY_READ,
+                              &hKey);
+  if (result != ERROR_SUCCESS) {
+    return false;
+  }
+
+  DWORD dataSize = (DWORD)len;
+  DWORD dataType;
+  result = RegQueryValueExA(hKey,
+                            "MachineGuid",
+                            NULL,
+                            &dataType,
+                            (LPBYTE)guid,
+                            &dataSize);
+
+  RegCloseKey(hKey);
+
+  return result == ERROR_SUCCESS && dataType == REG_SZ;
+}
+
+uint64_t ncclTestGetHostHash(const char* hostname) {
+  char hostHash[1024];
+
+  strncpy(hostHash, hostname, sizeof(hostHash));
+  int offset = strlen(hostHash);
+
+  char machineGuid[256];
+  if (getWindowsMachineGuid(machineGuid, sizeof(machineGuid))) {
+    strncpy(hostHash+offset, machineGuid, sizeof(hostHash)-offset-1);
+  }
+
+  hostHash[sizeof(hostHash)-1] = '\0';
+
+  return getHash(hostHash, strlen(hostHash));
+}
+
+int ncclTestGetHostname(char* name, size_t len) {
+  DWORD size = (DWORD)len;
+  if (!GetComputerNameA(name, &size)) {
+    return -1;
+  }
+  return 0;
+}
+
+int ncclTestGetPid() {
+  return (int)GetCurrentProcessId();
+}
+
+int ncclTestStrcasecmp(const char* s1, const char* s2) {
+  return _stricmp(s1, s2);
+}
+
+int ncclTestStrncasecmp(const char* s1, const char* s2, size_t n) {
+  return _strnicmp(s1, s2, n);
+}
+
+int ncclTestAsprintf(char** strp, const char* fmt, ...) {
+  va_list args;
+  va_start(args, fmt);
+  int size = _vscprintf(fmt, args);
+  va_end(args);
+  if (size < 0) {
+    return -1;
+  }
+
+  *strp = (char*)malloc(size + 1);
+  if (*strp == NULL) {
+    return -1;
+  }
+
+  va_start(args, fmt);
+  int result = vsprintf_s(*strp, size + 1, fmt, args);
+  va_end(args);
+  return result;
+}
+
+void ncclTestSetlinebuf(FILE* stream) {
+  setvbuf(stream, NULL, _IONBF, 0);
+}
+
+extern "C" {
+
+int optind = 1;
+int opterr = 1;
+int optopt = 0;
+char* optarg = NULL;
+
+static char* nextchar = NULL;
+
+int getopt_long(int argc, char* const argv[], const char* optstring,
+                const struct option* longopts, int* longindex) {
+  if (optind == 0) {
+    optind = 1;
+  }
+
+  optarg = NULL;
+
+  if (optind >= argc || argv[optind] == NULL) {
+    return -1;
+  }
+
+  const char* arg = argv[optind];
+
+  if (arg[0] == '-' && arg[1] == '-') {
+    const char* name = arg + 2;
+    const char* equals = strchr(name, '=');
+    size_t name_len = equals ? (equals - name) : strlen(name);
+
+    for (int i = 0; longopts[i].name != NULL; i++) {
+      if (strncmp(longopts[i].name, name, name_len) == 0 &&
+          strlen(longopts[i].name) == name_len) {
+        if (longindex) {
+          *longindex = i;
+        }
+
+        optind++;
+
+        if (longopts[i].has_arg == required_argument) {
+          if (equals) {
+            optarg = (char*)(equals + 1);
+          } else if (optind < argc) {
+            optarg = argv[optind++];
+          } else {
+            if (opterr) {
+              fprintf(stderr, "Option --%s requires an argument\n", longopts[i].name);
+            }
+            return '?';
+          }
+        } else if (longopts[i].has_arg == optional_argument) {
+          if (equals) {
+            optarg = (char*)(equals + 1);
+          }
+        }
+
+        if (longopts[i].flag) {
+          *(longopts[i].flag) = longopts[i].val;
+          return 0;
+        }
+        return longopts[i].val;
+      }
+    }
+
+    if (opterr) {
+      fprintf(stderr, "Unrecognized option: %s\n", arg);
+    }
+    optind++;
+    return '?';
+  }
+
+  if (arg[0] == '-' && arg[1] != '\0') {
+    if (nextchar == NULL || *nextchar == '\0') {
+      nextchar = (char*)(arg + 1);
+    }
+
+    char c = *nextchar++;
+    const char* opt = strchr(optstring, c);
+
+    if (opt == NULL) {
+      optopt = c;
+      if (opterr) {
+        fprintf(stderr, "Invalid option: -%c\n", c);
+      }
+      if (*nextchar == '\0') {
+        optind++;
+        nextchar = NULL;
+      }
+      return '?';
+    }
+
+    if (opt[1] == ':') {
+      if (*nextchar != '\0') {
+        optarg = nextchar;
+        optind++;
+        nextchar = NULL;
+      } else if (optind + 1 < argc) {
+        optarg = argv[++optind];
+        optind++;
+        nextchar = NULL;
+      } else {
+        optopt = c;
+        if (opterr) {
+          fprintf(stderr, "Option -%c requires an argument\n", c);
+        }
+        optind++;
+        nextchar = NULL;
+        return '?';
+      }
+    } else if (*nextchar == '\0') {
+      optind++;
+      nextchar = NULL;
+    }
+
+    return c;
+  }
+
+  return -1;
+}
+
+}
@@ -0,0 +1,55 @@
+#
+# Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved.
+#
+# See LICENSE.txt for license information
+#
+
+set(PERF_COMMON_SOURCES
+  common.cu
+  util.cu
+  timer.cc
+)
+
+set(COLLS
+  all_reduce
+  all_gather
+  broadcast
+  reduce_scatter
+  reduce
+  alltoall
+  alltoallv
+  gather
+  scatter
+  sendrecv
+  hypercube
+)
+
+foreach(COLL IN LISTS COLLS)
+  add_executable(${COLL}_perf ${COLL}.cu ${PERF_COMMON_SOURCES})
+  set_target_properties(${COLL}_perf PROPERTIES
+    RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}
+  )
+
+  target_include_directories(${COLL}_perf
+    PRIVATE
+      ${CMAKE_CURRENT_SOURCE_DIR}
+      ${CMAKE_SOURCE_DIR}/os
+      ${CMAKE_SOURCE_DIR}/verifiable
+  )
+
+  target_link_libraries(${COLL}_perf
+    PRIVATE
+      nccl_tests_options
+      nccl
+      CUDA::cudart
+      CUDA::cuda_driver
+      verifiable
+      nccl_test_os
+  )
+
+  if(MPI_FOUND)
+    target_compile_definitions(${COLL}_perf PRIVATE MPI_SUPPORT)
+    target_include_directories(${COLL}_perf PRIVATE ${MPI_CXX_INCLUDE_DIRS})
+    target_link_libraries(${COLL}_perf PRIVATE MPI::MPI_CXX)
+  endif()
+endforeach()
@@ -37,7 +37,13 @@ OBJ_FILES := $(SRC_FILES:%.cu=${DST_DIR}/%.o)
 BIN_FILES_LIST := all_reduce all_gather broadcast reduce_scatter reduce alltoall alltoallv scatter gather sendrecv hypercube
 BIN_FILES := $(BIN_FILES_LIST:%=${DST_DIR}/%_perf${NAME_SUFFIX})

-build: ${BIN_FILES}
+TEST_OS_SRCDIR := ../os
+TEST_OS_BUILDDIR := $(BUILDDIR)/os
+TEST_OS_LIB := $(TEST_OS_BUILDDIR)/libnccl_test_os.a
+NVCUFLAGS += -I$(TEST_OS_SRCDIR)
+NVLDFLAGS += -L$(TEST_OS_BUILDDIR) -lnccl_test_os
+
+build: os.build ${BIN_FILES}

 clean:
 	rm -rf ${DST_DIR}
@@ -46,6 +52,12 @@ TEST_VERIFIABLE_SRCDIR := ../verifiable
 TEST_VERIFIABLE_BUILDDIR := $(BUILDDIR)/verifiable
 include ../verifiable/verifiable.mk

+$(TEST_OS_LIB): os.build
+	@:
+
+os.%:
+	${MAKE} -C $(TEST_OS_SRCDIR) $* BUILDDIR=$(BUILDDIR)
+
 .PRECIOUS: ${DST_DIR}/%.o

 ${DST_DIR}/%.o: %.cu common.h util.h $(TEST_VERIFIABLE_HDRS)
@@ -64,12 +76,12 @@ ${DST_DIR}/timer.o: timer.cc timer.h
 	$(CXX) $(CXXFLAGS) -o $@ -c $<

 ifeq ($(DSO), 1)
-${DST_DIR}/%_perf$(NAME_SUFFIX): ${DST_DIR}/%.o ${DST_DIR}/common$(NAME_SUFFIX).o ${DST_DIR}/util$(NAME_SUFFIX).o ${DST_DIR}/timer.o $(TEST_VERIFIABLE_LIBS)
+${DST_DIR}/%_perf$(NAME_SUFFIX): ${DST_DIR}/%.o ${DST_DIR}/common$(NAME_SUFFIX).o ${DST_DIR}/util$(NAME_SUFFIX).o ${DST_DIR}/timer.o $(TEST_VERIFIABLE_LIBS) $(TEST_OS_LIB)
 	@printf "Linking  %-35s > %s\n" $< $@
 	@mkdir -p ${DST_DIR}
 	$(NVCC) -o $@ $(NVCUFLAGS) $^ -L$(TEST_VERIFIABLE_BUILDDIR) -lverifiable ${NVLDFLAGS} -Xlinker "--enable-new-dtags" -Xlinker "-rpath,\$$ORIGIN:\$$ORIGIN/verifiable"
 else
-${DST_DIR}/%_perf$(NAME_SUFFIX):${DST_DIR}/%.o ${DST_DIR}/common$(NAME_SUFFIX).o ${DST_DIR}/util$(NAME_SUFFIX).o ${DST_DIR}/timer.o $(TEST_VERIFIABLE_OBJS)
+${DST_DIR}/%_perf$(NAME_SUFFIX):${DST_DIR}/%.o ${DST_DIR}/common$(NAME_SUFFIX).o ${DST_DIR}/util$(NAME_SUFFIX).o ${DST_DIR}/timer.o $(TEST_VERIFIABLE_OBJS) $(TEST_OS_LIB)
 	@printf "Linking  %-35s > %s\n" $< $@
 	@mkdir -p ${DST_DIR}
 	$(NVCC) -o $@ $(NVCUFLAGS) $^ ${NVLDFLAGS}
@@ -77,4 +89,3 @@ endif

 clean_intermediates:
 	rm -f ${DST_DIR}/*.o $(TEST_VERIFIABLE_OBJS)
-
@@ -89,9 +89,7 @@ testResult_t AllGatherRunTest(struct threadArgs* args, int root, ncclDataType_t
  return testSuccess;
 }

-struct testEngine allGatherEngine = {
-  .getBuffSize = AllGatherGetBuffSize,
-  .runTest = AllGatherRunTest
+NCCL_WEAK struct testEngine ncclTestEngine = {
+  /* .getBuffSize = */ AllGatherGetBuffSize,
+  /* .runTest = */ AllGatherRunTest
 };
-
-#pragma weak ncclTestEngine=allGatherEngine
@@ -577,12 +577,13 @@ testResult_t AllReduceRunTest(struct threadArgs* args, int root, ncclDataType_t
  return testSuccess;
 }

-struct testEngine allReduceEngine = {
-  .getBuffSize = AllReduceGetBuffSize,
-  .runTest = AllReduceRunTest,
+NCCL_WEAK struct testEngine ncclTestEngine = {
+  /* .getBuffSize = */ AllReduceGetBuffSize,
+  /* .runTest = */ AllReduceRunTest,
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
+  /* .initCommConfig = */ nullptr,
+#endif
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
-  .getDevCommRequirements = AllReduceGetDevCommRequirements
+  /* .getDevCommRequirements = */ AllReduceGetDevCommRequirements
 #endif
 };
-
-#pragma weak ncclTestEngine=allReduceEngine
@@ -11,7 +11,9 @@
 #include "vector_types.h"
 #endif

+#if defined(NCCL_OS_LINUX)
 #pragma weak ncclAlltoAll
+#endif

 void AlltoAllGetCollByteCount(size_t *sendcount, size_t *recvcount, size_t *paramcount, size_t *sendInplaceOffset, size_t *recvInplaceOffset, size_t count, size_t eltSize, int nranks) {
  *paramcount = (count/nranks) & -(16/eltSize);
@@ -72,6 +74,7 @@ testResult_t AlltoAllGetDevCommRequirements(int deviceImpl, ncclDevCommRequireme
      }
      reqs->lsaBarrierCount = deviceCtaCount;
      return testSuccess;
+    #if defined(NCCL_OS_LINUX)
    case 3: // GinAlltoAllKernel
    case 4: // HybridAlltoAllKernel (LSA+GIN)
      if (commProperties.ginType == NCCL_GIN_TYPE_NONE) {
@@ -86,6 +89,7 @@ testResult_t AlltoAllGetDevCommRequirements(int deviceImpl, ncclDevCommRequireme
      reqs->ginForceEnable = true;
 #endif
      return testSuccess;
+    #endif
    default:
      return testNotImplemented;
  }
@@ -101,7 +105,7 @@ bool AlltoAllGetDevCommRequirements(int deviceImpl, ncclDevCommRequirements* req
    case 2: // NvlAlltoAllKernelOptimized
      reqs->lsaBarrierCount = deviceCtaCount;
      return true;
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7)
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7) && defined(NCCL_OS_LINUX)
    case 3: // GinAlltoAllKernel
    case 4: // HybridAlltoAllKernel (LSA+GIN)
      reqs->barrierCount = deviceCtaCount;
@@ -225,7 +229,7 @@ __global__ void NvlAlltoAllKernelOptimized(ncclWindow_t sendwin, size_t sendoffs
  bar.sync(ncclCoopCta(), cuda::memory_order_release);
 }

-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7)
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7) && defined(NCCL_OS_LINUX)
 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) {
  int ginContext = 0;
@@ -341,7 +345,7 @@ testResult_t AlltoAllRunColl(void* sendbuff, size_t sendoffset, void* recvbuff,
        TESTCHECK(testLaunchDeviceKernel(SPECIALIZE_KERNEL(NvlAlltoAllKernelOptimized, type, op), sendbuff, sendoffset, recvbuff, recvoffset, count, type, op, root, comm, stream));
        return testSuccess;
 #endif
-#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7)
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7) && defined(NCCL_OS_LINUX)
      case 3:
        TESTCHECK(testLaunchDeviceKernel(SPECIALIZE_KERNEL(GinAlltoAllKernel, type, op), sendbuff, sendoffset, recvbuff, recvoffset, count, type, op, root, comm, stream));
        return testSuccess;
@@ -391,12 +395,13 @@ testResult_t AlltoAllRunTest(struct threadArgs* args, int root, ncclDataType_t t
  return testSuccess;
 }

-struct testEngine alltoAllEngine = {
-  .getBuffSize = AlltoAllGetBuffSize,
-  .runTest = AlltoAllRunTest,
+NCCL_WEAK struct testEngine ncclTestEngine = {
+  /* .getBuffSize = */ AlltoAllGetBuffSize,
+  /* .runTest = */ AlltoAllRunTest,
+#if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
+  /* .initCommConfig = */ nullptr,
+#endif
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
-  .getDevCommRequirements = AlltoAllGetDevCommRequirements
+  /* .getDevCommRequirements = */ AlltoAllGetDevCommRequirements
 #endif
 };
-
-#pragma weak ncclTestEngine=alltoAllEngine
@@ -117,10 +117,10 @@
 #include <stdio.h>
 #include <ctype.h>
 #include <errno.h>
-#include <pthread.h>
+#include <atomic>
+#include <mutex>
 #include <stdint.h>
 #include <stdlib.h>
-#include <sys/stat.h>

 #define PRINT if (is_main_thread) printf

@@ -134,16 +134,16 @@ static int traffic_matrix_dim = 0;
 static double traffic_matrix_scale = 1.0; // scale factor applied to traffic matrix values
 static double distance_weighted_spread = 1.0; // 0.0 => uniform, 1.0 => fully distance-weighted (default)

-static pthread_mutex_t alltoallv_lock = PTHREAD_MUTEX_INITIALIZER;
+static std::mutex alltoallv_lock;
 static int alltoallv_inited = 0;

 // atomics used for error flag, since init can be multi-threaded
-static int alltoallv_error_set = 0;
+static std::atomic<int> alltoallv_error_set{0};
 static inline int AlltoAllvHasError() {
-  return __atomic_load_n(&alltoallv_error_set, __ATOMIC_RELAXED);
+  return alltoallv_error_set.load(std::memory_order_relaxed);
 }
 static inline void AlltoAllvSetError() {
-  __atomic_store_n(&alltoallv_error_set, 1, __ATOMIC_RELAXED);
+  alltoallv_error_set.store(1, std::memory_order_relaxed);
 }

 static void AlltoAllvParseEnv() {
@@ -175,7 +175,6 @@ static int AlltoAllvReadTrafficMatrix() {
  FILE* f = NULL;
  char* buf = NULL;
  size_t* tmp_matrix_data = NULL;
-  struct stat st;
  size_t fileLen = 0;
  char* p = NULL;
  char* end = NULL;
@@ -184,6 +183,7 @@ static int AlltoAllvReadTrafficMatrix() {
  size_t rows = 0, cols = 0, tmp_cols = 0;
  size_t n_elts = 0;
  size_t k = 0;
+  long pos = 0;
  int rc = 1;

  f = fopen(traffic_matrix_file, "rb");
@@ -192,11 +192,20 @@ static int AlltoAllvReadTrafficMatrix() {
    goto exit;
  }

-  if (fstat(fileno(f), &st) != 0) {
-    PRINT("Unable to stat alltoallv traffic matrix file %s.\n", traffic_matrix_file);
+  if (fseek(f, 0, SEEK_END) != 0) {
+    PRINT("Unable to seek alltoallv traffic matrix file %s.\n", traffic_matrix_file);
+    goto exit;
+  }
+  pos = ftell(f);
+  if (pos < 0) {
+    PRINT("Unable to tell alltoallv traffic matrix file size %s.\n", traffic_matrix_file);
+    goto exit;
+  }
+  fileLen = (size_t)pos;
+  if (fseek(f, 0, SEEK_SET) != 0) {
+    PRINT("Unable to rewind alltoallv traffic matrix file %s.\n", traffic_matrix_file);
    goto exit;
  }
-  fileLen = (size_t)st.st_size;

  buf = (char*)malloc(fileLen + 1);
  if (!buf) {
@@ -319,7 +328,7 @@ exit:
 }

 static void AlltoAllvInit(int nranks) {
-  pthread_mutex_lock(&alltoallv_lock);
+  std::lock_guard<std::mutex> lock(alltoallv_lock);

  if (!alltoallv_inited) {
    AlltoAllvParseEnv();
@@ -339,7 +348,6 @@ static void AlltoAllvInit(int nranks) {
    AlltoAllvSetError();
  }

-  pthread_mutex_unlock(&alltoallv_lock);
 }

 // get peer-to-peer bytes from traffic matrix
@@ -569,14 +577,13 @@ void AlltoAllvGetBuffSize(size_t *sendcount, size_t *recvcount, size_t count, in
    // if using traffic matrix, validate that the max buffer size is large enough
    size_t total_bytes_req = MAX(*sendcount, *recvcount);
    if (count < total_bytes_req) {
-      pthread_mutex_lock(&alltoallv_lock);
+      std::lock_guard<std::mutex> lock(alltoallv_lock);
      if (!AlltoAllvHasError()) {
        if (is_main_proc)
          printf("maxBytes (-e) must be at least %zu bytes as required by traffic matrix file %s (got %zu). Increase -e.\n",
                 total_bytes_req, traffic_matrix_file, count);
        AlltoAllvSetError();
      }
-      pthread_mutex_unlock(&alltoallv_lock);
      *sendcount = *recvcount = 0;
      return;
    }
@@ -616,9 +623,7 @@ testResult_t AlltoAllvRunTest(struct threadArgs* args, int root, ncclDataType_t
  return testSuccess;
 }

-struct testEngine alltoAllvEngine = {
-  .getBuffSize = AlltoAllvGetBuffSize,
-  .runTest = AlltoAllvRunTest
+NCCL_WEAK struct testEngine ncclTestEngine = {
+  /* .getBuffSize = */ AlltoAllvGetBuffSize,
+  /* .runTest = */ AlltoAllvRunTest
 };
-
-#pragma weak ncclTestEngine=alltoAllvEngine
@@ -106,9 +106,7 @@ testResult_t BroadcastRunTest(struct threadArgs* args, int root, ncclDataType_t
  return testSuccess;
 }

-struct testEngine broadcastEngine = {
-  .getBuffSize = BroadcastGetBuffSize,
-  .runTest = BroadcastRunTest
+NCCL_WEAK struct testEngine ncclTestEngine = {
+  /* .getBuffSize = */ BroadcastGetBuffSize,
+  /* .runTest = */ BroadcastRunTest
 };
-
-#pragma weak ncclTestEngine=broadcastEngine
@@ -5,25 +5,27 @@
 ************************************************************************/

 #include "common.h"
-#include <pthread.h>
 #include <cstdio>
 #include <type_traits>
 #include <limits>
-#include <getopt.h>
-#include <libgen.h>
 #include <string.h>
 #include <ctype.h>
+#include <mutex>
+#include <condition_variable>
+#include <thread>
+#include <vector>
 #include "cuda.h"
-#include <errno.h>     /* program_invocation_short_name */

 #include "util.h"
 #include "../verifiable/verifiable.h"

+#if defined(NCCL_OS_LINUX)
 #pragma weak ncclCommWindowRegister
 #pragma weak ncclCommWindowDeregister
 #pragma weak ncclDevCommCreate
 #pragma weak ncclDevCommDestroy
 #pragma weak ncclCommQueryProperties
+#endif

 #define DIVUP(x, y) \
    (((x)+(y)-1)/(y))
@@ -71,12 +73,17 @@ int test_ncclVersion = 0; // init'd with ncclGetVersion()
 #endif

 // For libnccl's < 2.13
+#if defined(NCCL_OS_LINUX)
 extern "C" __attribute__((weak)) char const* ncclGetLastError(ncclComm_t comm) {
  return "";
 }
+#elif defined(NCCL_OS_WINDOWS)
+extern "C" char const* ncclGetLastError(ncclComm_t comm);
+#endif

 int is_main_proc = 0;
 thread_local int is_main_thread = 0;
+static const char* programName = nullptr;

 // Command line parameter defaults
 int nThreads = 1;
@@ -155,7 +162,7 @@ static const char *getExtension(const char *path) {

 static output_file_type_t classifyOutputFile(const char *filename) {
  const char *extension = getExtension(filename);
-  if (extension != nullptr && strcasecmp(extension, "json") == 0) {
+  if (extension != nullptr && ncclTestStrcasecmp(extension, "json") == 0) {
    return JSON_FILE_OUTPUT;
  }

@@ -268,28 +275,28 @@ testResult_t InitData(void* data, const size_t count, size_t offset, ncclDataTyp

 void Barrier(struct threadArgs *args) {
  thread_local int epoch = 0;
-  static pthread_mutex_t lock[2] = {PTHREAD_MUTEX_INITIALIZER, PTHREAD_MUTEX_INITIALIZER};
-  static pthread_cond_t cond[2] = {PTHREAD_COND_INITIALIZER, PTHREAD_COND_INITIALIZER};
+  static std::mutex lock[2];
+  static std::condition_variable cond[2];
  static int counter[2] = {0, 0};

-  pthread_mutex_lock(&lock[epoch]);
+  std::unique_lock<std::mutex> ul(lock[epoch]);
  if(++counter[epoch] == args->nThreads)
-    pthread_cond_broadcast(&cond[epoch]);
+    cond[epoch].notify_all();

  if(args->thread+1 == args->nThreads) {
    while(counter[epoch] != args->nThreads)
-      pthread_cond_wait(&cond[epoch], &lock[epoch]);
+      cond[epoch].wait(ul);
    #ifdef MPI_SUPPORT
      MPI_Barrier(MPI_COMM_WORLD);
    #endif
    counter[epoch] = 0;
-    pthread_cond_broadcast(&cond[epoch]);
+    cond[epoch].notify_all();
  }
  else {
    while(counter[epoch] != 0)
-      pthread_cond_wait(&cond[epoch], &lock[epoch]);
+      cond[epoch].wait(ul);
  }
-  pthread_mutex_unlock(&lock[epoch]);
+  ul.unlock();
  epoch ^= 1;
 }

@@ -299,12 +306,12 @@ void Barrier(struct threadArgs *args) {
 template<typename T>
 void Allreduce(struct threadArgs* args, T* value, int average) {
  thread_local int epoch = 0;
-  static pthread_mutex_t lock[2] = {PTHREAD_MUTEX_INITIALIZER, PTHREAD_MUTEX_INITIALIZER};
-  static pthread_cond_t cond[2] = {PTHREAD_COND_INITIALIZER, PTHREAD_COND_INITIALIZER};
+  static std::mutex lock[2];
+  static std::condition_variable cond[2];
  static T accumulator[2];
  static int counter[2] = {0, 0};

-  pthread_mutex_lock(&lock[epoch]);
+  std::unique_lock<std::mutex> ul(lock[epoch]);
  if(counter[epoch] == 0) {
    if(average != 0 || args->thread == 0) accumulator[epoch] = *value;
  } else {
@@ -318,11 +325,11 @@ void Allreduce(struct threadArgs* args, T* value, int average) {
  }

  if(++counter[epoch] == args->nThreads)
-    pthread_cond_broadcast(&cond[epoch]);
+    cond[epoch].notify_all();

  if(args->thread+1 == args->nThreads) {
    while(counter[epoch] != args->nThreads)
-      pthread_cond_wait(&cond[epoch], &lock[epoch]);
+      cond[epoch].wait(ul);

    #ifdef MPI_SUPPORT
    if(average != 0) {
@@ -340,13 +347,13 @@ void Allreduce(struct threadArgs* args, T* value, int average) {

    if(average == 1) accumulator[epoch] /= args->totalProcs*args->nThreads;
    counter[epoch] = 0;
-    pthread_cond_broadcast(&cond[epoch]);
+    cond[epoch].notify_all();
  }
  else {
    while(counter[epoch] != 0)
-      pthread_cond_wait(&cond[epoch], &lock[epoch]);
+      cond[epoch].wait(ul);
  }
-  pthread_mutex_unlock(&lock[epoch]);
+  ul.unlock();

  *value = accumulator[epoch];
  epoch ^= 1;
@@ -451,7 +458,7 @@ testResult_t testStreamSynchronize(int ngpus, cudaStream_t* streams, ncclComm_t*
   }

   // We might want to let other threads (including NCCL threads) use the CPU.
-   if (idle) sched_yield();
+   if (idle) std::this_thread::yield();
  }
  free(done);
  return testSuccess;
@@ -567,8 +574,8 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
  Barrier(args);

 #if CUDART_VERSION >= 11030
-  cudaGraph_t graphs[args->nGpus];
-  cudaGraphExec_t graphExec[args->nGpus];
+  std::vector<cudaGraph_t> graphs(args->nGpus);
+  std::vector<cudaGraphExec_t> graphExec(args->nGpus);
  if (cudaGraphLaunches >= 1) {
    // Begin cuda graph capture
    for (int i=0; i<args->nGpus; i++) {
@@ -595,11 +602,11 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
  if (cudaGraphLaunches >= 1) {
    // End cuda graph capture
    for (int i=0; i<args->nGpus; i++) {
-      CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs+i));
+      CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs.data()+i));
    }
    // Instantiate cuda graph
    for (int i=0; i<args->nGpus; i++) {
-      CUDACHECK(cudaGraphInstantiate(graphExec+i, graphs[i], NULL, NULL, 0));
+      CUDACHECK(cudaGraphInstantiate(graphExec.data()+i, graphs[i], NULL, NULL, 0));
    }
    // Resync CPU, restart timing, launch cuda graph
    Barrier(args);
@@ -636,7 +643,7 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
  Barrier(args);

  int64_t wrongElts = 0;
-  static __thread int rep = 0;
+  static thread_local int rep = 0;
  rep++;
  for (int c = 0; c < datacheck; c++) {
      // Initialize sendbuffs, recvbuffs and expected
@@ -658,11 +665,11 @@ testResult_t BenchTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t
      if (cudaGraphLaunches >= 1) {
        // End cuda graph capture
        for (int i=0; i<args->nGpus; i++) {
-          CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs+i));
+          CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs.data()+i));
        }
        // Instantiate cuda graph
        for (int i=0; i<args->nGpus; i++) {
-          CUDACHECK(cudaGraphInstantiate(graphExec+i, graphs[i], NULL, NULL, 0));
+          CUDACHECK(cudaGraphInstantiate(graphExec.data()+i, graphs[i], NULL, NULL, 0));
        }
        // Launch cuda graph
        for (int i=0; i<args->nGpus; i++) {
@@ -897,13 +904,11 @@ testResult_t threadInit(struct threadArgs* args) {
  return testSuccess;
 }

-void* threadLauncher(void* thread_) {
-  struct testThread* thread = (struct testThread*)thread_;
+static void threadLauncher(struct testThread* thread) {
  thread->ret = thread->func(&thread->args);
-  return NULL;
 }
 testResult_t threadLaunch(struct testThread* thread) {
-  pthread_create(&thread->thread, NULL, threadLauncher, thread);
+  thread->thread = std::thread(threadLauncher, thread);
  return testSuccess;
 }

@@ -925,7 +930,18 @@ testResult_t run(); // Main function

 int main(int argc, char* argv[], char **envp) {
  // Make sure everyline is flushed so that we see the progress of the test
-  setlinebuf(stdout);
+  ncclTestSetlinebuf(stdout);
+  const char* slash = strrchr(argv[0], '/');
+  const char* backslash = strrchr(argv[0], '\\');
+  const char* sep = nullptr;
+  if (slash && backslash) {
+    sep = slash > backslash ? slash : backslash;
+  } else if (slash) {
+    sep = slash;
+  } else {
+    sep = backslash;
+  }
+  programName = sep ? sep + 1 : argv[0];

  #if NCCL_VERSION_CODE >= NCCL_VERSION(2,4,0)
    ncclGetVersion(&test_ncclVersion);
@@ -1182,7 +1198,7 @@ int main(int argc, char* argv[], char **envp) {
            "[-M,--memory <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]));
+          programName);
        return 0;
    }
  }
@@ -1227,7 +1243,7 @@ static bool parseInt(char *s, int *num) {
  while (*s && isspace(*s)) ++s;
  if (!*s) return false;

-  if (strncasecmp(s, "0b", 2) == 0)
+  if (ncclTestStrncasecmp(s, "0b", 2) == 0)
    *num = (int)strtoul(s + 2, &p, 2);
  else
    *num = (int)strtoul(s, &p, 0);
@@ -1247,9 +1263,9 @@ testResult_t run() {
 #ifdef MPI_SUPPORT
  MPI_Comm_size(MPI_COMM_WORLD, &totalProcs);
  MPI_Comm_rank(MPI_COMM_WORLD, &proc);
-  uint64_t hostHashs[totalProcs];
+  std::vector<uint64_t> hostHashs(totalProcs);
  hostHashs[proc] = getHostHash(hostname);
-  MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs, sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD);
+  MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs.data(), sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD);
  for (int p=0; p<totalProcs; p++) {
    if (p == proc) break;
    if (hostHashs[p] == hostHashs[proc]) localRank++;
@@ -1260,23 +1276,23 @@ testResult_t run() {
    color = proc & strtoul(splitMaskEnv, NULL, 16);
  } else if (splitMaskEnv = getenv("NCCL_TESTS_SPLIT")) {
    if (
-      (strncasecmp(splitMaskEnv, "AND", strlen("AND")) == 0 && parseInt(splitMaskEnv + strlen("AND"), &color)) ||
-      (strncasecmp(splitMaskEnv, "&", strlen("&")) == 0 && parseInt(splitMaskEnv + strlen("&"), &color))
+      (ncclTestStrncasecmp(splitMaskEnv, "AND", strlen("AND")) == 0 && parseInt(splitMaskEnv + strlen("AND"), &color)) ||
+      (ncclTestStrncasecmp(splitMaskEnv, "&", strlen("&")) == 0 && parseInt(splitMaskEnv + strlen("&"), &color))
    )
        color = proc & color;
    if (
-      (strncasecmp(splitMaskEnv, "OR", strlen("OR")) == 0 && parseInt(splitMaskEnv + strlen("OR"), &color)) ||
-      (strncasecmp(splitMaskEnv, "|", strlen("|")) == 0 && parseInt(splitMaskEnv + strlen("|"), &color))
+      (ncclTestStrncasecmp(splitMaskEnv, "OR", strlen("OR")) == 0 && parseInt(splitMaskEnv + strlen("OR"), &color)) ||
+      (ncclTestStrncasecmp(splitMaskEnv, "|", strlen("|")) == 0 && parseInt(splitMaskEnv + strlen("|"), &color))
    )
        color = proc | color;
    if (
-      (strncasecmp(splitMaskEnv, "MOD", strlen("MOD")) == 0 && parseInt(splitMaskEnv + strlen("MOD"), &color)) ||
-      (strncasecmp(splitMaskEnv, "%", strlen("%")) == 0 && parseInt(splitMaskEnv + strlen("%"), &color))
+      (ncclTestStrncasecmp(splitMaskEnv, "MOD", strlen("MOD")) == 0 && parseInt(splitMaskEnv + strlen("MOD"), &color)) ||
+      (ncclTestStrncasecmp(splitMaskEnv, "%", strlen("%")) == 0 && parseInt(splitMaskEnv + strlen("%"), &color))
    )
        color = proc % color;
    if (
-      (strncasecmp(splitMaskEnv, "DIV", strlen("DIV")) == 0 && parseInt(splitMaskEnv + strlen("DIV"), &color)) ||
-      (strncasecmp(splitMaskEnv, "/", strlen("/")) == 0 && parseInt(splitMaskEnv + strlen("/"), &color))
+      (ncclTestStrncasecmp(splitMaskEnv, "DIV", strlen("DIV")) == 0 && parseInt(splitMaskEnv + strlen("DIV"), &color)) ||
+      (ncclTestStrncasecmp(splitMaskEnv, "/", strlen("/")) == 0 && parseInt(splitMaskEnv + strlen("/"), &color))
    )
        color = proc / color;
  }
@@ -1291,7 +1307,7 @@ testResult_t run() {
  jsonIdentifyWriter(is_main_thread);

  size_t maxMem = ~0;
-  testResult_t report_result = writeDeviceReport(&maxMem, localRank, proc, totalProcs, color, hostname, program_invocation_short_name);
+  testResult_t report_result = writeDeviceReport(&maxMem, localRank, proc, totalProcs, color, hostname, programName);
  if(report_result != testSuccess) {
    return report_result;
  }
@@ -1315,11 +1331,11 @@ testResult_t run() {
  MPI_Bcast(&ncclId, sizeof(ncclId), MPI_BYTE, 0, mpi_comm);
  MPI_Barrier(MPI_COMM_WORLD); // Ensure Bcast is complete for HCOLL
 #endif
-  int gpus[nGpus*nThreads];
-  cudaStream_t streams[nGpus*nThreads];
-  void* sendbuffs[nGpus*nThreads];
-  void* recvbuffs[nGpus*nThreads];
-  void* expected[nGpus*nThreads];
+  std::vector<int> gpus(nGpus*nThreads);
+  std::vector<cudaStream_t> streams(nGpus*nThreads);
+  std::vector<void*> sendbuffs(nGpus*nThreads);
+  std::vector<void*> recvbuffs(nGpus*nThreads);
+  std::vector<void*> expected(nGpus*nThreads);
  size_t sendBytes, recvBytes;

  ncclTestEngine.getBuffSize(&sendBytes, &recvBytes, (size_t)maxBytes, (size_t)ncclProcs*nGpus*nThreads);
@@ -1333,7 +1349,7 @@ testResult_t run() {
      streams[i] = NULL;
    }
    else {
-      CUDACHECK(cudaStreamCreateWithFlags(streams+i, cudaStreamNonBlocking));
+      CUDACHECK(cudaStreamCreateWithFlags(streams.data()+i, cudaStreamNonBlocking));
    }
    int archMajor, archMinor;
    CUDACHECK(cudaDeviceGetAttribute(&archMajor, cudaDevAttrComputeCapabilityMajor, gpus[i]));
@@ -1363,18 +1379,14 @@ testResult_t run() {
  //if parallel init is not selected, use main thread to initialize NCCL
  ncclComm_t* comms = (ncclComm_t*)malloc(sizeof(ncclComm_t)*nThreads*nGpus);
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,19,0)
-  void* sendRegHandles[nThreads*nGpus];
-  void* recvRegHandles[nThreads*nGpus];
-  memset(sendRegHandles, 0, sizeof(sendRegHandles));
-  memset(recvRegHandles, 0, sizeof(recvRegHandles));
+  std::vector<void*> sendRegHandles(nThreads*nGpus);
+  std::vector<void*> recvRegHandles(nThreads*nGpus);
 #endif
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
-  ncclDevComm devComms[nThreads*nGpus];
+  std::vector<ncclDevComm> devComms(nThreads*nGpus);
 #endif
-  int64_t initGpuMem[nThreads];
-  int64_t bufferMemory[nThreads];
-  memset(initGpuMem, 0, sizeof(initGpuMem));
-  memset(bufferMemory, 0, sizeof(bufferMemory));
+  std::vector<int64_t> initGpuMem(nThreads);
+  std::vector<int64_t> bufferMemory(nThreads);
  if (!parallel_init) {
    // Capture the memory used by the GPUs before initializing the NCCL communicators
    int64_t* initFreeGpuMem = (int64_t*)calloc(nGpus*3, sizeof(int64_t));
@@ -1383,7 +1395,7 @@ testResult_t run() {
      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));
+    TESTCHECK(initComms(comms, nGpus*nThreads, ncclProc*nThreads*nGpus, ncclProcs*nThreads*nGpus, gpus.data(), ncclId));

     // Capture the memory used by the GPUs after initializing the NCCL communicators
     for (int g = 0; g < nGpus; ++g) {
@@ -1399,7 +1411,7 @@ testResult_t run() {
     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));
+       TESTCHECK(AllocateBuffs(sendbuffs.data()+i, sendBytes, recvbuffs.data()+i, recvBytes, expected.data()+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));
@@ -1456,7 +1468,7 @@ testResult_t run() {

       NCCLCHECK(ncclGroupStart());
       for (int i = 0; i < nGpus * nThreads; i++) {
-         NCCLCHECK(ncclDevCommCreate(comms[i], &reqs, devComms+i));
+         NCCLCHECK(ncclDevCommCreate(comms[i], &reqs, devComms.data()+i));
       }
       NCCLCHECK(ncclGroupEnd());
     }
@@ -1474,10 +1486,10 @@ testResult_t run() {
    free(initFreeGpuMem);
  }

-  int errors[nThreads];
-  double bw[nThreads];
-  int64_t devMemUsed[nThreads];
-  int bw_count[nThreads];
+  std::vector<int> errors(nThreads);
+  std::vector<double> bw(nThreads);
+  std::vector<int64_t> devMemUsed(nThreads);
+  std::vector<int> bw_count(nThreads);
  for (int t=0; t<nThreads; t++) {
    bw[t] = 0.0;
    errors[t] = bw_count[t] = 0;
@@ -1488,8 +1500,7 @@ testResult_t run() {

  writeResultHeader(report_cputime, report_timestamps);

-  struct testThread threads[nThreads];
-  memset(threads, 0, sizeof(struct testThread)*nThreads);
+  std::vector<struct testThread> threads(nThreads);

  for (int t=nThreads-1; t>=0; t--) {
    threads[t].args.minbytes=minBytes;
@@ -1504,40 +1515,40 @@ testResult_t run() {
    threads[t].args.nThreads=nThreads;
    threads[t].args.thread=t;
    threads[t].args.nGpus=nGpus;
-    threads[t].args.gpus=gpus+t*nGpus;
-    threads[t].args.sendbuffs = sendbuffs+t*nGpus;
-    threads[t].args.recvbuffs = recvbuffs+t*nGpus;
-    threads[t].args.expected = expected+t*nGpus;
+    threads[t].args.gpus=gpus.data()+t*nGpus;
+    threads[t].args.sendbuffs = sendbuffs.data()+t*nGpus;
+    threads[t].args.recvbuffs = recvbuffs.data()+t*nGpus;
+    threads[t].args.expected = expected.data()+t*nGpus;
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
-    threads[t].args.devComms = devComms+t*nGpus;
+    threads[t].args.devComms = devComms.data()+t*nGpus;
 #endif
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,19,0)
-    threads[t].args.sendRegHandles = sendRegHandles+t*nGpus;
-    threads[t].args.recvRegHandles = recvRegHandles+t*nGpus;
+    threads[t].args.sendRegHandles = sendRegHandles.data()+t*nGpus;
+    threads[t].args.recvRegHandles = recvRegHandles.data()+t*nGpus;
 #endif
    threads[t].args.ncclId = ncclId;
    threads[t].args.comms=comms+t*nGpus;
-    threads[t].args.streams=streams+t*nGpus;
+    threads[t].args.streams=streams.data()+t*nGpus;

-    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.errors=errors.data()+t;
+    threads[t].args.bw=bw.data()+t;
+    threads[t].args.bw_count=bw_count.data()+t;
+    threads[t].args.initGpuMem = initGpuMem.data() + t;
+    threads[t].args.bufferMemory = bufferMemory.data() + t;
+    threads[t].args.devMemUsed = devMemUsed.data() + t;

    threads[t].args.reportErrors = datacheck;

    threads[t].func = parallel_init ? threadInit : threadRunTests;
    if (t)
-      TESTCHECK(threadLaunch(threads+t));
+      TESTCHECK(threadLaunch(threads.data()+t));
    else
      TESTCHECK(threads[t].func(&threads[t].args));
  }

  // Wait for other threads and accumulate stats and errors
  for (int t=nThreads-1; t>=0; t--) {
-    if (t) pthread_join(threads[t].thread, NULL);
+    if (t) threads[t].thread.join();
    TESTCHECK(threads[t].ret);
    if (t) {
      errors[0] += errors[t];
@@ -1572,7 +1583,7 @@ testResult_t run() {
 #endif
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,28,0)
      if (deviceImpl) {
-        NCCLCHECK(ncclDevCommDestroy(comms[i], devComms+i));
+        NCCLCHECK(ncclDevCommDestroy(comms[i], devComms.data()+i));
      }
 #endif
      NCCLCHECK(ncclCommDestroy(comms[i]));
@@ -1597,7 +1608,7 @@ testResult_t run() {
  const double check_avg_bw = envstr ? atof(envstr) : -1;
  bw[0] /= bw_count[0];

-  writeResultFooter(errors, bw, check_avg_bw, program_invocation_short_name);
+  writeResultFooter(errors.data(), bw.data(), check_avg_bw, programName);
  if (memory_report) {
    memInfo_t memInfos[3];
    memInfos[0] = { initGpuMem[0], "Initialization" };
@@ -15,12 +15,13 @@
 #include <stdio.h>
 #include <cstdint>
 #include <algorithm>
+#include <thread>
 #ifdef MPI_SUPPORT
 #include "mpi.h"
 #endif
-#include <pthread.h>
 #include "nccl1_compat.h"
 #include "timer.h"
+#include "os.h"

 // For nccl.h < 2.13 since we define a weak fallback
 extern "C" char const* ncclGetLastError(ncclComm_t comm);
@@ -83,7 +84,7 @@ typedef enum {
    char hostname[1024];                            \
    getHostName(hostname, 1024);                    \
    printf(" .. %s pid %d: Test failure %s:%d\n",   \
-         hostname, getpid(),                        \
+         hostname, ncclTestGetPid(),                \
        __FILE__,__LINE__);                         \
    return r;                                       \
  }                                                 \
@@ -175,7 +176,7 @@ struct threadArgs {

 typedef testResult_t (*threadFunc_t)(struct threadArgs* args);
 struct testThread {
-  pthread_t thread;
+  std::thread thread;
  threadFunc_t func;
  struct threadArgs args;
  testResult_t ret;
@@ -188,10 +189,8 @@ extern testResult_t InitDataReduce(void* data, const size_t count, const size_t
 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 testResult_t AllocateBuffs(void **sendbuff, size_t sendBytes, void **recvbuff, size_t recvBytes, void **expected, size_t nbytes);

-#include <unistd.h>
-
 static void getHostName(char* hostname, int maxlen) {
-  gethostname(hostname, maxlen);
+  ncclTestGetHostname(hostname, maxlen);
  for (int i=0; i< maxlen; i++) {
    if (hostname[i] == '\0') {
      return;
@@ -203,46 +202,16 @@ static void getHostName(char* hostname, int maxlen) {
  }
 }

-#include <stdint.h>
-
-static uint64_t getHash(const char* string, size_t n) {
-  // Based on DJB2a, result = result * 33 ^ char
-  uint64_t result = 5381;
-  for (size_t c = 0; c < n; c++) {
-    result = ((result << 5) + result) ^ string[c];
-  }
-  return result;
-}
-
 /* Generate a hash of the unique identifying string for this host
 * that will be unique for both bare-metal and container instances
 * Equivalent of a hash of;
 *
- * $(hostname)$(cat /proc/sys/kernel/random/boot_id)
+ * $(hostname)$(cat /proc/sys/kernel/random/boot_id)       [Linux]
+ * $(hostname)$(MachineGuid from registry)                 [Windows]
 *
 */
-#define HOSTID_FILE "/proc/sys/kernel/random/boot_id"
 static uint64_t getHostHash(const char* hostname) {
-  char hostHash[1024];
-
-  // Fall back is the hostname if something fails
-  (void) strncpy(hostHash, hostname, sizeof(hostHash));
-  int offset = strlen(hostHash);
-
-  FILE *file = fopen(HOSTID_FILE, "r");
-  if (file != NULL) {
-    char *p;
-    if (fscanf(file, "%ms", &p) == 1) {
-        strncpy(hostHash+offset, p, sizeof(hostHash)-offset-1);
-        free(p);
-    }
-  }
-  fclose(file);
-
-  // Make sure the string is terminated
-  hostHash[sizeof(hostHash)-1]='\0';
-
-  return getHash(hostHash, strlen(hostHash));
+  return ncclTestGetHostHash(hostname);
 }

 #define HAVE_BF16 0
@@ -69,6 +69,16 @@ endif
 NVCUFLAGS  := -ccbin $(CXX) $(NVCC_GENCODE) $(CXXSTD) --extended-lambda
 CXXFLAGS   := $(CXXSTD)

+ifneq ($(OS),Windows_NT)
+NCCL_OS_LINUX := 1
+CXXFLAGS += -DNCCL_OS_LINUX
+NVCUFLAGS += -DNCCL_OS_LINUX
+else
+NCCL_OS_WINDOWS := 1
+CXXFLAGS += -DNCCL_OS_WINDOWS -DWIN32_LEAN_AND_MEAN -DNOMINMAX
+NVCUFLAGS += -DNCCL_OS_WINDOWS -DWIN32_LEAN_AND_MEAN -DNOMINMAX
+endif
+
 LDFLAGS    := -L${CUDA_LIB} -lcudart -lrt
 NVLDFLAGS  := -L${CUDA_LIB} -l${CUDARTLIB} -lrt

@@ -120,9 +120,7 @@ testResult_t GatherRunTest(struct threadArgs* args, int root, ncclDataType_t typ
  return testSuccess;
 }

-struct testEngine gatherEngine = {
-  .getBuffSize = GatherGetBuffSize,
-  .runTest = GatherRunTest
+NCCL_WEAK struct testEngine ncclTestEngine = {
+  /* .getBuffSize = */ GatherGetBuffSize,
+  /* .runTest = */ GatherRunTest
 };
-
-#pragma weak ncclTestEngine=gatherEngine
@@ -114,9 +114,7 @@ testResult_t HyperCubeRunTest(struct threadArgs* args, int root, ncclDataType_t
  return testSuccess;
 }

-struct testEngine hyperCubeEngine = {
-  .getBuffSize = HyperCubeGetBuffSize,
-  .runTest = HyperCubeRunTest
+NCCL_WEAK struct testEngine ncclTestEngine = {
+  /* .getBuffSize = */ HyperCubeGetBuffSize,
+  /* .runTest = */ HyperCubeRunTest
 };
-
-#pragma weak ncclTestEngine=hyperCubeEngine
@@ -108,9 +108,7 @@ testResult_t ReduceRunTest(struct threadArgs* args, int root, ncclDataType_t typ
  return testSuccess;
 }

-struct testEngine reduceEngine = {
-  .getBuffSize = ReduceGetBuffSize,
-  .runTest = ReduceRunTest
+NCCL_WEAK struct testEngine ncclTestEngine = {
+  /* .getBuffSize = */ ReduceGetBuffSize,
+  /* .runTest = */ ReduceRunTest
 };
-
-#pragma weak ncclTestEngine=reduceEngine
@@ -101,9 +101,7 @@ testResult_t ReduceScatterRunTest(struct threadArgs* args, int root, ncclDataTyp
  return testSuccess;
 }

-struct testEngine reduceScatterEngine = {
-  .getBuffSize = ReduceScatterGetBuffSize,
-  .runTest = ReduceScatterRunTest
+NCCL_WEAK struct testEngine ncclTestEngine = {
+  /* .getBuffSize = */ ReduceScatterGetBuffSize,
+  /* .runTest = */ ReduceScatterRunTest
 };
-
-#pragma weak ncclTestEngine=reduceScatterEngine
@@ -116,9 +116,7 @@ testResult_t ScatterRunTest(struct threadArgs* args, int root, ncclDataType_t ty
  return testSuccess;
 }

-struct testEngine scatterEngine = {
-  .getBuffSize = ScatterGetBuffSize,
-  .runTest = ScatterRunTest
+NCCL_WEAK struct testEngine ncclTestEngine = {
+  /* .getBuffSize = */ ScatterGetBuffSize,
+  /* .runTest = */ ScatterRunTest
 };
-
-#pragma weak ncclTestEngine=scatterEngine
@@ -123,12 +123,10 @@ testResult_t SendRecvRunTest(struct threadArgs* args, int root, ncclDataType_t t
  return testSuccess;
 }

-struct testEngine sendRecvEngine = {
-  .getBuffSize = SendRecvGetBuffSize,
-  .runTest = SendRecvRunTest,
+NCCL_WEAK struct testEngine ncclTestEngine = {
+  /* .getBuffSize = */ SendRecvGetBuffSize,
+  /* .runTest = */ SendRecvRunTest,
 #if NCCL_VERSION_CODE >= NCCL_VERSION(2,14,0)
-  .initCommConfig = SendRecvInitCommConfig,
+  /* .initCommConfig = */ SendRecvInitCommConfig,
 #endif
 };
-
-#pragma weak ncclTestEngine=sendRecvEngine
@@ -17,9 +17,11 @@

 #include "nccl.h"
 #include "util.h"
+#include "os.h"
 #include <assert.h>
 #include <errno.h>
 #include <string>
+#include <sstream>
 #include <iomanip>

 #define PRINT if (is_main_thread) printf
@@ -332,7 +334,7 @@ void jsonOutputInit(const char *in_path,
      return;
    }
    free(try_path);
-    if(asprintf(&try_path, "%s.%d", in_path, try_count++) == -1) {
+    if(ncclTestAsprintf(&try_path, "%s.%d", in_path, try_count++) == -1) {
      printf("# skipping json output; failed to probe destination\n");
      return;
    }
@@ -591,7 +593,7 @@ testResult_t writeDeviceReport(size_t *maxMem, int localRank, int proc, int tota
    CUDACHECK(cudaGetDeviceProperties(&prop, cudaDev));
    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);
+                      rank, color, ncclTestGetPid(), hostname, cudaDev, prop.pciDomainID, prop.pciBusID, prop.pciDeviceID, prop.name);
    }
    *maxMem = std::min(*maxMem, prop.totalGlobalMem);
  }
@@ -0,0 +1,23 @@
+#
+# Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved.
+#
+# See LICENSE.txt for license information
+#
+
+set(VERIFIABLE_LIB_SOURCES
+  verifiable.cu
+)
+
+add_library(verifiable STATIC ${VERIFIABLE_LIB_SOURCES})
+
+target_link_libraries(verifiable
+  PUBLIC
+    nccl_tests_options
+    nccl
+    CUDA::cudart
+)
+
+target_include_directories(verifiable
+  PUBLIC
+    ${CMAKE_CURRENT_SOURCE_DIR}
+)
@@ -41,7 +41,11 @@
 #include <cstdio>
 #include <cstdint>
 #include <cmath>
+#if defined(NCCL_OS_LINUX)
 #include <unistd.h>
+#elif defined(NCCL_OS_WINDOWS)
+#include <intrin.h>
+#endif

 using std::size_t;
 using std::int8_t;
@@ -495,23 +499,31 @@ __host__ __device__ uint64_t umul32hi(uint32_t a, uint32_t b) {
 __host__ __device__ uint64_t umul64hi(uint64_t a, uint64_t b) {
 #ifdef __CUDA_ARCH__
  return __umul64hi(a, b);
-#else
+#elif defined(NCCL_OS_LINUX)
  return uint64_t(__uint128_t(a)*__uint128_t(b) >> 64);
+#elif defined(NCCL_OS_WINDOWS)
+  return __umulh(a, b);
 #endif
 }

 __host__ __device__ int clz32(int x) {
 #ifdef __CUDA_ARCH__
  return __clz(x);
-#else
+#elif defined(NCCL_OS_LINUX)
  return x==0 ? 32 : __builtin_clz(x);
+#elif defined(NCCL_OS_WINDOWS)
+  unsigned long idx;
+  return _BitScanReverse(&idx, (unsigned long)x) ? 31 - (int)idx : 32;
 #endif
 }
 __host__ __device__ int clz64(long long x) {
 #ifdef __CUDA_ARCH__
  return __clzll(x);
-#else
+#elif defined(NCCL_OS_LINUX)
  return x==0 ? 64 : __builtin_clzll(x);
+#elif defined(NCCL_OS_WINDOWS)
+  unsigned long idx;
+  return _BitScanReverse64(&idx, (unsigned long long)x) ? 63 - (int)idx : 64;
 #endif
 }
 }
@@ -906,7 +918,7 @@ __host__ __device__ void genInput(
  // limit to two ranks contributing non-zero values. This way there is no ambiguity
  // of summation.
  int r = shuffleRank(rank_n, rank_me, rng);
-  uint64_t m = (rng*(r ? 0xbeef : 1)) & ((1ul<<FloatLayout<T>::mantissa_bits)-1);
+  uint64_t m = (rng*(r ? 0xbeef : 1)) & ((1ull<<FloatLayout<T>::mantissa_bits)-1);
  ans = r < 2 ? castTo<T>(1+m) : castTo<T>((uint64_t)0);
 }

@@ -916,8 +928,8 @@ __host__ __device__ void genOutput(
    std::false_type /*integral*/
  ) {
  shuffleRank(rank_n, -1, rng);
-  uint64_t m0 = (rng*(0 ? 0xbeef : 1)) & ((1ul<<FloatLayout<T>::mantissa_bits)-1);
-  uint64_t m1 = (rng*(1 ? 0xbeef : 1)) & ((1ul<<FloatLayout<T>::mantissa_bits)-1);
+  uint64_t m0 = (rng*(0 ? 0xbeef : 1)) & ((1ull<<FloatLayout<T>::mantissa_bits)-1);
+  uint64_t m1 = (rng*(1 ? 0xbeef : 1)) & ((1ull<<FloatLayout<T>::mantissa_bits)-1);
  if (rank_n == 1) {
    ans = castTo<T>(1+m0);
  } else {