Skip to main content
Back to Blog
AI/MLProgramming LanguagesSystem Administration
6 September 202615 min readUpdated 10 September 2026

CUDA Toolkit 13.4 Adds Windows on Arm Support and More Precise Shared-GPU Management

CUDA Toolkit 13.4 extends NVIDIA's GPU development platform with Windows on Arm support, preview support for the NVIDIA Rubin architecture, and new controls for managing shared...

By Hardware Team

CUDA Toolkit 13.4 extends NVIDIA's GPU development platform with Windows on Arm support, preview support for the NVIDIA Rubin architecture, and new controls for managing shared GPU resources. The release also updates CUDA Python, CCCL, CUDA Tile, NVIDIA Nsight tools, Compute Sanitizer, and core math libraries.

CUDA Toolkit 13.4 enhancements

Preview support for NVIDIA Rubin

CUDA Toolkit 13.4 provides functional preview support for the NVIDIA Rubin architecture, identified by compute capability 107. Developers can begin porting applications before Rubin support reaches general availability in a future CUDA Toolkit release. Rubin is NVIDIA's next-generation GPU architecture.

Multi-Process Service V3

Multi-Process Service (MPS) V3 adds a modernized control layer for automating and managing shared GPU resources. The update includes:

  • A scriptable command-line interface
  • Named server instances
  • Namespaces for organizing concurrent workloads
  • TOML configuration support
  • Streaming multiprocessor (SM) partition controls
  • cgroup-integrated GPU memory limits

These features allow compute performance, memory boundaries, and execution priority to be defined programmatically. MPS V3 is designed to work in containerized environments while maintaining resource isolation for individual processes.

CUDA Compute Fabric Transport

CUDA Compute Fabric Transport (CFT) provides a transport-oriented method for moving data across NVIDIA NVLink fabric. Rather than mapping every remote GPU allocation into a process's virtual address space, applications can target named logical endpoints with an endpoint ID and offset. They can then issue asynchronous put, get, and reduction operations directly from the GPU.

This model reduces virtual-address pressure in large multi-GPU systems and supports both unicast and multicast communication. CFT also reports completion and error status, allowing applications to detect, retry, or reroute failed fabric transfers.

CFT is available only through the CUDA Driver API and is intended primarily for communication-library developers. Most application developers can use higher-level libraries such as NVIDIA NCCL or NVSHMEM instead.

Locality domains

CUDA 13.4 adds programmatic access to locality domains. A locality domain is a portion of a GPU containing streaming multiprocessors and device memory. Applications can allocate device memory in a locality domain and create a green context with SM resources in that same domain.

Placing computation near the memory it accesses can improve performance on devices with multiple locality domains.

Querying unified-memory location

New API support for querying unified-memory residency gives performance-sensitive libraries and runtimes direct information about where managed or system-allocated data resides.

Unified memory simplifies heterogeneous programming, but high-performance software may still need locality information to avoid unnecessary page migrations, remote memory access, or inefficient staging. Residency queries allow applications and libraries to make more informed decisions about computation scheduling and data movement. The relevant runtime API is cudaMemGetLocationInfo.

Separating the CUDA driver and toolkit

CUDA SDK installers no longer bundle the NVIDIA driver. The appropriate nvidia-open driver or cuda-toolkit packages must be installed separately through a preferred package manager.

Coherent Driver-based Memory Management

On NVIDIA coherent platforms, including NVIDIA Grace Hopper, NVIDIA Grace Blackwell, and NVIDIA Vera Rubin, the driver now defaults to Coherent Driver-based Memory Management (CDMM) instead of NUMA.

NUMA remains supported and can be selected with a kernel module parameter. Because this is a node-wide setting that requires a driver reload or reboot, the mode should be selected before upgrading.

Compilers and NVCC

Supported host compiler compatibility now includes GCC 16 and Clang 22. The new SM_107 architecture target enables compilation for Rubin GPUs.

CUDA Python

CUDA Python expands Python access to core CUDA APIs and high-performance algorithms. CUDA Toolkit 13.4 updates development tools, memory management, CUDA graph workflows, and application portability.

