前言

第四章理解了GPU的调度机制和硬件架构，这一章进入性能优化的核心——内存。GPU计算能力强大，但数据供应跟不上就白搭。第五章讲解GPU的内存层次结构，重点是Shared Memory和Tiling技术。掌握这些，矩阵乘法性能可以提升10倍以上。

📦 配套资源：本系列文章配有完整的 GitHub 仓库，包含每章的练习题解答、CUDA 代码实现和详细注释。所有代码都经过测试，可以直接运行。

内存带宽：性能的天花板

问题的本质

回顾第三章的矩阵乘法：

1
2
3
4
5
6
7
8
9
10
11
12
__global__ void matMul(float *M, float *N, float *P, int width) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (row < width && col < width) {
        float sum = 0.0f;
        for (int k = 0; k < width; k++) {
            sum += M[row * width + k] * N[k * width + col];
        }
        P[row * width + col] = sum;
    }
}

计算 P[row][col] 需要读取 M 的一行和 N 的一列，共 2×width 个元素。每个元素 4 字节，width=1024 时：

1
2
3
每线程读取：2 × 1024 × 4 = 8192 字节
每线程计算：2 × 1024 = 2048 FLOP
算术强度：2048 / 8192 = 0.25 FLOP/Byte

现代 GPU 峰值 10+ TFLOPS，带宽 500+ GB/s，需要 20 FLOP/Byte 才能跑满计算单元。实际只有 0.25，GPU 大部分时间在等数据。

数据重复访问

更严重的是：同一数据被多个线程重复读取。

M[i][k] 被第 i 行的所有线程读取，N[k][j] 被第 j 列的所有线程读取。1024×1024 矩阵，每个元素被读 1024 次，但每次都从全局内存（DRAM）读。

这就是优化的切入点：让数据复用发生在快速存储上，而不是全局内存。

GPU 内存层次

层次结构

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
┌─────────────────────────────────────────────┐
│                   Host Memory               │   CPU DDR4/DDR5
│                   (GB 级别)                  │   ~50 GB/s
├─────────────────────────────────────────────┤
│              Global Memory (DRAM)           │   GPU 显存
│                  (GB 级别)                   │   ~500 GB/s
├─────────────────────────────────────────────┤
│    L2 Cache (SM 共享)                        │   几 MB
│                                              │   ~1-2 TB/s
├──────────────┬──────────────┬───────────────┤
│   Shared     │   Shared     │   Shared      │   每 SM
│   Memory     │   Memory     │   Memory      │   ~100 KB
│   L1 Cache   │   L1 Cache   │   L1 Cache    │   ~10 TB/s
├──────────────┼──────────────┼───────────────┤
│  Registers   │  Registers   │  Registers    │   每线程
│              │              │               │   ~几十 TB/s
└──────────────┴──────────────┴───────────────┘
     SM 0           SM 1           SM 2

核心规律：越靠近计算单元，容量越小，速度越快。

各级存储特点

存储类型 作用域 容量 延迟 程序员控制
寄存器 单线程 ~255个/线程 1周期 隐式（变量）
共享内存 Block内 ~100KB/SM ~20周期 显式
L1缓存 SM内 ~128KB/SM ~20周期 部分
L2缓存 全局 ~6MB ~200周期 无
全局内存 全局 ~数GB ~400周期 显式

关键洞察：

  • 寄存器最快，但容量有限，且只属于单个线程
  • 共享内存是程序员可控的 Block 级缓存，这是优化的主战场
  • 全局内存是唯一能容纳大数据的地方，但太慢

共享内存（Shared Memory）

基本概念

共享内存是 SM 上的可编程缓存：

  • Block 内所有线程共享：同 Block 的线程可以读写同一块共享内存
  • 生命周期与 Block 绑定：Block 结束，共享内存释放
  • 低延迟：约 20 周期，比全局内存快 20 倍

声明语法

静态分配（编译时确定大小）：

1
2
3
4
__global__ void kernel() {
    __shared__ float sharedData[256];  // Block 内共享
    // ...
}

动态分配（运行时确定大小）：

1
2
3
4
5
6
7
__global__ void kernel() {
    extern __shared__ float sharedData[];  // 大小由启动配置指定
    // ...
}

