cuda: build with CUDA 13 and target Blackwell by melonakos · Pull Request #3715 · arrayfire/arrayfire

melonakos · 2026-09-10T15:42:44Z

Builds the CUDA backend with CUDA 13 and adds Blackwell (sm_100 through sm_121) as a target. Fixes #3688 and #3695.

This is #3698 rebased onto master with three commits on top; it supersedes #3698 and #3696 and includes #3706. The commits on top fix what was still broken: CMake's bundled arch-selection function was shadowing ours, the default arch list still handed nvcc 13 a removed target (nvcc fatal : Unsupported gpu architecture 'compute_70'), CUDA 13 moved its Windows DLLs to bin/x64 so standalone installs found nothing, and JIT targets now come from nvrtcGetSupportedArchs instead of a hand-kept table.

Verified with a full afcuda build on CUDA 13.0 with MSVC. There is no NVIDIA GPU on the build machine, so nothing has run on hardware, and Linux is untested. Draft until someone with a Blackwell or Turing card under CUDA 13 runs the tests.

WilliamTambellini · 2026-10-13T22:49:43Z

 #      NAME: Fermi Kepler Maxwell Kepler+Tegra Kepler+Tesla Maxwell+Tegra Pascal Volta Turing Ampere
 #      NUM: Any number. Only those pairs are currently accepted by NVCC though:
-#            2.0 2.1 3.0 3.2 3.5 3.7 5.0 5.2 5.3 6.0 6.2 7.0 7.2 7.5 8.0 8.6 9.0
+#            2.0 2.1 3.0 3.2 3.5 3.7 5.0 5.2 5.3 6.0 6.2 7.0 7.2 7.5 8.0 8.6 8.9 9.0 10.0 10.3 11.0 12.0 12.1


would nt 10.0 be enough to run on blackwells?

WilliamTambellini

simple comment

graehl · 2026-11-03T01:09:31Z

Linux test report: CUDA 13.3 on an sm_120 Blackwell card

With the four small patches below on top of this PR (cd211d5), the CUDA test suite passes 134/134 on blackwell cuda 13.3 (AWS g7e ubuntu), both without cuDNN and with cuDNN 9.27 (AF_WITH_CUDNN=ON). Without the patches, CUDA 13.3 fails to build. Once the build is fixed, 26/134 tests pass, because every runtime-compiled (JIT) kernel fails to load.

Environment. Ubuntu 26.04, GCC 15.2, CMake 3.27.6, CUDA toolkit 13.3 (Thrust 3.3.4), driver 595.91.07 (open kernel module), NVIDIA RTX PRO 6000 Blackwell Server Edition (compute capability 12.0). CUDA backend and the unified library only (AF_BUILD_CPU=OFF, AF_BUILD_OPENCL=OFF), built once without cuDNN and once with Ubuntu's libcudnn9-cuda-13 9.27.0.42. CUDA_architecture_build_targets=Auto detected 12.0 and compiled -gencode arch=compute_120,code=sm_120.

Build failures with CUDA 13.3

thrust/system/cuda/detail/par.h no longer exists. CCCL 3.0 and 3.1 still ship it, and 3.2 removed it, which probably explains why the CUDA 13.0 build in the PR description succeeded. thrust_utils.hpp and kernel/regions.hpp include it. The public <thrust/system/cuda/execution_policy.h> defines thrust::cuda::par and also exists in older Thrust versions.
Without that header, thrust::pair (used in ThrustArrayFirePolicy.hpp) and thrust::distance (used in set.cu) are no longer declared indirectly. The patch includes <thrust/pair.h> and <thrust/distance.h> explicitly.
This one comes from GCC 15, not CUDA 13: common/deterministicHash.cpp uses uint8_t and needs <cstdint>.

JIT kernels fail when the toolkit is newer than the driver

Every JIT kernel failed with this error:

CU Link Error CUDA_ERROR_UNSUPPORTED_PTX_VERSION(222): ptxas application ptx input, line 9; fatal   : Unsupported .version 9.3; current version is '9.2'

NVRTC 13.3 emits PTX ISA 9.3, and driver 595.91 assembles at most 9.2. CUDA minor-version compatibility supports a toolkit newer than the driver for compiled GPU code (SASS), but not for PTX the driver must compile. compileModule always passes PTX to cuLinkAddData. Patch 4 changes this: when NVRTC's JIT target equals the device's own compute capability, NVRTC compiles with --gpu-architecture=sm_XY, and the cubin from nvrtcGetCUBIN (available since CUDA 11.1) loads directly. When the JIT target is clamped below the device, meaning the GPU is newer than NVRTC knows, the existing PTX + cuLink path still runs. That fallback was not exercised here.

