# CUDA
**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.

```c
__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.

## Concepts

 1. [[cuda-kernels|Kernels]]
 2. [[cuda-thread-hierarchy|Thread hierarchy]]
 3. [[cuda-warps|Warps]]
 4. [[cuda-warp-divergence|Warp divergence]]
 5. [[cuda-memory-model|Memory model]]
 6. [[cuda-global-memory|Global memory]]
 7. [[cuda-shared-memory|Shared memory]]
 8. [[cuda-constant-memory|Constant memory]]
 9. [[cuda-unified-memory|Unified memory]]
 10. [[cuda-memory-transfer|Memory transfer]]
 11. [[cuda-coalescing|Coalescing]]
 12. [[cuda-bank-conflicts|Bank conflicts]]
 13. [[cuda-synchronization|Synchronization]]
 14. [[cuda-atomics|Atomics]]
 15. [[cuda-reduction|Reduction]]
 16. [[cuda-streams|Streams]]
 17. [[cuda-events|Events]]
 18. [[cuda-occupancy|Occupancy]]
 19. [[cuda-error-handling|Error handling]]
 20. [[cuda-tensor-cores|Tensor cores]]
 21. [[cuda-example|Example]]
 22. [[cuda-overview|Overview]]