// 启动时指定共享内存大小
kernel<<<grid, block, sharedMemBytes>>>(args);

同步：__syncthreads()

共享内存是 Block 内共享的，需要同步保证数据一致性：

1
2
3
4
5
6
7
8
9
__shared__ float data[256];

// 阶段1：所有线程写入
data[threadIdx.x] = input[globalIdx];

__syncthreads();  // 等待所有线程完成写入

// 阶段2：所有线程读取（此时数据已就绪）
float val = data[(threadIdx.x + 1) % 256];

__syncthreads() 是栅栏同步：Block 内所有线程必须到达这一点，才能继续执行。

常见错误：

1
2
3
4
5
// 危险！条件内使用 syncthreads
if (threadIdx.x < 128) {
    data[threadIdx.x] = ...;
    __syncthreads();  // 只有部分线程执行，会死锁！
}

同步必须保证 Block 内所有活跃线程都执行到。

Tiling：分块处理

核心思想

既然全局内存慢但共享内存快，策略就是：

  1. 分块加载：将数据分成小块（Tile），逐块加载到共享内存
  2. 计算复用：在共享内存中完成该块的所有计算
  3. 移动窗口：处理下一块，直到完成

这样，每个数据从全局内存只读一次，但在共享内存中被多次使用。

Tiled 矩阵乘法

问题：计算 P = M × N，每个 P[i][j] = Σ M[i][k] × N[k][j]

朴素版本：每个线程独立读取整行和整列（大量重复读取）

Tiled 版本：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
┌───────────────┐   ┌───────────────┐
│   M 矩阵       │   │   N 矩阵       │
│  ┌───┐        │   │      ┌───┐    │
│  │Tile│ ────→ │   │      │Tile│   │
│  └───┘        │   │      └───┘    │
│               │   │        │      │
└───────────────┘   └────────│──────┘
                             ↓
                    ┌───────────────┐
                    │   P 矩阵       │
                    │      ┌───┐    │
                    │      │计算│    │
                    │      └───┘    │
                    └───────────────┘

步骤：

  1. 将 M 的一个 Tile 和 N 的一个 Tile 加载到共享内存
  2. Block 内所有线程使用共享内存中的数据进行部分计算
  3. 加载下一对 Tile，累加结果
  4. 重复直到完成

代码实现

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
#define TILE_WIDTH 16

__global__ void matMulTiled(float *M, float *N, float *P, int width) {
    // 共享内存声明
    __shared__ float Mds[TILE_WIDTH][TILE_WIDTH];
    __shared__ float Nds[TILE_WIDTH][TILE_WIDTH];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    // 计算该线程负责的 P 元素位置
    int row = by * TILE_WIDTH + ty;
    int col = bx * TILE_WIDTH + tx;
    
    float Pvalue = 0;
    
    // 分块循环
    for (int ph = 0; ph < width / TILE_WIDTH; ++ph) {
        
        // 协作加载 M 的 Tile
        Mds[ty][tx] = M[row * width + (ph * TILE_WIDTH + tx)];
        
        // 协作加载 N 的 Tile
        Nds[ty][tx] = N[(ph * TILE_WIDTH + ty) * width + col];
        
        __syncthreads();  // 确保 Tile 加载完成
        
        // 使用共享内存计算部分点积
        for (int k = 0; k < TILE_WIDTH; ++k) {
            Pvalue += Mds[ty][k] * Nds[k][tx];
        }
        
        __syncthreads();  // 确保计算完成再加载下一个 Tile
    }
    
    P[row * width + col] = Pvalue;
}

关键点解析

1. 协作加载

Block 内的线程分工加载 Tile：

1
Mds[ty][tx] = M[row * width + (ph * TILE_WIDTH + tx)];

每个线程加载一个元素，16×16 = 256 个线程加载 256 个元素。比单线程加载整个 Tile 高效得多。

2. 两次同步

1
2
3
__syncthreads();  // 第一次：确保数据加载完成
// ... 计算 ...
__syncthreads();  // 第二次：确保计算完成再覆盖共享内存

两次同步都必要：

  • 第一次：防止读到未加载的数据
  • 第二次：防止快线程覆盖慢线程还在用的数据

3. 内层循环

