Build and run the suite on Blackwell (sm_103) under CUDA 13.2

.github/workflows/test-build.yml

@@ -38,3 +38,41 @@ jobs:
         run: |
               echo "Built `ls bin/12.8/release | wc` Apps:"
               ls bin/12.8/release
+
+  # The accel-sim org publishes no CUDA 13 image, so this job starts from the
+  # stock NVIDIA devel image and installs what the minimal image would have
+  # provided. 13.2 is the toolkit the Blackwell port was developed against.
+  test-13-2:
+    runs-on: ubuntu-latest
+    container:
+     image: nvidia/cuda:13.2.1-devel-ubuntu24.04
+    env:
+      CUDA_INSTALL_PATH: /usr/local/cuda
+      BOOST_ROOT: /usr
+      DEBIAN_FRONTEND: noninteractive
+
+    steps:
+      - name: Install build dependencies
+        run: |
+             apt-get update
+             apt-get install -y --no-install-recommends \
+               build-essential cmake git ca-certificates libboost-dev zlib1g-dev \
+               python3
+
+      - name: Checkout repository with submodules
+        uses: actions/checkout@v4
+        with:
+          submodules: false
+
+      - name: Build Apps
+        run: |
+             git config --global --add safe.directory /__w/gpu-app-collection/gpu-app-collection
+             git submodule update --init -- src/cuda/cuda-samples
+             git submodule update --init -- src/cuda/cutlass-bench
+             /bin/bash test-build.sh ci
+
+      - name: Print Successful Apps
+        if: always()
+        run: |
+              echo "Built `ls bin/13.2/release | wc` Apps:"
+              ls bin/13.2/release

.gitignore

@@ -13,6 +13,16 @@ src/cuda/rodinia/3.1/cuda/nn/nn
 src/cuda/rodinia/3.1/cuda/particlefilter/particlefilter_float
 src/cuda/rodinia/3.1/cuda/particlefilter/particlefilter_naive
 src/cuda/rodinia/3.1/cuda/pathfinder/pathfinder
+src/cuda/rodinia/3.1/cuda/huffman/pavle
+src/cuda/rodinia/3.1/cuda/b+tree/b+tree.out
+src/cuda/rodinia/3.1/cuda/dwt2d/dwt2d
+src/cuda/rodinia/3.1/cuda/heartwall/heartwall
+src/cuda/rodinia/3.1/cuda/hybridsort/hybridsort
+src/cuda/rodinia/3.1/cuda/myocyte/myocyte.out
+src/cuda/rodinia/3.1/cuda/hotspot3D/3D
+src/cuda/rodinia/3.1/cuda/gaussian/gaussian
+src/cuda/rodinia/3.1/cuda/srad/srad_v1/srad
+src/cuda/rodinia/3.1/cuda/srad/srad_v2/srad
 4.2
 .venv/
 __pycache__/
@@ -24,4 +34,4 @@ tmp/
 src/cuda/GPU_Microbenchmark/ubench/**/*
 !src/cuda/GPU_Microbenchmark/ubench/**/*/
 !src/cuda/GPU_Microbenchmark/ubench/**/*.*
-!src/cuda/GPU_Microbenchmark/ubench/**/Makefile
+!src/cuda/GPU_Microbenchmark/ubench/**/Makefile

src/.gitignore

@@ -23,6 +23,8 @@ cuda/ispass-2009/WP/wsm5_gpu.f.cu
 cuda/ispass-2009/WP/wsm5_gpu.f.cu.cpp.ii
 cuda/ispass-2009/WP-*
 4.2/
+# built by setup_environment from the vendored sources alongside it
+cuda/common/sdk-compat/lib/
 cuda/lonestargpu-2.0/bin
 cuda/rodinia/3.1/cuda/huffman/pavle
 gpucomputingsdk_4.2.9_linux.run
@@ -70,3 +72,9 @@ cuda/GPU_Microbenchmark/shared_bw/shared_bw
 cuda/GPU_Microbenchmark/shared_lat/shared_lat
 cuda/rodinia/3.1/cuda/leukocyte/meschach_lib/confdefs.h
 cuda/rodinia/3.1/cuda/leukocyte/meschach_lib/config.log
