CUDA

CUDA (Compute Unified Device Architecture) is NVIDIA's platform and programming model for general-purpose computation on its graphics processing units. NVIDIA introduced CUDA in 2006 so developers could use GPU throughput without expressing computation through a graphics API.^[1] The platform includes language extensions and APIs, a compiler toolchain, runtime and driver interfaces, development tools, and optimized libraries. Software can use these facilities directly, but many applications reach CUDA through a library or a higher-level framework instead of containing custom GPU kernels.^[1]

CUDA is designed for heterogeneous systems in which host code runs on a CPU and device code runs on one or more GPUs. Its central abstraction is the kernel, a function launched across many GPU threads. CUDA organizes those threads into blocks and grids, provides several memory spaces, and supplies synchronization and asynchronous-execution mechanisms.^[2] CUDA is therefore broader than either the CUDA C++ language or the physical arithmetic units sometimes marketed as "CUDA cores." It is also distinct from any particular NVIDIA GPU generation.

NVIDIA's current release notes cover CUDA Toolkit 13.4, announced in September 2026; its general-availability build is versioned 13.4.1 and supersedes an earlier 13.4.0 developer preview.^[3][29] Since CUDA 11, toolkit components have carried independent version numbers, so the version of a bundled library or tool does not necessarily match the toolkit's release label.^[3]

History

Programmable graphics processors were used for non-graphics work before CUDA, but early programs generally had to map computations onto graphics concepts. Brook for GPUs, presented by Ian Buck and coauthors in 2004, was one attempt to provide a more direct model. Brook extended C with stream data types and operations, and its compiler and runtime mapped those operations to contemporary GPU hardware.^[4] It is an important predecessor in GPU computing, although CUDA developed into a separate hardware-software platform rather than simply adopting Brook unchanged.

NVIDIA introduced CUDA in 2006. The change gave developers a general-purpose programming interface that did not require graphics APIs, while retaining a throughput-oriented execution model suited to large numbers of similar operations.^[1] John Nickolls, Ian Buck, Michael Garland, and Kevin Skadron described the model in a 2008 paper, emphasizing a hierarchy of thread groups, shared memory, synchronization, and decomposition into independently schedulable blocks.^[5]

CUDA later became one of the major software paths for scientific computing and machine learning on NVIDIA GPUs. The 2012 AlexNet paper, for example, described an efficient GPU implementation used to train a large convolutional network for the ImageNet classification task.^[6] That result illustrates the importance of programmable GPU computing to deep learning, but it does not imply that every machine-learning system or accelerator uses CUDA.

The platform has expanded while retaining the host, device, kernel, grid, and block concepts. Later additions include unified memory facilities, graphs, cooperative groups, cluster-level features on supported devices, and a tile programming model documented in CUDA 13.3.^[2] These additions complement the traditional per-thread model and are not all available on every GPU or in every toolkit release.

Programming Model

Hosts, devices, and kernels

A CUDA application begins execution on a host CPU. The host program can allocate or register memory, move or map data, select devices, launch work on a GPU, and wait for results. Code executed by a GPU is device code. A device function invoked from the host as parallel work is called a kernel.^[2]

In CUDA C++, a conventional kernel is declared with the __global__ qualifier and launched with an execution configuration. A simplified vector-addition launch looks like this:

add_vectors<<<grid_size, block_size>>>(a, b, output, count);

The launch configuration specifies how many thread blocks the grid contains and how many threads each block contains. The launch is normally asynchronous with respect to the host, so host execution may continue before the kernel finishes.^[7] Error checking and synchronization must be placed deliberately rather than inferred from the source-code line order.

CUDA does not require every application developer to write kernels. Libraries can expose GPU implementations behind ordinary function calls, and frameworks can construct and schedule CUDA work on behalf of the user.^[1] Direct kernel programming is most useful when an application needs an operation that existing libraries do not provide, needs specialized fusion or data movement, or requires control over execution details.

Grids, blocks, and threads

A conventional kernel launch creates a grid of thread blocks. Grids and blocks may be one-, two-, or three-dimensional, which helps map work to vectors, images, volumes, or other structured data. Each thread can determine its position from built-in block and thread indices.^[2]

Threads in one block execute on a single streaming multiprocessor (SM). They can exchange data through block-scoped shared memory and synchronize with block-scoped barriers. A GPU normally has far fewer SMs than a large grid has blocks, so blocks are assigned to SMs as resources become available.^[2]

