CUDA Programming Model: Kernels, Threads, Blocks, and Grids

A CUDA program runs on two processors at once:

• The host is the CPU. It runs main(), manages files, allocates GPU memory, and launches work.
• The device is the GPU. It runs the parallel kernels you write.

They have separate memory spaces (host RAM vs device VRAM) and separate execution. Host code calls into device code through kernel launches, then continues asynchronously while the GPU works.

int main() {
  // host code here
  myKernel<<<grid, block>>>(d_arr);  // launch on device
  cudaDeviceSynchronize();           // wait for device
  // host code again
}

This split is the source of most CUDA bugs. Forget which side a pointer belongs to and you'll segfault or read garbage.

CUDA extends C++ with three function qualifiers that say where a function lives:

• __global__: a kernel. Called from the host, executes on the device. Must return void. This is the entry point for GPU work.
• __device__: a helper function callable only from device code. Inlined aggressively.
• __host__: a normal CPU function (the default if you write nothing).

You can stack __host__ __device__ to make a function compilable for both sides — useful for math utilities.

__device__ float square(float x) { return x * x; }

__global__ void apply(float* a, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) a[i] = square(a[i]);
}

Every kernel launch creates a grid. A grid is divided into blocks. Each block contains many threads. Three levels, three reasons:

• Threads are the unit of execution. Each runs the same kernel code with its own indices.
• Blocks are the unit of cooperation. Threads in the same block can share fast on-chip memory and synchronise with __syncthreads().
• Grids are the unit of the whole launch. Blocks don't talk to each other directly.

A block has up to 1024 threads (hardware limit). A grid can have billions of blocks. Pick blocks of 128-512 threads in practice — multiples of 32 (the warp size) so no warp is partial.

Inside a kernel, every thread sees four built-in variables that tell it which slice it owns:

• threadIdx.x/.y/.z — its position inside its block.
• blockIdx.x/.y/.z — its block's position inside the grid.
• blockDim.x/.y/.z — the size of each block.
• gridDim.x/.y/.z — the size of the grid (in blocks).

The canonical 1D index pattern:

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) out[i] = in[i] * 2.0f;

The if (i < N) guard matters: you almost always launch slightly more threads than you need (so the count divides evenly into blocks). The guard prevents out-of-bounds writes from the extras.

How the three levels nest.

CUDA borrows triple-angle-bracket syntax for kernel launches. The first argument is the grid shape, the second is the block shape.

// 1D: 1M elements, 256 threads per block
int N = 1 << 20;
int threads = 256;
int blocks  = (N + threads - 1) / threads;  // ceil-div
myKernel<<<blocks, threads>>>(d_arr, N);

Grid and block sizes can be int (1D) or dim3 (2D/3D). The ceiling-division trick (N + threads - 1) / threads is everywhere in CUDA code — it guarantees you cover all N elements even when N isn't a multiple of threads. Pair it with the if (i < N) guard inside the kernel.

When your data is 2D (images) or 3D (medical volumes, simulations), you can launch a matching grid. The hardware doesn't care — it's pure addressing convenience.

__global__ void blur(float* img, int W, int H) {
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  int y = blockIdx.y * blockDim.y + threadIdx.y;
  if (x >= W || y >= H) return;
  img[y * W + x] = /* ... */;
}

dim3 block(16, 16);                      // 256 threads
dim3 grid((W + 15) / 16, (H + 15) / 16);
blur<<<grid, block>>>(d_img, W, H);

The 16×16 block is a classic choice for image work — square, 256 threads, evenly divides most resolutions.

Threads in a block are not actually scheduled individually. The hardware bundles them into warps of 32 consecutive threads (by linearised index) and schedules warps. Two consequences:

• Pick block sizes that are multiples of 32. A block of 100 threads still uses 4 warps of 32 — the last 28 lanes are wasted on every instruction.
• Threads in a warp execute the same instruction at the same time. If they branch differently, the hardware serialises the paths (warp divergence).

Warps don't appear in the source code — you write threads, the hardware groups them. But your block size choice and your branch patterns are really warp-level decisions in disguise.

Putting the pieces together — kernel, launch, indices, guard:

__global__ void scaleByTwo(float* a, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) a[i] *= 2.0f;
}

int main() {
  const int N = 1'000'000;
  float* d_a;
  cudaMalloc(&d_a, N * sizeof(float));
  // ... copy input into d_a ...

  int threads = 256;
  int blocks  = (N + threads - 1) / threads;
  scaleByTwo<<<blocks, threads>>>(d_a, N);
  cudaDeviceSynchronize();

  cudaFree(d_a);
}

That's the skeleton of every basic CUDA program: allocate device memory, launch a kernel where each thread handles one element, synchronise, free.