CUDA 架构

1. 硬件结构

英伟达CUDA/GPU架构演变，以及不同架构的硬件能力：

硬件层次结构如下（以Fermi架构为例）：

• 一个GPU中包含若干个SM（Streaming Multiprocessor，流多处理器），对应上图中左边；
• 一个SM中包含32个CUDA Core（也叫SP），对应上图中右边；
• 一个CUDA Core中包含一个ALU，一个FPU。
• 有些SM中还包含Tensor Core，与CUDA Core协同参与计算。

1.1. 缓存层级

一个CUDA Core内部，包含：

• Register File：16K 32-bit寄存器文件；
• L0 I-Cache

一个SM内部，包含：

• Shared Memory：一个SM内部的threads可访问（一个Thread Block内的所有线程可访问）；
• L1 Cache：一个SM内部的threads可访问；
• Constant Cache；
• Register File：编译时分配给每个thread使用的寄存器文件。

注意：高级的 NVIDIA GPU 中，包含若干个 Subcore，比如 Ampere 架构的 GPU 中包含4个 Subcore，每个 Subcore 包含 32 个 CUDA Core 和 1 个 Tensor Core。如下图所示：https://developer.nvidia.com/blog/nvidia-ampere-architecture-in-depth/

整个GPU内部，包含：

• L2 Cache：所有SM共享访问；
• Global Memory：所有SM共享访问。

内存访问速度示意图：

注1：Shared Memory与L1 Cache共享片上内存，通过cudaFuncSetAttribute(kernel_name, cudaFuncAttributePreferredSharedMemoryCarveout, carveout);提示驱动分配多少给Shared Memory，但是尽量不要使用这个函数。

资料：

• CUDA: GPU内存架构示意。文章结尾有一些内存优化相关链接。
• huggingface – The Ultra-Scale Playbook:Training LLMs on GPU Clusters – A primer on GPUs
• 英伟达GPU架构总结

2. 编程模型

软件/硬件层次结构对应关系：

层级	说明	硬件对应
Grid	所有要执行的Block的集合	整个GPU
Block	一组线程，一起在同一个SM上执行	一个SM（流多处理器）
Thread	执行kernel代码的最小单位	一个CUDA Core

代码执行过程：

myKernel<<<100, 256>>>(data, n)  // 启动kernel

↓ 产生 Grid，包含100个Block

Block 0:
  ├── Thread 0  执行myKernel
  ├── Thread 1  执行myKernel
  ├── ...
  └── Thread 255 执行myKernel

Block 1:
  ├── Thread 0  执行myKernel
  ├── ...
  └── Thread 255 执行myKernel

...（总共100个Block）

在调度时，一个Block中的Theads只会分配到一个SM中执行，如果资源不允许，则需要执行多次调度循环，才能执行完这个Block中的所有Threads。

另一方面，多个Block可以分配到同一个SM中执行。即，如果资源允许，或者当前Block中的Theads由于访问延迟而阻塞时，SM可以调度其他Block中的Threads来执行。（比如一个Warp的寄存器写后读会产生24个时钟延迟，则需要分配24个Warp来掩盖延迟）

同一SM上可运行的Block数量 = min(
    ⌊最大thread数 / 每个Block的thread数⌋,
    ⌊共享内存大小 / 每个Block占用的shared memory⌋,
    ⌊寄存器总数 / (每个thread占用寄存器数 × 每个Block的thread数)⌋,
    硬件限制（通常8-16个Block）
)

查看SM资源利用率：

nvprof --metrics achieved_occupancy ./program
# occupancy = (实际运行的warp数) / (理论最大warp数)

Fermi同时执行最多16个kernel。

2.1. kernel 索引

变量	类型	说明
threadIdx	uint3	当前thread在其Block内的索引（0-based）
blockIdx	uint3	当前Block在Grid中的索引（0-based）
blockDim	dim3	当前Block的维度/大小（thread数量）
gridDim	dim3	当前Grid的维度/大小（Block数量）

<<<grid, block>>>中，grid表示Grid的维度/大小（Block数量），block表示Block的维度/大小（Thread数量）。

一维示例及执行分析：

__global__ void kernel1D(float *data) {
    // blockDim.x = 256（启动时指定）
    // blockIdx.x = 0, 1, 2, ... （当前Block在Grid中的位置）
    // threadIdx.x = 0, 1, 2, ..., 255（当前thread在Block中的位置）
    
    // 计算全局线程索引
    int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;
    
    data[globalIdx] = data[globalIdx] * 2;
}

