硬件适配基础：GPU、CPU、NPU架构概览
做算子融合，说白了就是跟硬件谈恋爱。你得摸清对方的脾气，知道它喜欢什么、讨厌什么。今天我就带大家看看这三个“对象”——GPU、CPU、NPU，它们各自的核心架构长什么样。
GPU架构概览：SM、Warp与Shared Memory
GPU这玩意儿，天生就是为并行而生的。我刚开始接触CUDA时，总觉得它就是个“很多核的CPU”，后来才发现完全不是那么回事。
GPU的核心执行单元叫SM（Streaming Multiprocessor）。你可以把SM想象成一个“小工厂”，里面有很多计算单元（CUDA Core），还有共享内存、寄存器文件这些资源。一个GPU卡上通常有几十到上百个SM。
那Warp是什么呢？它是GPU调度线程的基本单位。NVIDIA的GPU里，32个线程组成一个Warp。这32个线程执行同一条指令，但处理不同的数据——这就是SIMT（单指令多线程）。
关键点：Warp内的线程必须步调一致。如果出现分支发散（if-else），部分线程会等待，性能直接打折扣。我在优化一个图像处理算子时，就因为一个简单的条件判断，导致Warp利用率掉了30%。
再说Shared Memory。它是SM内部的一块高速缓存，由同一个Block内的所有线程共享。延迟比全局内存低一个数量级，但容量有限（通常几十KB）。
我个人的习惯是：能用Shared Memory就别用全局内存。比如做矩阵分块时，先把数据从全局内存搬到Shared Memory，计算完再写回去。这招在矩阵乘法里特别管用。
避坑指南：我曾经遇到过Shared Memory的Bank Conflict问题。简单说，就是多个线程同时访问同一个Bank的不同地址，导致串行化。解决办法是调整数据布局，或者用padding技巧。
GPU的内存层级大致是这样的：
| 层级 | 延迟 | 容量 | 作用域 |
|---|---|---|---|
| 寄存器 | ~1 cycle | 每个线程有限 | 线程私有 |
| Shared Memory | ~20 cycles | 几十KB | Block内共享 |
| L1 Cache | ~30 cycles | 几十KB | SM内 |
| L2 Cache | ~200 cycles | 几MB | 芯片级 |
| 全局内存 | ~400-800 cycles | 几GB到几十GB | 所有线程 |
CPU架构概览：Cache层级与SIMD
CPU和GPU的思路完全不同。CPU追求低延迟、高单核性能，适合处理复杂逻辑和串行任务。
Cache层级是CPU性能的关键。现代CPU通常有L1、L2、L3三级缓存。L1最快但最小（几十KB），L3最大但稍慢（几MB到几十MB）。
为什么Cache这么重要？因为内存访问的速度差距太大了。CPU从L1读数据只要几个时钟周期，但从主存读要几百个周期。你想想看，如果数据不在Cache里，CPU就得干等着。
实战经验：我在做算子融合时，经常要考虑数据的局部性。比如把两个连续的操作融合在一起，让中间结果留在L1 Cache里，而不是写回主存。这能省下大量带宽。
再说SIMD（单指令多数据）。CPU也有并行能力，不过是通过SIMD指令集实现的，比如x86的AVX、ARM的NEON。一条SIMD指令可以同时对多个数据执行相同操作。
举个例子，用AVX-512做向量加法：
// 假设两个float数组，长度是16的倍数
__m512 a = _mm512_load_ps(ptr_a);
__m512 b = _mm512_load_ps(ptr_b);
__m512 c = _mm512_add_ps(a, b);
_mm512_store_ps(ptr_c, c);这一下就处理了16个float。不过要注意，SIMD对数据对齐很敏感。我踩过坑：没对齐的数据加载，性能直接腰斩。
注意：CPU的SIMD宽度有限（AVX-512是512位），而且不适合处理不规则的数据结构。如果数据访问模式是随机的，SIMD反而帮倒忙。
NPU架构概览：脉动阵列与数据流
NPU（神经网络处理器）是专门为深度学习设计的。它跟GPU、CPU都不一样，核心思想是数据流驱动计算。
脉动阵列（Systolic Array）是NPU最经典的结构。想象一个二维的网格，每个格子是一个乘累加单元（MAC）。数据像血液一样在阵列里“脉动”流动，每个MAC只跟邻居通信。
这样做的好处是什么？数据复用率高。比如做矩阵乘法时，权重可以“驻留”在阵列里，输入数据流过去，输出结果流出来。这大大减少了内存访问。
核心优势：脉动阵列的能效比非常高。我见过一个NPU芯片，在同样算力下，功耗只有GPU的十分之一。这就是专用架构的魅力。
NPU的数据流方式也很有特色。常见的有三种：
- 权重固定（Weight Stationary）：权重留在MAC里，输入数据流动。适合卷积层。
- 输入固定（Input Stationary）：输入数据留在MAC里，权重流动。适合全连接层。
- 输出固定（Output Stationary）：部分和留在MAC里，输入和权重都流动。适合深度可分离卷积。
我建议你在做算子融合时，一定要考虑NPU的数据流特点。比如，把连续的卷积和激活函数融合，让数据在脉动阵列里一次处理完，避免中间结果写回DRAM。
个人经验：我曾经把一个BatchNorm融合进卷积里，在NPU上跑，速度提升了40%。原因就是减少了数据在阵列和内存之间的搬运次数。
下面这张图展示了三种架构的核心差异：
总结一下：GPU靠大量SM和Warp堆算力，CPU靠Cache和SIMD提效率，NPU靠脉动阵列和数据流降功耗。做算子融合时，你得根据目标硬件的特性，选择不同的融合策略。
比如在GPU上，我会优先考虑减少全局内存访问，多用Shared Memory。在CPU上，我会尽量让数据留在L1 Cache里，同时用SIMD加速。在NPU上，我会把连续的层融合成一个大的数据流操作。
嗯，这些就是硬件适配的基础。理解了这些，后面讲具体的融合策略时，你就能明白为什么有些融合在GPU上效果好，在CPU上却不行了。