The ordinary grid model requires blocks to be independently schedulable. CUDA does not guarantee the order in which blocks run, and a block generally cannot assume that another block has completed or is resident at the same time. This property lets the same grid run on GPUs with different numbers of SMs. It also means that an algorithm needing a global dependency commonly uses separate kernel launches, or an explicitly supported cooperative mechanism, rather than an ordinary in-kernel block barrier.^[2]

CUDA-capable devices have limits on block dimensions, threads per block, registers, shared memory, and other resources. Compute capability identifies an architectural feature level and helps determine such limits and supported instructions.^[8] It is not the same as a toolkit version: a toolkit can compile for several compute capabilities, while a GPU has a fixed compute capability.

Warps and SIMT execution

Within a block, CUDA groups threads into warps of 32. A warp follows the single-instruction, multiple-thread (SIMT) model: its threads execute the same kernel but retain their own registers, addresses, and logical control flow.^[2] The programming model permits threads to follow different branches. When lanes in a warp take different paths, some lanes are masked while the relevant path executes. This condition is called warp divergence and can reduce utilization.

SIMT is related to, but not identical with, fixed-width SIMD programming. CUDA exposes threads as separate logical entities, while the hardware issues work in warp-sized groups. Correct code should follow the documented memory and synchronization model instead of assuming undocumented scheduling behavior.

Warp awareness matters for performance. Adjacent threads often work best when they access memory in patterns the hardware can combine into efficient transactions. Block sizes that are multiples of the warp size also avoid a final warp with permanently unused lanes, although block size must be chosen together with register use, shared-memory use, and the amount of parallel work.^[2]

Thread-block clusters and tile programming

GPUs with compute capability 9.0 or later can optionally group blocks into thread-block clusters. Blocks in a cluster are scheduled within one graphics processing cluster and can use documented facilities for cluster-scoped synchronization and distributed shared memory.^[2] This is a qualified exception to the simpler rule that blocks in an ordinary grid cannot communicate directly. Code using clusters must check hardware and launch constraints rather than assuming the feature exists everywhere.

CUDA 13.3 also documents a tile programming model. A tile kernel expresses operations on multidimensional collections of values at block scope, and the compiler maps those operations to threads. SIMT kernels and tile kernels can coexist in one application and operate on the same device memory.^[2] Tile programming raises the abstraction level for suitable computations; it does not replace SIMT when per-thread control is needed.

Memory Model

CUDA exposes memory spaces with different locations, scopes, and performance characteristics. The exact capacities and cache arrangements vary by GPU architecture, so portable code queries limits instead of embedding values from a single product.^[2]

The name "local memory" refers to per-thread addressability, not necessarily fast on-chip storage. Register pressure can cause values to be placed in local memory, which is physically backed by device memory and may be much slower than a register. Shared memory is on-chip and explicitly managed by a block, but access patterns can still matter because of banking and synchronization.^[9]

Global-memory performance depends on access locality and transaction efficiency. When neighboring threads access suitably aligned neighboring addresses, the hardware can coalesce their requests. Scattered or redundant access may require more memory transactions. Caches can help some patterns, but they do not remove the need to understand data layout.^[9]

Host, device, managed, and system memory

Discrete systems commonly have system memory attached to the CPU and device memory attached to each GPU. Data may travel over PCIe or another supported interconnect. Some systems integrate CPU and GPU memory more closely. CUDA therefore offers more than one allocation and access model.^[10]

Explicit device allocation and copying gives the program direct control over placement and transfer timing. Pinned host memory can support asynchronous transfers and direct device access in appropriate configurations, but pinning consumes operating-system resources and should be used selectively. Mapped memory and zero-copy access can be useful when access is limited or transfer avoidance matters, but are not automatically faster than staging data in device memory.^[9]

Unified virtual addressing provides a common virtual address space for supported host and device allocations. Unified Memory, commonly allocated with cudaMallocManaged, lets the runtime and driver manage accessibility and movement of managed pages. System-allocated memory can also be directly accessible in some supported systems.^[10] These facilities simplify some programs, but they do not make all memory physically uniform. Page migration, faulting, oversubscription, processor access patterns, operating system, GPU architecture, and interconnect can materially affect behavior. Prefetching and placement advice are performance tools, not universal requirements.

Memory correctness also depends on synchronization and visibility. A pointer being valid in more than one processor's address space does not mean concurrent unsynchronized accesses are safe. Programs must use the documented CUDA memory model, stream dependencies, atomics, barriers, or system-level synchronization appropriate to the sharing pattern.

Asynchronous Execution and Coordination