This bears on the review question about building only 10.0. SASS compiled for sm_100 does not load on sm_120, so a 12.0 card would depend on PTX embedded for compute_100 and compiled by the driver. With a toolkit newer than the driver, that hits the same PTX version error.

Test harness and minor notes

The SuiteSparse URLs now redirect to a 403, so configuring with tests failed. I used the fixture bundle and the skip-if-present change from ci: restore build checks and CPU summation accuracy #3735.
CMake found an fmt 12 install, and the build broke on the non-const formatter::format overloads for dim3 and common::Version. Hiding that install so the bundled spdlog/fmt is used avoided it. fmt 12 support would need those overloads to be const.
The info banner shows Driver: for, because nvDriverVersion misreads the open-kernel-module line in /proc/driver/nvidia/version (NVIDIA UNIX Open Kernel Module for x86_64 595.91.07 ...). This only affects the displayed text.
cuDNN 9 loads only through the unversioned libcudnn.so: cudnnModule tries libcudnn.so.8* names first, so a system with only the cuDNN 9 runtime package (no -dev symlink) would not find it. The trace also prints 9.27 as 92.7.0, decoding the version with the cuDNN 8 scheme.

Patches (4 commits, src/ only, clang-format clean)

diff --git a/src/backend/common/deterministicHash.cpp b/src/backend/common/deterministicHash.cpp
index 2280d4cbb..930fcdaec 100644
--- a/src/backend/common/deterministicHash.cpp
+++ b/src/backend/common/deterministicHash.cpp
@@ -10,6 +10,7 @@
 #include <common/deterministicHash.hpp>
 
 #include <nonstd/span.hpp>
+#include <cstdint>
 #include <numeric>
 #include <string>
 
diff --git a/src/backend/cuda/ThrustArrayFirePolicy.hpp b/src/backend/cuda/ThrustArrayFirePolicy.hpp
index 339d3ea08..eb8c49f6a 100644
--- a/src/backend/cuda/ThrustArrayFirePolicy.hpp
+++ b/src/backend/cuda/ThrustArrayFirePolicy.hpp
@@ -13,6 +13,7 @@
 #include <memory.hpp>
 #include <platform.hpp>
 #include <thrust/memory.h>
