What are threadIdx, blockIdx, blockDim and gridDim in CUDA?

threadIdx, blockIdx, blockDim and gridDim are the four values every thread reads to work out which piece of the data it owns.

Four names, and the guide gives them two different types. threadIdx and blockIdx are uint3; blockDim and gridDim are dim3. Both are three unsigned values called .x, .y and .z, and the difference that matters is a default: "In C++11 and later, the default value of all components of dim3 is 1." That is why <<<8, 256>>> is legal when the parameters are dimension types. The integers convert to dim3, .y and .z come out 1, and a 1D launch is a 3D launch with two dimensions of size one. Nothing special-cases the 1D shape.

Two of the four vary per thread and two do not. threadIdx.x runs 0 to blockDim.x - 1 inside each block and restarts in the next one; blockIdx.x runs 0 to gridDim.x - 1 across the grid. blockDim and gridDim are your execution configuration handed back to you: every thread in the launch reads the same value, and you already knew it on the host. Treating them as four equal unknowns is what makes the formula feel arbitrary. There are two coordinates and two constants.

Which is why the most expensive typo in CUDA is gridDim where blockDim belongs. blockIdx.x * gridDim.x + threadIdx.x compiles, launches and returns cudaSuccess. Under <<<8, 256>>> it multiplies the block number by 8 instead of 256, so the largest index any of the 2048 threads can produce is 7 * 8 + 255, which is 311. Every thread piles into the first 312 elements of the array, 1688 elements are never touched, and most of the ones that are get written several times by threads that disagree. That is arithmetic, not a measurement, and you can check it on paper. What makes it survive review is the case where it cannot be caught: when the block count equals the block size, <<<256, 256>>>, gridDim.x and blockDim.x are the same number and the wrong formula gives the right answer.

The other trap is quieter. blockIdx.x and blockDim.x are both unsigned 32-bit, so their product is computed in 32 bits whatever you assign it to. Widen before the multiply, not after: blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x. Day 4's kernel does it that way for a reason day 8 makes concrete.

Measured

Numbers from the project's verification node, 2026-08-30: Tesla T4, compute capability 7.5, driver 595.84, toolkit CUDA 12.6 (V12.6.85), build line nvcc -O3 -arch=sm_75.

Day 4 ran a kernel over 2000 elements whose only job is to write its own blockIdx.x and threadIdx.x into the element it owns, so the output is the mapping itself rather than a result that could look plausible while being wrong. All four values, at element 1337, for the three launches it runs:

Launch	`gridDim.x`	`blockDim.x`	`blockIdx.x`	`threadIdx.x`
`<<<7, 256>>>`	7	256	5	57
`<<<8, 256>>>`	8	256	5	57
`<<<2, 1024>>>`	2	1024	1	313

The first two columns are the launch. The last two came off the device. Read the rows against each other and the split does the explaining: change gridDim.x alone, from 7 to 8, and element 1337 keeps the same owner, because the extra block sits past it. Change blockDim.x and the owner moves, because blockDim.x is the multiplier in the formula and gridDim.x is not.

The 7-block row is also the one that lost 208 elements, which no built-in variable can tell you about: a thread that was never started reads nothing. The check that catches it runs on the host.

Code

From code/day04-indexing/indexing.cu. The whole kernel, which reads three of the four built-ins and computes nothing else.

__global__ void writeOwner(int* blockOf, int* threadOf, size_t n) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        blockOf[i] = static_cast<int>(blockIdx.x);
        threadOf[i] = static_cast<int>(threadIdx.x);
    }
}

gridDim.x is the one it never reads, which is the point: a thread that indexes with the formula above does not need to know how many blocks were launched. The kernels that do need it are the ones with a grid-stride loop, where blockDim.x * gridDim.x is the step.

Diagram

index-tracer, preset 1d-basic, with the four values shown in a side panel that updates as you step: gridDim.x and blockDim.x fixed at the top, blockIdx.x and threadIdx.x changing with the selected thread.

Alt text: "A panel of four values beside a grid of twelve elements. Two of the values never change as you step from thread to thread; the two that do are the only thing that distinguishes one thread from another."

Related terms

Where you meet this

Day 4, grid, block and thread indexing, which owns this term and produced the table
Day 7, two-dimensional grids, where .y stops defaulting to 1
Day 8, bounds checks and grid-stride loops, the first kernel that needs gridDim
Day 21, what is a warp, for warpSize, the fifth built-in and the one with a different type
illegal memory access, where an index built from the wrong pair of these usually surfaces

Sources

CUDA Programming Guide, built-in variables: gridDim and blockDim are dim3, blockIdx and threadIdx are uint3, warpSize is int, and "In C++11 and later, the default value of all components of dim3 is 1": https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/cpp-language-extensions.html (checked 2026-08-30)
CUDA Programming Guide, Table 30, maximum block dimensions 1024 in x and y, 64 in z, and 1024 threads per block: https://docs.nvidia.com/cuda/cuda-programming-guide/05-appendices/compute-capabilities.html (checked 2026-08-30)
"CUDA, gridDim and blockDim", 120,836 views: https://stackoverflow.com/questions/16619274/cuda-griddim-and-blockdim (checked 2026-08-29)
CUDA Programming Guide, writing CUDA kernels, on launch configuration and block scheduling: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html (checked 2026-08-30)

Byline

Written by: unassigned. Reviewed by: unassigned. Written on: not set. Last checked: not set. Numbers captured 2026-08-30 on the project's verification node. This entry stays a draft until a named author and a different named reviewer sign it.