Introduction
In our last lesson, you successfully wrote and launched your first CUDA kernel, performing vector addition by launching a one-dimensional grid of threads. You learned how to assign a unique task to each thread using a simple formula to calculate a global ID from blockIdx.x and threadIdx.x. This was a crucial first step into the world of parallel computation.
This lesson builds directly on that foundation. We will generalize the concept of thread indexing to handle multi-dimensional data structures. Many problems in science, engineering, and especially AI—such as image processing or matrix multiplication—are inherently two- or three-dimensional. CUDA is designed to map its thread hierarchy naturally onto these problems. By the end of this lesson, you will be able to launch multi-dimensional grids and use CUDA's built-in variables to compute unique thread indices for both 1D and 2D layouts, setting the stage for more complex and realistic parallel algorithms.
From 1D to Multi-dimensional Grids
In the previous lesson, we treated the execution configuration <<<numBlocks, threadsPerBlock>>> as if it only took single integer values. In reality, these parameters, along with the built-in variables inside the kernel, are more flexible. CUDA uses a special dim3 struct type to define one, two, or three-dimensional configurations for both grids of blocks and blocks of threads.
The four built-in variables you've encountered are instances of these multi-dimensional types:
gridDim: The dimensions of the grid (e.g.,gridDim.x,gridDim.y,gridDim.z).blockDim: The dimensions of a thread block.blockIdx: The 3D index of the current block within the grid.threadIdx: The 3D index of the current thread within its block.
The following video provides an excellent overview of this concept, explaining how to use dim3 to launch kernels and how the built-in variables give each thread its unique coordinates in this multi-dimensional space.
This video, "CUDA Programming" by Tom Nurkkala, clearly explains the multi-dimensional nature of CUDA's execution model.
Please watch two segments: First, watch the part on runtime configuration, which introduces the dim3 data type and how it's used to define 1D, 2D, or 3D grids and blocks. Then, watch the crucial section on built-in variables. This explains the four variables (gridDim, blockDim, blockIdx, threadIdx) that kernels can access to understand their position in the execution hierarchy.
Thread Indexing: From 1D to 3D
With this multi-dimensional perspective, we can now derive the formulas to compute a unique, linear index for any thread, regardless of the grid's dimensionality. The core challenge is always the same: mapping a multi-dimensional coordinate (z, y, x) to a single offset in the computer's one-dimensional memory space. This process is often called "flattening" an index.
1D Grid of 1D Blocks (Review)
This is the case you're already familiar with from our vector addition example. The grid and blocks have dimensions only in x.
The formula is a direct application of the principle of calculating an offset:
2D Grid of 2D Blocks
Now, let's consider a more interesting case, such as processing a matrix or an image. It is natural to launch a 2D grid of 2D blocks to match the data's structure.
To process a matrix element at (row, col), a thread first needs to compute its unique global row and column. The logic is a straightforward extension of the 1D case, applied to both dimensions:
- Global Column:
int col = blockIdx.x * blockDim.x + threadIdx.x; - Global Row:
int row = blockIdx.y * blockDim.y + threadIdx.y;
The logic is identical: blockIdx.y * blockDim.y computes the starting row index of the current block, and threadIdx.y provides the thread's local row offset within that block.
Once a thread knows its (row, col) coordinates, it can compute the linear index to access the corresponding element in memory. For a matrix of width W stored in row-major order (the standard in C/C++), the formula is:
The following video walks through exactly this scenario for matrix multiplication, showing how to set up the 2D grid and use the indices inside the kernel.
CUDA Crash Course: Matrix Multiplication
Nick's "CUDA Crash Course: Matrix Multiplication" provides a perfect practical example of 2D indexing.
Watch from the introduction to 2D indexing. Pay close attention to how he visualizes the 2D grid and derives the formulas for the global row and column. The concrete example he works through is extremely helpful for building intuition.
A General and Performant Approach to Indexing
While calculating row and col separately and then combining them is intuitive, it can sometimes hide performance pitfalls related to memory access patterns. For your goal of becoming a kernel engineer, it's essential to understand the most robust and performant way to calculate a flattened index from first principles.
The key is to think of the 6D CUDA coordinate space (gridDim.z, gridDim.y, gridDim.x, blockDim.z, blockDim.y, blockDim.x) as a set of nested loops. The fastest-moving index in memory should correspond to the innermost loop, which for CUDA is threadIdx.x. A correct indexing formula ensures that adjacent threads in a warp (e.g., threads 0, 1, 2...) access adjacent memory locations. This is called coalesced memory access, a critical topic we will explore in the next module.
A general formula that respects this ordering can be constructed by repeatedly applying the rule new_offset = old_offset * dimension + index, starting from the slowest-moving dimension and working towards the fastest.
For a 2D grid of 2D blocks, the fully coalesced formula looks like this:
This expression looks intimidating, but it is just a mechanical nesting of the simple 1D formula. Each parenthesized group calculates a linear index at a higher level of the hierarchy, which then becomes the "block index" for the next level down.
The following resources provide a deep dive into this generalized approach, which will be invaluable for your understanding of low-level performance.
GPU Thread Hierarchy and Indexing |
This article provides clear tables and formulas for various grid and block configurations.
Focus on the section Thread Indexing - 2D Grid and 2D Block. Compare the formulas presented there with the nested formula I've shown above. You'll see they are mathematically equivalent ways of calculating a flattened index. The article breaks it into blockId and local threadId first, which can be easier to reason about.
The next resource introduces this "folding" concept from a C++ template metaprogramming perspective, which aligns with your interest in systems programming. It confirms that this nested approach is not only correct but also generates optimal code.
Folding a Global Linear Id from CUDA Dimensions and Indices
This blog post by Paul Keir offers a principled way to think about index calculation, connecting it to functional programming concepts and low-level performance.
Please read the following parts: Start from the beginning of the article to understand the motivation and the problem of calculating a global linear ID correctly. Read the section that explains why the simpler row * width + col approach can be bad for performance. This is a crucial insight. Finally, study the code in the section showing the CUDA C++ implementation. The cuda_global_linear_id<2>() function shows the exact nested formula we discussed, generalized for optimal performance.
To see this all come together in a comprehensive code example, we can return to the first video, which features a "who am I" kernel that calculates a unique ID in a full 3D grid, printing out all the intermediate steps.
Let's conclude with a detailed code walkthrough from the "CUDA Programming" video.
Watch the final, in-depth example from the "who am I" kernel. The presenter implements a 3D version of the nested indexing formula. He explains how blockId and threadId are calculated by "flattening" their respective 3D indices, and then combined to form the final global ID. The non-deterministic output he shows is also an important reminder that you cannot assume any execution order between blocks.
Conclusion
In this lesson, you have moved beyond simple 1D grids and learned how to map CUDA's execution hierarchy onto multi-dimensional problems. You can now configure and launch 1D and 2D grids and, more importantly, you understand the arithmetic required for each thread to calculate its unique global index within that grid.
Key Takeaways:
- Kernel launches can be configured with 1D, 2D, or 3D dimensions using the
dim3type. - The built-in variables
threadIdx,blockIdx,blockDim, andgridDimaredim3structs that provide each thread with its coordinates. - For 2D problems, you can compute a thread's global
(row, col)using the formulasrow = blockIdx.y * blockDim.y + threadIdx.yandcol = blockIdx.x * blockDim.x + threadIdx.x. - A general, nested formula
new_offset = old_offset * dimension + indexprovides a robust and performant way to calculate a "flattened" linear index for any grid dimensionality. - The order of operations in this formula is critical for achieving coalesced memory access, a cornerstone of GPU performance optimization.
In our next lesson, we will focus on robust error handling for CUDA API calls and kernel launches. Writing correct parallel code is challenging, and having a systematic way to detect and diagnose errors is a non-negotiable skill for any serious GPU programmer.