+#include <thrust/pair.h>
 #include <thrust/system/cuda/execution_policy.h>
 
 namespace arrayfire {
diff --git a/src/backend/cuda/compile_module.cpp b/src/backend/cuda/compile_module.cpp
index d7ee8182b..e983c0d22 100644
--- a/src/backend/cuda/compile_module.cpp
+++ b/src/backend/cuda/compile_module.cpp
@@ -262,8 +262,21 @@ Module compileModule(const string &moduleKey, span<const string> sources,
 
     int device       = getActiveDeviceId();
     auto computeFlag = getComputeCapability(device);
+    // When NVRTC supports the device's own architecture, let NVRTC emit SASS
+    // directly. PTX would have to be assembled by the driver, which rejects
+    // PTX newer than itself even though CUDA minor version compatibility
+    // allows a toolkit newer than the driver.
+#if CUDA_VERSION >= 11010
+    const cudaDeviceProp &prop = getDeviceProp(device);
+    const bool emitSASS =
+        computeFlag.first == prop.major && computeFlag.second == prop.minor;
+#else
+    const bool emitSASS = false;
+#endif
     array<char, 32> arch;
-    snprintf(arch.data(), arch.size(), "--gpu-architecture=compute_%d%d",
+    snprintf(arch.data(), arch.size(),
+             emitSASS ? "--gpu-architecture=sm_%d%d"
+                      : "--gpu-architecture=compute_%d%d",
              computeFlag.first, computeFlag.second);
     vector<const char *> compiler_options = {
         arch.data(),
@@ -296,35 +309,49 @@ Module compileModule(const string &moduleKey, span<const string> sources,
     NVRTC_COMPILE_CHECK(nvrtcCompileProgram(prog, compiler_options.size(),
                                             compiler_options.data()));
     auto compile_end = high_resolution_clock::now();
-    size_t ptx_size;
-    vector<char> ptx;
-    NVRTC_CHECK(nvrtcGetPTXSize(prog, &ptx_size));
-    ptx.resize(ptx_size);
-    NVRTC_CHECK(nvrtcGetPTX(prog, ptx.data()));
 
     char linkInfo[linkLogSize]  = {0};
     char linkError[linkLogSize] = {0};
 
-    CUlinkState linkState;
-    CUjit_option linkOptions[] = {
-        CU_JIT_INFO_LOG_BUFFER, CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES,
-        CU_JIT_ERROR_LOG_BUFFER, CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES,
-        CU_JIT_LOG_VERBOSE};
-
-    void *linkOptionValues[] = {
-        linkInfo, reinterpret_cast<void *>(linkLogSize), linkError,
-        reinterpret_cast<void *>(linkLogSize), reinterpret_cast<void *>(1)};
+    CUlinkState linkState = nullptr;
+    void *cubin           = nullptr;
+    size_t cubinSize      = 0;
+    vector<char> sass;
 
     auto link = high_resolution_clock::now();
-    CU_LINK_CHECK(cuLinkCreate(5, linkOptions, linkOptionValues, &linkState));
-    CU_LINK_CHECK(cuLinkAddData(linkState, CU_JIT_INPUT_PTX, (void *)ptx.data(),
-                                ptx.size(), moduleKey.c_str(), 0, NULL, NULL));
-
-    void *cubin = nullptr;
-    size_t cubinSize;
+#if CUDA_VERSION >= 11010
+    if (emitSASS) {
+        NVRTC_CHECK(nvrtcGetCUBINSize(prog, &cubinSize));
+        sass.resize(cubinSize);
+        NVRTC_CHECK(nvrtcGetCUBIN(prog, sass.data()));
+        cubin = sass.data();
+    }
+#endif
+    if (!emitSASS) {
+        size_t ptx_size;
+        vector<char> ptx;
+        NVRTC_CHECK(nvrtcGetPTXSize(prog, &ptx_size));
+        ptx.resize(ptx_size);
+        NVRTC_CHECK(nvrtcGetPTX(prog, ptx.data()));
+
+        CUjit_option linkOptions[] = {
+            CU_JIT_INFO_LOG_BUFFER, CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES,
+            CU_JIT_ERROR_LOG_BUFFER, CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES,
+            CU_JIT_LOG_VERBOSE};
+
+        void *linkOptionValues[] = {
+            linkInfo, reinterpret_cast<void *>(linkLogSize), linkError,
+            reinterpret_cast<void *>(linkLogSize), reinterpret_cast<void *>(1)};
+
+        CU_LINK_CHECK(
+            cuLinkCreate(5, linkOptions, linkOptionValues, &linkState));
+        CU_LINK_CHECK(cuLinkAddData(linkState, CU_JIT_INPUT_PTX,
+                                    (void *)ptx.data(), ptx.size(),
+                                    moduleKey.c_str(), 0, NULL, NULL));
+        CU_LINK_CHECK(cuLinkComplete(linkState, &cubin, &cubinSize));
+    }
 
     CUmodule modOut = nullptr;
-    CU_LINK_CHECK(cuLinkComplete(linkState, &cubin, &cubinSize));
     CU_CHECK(cuModuleLoadData(&modOut, cubin));
     auto link_end = high_resolution_clock::now();
 
@@ -391,7 +418,7 @@ Module compileModule(const string &moduleKey, span<const string> sources,
     }
 #endif
 
-    CU_LINK_CHECK(cuLinkDestroy(linkState));
+    if (linkState) { CU_LINK_CHECK(cuLinkDestroy(linkState)); }
     NVRTC_CHECK(nvrtcDestroyProgram(&prog));
 
     // skip --std=c++14 because it will stay the same. It doesn't
diff --git a/src/backend/cuda/kernel/regions.hpp b/src/backend/cuda/kernel/regions.hpp
index d1f1ed3b2..8e315442c 100644
--- a/src/backend/cuda/kernel/regions.hpp
+++ b/src/backend/cuda/kernel/regions.hpp
@@ -22,7 +22,7 @@
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/scan.h>
 #include <thrust/sort.h>
-#include <thrust/system/cuda/detail/par.h>
+#include <thrust/system/cuda/execution_policy.h>
 #include <thrust/transform_scan.h>
 
 static const int THREADS_X = 16;
diff --git a/src/backend/cuda/set.cu b/src/backend/cuda/set.cu
index d558d6e93..4d3a71f0d 100644
--- a/src/backend/cuda/set.cu
+++ b/src/backend/cuda/set.cu
@@ -16,6 +16,7 @@
 #include <af/dim4.hpp>
 
 #include <thrust/device_ptr.h>
+#include <thrust/distance.h>
 #include <thrust/set_operations.h>
 #include <thrust/sort.h>
 #include <thrust/unique.h>
diff --git a/src/backend/cuda/thrust_utils.hpp b/src/backend/cuda/thrust_utils.hpp
index 0646b934b..57d30a139 100644
--- a/src/backend/cuda/thrust_utils.hpp
+++ b/src/backend/cuda/thrust_utils.hpp
@@ -9,7 +9,7 @@
 
 #pragma once
 #include <ThrustArrayFirePolicy.hpp>
-#include <thrust/system/cuda/detail/par.h>
+#include <thrust/system/cuda/execution_policy.h>
 #include <thrust/version.h>
 #include <ThrustAllocator.cuh>

syurkevi · 2026-11-04T00:12:29Z

/cpu-tests all

syurkevi · 2026-11-04T00:12:43Z

/gpu-tests

github-actions · 2026-11-04T00:13:04Z

GPU tests requested by @syurkevi for 74baf54dbec623799158a58de0df2c75d7e941fd, the head of this pull request now: CUDA 12.9, CUDA 13.1 and OpenCL, chosen by the changed files: CUDA for CMakeLists.txt and OpenCL for CMakeLists.txt. The builds run first; the GPU jobs then wait for one maintainer approval. Later pushes are not part of this run. Run

github-actions · 2026-11-04T00:13:05Z

CPU tests (the accelerator builds are disabled in this repository) requested by @syurkevi for 74baf54dbec623799158a58de0df2c75d7e941fd, the head of this pull request now. Later pushes are not part of this run. Run

github-actions · 2026-11-04T00:55:35Z

GPU tests (CUDA 12.9, CUDA 13.1 and OpenCL) of 74baf54dbec623799158a58de0df2c75d7e941fd passed: CUDA 12.9: 133 of 133 passed; CUDA 13.1: 133 of 133 passed; OpenCL: 135 of 135 passed Run

CUDA 12.9: 133 of 133 passed
CUDA 13.1: 133 of 133 passed
OpenCL: 135 of 135 passed

github-actions · 2026-11-04T01:09:45Z

CPU tests of 74baf54dbec623799158a58de0df2c75d7e941fd passed Run

melonakos force-pushed the fix/cuda-13-build branch from 02dfe7e to 85ceb40 Compare September 10, 2026 18:16

melonakos mentioned this pull request Sep 11, 2026

cuda: serialise constant-memory uploads with their kernel launches #3725

Draft

WilliamTambellini reviewed Oct 13, 2026

View reviewed changes

roaffix mentioned this pull request Nov 2, 2026

Rework the CI: light checks on every push, CPU tests in the merge queue, GPU tests on request #3738

Merged

5 tasks done

syurkevi force-pushed the fix/cuda-13-build branch from cd211d5 to 0d3fd2e Compare November 3, 2026 05:01

syurkevi marked this pull request as ready for review November 3, 2026 05:04

syurkevi added 6 commits November 3, 2026 18:58

build updates for cuda 13

09e4cdb

fix fp64 aspect errors on devices without double support

6da6300

fix doxygen warnings and windows docs build

9a51ad3

cuda: load JIT kernels on drivers older than the toolkit

d9c18be

prepare 3.10.1 release with cuda 12.9 and 13.4 support

5069ca8

apply clang-format 15 to src, test and examples

74baf54

syurkevi force-pushed the fix/cuda-13-build branch from a0eed17 to 74baf54 Compare November 4, 2026 00:11

syurkevi approved these changes Nov 4, 2026

View reviewed changes

syurkevi merged commit 64c88d9 into master Nov 4, 2026
10 checks passed

Navigation Menu

Repository navigation

cuda: build with CUDA 13 and target Blackwell - #3715

cuda: build with CUDA 13 and target Blackwell#3715
syurkevi merged 6 commits into
masterfrom
fix/cuda-13-build

melonakos commented Sep 10, 2026 •

edited

Loading

Uh oh!

WilliamTambellini Oct 13, 2026

Uh oh!

WilliamTambellini left a comment

Uh oh!

graehl commented Nov 3, 2026 •

edited

Loading

Uh oh!

syurkevi commented Nov 4, 2026

Uh oh!

syurkevi commented Nov 4, 2026

Uh oh!

github-actions Bot commented Nov 4, 2026

Uh oh!

github-actions Bot commented Nov 4, 2026

Uh oh!

github-actions Bot commented Nov 4, 2026

Uh oh!

github-actions Bot commented Nov 4, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

Conversation

melonakos commented Sep 10, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

WilliamTambellini Oct 13, 2026

Choose a reason for hiding this comment

Uh oh!

WilliamTambellini left a comment

Choose a reason for hiding this comment

Uh oh!

graehl commented Nov 3, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Linux test report: CUDA 13.3 on an sm_120 Blackwell card

Build failures with CUDA 13.3

JIT kernels fail when the toolkit is newer than the driver

Test harness and minor notes

Uh oh!

syurkevi commented Nov 4, 2026

Uh oh!

syurkevi commented Nov 4, 2026

Uh oh!

github-actions Bot commented Nov 4, 2026

Uh oh!

github-actions Bot commented Nov 4, 2026

Uh oh!

github-actions Bot commented Nov 4, 2026

Uh oh!

github-actions Bot commented Nov 4, 2026

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

melonakos commented Sep 10, 2026 •

edited

Loading

graehl commented Nov 3, 2026 •

edited

Loading