int main() {
    kernel1D<<<100, 256>>>(data);  // 100个Block，每个Block 256个thread
}

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

Grid:
  Block 0:  blockIdx.x=0
    ├── Thread 0:  threadIdx.x=0, globalIdx=0
    ├── Thread 1:  threadIdx.x=1, globalIdx=1
    ├── ...
    └── Thread 255: threadIdx.x=255, globalIdx=255
    
  Block 1:  blockIdx.x=1
    ├── Thread 0:  threadIdx.x=0, globalIdx=256
    ├── Thread 1:  threadIdx.x=1, globalIdx=257
    ├── ...
    └── Thread 255: threadIdx.x=255, globalIdx=511
    
  ...
  
  Block 99: blockIdx.x=99
    ├── Thread 0:  threadIdx.x=0, globalIdx=25344
    └── ...

二维示例：

__global__ void kernel2D(float *matrix, int width) {
    // blockDim.x = 16, blockDim.y = 16
    // blockIdx.x = 0,1,2,... blockIdx.y = 0,1,2,...
    // threadIdx.x = 0,...,15  threadIdx.y = 0,...,15
    
    // 计算全局行列索引
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    int idx = y * width + x;
    matrix[idx] = matrix[idx] * 2;
}

int main() {
    dim3 blockDim(16, 16);      // 16x16=256个thread/block
    dim3 gridDim(10, 10);       // 10x10=100个block
    kernel2D<<<gridDim, blockDim>>>(matrix, width);
}

3. CUDA 中一些重要概念

3.1. 线程束分叉 Warp Divergence

当存在条件分支时，不同的线程执行不同的代码分支，需要串行执行两个分支。此时，当一些线程直线分支1时，其他线程的执行被浪费（执行结果不写入），导致一些clock浪费。

3.2. Wavefront

一次 Wavefront 是指一次打包的 load / store 操作。可能由于资源原因，或者地址不对齐等原因，导致 warp 的一次 request 可能需要分为多个 wavefront 来完成。比如理想情况下，warp 从 SMEM 一次请求 128 字节数据只需要一个 wavefront；如果发生 bank conflict，则需要多个 wavefront 来完成。

官方文档描述见Nsight Compute – Metrics Guide。并参考https://forums.developer.nvidia.com/t/reuse-of-l1-shared-memory-during-execution-of-consecutive-wavefronts/288610。

3.3. N-way Bank Conflict

示例：访问一个列主序矩阵(float)，warp 内每个线程使用指令 LDS.128 或 STS.128 一次访问一个 float4。则一次访问（32 * 16B = 512B）需要分成 4 个内存事物，分别是：T0 ~ T7、T8 ~ T15、T16 ~ T23、T24 ~ T31。

由于是列主序，T0 ~ T7 访问同一个bank 0，这个叫 8-way bank conflict，需要拆分成 8 个wavefront，产生 7 个 bank conflicts。warp 的一个 request，总共产生的 bank conflict 数量是：

bank_conflicts = 7 * 4 = 28

3.3. 计算强度 Arithmetic Intensity

在性能瓶颈分析过程中，需要确定瓶颈是计算能力还是内存带宽。计算强度定义为每次内存访问所执行的计算量，通常以 FLOPS/Byte 表示。定义公式：

计算强度 = 算术运算次数 / 访问的内存字节数

以 N * N 矩阵乘法为例，执行 N^3 次乘加运算，以及 N^2 * (N-1)次加法，访问 3 * N^2 个元素（A、B、C 矩阵），每个元素假设为 4 字节（float 类型），则计算强度为：

计算强度 = N^3 / (3 * N^2 * 4) = N / 12 (FLOPS/Byte)

Nsight Compute 相关资料

• 详细介绍：【CUDA调优指南】缓存&访存流程
• 深入理解 roofline 模型

4. 参考资料

• NVidia GPU指令集架构-寄存器
• CUDA Refresher: The CUDA Programming Model
• 深入解析 NVIDIA Hopper 架構
• CUDA C++ Best Practices Guide
• CUDA编程：基础与实践 pdf
• CUDA blogs

4.1. 手册等

• CUDA GPU Compute Capability：不同架构的计算能力对照表