4、共享内存与数据复用：利用共享内存减少全局内存访问
好，咱们接着聊共享内存。说实话，很多初学者觉得共享内存就是个“快一点的缓存”，但真正用好它，核心在于一个词——复用。
你想想看，全局内存访问一次要几百个周期，共享内存只要十几个周期。如果你能让一份数据被多个线程反复用，那性能提升就是几何级的。我当年优化一个矩阵乘法时，一开始没做复用，带宽利用率不到10%。后来加了共享内存做分块，直接干到了70%以上。嗯，那种感觉，很爽。
4.1 为什么需要数据复用？
先看一个简单场景。假设每个线程都要读一个数组的相邻元素：
// 朴素版本：每个线程独立读全局内存
__global__ void stencil_1d(float *in, float *out, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i > 0 && i < N-1) {
out[i] = (in[i-1] + in[i] + in[i+1]) / 3.0f;
}
}这里每个线程读3个元素。但相邻线程读的元素是重叠的——线程i和线程i+1都读了in[i]和in[i+1]。说白了，全局内存被反复读了多次。
我刚开始做这个优化时，觉得“反正GPU带宽大，多读几次无所谓”。结果一测，性能惨不忍睹。为什么？因为全局内存的带宽是共享的，你多读一次，别人就得等一次。
4.2 数据分块(Tiling)技术
数据分块，说白了就是把大问题切成小块，让每个块的数据能塞进共享内存。然后块内的线程协作，把数据一次性从全局内存搬到共享内存，后面所有计算都在共享内存上完成。
拿上面的1D stencil举例，分块后的代码长这样：
#define TILE_SIZE 32
#define RADIUS 1
__global__ void stencil_1d_tiled(float *in, float *out, int N) {
__shared__ float smem[TILE_SIZE + 2 * RADIUS];
int global_idx = blockIdx.x * TILE_SIZE + threadIdx.x;
int local_idx = threadIdx.x + RADIUS;
// 第一步：加载数据到共享内存（包括边界）
smem[local_idx] = in[global_idx];
if (threadIdx.x < RADIUS) {
smem[local_idx - RADIUS] = in[global_idx - RADIUS];
smem[local_idx + TILE_SIZE] = in[global_idx + TILE_SIZE];
}
__syncthreads();
// 第二步：从共享内存计算
if (global_idx >= RADIUS && global_idx < N - RADIUS) {
out[global_idx] = (smem[local_idx-1] + smem[local_idx] + smem[local_idx+1]) / 3.0f;
}
}这里有个关键点：__syncthreads()。我见过太多人忘记加这个同步，结果线程B读共享内存时，线程A还没写完。数据全是乱的。嗯，这个坑我踩过不止一次。
核心思想：全局内存只读一次，共享内存反复用。每个数据元素被加载到共享内存后，可以被块内所有线程访问多次。
4.3 Tile大小对性能的影响
Tile大小不是越大越好，也不是越小越好。这里有个微妙的平衡。
| Tile大小 | 优点 | 缺点 |
|---|---|---|
| 16x16 | 占用共享内存少，占用寄存器少 | 数据复用率低，带宽利用率一般 |
| 32x32 | 复用率高，计算密度大 | 共享内存占用大，可能降低occupancy |
| 64x64 | 复用率极高 | 共享内存不够用，block数太少，隐藏延迟能力差 |
我个人习惯的做法是：先算共享内存上限。比如一个SM有48KB共享内存，每个block用16KB，那最多同时跑3个block。如果每个block用32KB，那就只能跑1个block。occupancy从100%掉到33%，延迟隐藏能力大打折扣。
我曾经优化一个2D卷积，一开始用了64x64的tile，觉得复用率很高。结果一测，性能还不如32x32。为什么？因为occupancy太低，内存延迟藏不住，计算单元经常空转。后来改成32x32，occupancy上去了，性能反而提升了30%。
经验法则：Tile大小选择，优先保证occupancy不低于50%。然后在这个前提下，尽量增大tile以提高数据复用率。具体数值需要实测，不同GPU架构的最佳值不同。
4.4 避坑指南
我总结几个常见的坑，都是我在项目中踩过的：
- 边界处理不当：分块时，边界元素需要额外加载。比如stencil半径为1，每个tile需要多加载左右各1个元素。忘记处理边界，结果就是数组越界，数据全错。
- bank conflict：共享内存有32个bank，如果多个线程访问同一个bank的不同地址，就会产生冲突。我遇到过一种情况：tile大小设为32，但stride也是32，结果所有线程都访问同一个bank，性能直接腰斩。解决办法是加padding，比如把
smem[32][32]改成smem[32][33]。 - 寄存器溢出：tile太大时，每个线程需要的寄存器变多。如果寄存器不够用，编译器会把数据 spill 到本地内存（其实还是全局内存），性能反而下降。我建议用
--ptxas-options=-v编译选项查看寄存器使用量。
注意：共享内存不是免费的午餐。它占用了SM的资源，如果使用不当，反而会降低性能。我的建议是：先做profiling，找到真正的瓶颈是计算还是访存。如果访存是瓶颈，再考虑用共享内存优化。
4.5 实战建议
如果你刚开始接触共享内存优化，我建议从简单的1D stencil开始。先写一个朴素版本，测一下带宽。然后加上共享内存分块，再测一下。对比两个版本的性能，你就能直观感受到共享内存的威力。
然后可以尝试2D矩阵乘法，这是最经典的tiling案例。我记得第一次把矩阵乘法优化到接近理论峰值时，那种成就感，嗯，比打游戏通关还爽。
最后说一句：不要迷信理论计算。理论带宽和实际带宽之间，往往差着一个数量级。唯一能信的是profiler的数据。用ncu或者nvprof看看L1/L2命中率、共享内存bank conflict、occupancy这些指标，然后针对性地优化。
好了，这一章就到这里。下一章我们会聊bank conflict的深入分析与优化，到时候我会分享一个我花了三天才解决的bank conflict案例，保证让你印象深刻。