> For the complete documentation index, see [llms.txt](https://handbook.modular.com/llms.txt).
> Markdown versions of all pages are available by appending .md to any URL.

import LinkList from '@site/src/components/LinkList'; import
WarpDivergenceVisualizer from '@site/src/components/WarpDivergenceVisualizer';

# GPU threads, warps, blocks, and grids

A GPU kernel exposes a large amount of parallel work. The execution model
organizes that work into a hierarchy of threads, warps, thread blocks, and
grids. Each level answers a different question: what one logical worker does,
which workers execute together, which workers can cooperate, and how the full
kernel launch covers the input.

## Threads

A **thread**, also known as a **work unit** on AMD GPUs, is the smallest logical
unit of execution within a kernel function. Every thread runs the same kernel
code, but built-in coordinates let each thread select different data. This model
is called SIMT: Single Instruction, Multiple Threads.

You don’t manually assign each thread what to do. Instead, each thread typically
combines the position within a block with the block position to compute which
part of the input to process. For more information, see
[grids and indexing](#grids-and-indexing) below.

The same kernel body can process a vector with thousands or millions of
elements. Threads can also use two- or three-dimensional coordinates, which map
naturally to matrices, images, and tiled computations.

## Warps

A **warp**, also known as a **wavefront** on AMD GPUs, is a subset of threads
from a thread block that execute together. On NVIDIA GPUs a warp is 32 threads.
AMD wavefronts are traditionally 64, and 32 or 64 on RDNA.

The warp is the actual scheduling unit. It is what gets presented to a warp
scheduler, and a warp can only be issued one instruction at a time. When a warp
receives an instruction, the active threads of it each execute that instruction
on their own registers and data. Different warps are independent and can execute
different instructions at the same time.

Threads in a warp don't need to produce the same result, but control flow is
most efficient when they follow the same path. If a branch sends some threads
one way and others another way, the warp executes the paths separately. Each
has a different active-thread mask so that only the threads on the current path
are active. This is called **warp divergence**, and it lowers efficiency when
the divergent paths contain substantial work, since threads off the current
path sit idle.

## Thread blocks

A **thread block (or block)**, also known as a **workgroup** on AMD GPUs, is a
subset of threads within a grid, which is the top-level organizational structure
of the threads executing a kernel function. As the primary building block for
workload distribution, thread blocks serve multiple crucial purposes:

1. They break down the overall workload (managed by the grid) of a kernel
   function into smaller, more manageable portions that can be processed
   independently. This division allows for better resource utilization and
   scheduling flexibility across multiple
   [streaming multiprocessors (SMs)](https://handbook.modular.com/kernel-optimization/gpu-architecture-fundamentals/streaming-multiprocessors.md)
   in the GPU.
2. Thread blocks provide a scope for threads to collaborate through
   [shared memory](https://handbook.modular.com/kernel-optimization/gpu-architecture-fundamentals/gpu-memory.md#shared-memory-smem-and-l1-cache)
   and synchronization primitives, enabling efficient parallel algorithms and
   data sharing patterns.
3. Thread blocks help with scalability by allowing the same program to
   run efficiently across different GPU architectures, as the hardware can
   automatically distribute blocks based on available resources.

You can specify the number of thread blocks in a grid and how they are arranged
across one, two, or three dimensions. Each block within the grid is assigned a
unique block index that determines the position within the grid. Similarly, you
also specify the number of threads per thread block and how they are arranged
across one, two, or three dimensions. Current NVIDIA and AMD GPUs commonly allow
up to 1024 threads per block, subject to device-specific dimension, register,
and shared-memory limits.

The GPU assigns each thread block within the grid to a streaming multiprocessor
(SM), where the block generally remains resident until it completes; blocks are
not migrated between SMs. Several blocks can reside on one SM at the same time
when registers, shared memory, and scheduler slots allow.

Threads within a block can share data through shared memory and synchronize
using built-in mechanisms, but they cannot directly communicate with threads in
other blocks.

## Grids and indexing

A **grid** is the top-level organizational structure that defines a kernel
launch—it contains all the thread blocks, which contain all the execution
threads.

To write a GPU kernel, you must specify the work to perform by creating a grid.
On an NVIDIA GPU, the CUDA programming model allows you to define the kernel’s
launch dimensions and coordinates with these values:

- `gridDim`: Grid dimensions, namely the number of blocks in each dimension.
- `blockDim`: Block dimensions, namely the number of threads in each dimension
  of a block.
- `blockIdx`: The current block position within the grid.
- `threadIdx`: The current thread position within the block.

A kernel combines these values, typically by computing a global thread ID, to
determine which element or portion of data each thread should process.

## Writing a kernel function

Now that you understand how GPU work is organized into threads, warps,
thread blocks, and grids, you can better understand kernel code.

Here’s a simple CUDA kernel that performs a vector addition, where every
thread handles one element:

```cpp
__global__ void vecAdd(float* A, float* B, float* C, int vectorLength)
{
    int workIndex = threadIdx.x + blockDim.x * blockIdx.x;
    if (workIndex < vectorLength)
    {
        C[workIndex] = A[workIndex] + B[workIndex];
    }
}
```

The bounds check here is important because the grid rarely divides the input
evenly. The launch itself uses an **execution configuration** between triple
angle brackets, where the first value sets the number of blocks in the grid and
the second sets the number of threads in each block:

```cpp
int threads = 256;
int blocks = (vectorLength + threads - 1) / threads;

vecAdd<<<blocks, threads>>>(devA, devB, devC, vectorLength);
```

The grid usually contains more blocks than the GPU can run at once. As blocks
finish, the GPU assigns new blocks to free SM capacity. This scheduling model
lets one launch scale across GPUs with different SM counts without changing the
kernel.

However, the CUDA code above only works for NVIDIA GPUs. If you want to
program an AMD GPU, you might instead use AMD’s native ROCm software. Both
frameworks use C/C++ as the core programming language, with some extra
keyword and syntax extensions.

Alternatively, [MAX](https://www.modular.com/open-source/max) provides
a hardware-agnostic programming model for GPUs and other accelerators,
using the [Mojo programming language](https://mojolang.org/). This
framework provides the same grid-and-block programming model described
above, except the code can target NVIDIA GPUs, AMD GPUs, Apple silicon,
and more. For example, here’s the same vector addition function written
in Mojo:

```mojo
from std.gpu import block_dim, block_idx, thread_idx
from max.gpu.host import DeviceContext

def kernel():
    var i = block_idx.x * block_dim.x + thread_idx.x
    # Process element i.

def main() raises:
    var ctx = DeviceContext()
    # Same geometry as the CUDA launch above.
    ctx.enqueue_function[kernel](https://handbook.modular.com/kernel-optimization/gpu-architecture-fundamentals/grid_dim=4096, block_dim=256.md)
    ctx.synchronize()
```

`grid_dim` and `block_dim` correspond to the two CUDA launch values, and the
index arithmetic is identical. Mojo also exposes `global_idx`, a shorthand that
computes `block_idx.x * block_dim.x + thread_idx.x` for you, so the same line
could be written as `var i = global_idx.x`.

:::note
The choice of programming model shows up when you customize kernels inside an
inference framework. vLLM and SGLang both build on the vendor toolchains. Their
kernels are CUDA C++ for NVIDIA and ROCm/HIP for AMD, alongside
[Triton](https://handbook.modular.com/kernel-optimization/kernel-optimization-tools.md#triton-and-python-kernel-dsls)
and other kernel libraries for some operations. A kernel you write for one
vendor generally needs a separate variant, and separate tuning, for the other.
MAX takes a different approach: the same Mojo kernel implementation can target
both NVIDIA and AMD GPUs.
:::

To learn about programming NVIDIA GPUs, see the
[CUDA Programming Guide](https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html).

To learn about programming any GPU, see the
[Mojo GPU programming guide](https://max.modular.com/gpu/intro-tutorial/).

## How the hierarchy affects kernel performance

Launch geometry changes how a kernel uses the hardware:

- A block often needs enough threads to provide several warps, but an oversized
  block can consume too many registers or too much shared memory.
- Adjacent threads should access adjacent global memory addresses when possible.
  This lets the GPU combine memory requests into efficient transactions.
- Branches should keep threads within a warp on the same path when the
  algorithm allows.
- A grid should expose enough independent blocks to keep all SMs busy.

The best geometry depends on the work per thread and the resources required by
each block. Resource use also affects
[occupancy](https://handbook.modular.com/kernel-optimization/gpu-architecture-fundamentals/streaming-multiprocessors.md#occupancy).
Fixed rules such as “always use 256 threads” are starting points, not universal
answers.