CUDA Toolkit 13.4 Update 1 - Release Notes

1. Overview

Welcome to the release notes for NVIDIA® CUDA® Toolkit 13.4 Update 1. This release includes enhancements and fixes across the CUDA Toolkit and its libraries.

This documentation is organized into two main sections:

  • CUDA Platform

    Focuses on the core CUDA infrastructure including component versions, driver compatibility, compiler/runtime features, issues, and deprecations.

  • CUDA Libraries

    Covers the specialized computational libraries with their feature updates, performance improvements, API changes, and version history across CUDA 13.x releases.

2. CUDA Platform

2.1. CUDA Toolkit Major Components

For CUDA 13.4 Update 1, the table below indicates the versions:

Table 1 CUDA 13.4 Update 1 Component Versions

Component Name

Version Information

Supported Architectures

Supported Platforms

CUDA C++ Core Compute Libraries

Thrust

3.4.3

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUB

3.4.3

libcu++

3.4.3

Cooperative Groups

13.3.4.3.1

CUDA Compatibility Package (Orin)

13.4.47145772

arm64-sbsa

Linux

CUDA Application Compiler (crt)

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA Compilation Optimizer (ctadvisor)

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA Runtime (cudart)

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA culibos

13.4.92

x86_64, arm64-sbsa

Linux

CUDA cuobjdump

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUPTI

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA cuxxfilt (demangler)

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA Documentation

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA GDB

13.4.92

x86_64, arm64-sbsa

Linux

CUDA NVCC

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA nvdisasm

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA NVML Headers

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA nvprune

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA NVRTC

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA NVTX

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA OpenCL

13.4.92

x86_64, arm64 (Windows)

Linux, Windows

CUDA Profiler API

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA Sandbox dev

13.4.92

x86_64, arm64-sbsa

Linux

CUDA Compute Sanitizer API

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA TILE-IR AS

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA cuBLAS

13.8.0.4

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA cuDLA

13.4.92

arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA cuFFT

12.4.0.43

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA cuFile

1.19.1.55

x86_64, arm64-sbsa

Linux

CUDA cuobjclient

1.3.1.55

x86_64, arm64-sbsa

Linux

CUDA cuRAND

10.4.4.72

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA cuSOLVER

12.3.4.7

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA cuSPARSE

12.8.6.72

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA NPP

13.2.0.58

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA nvFatbin

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA nvJitLink

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA nvJPEG

13.2.3.58

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA nvptxcompiler

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

CUDA nvvm

13.4.92

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

Nsight Compute

2026.3.1.2

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

Nsight Systems

2026.3.2.476

x86_64, arm64-sbsa, arm64 (Windows)

Linux, Windows

Nsight Visual Studio Edition (VSE)

2026.3.0.26187

x86_64, arm64 (Windows)

Windows

Visual Studio Integration

13.4.92

x86_64, arm64 (Windows)

Windows

2.2. CUDA Driver

Note

The NVIDIA driver is no longer bundled with the CUDA Toolkit – on Windows starting with CUDA 13.1, and on Linux starting with CUDA 13.4. Download and install the appropriate driver from the official NVIDIA Driver Downloads page.

Running a CUDA application requires a system with at least one CUDA-capable GPU and a driver that is compatible with the CUDA Toolkit. For more information about various GPU products that are CUDA-capable, visit https://developer.nvidia.com/cuda/gpus.

The NVIDIA driver branch corresponding to each CUDA Toolkit release is shown below. Update releases within a CUDA minor version use the same driver branch.

Table 2 CUDA Toolkit and Corresponding Driver Branch

CUDA Toolkit

Corresponding Driver Branch

CUDA 13.4

R615

CUDA 13.3

R610

CUDA 13.2

R595

CUDA 13.1

R590

CUDA 13.0

R580

Note

Existing CUDA 13.x applications run on drivers >=580 under CUDA minor version compatibility. CUDA 13.4 new features and newly enabled platforms require an R615 or later driver that supports them. The Windows driver for the RTX Spark device is 616.41 or later.

The CUDA driver is backward compatible: applications compiled against a particular CUDA Toolkit version continue to work on subsequent (later) driver releases. In addition, CUDA minor version compatibility allows applications to run on a driver older than the corresponding driver branch, within the ranges shown below. The installed driver must meet or exceed the minimum required version for the CUDA Toolkit. For details, see the CUDA Compatibility Guide and CUDA Compatibility and Upgrades.

Table 3 CUDA Toolkit and Minimum Required Driver Version for CUDA Minor Version Compatibility

CTK Version

Driver Range for Minor Version Compatibility

Min

Max

13.x

>= 580

N/A

12.x

>= 525

< 580

11.x

>= 450

< 525

CUDA 11.0 shipped with earlier driver versions. Minor-version compatibility across the CUDA 11.x family requires driver version 450.80.02 or later on Linux, or 452.39 or later on Windows.

Older CUDA versions (12.9 and earlier)
Table 4 Corresponding Driver Versions (CUDA 12.9 and earlier)

CUDA Toolkit

Corresponding Driver Version

Linux x86_64 Driver Version

Windows x86_64 Driver Version

CUDA 12.9 Update 1

>=575.57.08

>=576.57

CUDA 12.9 GA

>=575.51.03

>=576.02

CUDA 12.8 Update 1

>=570.124.06

>=572.61

CUDA 12.8 GA

>=570.26

>=570.65

CUDA 12.6 Update 3

>=560.35.05

>=561.17

CUDA 12.6 Update 2

>=560.35.03

>=560.94

CUDA 12.6 Update 1

>=560.35.03

>=560.94

CUDA 12.6 GA