cuda.core 1.1.0

Following CUDA Python 1.0, cuda.core 1.1.0 expands the stable Pythonic CUDA API with:

  • Texture and surface programming
  • More control over managed memory
  • Improved CUDA graph integration
  • Complete type information for development tools and agents

Texture and surface programming

The new cuda.core.texture module provides Python APIs for CUDA texture and surface memory. OpaqueArray and MipmappedArray represent hardware-laid-out GPU allocations. TextureObject supports bindless, hardware-filtered kernel reads, while SurfaceObject supports typed kernel-side loads and stores.

from cuda.core import Device
from cuda.core.texture import (
    OpaqueArrayOptions,
    ResourceDescriptor,
    TextureObjectOptions,
)
from cuda.core.typing import ArrayFormatType, FilterModeType

dev = Device()
dev.set_current()
stream = dev.create_stream()

with dev.create_opaque_array(
    OpaqueArrayOptions(
        shape=(1024, 1024),
        format=ArrayFormatType.FLOAT32,
        num_channels=1,
    )
) as array:
    array.copy_from(image, stream=stream)

resource = ResourceDescriptor.from_opaque_array(array)
options = TextureObjectOptions(filter_mode=FilterModeType.LINEAR)

with dev.create_texture_object(
    resource=resource,
    options=options,
) as texture:
    # Pass texture.handle to a CUDA C++ kernel.
    run_kernel(texture.handle)

NUMA-aware managed memory

ManagedMemoryResource.allocate() now returns a ManagedBuffer with a property-based interface for CUDA memory advice. Applications can configure read-mostly data, preferred placement, and processor access.

The new Host type complements Device when specifying memory locations. It can represent any host memory, a specific NUMA node, or the NUMA node associated with the calling thread.

from cuda.core import Device, Host, ManagedMemoryResource
from cuda.core.utils import prefetch_batch

dev = Device()
dev.set_current()
stream = dev.create_stream()
mr = ManagedMemoryResource()

weights = mr.allocate(weights_nbytes, stream=stream)
output = mr.allocate(output_nbytes, stream=stream)

weights.read_mostly = True
weights.preferred_location = dev
weights.accessed_by.add(dev)

prefetch_batch(stream, [weights, output], dev)

## Launch GPU work, then move the result to host memory.
output.prefetch(Host(), stream=stream)
stream.sync()

Development and graph workflows

cuda.core 1.1 provides .pyi type stubs for every public API. IDEs and coding agents can use these stubs to access type information, function signatures, and return types.

GraphBuilder.graph_definition exposes a captured graph as a GraphDefinition. Developers can combine stream capture with explicit graph construction, including inspecting or extending a captured graph.

Other changes include device-specific NVLink enumeration, expanded green-context workqueue configuration, path-like inputs for Program and ObjectCode, and a public Buffer.size property. The release also strengthens IPC validation, free-threaded Python correctness, and CUDA process checkpoint restoration.

cuda.compute 1.1

cuda.compute provides Python access to NVIDIA CUDA Core Compute Libraries (CCCL) algorithms, including sort, scan, reduce, and transform operations.

cuda.compute 1.1 adds ahead-of-time compilation of algorithm objects for multiple GPU architectures, including on build systems without a GPU. ProxyArray and ProxyValue describe argument types without allocating device memory. serialize() creates an artifact that can be stored and deployed, while deserialize() restores the algorithm on the target system without recompiling it.

The following example compiles a reduction for sm_80 and sm_90 without requiring a GPU, then saves the result:

import numpy as np
from cuda.compute import (
    OpKind,
    ProxyArray,
    ProxyValue,
    make_reduce_into,
    serialize,
)

reducer = make_reduce_into(
    d_in=ProxyArray(np.int32),
    d_out=ProxyArray(np.int32),
    op=OpKind.PLUS,
    h_init=ProxyValue(np.int32),
    compute_capability=[80, 90],  # Build for sm_80 and sm_90.
)

with open("reduce.cclb", "wb") as file:
    file.write(serialize(reducer))

