Frontiers of Systems Programming

CUDA

/ KOO-duh /

Once you have a GPU and want to run your own computation on it, you need a way to write that computation, ship it to the device, and launch thousands of copies of it. CUDA is NVIDIA's platform and programming model for doing exactly that on NVIDIA GPUs. You write a small function called a kernel in a C/C++ dialect, mark it to run on the device, and then ask the GPU to run it across a grid of threads.

The model has a specific vocabulary worth getting right. A kernel is the function each thread runs. You launch it with a configuration like myKernel<<<numBlocks, threadsPerBlock>>>(...), which creates a grid made of blocks, each block made of threads; every thread learns its own coordinates (blockIdx, threadIdx) and uses them to pick which data element it owns - thread number i handles c[i]. Underneath, the hardware does not schedule threads one at a time: it runs them in lockstep groups of 32 called a warp, where all 32 threads execute the same instruction together. Threads in a block can cooperate through fast on-chip shared memory and synchronize with a barrier (__syncthreads()), while different blocks run independently and may not even be live at the same time. The host (CPU) code allocates device memory with cudaMalloc(), copies inputs over with cudaMemcpy(), launches the kernel, and copies results back.

It matters because CUDA, more than anything else, made general-purpose GPU computing mainstream and remains the dominant ecosystem for deep learning and high-performance computing, with a deep stack of libraries on top. Be honest about two things. First, CUDA is proprietary and NVIDIA-only; portable alternatives like OpenCL and SYCL exist precisely because of that lock-in. Second, getting real speed is not automatic: you have to keep warps from diverging on branches, lay out memory so a warp's accesses coalesce into few transactions, and overlap data transfers with compute - 'I wrote a kernel' and 'it is fast' are very different milestones.

__global__ void add(const float *a, const float *b, float *c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; // this thread's index if (i < n) c[i] = a[i] + b[i]; // each thread does one element } // host: add<<<(n+255)/256, 256>>>(a, b, c, n); // grid of blocks of 256 threads

A CUDA kernel: each of thousands of threads computes its index and handles exactly one array element. The host launches a grid of blocks.

CUDA is proprietary and runs only on NVIDIA hardware - that is why portable options (OpenCL, SYCL) exist. And writing a kernel that compiles is easy; making it fast requires avoiding warp divergence, coalescing memory accesses, and hiding the CPU-GPU transfer cost - none of which the compiler does for you.

Also called
Compute Unified Device ArchitectureNVIDIA CUDA