1
2
3
for (int k = 0; k < TILE_WIDTH; ++k) {
    Pvalue += Mds[ty][k] * Nds[k][tx];
}

这个循环只访问共享内存，没有全局内存访问。这是性能提升的来源。

性能分析

朴素版本：

  • 每线程读取全局内存：2 × width 次
  • 总全局内存访问：width³ × 2（读）+ width²（写）

Tiled 版本：

  • 每 Tile 阶段：Block 读 2 × TILE_WIDTH² 个元素
  • 共 width/TILE_WIDTH 个阶段
  • 每线程贡献：2 × width 次（与朴素相同？不对！）

关键差异：在 Tiled 版本中，每个全局内存读取被 TILE_WIDTH 个线程共享使用。

1
2
全局内存访问减少倍数 = TILE_WIDTH
算术强度提升 = TILE_WIDTH 倍

TILE_WIDTH = 16 时，算术强度从 0.25 提升到 4 FLOP/Byte。TILE_WIDTH = 32 时可达 8 FLOP/Byte。

边界处理

上面的代码假设 width 是 TILE_WIDTH 的倍数。实际需要处理边界：

1
2
3
4
5
6
7
8
9
10
// 带边界检查的加载
if (row < width && (ph * TILE_WIDTH + tx) < width)
    Mds[ty][tx] = M[row * width + (ph * TILE_WIDTH + tx)];
else
    Mds[ty][tx] = 0;  // 越界元素填0

if ((ph * TILE_WIDTH + ty) < width && col < width)
    Nds[ty][tx] = N[(ph * TILE_WIDTH + ty) * width + col];
else
    Nds[ty][tx] = 0;

填 0 不影响加法结果，是处理边界的常用技巧。

内存访问模式

合并访问（Coalesced Access）

GPU 内存控制器按 32 字节或 128 字节 的事务读写数据。如果 Warp 中的线程访问连续地址，可以合并成一次事务：

好的访问模式：

1
2
// 线程 0,1,2,...,31 访问连续地址
data[threadIdx.x];  // 一次 128B 事务

差的访问模式：

1
2
// 线程 0,1,2,...,31 访问跨步地址
data[threadIdx.x * 32];  // 32 次事务！

矩阵访问的陷阱：

1
2
3
// P[i][j] 遍历
// 按行遍历（好）：data[row * width + col]，col 连续
// 按列遍历（差）：data[row * width + col]，row 连续（跨步 = width）

共享内存 Bank 冲突

共享内存分成 32 个 Bank，每个 Bank 宽度 4 字节。不同 Bank 可以同时访问，但同一 Bank 的不同地址会串行化。

Bank 映射：

1
2
3
4
5
地址 0,32,64,...  → Bank 0
地址 4,36,68,...  → Bank 1
地址 8,40,72,...  → Bank 2
...
地址 124,156,... → Bank 31

无冲突：

1
2
data[threadIdx.x];      // 32 线程访问 32 个 Bank
data[threadIdx.x * 2];  // 跨步 2，访问 Bank 0,2,4,...（无冲突）

有冲突：

1
data[threadIdx.x * 32]; // 所有线程访问 Bank 0！32-way 冲突

特例——广播：

1
data[0];  // 所有线程读同一地址，硬件广播，无冲突

矩阵转置的 Bank 冲突

考虑共享内存中 16×16 的矩阵：

1
2
3
4
5
6
7
__shared__ float tile[16][16];

// 加载（列访问）
tile[ty][tx] = input[row * width + col];  // 无冲突

// 存储（行访问）
output[col * width + row] = tile[tx][ty];  // 有冲突！

tile[tx][ty] 使得线程 0,1,2,…,15 分别访问 tile[0][ty], tile[1][ty], …，这些元素地址为 ty, ty+16, ty+32, …，步长 16×4 = 64 字节 = 16 个 Bank。部分线程会访问同一 Bank。

解决方案——Padding：

1
2
3
4
__shared__ float tile[16][17];  // 多加一列

// 现在 tile[i][j] 的地址是 i*17 + j
// 步长变成 17×4 = 68 字节，不再对齐

Padding 打破了 Bank 对齐，消除冲突。代价是浪费一点共享内存。

常量内存（Constant Memory）

