NVIDIA CUDA 13.3 Enhances GPU Development with Tile Programming in C++, Compiler Autotuning, and Python Updates NVIDIA released CUDA 13.3, introducing tile programming in C++ that automates low-level GPU management for optimized kernel development across all supported architectures. The update also includes CUDA Python 1.0 with green contexts and process checkpointing, plus the CompileIQ compiler autotuning framework that delivers up to 15% speedup on critical kernels like GEMM and attention. NVIDIA CUDA 13.3 brings new capabilities and performance optimizations to developers across the CUDA ecosystem. The launch of NVIDIA CUDA Tile programming in C++ , enables high-level, tile-based kernel development that automatically manages complex low-level GPU details for optimal performance and portability. Additionally, CUDA Tile programming is now supported on Compute Capability 9.0 NVIDIA Hopper GPUs in addition to all other supported GPU architectures. We are also releasing CUDA Python 1.0, solidifying the support and stability of the CUDA Python SW ecosystem, and introducing critical features like green contexts and process checkpointing. For performance enthusiasts, the newly launched NVIDIA CompileIQ compiler auto-tuning framework delivers up to a 15% speedup on critical kernels like GEMM and attention. This release also features official C++23 support in NVCC, expanded tensor interoperability with DLPack/mdspan in CCCL 3.3, and numerous updates to the math libraries cuBLAS, cuSPARSE, cuSOLVER and profiling tools Nsight Compute and Nsight Systems . Release of CUDA Tile C++ With the release of CUDA 13.3, CUDA Tile support is extended to C++, enabling the large existing C++ codebase and developer base to create highly-optimized GPU tile kernels. This model automates parallelism, memory movement, asynchrony, and other low-level details, resulting in C++ code that is portable across NVIDIA GPU architectures. For more information, check out our blog post https://developer.nvidia.com/blog/develop-high-performance-gpu-kernels-in-cpp-with-nvidia-cuda-tile/ . Release of CUDA Python 1.0 CUDA Python is a set of libraries that expose CUDA to the Python programming language. By providing the 1.0 release, we are committing to semantic versioning: ensuring breaking API changes only during major-version releases. Minor releases add features and patch releases are bug fixes. Any public API scheduled for removal is first deprecated in a minor release with a clear replacement path. The following is more information on the software components included in CUDA Python 1.0. library | description | next major version | cuda.binding | Low-level Python bindings to CUDA C APIs. | 13.3.0 | cuda.core | Pythonic access to CUDA Runtime and other core functionality | 1.0.0 | cccl-cuda | Pythonic access to CCCL parallel algorithms and easy access to CCCL’s highly efficient and customizable parallel algorithms | 1.0.0 | cuda-pathfinder | Utilities for locating CUDA components installed in the user’s Python environment | 1.6 | cuda.coop is also available in the cuda-cccl package under the experimental namespace, which is subject to API changes. cuda.coop provides the reusable block-wide and warp-wide device primitives for use within Numba CUDA kernels. cuda.core is now stable cuda.core provides a Pythonic interface to the CUDA runtime, including devices, streams, programs, linkers, memory resources, and graphs. Version 1.0 consolidates APIs that have been stabilizing over the previous release cycles into a single supported surface. At the same time, we added support for green contexts, CUDA checkpointing, and more. Green contexts: Split a GPU’s SMs into disjoint partitions, each with its own context and streams, so latency-sensitive kernels are shielded from long-running throughput kernels in the same process. Process checkpointing : Snapshot the full CUDA state of a running process—including device allocations, streams, context—and restore it later. Unlocks CRIU-style workflows for GPU processes: fault-tolerant long jobs, preemption and migration on shared clusters, and fast warm-start of inference workers. Only available in Linux. Inter-process sharing IPC : Share GPU memory across Python processes without copying through the host. One process allocates, and others map the same physical VRAM into their own address space. Ideal for multi-process ML serving and zero-copy producer/consumer pipelines. The following are quick examples of how to use cuda.core APIs. python from cuda.core import Device, Stream, Program, ProgramOptions, LaunchConfig, launch pick and activate a GPU dev = Device dev.set current create a CUDA stream stream = dev.create stream NVRTC compile + lookup prog = Program src, code type="c++", options = ProgramOptions arch=f"sm {dev.arch}" kernel = prog.compile "cubin" .get kernel "my kernel" launch a kernel launch stream, LaunchConfig grid=64, block=256 , kernel, args JIT-LTO linking from cuda.core import Linker, LinkerOptions module = Linker obj1, obj2 , options=LinkerOptions arch=f"sm {dev.arch}" .link "cubin" NVRTC precompiled headers from cuda.core import ProgramOptions opts = ProgramOptions std="c++17", arch=f"sm {dev.arch}", create pch=True, pch dir="/tmp/pch" Memory resources, incl. NUMA-aware pools from cuda.core import DeviceMemoryResource, PinnedMemoryResource, PinnedMemoryResourceOptions, ManagedMemoryResource, ManagedMemoryResourceOptions NUMA-pinned host memory pinned = PinnedMemoryResource PinnedMemoryResourceOptions numa id=0 CUDA graphs: stream capture and explicit construction from cuda.core.graph import GraphBuilder, GraphDef gb = stream.create graph builder gb.begin building graph = gb.end building .complete graph.launch stream gdef = GraphDef gdef.add kernel node kernel, LaunchConfig grid=64, block=256 , args=args IPC: share GPU memory across Python processes from cuda.core import DeviceMemoryResource, DeviceMemoryResourceOptions mr = DeviceMemoryResource dev, options=DeviceMemoryResourceOptions max size=1 << 20, ipc enabled=True buffer = mr.allocate nbytes buffer is picklable and can be sent over mp.Queue Green contexts: partition SMs into disjoint groups from cuda.core import ContextOptions, SMResourceOptions sm = dev.resources.sm long grp, crit grp = sm.split SMResourceOptions count= sm.sm count - 16, 16 0 ctx crit = dev.create context ContextOptions resources= crit grp s crit = ctx crit.create stream Process checkpoint / restore Linux from cuda.core import checkpoint proc = checkpoint.Process os.getpid proc.lock timeout ms=5000 proc.checkpoint proc.restore proc.unlock device allocations and context are restored TMA / TensorMapDescriptor from cuda.core import StridedMemoryView, TensorMapDescriptor tmap = StridedMemoryView tensor .as tensor map box shape= 128, DLPack-friendly strided views from cuda.core.utils import StridedMemoryView view = StridedMemoryView torch tensor ; capsule = view. dlpack System info NVML from cuda.core import system print system.num devices, system.driver version cuda.bindings.nvml from cuda.bindings import nvml nvml.init name = nvml.device get name nvml.device get handle by index v2 0 cuda.bindings.nvfatbin from cuda.bindings import nvfatbin handle = nvfatbin.create CCCL Python release 1.0.0: cuda.compute cuda.compute brings the CUDA Core Compute Libraries CCCL ’s highly tuned parallel algorithms—sort, scan, reduce, transform, unique, histogram, top-k, and more—to Python as host-callable building blocks. Changes since the last release include: - Python lambdas can be used as algorithm operators, reducing boilerplate for simple reductions, scans, transforms, and predicates. - Algorithms support operators with side effects state , enabling use cases like running accumulators and conditional transforms. - New cuda.compute.upper bound and cuda.compute.lower bound APIs expose CUB’s parallel binary search to Python. - Consolidated caching across all algorithms for faster repeated invocations. python import cuda.compute from cuda.compute import OpKind d input = cp.arange 1, 1 000 001, dtype=cp.int32 d output = cp.empty 1, dtype=cp.int32 h init = np.array 0 , dtype=np.int32 cuda.compute.reduce into d input, d output, OpKind.PLUS, d input.size, h init cuda.compute.reduce into d input, d output, lambda a, b: a if a b else b, d input.size, h init, cuda.coop exposes CCCL’s warp-wide and block-wide cooperative primitives for use inside Numba CUDA kernels. At the moment, this module is under the experimental namespace and may have API changes that don’t follow semantic versioning. python from numba import cuda from cuda.coop. experimental import block, warp THREADS = 128 block sum = coop.block.make sum numba.int32, THREADS @cuda.jit link=block sum.files def reduce kernel data, out : Each thread contributes one element to the block-wide reduction total = block sum data cuda.threadIdx.x if cuda.threadIdx.x == 0: out 0 = total h in = np.ones THREADS, dtype=np.int32 d in = cuda.to device h in d out = cuda.device array 1, dtype=np.int32 reduce kernel 1, THREADS d in, d out assert d out.copy to host 0 == THREADS 128 New Numba CUDA MLIR backend Numba CUDA MLIR is a new Numba-compatible kernel generator for Python, written from the ground up on top of MLIR and the modern NVVM toolchain. It preserves the familiar @cuda.jit programming model from Numba-CUDA while delivering lower compile latency, better diagnostics, and a cleaner path to target new GPU architectures and features as they land in the NVVM stack. Numba CUDA MLIR can be used as a drop-in replacement for numba.cuda by simply replacing the import statement: python Before from numba import cuda After from numba cuda mlir import cuda @cuda.jit def vector add a, b, out : i = cuda.grid 1 if i < out.shape 0 : out i = a i + b i Beyond existing Numba-CUDA compatibility Numba CUDA MLIR also features: Faster JIT compile https://github.com/NVIDIA/numba-cuda-mlir/blob/main/tests/benchmarks/README.md . Across a suite of real kernels vector add, softmax, Cholesky, attention, Black-Scholes, FFT, matmul , warm JIT compile times are ~1.4x faster on geomean and up to ~2x faster on individual kernels versus Numba-CUDA. Lower launch latency https://github.com/NVIDIA/numba-cuda-mlir/blob/main/tests/benchmarks/launch latency ubench/README.md . Host-side kernel dispatch overhead drops by roughly 2-3.5x for typical kernels and up to ~17x for kernels with many scalar arguments, where argument packing previously dominated. You can test Numba CUDA MLIR 0.3 by installing it from PyPI numba-cuda-mlir cu13 and follow its development on GitHub. Try CUDA Python today Install the CUDA Python stack directly from PyPI: pip install cuda-python cuda-cccl numba-cuda-mlir cu13 This pulls in cuda.bindings 13.3.0 , cuda.core 1.0.0 , cuda.compute 1.0.0 , along with cuda-pathfinder for library discovery. CompileIQ launched A new compiler auto-tuning framework for maximum performance on GPU kernels called CompileIQ, launches with CUDA 13.3. GPU compilers apply generic optimization heuristics that are broadly effective but aren’t necessarily optimal for specific kernels. CompileIQ flips this dynamic by using evolutionary and genetic algorithms to generate specialized compiler configurations custom-tailored to each kernel. This unlocks extra performance. For example, for critical kernels like GEMM and attention, which account for over 90% of LLM inference compute, CompileIQ delivers up to a 15% speedup on already-optimized Triton attention and CUTLASS GEMM kernels. Read more about CompileIQ, including how it works and how to use it, in this blog post. Math libraries Core CUDA math libraries in CUDA 13.3 include several new features and notable performance improvements available, including: - cuSPARSE: - Support for CSC format in SpSV and SpSM. - Support for mixed precision in SpMVOp. - Support for mixed index type 64-bit offset, 32-bit index CSR matrix in SpMvOp computation - Improved cusparseSpMVOp createDescr performance by 2.5x. - Introduced new API SPMVOP ALG1, which supports: - Updating matrix values while maintaining the same sparsity pattern. - Optimized buffer size. - Reduced preprocess overhead. - cuBLAS: - CUDA green context support. - Performance improvement to FP4 matmuls on NVIDIA Blackwell Ultra. - Performance improvement to TF32 matmuls on NVIDIA Blackwell and Blackwell Ultra. - SYMV performance improvements for NVIDIA Hopper, Blackwell, and Blackwell Ultra. - Improved user experience for FP64 emulated matmuls by enforcing a fixed workspace size that is constant across the problem space. - cuSOLVER: - A 64-bit interface cusolverDnXpolar exposes the QDWH algorithm implementation for polar decomposition in cuSOLVERDn - A 64-bit interface cusolverDnXstedc , which computes the eigenvalues and, optionally, eigenvectors of a symmetric tridiagonal matrix using the divide and conquer method - Performance improvements for cusolverDnXgeev with eigenvectors by moving the eigenvector post-processing from the host to the device. - A 64-bit interface - Public 64-bit interface cusolverDnXpolar , which exposes the QDWH algorithm implementation for polar decomposition in cuSOLVERDn available in 13.2 U1 . - Public 64-bit interface cusolverDnXstedc , which computes the eigenvalues and, optionally, eigenvectors of a symmetric tridiagonal matrix using the divide and conquer method available in 13.2 U1 . - Performance improvements for cusolverDnXgeev with eigenvectors by moving the eigenvector post-processing from the host to the device. cusolverDn D,Z syevj uses low-precision preconditioning, which typically improves the time-to-solution by 20% for mid-sized and large matrices on B200, and by even more on GPUs with a large FP32: FP64 ratio. CCCL CUDA 13.3 ships with CCCL 3.3. Highlights include DLPack/mdspan interoperability, a comprehensive random number distribution library, new search and segmented scan algorithms, and a flexible N-to-M transform. Tensor interoperability Deep learning frameworks speak in tensors, but CUDA C++ code often has to work one level lower—raw pointers, shapes, strides, and hand-written indexing. CCCL makes it easier to preserve that tensor structure across the boundary between Python frameworks and CUDA C++. With DLPack interoperability, tensors from frameworks such as PyTorch, JAX, and CuPy can be converted into cuda::std::mdspan views with cuda::to device mdspan https://nvidia.github.io/cccl/unstable/libcudacxx/extended api/mdspan/dlpack to mdspan.html conversion-functions for use in C++ kernels, and cuda::std::mdspan views can be converted back to DLPack with . https://nvidia.github.io/cccl/unstable/libcudacxx/extended api/mdspan/dlpack to mdspan.html conversion-functions cuda::to dlpack tensor CCCL also extends this tensor-view model inside kernels with cuda::shared memory mdspan https://nvidia.github.io/cccl/unstable/libcudacxx/extended api/mdspan/shared memory accessor.html . Instead of treating shared memory as a flat buffer, developers can create multi-dimensional views over shared-memory tiles, making indexing clearer and less error-prone. The shared-memory specialization also provides address-space safety checks and guarantees shared-memory load/store instructions. Random number distributions CCCL 3.3 adds a comprehensive set of device-compatible random distributions to