>=560.28.03

>=560.76

CUDA 12.5 Update 1

>=555.42.06

>=555.85

CUDA 12.5 GA

>=555.42.02

>=555.85

CUDA 12.4 Update 1

>=550.54.15

>=551.78

CUDA 12.4 GA

>=550.54.14

>=551.61

CUDA 12.3 Update 1

>=545.23.08

>=546.12

CUDA 12.3 GA

>=545.23.06

>=545.84

CUDA 12.2 Update 2

>=535.104.05

>=537.13

CUDA 12.2 Update 1

>=535.86.09

>=536.67

CUDA 12.2 GA

>=535.54.03

>=536.25

CUDA 12.1 Update 1

>=530.30.02

>=531.14

CUDA 12.1 GA

>=530.30.02

>=531.14

CUDA 12.0 Update 1

>=525.85.12

>=528.33

CUDA 12.0 GA

>=525.60.13

>=527.41

CUDA 11.8 GA

>=520.61.05

>=520.06

CUDA 11.7 Update 1

>=515.48.07

>=516.31

CUDA 11.7 GA

>=515.43.04

>=516.01

CUDA 11.6 Update 2

>=510.47.03

>=511.65

CUDA 11.6 Update 1

>=510.47.03

>=511.65

CUDA 11.6 GA

>=510.39.01

>=511.23

CUDA 11.5 Update 2

>=495.29.05

>=496.13

CUDA 11.5 Update 1

>=495.29.05

>=496.13

CUDA 11.5 GA

>=495.29.05

>=496.04

CUDA 11.4 Update 4

>=470.82.01

>=472.50

CUDA 11.4 Update 3

>=470.82.01

>=472.50

CUDA 11.4 Update 2

>=470.57.02

>=471.41

CUDA 11.4 Update 1

>=470.57.02

>=471.41

CUDA 11.4.0 GA

>=470.42.01

>=471.11

CUDA 11.3.1 Update 1

>=465.19.01

>=465.89

CUDA 11.3.0 GA

>=465.19.01

>=465.89

CUDA 11.2.2 Update 2

>=460.32.03

>=461.33

CUDA 11.2.1 Update 1

>=460.32.03

>=461.09

CUDA 11.2.0 GA

>=460.27.03

>=460.82

CUDA 11.1.1 Update 1

>=455.32

>=456.81

CUDA 11.1 GA

>=455.23

>=456.38

CUDA 11.0.3 Update 1

>= 450.51.06

>= 451.82

CUDA 11.0.2 GA

>= 450.51.05

>= 451.48

CUDA 11.0.1 RC

>= 450.36.06

>= 451.22

CUDA 10.2.89

>= 440.33

>= 441.22

CUDA 10.1 (10.1.105 general release, and updates)

>= 418.39

>= 418.96

CUDA 10.0.130

>= 410.48

>= 411.31

CUDA 9.2 (9.2.148 Update 1)

>= 396.37

>= 398.26

CUDA 9.2 (9.2.88)

>= 396.26

>= 397.44

CUDA 9.1 (9.1.85)

>= 390.46

>= 391.29

CUDA 9.0 (9.0.76)

>= 384.81

>= 385.54

CUDA 8.0 (8.0.61 GA2)

>= 375.26

>= 376.51

CUDA 8.0 (8.0.44)

>= 367.48

>= 369.30

CUDA 7.5 (7.5.16)

>= 352.31

>= 353.66

CUDA 7.0 (7.0.28)

>= 346.46

>= 347.62

2.3. New Features

2.3.1. CUDA Platform

  • None

2.3.2. CUDA Developer Tools

For details on new features, improvements, and bug fixes, see the changelogs for:

2.3.3. CUDA Compiler

2.3.4. CUDA C++ Core Libraries (CCCL)

  • None

2.3.5. CUDA Python

  • None

2.3.6. CUDA Tile

  • None

2.3.7. CUDA Tile IR

  • None

2.4. Resolved Issues

2.4.1. General CUDA

  • Fixed an issue where applications statically built with, or dynamically linking to, an older CUDA Runtime (down to CUDA 11.2) with the R615 driver could report benign public error messages indicating that IMEX channels were incorrectly set up.

2.4.2. CUDA Compiler

  • Fixed an issue where kernels using the fabric.try_put.tensor, fabric.try_red.tensor, or fabric.try_atom PTX instructions did not set the fabric-write grid attribute, which could leave fabric writes unfenced at grid completion when the kernel was launched in a CUDA graph.

2.4.3. CUDA Tools

  • None

2.5. Known Issues

2.5.1. CUDA Platform

  • None

2.5.2. CUDA Compiler

  • None

2.6. Deprecated or Dropped Features

2.6.1. CUDA Platform

  • Legacy Nsight Eclipse Edition plugins are no longer delivered in CUDA Toolkit packages beginning with CUDA 13.3.

  • Python 3.10 support is deprecated across the CUDA Python 13.4 packages.

2.6.2. Architectures

  • CUDA 14.0 will move to Armv8.2-A as the minimum supported architecture for ARM64-SBSA.

2.6.3. Operating Systems

  • None

2.6.4. CUDA Toolchains

  • None

3. CUDA Libraries

This section covers CUDA Libraries release notes for 13.x releases.

Note

Documentation will be updated to accurately reflect supported C++ standard libraries for CUDA Math Libraries.

3.1. cuBLAS Library

