NVIDIA’s CUDA 13.5 takes Vera Rubin to general availability without TCC mode

CUDA 13.5 brings Vera Rubin to general availability and adds MPS v3 sharing controls. On Windows, Rubin supports WDDM and MCDM but not TCC mode.

NVIDIA says CUDA Toolkit 13.5 moves Vera Rubin support from early developer access to general availability, ready for enterprise production. In its developer blog post from November 4, the company adds that Windows on Arm support, generally available since CUDA 13.4, now extends to DGX Station for Windows.

On Windows, Rubin runs under WDDM or MCDM and TCC mode is not supported

The change most likely to break an existing setup is a driver-model one. Starting with Rubin, CUDA supports the Windows Driver Display Model (WDDM) and Microsoft’s Compute Driver Model (MCDM), which is the compute-only option. Tesla Compute Cluster (TCC) mode is not supported on Rubin and later GPUs, so NVIDIA tells developers whose applications or deployment tools assume TCC to switch to WDDM or MCDM for their target system.

MPS v3 adds weighted timeslicing and lets idle GPU partitions lend their cores to busy ones

The scheduling work sits in Multi-Process Service (MPS) v3, which controls how several workloads share one GPU. Weighted timeslicing sets a base timeslice and scales it per MPS server, so a server hosting a heavier workload gets longer compute turns. Load balancing lets a workload use streaming multiprocessors (SMs) beyond its base partition when they are free: borrowing is switched on with --max-chunks, and a partition with no attached clients can lend its SMs with --shareable=true. MPS reclaims borrowed SMs before reassigning them, so applications keep seeing a stable base partition size. A third addition maps a cgroup v2 subtree to an MPS namespace, which routes clients connecting through the daemon’s main socket automatically.

Locality domains get a related option. By default each localized partition has the same guaranteed SM count, which leaves the extra SMs outside any partition. A new flag includes them, at the cost of partitions of unequal size, and discovery mode in cudaDevSmResourceSplit() reports each partition’s SM count so work can be divided to match.

CUDA Tile kernels can now reach other NVLink GPUs and call SIMT library functions, the latter as a preview

CUDA Tile kernels can use unicast pointers to read memory on another NVLink-connected GPU, and multicast pointers with dedicated multimem operations for reductions and stores across GPUs on NVSwitch fabrics with NVLink SHARP. The preview of Tile and SIMT interoperability removes a round trip: before 13.5 the two styles could only be combined across kernel boundaries, with data passing through global memory, and now a Tile kernel can call SIMT library functions mid-kernel on data already in registers or shared memory. The preview covers library functions such as those in NVSHMEM and CCCL Coop only, and calling a developer’s own __device__ functions is planned for a later release.

The post quotes its speedups for data center parts and does not mention GeForce or RTX cards

NVIDIA says NVFP4 and MXFP8 matrix multiplication is faster on Vera Rubin, without giving a figure, that FP64 symmetric updates run up to 1.5 times faster on the B200, and that sparse CSR matrix products run 1.65 to 2.27 times faster across data center GPUs. None of the post’s numbers are for a card in a desktop. For someone running CUDA on a GeForce GPU, the release notes may say more about what 13.5 changes than this announcement does.