Understanding GPU Thread Hierarchy

The CUDA Hierarchy of Parallelism

When you launch a CUDA program, you are not manually assigning individual tasks to the thousands of cores. Instead, you define a single function, called a kernel, and instruct the GPU to execute it using a massive number of threads. To manage this army of threads, CUDA organizes them into a three-tiered hierarchy.

Let's begin with a video that introduces these concepts at a high level.

As the video explains, the hierarchy consists of:

1. Grid: A grid encompasses all threads for a given kernel launch. It is the largest unit in the hierarchy and can be thought of as the entire computational task you've sent to the GPU. A grid is composed of a collection of thread blocks.
2. Thread Block: A block is a group of threads. The key feature of a block is that all threads within it are guaranteed to execute on the same Streaming Multiprocessor (SM). This co-residency allows them to cooperate by sharing data through a fast, on-chip memory (called shared memory, which we'll cover soon) and synchronizing their execution.
3. Thread: A thread is the smallest and most fundamental unit of execution. Each thread executes an instance of your kernel function. It has its own private set of registers and local memory.

The warp, which we learned about in the previous lesson, fits neatly into this model. A warp is a group of 32 threads within a thread block that execute the same instruction at the same time (the SIMT model). The hardware schedules and executes work at the warp level.

This hierarchy can be visualized as a nested structure.

Mapping Software Abstractions to Physical Hardware

Your interest in low-level engineering makes it vital to connect this software model back to the GPU architecture we previously studied. The mapping is direct and intuitive:

• A Kernel Grid is launched on the entire GPU Device. The CUDA runtime is responsible for distributing the blocks of the grid to the available SMs on the device.
• A Thread Block is assigned to and runs its entire lifetime on a single SM. This is a critical constraint. An SM can often run multiple blocks concurrently if it has sufficient resources (registers, shared memory).
• A Warp within a block is the fundamental unit of scheduling managed by the SM's Warp Schedulers.
• A Thread is logically processed by an execution unit like a CUDA Core.

This software-to-hardware mapping is the foundation of GPU performance.

For a deeper dive into this connection and its performance implications, the following reading is excellent.

As the text highlights, since a thread block resides on a single SM, and an SM has multiple warp schedulers, you must launch blocks with enough threads (i.e., enough warps) to keep all those schedulers busy. This is a primary principle for achieving high hardware utilization.

Navigating the Hierarchy: Thread Indexing

To perform useful work, each thread needs a unique identity. For example, if you are processing an image, you need a way to assign each thread to a specific pixel. CUDA provides this capability through a set of built-in variables that a thread can query from within a kernel.

The organization of grids and blocks can be one, two, or three-dimensional. This is purely for programmer convenience and makes it easy to map threads to multi-dimensional data structures like matrices (2D) or volumes (3D).

The key built-in variables are:

• gridDim: The dimensions of the grid (how many blocks in each dimension).
• blockDim: The dimensions of the thread block (how many threads in each dimension).
• blockIdx: The index of the current thread block within the grid.
• threadIdx: The index of the current thread within its block.

Each of these is a dim3 type, meaning it has .x, .y, and .z components. For a 1D grid of 1D blocks, you would only use the .x component.

The most common task inside a kernel is to use these variables to compute a single, unique global index for each thread. In a 1D case, the pattern is simple and fundamental:

int global_thread_id = blockIdx.x * blockDim.x + threadIdx.x;

Let's break this down:

• blockIdx.x * blockDim.x: This calculates the starting index of the current block. For example, if each block has 256 threads (blockDim.x = 256), then block 0 starts at index 0, block 1 starts at index 256, block 2 starts at index 512, and so on.
• + threadIdx.x: This adds the thread's local index within its block to find its unique position within the entire grid.

The following video provides an excellent, step-by-step walkthrough of this exact calculation.

For 2D data, the logic extends naturally to both dimensions:

int global_idx_x = blockIdx.x * blockDim.x + threadIdx.x;
int global_idx_y = blockIdx.y * blockDim.y + threadIdx.y;

This allows a thread to easily find its corresponding (x, y) coordinate in a matrix or image, which is a pattern you will see in almost every 2D kernel.

Conclusion

In this lesson, we have defined the essential software abstractions that allow programmers to command the immense power of a GPU. You now have a mental model that connects the logical thread hierarchy to the physical hardware.

Key Takeaways:

• CUDA programs organize work using a hierarchy of grids, thread blocks, warps, and threads.
• A grid represents the entire kernel launch, a block is a group of cooperating threads, a warp is the unit of hardware scheduling, and a thread is the smallest unit of execution.
• This software hierarchy maps directly to the hardware: grids run on the GPU, blocks on a single SM, and warps are managed by warp schedulers.
• Threads use built-in variables (blockIdx, threadIdx, etc.) to compute a unique global index, which allows them to partition data and work independently. This indexing is the cornerstone of writing correct and scalable kernels.

In our previous lesson, we saw the physical processors (SMs). Today, we saw how work is organized for them (the thread hierarchy). In our next lesson, we will explore the final piece of the core puzzle: the GPU memory hierarchy. We'll investigate the different types of memory available to these threads—from ultra-fast registers to high-capacity global memory—and understand how their characteristics dictate kernel performance.