3.1.1. cuBLAS: Release 13.4 Update 1

  • New Features

    • Emulated FP64 matrix multiplications:

      • Improved ZGEMM peak performance by up to 25% on Rubin GPUs.

      • When the per-handle workspace is insufficient, fixed-point FP64 emulation now allocates its temporary workspace from a cuBLAS-managed, per-device CUDA memory pool that retains memory across stream synchronizations. This avoids repeated allocation overhead without requiring users to tune the default memory pool.

  • Known Issues

    • When running NVIDIA Compute Sanitizer (versions up to CUDA Toolkit 13.4) with cuBLAS-linked workloads, the tool may report false-positive errors regarding invalid global reads or out-of-bounds (OOB) memory accesses inside cuBLAS accelerated kernels. This is a known issue caused by a limitation in Compute Sanitizer’s ability to accurately identify the PTX semantics of the asynchronous global-to-shared memory copy instruction cp.async.cg.shared.global [dst], [src], cp-size, src-size when src-size < cp-size. In this specific scenario, where the number of bytes to copy exceeds the source buffer size, the hardware guarantees that the remaining bytes (cp-size - src-size) in the destination shared memory are automatically padded with zeros, ensuring that the memory access is neither invalid nor OOB. Hence, Compute Sanitizer reports for these specific cases can be safely ignored. [6196059]

    • cuBLASLt Grouped GEMM with per-batch tensor-wide scales causes an invalid memory access for groups where m > 0, n > 0, and k = 0. As a workaround, always pass valid scale pointers for each group. This issue was introduced in CUDA Toolkit 13.1 (cuBLAS 13.2.0). [CUB-10481]

  • Resolved Issues

    • Fixed an issue where cuBLASLt Grouped GEMM operations with CUBLAS_POINTER_MODE_HOST could lead to incorrect results on Hopper GPUs when the number of waves was larger than 2. This issue was introduced in CUDA Toolkit 13.2 Update 1 (cuBLAS 13.4.0). [6681084]

    • Fixed an issue where cuBLASLt Grouped GEMM operations could lead to incorrect results on B300 and Rubin GPUs when the input matrices used the NVFP4 data type, the number of waves was larger than 3, and the algorithm attribute CUBLASLT_ALGO_CONFIG_STAGES_ID was CUBLASLT_MATMUL_STAGES_768xAUTO. This issue was introduced in CUDA Toolkit 13.4 (cuBLAS 13.7.0). [6681084]

    • Fixed an issue where cublasLtMatmul() could return incorrect output when using split-K if the number of K splits (SPLITK_NUM) did not evenly divide the K dimension and the floor division of K / SPLITK_NUM was an exact multiple of the stage size. This affected only algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66 on GPUs with compute capability 9.0, 10.x, and 11.x. This issue was introduced in CUDA Toolkit 12.6 Update 2 (cuBLAS 12.6.3). [6580689]

    • Fixed an issue where cublasLtMatmul() could produce incorrect results on B300 and Rubin GPUs when the input matrices used the NVFP4 data type. This affected only algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66 and the algorithm attribute CUBLASLT_ALGO_CONFIG_STAGES_ID set to CUBLASLT_MATMUL_STAGES_768xAUTO. This issue was introduced in CUDA Toolkit 13.3 (cuBLAS 13.5.1). [CUB-10573]

