内存访问优化策略：数据复用、预取、对齐与Bank冲突避免
各位同学，今天我们来聊聊内存访问优化。说实话，这块内容是我在AI芯片领域摸爬滚打这么多年，觉得最值得花时间讲透的。你想想看，算力再强，数据喂不进去，一切都是白搭。我见过太多项目，算子融合做得漂亮，结果卡在内存访问上，性能直接腰斩。
内存访问优化，说白了就是四个核心策略：数据复用（Tiling）、数据预取（Prefetch）、内存对齐（Coalescing）、Bank冲突避免。咱们一个一个拆开讲。
1. 数据复用（Tiling）—— 把数据“圈”在缓存里
先问大家一个问题：为什么需要Tiling？
我刚开始做AI芯片编译器时，遇到一个矩阵乘法算子。输入矩阵是1024x1024，输出也是1024x1024。直接算，数据量太大，缓存根本装不下。每次计算都要从DDR里反复搬运数据，带宽利用率低得可怜。
Tiling的思路很简单：把大矩阵切成小块，让每个小块能完整放进缓存里。这样，计算过程中数据就在缓存里反复用，不用频繁去DDR“取货”。
核心原则：Tile的大小要刚好能塞进缓存，同时保证计算密度足够高。太小了，搬运开销占比大；太大了，缓存装不下，反而触发换入换出。
我个人的习惯是，先看目标芯片的L1缓存大小。比如L1是32KB，那Tile的尺寸就控制在16KB左右，留一半给其他数据。然后根据算子的计算模式，选择最合适的Tile形状。比如矩阵乘法，我通常选方形Tile，因为访存模式最规整。
举个例子，一个简单的矩阵乘法Tiling伪代码：
// 原始矩阵乘法
for (i = 0; i < N; i++)
for (j = 0; j < N; j++)
for (k = 0; k < N; k++)
C[i][j] += A[i][k] * B[k][j];
// Tiling后的矩阵乘法
for (ii = 0; ii < N; ii += TILE_SIZE)
for (jj = 0; jj < N; jj += TILE_SIZE)
for (kk = 0; kk < N; kk += TILE_SIZE)
for (i = ii; i < ii + TILE_SIZE; i++)
for (j = jj; j < jj + TILE_SIZE; j++)
for (k = kk; k < kk + TILE_SIZE; k++)
C[i][j] += A[i][k] * B[k][j];
你看，外层循环多了三个维度，但每个内层循环的数据量都控制在TILE_SIZE范围内。这样，A和B的子块就能常驻缓存，反复使用。
避坑指南：我曾经在一个项目中，Tile尺寸选得太大，结果缓存频繁miss，性能反而比没做Tiling还差。后来我加了一个profiling步骤，先跑一遍小数据，看看缓存命中率，再调整Tile大小。嗯，这个习惯我一直保留到现在。
2. 数据预取（Prefetch）—— 让数据“提前到位”
数据预取，说白了就是提前把数据从慢速存储搬到快速存储。比如，当前正在计算第i个数据时，提前把第i+1个数据从DDR搬到缓存里。这样，计算单元不用等数据，流水线就能跑满。
我记得有一次优化一个卷积算子，计算量不大，但访存模式很随机。每次计算都要等数据，流水线经常stall。后来我加了软件预取指令，提前把下一轮计算需要的数据搬进来。效果立竿见影，性能提升了30%以上。
预取有两种方式：
- 硬件预取：芯片自动识别访存模式，提前加载数据。优点是透明，缺点是只能处理规则模式。
- 软件预取：编译器或程序员手动插入预取指令。优点是灵活，能处理不规则模式，缺点是需要人工调优。
我个人更倾向于软件预取，尤其是在AI芯片上。因为AI算子的访存模式往往很复杂，硬件预取经常“猜错”。
一个典型的软件预取代码示例：
for (i = 0; i < N; i++) {
// 预取下一轮需要的数据
__builtin_prefetch(&A[i + PREFETCH_DIST], 0, 1);
// 当前计算
result += A[i] * B[i];
}
这里PREFETCH_DIST是预取距离，需要根据内存延迟和计算时间来确定。太近了，数据还没到；太远了，可能把缓存污染了。
注意：预取不是万能的。如果数据依赖关系太复杂，或者访存模式完全随机，预取反而会浪费带宽。我曾经在一个稀疏矩阵计算中尝试预取，结果带宽被预取指令占满，正常计算反而变慢了。所以，预取前一定要分析访存模式。
3. 内存对齐（Coalescing）—— 让访存“合并”起来
内存对齐，或者说访存合并，是GPU和AI芯片上非常重要的优化手段。你想想看，如果每个线程都去访问不同地址，而且地址不连续，那内存控制器就要为每个线程单独发起一次访存请求。带宽利用率极低。
合并访问的核心思想是：让连续的线程访问连续的内存地址。这样，多个访存请求就能合并成一次大的突发传输，带宽利用率大幅提升。
举个例子，在CUDA编程中：
// 非合并访问：线程0访问A[0]，线程1访问A[100]，线程2访问A[200]...
float val = A[threadIdx.x * 100];
// 合并访问：线程0访问A[0]，线程1访问A[1]，线程2访问A[2]...
float val = A[threadIdx.x];
你看，第二种写法，所有线程访问的地址是连续的，内存控制器可以一次把整个cache line搬进来。第一种写法，每个线程访问的地址间隔100个元素，每次都要单独访存，效率差很多。
我建议，在设计数据结构时，尽量把数据按访问顺序连续排列。比如，结构体数组（AoS）和数组结构体（SoA）的选择：
| 布局方式 | 说明 | 合并访问效果 |
|---|---|---|
| AoS（结构体数组） | struct { float x, y, z; } points[N]; | 如果只访问x，则每个结构体只取一个float，浪费带宽 |
| SoA（数组结构体） | float x[N], y[N], z[N]; | 连续访问x数组，完全合并，带宽利用率高 |
在AI芯片上，我通常推荐SoA布局。尤其是做算子融合时，多个算子的中间结果如果能按SoA组织，访存效率会高很多。
经验之谈：我曾经优化过一个Transformer模型，原始代码用的是AoS布局，每个token的embedding放在一个结构体里。改成SoA后，访存带宽利用率从40%提升到85%，推理速度直接快了1.5倍。嗯，这个优化其实很简单，但效果惊人。
4. Bank冲突避免—— 别让“车道”堵死
Bank冲突是共享内存（Shared Memory）特有的问题。共享内存被分成多个Bank，每个Bank可以独立访问。如果多个线程同时访问同一个Bank的不同地址，就会发生Bank冲突，导致访问串行化。
你想想看，共享内存就像一个有多个车道的公路。每个Bank是一个车道。如果所有线程都去同一个车道，那只能一个一个过，速度自然慢。
避免Bank冲突的常用方法：
- 填充（Padding）：在数组的每一行末尾加几个无用元素，让相邻行的相同列落在不同Bank上。
- 改变访问模式：比如，把按列访问改成按行访问，或者用转置的方式重新组织数据。
- 使用向量化访问：一次读取多个元素，减少Bank冲突的概率。
举个例子，一个2D数组的Bank冲突问题：
// 假设共享内存有32个Bank，每个Bank宽度4字节
// 数组定义：__shared__ float data[32][32];
// 冲突访问：线程i访问data[i][j]，如果j相同，则所有线程访问同一Bank
float val = data[threadIdx.x][j]; // 32路Bank冲突！
// 无冲突访问：线程i访问data[i][i]，每个线程访问不同Bank
float val = data[threadIdx.x][threadIdx.x]; // 无冲突
在实际项目中，我经常用填充法。比如，把数组定义成data[32][33]，每行多一个元素。这样，原本冲突的访问模式就变成了无冲突。
避坑指南：我曾经在一个卷积算子的共享内存优化中，忽略了Bank冲突。结果共享内存带宽利用率只有理论值的1/4。后来我用填充法，每行加了一个float的padding，Bank冲突完全消除，性能提升了3倍。所以，别小看这个padding，它可能是性能瓶颈的关键。
知识体系总览
下面这张图，是我自己总结的内存访问优化知识体系。你可以把它当作一个快速参考。
好了，以上就是内存访问优化的四个核心策略。说实话，这些策略单独拿出来都不难，难的是把它们组合起来，针对具体的算子做调优。我个人的经验是，先做profiling，找到瓶颈，再对症下药。别一上来就全上，反而可能互相干扰。
希望今天的分享对你有帮助。如果你在实际项目中遇到内存访问相关的问题，欢迎随时交流。
公众号:蓝海资料掘金营，微信deep3321