A CUDA stream is an ordered queue of operations such as kernel launches, memory operations, and event commands. Operations in one stream follow stream order. Operations in different streams may overlap when dependencies, resource availability, memory type, and hardware capabilities permit.^[7] Multiple streams therefore make concurrency possible, but do not guarantee that two operations will execute simultaneously.

Events record progress in a stream. Other streams can wait on an event, and the host can query or wait for it. Events are also commonly used for device-side timing because a host timer alone may measure submission overhead rather than completion of asynchronous work.^[7]

Synchronization exists at several scopes:

• A block barrier coordinates participating threads within a block.
• Warp-level primitives coordinate lanes under their documented participation rules.
• Cooperative Groups provides explicit group abstractions, including supported grid- and cluster-level patterns.
• Stream order and events express dependencies between queued operations.
• Device or stream synchronization makes the host wait for specified work.
• Separate kernel launches provide a natural grid-wide phase boundary when issued with the required dependency.

Choosing an unnecessarily broad synchronization point can prevent useful overlap. Choosing one that is too narrow can produce races or stale reads. Correctness comes first; concurrency should be introduced only after dependencies are explicit and tested.

CUDA Graphs provide another submission mechanism. A graph records operations and dependencies so an application can instantiate and launch the graph repeatedly. Graphs can reduce repeated launch-management overhead in workloads with a stable dependency structure, but graph construction, update rules, and supported nodes add their own constraints.^[11]

Compilation and APIs

NVCC, PTX, and native code

The nvcc compiler driver separates host code from CUDA device code and coordinates a supported host compiler with NVIDIA's device compilation tools. Device code can be represented as Parallel Thread Execution (PTX), a virtual instruction set, or as native binary code for a target GPU architecture, commonly called a cubin.^[12]

An application can embed native code for several compute capabilities plus PTX for a virtual architecture. When native code for the installed GPU is available, the driver can load it directly. When compatible PTX is present, the driver can just-in-time compile it for a later GPU architecture. Keeping PTX can extend hardware forward compatibility, but it may add startup compilation time and requires a driver new enough to understand that PTX version.^[12][13]

Compilation targets are an engineering choice. Shipping only code for one SM target narrows the supported hardware. Shipping many native targets increases artifact size. Shipping PTX provides another compatibility path but is subject to driver and feature constraints. Build systems should select targets from the actual deployment fleet rather than assuming one flag is best for every application.

Runtime and driver APIs

The CUDA Runtime API provides device management, memory, streams, events, kernel launch support, and related services. It normally manages context initialization automatically. The lower-level Driver API exposes explicit contexts, modules, and other controls. NVIDIA documents the runtime as implemented on top of the driver API and notes that most applications can reach full performance without calling the driver API directly.^[14]

The two APIs can interoperate. An application may use a driver-created context with runtime calls, allocate memory through either interface, and invoke runtime-based libraries from driver-API code.^[14] Direct driver use is valuable for language runtimes, just-in-time compilers, plugin systems, and applications that need explicit module or context management. It is not intrinsically faster merely because it is lower level.

CUDA also has APIs and language bindings beyond CUDA C++. NVIDIA documents CUDA Python facilities, and other projects provide bindings or compilation paths from Python, Fortran, and domain-specific languages. These interfaces differ in coverage and release cadence, so their own documentation is authoritative for supported features.

Toolkit, Libraries, and Development Tools

The CUDA Toolkit combines compiler and runtime components with headers, libraries, command-line utilities, samples, and debugging or profiling tools. A display or data-center driver is a separate system component even when an installer can offer it alongside a toolkit. Release notes and platform installation guides specify supported operating systems, host compilers, drivers, and package methods for each release.^[3][15][16]

The toolkit's general-purpose libraries cover common numerical and data-processing operations:

These libraries let applications use architecture-tuned implementations without writing each kernel from scratch.^[17] They are separate components with their own APIs, supported configurations, numerical behavior, and version histories.

For deep learning, cuDNN provides GPU implementations of operations including attention, convolution, matrix multiplication, normalization, pooling, pointwise operations, and multi-operation fusion.^[18] NCCL provides topology-aware collective communication primitives for multiple GPUs and machines. It handles communication operations such as aggregation and data exchange, not model definition, scheduling, or the rest of a distributed training system.^[19]

Nsight Systems traces CPU and GPU activity at system scope, while Nsight Compute analyzes CUDA kernels using hardware performance metrics.^[26][27] Compute Sanitizer checks classes of memory, race, initialization, and synchronization errors.^[28] Tool availability and feature coverage vary by host platform and release, so developers should consult the matching tool documentation rather than assuming every utility is bundled identically.