3.1.2. cuBLAS: Release 13.4

  • New Features

    • Emulated FP64 matrix multiplications:

      • Emulated FP64 matrix multiplications now leverage the Ozaki-II scheme when it provides a performance benefit over the Ozaki-I scheme; the Ozaki-II scheme is supported on NVIDIA Ampere and newer GPUs.

      • On B200 and RTX PRO 6000 Blackwell Server Edition GPUs, this enables up to 175 and 45 TFLOPS of emulated DGEMM performance, respectively, and up to 295 and 70 TFLOPS of emulated ZGEMM performance, respectively, by leveraging the 2M algorithm.

      • Emulated FP64 matrix multiplications now support Rubin GPUs (compute capability 10.7), leveraging the Ozaki-I and Ozaki-II schemes and the new TI16 type via the tcgen05.mma instruction to achieve up to 212 TFLOPS of emulated DGEMM performance. Emulated ZGEMM can reach up to 301 TFLOPS, with further improvements to come.

    • cuBLASLt adds experimental support for an alternative scaling-factor layout for MXFP8 matmuls through the CUBLASLT_MATMUL_MATRIX_SCALE_VEC32_MN_K4_UE8M0 and CUBLASLT_MATMUL_MATRIX_SCALE_VEC128_MN_K4_UE8M0 scaling modes. For more information, see the cuBLAS documentation.

    • Added support for the NVIDIA Rubin (compute capability 10.7) GPU architecture.

    • cuBLASLt Grouped GEMM performance is improved on Blackwell data center GPUs through dynamic scheduling of matrix computations. Improvements of up to 20% can be seen for Grouped GEMM calls with a large number of groups (for example, 32). [5913842] [5992105] [CUB-9925]

  • Known Issues

    • When running NVIDIA Compute Sanitizer (versions up to CUDA Toolkit 13.4) with cuBLAS-linked workloads, the tool may report false-positive errors regarding invalid global reads or out-of-bounds (OOB) memory accesses inside cuBLAS accelerated kernels. This is a known issue caused by a limitation in Compute Sanitizer’s ability to accurately identify the PTX semantics of the asynchronous global-to-shared memory copy instruction cp.async.cg.shared.global [dst], [src], cp-size, src-size when src-size < cp-size. In this specific scenario, where the number of bytes to copy exceeds the source buffer size, the hardware guarantees that the remaining bytes (cp-size - src-size) in the destination shared memory are automatically padded with zeros, ensuring that the memory access is neither invalid nor OOB. Hence, Compute Sanitizer reports for these specific cases can be safely ignored. [6196059]

    • cuBLASLt Grouped GEMM with per-batch tensor-wide scales causes an invalid memory access for groups where m > 0, n > 0, and k = 0. As a workaround, always pass valid scale pointers for each group. This issue was introduced in CUDA Toolkit 13.1 (cuBLAS 13.2.0). [CUB-10481]

    • cublasLtMatmul() can return incorrect output when using split-K if the number of K splits (SPLITK_NUM) does not evenly divide the K dimension and the floor division of K / SPLITK_NUM is an exact multiple of the stage size. This affects only algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66 on GPUs with compute capability 9.0, 10.x, and 11.x. As a workaround, use cublasLtMatmulAlgoConfigGetAttribute() to query the number of K splits (CUBLASLT_ALGO_CONFIG_SPLITK_NUM) and the stage size (derived from CUBLASLT_ALGO_CONFIG_STAGES_ID), then use cublasLtMatmulAlgoConfigSetAttribute() to set a value that evenly divides the K dimension, or for which K / SPLITK_NUM is not an exact multiple of the stage size. This issue was first introduced in CUDA Toolkit 12.6 Update 2 (cuBLAS 12.6.3). [6580689]

    • cuBLASLt Grouped GEMM operations with CUBLAS_POINTER_MODE_HOST can lead to incorrect results on Hopper GPUs when the number of waves is larger than 2. As a workaround, pass device alpha and beta using CUBLAS_POINTER_MODE_DEVICE. This issue was introduced in CUDA Toolkit 13.2 Update 1 (cuBLAS 13.4.0). [6681084]

    • cuBLASLt Grouped GEMM operations can lead to incorrect results on B300 and Rubin GPUs when the input matrices use the NVFP4 data type, the number of waves is larger than 3, and the algorithm attribute CUBLASLT_ALGO_CONFIG_STAGES_ID is CUBLASLT_MATMUL_STAGES_768xAUTO. As a workaround, skip such candidates when using cublasLtMatmulAlgoGetHeuristic(). [6681084]

    • cublasLtMatmul() can produce incorrect results on B300 and Rubin GPUs when the input matrices use the NVFP4 data type. This affects only algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66 and the algorithm attribute CUBLASLT_ALGO_CONFIG_STAGES_ID is CUBLASLT_MATMUL_STAGES_768xAUTO. [CUB-10573]

  • Resolved Issues

    • Between CUDA Toolkit 13.3 Update 1 and 13.4, cuBLAS released an independent patch release (cuBLAS 13.6.1) that resolves several issues. Refer to the associated cuBLAS patch release notes for details.

    • Fixed an issue where cublasLtMatmulAlgoGetHeuristic() could return no algorithms, and cublasLtMatmul() could return CUBLAS_STATUS_NOT_SUPPORTED, for some NVFP4 matmuls with NVFP4 output and the CUBLASLT_EPILOGUE_BIAS epilogue on GPUs with compute capability 12.x. [CUB-10222]

    • Fixed an issue where cublasLtMatmulAlgoGetHeuristic() returned CUBLAS_STATUS_INTERNAL_ERROR and invalidated an in-flight CUDA graph capture when called while a non-default blocking stream was being captured. This issue was introduced in CUDA Toolkit 13.3 (cuBLAS 13.5.1). [6288786]

    • Fixed an issue where cublas<t>gemv() with trans equal to CUBLAS_OP_T or CUBLAS_OP_C could perform an illegal memory access when the number of elements addressed by the output vector (n * incy) exceeded the 32-bit index range. cublasZgemv() could additionally return CUBLAS_STATUS_NOT_SUPPORTED when n * lda exceeded that range. [6207926]

    • Fixed an issue where cublasZgemv() with trans equal to CUBLAS_OP_C and incy greater than 1 produced incorrect results when FP64 fixed-point emulation was enabled (CUBLAS_FP64_EMULATED_FIXEDPOINT_MATH with CUBLAS_EMULATION_STRATEGY_EAGER). [6207926]