+
+# generated by running the benchmarks rather than building them
+cuda/parboil/benchmarks/**/run/
+__pycache__/
+cuda/rodinia/3.1/cuda/kmeans/kmeans
+cuda/rodinia/3.1/cuda/mummergpu/src/obj/

src/cuda/DeepBench/code/nvidia/Makefile

@@ -27,17 +27,17 @@ all: gemm conv rnn
 
 gemm:
 	$(MKDIR) $(BIN_DIR) 
-	$(CUDA_PATH)/bin/$(NVCC) gemm_bench.cu -DUSE_TENSOR_CORES=0 -DPAD_KERNELS=$(PAD_KERNELS) -o $(BIN_DIR)/gemm_bench -I $(KERNELS_DIR) -I $(CUDA_PATH)/include -L $(BLAS_PATH) -l$(BLAS_LIBRARY) -L $(CUDA_LIB64) -lcurand $(CUDA_CPPFLAGS) -std=c++14
+	$(CUDA_PATH)/bin/$(NVCC) gemm_bench.cu -DUSE_TENSOR_CORES=0 -DPAD_KERNELS=$(PAD_KERNELS) -o $(BIN_DIR)/gemm_bench -I $(KERNELS_DIR) -I $(CUDA_PATH)/include -L $(BLAS_PATH) -l$(BLAS_LIBRARY) -L $(CUDA_LIB64) -lcurand $(CUDA_CPPFLAGS) -std=c++17
 	$(CUDA_PATH)/bin/$(NVCC) gemm_bench.cu -DUSE_TENSOR_CORES=1 -DPAD_KERNELS=$(PAD_KERNELS) -o $(BIN_DIR)/gemm_bench-tencore -I $(KERNELS_DIR) -I $(CUDA_PATH)/include -L $(BLAS_PATH) -l$(BLAS_LIBRARY) -L $(CUDA_LIB64) -lcurand $(CUDA_CPPFLAGS) 
 
 conv:
 	$(MKDIR) $(BIN_DIR)
-	$(CUDA_PATH)/bin/$(NVCC) conv_bench.cu -DUSE_TENSOR_CORES=0 -DPAD_KERNELS=$(PAD_KERNELS) -o $(BIN_DIR)/conv_bench -I $(KERNELS_DIR) -I $(CUDA_PATH)/include -I $(CUDNN_PATH)/include/ -L $(CUDNN_PATH)/lib64/ -L $(CUDA_LIB64) -lcurand -lcudnn $(CUDA_CPPFLAGS) -std=c++14
+	$(CUDA_PATH)/bin/$(NVCC) conv_bench.cu -DUSE_TENSOR_CORES=0 -DPAD_KERNELS=$(PAD_KERNELS) -o $(BIN_DIR)/conv_bench -I $(KERNELS_DIR) -I $(CUDA_PATH)/include -I $(CUDNN_PATH)/include/ -L $(CUDNN_PATH)/lib64/ -L $(CUDA_LIB64) -lcurand -lcudnn $(CUDA_CPPFLAGS) -std=c++17
 	$(CUDA_PATH)/bin/$(NVCC) conv_bench.cu -DUSE_TENSOR_CORES=1 -DPAD_KERNELS=$(PAD_KERNELS) -o $(BIN_DIR)/conv_bench-tencore -I $(KERNELS_DIR) -I $(CUDA_PATH)/include -I $(CUDNN_PATH)/include/ -L $(CUDNN_PATH)/lib64/ -L $(CUDA_LIB64) -lcurand -lcudnn $(CUDA_CPPFLAGS) 
 
 rnn:
 	$(MKDIR) $(BIN_DIR)
