NVIDIA CUDA 13.5 Brings NVIDIA Vera Rubin GA Support

As artificial intelligence and accelerated computing continue to push the boundaries of what’s possible, the latest NVIDIA CUDA release arrives to unlock unprecedented performance and scale.

NVIDIA CUDA 13.5 delivers general availability (GA) across major new platforms. NVIDIA Vera Rubin support moves from early developer access to general availability (GA), ready for enterprise production. Windows on Arm support, GA since CUDA 13.4, now extends to NVIDIA DGX Station for Windows.

CUDA 13.5 also introduces more customization of how GPU work is scheduled via locality domains and NVIDIA Multi-Process Service (MPS). And CUDA 13.5 provides innovations across libraries, compilers, NVIDIA CUDA Tile, and developer tools.

This post will showcase the highlights of CUDA 13.5.

CUDA 13.5 enhancements

Some notable enhancements that come with CUDA 13.5 include:

CUDA 13.5 gives applications a choice between uniform streaming multiprocessor (SM) partitions across locality domains or partitions that include all available SMs. By default, localized green-context partitions use the same guaranteed SM count per domain, leaving supplemental SMs outside those partitions.

Applications can use discovery mode in the cudaDevSmResourceSplit() API to obtain each partition’s SM count and distribute work accordingly. This allows you to choose between the uniformity and predictability of uniform SM size for each locality domain, or to choose the maximum number of SMs that each locality domain has, and utilize as much of the GPU hardware as possible. For more info see the CUDA Runtime API.

MPS v3 gives you more control over how workloads share a GPU, from scheduling individual servers to assigning resources to groups of clients.

Weighted timeslicing lets you give longer compute turns to servers hosting workloads that need them. Set a base timeslice, scale it per MPS server, and independently tune how aggressively the GPU switches between workloads.

SM partition load balancing lets workloads use capacity beyond their base partition when resources are available. Enable borrowing with --max-chunks, and allow a partition to lend its SMs while it has no attached clients with --shareable=true. These settings are independent, so a partition can borrow, lend, or do both. MPS reclaims borrowed resources before reassigning them, while applications continue to see a stable base partition size.

Cgroup-based routing connects your workload launcher’s process groupings to MPS resource policies. Map a cgroup v2 subtree to an MPS namespace, and clients connecting through the daemon’s main socket are routed to that namespace automatically.

Per-namespace SM partition assignments let you choose a partition for each device once, then apply that selection to every client routed to the namespace. You no longer need to configure CUDA_MPS_SM_PARTITION separately for those clients. Namespace assignments take precedence, while devices without an assignment retain the environment-variable fallback.

Polling-based memory accounting provides a fallback when cgroup threshold notifications are missing or unreliable. During eviction decisions, MPS can read current usage and limits directly instead of relying on cached notification state. To force polling, run nvidia-cuda-mps-control memacct set force_polling=true before memory accounting starts. This adds work per eviction decision but removes its dependence on notification delivery.

CUDA 13.5 adds CUDA Graph support for batched 1D memory copies. Applications can now include existing cudaMemcpyBatchAsync calls in stream capture without changing the batched-copy call syntax, allowing multiple transfers to be captured and replayed as part of a reusable graph workflow. Applications can also add batched-copy nodes explicitly and update compatible copy addresses and sizes without changing the batch structure.

Starting with the NVIDIA Rubin architecture, CUDA supports Windows Driver Display Model (WDDM) and Microsoft Compute Driver Model (MCDM), with MCDM providing the compute-only driver model. Tesla Compute Cluster (TCC) is not supported on Rubin and later GPUs. Applications and deployment tools that assume TCC mode should use the supported WDDM or MCDM configuration for their target system.

CUDA 13.5 adds support for multi-GPU communication directly within CUDA Tile kernels. Unicast pointers let a kernel access memory on another NVLink-connected GPU, while multicast pointers and dedicated multimem operations support reductions and stores across multiple GPUs on NVIDIA NVSwitch fabrics with NVIDIA NVLink SHARP support. Applications prepare the shared-memory mappings before launch, then combine communication with computation inside the kernel, enabling workloads such as fused matrix multiplication and cross-GPU reduction.