特点

  • 只读：Kernel 内不能写
  • 缓存优化：有专用缓存，广播访问效率高
  • 容量有限：64 KB

适用场景

所有线程读相同数据（如卷积核、变换矩阵）：

1
2
3
4
5
6
7
8
9
10
11
12
__constant__ float kernel[9];  // 声明

// Host 端写入
cudaMemcpyToSymbol(kernel, h_kernel, 9 * sizeof(float));

// Kernel 内使用
__global__ void conv(...) {
    float sum = 0;
    for (int i = 0; i < 9; i++) {
        sum += data[i] * kernel[i];  // 所有线程读相同 kernel[i]
    }
}

如果每个线程读不同地址，常量内存反而更慢（串行化）。

寄存器优化

寄存器的重要性

寄存器是最快的存储，单周期延迟。但数量有限（每 SM 约 64K 个），过度使用会导致：

  1. 寄存器溢出（Spilling）：溢出到 Local Memory（实际是全局内存），极慢
  2. 占用率下降：每线程用更多寄存器，SM 能容纳的线程数减少

查看寄存器使用

1
nvcc --ptxas-options=-v kernel.cu

输出：

1
ptxas info: Used 32 registers, ...

控制策略

编译器提示：

1
2
3
4
__global__ void __launch_bounds__(256, 4) kernel(...) {
    // 告诉编译器：每 Block 256 线程，每 SM 至少 4 个 Block
    // 编译器据此优化寄存器分配
}

编译选项：

1
nvcc -maxrregcount=32 kernel.cu  # 限制每线程最多 32 个寄存器

权衡：限制过严可能导致溢出，限制过松可能降低占用率。需要实测。

数据局部性优化清单

空间局部性

定义：访问的数据在内存中相邻

优化手段：

  • 连续访问，利用合并
  • 数据布局优化（AoS → SoA）
1
2
3
4
5
6
7
8
9
// Array of Structures（差）
struct Particle { float x, y, z; };
Particle particles[N];
particles[i].x;  // 跨步 12 字节

// Structure of Arrays（好）
struct Particles { float x[N], y[N], z[N]; };
Particles p;
p.x[i];  // 连续访问

时间局部性

定义：同一数据短期内被多次访问

优化手段：

  • Tiling 到共享内存/寄存器
  • 循环分块
1
2
3
4
5
6
7
8
9
10
11
12
// 无时间局部性
for (int k = 0; k < N; k++) {
    C[i][j] += A[i][k] * B[k][j];  // A、B 每次从全局内存读
}

// 有时间局部性（Tiled）
for (int tile = 0; tile < N/TILE; tile++) {
    // 加载 tile 到共享内存
    for (int k = 0; k < TILE; k++) {
        C[i][j] += As[ty][k] * Bs[k][tx];  // 从共享内存读，复用
    }
}

实战：优化后的矩阵乘法性能

以 1024×1024 矩阵为例：

版本 全局内存访问 算术强度 相对性能
朴素 2×10⁹ 次 0.25 1×
Tiled (16×16) 1.25×10⁸ 次 4 ~8×
Tiled (32×32) 6.25×10⁷ 次 8 ~12×
+ 寄存器优化 更少 16+ ~20×

实际提升取决于具体 GPU 和问题规模，但 10× 以上是常见的。

小结

第五章是性能优化的核心：

内存层次认知：寄存器→共享内存→L2→全局内存，速度差 100 倍以上。写高性能代码就是让热数据留在快存储。

共享内存本质：程序员可控的 Block 级缓存。声明简单，但需要正确同步。两次 __syncthreads() 别忘。

Tiling 核心：分块加载，块内复用。矩阵乘法从 0.25 提升到 8+ FLOP/Byte，这是实打实的 10× 加速。

访问模式：

  • 全局内存要合并访问
  • 共享内存要避免 Bank 冲突
  • 必要时用 Padding

局部性原则：空间局部性（连续访问），时间局部性（重复使用）。Tiling 同时利用了两者。

掌握了内存优化，GPU 的计算能力才能真正发挥。下一章会学习更多计算模式——卷积、规约、前缀和，这些都需要精心设计的内存访问策略。

参考资料：

本文 GitHub 仓库: https://github.com/psmarter/PMPP-Learning