CUDA is an NVIDIA technology that allows thousands of video card cores to be used for solving non-trivial computational problems. The graphics processor performs many operations simultaneously rather than sequentially, which significantly speeds up the processing of huge data arrays compared to the central processing unit.
CUDA is used wherever massive volumes of information need to be processed. Key areas include deep learning and neural network training, scientific modeling of physical and chemical processes, computer vision, real-time video stream processing, and financial analysis. The architecture is critically important for autonomous vehicles, medical imaging, and computer graphics tasks requiring non-game rendering.
A typical problem is the complexity of debugging data race conditions and deadlocks when parallelizing code. The programmer must consider the memory hierarchy: frequent data transfers between system RAM and video memory create a performance bottleneck. Suboptimal thread configuration leads to idle compute units, and mismatched CUDA architecture versions often cause hidden compatibility errors at the compilation stage.
How CUDA works
The principle of operation is based on the SIMT architecture (single instruction, multiple threads), which differs from classical SIMD instructions of central processors in its flexibility. The central processor sends a kernel (function) to the GPU, where the scheduler distributes tasks among streaming multiprocessors. Unlike OpenGL or DirectX compute shaders, CUDA provides direct access to shared memory and allows arbitrary algorithms to be implemented without being tied to the graphics pipeline. Compared to OpenCL, CUDA technology offers a richer set of libraries (cuBLAS, cuDNN) and better integration with proprietary profiling tools. Groups of 32 threads (warps) execute instructions synchronously. The key efficiency mechanism is the hiding of global video memory access latency: while one warp is waiting for data, the scheduler instantly switches the multiprocessor to execute another warp that is ready to run, reducing pipeline stalls to zero without the costly context switches characteristic of CPUs.
CUDA functionality
- Massive parallelism programming model. CUDA provides an extension of the C/C++ language that allows thousands of lightweight threads to be launched on the GPU simultaneously. Threads are grouped into blocks forming a grid, ensuring scaling from hundreds to millions of cores without changing the application code.
- Thread execution hierarchy. The programmer defines a kernel executed by an array of threads. Each thread has unique indices:
threadIdx,blockIdx,blockDim, andgridDim. These built-in variables allow a thread to identify its position in the global problem space for computing the addresses of the data being processed. - Thread grouping into warps. The hardware scheduler groups 32 threads into a warp, the minimum execution unit on multiprocessors. All threads in a warp execute one instruction simultaneously under the SIMT architecture. Divergence of warp threads into different code branches leads to sequential execution of paths and reduced performance.
- Built-in vector data types. For efficient work with graphics and simulation data, CUDA defines short vector types:
float2,float4,char4,double2. Access to them is hardware-aligned, and vector operations are translated into a single instruction, increasing the throughput of the memory subsystem. - Kernel function specifiers. The
__global__specifier denotes a kernel function called from the host and executed on the device. The__device__specifier indicates a function called and executed only on the GPU. The__host__specifier paired with__device__allows the compiler to generate code versions for CPU and GPU simultaneously. - Virtual architectures and JIT compilation. The NVCC compiler separates the source code into host and device parts. Generating intermediate PTX representation for a virtual architecture allows the driver’s JIT compiler to adapt the code to the actual hardware, ensuring compatibility of new GPUs with previously compiled applications.
- Dynamic parallelism. Starting with the Kepler architecture, kernels can spawn child kernels without returning control to the host. Such nested parallelism allows recursive algorithms and adaptive computational refinement to be implemented directly on the GPU side, minimizing data transfer over the PCI Express bus.
- Memory hierarchy management. A CUDA kernel operates with registers, local memory, shared memory (
__shared__), global memory, constant memory (__constant__), and texture memory. Manual allocation is critically important: registers are the fastest resource, shared memory serves as a programmable L1 cache for data exchange within a block. - Atomic operations. For correct simultaneous data modification from different threads without locks, functions from the
atomicAdd,atomicCAS, andatomicExchfamilies are used. They guarantee the execution of a read-modify-write operation as an indivisible transaction in shared or global memory, ensuring state consistency. - Asynchronous stream parallelism. Using different streams allows independent operations to be overlapped: while one kernel is computing, another kernel can copy data over the bus or use a different compute unit. Creating non-blocking streams maximizes the utilization of computational resources and bandwidth.
- System event profiling. The CUPTI interface provides access to low-level hardware counters and driver activity tracing. Profilers based on this API collect metrics on multiprocessor occupancy, cache misses, and warp serialization to find bottlenecks in kernel code.
- Unified memory mechanism. Unified Memory technology creates a single virtual address space for CPU and GPU. The programmer uses a regular pointer, and the driver and hardware page controller automatically migrate data to the processor that accesses it, simplifying the porting of complex data structures.
- Pipelined asynchronous prefetching. The
cuda::memcpy_asyncextension in cooperative groups allows data loading from global memory to shared memory to be initiated without the involvement of thread computational instructions. The thread can continue computations with already available data, and after the asynchronous operation completes, use a barrier for synchronization. - Cooperative groups. This programming interface abstracts the level of synchronization, allowing collective operations to be explicitly specified over a group of threads ranging in size from a warp to the entire grid. Functions like
coalesced_grouporgrid_groupwith async()call ensure portability of barrier synchronization across different GPU architectures. - Constant broadcasting via cache. Data marked as
__constant__is read-only accessible to all threads and cached in dedicated constant memory. If all threads in a warp read the same address, the cache broadcasts the value in a single cycle, critically accelerating access to coefficient tables or algorithm parameters. - Texture memory and hardware interpolation. Texture memory blocks are cached taking two-dimensional spatial locality into account. The hardware texture sampling unit performs bilinear filtering and coordinate normalization for free, which is indispensable in image processing where pixel addressing is often done with floating-point values.
- cuDNN deep neural network library. This library implements standard primitives for convolution, pooling, activations, and normalization, optimally tuned considering tensor sizes. It automatically selects a convolution strategy (Winograd, FFT, implicit gemm), choosing the configuration that ensures minimal execution time on the target GPU.
- Tensor cores for matrix arithmetic. Starting with Volta, specialized Tensor Core units are integrated into multiprocessors. They perform fused multiply-add of
4x4matrices in a single cycle. In mixed precision mode, input matrices are supplied inFP16format, and accumulation occurs inFP32, dramatically accelerating neural network training. - Tensor Core (Hardware matrix multiplication with accumulation)Tensor (Multidimensional container for numerical data)FP16 (Compact representation of Floating-Point numbers)
- CUB parallel algorithms library. The adaptive low-level library provides occupancy-optimized blocks for basic operations: reduction, scan (prefix sum), and histogramming. CUB automatically tunes the work distribution strategy to the input data size and target GPU characteristics, hiding the complexity of fine kernel tuning.
- Interprocess memory sharing. CUDA IPC technology allows exporting allocated memory descriptors from one process and importing them into another. This enables multiple CPU processes to directly exchange data on the GPU without copying through the host, accelerating high-load inference services in a containerized environment.
- IPC (Instructions per cycle)
- Multi-node scaling with GPUDirect. The RDMA interface allows InfiniBand or RoCE network adapters to read and write directly to GPU global memory. Bypassing the CPU and OS kernel buffers, data is transmitted over the network with minimal latency, ensuring linear scalability in large clusters for distributed training tasks.
Comparisons
- CUDA vs OpenCL. CUDA is a proprietary NVIDIA platform with deep hardware integration, providing maximum performance through the NVCC compiler and optimized
cuBLASandcuDNNlibraries. OpenCL offers an open standard for cross-platform computing on GPUs and FPGAs from various vendors, but suffers from averaged performance and a more verbose programming interface compared to CUDA. - CUDA vs DirectCompute. DirectCompute functions as a component of the DirectX API, targeted exclusively at the Windows environment and primarily oriented towards gaming and multimedia tasks using HLSL. CUDA provides access to the full command set of the NVIDIA architecture without graphics pipeline restrictions, supporting low-level shared memory operations and asynchronous data transfer outside the rendering framework.
- CUDA vs ROCm. ROCm is an open-source AMD initiative offering an HIP compiler for direct translation of CUDA code to Radeon Instinct and Radeon Pro GPUs. Despite syntactic proximity, ROCm lags in driver stability and ecosystem maturity, whereas CUDA guarantees strict computational determinism and exclusive access to tensor cores for accelerating matrix operations.
- CUDA vs SYCL. SYCL is based on standard C++ and uses Data Parallel C++ abstractions, allowing a single codebase to run on NVIDIA, Intel, and AMD through unified shared memory mechanisms. In comparison, CUDA relies on language extensions requiring a separate compilation stage but provides finer control over thread hierarchy, texture memory, and precise profiling through the Nsight toolkit.
- CUDA vs WebGPU. WebGPU implements GPU computing in the browser via an API designed for web environment isolation, with shader compilation into SPIR-V and command buffer management. CUDA provides orders of magnitude greater throughput through direct interaction with host memory and NVLink, making it indispensable for computationally intensive simulations, whereas WebGPU is suitable for lightweight model inference on the client.
OS and driver support
CUDA operates through a two-component software infrastructure: a user-mode driver (libcuda.so or nvcuda.dll library) and a kernel-mode driver installed as part of the NVIDIA Display Driver package. Support covers 64-bit versions of Windows (WDDM driver), Linux (with an open kernel module and proprietary GSP microcontroller on the GPU), and macOS (discontinued after CUDA 10.2, driver no longer supplied). The kernel-mode driver is responsible for video memory allocation, context management, and queue scheduling, while the user-mode library translates CUDA API calls into ioctl commands to the device through the software interface, ensuring unified operation of compute kernels across different operating systems.
Security
Computation protection is implemented through the CUDA context model, where each process receives an isolated virtual GPU address space, preventing the reading of another process’s memory without explicit use of interprocess sharing (IPC) with permission checks through descriptors exported by the driver. Hardware support for Unified Memory with the page migration mechanism on GPUs with Hopper architecture and newer utilizes the Multi-Instance GPU (MIG) hardware isolation unit, allowing a single GPU to be physically partitioned into independent instances with hardware-enforced memory bandwidth and cache isolation, while access control to contexts and synchronization are ensured by trusted driver code, excluding unauthorized command buffer modification.
Logging
The logging mechanism in CUDA is based on the CUDA Profiling Tools Interface (CUPTI) and built-in callback functions registered in the context via the cuptiSubscribe API. When kernel execution errors or memory faults occur, the driver writes a data structure to the system log, simultaneously generating a CUPTI_EVENT_KERNEL_ERROR event, which the user subscription intercepts and outputs to the console or a file. Low-level tracing is implemented through the NVIDIA Tools Extension SDK library, which injects logging agents into the GPU command stream without stopping execution: events are written to a GPU ring buffer, from where they are asynchronously retrieved by the host upon calling cuptiActivityFlushAll and serialized into a structured log without affecting the critical execution path.
Limitations
Each specific GPU has a grid dimension limit (maximum 2^31-1 blocks along the X axis), a limit on the number of threads per block (1024 for most architectures), and a shared memory volume per block limit (up to 48 KB configurable plus 164 KB per SM in the Blackwell architecture), requiring the developer to manually set the kernel launch configuration in triple angle brackets with verification via cudaGetDeviceProperties. The programming model imposes a requirement to match the compute capability when compiling PTX into SASS for the target architecture (the -arch flag), and the use of dynamic parallelism or Unified Memory is limited to GPUs with Compute Capability 3.5 and above, which is strictly controlled by the JIT compiler at module loading time.
History and development
CUDA debuted in 2006 with the G80 GPU and CUDA 1.0 driver, introducing the SIMT programming model and C language extension for GPU programming without graphics APIs. Revolutionary changes included the emergence of Fermi (2010) with a cache hierarchy and ECC support, Kepler (2012) with dynamic parallelism allowing GPU kernels to spawn child kernels, and Volta (2017) with tensor cores for mixed precision. The modern stage of development (Hopper and Blackwell, 2022–2025) brought clusters of streaming multiprocessors with distributed shared memory within the cluster, an asynchronous programming model through thread-level barriers, and support for the Transformer Engine, performing FP8 computations with dynamic coefficient scaling in hardware.