GPU Architecture for CUDA

线程模型

NVIDIA GPU 是 SIMT (Single Instruction Multiple Thread) 模型，每一个指令可以被多个线程执行．

从底层的执行模型来讲，GPU 对线程的组织分为三层：

1. Grid, a collection of blocks
2. Block, a collection of threads
3. Thread

每一个 block 内的线程会被分配一个 threadIdx．对于一个线程，其 unique global index 可以表示为 blockIdx * blockDim + threadIdx．

对于 GPU 编程来说，CPU 中靠循环解决的问题，在 GPU 编程中靠多线程并行解决．当我们调用 CUDA kernel 时，需要指定 blockDim, gridDim，分别表示一个 block 包含多少 threads 和一个 grid 包含多少 blocks．

这里需要说明的是，block dim 和 grid dim 本身就可以是二维，甚至是三维的．其便利之处在于，我们可以直接 issue 二维的线程来更加直观地对应多维数据的索引．例如对于矩阵乘法 $C=AB$，我们可以 issue 二维的线程，每一个线程计算单个 $C_{ij}$，对应关系更加清晰．

1
2

row = blockIdx.y * blockDim.y + threadIdx.y;
col = blockIdx.x * blockDim.x + threadIdx.x;