CUDA Programming Model: Grids, Blocks, Threads, and Hardware Mapping

Introduction

Welcome back. In the previous module, we worked our way down the hardware stack, from the high-level differences between CPUs and GPUs to the specifics of the memory systems they rely on. You learned about throughput-oriented design, the SIMT execution model, and the critical trade-offs between memory technologies like HBM and GDDR.

Now, we pivot from the "what" of the hardware to the "how" of programming it. This lesson begins our exploration of the CUDA programming model, which provides the software abstractions needed to command the massive parallelism of a GPU. Our goal is to answer a fundamental question: How do you express a problem so that thousands of hardware cores can work on it simultaneously?

Your learning outcome for this lesson is to describe the core concepts of the CUDA programming model: grids, blocks, and threads, and their mapping to the hardware. We will connect the software abstractions you write in code directly to the hardware components we've already studied, like the Streaming Multiprocessor (SM).

1. A Tale of Two Processors: Host and Device

The CUDA programming model begins with the assumption of a heterogeneous system. This means your application code runs across at least two distinct types of processors:

• The Host: The CPU and its main memory (host memory). Your application always starts and ends its execution here. The host is responsible for overall control flow, I/O, and launching work on the GPU.
• The Device: The GPU and its dedicated memory (device memory, e.g., HBM or GDDR). The device acts as a powerful coprocessor, executing computationally intensive parallel tasks.

A function written to be executed on the device is called a kernel. The act of telling the GPU to run a kernel is called a kernel launch. When the host launches a kernel, it's not just starting one function call; it's instructing the device to execute that function potentially millions of times in parallel. This army of parallel executions is organized through a clear hierarchy.

2. The Software Hierarchy: Grids, Blocks, and Threads

To manage the complexity of launching millions of parallel operations, CUDA organizes them into a three-level hierarchy. This structure is what you, the programmer, will use to define the scope and organization of the parallel work.

Let's break down each level:

• Thread: The most granular unit of execution. You write the kernel code from the perspective of a single thread. Each thread has a unique ID and will execute the same code but can operate on different data.
• Thread Block (or Block): A group of threads. Threads within the same block can cooperate by sharing data through a fast, on-chip memory space (called "shared memory") and synchronizing their execution. A block is a 1D, 2D, or 3D arrangement of threads.
• Grid: A collection of all the thread blocks for a single kernel launch. The grid is also a 1D, 2D, or 3D arrangement of blocks.

This hierarchy isn't just a logical grouping; it maps directly to how the GPU hardware is designed and how it executes code. To see how these concepts translate into actual code, let's watch a segment from a Stanford lecture on parallel computing.

As the video explains, when you launch a kernel, you specify the dimensions of the grid (the number of blocks) and the dimensions of the blocks (the number of threads per block). Inside the kernel, each thread can then compute its unique global index using built-in variables. For a simple 1D grid of 1D blocks, this looks like:

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

This single line is the key to partitioning a large problem. Each of the thousands of threads executes this same line of code but gets a different idx, allowing it to work on a unique element of an array, for example.

3. Mapping the Software Model to GPU Hardware

Now for the crucial part of this lesson: connecting these software abstractions to the physical hardware we've been studying. Understanding this mapping is essential for reasoning about performance.

Synthesizing from the guide and other resources, here are the key mappings:

Grid ➔ Device

A grid represents the total workload for a kernel. The GPU's hardware scheduler is responsible for distributing all the thread blocks in the grid across all the available Streaming Multiprocessors (SMs) on the device. This is what makes CUDA code scalable: the same code launching a large grid can run on a GPU with 8 SMs or 80 SMs. The hardware automatically keeps its resources fed with work from the grid.

Thread Block ➔ Streaming Multiprocessor (SM)

This is the most important mapping to understand. All threads within a single thread block are guaranteed to execute on a single SM. This has profound implications:

1. Cooperation: Because they are on the same physical processor core, threads in a block can efficiently share data using the SM's dedicated shared memory.
2. Synchronization: Threads in a block can synchronize their execution. A call to __syncthreads() acts as a barrier, ensuring that no thread in the block proceeds past that point until all threads have reached it.
3. Independence: There is no guaranteed synchronization or execution order between different blocks. A block might run to completion before another one even starts. This means thread blocks must be independent units of work that do not rely on each other.

To visualize the hardware that executes a thread block, look at the following diagram of an SM.

Thread ➔ Warp ➔ CUDA Core

The mapping for an individual thread is more nuanced. While you program a single thread, the hardware executes threads in groups of 32 called warps.

• A warp is the fundamental unit of scheduling and execution on an SM.
• All 32 threads in a warp execute the same instruction at the same time in lock-step. This is the hardware implementation of the Single Instruction, Multiple Thread (SIMT) model we discussed in Module 1.

When the warp scheduler on the SM picks an instruction to execute, it issues it to an entire warp. If the instruction is a floating-point addition, 32 CUDA cores will perform that addition for the 32 threads in the warp. This is how GPUs achieve such massive parallelism.

The video below explains this hardware execution model, including the concept of "implicit SIMD" and how the GPU scheduler uses many active warps to hide latency.

A critical consequence of this warp-based execution is warp divergence. If an if-else statement causes threads within the same warp to take different paths, the hardware must execute both paths serially. While one group of threads executes the if block, the other threads in the warp are idle (masked off), and vice-versa for the else block. This effectively halves the utilization of the hardware for that portion of code and is a common source of performance loss.

For a final clarification of the SIMT concept:

Conclusion

In this lesson, you have learned the fundamental programming abstractions of CUDA and how they map onto the GPU's physical hardware. This bridge between software and hardware is the cornerstone of effective GPU programming.

Key Takeaways:

• The CUDA model uses a host (CPU) to launch kernels on a device (GPU).
• Parallel work is organized into a hierarchy: Grids contain Thread Blocks, and Thread Blocks contain Threads.
• This software hierarchy maps directly to the hardware:
  • A Grid is a complete workload distributed by the GPU scheduler across all SMs on the device.
  • A Thread Block is a unit of cooperation whose threads all run on a single SM, allowing them to use shared memory and synchronize.
  • Threads are executed in groups of 32 called Warps, which are the fundamental unit of scheduling and execute in a SIMT fashion.

Preview of the Next Lesson:

Now that you understand the "what" and "where" of thread execution, the next logical questions are "how many?" and "what if?". In the next lesson, we will delve into the performance implications of this model. We will explore the concepts of kernel launches, thread divergence, and occupancy, and learn how to choose the right grid and block sizes to maximize the utilization of the GPU's powerful hardware.