3、算子访存模式分析：连续访问、随机访问、广播、归约
各位同学，咱们今天聊点实在的。算子跑不快，十有八九是内存带宽在拖后腿。我见过太多团队，算子逻辑写得天衣无缝，一上NPU就卡成PPT。为什么？说白了，就是没搞懂访存模式对带宽的消耗差异。
我个人习惯，拿到一个新算子，第一件事不是看计算逻辑，而是先画它的访存模式图。你想想看，NPU的计算单元再快，数据喂不进去也是白搭。今天我就把这四种最常见的访存模式——连续、随机、广播、归约——掰开了揉碎了讲清楚。
核心观点：访存模式决定了你能榨出多少有效带宽。连续访问是亲儿子，随机访问是败家子，广播和归约则是双刃剑。
3.1 连续访问：带宽利用率的天花板
连续访问，就是按内存地址顺序挨个读。这是NPU最喜欢的模式，没有之一。为什么？因为硬件预取器能完美预测下一个要读的地址，DDR的burst传输也能拉满。
我在项目中遇到过，一个简单的向量加法，连续读写，带宽利用率能跑到理论峰值的95%以上。但你要是稍微跳个步长，立马掉到60%。
// 连续访问示例：向量加法
for (int i = 0; i < N; i++) {
C[i] = A[i] + B[i]; // A、B、C都是连续地址
}| 访存模式 | 带宽利用率 | 典型场景 |
|---|---|---|
| 连续访问 | 90%-100% | 向量运算、矩阵乘法连续维 |
| 跨步访问 | 30%-60% | 矩阵转置、图像通道分离 |
| 随机访问 | <10% | 稀疏矩阵、索引查找 |
3.2 随机访问：带宽杀手，能避则避
随机访问，就是地址毫无规律地跳来跳去。每次访存都要重新打开行、列，DDR的预充电和激活延迟能把性能打骨折。
我曾经优化过一个稀疏卷积算子，原始实现就是随机访问，带宽利用率只有可怜的3%。后来改成先对索引排序，再连续读取，直接干到了70%。
避坑指南：我曾经踩过一个坑——在NPU上用哈希表做特征查找。每次查表都是随机访问，带宽直接崩了。后来改成连续内存的二分查找，虽然计算量大了点，但带宽利用率翻了好几倍。
// 随机访问示例：稀疏矩阵的CSR格式读取
for (int i = 0; i < nnz; i++) {
int col = col_indices[i]; // col_indices是随机分布的
float val = values[i];
output[col] += val; // 对output的访问也是随机的
}3.3 广播：一次读取，多处使用
广播模式，就是同一个数据被多个计算单元同时使用。比如卷积里的权重，或者Batch Normalization里的均值和方差。
NPU对广播有专门的硬件支持——通常会把广播数据缓存在共享内存或者寄存器文件中。但要注意，广播的数据量不能太大，否则会挤占其他数据的缓存空间。
我的经验：广播数据尽量放在片上缓存，别让它反复去DDR读。我习惯把卷积的权重提前搬到共享内存，这样每个像素计算时都是片上访问，带宽消耗几乎为零。
// 广播示例：卷积中的权重复用
for (int oh = 0; oh < OH; oh++) {
for (int ow = 0; ow < OW; ow++) {
float sum = 0;
for (int kh = 0; kh < KH; kh++) {
for (int kw = 0; kw < KW; kw++) {
// weight[kh][kw] 被所有输出像素广播使用
sum += input[oh+kh][ow+kw] * weight[kh][kw];
}
}
output[oh][ow] = sum;
}
}3.4 归约：多对一的带宽压力
归约模式，就是多个数据经过运算合并成一个。比如求和、求最大值。归约的带宽消耗很特殊——读的时候是连续或随机的，但写的时候只有一个目标地址。
这里有个坑：多个线程同时写同一个地址，会产生写冲突。NPU通常用原子操作或者树形归约来解决，但原子操作会串行化，带宽利用率直线下降。
我记得有一次优化LayerNorm，原始的归约实现是每个线程都去写全局内存的同一个位置，结果带宽利用率不到5%。后来改成先用共享内存做局部归约，再全局归约，带宽利用率提升到了60%。
// 归约示例：向量求和（低效版本）
float sum = 0;
for (int i = 0; i < N; i++) {
sum += data[i]; // 连续读，但只有一个写目标
}
// 高效版本：分块归约
float block_sum[BLOCK_SIZE];
for (int i = tid; i < N; i += BLOCK_SIZE) {
block_sum[tid] += data[i];
}
__syncthreads();
// 树形归约
for (int stride = BLOCK_SIZE/2; stride > 0; stride >>= 1) {
if (tid < stride) {
block_sum[tid] += block_sum[tid + stride];
}
__syncthreads();
}3.5 四种模式的带宽消耗对比
咱们直接上数据，看看这四种模式在典型NPU上的表现差异：
| 访存模式 | 读带宽利用率 | 写带宽利用率 | 典型延迟 | 优化难度 |
|---|---|---|---|---|
| 连续访问 | 95% | 95% | 低 | 低 |
| 随机访问 | 5%-15% | 5%-15% | 极高 | 高 |
| 广播 | 一次读取，N次使用 | N/A | 低（片上缓存） | 中 |
| 归约 | 90%（连续读） | 10%-30%（写冲突） | 中 | 中 |
总结一下：
- 连续访问是王道，能连续就别跳步
- 随机访问是毒药，能排序就别乱读
- 广播要善用片上缓存，别让数据反复从DDR来
- 归约要分块处理，避免全局写冲突
嗯，这四种模式基本覆盖了NPU上90%的算子访存场景。你下次拿到一个新算子，先对着这四种模式画个访存图，带宽瓶颈在哪，一眼就能看出来。
好了，这一节就到这。记住，分析算子性能，先看访存模式。模式对了，带宽自然就上去了。