3.1.3. cuBLAS: Release 13.3 Update 1

  • New Features

    • The TMA-based kernel (Hopper and newer) now accelerates DSYMV in addition to the already-enabled SSYMV. The 16-byte alignment requirement for the A pointer was dropped for this kernel, and support for atomics was added through cublasSetAtomicsMode(). The geomean speedup across architectures and datatypes is 1.3x, and up to 5.9x.

  • Known Issues

    • Non-default epilogues are unintentionally allowed for cublasLtMatmul() with int8 inputs using regular data ordering and scale type CUDA_R_32F. This is an undocumented and lightly tested feature that users are discouraged from using, and it is planned for removal in the next major release. [CUB-10067]

    • In cuBLASLt, the heuristics for the Grouped GEMM API return sub-optimal algorithms when the C and D matrices use CUBLASLT_ORDER_ROW ordering. As a workaround, swap CUBLASLT_MATMUL_PREF_GROUPED_DESC_D_AVERAGE_ROWS and CUBLASLT_MATMUL_PREF_GROUPED_DESC_D_AVERAGE_COLS in the preferences before calling cublasLtMatmulAlgoGetHeuristic(). [6335555]

    • Multiple cuBLASLt Grouped GEMM operations reusing the same workspace on Hopper GPUs may lead to hangs or unspecified launch failures. This issue was introduced in CUDA Toolkit 13.2 Update 1 (cuBLAS 13.4.0). [6456362]

    • Calling cublasLtMatmulAlgoGetHeuristic() while a non-default blocking stream is being captured returns CUBLAS_STATUS_INTERNAL_ERROR. This issue was introduced in CUDA Toolkit 13.3 (cuBLAS 13.5.1). [6288786]

    • GEMV-like operations (e.g., cublas<t>gemv(), or cublasLtMatmul() with M or N equal to 1) may return CUBLAS_STATUS_NOT_SUPPORTED for certain shapes and workspace configurations. This issue was introduced in CUDA Toolkit 13.3 (cuBLAS 13.5.1).

    • Strided batched GEMM operations using broadcast operands may perform out-of-bounds memory reads on GPUs with compute capability 12.0 or 12.1 when an input matrix shared across batches (via zero or overlapping strides) is smaller than a few hundred KB. Numerical accuracy is unaffected, but these invalid accesses can trigger compute-sanitizer warnings or, in rare instances, illegal memory access errors. This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [6040940] [5996751]

    • GEMM kernels on GPUs with compute capability 10.x or 12.x may access device alpha and beta pointers before calling cudaGridDependencySynchronize(). This can result in a WAR hazard if the preceding PDL kernel produces alpha and beta values on device after calling cudaTriggerProgrammaticLaunchCompletion(). This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [CUB-10409]

    • GEMV-like operations (e.g., cublas<t>gemv(), or cublasLtMatmul() with M or N equal to 1) may produce incorrect results on GPUs with compute capability 12.0 or 12.1 when the matrix is transposed, its non-accumulation dimension is greater than 3145680, and beta is not equal to 0. This issue was introduced in CUDA Toolkit 13.3 Update 1 (cuBLAS 13.6.0). [CUB-10445]

    • cuBLASLt Grouped GEMM operations with CUBLAS_POINTER_MODE_HOST can lead to incorrect results on Hopper GPUs when the number of waves is larger than 2. As a workaround, pass device alpha and beta using CUBLAS_POINTER_MODE_DEVICE. This issue was introduced in CUDA Toolkit 13.2 Update 1 (cuBLAS 13.4.0). [6681084]

  • Resolved Issues

    • Fixed an issue where cublasXt<t>spmm() could produce incorrect results with m greater than 46340. [6155165]

    • Fixed an issue where cublasLtMatmul() could run an unsupported combination of data types: an FP32-like compute type with FP32 C and D and non-FP32 A and B, in which case A and B are incorrectly interpreted as FP32 matrices. [CUB-9942]

    • Fixed an issue where cublasLtMatmul() returned CUBLAS_STATUS_NOT_SUPPORTED for FP8 Grouped GEMM problems with scale modes CUBLASLT_MATMUL_MATRIX_SCALE_VEC128_32F and CUBLASLT_MATMUL_MATRIX_SCALE_BLK128x128_32F on Hopper GPUs. [CUB-10031]

    • Fixed an issue where cublasLtMatmul() with int8 inputs and scale type CUDA_R_32I would allow non-default epilogues on Blackwell GPUs with compute capability 10.x and return incorrect results. The correct behavior is to disallow all but the default epilogue, as documented. [CUB-10066]

3.1.4. cuBLAS: Release 13.3

  • New Features

    • Enabled memory-parsimonious tiling for FP64 emulated matrix multiplications. This improvement ensures that the workspace memory budget no longer exceeds 8 GB.

    • Added support for CUDA Green contexts.

    • Improved FP4 matrix multiplication performance on Blackwell Ultra GPUs by a geometric mean of 5% across a wide range of problems, with up to 7% speedup for some small problems.

    • Improved TF32 matrix multiplication performance on Blackwell and Blackwell Ultra GPUs by a geometric mean of 27% across a wide range of problems and layouts, with up to 3.5x speedup for some small problems.

    • Improved TF32 TN matrix multiplication performance on Hopper GPUs by a geometric mean of 11% across a wide range of problems, with up to 40% speedup for some small problems.

    • Improved SYMV performance with TMA-based acceleration for Hopper, Blackwell, and Blackwell Ultra kernels with up to 27% geomean speedup.

  • Known Issues

    • Multiple cuBLASLt Grouped GEMM operations reusing the same workspace on Hopper GPUs may lead to hangs or unspecified launch failures. This issue was introduced in CUDA Toolkit 13.2 Update 1 (cuBLAS 13.4.0). [6456362]

    • cuBLASLt Grouped GEMM operations with CUBLAS_POINTER_MODE_HOST can lead to incorrect results on Hopper GPUs when the number of waves is larger than 2. As a workaround, pass device alpha and beta using CUBLAS_POINTER_MODE_DEVICE. This issue was introduced in CUDA Toolkit 13.2 Update 1 (cuBLAS 13.4.0). [6681084]

    • Calling cublasLtMatmulAlgoGetHeuristic() while a non-default blocking stream is being captured returns CUBLAS_STATUS_INTERNAL_ERROR. This issue was introduced in CUDA Toolkit 13.3 (cuBLAS 13.5.1). [6288786]

    • GEMV-like operations (e.g., cublas<t>gemv(), or cublasLtMatmul() with M or N equal to 1) may return CUBLAS_STATUS_NOT_SUPPORTED for certain shapes and workspace configurations. This issue was introduced in CUDA Toolkit 13.3 (cuBLAS 13.5.1).

    • Strided batched GEMM operations using broadcast operands may perform out-of-bounds memory reads on GPUs with compute capability 12.0 or 12.1 when an input matrix shared across batches (via zero or overlapping strides) is smaller than a few hundred KB. Numerical accuracy is unaffected, but these invalid accesses can trigger compute-sanitizer warnings or, in rare instances, illegal memory access errors. This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [6040940] [5996751]

    • GEMM kernels on GPUs with compute capability 10.x or 12.x may access device alpha and beta pointers before calling cudaGridDependencySynchronize(). This can result in a WAR hazard if the preceding PDL kernel produces alpha and beta values on device after calling cudaTriggerProgrammaticLaunchCompletion(). This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [CUB-10409]