CUDA 13.5 allows CUDA Tile kernels to call SIMT library device functions directly. Before this release, Tile and SIMT library code could only be combined across kernel boundaries, with data passing through global memory between them. With this preview, kernels written in cuTile Python or CUDA Tile C++ can call device functions from SIMT libraries in the middle of their own code, on data that is already in registers or shared memory. These libraries include NVIDIA NVSHMEM, NVIDIA CCCL Coop, NVIDIA Warp and NVIDIA MathDx, which covers NVIDIA cuFFTDx, NVIDIA cuBLASDx, NVIDIA cuSolverDx, and NVIDIA nvCOMPDx, and this interop support will be enabled by each of these libraries soon.

A tile library (.tilelib) packages interface and implementation information for SIMT library functions. The compiler partitions each tile across threads the way the library expects, communicates any shared memory and synchronization requirements, and produces a single kernel. To enable the preview in C++, compile with nvcc --tile-preview simt and pass the tilelib to the tile compiler. In cuTile Python, use cuda.tile_preview.foreign_call, provided by the cuda.tile_preview package.

In 13.5, the Tile-SIMT interop preview covers library functions only. Calling your own __device__ functions from tile code is planned for a future release. For more information, see the documentation.

NVIDIA CUDA Python continues to expand Pythonic access to CUDA programming. The CUDA Python core library cuda.core 1.2 introduces batched memory copies, Programmatic Dependent Launch (PDL), simpler runtime compilation, and more flexible CUDA Graph updates.

The new cuda.core.utils.copy_batch() submits multiple buffer copies through a single CUDA driver call. CopyOptions provides control over source access ordering, memory-location hints, and overlap behavior, giving applications more control over data movement.

LaunchConfig now exposes PDL through the programmatic_stream_serialization option. On GPUs with compute capability 9.0 or newer, cooperating kernels can overlap execution within the same CUDA stream, allowing a dependent kernel to begin independent work before its predecessor completes. This gives Python applications access to inter-kernel overlap while retaining explicit control over launch configuration. See the LaunchConfig documentation.

Applications can update graph-definition nodes before future instantiations and modify executable graph nodes through graph[node] views without reinstantiating the graph. Supported updates include kernel parameters, memory copies, memory sets, child graphs, events, and host callbacks. These capabilities simplify reusing graphs as application parameters change.

CUDA Toolkit 13.5 brings CCCL 3.5, with public tuning APIs for CUB, reproducible floating-point scans, a new batch Top-K selection and reductions that consume problem sizes directly from GPU memory. These additions give developers more control over algorithm execution and simplify pipelines with data-dependent workloads.

CCCL 3.5 introduces public tuning policies for CUB device-wide algorithms. Developers can customize parameters such as threads per block, items per thread, vectorization, and algorithm selection through cuda::execution::tune(...) in an execution environment.

The following example shows how a policy selector chooses tuning parameters for a GPU compute capability. This selector returns cub::ReducePolicy for device-wide reductions; a scan selector would return cub::ScanPolicy. Applications that do not supply custom policies continue to use CUB’s default tuning. For examples, see the CUB tuning documentation.

Top-K selection is an important operation in AI workloads, including token selection during text generation and selecting relevant tokens for sparse attention. The new cub::DeviceBatchedTopK selects the smallest or largest K items independently across many segments, supporting keys and key-value pairs, variable segment sizes, and a K value that can vary by segment. NVIDIA has developed this algorithm with AI workloads in mind. CUB’s batched Top-K is integrated into FlashInfer, a library of GPU kernels for LLM inference.

cub::DeviceScan adds support for run-to-run reproducibility for floating-point summation. Now you can request a fixed reduction order through the execution environment, producing repeatable results on the same GPU with the same input and build configuration.

The following example requests reproducible results for an inclusive sum. See the cub::DeviceScan documentation for details.

CUB reductions can now consume an element count stored in device-accessible memory using cuda::args::deferred. A preceding kernel or selection operation can produce that count, and the reduction can read it in stream order without transferring it to the host. This also enables captured reductions to process a different element count on each CUDA Graph replay without updating or recapturing the graph.

See the cub::DeviceReduce documentation for details.

