When programming a GPU, we organize work using threads and blocks. A thread is the smallest unit of parallel work. All threads run the same code, but each thread has a unique index it uses to operate on a different piece of data. A block is a programmer-defined group of threads, and each block is assigned to one SM.

__global__ void add_one(float* data) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    data[i] = data[i] + 1;
}

This is written in CUDA C, the language used to write GPU kernels. Do not worry if you are not familiar with it. The key idea is in the index calculation, not the syntax. __global__ marks the function as a GPU kernel, a function that runs on every thread simultaneously. Each thread computes its own i using three built-in values: blockIdx.x is the block's position in the grid, blockDim.x is the number of threads per block, and threadIdx.x is the thread's position within its block. Together they give every thread a globally unique index, so each thread adds 1 to a different element of data.

When the GPU executes a block, it divides the block's threads into warps of 32 and schedules them on the SM. The programmer thinks in terms of blocks; the hardware executes in terms of warps.

All threads within a block share the SM's shared memory and can synchronize with each other. This is the key connection between the programming model and the hardware: a block maps to one SM, its threads map to warps, and its data lives in that SM's shared memory.