3.1.5. cuBLAS: Release 13.2 Update 2

  • Known Issues

    • Multiple cuBLASLt Grouped GEMM operations reusing the same workspace on Hopper GPUs may lead to hangs or unspecified launch failures. This issue was introduced in CUDA Toolkit 13.2 Update 1 (cuBLAS 13.4.0). [6456362]

    • Strided batched GEMM operations using broadcast operands may perform out-of-bounds memory reads on GPUs with compute capability 12.0 or 12.1 when an input matrix shared across batches (via zero or overlapping strides) is smaller than a few hundred KB. Numerical accuracy is unaffected, but these invalid accesses can trigger compute-sanitizer warnings or, in rare instances, illegal memory access errors. This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [6040940] [5996751]

    • GEMM kernels on GPUs with compute capability 10.x or 12.x may access device alpha and beta pointers before calling cudaGridDependencySynchronize(). This can result in a WAR hazard if the preceding PDL kernel produces alpha and beta values on device after calling cudaTriggerProgrammaticLaunchCompletion(). This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [CUB-10409]

  • Resolved Issues

    • Fixed an issue where cublasLtMatmul() ignored the tensor-wide scaling value specified by CUBLASLT_MATMUL_DESC_D_SCALE_POINTER for NVFP4 matrix multiplications with NVFP4 output, resulting in incorrect results. This affected algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66 on GPUs with compute capability 10.x and 11.x. This issue was introduced in CUDA Toolkit 13.2 Update 1 (cuBLAS 13.4.0). [6059292]

3.1.6. cuBLAS: Release 13.2 Update 1

Note

CUDA Toolkit 13.2 Update 1 contains a critical cuBLAS bug for an issue where cublasLtMatmul() could ignore tensor-wide scaling for NVFP4 matrix multiplications, resulting in incorrect results. Please see the cuBLAS patch release notes for an available cuBLAS patch (13.4.1) to resolve this issue.

  • New Features

    • Extended the experimental Grouped GEMM API in cuBLASLt to support NVFP4 inputs and bias epilogues on Blackwell GPUs with Compute Capability 10.x and 11.0. Grouped GEMM NVFP4 support currently uses only MMA tile sizes with K equal to 64.

    • Extended the experimental Grouped GEMM API in cuBLASLt to support BF16, FP16, and FP8 input data types with BF16, FP16, and FP32 output data types on Hopper GPUs. For FP8 inputs, tensorwide scaling and block scaling (VEC128 and BLK128x128) are supported.

    • Improved Grouped GEMM performance on Blackwell GPUs, providing up to 20% higher performance for large problem sizes where the matrices exceed the L2 cache size.

  • Known Issues

    • cublasLtMatmul() ignores the tensor-wide scaling value provided by CUBLASLT_MATMUL_DESC_D_SCALE_POINTER for NVFP4 matrix multiplications with NVFP4 output, leading to incorrect results. This affects algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66 on GPUs with compute capability 10.x and 11.x. This issue was introduced in CUDA Toolkit 13.2 Update 1. [6059292]

    • Multiple cuBLASLt Grouped GEMM operations reusing the same workspace on Hopper GPUs may lead to hangs or unspecified launch failures. This issue was introduced in CUDA Toolkit 13.2 Update 1 (cuBLAS 13.4.0). [6456362]

    • cuBLASLt Grouped GEMM operations with CUBLAS_POINTER_MODE_HOST can lead to incorrect results on Hopper GPUs when the number of waves is larger than 2. As a workaround, pass device alpha and beta using CUBLAS_POINTER_MODE_DEVICE. This issue was introduced in CUDA Toolkit 13.2 Update 1 (cuBLAS 13.4.0). [6681084]

    • Strided batched GEMM operations using broadcast operands may perform out-of-bounds memory reads on GPUs with compute capability 12.0 or 12.1 when an input matrix shared across batches (via zero or overlapping strides) is smaller than a few hundred KB. Numerical accuracy is unaffected, but these invalid accesses can trigger compute-sanitizer warnings or, in rare instances, illegal memory access errors. This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [6040940] [5996751]

    • GEMM kernels on GPUs with compute capability 10.x or 12.x may access device alpha and beta pointers before calling cudaGridDependencySynchronize(). This can result in a WAR hazard if the preceding PDL kernel produces alpha and beta values on device after calling cudaTriggerProgrammaticLaunchCompletion(). This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [CUB-10409]

  • Resolved Issues

    • Fixed an issue in cublasLtMatmulAlgoGetHeuristic() that could result in no algorithm candidates being returned for Grouped GEMM on Blackwell GPUs. [CUB-9657]