CUDA Toolkit 13.4

CUDA Toolkit 13.4 was announced on September 9, 2026. Its general-availability build is versioned 13.4.1 and supersedes a 13.4.0 developer preview; the third digit is a build number rather than an update release.^[3] The release covers a new host platform, early enablement for an unreleased GPU architecture, new controls for dividing one GPU between workloads, and separately versioned updates to the Python bindings, the C++ core libraries, the compilers and the profiling tools.^[29]

Windows on Arm and the Rubin preview

CUDA applications have run on Arm CPUs for years through Linux, and 13.4 extends application development to Windows on Arm. The release notes record the change as support for Windows on Arm on RTX Spark devices, and the component table marks most toolkit libraries as shipping for arm64 (Windows) alongside x86_64 and arm64-sbsa.^[3][29] NVIDIA's core math libraries add Windows on Arm support for what the company calls the N1X laptop ecosystem.^[29]

The release also adds functional support for the NVIDIA Rubin GPU architecture as a preview, so that developers can begin porting applications before CUDA support for Rubin reaches general availability in a later toolkit.^[29] NVCC gains an SM_107 architecture target, and the release notes identify Rubin as compute capability 10.7.^[3] The compute-capability appendix places 10.7 in the same family as compute capabilities 10.0 and 10.3 for the family-specific compute_100f target.^[8] The announcement blog writes the same figure as "compute capability 107", matching the spelling of the compiler target. NVIDIA's release notes are explicit about the status: Vera Rubin support in CUDA Toolkit 13.4 is a developer preview and is not intended for benchmarking, performance analysis, or production deployment.^[3]

Sharing a GPU: the Multi-Process Service

The Multi-Process Service (MPS) is a runtime service that lets CUDA work from several processes share one GPU cooperatively. It has three parts: a control daemon (nvidia-cuda-mps-control) that starts and stops servers and brokers client connections, a server process that owns the GPU scheduling resources and acts as the clients' shared connection to the device, and a client runtime built into the CUDA driver library that ordinary applications use transparently.^[30] MPS is worth deploying when no single process generates enough work to saturate the GPU, a common situation in strong-scaling runs where the problem size is fixed and the work per process shrinks; MPS lets kernels from different processes run concurrently instead of being serialized.^[33] It also provides memory and SM partitioning, priority, and dynamic resource adjustment.^[30] NVIDIA documents MPS as supported on Linux and QNX only, and only for 64-bit applications.^[33]

CUDA 13.4 introduces MPS V3, an opt-in control-daemon interface selected with the -p 3 flag or the CUDA_MPS_PROTOCOL_VERSION environment variable. It replaces the interactive shell of the legacy V2 interface with a scriptable module verb command syntax, supports multiple concurrently running named servers instead of one per user, adds namespaces that subdivide a server's resources and route clients to them, and reads a TOML file that defines servers, namespaces, partitions and features at daemon startup.^[34] The release notes additionally list cgroup-integrated device-memory limits and per-container time slicing for GPU fractionalization, with allocation and memory-reporting APIs, NVML and nvidia-smi all honoring the configured limit.^[3] Legacy MPS V2 remains the default and continues to work unchanged.^[34]

Compute Fabric Transport and memory locality

CUDA Compute Fabric Transport (CFT) is a transport-centric interface for moving data between GPUs across an NVLink fabric. The established approach maps a peer's physical memory into the local virtual address space, after which kernels read and write it with ordinary loads and stores. CFT instead exposes a logical endpoint: a named transport object identified by a 32-bit endpoint id, with a target inside it addressed by that id together with a 64-bit offset. Asynchronous put, get, atomic and reduction operations are then issued against that pair directly from the GPU.^[29][31]

NVIDIA gives two reasons for the design. A virtual address provides no channel for reporting a transient fabric error such as packet loss or a link reset, so a failing remote access faults or kills the process and a kernel has no way to retry or reroute. A virtual mapping can also only name memory, while an endpoint id names a resource, although memory is currently the only kind of resource that logical endpoints expose.^[31] An imported endpoint stays valid when the owner rebinds the memory behind it, which a mapped allocation does not.^[31] CFT is available only through the CUDA driver API and is aimed at authors of communication libraries; NVIDIA directs most developers to NCCL or NVSHMEM instead.^[29][31]

Two smaller memory changes accompany it. CUDA 13.4 exposes locality domains, which are portions of a GPU containing both streaming multiprocessors and device memory. An application can allocate device memory in a locality domain and create a green context with SM resources in the same domain, keeping computation near the memory it reads.^[29] New cuMemGetLocationInfo and cudaMemGetLocationInfo calls report where a managed or system-allocated unified-memory allocation currently resides, so libraries and runtimes can choose compute and communication resources from measured locality instead of assumption.^[3][29]

