硬件架构基础：CPU与GPU的架构差异、内存层次结构、计算单元与线程调度模型
各位同学好，我是老张。今天咱们聊点硬核的——硬件架构。说实话，我见过太多做算子优化的同学，一上来就调代码、改循环，结果性能死活上不去。为什么？因为你根本不了解你的对手是谁。
CPU和GPU，这俩家伙虽然都是芯片，但脾气秉性完全不同。你拿优化CPU那套去搞GPU，就像让短跑运动员去举重——不是不行，但肯定别扭。今天我就把这两位的底细给你扒清楚。
一、CPU vs GPU：一个像管家，一个像包工头
先问个问题：为什么CPU核心那么少，GPU核心那么多？
CPU的设计哲学是「快且全能」。它要处理操作系统调度、中断响应、分支预测、乱序执行……说白了，它得是个多面手。我早期做嵌入式开发时，一个四核CPU跑Linux，既要处理网络包，又要渲染界面，还得跑数据库——它都能扛住。但代价是什么？每个核心都极其复杂，面积大、功耗高。
GPU就不一样了。它就是个包工头，只管一件事：并行计算。你想想看，渲染一张1920×1080的图片，需要处理200多万个像素点。如果让CPU一个个算，黄花菜都凉了。GPU的做法是：招几千个廉价劳动力，每人算一个像素，同时开工。
核心差异一句话总结：
- CPU：少量复杂核心（4-16个），擅长串行任务、复杂逻辑、低延迟
- GPU：大量简单核心（几千个），擅长并行任务、简单运算、高吞吐
我在项目中遇到过一件事：有个同事把一段矩阵乘法直接扔到CPU上跑，用了OpenMP多线程，优化了三天，性能提升30%。后来我建议他试试GPU，用CUDA重写了一遍，性能直接翻了20倍。他当时就懵了——不是他代码写得不好，是工具没选对。
二、内存层次结构：数据搬运的艺术
做算子优化，说白了就是在跟内存打交道。你写的每一行代码，最终都会变成数据在内存和计算单元之间的搬运。搬运得快，性能就好；搬运得慢，再牛的算法也白搭。
咱们先看一张图，这是我手绘的内存层次结构：
这张图你得刻在脑子里。从上往下，容量越来越大，但速度越来越慢。寄存器访问只要1个时钟周期，L1 Cache大概10个周期，L2 Cache上百个周期，到了HBM/DDR就是上千个周期。
我的经验之谈：
做算子优化时，我习惯先问自己三个问题：
- 数据能不能塞进寄存器？能就别放L1。
- 数据能不能复用？能就别反复从HBM读。
- 线程之间有没有共享数据？有就用Shared Memory。
这三个问题想清楚了，性能至少提升50%。
三、计算单元与线程调度模型
CPU的线程调度，大家应该都熟悉。一个四核CPU，开四个线程，每个线程跑在一个核心上。操作系统负责分配，线程切换有开销，但可控。
GPU的调度模型完全不同。拿NVIDIA的GPU举例，它里面有很多个SM（Streaming Multiprocessor）。每个SM里又有几十个CUDA Core。你写一个kernel函数，会启动成千上万个线程。这些线程被组织成Block，Block再被分配到SM上执行。
关键点来了：GPU的线程切换几乎是零开销的。为什么？因为GPU用「大量线程」来隐藏延迟。当一个线程在等内存数据时，SM立刻切换到另一个就绪的线程继续算。你想想看，如果有一万个线程在排队，内存延迟再高，计算单元也闲不下来。
CPU vs GPU 线程调度对比：
| 特性 | CPU | GPU |
|---|---|---|
| 线程数量 | 几十个 | 成千上万 |
| 切换开销 | 高（几百个周期） | 几乎为零 |
| 调度单位 | 线程 | Warp（32线程一组） |
| 隐藏延迟方式 | Cache + 分支预测 | 大量线程并行 |
我记得刚接触CUDA时，写了一个简单的向量加法。启动了一万个线程，心想这肯定快。结果一测，比CPU还慢。后来才发现，我每个线程只算了一个数，数据根本没复用，全在HBM里来回读。这就是典型的「没理解内存层次结构」的坑。
避坑指南：
我曾经犯过一个低级错误：在GPU kernel里用了大量的if-else分支。结果一个Warp里的32个线程走了不同的分支，导致串行执行，性能直接掉了90%。
记住：GPU最怕分支发散。同一个Warp里的线程，最好走相同的代码路径。如果实在避免不了，就尽量让分支条件基于线程ID，这样至少能保证相邻线程走同一条路。
四、实战视角：这些知识怎么用？
说了这么多理论，咱们落地到实际优化中。假设你要优化一个矩阵乘法算子：
// 朴素实现
for (int i = 0; i < M; i++) {
for (int j = 0; j < N; j++) {
float sum = 0.0f;
for (int k = 0; k < K; k++) {
sum += A[i * K + k] * B[k * N + j];
}
C[i * N + j] = sum;
}
}这段代码在CPU上跑，如果矩阵不大，还行。但到了GPU上，你得考虑：
- 数据分块：把大矩阵切成小块，让每个Block处理一块。这样数据能塞进Shared Memory，避免反复读HBM。
- 内存合并访问：相邻线程访问连续内存地址。GPU的显存控制器喜欢连续访问，跳着读会降低带宽利用率。
- 减少同步：Block内部的__syncthreads()是必要的，但别滥用。同步一次，所有线程都得等，性能就掉一截。
我优化过一个类似的算子，从最初的200 GFLOPS优化到了1.2 TFLOPS。怎么做到的？说白了就是三板斧：分块、合并访问、减少同步。但每一板斧背后，都是对硬件架构的深刻理解。
一个小技巧：
拿到一个新算子，我习惯先用NVIDIA的Nsight或者AMD的ROCProfiler跑一下性能分析。看看是计算瓶颈还是访存瓶颈。如果是访存瓶颈，那就优化数据布局；如果是计算瓶颈，那就看看能不能用更快的指令（比如FMA）。
别凭感觉优化，让数据说话。
好了，这一章的内容就到这里。硬件架构是算子优化的地基，地基不稳，上面盖再高的楼也是危房。下一章咱们会深入聊聊「算子性能分析方法论」，到时候我会带大家手把手分析一个实际算子的性能瓶颈。
想要更多资料，联系可可