3.1.7. cuBLAS: Release 13.2

  • New Features

    • Extended the experimental Grouped GEMM API in cuBLASLt to support MXFP8 inputs on GPUs with Compute Capability 10.x and 11.0.

    • Added control over special-case handling in FP32 emulation via the environment variable CUBLAS_EMULATION_SPECIAL_VALUES_SUPPORT_MASK. Setting CUBLAS_EMULATION_SPECIAL_VALUES_SUPPORT_MASK=0 can improve performance for applications that do not require preservation of infinity and NaN values, without requiring code changes. For more information, see the cudaEmulationSpecialValuesSupport_t documentation.

    • Added FP64 fixed-point emulation support to the cublas[D|Z]syrk, cublas[D|Z]syr2k, cublasZherk, and cublasZher2k routines. When the math mode is set to CUBLAS_FP64_EMULATED_FIXEDPOINT_MATH, cuBLAS will automatically use FP64 emulation for sufficiently large SYRK and HERK problems. Current support is limited to GPUs with Compute Capability 10.0.

    • Improved performance on RTX PRO 6000 GPUs, delivering up to 20% speedup for FP8, FP16/BF16, TF32, and INT8 precisions.

    • Improved GEMM performance on DGX Spark systems for MXFP8 and NVFP4 data types in large M and N problem sizes, with up to 3× performance improvement for selected matrix shapes.

  • Known Issues

    • On Blackwell GPUs, FP64 fixed-point emulation kernels may produce incorrect results or experience data corruption when executed concurrently with third-party kernels that allocate tensor memory. [CUB-9633]

    • Strided batched GEMM operations using broadcast operands may perform out-of-bounds memory reads on GPUs with compute capability 12.0 or 12.1 when an input matrix shared across batches (via zero or overlapping strides) is smaller than a few hundred KB. Numerical accuracy is unaffected, but these invalid accesses can trigger compute-sanitizer warnings or, in rare instances, illegal memory access errors. This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [6040940] [5996751]

    • GEMM kernels on GPUs with compute capability 10.x or 12.x may access device alpha and beta pointers before calling cudaGridDependencySynchronize(). This can result in a WAR hazard if the preceding PDL kernel produces alpha and beta values on device after calling cudaTriggerProgrammaticLaunchCompletion(). This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [CUB-10409]

  • Resolved Issues

    • Fixed an issue in cublasLtMatmul that could lead to incorrect results when it ran concurrently with another kernel that uses Tensor Memory. This issue only affected algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66 on GPUs with Compute Capability 10.x and 11.x, and existed since cuBLAS 12.8. [5807900]

    • Fixed an issue in cublasLtMatmul that could lead to incorrect results or invalid memory access errors for large leading dimensions, specifically when the product of the data type size and the leading dimension of a matrix exceeded the bounds of a signed 32 bit integer. This issue affected GPUs with Compute Capability 9.0, 10.x, or 11.0, existed since cuBLAS 12.6 Update 2, and only affected algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66. [CUB-9572]

    • Fixed an issue in the cuBLASLt Matmul API that could cause FP8 kernels to hang on GPUs with Compute Capability 9.0 when beta != 0 and scale_C = 0. This issue only affected algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66. [CUB-9627]

    • Fixed an issue in the cuBLASLt Grouped GEMM API that ignored groups with k = 0, leading to incorrect results. This issue existed since CUDA 13.1. [CUB-9529]

    • Fixed an issue in the cuBLASLt Matmul API that could cause incorrect results when C broadcasting was used (LDC = 0). [5845724]

    • Added missing checks for matrix pointer alignment in the cublasLtMatmul API. [CUB-9577] [CUB-9599] [CUB-9585]

    • Fixed an issue in cublasLtMatmul that could lead to incorrect results for NVFP4 precision on B300 and GB300 GPUs when the m dimension was not a multiple of 64. [CUB-9577]

    • Fixed an issue in cublasLtMatmul that could lead to incorrect results for NVFP4 precision on future GPUs, impacting future hardware compatibility. [CUB-9570]

    • Fixed an issue in GEMM and Matmul APIs with BF16 and FP16 inputs on DGX Spark and FP8 inputs on GeForce that could potentially cause illegal memory accesses. [5846563]

    • Fixed an issue in cuBLASLt to enable CUBLASLT_EPILOGUE_BGRADA and CUBLASLT_EPILOGUE_BGRADB epilogues when the C matrix CUBLASLT_MATRIX_LAYOUT_ORDER was set to CUBLASLT_ORDER_ROW. [4617436]

    • Fixed an integer overflow bug in complex, emulated FP64 matrix multiplication. The affected routines include cublasZgemm, cublasZtrsm, cublasGemmEx, and cublasLtMatmul. The overflow occurred when 2*m*n + m exceeded UINT_MAX, where m is the number of rows of op(A) and C, and n is the number of columns of op(B) and C. [5720478]

    • Improved GB200 and B200 performance for MXFP8 and NVFP4 precisions when M and N were less than or equal to 32. [CUB-9646]

3.1.8. cuBLAS: Release 13.1 Update 1

  • Known Issues

    • The cuBLASLt Grouped GEMM API ignores groups with k = 0, which can lead to incorrect results. As a workaround, initialize output matrices D with beta*C for all groups, and then compute Grouped GEMM as D += A*B so the result for groups with k = 0 is computed properly. This issue applies to the experimental cuBLASLt Grouped GEMM API introduced in CUDA 13.1. [CUB-9529]

    • Complex FP64 GEMM routines using fixed-point emulation can produce incorrect results when matrix dimensions are large enough that m*n > 2^31 due to integer overflow in an address calculation. [5720478]

    • cublasLtMatmul() may produce incorrect results when run concurrently with another kernel that uses Tensor Memory. This issue affects only algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66 on GPUs with compute capability 10.x and 11.x, and has existed since cuBLAS 12.8. [5807900]

    • Strided batched GEMM operations using broadcast operands may perform out-of-bounds memory reads on GPUs with compute capability 12.0 or 12.1 when an input matrix shared across batches (via zero or overlapping strides) is smaller than a few hundred KB. Numerical accuracy is unaffected, but these invalid accesses can trigger compute-sanitizer warnings or, in rare instances, illegal memory access errors. This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [6040940] [5996751]

    • cuBLAS GEMM kernels on GPUs with compute capability 10.x or 12.x may access device alpha and beta pointers before calling cudaGridDependencySynchronize(). This can result in a WAR hazard if the preceding PDL kernel produces alpha and beta values on device after calling cudaTriggerProgrammaticLaunchCompletion(). This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [CUB-10409]

  • Resolved Issues

    • Fixed an issue where fixed point emulation with 7 mantissa bits or less could trigger unspecified launch failures. [5692684]

    • Fixed an issue where cublasLtMatmul with FP8 arguments and CUBLASLT_MATMUL_MATRIX_SCALE_SCALAR_32F scaling mode (default) incorrectly required scaling factor addresses to be 16-byte aligned. This issue existed since cuBLAS 12.9. [5728938]

