CUDA 13.5 makes Vera Rubin support production-ready, with Windows driver changes

Nvidia has released CUDA Toolkit 13.5, which, according to the company’s developer blog, moves support for the Vera Rubin GPU architecture from early developer access to general availability, “ready for enterprise production.”

The release also extends Windows on Arm support to DGX Station for Windows and adds new scheduling, communication and profiling features.

Rubin on Windows drops TCC mode

The change most likely to break existing deployments is a driver-model one. Starting with Rubin, CUDA supports the Windows Driver Display Model (WDDM) and the Microsoft Compute Driver Model (MCDM), the latter being the compute-only option. Nvidia says the Tesla Compute Cluster (TCC) mode is not supported on Rubin and later GPUs, so deployment tools that assume TCC need to target WDDM or MCDM instead.

On the scheduling side, locality domains get a new option. By default, localized green-context partitions use the same guaranteed SM count in every domain, which leaves some streaming multiprocessors unused. The new cudaDevSmResourceGroupLocalityDomainSupplementalSms flag, combined with cudaDevSmResourceGroupLocalityDomainId, pulls those extra SMs into each partition. The partitions can then differ in size, and applications can read each one’s SM count through discovery mode in cudaDevSmResourceSplit().

Multi-Process Service v3 adds weighted timeslicing, SM partition load balancing (borrowing is enabled with --max-chunks, lending with --shareable=true), cgroup v2 routing to MPS namespaces, per-namespace SM partition assignments and a polling fallback for memory accounting, switched on with nvidia-cuda-mps-control memacct set force_polling=true.

CUDA Graphs now accept batched 1D copies, so existing cudaMemcpyBatchAsync calls can be captured in stream capture without changing syntax.

Tile kernels can talk to other GPUs and call SIMT libraries

CUDA Tile gains multi-GPU communication inside kernels. Unicast pointers reach memory on NVLink-connected GPUs, while multicast pointers and multimem operations handle reductions and stores across NVSwitch fabrics with NVLink SHARP. Nvidia names fused matrix multiplication with cross-GPU reduction as an example use.

A public preview also lets cuTile Python and CUDA Tile C++ kernels call device functions from SIMT libraries mid-kernel. It is enabled with nvcc --tile-preview simt in C++ or cuda.tile_preview.foreign_call in Python. The libraries listed are NVSHMEM, CCCL Coop, Warp and MathDx, but Nvidia says each still has to enable the interop “soon”. Calling your own __device__ functions is planned for a later release.

Other pieces of the release:

• CCCL 3.5 adds public CUB tuning through cuda::execution::tune(...), run-to-run reproducible floating-point scans, and cub::DeviceBatchedTopK, which Nvidia says is integrated into FlashInfer.
• cuda.core 1.2 adds copy_batch(), Programmatic Dependent Launch through LaunchConfig, and graph updates without reinstantiation.
• Math libraries: Nvidia reports up to 1.5x faster FP64 symmetric updates on B200 in cuBLAS, and 1.65x to 2.27x faster CSR SpMM in cuSPARSE.
• Tools: a public preview of the Nsight Perf SDK for AI, Nsight Compute 2026.4 with a timeline page, and CUDA-GDB stops on driver log events.

What the release means for teams on Rubin hardware

The practical effect depends on where a team sits in the stack. Developers writing kernels for Rubin now have a toolkit Nvidia calls production-ready, and the cuBLAS changes (retuned NVFP4 and MXFP8 matmuls, including mixture-of-experts training shapes) are meant to help without per-shape autotuning. Those Rubin performance claims are Nvidia’s own, and the post gives no figures for them.

Operators who share GPUs are the other group affected. MPS v3’s borrowing and lending lets one partition use idle capacity while applications still see a stable base size, and cgroup routing ties it to launchers that already group processes. Windows shops with TCC-based tooling will have to rework it before moving to Rubin.

The surrounding software is not all caught up, though. In a Kubernetes post published two days earlier, Nvidia said its Vera Rubin recipe in AI Cluster Runtime v1.0.0 ships as Preview, not Supported. The same post said inference performance validation was deferred because vLLM’s published images do not yet include compiled Rubin kernels, which need CUDA 13.4 or later. It also reported a cross-node NCCL all-reduce reaching 744.71 GB/s of bus bandwidth at 16 GiB in a training recipe test. That points to a gap between toolkit availability and downstream frameworks shipping builds for it.

On the hardware side, Nvidia said in an October 27 press release that Vera Rubin NVL72 systems and Vera CPUs are generally available on Oracle Cloud Infrastructure. Customers renting Rubin there therefore have a matching GA toolkit. The CUDA post does not say which cloud providers or OEM systems have Rubin in volume, and it gives no benchmarks comparing 13.5 with 13.4 on Rubin.

Where to get it and what is still missing

The toolkit is available to download now, and Nvidia points readers to the CUDA 13.5 release notes for supported platforms and compatibility. Windows on Arm support has been generally available since CUDA 13.4. The Nsight Perf SDK for AI and the Tile-SIMT interop are previews, and cusparseSpMVGeneric and cusparseSpMatConvert are marked experimental.

Anyone who parses cuobjdump symbol dumps should check their scripts, since --dump-elf-symbols output now uses four descriptive columns (Type, Bind, Vis, Other) instead of raw hexadecimal fields. The new --precompile option in nvprune can compile PTX and TileIR to native SASS during packaging to avoid JIT overhead, but Nvidia advises validating the result on every target architecture.