The Thread Hierarchy

Limits (most GPUs): max 1024 threads per block; block dims up to (1024, 1024, 64); grid up to (2³¹−1, 65535, 65535).

5.1 Built-in variables

thread index within its block block index within the grid number of threads per block number of blocks in the grid

5.2 Launch configuration

Limits (most GPUs): max 1024 threads per block; block dims up to (1024, 1024, 64); grid up to (2³¹−1, 65535, 65535).

kernel<<<gridDim, blockDim, sharedMemBytes, stream>>>(args...);

dim3 block(16, 16);                                   // 256 threads per block (2-D)
dim3 grid((width + 15) / 16, (height + 15) / 16);     // enough blocks to cover the image
kernel2D<<<grid, block>>>(d_img, width, height);

5.3 Computing indices

// 1-D
int i = blockIdx.x * blockDim.x + threadIdx.x;

// 2-D (row-major image)
int x = blockIdx.x * blockDim.x + threadIdx.x;   // column
int y = blockIdx.y * blockDim.y + threadIdx.y;   // row
if (x < width && y < height) {
    int idx = y * width + x;
    out[idx] = 255 - in[idx];                      // invert image
}

// 3-D volume
int z = blockIdx.z * blockDim.z + threadIdx.z;
int idx3 = (z * ny + y) * nx + x;

5.4 Choosing block sizes

Use a multiple of 32 (warp size): 128, 256, 512 are common starting points.

Enough blocks to fill all SMs (≥ several blocks per SM).

Tune with profiling and the occupancy API (§20).

5.5 Matrix addition (2-D example)

__global__ void matAdd(const float* A, const float* B, float* C, int rows, int cols) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    if (row < rows && col < cols) {
        int i = row * cols + col;
        C[i] = A[i] + B[i];
    }
}