# CUDA Overview

Quick reference for [[cuda|CUDA]] qualifiers, launch syntax, the runtime API, and tooling.

## Qualifiers and launch syntax

```c
__global__ void f(...)          // kernel: called from host, runs on device
__device__ int  g(...)          // device function: callable only from device code
__host__   int  h(...)          // host function (the default)
__shared__ float tile[256];     // per-block scratchpad, shared by threads in a block
__constant__ float coeff[16];   // read-only, broadcast-optimised device memory

f<<<blocks, threads>>>(...)                  // basic launch
f<<<blocks, threads, shmem_bytes>>>(...)     // with dynamic shared memory
f<<<blocks, threads, shmem_bytes, stream>>>  // on a specific stream

threadIdx.x   // thread index within its block
blockIdx.x    // block index within the grid
blockDim.x    // threads per block
gridDim.x     // blocks per grid

__syncthreads()   // barrier across all threads in a block
```

## Runtime API

```c
cudaMalloc(&ptr, bytes)              // allocate device memory
cudaMallocManaged(&ptr, bytes)       // allocate unified memory, migrated on demand
cudaFree(ptr)                        // release device memory
cudaMemcpy(dst, src, bytes, kind)    // copy; kind is cudaMemcpyHostToDevice etc.
cudaMemcpyAsync(...)                 // same, queued on a stream

cudaDeviceSynchronize()              // block the host until the device is idle
cudaStreamCreate(&stream)            // create a stream for concurrent work
cudaStreamSynchronize(stream)        // block until one stream drains

cudaEventCreate(&ev)                 // create a timing event
cudaEventRecord(ev, stream)          // mark a point in the stream
cudaEventElapsedTime(&ms, a, b)      // milliseconds between two recorded events

cudaGetLastError()                   // fetch and clear the last error
cudaGetErrorString(err)              // human-readable message for an error code
cudaGetDeviceProperties(&prop, dev)  // query core count, bandwidth, limits
```

## Environment and tooling

^ Name ^ Kind ^ Description ^
| `CUDA_VISIBLE_DEVICES` | env | Restrict which GPUs the process can see, e.g. `0,2` |
| `CUDA_LAUNCH_BLOCKING` | env | Set to `1` to make launches synchronous, so errors surface at the launch site |
| `-arch=sm_XX` | nvcc | Target a specific compute capability instead of the default |
| `--ptxas-options=-v` | nvcc | Print per-kernel register and shared memory usage |
| `-lineinfo` | nvcc | Emit line numbers for profiler and sanitiser output |
| `nvidia-smi` | tool | Device inventory, utilisation, memory use, running processes |
| `ncu` | tool | Nsight Compute: per-kernel counters, occupancy, bottleneck analysis |
| `nsys` | tool | Nsight Systems: timeline of kernels, transfers, and host activity |
| `compute-sanitizer` | tool | Race and out-of-bounds checker, the CUDA equivalent of Valgrind |