CUDA Toolkit 13.5 introduces the public preview of the NVIDIA Nsight Perf SDK for AI, bringing always-on profiling to long-running training and inference workloads. The SDK preserves selected GPU and CPU performance metrics and execution context from the original run, helping developers investigate intermittent performance issues that can be difficult to reproduce. Developers can explore the collected data in Perfetto. Learn more here.

NVIDIA Nsight Compute 2026.4 introduces a dedicated timeline page in NVIDIA Nsight Compute, making it easier to explore GPU workload performance over time. The page brings time-resolved profiling data into a central view to help developers identify workload phases, transient stalls, and shifts in hardware utilization that aggregate metrics can hide. Developers can switch between timelines and access integrated function stats analysis without leaving the page, keeping related performance insights together. Direct links from the details page simplify navigation to the relevant timeline, while built-in guidance helps users get started when a report contains no timeline data.

Figure 1. Nsight Compute Timeline page with function statistics

CUDA Toolkit 13.5 adds support in CUDA-GDB for stopping execution on CUDA Driver log events, helping developers investigate errors when they occur. Configurable rules filter events by thread, severity, or message pattern, so developers can focus on relevant diagnostics. When a matching event stops execution, developers can inspect the host-side call stack to trace the triggering code and identify the root cause.

This release also supports GPU core dump generation while CUPTI tracing is enabled. Developers can capture GPU state when an exception occurs and inspect the dump later using a supported debugger, without having to reproduce the failure with tracing disabled.

A new NVIDIA Compute Sanitizer agent skill helps AI coding agents choose the right diagnostics and interpret CUDA memory, race-condition, and synchronization errors.

NVIDIA Nsight Cloud now supports OpenTelemetry trace integration and profiling configuration for NVIDIA Dynamo deployments through Nsight Operator, helping developers investigate AI workload performance at scale.

Core CUDA math libraries in CUDA 13.5 now have performance improvements on Vera Rubin. cuBLAS extends Ozaki-II scheme to FP64 symmetric rank-k updates. Other new features and notable performance improvements include:

Faster GEMMs on NVIDIA Vera Rubin: Updated heuristics and kernel tuning improve out-of-the-box NVFP4 and MXFP8 matmul performance on Vera Rubin, including mixture-of-experts (MoE) LLM training shapes. Training and inference workloads get higher throughput without per-shape autotuning.

Ozaki-II for FP64 symmetric rank-k updates: The Ozaki-II scheme now supports {D,Z}SYR{K,2K,KX} and ZHER{K,2K,KX} on NVIDIA Ampere and newer GPUs. As a result, applications that exercise Cholesky factorizations and least-squares solvers built on cuSOLVER now get Tensor Core acceleration for these routines as well.

Faster native FP64 symmetric updates on NVIDIA Blackwell: SYR2K, HER2K, SYRKX, and HERKX in double and double-complex precision are up to 1.5x faster on NVIDIA B200. Scientific codes that use native FP64 symmetric updates, such as eigensolvers and matrix factorizations, speed up with no code changes.

Optimized BF16 kernels for unaligned shapes: BF16 matmuls whose dimensions aren’t multiples of eight now use the same optimized kernels as FP16 on NVIDIA Hopper and newer GPUs. Models with irregular layer sizes, including convolutions dispatched through NVIDIA cuDNN, no longer pay a penalty for choosing BF16 over FP16.

Windows on Arm support for Blackwell GPUs: cuBLAS adds functional support for Blackwell GPUs, including GB300 workstations, on Windows on Arm. Developers on Arm-based Windows systems can run cuBLAS-accelerated applications without switching platforms.

Faster FP32 and FP64 CSR SpMM with ALG3: Deliver 1.65x to 2.27x faster performance across data center GPUs, accelerating sparse matrix–dense matrix multiplication.

New SpSV ALG2 algorithm: Deliver 2.7x faster solve performance than ALG1, measured by geometric mean, accelerating workloads that repeatedly solve sparse triangular systems.

Faster SpSV analysis: Achieve a 2.79x geometric-mean speedup in the analysis phase, reducing setup time for sparse triangular solves.

Lower SpSV memory usage: Reduce the memory footprint of sparse triangular solves, leaving more GPU memory available for application data.