3.1.9. cuBLAS: Release 13.1

  • New Features

    • Introduced experimental support for grouped GEMM in cuBLASLt. Users can create a matrix with grouped layout using cublasLtGroupedMatrixLayoutCreate or cublasLtGroupedMatrixLayoutInit, where matrix shapes are passed as device arrays. cublasLtMatmul now accepts matrices with grouped layout, in which case matrices are passed as a device array of pointers, where each pointer is a separate matrix that represents a group with its own shapes. Initial support covers A/B types FP8 (E4M3/E5M2), FP16, and BF16, with C/D types FP16, BF16, and FP32; column-major only, default epilogue, 16-byte alignment; requires GPUs with compute capability 10.x or 11.0.

      In addition, the following experimental features were added as part of grouped GEMM:

      • Per-batch tensor-wide scaling for FP8 inputs, enabled by the new cublasLtMatmulDescAttributes_t entry CUBLASLT_MATMUL_MATRIX_SCALE_PER_BATCH_SCALAR_32F.

      • Per-batch device-side alpha and beta, enabled by the new cublasLtMatmulDescAttributes_t entries CUBLASLT_MATMUL_DESC_ALPHA_BATCH_STRIDE and CUBLASLT_MATMUL_DESC_BETA_BATCH_STRIDE.

    • Improved performance on NVIDIA DGX Spark for CFP32 GEMMs. [5514146]

    • Added sm_121 DriveOS support.

    • Improved performance on Blackwell (sm_100 and sm_103) via heuristics tuning for FP32 GEMMs whose shapes satisfy M, N >> K. [CUB-8572]

    • Improved performance of FP16, FP32, and CFP32 GEMMs on Blackwell Thor.

  • Resolved Issues

    • Fixed missing memory initialization in cublasCreate() that could result in emulation environment variables being ignored. [CUB-9302]

    • Removed unnecessary overhead related to loading kernels on GPUs with compute capability 10.3. [5547886]

    • Fixed FP8 matmuls potentially failing to launch on multi-device Blackwell GeForce systems. [CUB-9487]

    • Added stricter checks for in-place matmul to prevent invalid use cases (C == D is allowed if and only if Cdesc == Ddesc). As a side effect, users are no longer able to use D as a dummy pointer for C when using CUBLASLT_POINTER_MODE_DEVICE with beta = 0. However, a distinct dummy pointer may still be passed. The stricter checking was added in CUDA Toolkit 13.0 Update 2. [5471880]

    • Fixed cublasLtMatmul with INT8 inputs, INT32 accumulation, and INT32 outputs potentially returning CUBLAS_STATUS_NOT_SUPPORTED when dimension N is larger than 65,536 or when batch count is larger than 1. [5541380]

    • Added validation for batched matmul to reject invalid configurations where the batch counts differ (Adesc batch count != Bdesc batch count). [5645772]

  • Known Issues

    • The Grouped GEMM cuBLASLt API ignores groups with k = 0, which can lead to incorrect results. As a workaround, initialize each output matrix D with beta * C for all groups before the call, then compute Grouped GEMM as D += A * B so that the result for groups with k = 0 is preserved. This issue applies to the experimental Grouped GEMM cuBLASLt API released in CUDA 13.1. [CUB-9529]

    • cublasLtMatmul() may produce incorrect results when run concurrently with another kernel that uses Tensor Memory. This issue affects only algorithms with CUBLASLT_ALGO_CONFIG_ID equal to 66 on GPUs with compute capability 10.x and 11.x, and has existed since cuBLAS 12.8. [5807900]

    • Strided batched GEMM operations using broadcast operands may perform out-of-bounds memory reads on GPUs with compute capability 12.0 or 12.1 when an input matrix shared across batches (via zero or overlapping strides) is smaller than a few hundred KB. Numerical accuracy is unaffected, but these invalid accesses can trigger compute-sanitizer warnings or, in rare instances, illegal memory access errors. This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [6040940] [5996751]

    • cuBLAS GEMM kernels on GPUs with compute capability 10.x or 12.x may access device alpha and beta pointers before calling cudaGridDependencySynchronize(). This can result in a WAR hazard if the preceding PDL kernel produces alpha and beta values on device after calling cudaTriggerProgrammaticLaunchCompletion(). This issue was introduced in CUDA Toolkit 13.0 Update 2 (cuBLAS 13.1.0). [CUB-10409]

3.1.10. cuBLAS: Release 13.0 Update 2

  • New Features

    • Enabled opt-in fixed-point emulation for FP64 matmuls (D/ZGEMM) which improves performance and power-efficiency. The implementation follows the Ozaki-1 Scheme and leverages an automatic dynamic precision framework to ensure FP64-level accuracy. See here for more details on fixed-point emulation along with the table of supported compute-capabilities and the CUDA library samples for example usages.

    • Improved performance on NVIDIA DGX Spark for FP16/BF16 and FP8 GEMMs.

    • Added support for BF16x9 FP32 emulation to cublas[SC]syr[2]k and cublasCher[2]k routines. With the math mode set to