内存层次结构

本文档深入讲解 GPU 内存层次结构及其对 GEMM 优化的影响。

---

GPU 内存层次结构概览

┌─────────────────────────────────────────────────────────────┐
│                      CPU Host Memory                         │
│                    (DDR4/DDR5, 大容量)                        │
│                      延迟: ~100-300 ns                        │
└─────────────────────────────────────────────────────────────┘
                              │
                              │ PCIe 4.0/5.0 (~32 GB/s)
                              ▼
┌─────────────────────────────────────────────────────────────┐
│                    GPU Global Memory                         │
│                    (GDDR6/HBM, 16-80 GB)                     │
│                      延迟: ~400-800 cycles                    │
│                      带宽: 400-900 GB/s                       │
└─────────────────────────────────────────────────────────────┘
                              │
                              │
┌─────────────────────────────────────────────────────────────┐
│                      L2 Cache                                │
│                    (6-50 MB, 共享)                           │
│                      延迟: ~100-200 cycles                    │
└─────────────────────────────────────────────────────────────┘
                              │
                              │
┌─────────────────────────────────────────────────────────────┐
│                    Shared Memory                             │
│                    (48-164 KB/SM, 用户管理)                   │
│                      延迟: ~20-30 cycles                      │
│                      带宽: ~数 TB/s                           │
└─────────────────────────────────────────────────────────────┘
                              │
                              │
┌─────────────────────────────────────────────────────────────┐
│                   Register File                              │
│                    (64K 32-bit/SM)                           │
│                      延迟: ~1 cycle                           │
│                      带宽: ~数十 TB/s                         │
└─────────────────────────────────────────────────────────────┘

---

各级内存详解

1. 全局内存 (Global Memory)

特点：

• 容量最大（16-80 GB）
• 延迟最高（400-800 cycles）
• 所有线程可访问
• 通过 PCIe 与 Host 通信

访问模式：

// 全局内存访问
__global__ void kernel(float* global_data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = global_data[idx];  // 全局内存加载
    global_data[idx] = val * 2.0f; // 全局内存存储
}

合并访存 (Coalesced Access)：

当同一线程束内的线程访问连续地址时，访问可以合并：

// ✅ 合并访存：线程访问连续地址
int idx = threadIdx.x;
float val = data[idx];  // 合并为一次 128-byte 事务

// ❌ 非合并访存：线程访问跨度地址
int idx = threadIdx.x * stride;  // stride > 1
float val = data[idx];  // 多次独立事务

2. 共享内存 (Shared Memory)

特点：

• 用户管理的缓存
• 低延迟（~20-30 cycles）
• 高带宽（~数 TB/s）
• 线程块内共享

Bank Conflict：

共享内存被划分为 32 个 bank，每个 bank 宽度 4 字节：

__shared__ float data[32][32];

// ✅ 无 bank conflict：不同 bank
float val = data[threadIdx.y][threadIdx.x];

// ❌ bank conflict：同一线程束访问同一 bank
float val = data[threadIdx.x][threadIdx.x];  // 32-way conflict

// ✅ 添加 padding 避免 conflict
__shared__ float data[32][33];  // +1 padding
float val = data[threadIdx.x][threadIdx.x];  // 无 conflict

在 GEMM 中的应用：

__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];

// 加载到共享内存
As[ty][tx] = A[row * K + k + tx];
Bs[ty][tx] = B[(k + ty) * N + col];

__syncthreads();  // 线程块同步

// 从共享内存计算
for (int i = 0; i < BLOCK_SIZE; i++) {
    sum += As[ty][i] * Bs[i][tx];
}

3. 寄存器 (Register)

特点：

• 最快的存储（~1 cycle）
• 线程私有
• 数量有限（64K 32-bit/SM）
• 溢出会降低性能

寄存器分块优化：

// 每个线程计算 TILE_M × TILE_N 的输出块
float regA[TILE_M];  // 寄存器存储 A 的行
float regB[TILE_N];  // 寄存器存储 B 的列
float regC[TILE_M][TILE_N] = {0};  // 累加器

for (int k = 0; k < K; k += BLOCK_SIZE) {
    // 加载到寄存器
    for (int i = 0; i < TILE_M; i++) {
        regA[i] = As[ty * TILE_M + i][tk];
    }
    for (int j = 0; j < TILE_N; j++) {
        regB[j] = Bs[tk][tx * TILE_N + j];
    }
    
    // 寄存器内计算（无内存访问）
    for (int i = 0; i < TILE_M; i++) {
        for (int j = 0; j < TILE_N; j++) {
            regC[i][j] += regA[i] * regB[j];
        }
    }
}

---

内存访问优化策略

策略 1：数据复用

原理： 每次加载的数据参与多次计算

朴素 GEMM：
- 每个输出元素需要 2K 次内存访问
- 总访问量：2 × M × N × K

分块 GEMM：
- 每个数据块加载一次，复用 BLOCK_SIZE 次
- 总访问量减少为：2 × M × N × K / BLOCK_SIZE

策略 2：延迟隐藏

原理： 计算与内存访问重叠

// 双缓冲：加载和计算重叠
__shared__ float buffer[2][BLOCK_SIZE][BLOCK_SIZE];
int current = 0, next = 1;

for (int k = 0; k < K; k += BLOCK_SIZE) {
    // 异步加载下一个块
    load_async(buffer[next], k + BLOCK_SIZE);
    
    // 计算当前块
    compute(buffer[current]);
    
    // 同步并交换缓冲区
    __syncthreads();
    swap(current, next);
}

策略 3：向量化加载

原理： 使用 SIMD 指令一次加载多个数据

// 使用 float4 一次加载 4 个 float
float4* A4 = reinterpret_cast<float4*>(A_shared);
float4 a = A4[tx];  // 一次加载 4 个 float

// 展开后计算
sum0 += a.x * b.x;
sum1 += a.y * b.y;
sum2 += a.z * b.z;
sum3 += a.w * b.w;

---

性能分析工具

Nsight Compute 内存分析

# 分析内存吞吐量
ncu --metrics gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed \
    ./benchmark

# 分析共享内存 bank conflict
ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum \
    ./benchmark

关键指标

指标	含义	目标值
DRAM Throughput	全局内存吞吐量	> 80%
Shared Memory Bank Conflicts	Bank conflict 次数	接近 0
L2 Cache Hit Rate	L2 缓存命中率	越高越好
Register Usage	每线程寄存器数	< 255

---

不同 GPU 架构的差异

架构	共享内存/SM	寄存器/SM	L2 Cache	特点
Volta (V100)	96 KB	64K	6 MB	引入 Tensor Core
Ampere (A100)	164 KB	64K	40 MB	大 L2, TF32
Hopper (H100)	228 KB	64K	50 MB	FP8, TMA
Ada (RTX 4090)	100 KB	64K	6 MB	消费级

优化建议：

• 针对目标架构调整 BLOCK_SIZE
• 利用架构特性（如 Tensor Core）
• 使用 AutoTuner 自动搜索最优参数

---

参考资料

1. CUDA C Best Practices Guide - NVIDIA
2. Programming Massively Parallel Processors - Kirk & Hwu
3. CUDA Optimization Guide - NVIDIA