Driver packaging and the CDMM default

CUDA SDK installers no longer bundle the NVIDIA driver. The release notes date the change to CUDA 13.1 on Windows and CUDA 13.4 on Linux; the appropriate open-kernel-module driver or toolkit packages are installed separately through a package manager.^[3][29] The separately released R615 driver packages that correspond to CUDA 13.4 also drop the proprietary kernel modules, so supported Linux systems use the NVIDIA open kernel modules.^[3]

The R615 driver also changes a default that matters on NVIDIA's hardware-coherent platforms, where CPU and GPU memory sit in one coherence domain. The driver now defaults to Coherent Driver-based Memory Management (CDMM) instead of onlining GPU memory to the operating system as a NUMA node; NVIDIA names Grace Hopper, Grace Blackwell and Vera Rubin as examples of the affected platforms.^[29] NUMA mode remains fully supported and is selected with the NVreg_CoherentGPUMemoryMode kernel module parameter. The choice is node-wide and takes effect only after a driver reload or reboot, so NVIDIA advises selecting the mode before upgrading.^[3][29]

CUDA Python, CCCL, and tools

The stable Pythonic CUDA API added a large batch of features in cuda.core 1.1.0, released in July 2026 and highlighted in NVIDIA's 13.4 announcement; the 13.4 release notes list cuda.core 1.1.1 as the version shipping alongside the toolkit.^[3][32] It adds a cuda.core.texture module covering texture and surface programming, with OpaqueArray and MipmappedArray for hardware-laid-out allocations, TextureObject for bindless hardware-filtered reads and SurfaceObject for typed loads and stores. Managed-memory allocations now expose CUDA memory advice as properties, including read-mostly data, preferred location and accessing processors, with a Host type that can name any host memory, a particular NUMA node, or the node associated with the calling thread. The release also ships .pyi type stubs for every public API and exposes a captured graph as a GraphDefinition so that stream capture can be combined with explicit graph construction.^[29][32] The companion cuda.compute 1.1 adds ahead-of-time compilation of algorithm objects for several compute capabilities at once, including on build systems without a GPU, using proxy descriptions of argument types plus serialize() and deserialize().^[29]

CUDA 13.4 ships CCCL 3.4. A warp-specialized cub::DeviceScan implementation for Blackwell GPUs uses the Tensor Memory Accelerator to overlap data movement with computation while cutting synchronization overhead; NVIDIA reports that the new cub::DeviceScan::Sum reaches up to 92 percent memory-bandwidth utilization in its benchmarks, against up to roughly 50 percent for the previous implementation, across the data types tested.^[29] CUB device-wide algorithms gain single-call overloads that obtain temporary storage from a memory resource supplied through an execution environment, removing the older query-then-call pattern while leaving the two-phase APIs in place for code that needs explicit storage management.^[3][29] CCCL also adds cub::WarpReduceBatched for reducing several independent batches across one warp, and a set of C++ Standard Library parallel algorithms in cuda::std selected with the cuda::execution::gpu execution policy.^[3][29]

Tooling moved with the toolkit. Nsight Systems 2026.5.1 adds support for CUDA 13.4, Rubin GPUs and Windows on Arm, along with NIC metric collection through the NVIDIA DOCA Telemetry Service and an NCCL straggler-analysis recipe. Nsight Compute 2026.3 adds Tile IR support for CUDA Tile workloads. NVIDIA also released Nsight Python 1.0, a decorator and context-manager interface that automates kernel benchmarking and architectural metric collection from a Python script.^[29]

Use in Machine Learning and Scientific Computing

High-level frameworks can place tensor operations on CUDA devices and call CUDA libraries or generated kernels. PyTorch documents CUDA tensor placement, streams, memory management, graphs, and precision behavior.^[20] JAX publishes installation paths for NVIDIA GPU support using CUDA.^[21] Framework support should be checked against the framework's current compatibility documentation rather than inferred from the CUDA Toolkit version alone.

This layered design means that several statements can all be true:

• A researcher may use CUDA hardware acceleration without writing CUDA C++.
• A framework can dispatch some operations to cuDNN or cuBLAS and generate others itself.
• A workload can contain both CPU and GPU phases.
• The same framework can offer non-CUDA backends.
• An operation can fall back, be recompiled, or be unsupported depending on device, data type, package build, and software versions.

