NVIDIA CUDA 13.5 Brings NVIDIA Vera Rubin GA Support NVIDIA released CUDA 13.5, moving NVIDIA Vera Rubin support from early developer access to general availability for enterprise production and extending Windows on Arm support, GA since CUDA 13.4, to NVIDIA DGX Station for Windows. The toolkit adds locality domain controls, including the cudaDevSmResourceGroupLocalityDomainSupplementalSms flag with cudaDevSmResourceGroupLocalityDomainId and a discovery mode in the cudaDevSmResourceSplit() API, letting applications use all device SMs across partitions. CUDA 13.5 also ships MPS v3 with weighted timeslicing, SM partition load balancing via --max-chunks and --shareable=true, cgroup-based routing, and per-namespace SM partition assignments. 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 https://developer.nvidia.com/blog/cuda-toolkit-13-4-adds-windows-on-arm-support-and-greater-control-over-shared-gpus/ , 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 https://developer.nvidia.com/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: Locality domain enhancements 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. The new cudaDevSmResourceGroupLocalityDomainSupplementalSms flag, combined with cudaDevSmResourceGroupLocalityDomainId , includes these additional SMs within each selected domain. This enables applications to use all device SMs when partitioning across locality domains, although the resulting partitions may differ in size. 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 https://docs.nvidia.com/cuda/cuda-runtime-api/cuda runtime api/group CUDART EXECUTION CONTEXT.html?highlight=cudadevsmresourcesplit CPPv422cudaDevSmResourceSplitP15cudaDevResourcejPK15cudaDevResourceP15cudaDevResourcejP28cudaDevSmResourceGroupParams . NVIDIA MPS 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 Weighted timeslicing https://docs.nvidia.com/deploy/mps/latest/mpsv3-interface.html 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 SM partition load balancing https://docs.nvidia.com/deploy/mps/latest/mpsv3-interface.html 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 Cgroup-based routing https://docs.nvidia.com/deploy/mps/latest/mpsv3-namespaces.html cgroup-matching 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. SM partition assignments Per-namespace SM partition assignments https://docs.nvidia.com/deploy/mps/latest/mpsv3-namespaces.html 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 fallback for memory accounting Polling-based memory accounting https://docs.nvidia.com/deploy/mps/latest/mpsv3-memory-partitioning.html detection-modes-and-polling 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. Batched 1d copies in CUDA graphs 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. Windows driver modes on NVIDIA Rubin 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 Tile updates Multi-gpu communication in CUDA Tile CUDA 13.5 adds support for multi-GPU communication directly within CUDA Tile kernels https://docs.nvidia.com/cuda/tile-ir/latest/sections/prog model.html multi-gpu-communication . 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. Tile to single instruction multi-thread SIMT interoperability public preview 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 https://docs.nvidia.com/nvshmem/api/latest/index.html , NVIDIA CCCL Coop https://nvidia.github.io/cccl/unstable/python/coop.html , NVIDIA Warp https://nvidia.github.io/warp/stable/ and NVIDIA MathDx https://docs.nvidia.com/cuda/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 https://docs.nvidia.com/cuda/tile-ir/latest/sections/prog model.html simt-interoperability . NVIDIA CUDA Python 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. Batched memory copies 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. The following example submits two independent copies using existing source and destination buffers: python from cuda.core.utils import copy batch Each source and destination pair must have matching sizes. Copies within the batch must not alias. copy batch stream, srcs= src a, src b , dsts= dst a, dst b , stream.sync PDL from Python 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 https://nvidia.github.io/cuda-python/cuda-core/1.2.0/generated/cuda.core.LaunchConfig.html . python from cuda.core import LaunchConfig, launch config = LaunchConfig grid=num blocks, block=256, programmatic stream serialization=True, The kernel must implement the PDL synchronization protocol. launch stream, config, dependent kernel, kernel args More flexible CUDA Graph updates 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. python from cuda.core import LaunchConfig from cuda.core.graph import GraphDefinition Define a graph with one kernel node. graph def = GraphDefinition node = graph def.launch LaunchConfig grid=1, block=256 , kernel, output a, Instantiate once and run with the original output buffer. graph = graph def.instantiate graph.launch stream stream.sync Update the executable node's arguments without reinstantiating. graph node .update args= output b, The same graph now launches the kernel with output b. graph.launch stream stream.sync CCCL 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. Customize CUB performance with public tuning APIs 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 https://nvidia.github.io/cccl/unstable/cub/tuning.html . struct ReduceTuning { host device constexpr cub::ReducePolicy operator cuda::compute capability cc const { auto pass = cub::ReducePassPolicy{ .threads per block = 256, // Use 20 items per thread for compute capability 10.0 // and newer, and 16 for older GPUs. .items per thread = cc = cuda::compute capability{10, 0} ? 20 : 16, .vec size = 4, .reduce algorithm = cub::BLOCK REDUCE WARP REDUCTIONS, .load modifier = cub::LOAD LDG }; // Provide tuning for both multi-tile and single-tile reductions. return {.multi tile = pass, .single tile = pass}; } }; // Attach the selector to the execution environment using the public tuning API. auto env = cuda::std::execution::env{ cuda::stream ref{stream}, cuda::execution::tune ReduceTuning{} }; // Keep using the public device-wide API. CUB applies the custom policy // supplied through the environment. auto status = cub::DeviceReduce::Sum d input, d output, num items, env ; Batched Top-K selection 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. constexpr int segment size = 4; constexpr int k = 2; auto segments in = cuda::make strided iterator cuda::make counting iterator d input , segment size ; auto segments out = cuda::make strided iterator cuda::make counting iterator d output , k ; auto env = cuda::std::execution::env{ ... }; auto status = cub::DeviceBatchedTopK::MaxKeys segments in, segments out, cuda::args::constant