Add support for CUDA 13.1 and Blackwell GPUS

Update CUDA and driver mappings. Update cuFFT and cuSolver library versions. Fixes for depricated interfaces.
melonakos · Mar 7, 2026 · Sep 10, 2026 · Sep 10, 2026 · Sep 10, 2026 · Sep 10, 2026
commit 85ceb407486d7bcea5ac9011f6c6b1188878c7cb
diff --git a/CMakeModules/select_compute_arch.cmake b/CMakeModules/select_compute_arch.cmake
@@ -7,7 +7,7 @@
 #      ARCH_AND_PTX : NAME | NUM.NUM | NUM.NUM(NUM.NUM) | NUM.NUM+PTX
 #      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
 #      Returns LIST of flags to be added to CUDA_NVCC_FLAGS in ${out_variable}
 #      Additionally, sets ${out_variable}_readable to the resulting numeric list
 #      Example:
@@ -111,6 +111,18 @@ if(CUDA_VERSION VERSION_GREATER_EQUAL "12.0")
   list(REMOVE_ITEM CUDA_ALL_GPU_ARCHITECTURES "3.5" "3.7")
 endif()
 
+if(CUDA_VERSION VERSION_GREATER_EQUAL "12.8")
+  list(APPEND CUDA_KNOWN_GPU_ARCHITECTURES "Blackwell")
+  list(APPEND CUDA_COMMON_GPU_ARCHITECTURES "12.0")
+  list(APPEND CUDA_ALL_GPU_ARCHITECTURES "10.0" "10.3" "11.0" "12.0")
+
+  set(_CUDA_MAX_COMMON_ARCHITECTURE "12.0+PTX")
+  set(CUDA_LIMIT_GPU_ARCHITECTURE "12.0")
+
+  list(REMOVE_ITEM CUDA_COMMON_GPU_ARCHITECTURES "5.0" "5.3" "6.0" "6.1")
+  list(REMOVE_ITEM CUDA_ALL_GPU_ARCHITECTURES "5.0" "5.2" "5.3" "6.0" "6.1" "6.2")
+endif()
+
 list(APPEND CUDA_COMMON_GPU_ARCHITECTURES "${_CUDA_MAX_COMMON_ARCHITECTURE}")
 
 # Check with: cmake -DCUDA_VERSION=7.0 -P select_compute_arch.cmake
@@ -229,7 +241,7 @@ function(CUDA_SELECT_NVCC_ARCH_FLAGS out_variable)
       set(add_ptx TRUE)
       set(arch_name ${CMAKE_MATCH_1})
     endif()
-    if(arch_name MATCHES "^([0-9]\\.[0-9](\\([0-9]\\.[0-9]\\))?)$")
+    if(arch_name MATCHES "^([0-9]+\\.[0-9]+(\\([0-9]+\\.[0-9]+\\))?)$")
       set(arch_bin ${CMAKE_MATCH_1})
       set(arch_ptx ${arch_bin})
     else()

diff --git a/src/backend/cuda/CMakeLists.txt b/src/backend/cuda/CMakeLists.txt
@@ -858,7 +858,13 @@ if(AF_INSTALL_STANDALONE)
   endif()
 
   if(WIN32 OR NOT AF_WITH_STATIC_CUDA_NUMERIC_LIBS)
-    if(CUDA_VERSION_MAJOR VERSION_EQUAL 12)
+    if(CUDA_VERSION_MAJOR VERSION_EQUAL 13)
+        if(CUDA_VERSION_MINOR VERSION_GREATER_EQUAL 1)
+            afcu_collect_libs(cufft LIB_MAJOR 12 LIB_MINOR 1)
+        else()
+            afcu_collect_libs(cufft LIB_MAJOR 12 LIB_MINOR 0)
+        endif()
+    elseif(CUDA_VERSION_MAJOR VERSION_EQUAL 12)
         afcu_collect_libs(cufft LIB_MAJOR 11 LIB_MINOR 3)
     elseif(CUDA_VERSION_MAJOR VERSION_EQUAL 11)
         afcu_collect_libs(cufft LIB_MAJOR 10 LIB_MINOR 4)
@@ -869,7 +875,9 @@ if(AF_INSTALL_STANDALONE)
     if(CUDA_VERSION VERSION_GREATER 10.0)
       afcu_collect_libs(cublasLt)
     endif()