Large training jobs often combine computation libraries with communication libraries. A framework may use cuDNN or generated kernels for local tensor operations and NCCL collectives for gradients or parameters across GPUs. CUDA streams and events then help express local ordering and overlap. This is an ecosystem of components, not one monolithic library.

CUDA is also used in numerical simulation, linear algebra, signal and image processing, computational chemistry, medical imaging, analytics, and other parallel workloads. The available libraries demonstrate broad domain coverage, but whether a particular application benefits depends on its algorithm and implementation.^[17] Tasks with substantial independent work and reusable data on the GPU are generally better candidates than short serial tasks dominated by transfer or launch overhead.

Performance Engineering

CUDA supplies mechanisms for GPU execution; it does not guarantee a particular speedup. A defensible comparison fixes the algorithm, required accuracy, input size, CPU and GPU implementations, hardware, software versions, warm-up policy, transfer accounting, and timing boundaries. Comparisons that omit these details can attribute algorithmic or implementation differences to the platform.

NVIDIA's best-practices guidance emphasizes measuring the application, locating hotspots, and applying changes iteratively.^[9] Common performance factors include:

• Available parallelism: The grid needs enough ready work to use the GPU. A small launch may leave much of the device idle.
• Data movement: Transfers between host and device can dominate short computations. Keeping reusable data on the device or overlapping independent work may help.
• Memory access: Coalesced global-memory access and effective data reuse reduce wasted traffic. Shared memory can help, but copying into it is worthwhile only when reuse or access transformation repays the cost.
• Control flow: Warp divergence can leave lanes inactive. Divergence matters when it occurs within a warp, not merely because different blocks take different paths.
• Resource use: Registers and shared memory consumed by each block affect how many blocks and warps can be resident. Maximum occupancy is not automatically maximum performance.
• Arithmetic behavior: Precision, operation ordering, fused operations, and specialized hardware can change both speed and numerical results.
• Launch and synchronization overhead: Many tiny kernels or frequent host waits can consume time that a graph, batching, fusion, or asynchronous pipeline might reduce.

Profiling should be hypothesis-driven. A high memory-throughput percentage suggests different work than a dependency stall or low launch occupancy. Hardware counters also need context; one metric rarely identifies a correction by itself. Performance tuning that relies on undocumented behavior may break on a later architecture even when the program appeared to work on the original GPU.

Correctness testing must include asynchronous errors. A kernel launch can return control before execution fails, so an error may surface at a later synchronization point. Development builds often add explicit launch checks and synchronization to localize problems, then remove unnecessary global waits after the code is correct.

Numerical results can differ from CPU results or across GPU implementations because floating-point operations are not generally associative and parallel reductions can use different orders. A performance optimization that changes precision or accumulation order must be evaluated against the application's accuracy requirements, not treated as a transparent substitution.

Compatibility and Deployment

CUDA compatibility has several directions that are easy to confuse:

• Driver backward compatibility: A newer NVIDIA driver can generally run applications built with an older CUDA Toolkit.
• Minor-version compatibility: Beginning with CUDA 11, an application built within a major toolkit family can run on a driver meeting that family's minimum requirement, with limitations.
• PTX hardware forward compatibility: Compatible PTX can be compiled by a sufficiently new driver for a later GPU architecture.
• Application compatibility: Libraries, frameworks, extensions, and custom kernels still impose their own supported versions and architecture targets.

Minor-version compatibility is not unconditional. NVIDIA documents that newer toolkit features may require a newer driver, PTX generated by a newer toolchain cannot be consumed by an older driver that does not understand it, and applications must use appropriate architecture targets.^[13] Library packages can also depend on specific versions of other libraries.

The "CUDA Version" shown by tools such as nvidia-smi describes a driver capability, not necessarily the toolkit installed in a development environment or container. Deployment diagnosis should distinguish the host driver, user-space runtime and libraries, toolkit used to build an application, and GPU architecture.

Containers package user-space components but normally rely on the host's NVIDIA driver. A container image therefore does not make driver requirements disappear. Cluster operators commonly standardize driver branches, publish supported container families, and compile application artifacts for a deliberate set of compute capabilities.

Platform support is release-specific. At the cutoff, NVIDIA maintained separate installation guides for Linux and Microsoft Windows.^[15][16] The safest procedure is to select a toolkit release, then verify its release notes, driver requirements, host compiler support, operating-system matrix, and target GPU architectures as one configuration. Copying an old compatibility table into an application document is less reliable than checking the matching release documentation.

Licensing and Portability