-	$(CUDA_PATH)/bin/$(NVCC) rnn_bench.cu -DUSE_TENSOR_CORES=0 -o $(BIN_DIR)/rnn_bench -I $(KERNELS_DIR) -I $(CUDA_PATH)/include -I $(CUDNN_PATH)/include/ -L $(CUDNN_PATH)/lib64/ -L $(CUDA_LIB64) -lcurand -lcudnn $(CUDA_CPPFLAGS) -std=c++14
+	$(CUDA_PATH)/bin/$(NVCC) rnn_bench.cu -DUSE_TENSOR_CORES=0 -o $(BIN_DIR)/rnn_bench -I $(KERNELS_DIR) -I $(CUDA_PATH)/include -I $(CUDNN_PATH)/include/ -L $(CUDNN_PATH)/lib64/ -L $(CUDA_LIB64) -lcurand -lcudnn $(CUDA_CPPFLAGS) -std=c++17
 	$(CUDA_PATH)/bin/$(NVCC) rnn_bench.cu -DUSE_TENSOR_CORES=1 -o $(BIN_DIR)/rnn_bench-tencore -I $(KERNELS_DIR) -I $(CUDA_PATH)/include -I $(CUDNN_PATH)/include/ -L $(CUDNN_PATH)/lib64/ -L $(CUDA_LIB64) -lcurand -lcudnn $(CUDA_CPPFLAGS) 
 
 all_reduce: nccl_single nccl_mpi

src/cuda/DeepBench/code/nvidia/tensor.h

@@ -8,6 +8,7 @@
 
 #include <thrust/device_ptr.h>
 #include <thrust/fill.h>