-    if(CUDA_VERSION_MAJOR VERSION_EQUAL 12)
+    if(CUDA_VERSION_MAJOR VERSION_EQUAL 13)
+        afcu_collect_libs(cusolver LIB_MAJOR 12 LIB_MINOR 0)
+    elseif(CUDA_VERSION_MAJOR VERSION_EQUAL 12)
         afcu_collect_libs(cusolver LIB_MAJOR 11 LIB_MINOR 7)
     else()
         afcu_collect_libs(cusolver)
@@ -879,7 +887,9 @@ if(AF_INSTALL_STANDALONE)
       afcu_collect_libs(nvJitLink)
     endif()
   elseif(NOT ${use_static_cuda_lapack})
-    if(CUDA_VERSION_MAJOR VERSION_EQUAL 12)
+    if(CUDA_VERSION_MAJOR VERSION_EQUAL 13)
+        afcu_collect_libs(cusolver LIB_MAJOR 12 LIB_MINOR 0)
+    elseif(CUDA_VERSION_MAJOR VERSION_EQUAL 12)
         afcu_collect_libs(cusolver LIB_MAJOR 11 LIB_MINOR 7)
     else()
         afcu_collect_libs(cusolver)

diff --git a/src/backend/cuda/cufft.cu b/src/backend/cuda/cufft.cu
@@ -38,19 +38,25 @@ const char *_cufftGetResultString(cufftResult res) {
 
         case CUFFT_UNALIGNED_DATA: return "cuFFT: unaligned data (deprecated)";
 
+#if CUDA_VERSION < 13000
         case CUFFT_INCOMPLETE_PARAMETER_LIST:
             return "cuFFT: call is missing parameters";
+#endif
 
         case CUFFT_INVALID_DEVICE:
             return "cuFFT: plan execution different than plan creation";
 
+#if CUDA_VERSION < 13000
         case CUFFT_PARSE_ERROR: return "cuFFT: plan parse error";
+#endif
 
         case CUFFT_NO_WORKSPACE: return "cuFFT: no workspace provided";
 
         case CUFFT_NOT_IMPLEMENTED: return "cuFFT: not implemented";
 
+#if CUDA_VERSION < 13000
         case CUFFT_LICENSE_ERROR: return "cuFFT: license error";
+#endif
 
 #if CUDA_VERSION >= 8000
         case CUFFT_NOT_SUPPORTED: return "cuFFT: not supported";

diff --git a/src/backend/cuda/device_manager.cpp b/src/backend/cuda/device_manager.cpp
@@ -101,6 +101,8 @@ static const int jetsonComputeCapabilities[] = {
 
 // clang-format off
 static const cuNVRTCcompute Toolkit2MaxCompute[] = {
+    {13010, 9, 0, 0},
+    {13000, 9, 0, 0},
     {12090, 9, 0, 0},
     {12080, 9, 0, 0},
     {12070, 9, 0, 0},
@@ -147,6 +149,8 @@ struct ComputeCapabilityToStreamingProcessors {
 // clang-format off
 static const ToolkitDriverVersions
     CudaToDriverVersion[] = {
+        {13010, 580.65f, 580.65f},
+        {13000, 580.65f, 580.65f},
         {12090, 525.60f, 528.33f},
         {12080, 525.60f, 528.33f},
         {12070, 525.60f, 528.33f},
@@ -598,9 +602,15 @@ DeviceManager::DeviceManager()
                 AF_TRACE("Unsuppored device: {}", dev.prop.name);
                 continue;
             } else {
+                int clockRate;
+                #if CUDA_VERSION < 13000
+                clockRate = dev.prop.clockRate;
+                #else
+                CUDA_CHECK(cudaDeviceGetAttribute(&clockRate, cudaDevAttrClockRate, i));
+                #endif
                 dev.flops = static_cast<size_t>(dev.prop.multiProcessorCount) *
                             compute2cores(dev.prop.major, dev.prop.minor) *
-                            dev.prop.clockRate;
+                            clockRate;
                 dev.nativeId = i;
                 AF_TRACE(
                     "Found device: {} (sm_{}{}) ({:0.3} GB | ~{} GFLOPs | {} "

diff --git a/src/backend/cuda/kernel/regions.hpp b/src/backend/cuda/kernel/regions.hpp
@@ -341,7 +341,11 @@ __global__ static void update_equiv(arrayfire::cuda::Param<T> equiv_map,
 }
 
 template<typename T>
+#if CUDA_VERSION < 13000
 struct clamp_to_one : public thrust::unary_function<T, T> {
+#else
+struct clamp_to_one {
+#endif
     __host__ __device__ T operator()(const T& in) const {
         return (in >= (T)1) ? (T)1 : in;
     }