CCCL 3.4

CUDA Toolkit 13.4 includes CCCL 3.4, which adds a faster cub::DeviceScan implementation for NVIDIA Blackwell GPUs, single-call APIs across CUB device-wide algorithms, batched warp reductions, and parallel C++ Standard Library algorithms in cuda::std.

Faster device-wide scans on Blackwell

A new warp-specialized implementation of cub::DeviceScan for Blackwell uses the Tensor Memory Accelerator (TMA) to overlap memory movement and computation while reducing synchronization overhead.

In benchmark results on an NVIDIA Blackwell GPU, the new cub::DeviceScan::Sum implementation reaches up to 92% memory-bandwidth utilization across the tested data types, compared with around 50% in a previous implementation. It is optimized for large scan workloads and retains fallbacks for unsupported architectures, data types, iterators, and toolchains.

Single-call CUB APIs

CCCL 3.4 completes the rollout of environment-based, single-call overloads across CUB device-wide algorithms. Previously, applications generally called a CUB algorithm once to determine temporary-storage requirements, allocated that storage, and called the algorithm again to perform the operation.

The new overloads obtain temporary storage from a memory resource supplied through an execution environment:

auto device = cuda::devices[0];
auto stream = cuda::stream{device};
auto pool   = cuda::device_default_memory_pool(device);

auto env = cuda::std::execution::env{
    cuda::stream_ref{stream},
    pool
};

cub::DeviceReduce::Sum(d_input, d_output, num_items, env);

The traditional two-phase APIs remain available for applications that require explicit storage management.

Batched warp reductions

CCCL adds cub::WarpReduceBatched, a warp-wide collective for reducing multiple independent batches of values distributed across a warp. Processing the batches together reduces shuffle operations and increases the amount of useful work performed by each warp.

Parallel C++ Standard Library algorithms on the GPU

CUDA 13.4 introduces the C++ Standard Library parallel-algorithm model in cuda::std. Developers can invoke algorithms including copy_if, find_if, merge, reduce, transform, and scan operations with the cuda::execution::gpu execution policy.

#include <cuda/std/algorithm>
#include <cuda/std/execution>

struct is_positive
{
    __host__ __device__
    bool operator()(int value) const
    {
        return value > 0;
    }
};

cuda::std::copy_if(
    cuda::execution::gpu,
    d_first,
    d_last,
    d_output,
    is_positive{}
);

The algorithms operate on device-accessible ranges and use CCCL and CUB implementations. They provide a standard interface for GPU execution while retaining CUDA-specific features such as streams and memory resources through customizable execution policies.

Programmatic Dependent Launch in CUDA Tile IR

Support for Programmatic Dependent Launch (PDL) in CUDA Tile IR enables inter-kernel overlap on the same CUDA stream. A dependent kernel can begin execution before its predecessor has completed.

New views in CUDA Tile C++

CUDA Tile C++ adds views for loading and storing data:

  1. Strided view: Creates statically sized data chunks whose spacing is determined by a compile-time stride. This supports access patterns commonly found in stencil-like operations.
  2. Gather scatter view: Provides access to non-adjacent array chunks, supporting sparse access patterns.

Developer tools

Nsight Python

Nsight Python 1.0 is a Python kernel-profiling interface for automating performance analysis across multiple kernel configurations with NVIDIA Nsight Tools. A decorator and context manager support kernel benchmarking, architectural metric collection, GPU-throttling prevention, and performance visualization in a single script.

NVIDIA Nsight Compute

Nsight Compute 2026.3 adds Tile IR support for CUDA Tile workloads. Developers can inspect Tile IR on the source page and correlate it with CUDA Tile source and generated code. The release also improves register-spill information for OptiX workloads and enhances Nsight Copilot.

NVIDIA Nsight Systems

Nsight Systems 2026.5.1 expands platform coverage and workload visibility across CUDA, CPUs, AI frameworks, networking, and storage. The web release supports CUDA 13.4, Rubin GPUs, and Windows on Arm. It also projects NVTX ranges into an "All Streams" hierarchy, demangles cuTile names, and displays CUDA workloads submitted through CiG streams on the timeline.