+#include <thrust/copy.h>
 
 template <typename T>
 class Tensor {

src/cuda/GPU_Microbenchmark/common/arch_filter.mk

@@ -0,0 +1,11 @@
+# Drop any architecture the installed toolkit will not accept, instead of
+# failing the build. These ubenches pin an explicit ARCH list (TMA and GMMA
+# need the accelerated targets), and those lists drift: CUDA 13 removed
+# sm_101, for example. The accelerated "a" variants are matched on their
+# base number, so sm_90a survives as long as compute_90 is supported.
+#
+# Include this after setting ARCH and before deriving NVCC_FLAGS from it.
+NVCC_SUPPORTED_ARCH := $(shell nvcc --list-gpu-arch 2>/dev/null | sed -e 's/compute_//')
+ifneq ($(strip $(NVCC_SUPPORTED_ARCH)),)
+ARCH := $(foreach a,$(ARCH),$(if $(filter $(patsubst %a,%,$(patsubst sm_%,%,$(a))),$(NVCC_SUPPORTED_ARCH)),$(a)))
+endif

src/cuda/GPU_Microbenchmark/common/common.mk

@@ -1,7 +1,16 @@
 BASE_DIR := $(shell pwd)
 BIN_DIR := $(BASE_DIR)/../../../bin/
 
+# Some ubenches need a minimum ISA (cp.async is sm_80+, wgmma sm_90+). They set
+# MIN_ARCH before including this file and the -gencode flags for anything older
+# are dropped, so one such benchmark cannot fail a build that still targets the
+# older architectures for everything else.
+MIN_ARCH ?=
+ifeq ($(strip $(MIN_ARCH)),)
 CUOPTS =  $(CUDA_CPPFLAGS)
+else
+CUOPTS = $(shell echo '$(CUDA_CPPFLAGS)' | tr ' ' '\n' | awk -v m=$(MIN_ARCH) '/compute_/ { a=$$0; sub(/.*compute_/,"",a); sub(/[^0-9].*/,"",a); if (a+0 < m+0) next } { printf "%s ", $$0 }')
+endif
 
 CC := nvcc
 

src/cuda/GPU_Microbenchmark/hw_def/common/gpuConfig.h

@@ -282,6 +282,9 @@ inline unsigned initializeDeviceProp(unsigned deviceID, int argc, char *argv[])
 
         int clockRateKHz;
         cudaDeviceGetAttribute(&clockRateKHz, cudaDevAttrClockRate, deviceID);
+        // cudaDeviceProp::memoryClockRate went the same way as ::clockRate in CUDA 13
+        int memClockRateKHz;
+        cudaDeviceGetAttribute(&memClockRateKHz, cudaDevAttrMemoryClockRate, deviceID);
 
         // core stats
         config.SM_NUMBER = deviceProp.multiProcessorCount;
@@ -310,7 +313,7 @@ inline unsigned initializeDeviceProp(unsigned deviceID, int argc, char *argv[])
 
         // memory
         config.MEM_SIZE = deviceProp.totalGlobalMem;
-        config.MEM_CLK_FREQUENCY = deviceProp.memoryClockRate * 1e-3f;
+        config.MEM_CLK_FREQUENCY = memClockRateKHz * 1e-3f;
         config.MEM_BITWIDTH = deviceProp.memoryBusWidth;
         config.CLK_FREQUENCY = clockRateKHz * 1e-3f;
 

src/cuda/GPU_Microbenchmark/ubench/core/MaxFlops_gmma/Makefile

@@ -9,6 +9,7 @@ INCLUDE += -I$(GPUAPPS_ROOT)/src/cuda/cutlass-bench/include -I./
 
 # GMMA is only supported in sm_90a
 ARCH?=sm_90a
+include ../../../common/arch_filter.mk
 # Unset the CUDA_CPPFLAGS which is set based on CUDA version
 CUDA_CPPFLAGS=
 # Generate code for both sm_XXX and compute_XXX (SASS and PTX)

src/cuda/GPU_Microbenchmark/ubench/core/config_tensor/Makefile

@@ -9,4 +9,10 @@ SRC = config_tensor.cu
 
 EXE = config_tensor
 
+# <mma.h> only declares nvcuda::wmma for __CUDA_ARCH__ >= 700, so a gencode
+# list that still reaches back to Maxwell/Pascal fails here with "wmma is not
+# a class or namespace name". The GENCODE_SMxx assignments above predate
+# CUDA_CPPFLAGS and never applied to it.
+MIN_ARCH = 70
+
 include ../../../common/common.mk

src/cuda/GPU_Microbenchmark/ubench/core/lat_gmma/Makefile

@@ -9,6 +9,7 @@ INCLUDE += -I$(GPUAPPS_ROOT)/src/cuda/cutlass-bench/include -I./
 
 # GMMA is only supported in sm_90a
 ARCH?=sm_90a
+include ../../../common/arch_filter.mk
 # Unset the CUDA_CPPFLAGS which is set based on CUDA version
 CUDA_CPPFLAGS=
 # Generate code for both sm_90a and compute_90a (SASS and PTX)

src/cuda/GPU_Microbenchmark/ubench/core/tensor_bw_half/Makefile

@@ -9,4 +9,10 @@ SRC = tensor_bw_half.cu
 
 EXE = tensor_bw_half
 
+# <mma.h> only declares nvcuda::wmma for __CUDA_ARCH__ >= 700, so a gencode
+# list that still reaches back to Maxwell/Pascal fails here with "wmma is not
+# a class or namespace name". The GENCODE_SMxx assignments above predate
+# CUDA_CPPFLAGS and never applied to it.
+MIN_ARCH = 70
+
 include ../../../common/common.mk

src/cuda/GPU_Microbenchmark/ubench/core/tensor_lat_half/Makefile

@@ -9,4 +9,10 @@ SRC = tensor_lat_half.cu
 
 EXE = tensor_lat_half
 
+# <mma.h> only declares nvcuda::wmma for __CUDA_ARCH__ >= 700, so a gencode
+# list that still reaches back to Maxwell/Pascal fails here with "wmma is not
+# a class or namespace name". The GENCODE_SMxx assignments above predate
+# CUDA_CPPFLAGS and never applied to it.
+MIN_ARCH = 70
+
 include ../../../common/common.mk

src/cuda/GPU_Microbenchmark/ubench/mem/mem_config/mem_config.cu

@@ -12,7 +12,7 @@ int main(int argc, char *argv[])
   snprintf(msg, sizeof(msg), "Global memory size = %.0f GB\n",
            static_cast<float>(deviceProp.totalGlobalMem / 1073741824.0f));
   std::cout << msg;
-  std::cout << "Memory Clock rate = " << deviceProp.memoryClockRate * 1e-3f
+  std::cout << "Memory Clock rate = " << config.MEM_CLK_FREQUENCY
             << " Mhz\n";
   std::cout << "Memory Bus Width = " << deviceProp.memoryBusWidth << " bit\n";
   std::cout << "Memory type = " << dram_model_str[DRAM_MODEL] << "\n";
@@ -38,7 +38,7 @@ int main(int argc, char *argv[])
               << dram_model_freq_ratio[DRAM_MODEL] << std::endl;
 
     // timing
-    float device_freq_MHZ = (deviceProp.memoryClockRate * 1e-3f * 2) /
+    float device_freq_MHZ = (config.MEM_CLK_FREQUENCY * 2) /
                             dram_model_freq_ratio[DRAM_MODEL];
     if (DRAM_MODEL == dram_model::HBM)
     {

src/cuda/GPU_Microbenchmark/ubench/shd/shared_cp_async/Makefile

@@ -4,4 +4,7 @@ EXE = cp_async
 
 NVCC_FLGAS = -Xptxas -dlcm=cg -Xptxas -dscm=wt
 
+# ldgsts.cu uses cp.async, which ptxas rejects below sm_80
+MIN_ARCH = 80
+
 include ../../../common/common.mk

src/cuda/GPU_Microbenchmark/ubench/shd/shared_cp_async/ldgsts.cu

@@ -93,6 +93,10 @@ int main(int argc, char** argv) {
     // ==============================
     cudaDeviceProp prop;
     cudaGetDeviceProperties(&prop, 0);
+    // clockRate/memoryClockRate were removed from cudaDeviceProp in CUDA 13
+    int smClockKHz = 0, memClockKHz = 0;
+    cudaDeviceGetAttribute(&smClockKHz, cudaDevAttrClockRate, 0);
+    cudaDeviceGetAttribute(&memClockKHz, cudaDevAttrMemoryClockRate, 0);
 
     size_t l2_size_bytes = prop.l2CacheSize;
     if (l2_size_bytes == 0) {
@@ -175,7 +179,7 @@ for (size_t b = 0; b < num_blocks; b++) {
 double avg_cycles = static_cast<double>(total_cycles) / num_blocks;
 
 // GPU frequency (kHz → Hz)
-double gpu_clock_hz = static_cast<double>(prop.clockRate) * 1000.0;
+double gpu_clock_hz = static_cast<double>(smClockKHz) * 1000.0;
 
 // Time in seconds
 double time_sec = avg_cycles / gpu_clock_hz;
@@ -192,8 +196,8 @@ double bw_gbs   = bytes_moved / time_sec / 1e9;
 double bytesclk = bytes_moved / avg_cycles;   // Bytes per GPU cycle
 
 std::cout << "---------------------------------\n";
-std::cout << "SM Clock               = " << prop.clockRate / 1000.0 << " MHz\n";
-std::cout << "Memory Clock           = " << prop.memoryClockRate / 1000.0 << " MHz\n";
+std::cout << "SM Clock               = " << smClockKHz / 1000.0 << " MHz\n";
+std::cout << "Memory Clock           = " << memClockKHz / 1000.0 << " MHz\n";
 std::cout << "Avg cycles (per block) = " << avg_cycles << "\n";
 std::cout << "Time (s)               = " << time_sec << "\n";
 std::cout << "Bytes moved            = " << bytes_moved / (1024.0*1024*1024) << " GB\n";

src/cuda/GPU_Microbenchmark/ubench/system/deviceQuery/Makefile

@@ -2,7 +2,10 @@
 SRC = deviceQuery.cpp
 
 EXE = deviceQuery
-INCLUDE :=  -I../../../../cuda-samples/Common/
+# cuda-samples is cloned by a later target in the top-level Makefile, so it is
+# not there yet when this builds. The only thing needed from it here is
+# _ConvertSMVer2Cores, which the vendored sdk-compat headers also provide.
+INCLUDE :=  -I../../../../cuda-samples/Common/ -I$(GPUAPPS_SDK_INC)
 NVCC_FLGAS =
 
 include ../../../common/common.mk

src/cuda/GPU_Microbenchmark/ubench/system/deviceQuery/deviceQuery.cpp

@@ -32,6 +32,10 @@ int main(int argc, char **argv) {
     cudaSetDevice(dev);
     cudaDeviceProp deviceProp;
     cudaGetDeviceProperties(&deviceProp, dev);
+    // clockRate/memoryClockRate were removed from cudaDeviceProp in CUDA 13
+    int smClockKHz = 0, memClockKHz = 0;
+    cudaDeviceGetAttribute(&smClockKHz, cudaDevAttrClockRate, dev);
+    cudaDeviceGetAttribute(&memClockKHz, cudaDevAttrMemoryClockRate, dev);
 
     // device
     printf("  Device : \"%s\"\n\n", deviceProp.name);
@@ -40,7 +44,7 @@ int main(int argc, char **argv) {
 
     // core
     printf("  GPU Max Clock rate                             : %.0f MHz \n",
-           deviceProp.clockRate * 1e-3f);
+           smClockKHz * 1e-3f);
     printf("  Multiprocessors Count                       : %d\n",
            deviceProp.multiProcessorCount);
     printf("  Maximum number of threads per multiprocessor: %d\n",
@@ -81,7 +85,7 @@ int main(int argc, char **argv) {
              static_cast<float>(deviceProp.totalGlobalMem / 1073741824.0f));
     printf("%s", msg);
     printf("  Memory Clock rate                           : %.0f Mhz\n",
-           deviceProp.memoryClockRate * 1e-3f);
+           memClockKHz * 1e-3f);
     printf("  Memory Bus Width                            : %d bit\n",
            deviceProp.memoryBusWidth);
 

src/cuda/GPU_Microbenchmark/ubench/tma/mbarrier/Makefile

@@ -3,6 +3,8 @@ SRC = mbarrier.cu
 EXE = mbarrier
 
 ARCH?=sm_80 sm_90a sm_100a sm_101 sm_120 
+include ../../../common/arch_filter.mk
+
 # Unset the CUDA_CPPFLAGS which is set based on CUDA version
 CUDA_CPPFLAGS=
 

src/cuda/GPU_Microbenchmark/ubench/tma/tma_bulk/Makefile

@@ -4,6 +4,7 @@ EXE = tma_bulk
 
 # TMA is supported on SM_90a and above
 ARCH?=sm_90a sm_100a sm_101 sm_120 
+include ../../../common/arch_filter.mk
 # Unset the CUDA_CPPFLAGS which is set based on CUDA version
 # but TMA is only supported on SM_90a and above
 CUDA_CPPFLAGS=

src/cuda/GPU_Microbenchmark/ubench/tma/tma_bulk/tma_bulk.cu

@@ -1,6 +1,7 @@
 #include <math.h>
 #include <stdio.h>
 #include <stdlib.h>
+#include <getopt.h>   // getopt/optarg; no longer pulled in transitively
 #include <cuda/barrier>
 #include <cuda/ptx>
 

src/cuda/GPU_Microbenchmark/ubench/tma/tma_tensor/Makefile

@@ -4,6 +4,7 @@ EXE = tma_tensor
 
 # TMA is supported on SM_90a and above
 ARCH?=sm_90a sm_100a sm_101 sm_120 
+include ../../../common/arch_filter.mk
 # Unset the CUDA_CPPFLAGS which is set based on CUDA version
 # but TMA is only supported on SM_90a and above
 CUDA_CPPFLAGS=

src/cuda/GPU_Microbenchmark/ubench/tma/tma_tensor/tma_tensor.cu

@@ -5,6 +5,7 @@
 #include <math.h>
 #include <stdio.h>
 #include <stdlib.h>
+#include <getopt.h>   // getopt/optarg; no longer pulled in transitively
 #include <stdint.h>
 #include <unordered_map>
 using barrier = cuda::barrier<cuda::thread_scope_block>;