Expanded mixed-index support: Use 64-bit offsets with 32-bit indices across SpMV (CSR, SELL), SpMM (CSR), SpSV (SELL), and sparse-to-dense/dense-to-sparse conversions (CSR, CSC), enabling larger sparse matrices while keeping index memory usage lower than with fully 64-bit indexing.

Mixed-precision CSR SpMV: Combine FP16 or BF16 matrix values with FP32 input/output vectors and FP32 computation in cusparseSpMV, reducing matrix storage requirements while retaining FP32 accumulation.

New experimental cusparseSpMVGeneric API: Perform SpMV on symmetric matrices while storing only the upper or lower triangle, reducing storage requirements by avoiding duplicate off-diagonal entries.

New experimental cusparseSpMatConvert API: Convert between CSR and CSC formats, making it easier to adapt sparse matrix layouts to different operations.

Faster graph coloring with csrcolor: Achieve approximately 3x faster performance, measured by geometric mean, reducing time spent in sparse matrix coloring.

GPU Hybrid decode throughput: Decode throughput for NVJPEG_BACKEND_GPU_HYBRID backend improves by 10% to 600%, depending on the image size, and whether standard Huffman table is used.

CUDA 13.5 introduces a --precompile option in nvprune, allowing developers to compile intermediate device code (PTX, TileIR) to native SASS/ELF during the prune step.

This eliminates runtime JIT overhead for downstream consumers who deploy fat binaries to known target architectures (e.g. cloud providers, package distributors).

You should select targets that match your deployment environment and validate the resulting artifacts across the GPU architectures you intend to support.

CUDA 13.5 extends advanced controls file (ACF) support to the PTX Compiler and NVVM libraries. This allows compilation workflows built on these libraries to apply CompileIQ-generated compiler controls.

CUDA 13.5 expands PTX and device-intrinsic capabilities to help developers optimize data movement and mathematical operations.

Tiled computations frequently encounter tensor boundaries where part of a tile falls outside the tensor’s dimensions. Handling these ragged tiles efficiently is important for workloads such as MoE, matrix multiplication, and convolutions.

CUDA 13.5 extends the PTX tensormap.replace instruction so kernels can update a tensor map’s Tensor Memory Accelerator (TMA) out-of-bounds addressing mode directly. This gives developers control over addressing behavior when reusing tensor-map descriptors and can eliminate a performance-costly workaround for illegal-address failures in affected workloads.

CUDA 13.5 adds the __sin2pif, __cos2pif, and __sincos2pif device intrinsics for fast approximate single-precision trigonometric calculations with inputs expressed in cycles.

This capability is useful for phase-based computations that can use fast approximate trigonometric results.

Corresponding PTX .2pi variants of the sin and cos instructions avoid an additional angle-conversion operation for workloads that already use cycle-based phase.

CUDA 13.5 improves ELF symbol-table output in cuobjdump, making it easier to understand the functions and variables contained in CUDA binaries.

Table 1. Symbol attributes and CUDA-specific metadata

Previously, symbol information and CUDA-specific attributes appeared as raw hexadecimal fields that developers had to decode manually. The updated output presents descriptive values in four columns, following conventions familiar to users of readelf -s:

This makes it easier to distinguish kernel entry points from device functions and inspect attributes associated with global, constant, shared, and managed variables. When multiple CUDA attributes apply, the output lists their names together. Unrecognized attribute values remain visible as UNKNOWN(0xNN).

The enhancement applies to the existing --dump-elf and --dump-elf-symbols options. For example, developers can inspect a CUDA binary with:

No additional option is required to obtain the descriptive output. Existing information—including symbol names, values, sizes, and section indices—remains available.

If you use scripts to parse these symbol-table dumps, update the script parsers to accommodate the new columns and named attribute values.

Get started with CUDA Toolkit 13.5

CUDA Toolkit 13.5 provides GA support for NVIDIA Vera Rubin GPU architecture, updated GPU resource-management and communication capabilities, and enhancements across CUDA Python, CCCL, CUDA Tile, NVIDIA Nsight developer tools, and core math libraries.

Download CUDA Toolkit 13.5 and review the CUDA Toolkit 13.5 release notes for the complete list of features, supported platforms, and compatibility information.

Jackson Marusarz, Mahender Hari, Mridula Prakash, Daniel Rodriguez, and Sandeep Hiremath contributed to the work in this blog.