CUDA Python Cheat Sheet

CUDA Python cheat sheet

CUDA Python (cuda-python) is NVIDIA's official Python interface to CUDA. It's now a metapackage of independently versioned parts: cuda.bindings — thin, 1:1 Python bindings to the CUDA Driver, Runtime, and NVRTC C APIs; and cuda.core — a higher-level, Pythonic layer (Device, Stream, Program, launch) that hides the boilerplate. Use it to compile and launch your own CUDA C++ kernels, manage device memory, and interop with CuPy/Numba/PyTorch. Targets cuda-python 13.x (CUDA 12/13), Python 3.9+.

setup & device run kernels (cuda.core) low-level bindings streams & interop math libs & ecosystem gotcha ★ most common

Verified 2026-08-31 against nvidia.github.io/cuda-python (cuda-python 13.x = cuda.bindings + cuda.core; CUDA 12.x & 13.x). cuda.core lives under the cuda.core.experimental namespace (API still stabilizing). Not a kernel-writing DSL — for that see Numba/Warp; this drives CUDA from Python.

Outline

Reach for cuda.core for a clean Device/Stream/launch workflow; drop to cuda.bindings when you need an exact C API call. Interop lets you share memory with CuPy/Torch instead of copying.

1Install & the splitmetapackage

pip install cuda-python★
The umbrella metapackage — pulls in cuda-bindings + cuda-core. Needs an NVIDIA GPU + driver; the CUDA toolkit libs come via pip wheels.
pip install cuda-core cuda-bindings # install parts independently
Each subpackage is versioned separately now, so you can pin just what you use.
from cuda.core.experimental import Device, Program, LaunchConfig, launch from cuda.bindings import driver, runtime, nvrtc★
Two import roots: high-level cuda.core.experimental, low-level cuda.bindings.*.
# old: from cuda import cuda, cudart, nvrtc (pre-12.x layout)migration
The flat from cuda import cuda layout was replaced by cuda.bindings.driver / .runtime / .nvrtc. Update old code.

2Device & contextcuda.core

dev = Device(0) dev.set_current()★
Device wraps a GPU + its context. set_current() makes it active for this thread.
dev.name · dev.compute_capability · dev.properties
Query the device. Device().device_id; iterate all GPUs with range(Device.num_devices)-style helpers.
stream = dev.create_stream()★
A Stream is the ordered queue you submit compiles, copies, and kernel launches to (card 9).
arch = "".join(str(i) for i in dev.compute_capability) # e.g. "90"
You'll pass the arch (sm_90) to the compiler in card 3.

3Compile a kernelNVRTC, made easy

code = r''' extern "C" __global__ void add(float* a, float* b, float* c, int n){ int i = blockIdx.x*blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i]; }'''★
Write the kernel as a CUDA C++ string. Use extern "C" so the name isn't mangled.
from cuda.core.experimental import Program, ProgramOptions prog = Program(code, code_type="c++", options=ProgramOptions(arch="sm_90"))★
Program compiles source via NVRTC. Set target arch through ProgramOptions.
mod = prog.compile("cubin") add = mod.get_kernel("add")★
Compile to a module ("cubin" or "ptx"), then grab the kernel handle by name.

4Launch config & launchgrid × block

block = 256 grid = (n + block - 1) // block cfg = LaunchConfig(grid=grid, block=block)★
LaunchConfig replaces the <<<grid, block>>> syntax. Add shmem_size= for dynamic shared memory.
launch(stream, cfg, add, a_ptr, b_ptr, c_ptr, np.int32(n))★
Submit the kernel to a stream with its arguments (device pointers + scalars). It runs asynchronously.
stream.sync()★
Block until the stream's queued work (launch, copies) finishes before reading results back.
# pass scalars with the right dtype (np.int32/float32), not python intgotcha
Kernel arg types must match the C++ signature exactly — wrap scalars with numpy dtypes.

5Memory & buffersallocate & copy

buf = dev.allocate(n * 4) # device Buffer, n float32s★
Device.allocate() returns a device Buffer. buf.handle is the raw pointer to pass to launch.
buf.copy_from(host_array, stream=stream) buf.copy_to(host_out, stream=stream)★
Async host↔device transfers on a stream. Buffers follow the memory-resource model for pooled allocation.
buf.close() # or use as a context manager
Free device memory explicitly (or scope it with with).

6Runtime APIcudart, in Python

from cuda.bindings import runtime as cudart err, = cudart.cudaSetDevice(0)★
The Runtime API mirrors cudart 1:1. Every call returns a tuple whose first element is an error code.
err, dptr = cudart.cudaMalloc(n * 4) err, = cudart.cudaMemcpy(dptr, hptr, n*4, cudart.cudaMemcpyKind.cudaMemcpyHostToDevice)★
Classic malloc/memcpy flow. Enums are nested (cudaMemcpyKind.*). cudaFree, cudaDeviceSynchronize, cudaMemset.
err, count = cudart.cudaGetDeviceCount() err, props = cudart.cudaGetDeviceProperties(0)
Device enumeration & properties, exactly as in C.

7Driver & NVRTCthe low road

from cuda.bindings import driver, nvrtc driver.cuInit(0) err, ctx = driver.cuCtxCreate(0, dev)
The Driver API (cu*): explicit context/module/stream management — the manual version of what cuda.core automates.
err, prog = nvrtc.nvrtcCreateProgram(src, b"k.cu", 0, [], []) nvrtc.nvrtcCompileProgram(prog, len(opts), opts)
NVRTC runtime compilation at the C level. Fetch PTX with nvrtcGetPTX, load via cuModuleLoadData + cuModuleGetFunction.
driver.cuLaunchKernel(func, gx,gy,gz, bx,by,bz, shmem, stream, args, 0)
The raw launch. cuda.core.launch exists precisely so you rarely write this.
from cuda.bindings import nvjitlink, nvvm # link/IR tooling
Bindings also cover nvJitLink, NVVM, nvFatbin, cuFile (GPUDirect Storage), and NVML.

8Error checkingcheck every call

def CHECK(res): err, *vals = res if err != cudart.cudaError_t.cudaSuccess: raise RuntimeError(cudart.cudaGetErrorString(err)) return vals[0] if len(vals)==1 else vals★
Bindings never raise — you must check the returned error code. A tiny helper keeps call sites clean.
ptr = CHECK(cudart.cudaMalloc(n*4))
Wrap every binding call. For the driver API compare against driver.CUresult.CUDA_SUCCESS.
# silent wrong results usually = an unchecked error codegotcha
cuda.core raises Python exceptions on failure; the raw bindings do not — don't skip the check.

Worth memorizing

cuda-pythonmetapackage = cuda.bindings + cuda.core

cuda.core.experimentalhigh-level Device/Stream/launch

cuda.bindings.{driver,runtime,nvrtc}1:1 C APIs

Device(0).set_current()select a GPU

Program(...).compileNVRTC-compile CUDA C++ at runtime

LaunchConfig + launchreplaces <<<grid,block>>>

dev.allocate / copy_fromdevice Buffer + async copies

stream.sync()wait before reading results

bindings return (err, ...)check the error code yourself

__cuda_array_interface__ / DLPackzero-copy interop

nvmath-pythoncuBLAS/cuFFT/cuSOLVER, pythonic

write kernels in Python?Numba @cuda.jit / Warp, not this

CUDA Python cheat sheet

Outline

Setup & device

Run kernels (cuda.core)

Low-level bindings

Streams & interop

Math libs & ecosystem

Setup & Device

Run Kernels (cuda.core)

Low-level Bindings

Streams & Interop

Math Libs & Ecosystem

Worth memorizing