Thread hierarchy

Thread index visualizer

Drag gridDim and blockDim, then click any thread to see which element it owns. 1D and 2D grids, with out-of-range threads shown live.

Dimensions

Launch config

blockDim.x64

Threads per block. Always a multiple of 32.

gridDim.x4

Number of blocks. Capped at 8 here so the picture stays readable.

n (element count)200

Real data length. Threads past n are rejected by the bounds check.

Threads launched

256

Idle threads

Coverage is complete. Blocks needed: ceil(200 / 64) = 4, current gridDim = 4.

Click any thread to see its index mathOut of range (rejected by if)warp 0warp 1 …

blockIdx.x = 0global 0 … 63

blockIdx.x = 1global 64 … 127

blockIdx.x = 2global 128 … 191

blockIdx.x = 3global 192 … 255

Selected thread

blockIdx.x

threadIdx.x

warp (in block)

lane (in warp)

int i = 1 * 64 + 5 = 69i < n, handles c[69]

Matching code

CUDA C++

1__global__ void vecAdd(const float* a, const float* b, float* c, int n) {2    int i = blockIdx.x * blockDim.x + threadIdx.x;3    if (i < n) {4        c[i] = a[i] + b[i];5    }6}7 8int n = 200;9int threads = 64;10int blocks = (n + threads - 1) / threads;   // = 411vecAdd<<<blocks, threads>>>(d_a, d_b, d_c, n);