CUDA Exam Mastery and Performance
Multidimensional Coordinate Mapping
From Lines to Grids
Many computational problems, like processing images or multiplying matrices, are naturally two-dimensional or even three-dimensional. Yet, a GPU's global memory is a single, continuous, one-dimensional line of addresses. This creates a fundamental challenge: how do we map a thread's position in a 2D or 3D grid to a unique spot in that flat memory space?
The solution is to calculate a unique linear index for each thread. We'll start with the most common case, a 2D grid, which is perfect for tasks involving images or matrices.
The 2D Mapping Formula
When you launch a 2D grid of threads, CUDA provides built-in variables like threadIdx and blockIdx with .x and .y components. This gives each thread a local coordinate within its block and a block coordinate within the grid. To find a thread's unique global position, we need to combine these coordinates.
First, we calculate the thread's global column index (globalX) and global row index (globalY).
int globalX = blockIdx.x * blockDim.x + threadIdx.x;
int globalY = blockIdx.y * blockDim.y + threadIdx.y;
Now we have a unique (globalX, globalY) coordinate for every thread. The next step is to map this 2D coordinate to a single 1D index. Most programming languages, including C and C++, store multidimensional arrays in what's known as . Imagine reading a book: you read all the characters in the first row from left to right, then move to the next row and do the same. Memory is accessed in the same pattern.
To get the linear index, we count how many full rows we've passed and then add our column position within the current row.
The alternative is , where data is stored column by column. This is common in languages like Fortran and MATLAB, but for CUDA C++, you'll almost always use row-major.
Expanding to 3D
The logic extends cleanly to three dimensions, which is useful for volumetric data or simulations in 3D space. First, we find the global (x, y, z) coordinates for each thread.
int globalX = blockIdx.x * blockDim.x + threadIdx.x;
int globalY = blockIdx.y * blockDim.y + threadIdx.y;
int globalZ = blockIdx.z * blockDim.z + threadIdx.z;
To linearize this 3D coordinate, we think of the data as a stack of 2D slices. We first calculate how many full slices we are from the front, then find our position within the current slice using the 2D formula.
Handling Boundaries
It's rare that your data's dimensions will be a perfect multiple of your block dimensions. For example, you might want to process a 1000x1000 image using 16x16 thread blocks. This mismatch means your kernel will launch threads that fall outside the actual data boundary.
If these extra threads try to access memory, they will read or write out of bounds, leading to errors or corrupted data.
The solution is a simple but critical boundary check. After calculating a thread's global coordinates, we add a guard to ensure it's within the valid data range before doing any work.
__global__ void processData(float* data, int width, int height) {
// Calculate global thread coordinates
int globalX = blockIdx.x * blockDim.x + threadIdx.x;
int globalY = blockIdx.y * blockDim.y + threadIdx.y;
// Boundary check
if (globalX < width && globalY < height) {
// Calculate linear index
int index = globalY * width + globalX;
// Do work on data[index]
data[index] = ...;
}
}
This check ensures that only threads corresponding to actual data points perform computations, preventing memory errors and making your kernels robust.
What is the primary reason for calculating a linear index for each thread in a CUDA kernel?
In a 2D grid processing an image of width 1920, a thread has the global coordinate globalX = 100 and globalY = 5. Using standard row-major order, what is its linear index?
With these mapping techniques, you can efficiently apply the massive parallelism of a GPU to complex, multidimensional problems.