CPU metric sets group related hardware counters for collection in a single pass. This supports progressive bottleneck analysis through Topdown metric sets. For PyTorch workloads, the --pytorch=functions-trace-shapes option adds information such as tensor shapes and training parameters to traced functions. Developers can choose between the additional detail of shape tracing and the lower overhead of the existing function-tracing option.

Network, storage, and cluster profiling

Network profiling adds high-frequency NIC metric collection through the NVIDIA DOCA Telemetry Service. Developers can correlate traffic, congestion notifications, and send waits with application activity without requiring elevated privileges.

A new NCCL straggler analysis recipe examines collective timing to identify ranks that repeatedly delay communicator progress.

Storage profiling adds an S3 access summary analysis recipe that aggregates access patterns and I/O statistics across processes and hosts. The analysis can identify hot buckets and objects, frequent small transfers, and workload imbalances.

NVIDIA SCADA metrics profiling adds counters and histograms from the SCaled Accelerated Data Access storage architecture to the timeline, allowing storage-server activity to be correlated with GPU and CPU events.

For multi-node and cluster analysis, the experimental vClock plugin improves report alignment without changing system clocks or requiring privileged access when high-precision synchronization, such as PTP, is unavailable.

Binary payloads and plugin development

NVTX binary payloads can be exported as dynamic relational tables, with payload fields represented as columns for downstream analysis in SQLite, Arrow, and Arrow/Parquet formats.

The Nsight Systems plugin framework adds an initialization stage, a process-exit callback API, and plugin-library loading in subprocesses. These features allow collection to begin before an application starts, capture data generated during shutdown, and extend profiling across child processes.

NVIDIA Nsight Cloud

Nsight Cloud supports viewing and analyzing profiling reports on remote, headless systems. Nsight Operator adds improvements for analysis, OpenTelemetry, and NVIDIA Dynamo, along with a new documentation site.

NVIDIA Nsight AI

Nsight AI adds specialized AI assistance to accelerated-computing workflows. The NVIDIA-hosted CUDA MCP Server connects supported AI coding agents to current CUDA documentation and code examples. The open-source Nsight Copilot Blueprint provides a self-hosted CUDA AI backend for teams that operate it in their own environment.

NVIDIA Compute Sanitizer

Compute Sanitizer improves shared-memory out-of-bounds detection through compile-time patching on Hopper and newer architectures. Initcheck adds Batched memcpy async support, while racecheck adds support for per-cluster-block filtering.

NVIDIA Core Math Libraries

Core math libraries in CUDA Toolkit 13.4 add functional support for the Rubin GPU architecture and Windows on Arm support for the N1X Laptop ecosystem.

Updates to cuBLAS include:

  • Improved double-precision performance through fixed-point emulation using the Ozaki-II scheme.
  • Dynamic scheduling of Grouped GEMM computations across streaming multiprocessors on Blackwell data center GPUs. The scheduling is intended to reduce load imbalance in common MoE workloads and can improve performance for calls with many groups, such as 32 groups, as well as concurrent Grouped GEMM operations.
  • Experimental cuBLASLt scaling modes, CUBLASLT_MATMUL_MATRIX_SCALE_VEC32_MN_K4_UE8M0 and CUBLASLT_MATMUL_MATRIX_SCALE_VEC128_MN_K4_UE8M0, for an alternative scaling-factor layout supporting A and B FP8-precision tensors. The modes group scaling factors in sets of four and store them in M-major or N-major layouts. Padding is added when the major dimension is not divisible by four.

CUDA Toolkit 13.4 overview

CUDA Toolkit 13.4 adds Windows on Arm support, preview support for NVIDIA Rubin, more detailed GPU resource-management and communication controls, and updates across CUDA Python, CCCL, CUDA Tile, NVIDIA Nsight tools, Compute Sanitizer, and core math libraries. The release also includes support and compatibility changes documented in the CUDA Toolkit 13.4 release notes.