CUDA Thread Hierarchy Notes

CUDA Thread Hierarchy Notes

CUDA organizes work into a simple hierarchy: threads, warps, blocks, and grids.

1. Thread

A thread is the smallest execution unit in CUDA. Each thread runs one instance of the kernel.

Each thread has its own:

• threadIdx, which tells the thread its location inside the block
• private registers
• private local variables

2. Warp

A warp is a group of 32 threads. The GPU schedules and executes work in warps.

For example, a block with 256 threads contains 8 warps:

Threads inside the same warp execute together. If threads in one warp take different branches, the warp diverges. Divergence can reduce performance because the GPU may need to execute each branch path separately.

3. Block

A block, also called a thread block, is a group of threads. Threads in the same block can cooperate with each other.

Threads inside the same block can:

• share low-latency shared memory
• synchronize using __syncthreads()
• work together on the same part of a problem

Example kernel launch:

kernel<<<100, 256>>>();

This launch creates 100 blocks, and each block has 256 threads. Inside each block:

A block runs entirely on one Streaming Multiprocessor, or SM. The block is not split across multiple SMs.

4. Grid

A grid is the full collection of blocks launched by one kernel call.

kernel<<<100, 256>>>();

This means:

The total number of threads is:

CUDA Built-in Variables

CUDA gives every thread built-in variables that describe where it is located in the launch.

threadIdx

threadIdx is the thread's index inside its block.

blockDim

blockDim gives the number of threads in each block along each dimension.

For example:

dim3 block(16, 16);

Inside the kernel, this means:

The number of threads per block is:

blockIdx

blockIdx is the block's index inside the grid.

gridDim

gridDim gives the number of blocks in the grid along each dimension.

For example:

dim3 grid(64, 64);

Inside the kernel, gridDim.x is 64 and gridDim.y is 64. The grid contains:

Global Thread Index

For a 1D kernel, each thread usually computes one global index. That index tells the thread which array element to process.

Example:

Then:

2D Thread Index Example

For images or matrices, a 2D block and 2D grid often make the indexing easier to understand.

dim3 block(16, 16);
dim3 grid(64, 64);

Each block has:

The grid has:

The total number of threads is:

A 2D kernel often computes the global row and column like this:

Hardware Mapping Reference

The CUDA hierarchy is logical, but it maps to GPU hardware in a useful way.

Core Hardware Rules

• A warp always contains exactly 32 threads.
• A block is made of threads.
• A grid is made of blocks.
• Threads in the same block can synchronize.
• Blocks cannot directly synchronize with each other inside a normal kernel.
• A block runs entirely on one SM.
• Different blocks can run on different SMs.

Simple Analogy

CUDA Function Declarations and Device/Host Distinction

CUDA extends C/C++ with function specifiers. These keywords say where a function runs and where it can be called from.

Important Rules

• __global__ marks a function as a CUDA kernel.
• A __global__ kernel is launched with triple angle brackets, like kernel<<<grid, block>>>();.
• A __global__ function must return void.
• CUDA specifiers use double underscores before and after the word, such as __global__.
• __host__ and __device__ can be used together, such as __host__ __device__ void CommonMath(). This tells the compiler to generate both CPU and GPU versions.
• If no specifier is given, CUDA treats the function as a normal host function.

Compiling a CUDA Program

CUDA programs are usually stored in .cu files and compiled with NVIDIA's CUDA compiler driver, nvcc.

• Source separation: nvcc separates host code from device code.
• Host code: CPU code is passed to a normal C/C++ compiler, such as gcc, clang, or cl.exe.
• Device code: GPU code is compiled into PTX, which stands for Parallel Thread Execution.
• Hardware mapping: PTX is translated into lower-level GPU instructions called SASS. This can happen before runtime or just-in-time at runtime.

Practical Application: Color-to-Grayscale Conversion

An RGB image stores three color channels: red, green, and blue. A grayscale image stores one intensity value per pixel.

A common weighted conversion is:

The CUDA kernel below maps each thread to one pixel:

#define CHANNELS 3 // Red, Green, Blue

__global__ void colorConvert(unsigned char* grayImage,
                             unsigned char* rgbImage,
                             int width, int height)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;

    if (x < width && y < height) {
        int grayOffset = y * width + x;
        int rgbOffset  = grayOffset * CHANNELS;

        unsigned char r = rgbImage[rgbOffset];
        unsigned char g = rgbImage[rgbOffset + 1];
        unsigned char b = rgbImage[rgbOffset + 2];

        grayImage[grayOffset] =
            (unsigned char)(0.21f * r + 0.71f * g + 0.07f * b);
    }
}

Thread Block Scheduling and Hardware Limits

Thread blocks are designed to be independent. Because of this, the GPU can schedule blocks in any order across available SMs.

Hardware Capacity and Resource Limits

Each SM has limits on how many blocks and threads can be active at the same time. For the example below, assume these limits:

• Maximum blocks per SM: 8 blocks
• Maximum threads per SM: 1536 threads
• Maximum threads per block: 1024 threads

Block size affects occupancy. Occupancy is how much of the SM's available thread capacity is being used.

Warps and Control Flow Divergence Mechanics

Linearization Layout Mapping Rules

CUDA can define blocks in 1D, 2D, or 3D, but the hardware still forms warps from a linear ordering of threads.

• Thread coordinates are linearized in row-major order.
• The x dimension increments fastest, followed by y, then z.
• The linear thread order is split into groups of 32 threads.
• Warp 0 contains linear thread positions 0 through 31. Warp 1 contains positions 32 through 63, and so on.

Warp Control Divergence

Threads in the same warp normally execute the same instruction together. If an if statement sends some threads down one branch and other threads down another branch, the warp diverges.

• The GPU runs one branch while masking off the threads that should not take that branch.
• Then it switches masks and runs the other branch.
• This makes the branch paths run sequentially instead of fully in parallel.

Visual examples

Quantifying the boundary check performance penalty

Consider a 1D vector addition kernel processing n = 1000 elements using 256 threads per block:

__global__ void vecAddKernel(float* A, float* B, float* C, int n) {
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

The launch needs four blocks:

Conclusion: Only one warp out of 32 total warps diverges. The other 31 warps execute normally. This is why simple boundary checks are usually acceptable, especially when only the final block is partially outside the data range.