3. 数据复用与局部性优化：时间局部性与空间局部性、Tiling技术（分块）、Shared Memory手动缓存、Register重用来减少显存访问
大家好，我是你们的老朋友。今天我们来聊聊显存优化的核心——数据复用与局部性。说实话，很多同学在写GPU程序时，一上来就想着怎么并行，怎么把线程堆满。但往往忽略了最关键的一点：数据到底是怎么在芯片里流动的。
我见过太多项目，明明计算量不大，却因为显存访问太频繁，导致性能直接腰斩。嗯，这其实是个很常见的坑。今天我们就把它彻底讲透。
3.1 时间局部性与空间局部性：两个基本概念
先问大家一个问题：为什么CPU和GPU都要搞缓存？说白了，就是访存太慢了。从显存读一次数据，可能要几百个时钟周期。而计算一个浮点乘加，只需要几个周期。这中间的差距，就是我们需要优化的空间。
时间局部性：如果一个数据被访问过，那么它很可能在不久的将来再次被访问。举个例子，你在循环里反复读同一个变量，这就是时间局部性。
空间局部性：如果一个地址被访问了，那么它附近的地址很可能也会被访问。比如你遍历一个数组，相邻元素就是空间局部性。
核心观点：GPU的优化，本质上就是利用这两种局部性，把数据尽量往靠近计算单元的地方搬。越靠近计算单元，访问速度越快，功耗也越低。
我在项目中遇到过这样一个案例：一个矩阵乘法算子，原始版本直接读写全局显存，带宽利用率只有20%。后来我们做了分块和共享内存缓存，带宽利用率直接飙到85%以上。你想想看，这差距有多大？
3.2 Tiling技术（分块）：把大问题切成小问题
Tiling，说白了就是分块。把一个大矩阵切成若干个小块，每次只处理一小块。这样，这个小块的数据就可以完全放到共享内存里，反复使用。
为什么要分块？因为共享内存太小了。以NVIDIA A100为例，每个SM只有192KB的共享内存。你不可能把整个矩阵都塞进去。但你可以把矩阵切成16x16、32x32的小块，每个块刚好能放进共享内存。
来看一个简单的例子：
// 没有分块的矩阵乘法（直接访问全局显存）
__global__ void matmul_naive(float *A, float *B, float *C, int N) {
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;
for (int k = 0; k < N; k++) {
sum += A[row * N + k] * B[k * N + col];
}
C[row * N + col] = sum;
}
// 使用Tiling的矩阵乘法（利用共享内存）
#define TILE_SIZE 16
__global__ void matmul_tiled(float *A, float *B, float *C, int N) {
__shared__ float As[TILE_SIZE][TILE_SIZE];
__shared__ float Bs[TILE_SIZE][TILE_SIZE];
int row = blockIdx.y * TILE_SIZE + threadIdx.y;
int col = blockIdx.x * TILE_SIZE + threadIdx.x;
float sum = 0.0f;
for (int tile = 0; tile < N / TILE_SIZE; tile++) {
// 加载数据到共享内存
As[threadIdx.y][threadIdx.x] = A[row * N + tile * TILE_SIZE + threadIdx.x];
Bs[threadIdx.y][threadIdx.x] = B[(tile * TILE_SIZE + threadIdx.y) * N + col];
__syncthreads();
// 计算当前tile
for (int k = 0; k < TILE_SIZE; k++) {
sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
}
__syncthreads();
}
C[row * N + col] = sum;
}个人经验：分块大小不是越大越好。我建议你根据共享内存大小和寄存器数量来调。一般来说，16x16或32x32是比较稳妥的选择。太大容易导致共享内存溢出，太小又无法充分利用局部性。
3.3 Shared Memory手动缓存：把控制权握在自己手里
共享内存，你可以把它理解成程序员手动管理的缓存。L1缓存是硬件自动管理的，你没法控制。但共享内存是你可以显式控制的。这意味着什么？意味着你可以精确地决定哪些数据要缓存，缓存多久，什么时候替换。
为什么要手动管理？因为硬件缓存有时候不够聪明。举个例子，如果你的数据访问模式是跨步的（stride很大），L1缓存的预取机制可能完全失效。这时候，手动用共享内存来缓存，效果会好得多。
我曾经优化过一个3D卷积算子。原始版本直接读全局显存，性能惨不忍睹。后来我们用共享内存缓存了输入特征图的多个通道，每次计算时直接从共享内存读，避免了重复的全局访存。最终性能提升了3倍多。
避坑指南：我曾经犯过一个错误——在共享内存里放了太多数据，导致每个SM能同时运行的block数量减少。共享内存是有限的，你用得越多，SM上能同时跑的block就越少。这反而会降低吞吐。所以，共享内存不是越多越好，要平衡使用。
3.4 Register重用：最极致的优化
寄存器是GPU里最快的存储单元。没有之一。访问寄存器只需要1个时钟周期，而访问共享内存需要几十个周期，访问全局显存需要几百个周期。所以，能用寄存器就别用共享内存，能用共享内存就别用全局显存。
寄存器重用，说白了就是把反复使用的数据尽量放在寄存器里。比如在矩阵乘法中，每个线程计算一个输出元素时，会累加很多次。这个累加变量就应该放在寄存器里，而不是每次从显存读。
来看一个更极致的例子：
// 寄存器重用的卷积计算
__global__ void conv_register_reuse(float *input, float *weight, float *output,
int H, int W, int K) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
// 把权重加载到寄存器（每个线程只加载自己需要的部分）
float w[9]; // 假设3x3卷积
for (int i = 0; i < 3; i++) {
for (int j = 0; j < 3; j++) {
w[i * 3 + j] = weight[i * 3 + j];
}
}
// 计算输出，累加变量也在寄存器
float sum = 0.0f;
for (int i = 0; i < 3; i++) {
for (int j = 0; j < 3; j++) {
int in_x = x + j - 1;
int in_y = y + i - 1;
if (in_x >= 0 && in_x < W && in_y >= 0 && in_y < H) {
sum += input[in_y * W + in_x] * w[i * 3 + j];
}
}
}
output[y * W + x] = sum;
}关键点：寄存器数量是有限的。每个线程最多能用255个寄存器（NVIDIA架构）。如果你用太多，编译器会 spill 到 local memory（实际上就是全局显存），那就得不偿失了。所以，寄存器重用要适度，别贪心。
3.5 三种技术的对比与选择
| 技术 | 存储位置 | 访问延迟 | 容量 | 适用场景 |
|---|---|---|---|---|
| Register重用 | 寄存器 | 1 cycle | ~255个/线程 | 频繁使用的标量或小数组 |
| Shared Memory缓存 | 共享内存 | ~30 cycles | ~48-192KB/SM | 线程间共享的数据块 |
| Tiling分块 | 共享内存+全局显存 | 取决于实现 | 无限制 | 大矩阵、大张量操作 |
我个人习惯是：先做Tiling，再用Shared Memory缓存，最后用Register优化热点。这个顺序能让你在大多数场景下获得不错的收益。
3.6 知识体系总览
下面这张图，是我自己总结的数据复用优化框架。你可以把它当作一个检查清单，每次优化时对照着看：
这张图里，我把三种技术按照它们利用的局部性类型做了分类。你可能会问：为什么Tiling属于空间局部性？因为分块后，每个块内的数据在内存中是连续的，访问时自然就利用了空间局部性。而Register重用更多是利用时间局部性——同一个数据反复使用。
我的建议：刚开始做优化时，别想着一步到位。先做Tiling，把大问题变小。然后加Shared Memory缓存，把热点数据放进去。最后再看哪些数据可以放寄存器。每一步都能带来可观的收益，而且每一步的改动都不大。
好了，这一章的内容就到这里。数据复用和局部性优化，说白了就是让数据离计算更近。你想想看，如果每次计算都要去显存读数据，那再多的计算单元也是白搭。记住这个核心思想，后面的优化技巧都是围绕它展开的。
想要更多资料，联系可可