GPU硬件架构

📅 4周详细学习计划

1. SM (Streaming Multiprocessor) 架构

1.1 SM核心组件

• CUDA Cores：1024个单精度FP32运算单元，每个周期可执行一次FMA
• Tensor Cores：4个矩阵运算单元，每个周期执行64次FMA（FP16×FP16→FP32）
• Warp Scheduler：4个调度器，每个管理一个Warp（32线程）
• Register File：256KB寄存器文件，每线程最多255个寄存器
• Shared Memory/L1 Cache：总计256KB，可配置比例（如164KB Shared + 92KB L1）

1.2 SM工作原理

• 指令流水线：16级流水线，支持指令级并行
• Thread Coarsening：一个Warp执行32个线程的相同指令
• 延迟隐藏：通过Warp切换隐藏内存延迟（~400周期）

2. Warp执行模型

2.1 Warp调度

• 就绪Warp：调度器选择已准备好的Warp执行（无依赖）
• 延迟隐藏：当一个Warp等待内存时，切换到另一个Warp
• 占用率影响：更多活跃Warp = 更好的延迟隐藏

2.2 分支发散 (Branch Divergence)

• 问题：Warp内线程走不同分支时，串行执行所有分支
• 影响：性能下降与分支复杂度成正比
• 优化：重组数据使同一Warp线程执行相同分支

// ❌ 分支发散
if (threadIdx.x < 16) {
    do_something_A();  // 前16个线程执行
} else {
    do_something_B();  // 后16个线程执行
}

// ✅ 避免发散 - 使用掩码
unsigned int mask = (threadIdx.x < 16) ? 0xFFFF : 0x0000;
// 使用__ballot_sync()获取线程投票

2.3 Warp协作通信

• __shfl_sync()：Warp内线程直接交换数据（寄存器级）
• __ballot_sync()：收集线程投票（返回32位掩码）
• __any_sync() / __all_sync()：条件判断

// Warp Shuffle实现归约
__device__ float warp_reduce_sum(float val) {
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xFFFFFFFF, val, offset);
    }
    return val;
}

3. Occupancy（占用率）详解

• 定义：Occupancy = 活跃Warp数 / SM最大支持Warp数
• 计算因素：寄存器使用量、Shared Memory使用量、Block大小
• 目标：通常50%-100%是理想范围
• 权衡：高Occupancy不一定高性能（需平衡其他因素）

// Occupancy计算示例 (A100)
// A100: 每SM 2048线程, 64个Warp, 65536寄存器

// Case 1: 低寄存器使用
// 每线程32寄存器 → 2048*32=65536 → 100% Occupancy

// Case 2: 高寄存器使用
// 每线程64寄存器 → 1024*64=65536 → 50% Occupancy

// 使用__launch_bounds__控制寄存器
__global__ void __launch_bounds(256, 2)  // maxThreadsPerBlock=256, minBlocksPerSM=2
my_kernel() {
    // 编译器会优化寄存器使用以达到目标Occupancy
}

4. Tensor Core详解

4.1 Tensor Core支持的精度

• FP16：半精度，A100/H100支持
• BF16：Brain Float，范围与FP32相同
• TF32：TensorFloat-32，FP32范围+FP16精度
• INT8：8位整数，推理常用
• FP8 (H100)：新增精度，推理训练均可

4.2 WMMA API使用

// WMMA (Warp Matrix Multiply Accumulate) 示例
#include <mma.h>

using namespace nvcuda;

__global__ void tensor_core_gemm(half *A, half *B, float *C) {
    // 声明片段
    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;
    
    // 加载矩阵
    wmma::load_matrix_sync(a_frag, A, 16);
    wmma::load_matrix_sync(b_frag, B, 16);
    
    // 执行矩阵乘加
    wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);
    
    // 存储结果
    wmma::store_matrix_sync(C, c_frag, 16, wmma::mem_row_major);
}

5. 显存层次详解

5.1 HBM vs GDDR

• HBM (High Bandwidth Memory)：高带宽，低功耗，用于A100/H100
• GDDR6X：高带宽，用于游戏GPU（如RTX 3090）
• 带宽差异：A100: 2TB/s vs RTX 3090: 936GB/s

6. NVLink / NVSwitch

• NVLink 3.0：A100支持，每GPU 600GB/s双向带宽
• NVSwitch：全连接交换机，支持8 GPU全互联
• 应用场景：多卡推理、分布式训练、AllReduce通信
• 编程模型：NCCL库封装，透明处理多卡通信