The CUDA Toolkit is distributed under NVIDIA's SDK agreement. The agreement at the cutoff licensed development of applications for use in systems with NVIDIA GPUs and identified which toolkit portions could be redistributed.^[22] It also recognized bundled NVIDIA and third-party components with separate legal notices or licenses. For that reason, describing every file in the CUDA ecosystem with a single open-source or closed-source label is imprecise. Developers distributing an application need to review the agreement and the notices for the components they ship.

CUDA source code also carries a technical portability cost. CUDA C++ resembles standard C++, but kernels, runtime calls, library APIs, build targets, and performance assumptions are platform-specific. Separating domain logic from backend code, using stable library interfaces, and testing more than one target can reduce migration cost, but cannot guarantee equivalent behavior or performance.

Several alternatives address different portability goals:

• OpenCL is a royalty-free, cross-platform standard for parallel programming across heterogeneous systems.^[23]
• SYCL is a royalty-free C++ programming model for heterogeneous devices. Implementations map the specification to supported backends and hardware.^[24]
• AMD HIP is a C++ runtime and kernel language with CUDA-like concepts and tools for porting CUDA code. Similar APIs can reduce mechanical changes, but code may still need unsupported-feature work, library substitutions, and hardware-specific tuning.^[25]
• Higher-level frameworks and compiler systems can hide some backend differences, but only for operations and devices that the chosen backend implements.

These options should not be ranked by a generic API-level speed claim. Results depend on the implementation, compiler, libraries, workload, device, and tuning. Portability and peak performance are also not binary opposites: an application can keep portable high-level logic while using optimized backend-specific kernels for selected operations.

Limitations

CUDA's main technical limitation is that its licensed target and execution stack are tied to NVIDIA GPUs.^[22] Applications that depend directly on CUDA APIs or libraries need an NVIDIA-compatible deployment path, or a separate port for other hardware.

Efficient custom kernels require knowledge beyond ordinary sequential programming. Developers must reason about many concurrent threads, memory visibility, synchronization scope, asynchronous failure, resource limits, and architecture-dependent performance. Libraries and frameworks reduce this burden for standard operations, but they do not remove it from custom extensions or system integration.

Hardware and software support changes over time. A new toolkit can deprecate old compilation targets, require a newer driver for some features, or change host-toolchain support. Conversely, selecting an old toolkit to retain an old GPU target can prevent use of newer language, compiler, or library features. Long-lived applications need an explicit support matrix and reproducible builds.

Finally, GPU acceleration is not suitable for every workload. Serial dependencies, small problem sizes, irregular access, limited device memory, data-transfer cost, and latency requirements can make a CPU or another accelerator more appropriate. CUDA is a set of tools for expressing and executing supported parallel work, not evidence that any particular program will be faster.

See also

• NVIDIA H100
• TensorRT
• OpenAI Triton
• NVIDIA Vera Rubin

References

