CUDA is NVIDIA's platform for running general-purpose code on GPUs. You write a function called a kernel, and the GPU runs it in thousands of threads at once. It consists of a language extension to C and C++ (compiled with nvcc), a runtime library, and a driver. Where OpenMP parallelises across CPU cores that already share your address space, and MPI runs separate processes with separate memory, CUDA targets a separate device with its own memory. You have to move data there and back explicitly.

The execution model is SIMT (single instruction, multiple threads). A kernel launch creates a grid of thread blocks, and each block contains up to 1024 threads. Threads are scheduled in groups of 32 called warps, and every thread in a warp executes the same instruction at the same time. Unlike OpenMP, where the runtime picks the thread count for you, the launch geometry is yours to choose on every call.

__global__ void hello(void) {
    printf("block %d thread %d\n", blockIdx.x, threadIdx.x);
}
 
int main(void) {
    hello<<<2, 4>>>();       // 2 blocks of 4 threads = 8 threads total
    cudaDeviceSynchronize(); // wait for the GPU before exiting
    return 0;
}

The <<<blocks, threads>>> syntax is the kernel launch. Inside the kernel, blockIdx, blockDim, and threadIdx tell each thread which piece of the work it owns. The usual pattern is int i = blockIdx.x * blockDim.x + threadIdx.x, which flattens the hierarchy back into a single global index. Compile with nvcc -o prog prog.cu.

A kernel launch is asynchronous. Control returns to the CPU immediately, before the GPU has run anything, which is why the cudaDeviceSynchronize() above is needed for the printf output to appear. This is also the source of the most common beginner bug: a kernel that fails to launch reports no error at the launch site, so a broken program can exit with status 0 and no output at all. See CUDA error handling.

A GPU is not faster at any individual operation. It wins by having enormous latency-hiding capacity, so it needs enough parallel work to keep thousands of threads resident. Below that threshold the CPU is usually faster, and for small problems the PCIe transfer alone can cost more than the whole computation. Whether a kernel is limited by arithmetic or by memory bandwidth is the question the Roofline model exists to answer, and on GPUs the answer is almost always memory.

1. Kernels
2. Thread hierarchy
3. Warps
4. Warp divergence
5. Memory model
6. Global memory
7. Shared memory
8. Constant memory
9. Unified memory
10. Memory transfer
11. Coalescing
12. Bank conflicts
13. Synchronization
14. Atomics
15. Reduction
16. Streams
17. Events
18. Occupancy
19. Error handling
20. Tensor cores
21. Example
22. Overview