No history yet

1D Indexing Formulas

Finding Your Place in the Grid

When you launch a CUDA kernel, you unleash thousands of threads. To do useful work, each thread needs to answer a simple question: "Which piece of data am I supposed to process?" The answer lies in its unique global index.

Imagine a large array of data in memory. We want to assign each thread to a single element in that array. To do this, every thread calculates its position within the entire grid of threads. This calculation relies on three built-in variables provided by the CUDA runtime.

Every GPU thread needs to know which element of the array it should process.

Let's think of the grid of threads like seats in a movie theater arranged in long, single rows.

  1. threadIdx.x: This is your seat number within your specific row. It's a local ID, starting from 0. If you're in the 5th seat, your threadIdx.x is 4.
  2. blockIdx.x: This is the row number itself. It’s a unique ID for the block (the row) you're in, also starting from 0. The first row is blockIdx.x = 0, the second is blockIdx.x = 1, and so on.
  3. blockDim.x: This tells you how many seats are in any given row. This value is the same for every block in your grid. For example, every row might have 256 seats.

The Indexing Formula

With these three variables, any thread can find its unique global index in a 1D grid using a straightforward formula.

intindex=blockIdx.x∗blockDim.x+threadIdx.x;int index = blockIdx.x * blockDim.x + threadIdx.x;

Let’s break it down with our theater analogy. Suppose each row (blockDim.x) has 100 seats. If you are in row #3 (blockIdx.x = 2) and seat #15 (threadIdx.x = 14), your global seat number is calculated as:

index = 2 * 100 + 14 = 214

This means you are the 215th person in the theater (since the index is 0-based).

Putting It into Practice

This formula is the heart of most simple CUDA kernels. In code, it allows each thread to pick its unique element from an array stored in global memory. Here’s a basic kernel that adds two vectors, a and b, and stores the result in c.

__global__ void add_vectors(float *a, float *b, float *c, int n) {
    // Calculate the global thread index.
    int index = blockIdx.x * blockDim.x + threadIdx.x;

    // A 'stride' is sometimes used for more complex scenarios,
    // but for a simple 1:1 mapping, it's just the total number of threads.
    int stride = gridDim.x * blockDim.x;

    // Use a loop with the stride to ensure all elements are processed,
    // even if there are more elements than threads.
    for (int i = index; i < n; i += stride) {
        c[i] = a[i] + b[i];
    }
}

The key line is int index = blockIdx.x * blockDim.x + threadIdx.x;. Once calculated, index is used to access the correct elements in the arrays a, b, and c. The for loop is a common pattern that makes the kernel more robust, allowing you to process arrays of any size, even if they're larger than the total number of threads in your grid.

By ensuring every thread operates on a different index, you achieve massive parallelism. Each thread performs its addition independently and simultaneously, which is the source of the GPU's immense processing power(). Mastering this simple indexing formula is the first major step in writing effective CUDA programs.

Quiz Questions 1/4

What is the primary purpose of calculating a unique global index for each thread in a CUDA kernel?

Quiz Questions 2/4

A kernel is launched with blocks containing 512 threads each. A thread has a blockIdx.x of 10 and a threadIdx.x of 20. What is its global index?