第一章：硬件基础与性能模型
各位同学好，我是你们的老朋友。今天咱们聊聊算子优化的根基——硬件基础与性能模型。说实话，我见过太多人一上来就撸代码优化，结果跑出来的性能还不如编译器默认生成的。为什么？说白了，你不了解硬件，就不知道瓶颈在哪。
这一章我会从GPU和CPU的架构讲起，再引入Roofline模型这个分析利器。嗯，内容有点多，但都是干货。我尽量用大白话讲清楚。
1.1 GPU架构概览
先聊聊GPU。很多人觉得GPU就是一堆计算单元堆在一起，其实没那么简单。我刚开始接触CUDA时，也以为把循环扔到GPU上就能快，结果被现实狠狠教育了一顿。
1.1.1 SM（Streaming Multiprocessor）
SM是GPU的核心计算单元。你可以把它想象成一个"小工厂"——里面有多个计算核心、共享内存、寄存器文件、调度器等。每个SM独立执行任务，但彼此之间通过全局内存通信。
我在项目中遇到过一个问题：某个算子占用寄存器太多，导致每个SM能同时运行的线程块数量骤降。性能直接腰斩。后来我学乖了，写算子前先查一下目标GPU的SM资源限制。
关键点：SM的资源（寄存器、共享内存）是有限的。线程块占用的资源越多，SM能同时容纳的线程块就越少，并行度就越低。
1.1.2 Warp
Warp是GPU执行的基本单位。NVIDIA GPU中，32个线程组成一个warp。这32个线程执行相同的指令，但处理不同的数据——这就是SIMT（单指令多线程）模型。
你想想看，如果warp内32个线程走了不同的分支（比如if-else），会发生什么？嗯，所有线程都会执行所有分支，只是不满足条件的线程被"屏蔽"了。这就是warp divergence，性能杀手之一。
避坑指南：我曾经在一个图像处理算子中，让每个线程处理不同颜色的像素，结果分支判断导致warp divergence严重。后来改成让连续线程处理连续像素，性能提升了3倍。
1.1.3 Memory Hierarchy
GPU的内存层次结构，说白了就是"越靠近计算单元，越快但越小"。从快到慢依次是：
- 寄存器：每个线程私有，最快，但数量有限
- 共享内存：同一个线程块内共享，可编程控制
- L1/L2缓存：硬件自动管理，对开发者透明
- 全局内存：所有线程可访问，但延迟最高
我记得有个矩阵乘法的优化案例，核心思路就是把频繁访问的数据从全局内存搬到共享内存。就这么一个改动，性能提升了10倍以上。你想想看，数据搬运的代价有多大。
注意：共享内存虽然快，但容量很小（通常几十KB）。用的时候要精打细算，别一股脑全塞进去。
1.2 CPU架构概览
说完GPU，咱们再看看CPU。虽然现在很多计算都往GPU上搬，但CPU依然是绕不开的。很多算子的前处理、后处理还是在CPU上跑的。
1.2.1 SIMD
SIMD（单指令多数据）是CPU加速的利器。简单说，就是一条指令同时处理多个数据。x86平台的SSE/AVX、ARM平台的NEON都是SIMD指令集。
我刚开始写SIMD代码时，总觉得编译器会自动向量化。后来发现，编译器没那么智能。比如循环中有依赖关系、指针别名等问题，编译器就怂了。这时候就得手动写intrinsic或者内联汇编。
// 一个简单的AVX2向量化示例
// 计算两个float数组的逐元素和
#include <immintrin.h>
void vec_add(float* a, float* b, float* c, int n) {
int i;
for (i = 0; i < n - 8; i += 8) {
__m256 va = _mm256_loadu_ps(&a[i]);
__m256 vb = _mm256_loadu_ps(&b[i]);
__m256 vc = _mm256_add_ps(va, vb);
_mm256_storeu_ps(&c[i], vc);
}
// 处理剩余元素
for (; i < n; i++) {
c[i] = a[i] + b[i];
}
}小技巧：写SIMD代码时，记得保证数据对齐。不对齐的访存虽然也能工作，但性能会下降不少。我习惯用_aligned_malloc或者posix_memalign来分配对齐内存。
1.2.2 Cache
CPU的Cache是另一个性能关键点。L1最快但最小（通常32KB），L2稍大但慢一些，L3更大但更慢。如果数据在L1中命中，延迟只有几个时钟周期；如果跑到主存去取，那就是上百个周期。
你想想看，一个简单的循环，如果数据访问模式是连续的（空间局部性好），Cache命中率就高。如果跳来跳去（比如按列访问二维数组），Cache就频繁miss，性能直接崩。
核心原则：尽量让数据访问具有空间局部性和时间局部性。说白了，就是访问相邻的数据，并且重复使用已经加载到Cache中的数据。
1.3 Roofline模型与性能分析
好了，硬件基础讲完了。但光知道硬件还不够，你得知道怎么分析性能瓶颈。Roofline模型就是干这个的。
1.3.1 什么是Roofline模型
Roofline模型把性能（FLOPS）和计算强度（FLOP/Byte）画在一张图上。横轴是计算强度，纵轴是性能。天花板由两个因素决定：
- 计算上限：硬件的峰值FLOPS
- 带宽上限：硬件的内存带宽 × 计算强度
这两条线组成一个"屋顶"形状，所以叫Roofline。你的算子性能点落在屋顶下方，离屋顶越近，说明优化得越好。
核心思想：如果算子的计算强度低（比如向量加法），瓶颈在带宽，属于"内存受限"；如果计算强度高（比如矩阵乘法），瓶颈在计算，属于"计算受限"。
1.3.2 如何用Roofline指导优化
我在项目中用Roofline模型分析过一个卷积算子。发现它的计算强度只有0.5 FLOP/Byte，明显是带宽受限。于是我做了两件事：
- 用共享内存做数据复用，减少全局内存访问
- 调整线程块大小，提高L1/L2缓存命中率
优化后，计算强度提升到了2.0 FLOP/Byte，性能从屋顶的"带宽段"移到了"计算段"。嗯，效果立竿见影。
避坑指南：我曾经犯过一个错误——盲目增加计算强度。比如在带宽受限的算子中强行做计算复用，结果计算强度上去了，但带宽利用率反而下降了。记住，Roofline模型是指导，不是教条。
1.3.3 绘制Roofline图
下面我用SVG画一个简单的Roofline模型图，帮你直观理解：
这张图里，红色斜线代表带宽上限，绿色水平线代表计算上限。向量加法落在带宽受限区，矩阵乘法落在计算受限区。优化方向完全不同。
1.3.4 性能分析工具
光有模型还不够，你得有工具来获取实际数据。我常用的工具有：
| 工具 | 用途 | 平台 |
|---|---|---|
| nvprof / nsys | GPU性能分析，获取kernel耗时、带宽利用率 | NVIDIA GPU |
| perf | CPU性能分析，获取Cache miss、分支预测等 | Linux |
| Intel VTune | CPU/GPU综合性能分析 | Intel平台 |
| rocprof | AMD GPU性能分析 | AMD ROCm |
个人习惯：我一般先用Roofline模型做理论分析，确定瓶颈方向。然后用性能工具采集实际数据，验证分析结果。最后针对瓶颈做优化。这个流程屡试不爽。
小结
这一章我们聊了GPU的SM、Warp、内存层次结构，CPU的SIMD和Cache，以及Roofline模型。这些东西是算子优化的基础。你想想看，不了解硬件，你怎么知道瓶颈在哪？不了解性能模型，你怎么知道优化方向对不对？
嗯，内容确实不少。但别急，后面我们会一步步把这些知识用到实战中。记住，理论是死的，但人是活的。多动手、多分析、多总结，你也能成为算子优化高手。
想要更多资料，联系可可