1. ^1 ^2 ^3 ^4NVIDIA, "Introduction," CUDA Programming Guide, updated September 9, 2026. docs.nvidia.com/...introduction
2. ^1 ^2 ^3 ^4 ^5 ^6 ^7 ^8 ^9 ^10 ^11NVIDIA, "Programming Model," CUDA Programming Guide, updated September 9, 2026. docs.nvidia.com/...programming-model
3. ^1 ^2 ^3 ^4 ^5 ^6 ^7 ^8 ^9 ^10 ^11 ^12 ^13 ^14 ^15 ^16 ^17 ^18 ^19 ^20 ^21NVIDIA, "CUDA Toolkit 13.4 Release Notes," 2026. docs.nvidia.com/...cuda-toolkit-release-notes
4. ^Ian Buck, Tim Foley, Daniel Horn, Jeremy Sugerman, Kayvon Fatahalian, Mike Houston, and Pat Hanrahan, "Brook for GPUs: Stream Computing on Graphics Hardware," ACM Transactions on Graphics 23(3), 2004. graphics.stanford.edu/...brookgpu.pdf
5. ^John Nickolls, Ian Buck, Michael Garland, and Kevin Skadron, "Scalable Parallel Programming with CUDA," ACM Queue 6(2), 2008. research.nvidia.com/...e-parallel-programming-cuda
6. ^Alex Krizhevsky, Ilya Sutskever, and Geoffrey E. Hinton, "ImageNet Classification with Deep Convolutional Neural Networks," Advances in Neural Information Processing Systems 25, 2012. proceedings.neurips.cc/...436e924a68c45b-Paper.pdf
7. ^1 ^2 ^3NVIDIA, "Asynchronous Execution," CUDA Programming Guide, updated September 9, 2026. docs.nvidia.com/...asynchronous-execution
8. ^1 ^2NVIDIA, "Compute Capabilities," CUDA Programming Guide, updated September 9, 2026. docs.nvidia.com/...compute-capabilities
9. ^1 ^2 ^3 ^4NVIDIA, "CUDA C++ Best Practices Guide 13.4," 2026. docs.nvidia.com/...cuda-c-best-practices-guide
10. ^1 ^2NVIDIA, "Unified and System Memory," CUDA Programming Guide, updated September 9, 2026. docs.nvidia.com/...understanding-memory
11. ^NVIDIA, "CUDA Graphs," CUDA Programming Guide, updated September 9, 2026. docs.nvidia.com/...cuda-graphs
12. ^1 ^2NVIDIA, "NVCC: The NVIDIA CUDA Compiler," CUDA Programming Guide, updated September 9, 2026. docs.nvidia.com/...nvcc
13. ^1 ^2NVIDIA, "Minor Version Compatibility," CUDA Compatibility Guide, 2026. docs.nvidia.com/...minor-version-compatibility
14. ^1 ^2NVIDIA, "The CUDA Driver API," CUDA Programming Guide, updated September 9, 2026. docs.nvidia.com/...driver-api
15. ^1 ^2NVIDIA, "CUDA Installation Guide for Linux 13.4," 2026. docs.nvidia.com/...cuda-installation-guide-linux
16. ^1 ^2NVIDIA, "CUDA Installation Guide for Microsoft Windows 13.4," 2026. docs.nvidia.com/...llation-guide-microsoft-windows
17. ^1 ^2NVIDIA, "CUDA Libraries Documentation," 2026. docs.nvidia.com/cuda-libraries
18. ^NVIDIA, "NVIDIA cuDNN," updated September 2, 2026. docs.nvidia.com/...latest
19. ^NVIDIA, "NVIDIA Deep Learning NCCL Documentation," 2026. docs.nvidia.com/...nccl
20. ^PyTorch Contributors, "CUDA Semantics," PyTorch documentation, updated July 17, 2026. docs.pytorch.org/...cuda
21. ^JAX Contributors, "Installation: NVIDIA GPU," JAX documentation, 2026. docs.jax.dev/...installation
22. ^1 ^2NVIDIA, "License Agreement for NVIDIA Software Development Kits," CUDA Toolkit EULA, updated January 26, 2026. docs.nvidia.com/...eula
23. ^Khronos Group, "OpenCL: The Open Standard for Parallel Programming of Heterogeneous Systems," 2026. khronos.org/opencl
24. ^Khronos Group, "SYCL 2020 Specification, Revision 11," 2025. registry.khronos.org/...sycl-2020
25. ^Advanced Micro Devices, "What Is HIP?," HIP documentation, 2026. rocm.docs.amd.com/...what_is_hip
26. ^NVIDIA, "Nsight Systems User Guide," 2026. docs.nvidia.com/...UserGuide
27. ^NVIDIA, "Nsight Compute Documentation," 2026. docs.nvidia.com/...NsightCompute
28. ^NVIDIA, "Compute Sanitizer Documentation," 2026. docs.nvidia.com/...ComputeSanitizer
29. ^1 ^2 ^3 ^4 ^5 ^6 ^7 ^8 ^9 ^10 ^11 ^12 ^13 ^14 ^15 ^16 ^17 ^18 ^19 ^20 ^21 ^22 ^23 ^24 ^25 ^26Jonathan Bentz, "CUDA Toolkit 13.4 Adds Windows on Arm Support and Greater Control over Shared GPUs," NVIDIA Technical Blog, September 9, 2026. developer.nvidia.com/...r-control-over-shared-gpus
30. ^1 ^2NVIDIA, "Multi-Process Service," MPS documentation, updated September 9, 2026. docs.nvidia.com/...latest
31. ^1 ^2 ^3 ^4 ^5NVIDIA, "Compute Fabric Transport," CUDA Programming Guide, 2026. docs.nvidia.com/...compute-fabric-transport
32. ^1 ^2 ^3NVIDIA, "cuda.core 1.1.0 Release Notes," CUDA Python documentation, 2026. nvidia.github.io/...1.1.0-notes
33. ^1 ^2NVIDIA, "When to Use MPS," MPS documentation, 2026. docs.nvidia.com/...when-to-use-mps
34. ^1 ^2NVIDIA, "MPS v3 Interface," MPS documentation, 2026. docs.nvidia.com/...mpsv3-interface