第二章:硬件基础与计算模型

说实话,做算子优化这些年,我最大的感触是——不懂硬件,你写的代码就像在黑暗中开枪。你可能打中了,但更多时候是浪费子弹。今天我们就来聊聊CPU、GPU、NPU这三兄弟到底有啥不一样,以及它们的存储层次和计算单元是怎么工作的。

2.1 CPU vs GPU vs NPU:三个不同的世界

先说说我个人的理解。CPU像个全能选手,什么活都能干,但一次干不了太多。GPU像个流水线工人,擅长同时干大量重复的活。NPU呢?它是个专才,只干一件事——神经网络计算,但干得特别快。

特性 CPU GPU NPU
核心数量 4-16个(高性能) 数千个 专用计算阵列
控制单元 复杂(分支预测、乱序执行) 简单 极简
缓存 大(L1/L2/L3) 中等(L1/L2) 小(片上SRAM)
典型应用 操作系统、通用计算 图形渲染、并行计算 AI推理、训练

我在项目中遇到过一个问题:用CPU跑一个3x3的卷积,结果比GPU还快。为什么?因为卷积核太小,数据搬运的开销远大于计算本身。你想想看,GPU启动一次kernel就要好几微秒,而CPU几纳秒就搞定了。所以别迷信硬件,得看场景。

2.2 存储层次:数据搬运的艺术

算子优化说白了,就是在和数据搬运做斗争。我经常跟团队说:计算是免费的,数据搬运才是成本。

2.2.1 CPU的存储层次

CPU的缓存设计很讲究。L1 Cache一般32KB-64KB,延迟1-2个时钟周期。L2 Cache 256KB-1MB,延迟10-20个时钟周期。L3 Cache几MB到几十MB,延迟30-50个时钟周期。再往下就是DDR内存了,延迟几百个时钟周期。

关键点:CPU的缓存是透明的,你不需要手动管理。但GPU和NPU不一样,很多时候你得自己管数据在哪。

2.2.2 GPU的存储层次

GPU的存储结构更复杂。以NVIDIA的架构为例:

  • 寄存器:每个线程私有,最快
  • 共享内存(Shared Memory):类似CPU的L1,但可编程控制
  • L1/L2 Cache:硬件管理,但L1可以配置为共享内存
  • HBM(高带宽内存):GPU的主存,带宽极高(可达2TB/s以上)

我记得有一次优化一个矩阵乘法算子,发现瓶颈不在计算,而在共享内存的bank conflict。说白了就是多个线程同时访问同一个bank,导致串行化。调整了一下数据布局,性能直接翻倍。

2.2.3 NPU的存储特点

NPU的存储设计更极端。它通常只有片上SRAM(几MB到几十MB),没有复杂的缓存层次。为什么?因为神经网络的计算模式很规律,数据流可以提前规划好。

我的建议:做NPU算子优化时,重点关注数据复用。把权重和输入数据尽量留在片上,减少和DDR的交互。我曾经在某个NPU上做优化,把数据搬运次数从10次降到2次,性能提升了4倍。

2.3 计算单元:SIMD、Tensor Core、脉动阵列

计算单元是算子的心脏。不同硬件有不同的设计哲学。

2.3.1 SIMD(单指令多数据)

CPU和GPU都支持SIMD。CPU的SIMD指令集有SSE、AVX等,一次可以处理128位、256位甚至512位的数据。GPU的SIMD叫Warp(NVIDIA)或Wavefront(AMD),32个或64个线程一起执行同一条指令。

// CPU SIMD示例:使用AVX2做向量加法
__m256 a = _mm256_load_ps(ptr_a);
__m256 b = _mm256_load_ps(ptr_b);
__m256 c = _mm256_add_ps(a, b);
_mm256_store_ps(ptr_c, c);

嗯,这里要注意:SIMD对数据对齐很敏感。不对齐的数据加载会慢很多。我踩过这个坑,当时一个算子性能上不去,查了半天发现是数据没对齐到32字节。

2.3.2 Tensor Core

Tensor Core是NVIDIA从Volta架构开始引入的专用计算单元。它专门做矩阵乘加运算(D = A * B + C),一次可以处理4x4的矩阵运算。到了Ampere架构,甚至可以处理稀疏矩阵。

性能对比:一个Tensor Core的吞吐量相当于几十个CUDA Core。所以做矩阵乘法时,能用Tensor Core就别用普通指令。

我建议你在写算子时,优先考虑能不能用Tensor Core。比如卷积可以转化为矩阵乘法(im2col),然后就能用Tensor Core加速了。当然,im2col本身有数据重排的开销,需要权衡。

2.3.3 NPU的脉动阵列

NPU的计算单元通常是脉动阵列(Systolic Array)。数据像流水一样在阵列中流动,每个计算单元只做简单的乘加操作。这种设计的好处是:

  • 数据复用率高,减少搬运
  • 控制逻辑简单,面积小
  • 可以做到很高的频率

但缺点也很明显:灵活性差。如果算子的计算模式不匹配脉动阵列,性能会大打折扣。

2.4 带宽与延迟:算子的命门

带宽和延迟是衡量存储系统性能的两个关键指标。带宽决定了你能搬多少数据,延迟决定了你搬一次数据要等多久。

存储层次 典型带宽 典型延迟
L1 Cache ~1TB/s 1-2 ns
L2 Cache ~500GB/s 5-10 ns
HBM ~2TB/s 100-200 ns
DDR4 ~50GB/s 50-100 ns

你可能会问:为什么HBM延迟比DDR还高,但带宽却高那么多?因为HBM通过堆叠和宽接口实现了高带宽,但物理距离远导致延迟高。所以HBM适合流式访问,不适合随机访问。

避坑指南:我曾经在优化一个算子时,用了大量的随机内存访问,结果HBM的带宽利用率只有10%。后来改成连续访问,带宽利用率到了80%以上。记住:连续访问是王道。

2.5 知识体系总览

下面这张图总结了本章的核心内容。我建议你把它打印出来贴在工位上,做算子优化时随时参考。

算子优化硬件基础知识体系 CPU GPU NPU 存储层次 寄存器 → L1 Cache → L2 Cache → L3 Cache → HBM/DDR 延迟递增 · 容量递增 · 带宽递减 计算单元 SIMD (CPU/GPU) Tensor Core (GPU) 脉动阵列 (NPU) 关键指标 带宽 (GB/s) 延迟 (ns)

这张图把CPU、GPU、NPU的存储层次和计算单元放在一起对比。你会发现,虽然三者差异很大,但核心思想是一样的:让数据靠近计算单元,让计算单元满负荷工作。

好了,这一章的内容就到这里。记住我反复强调的那句话:计算是免费的,数据搬运才是成本。下一章我们会深入具体的算子优化技巧,到时候这些硬件知识都会用上。

课后练习:打开你的代码,看看有没有频繁的随机内存访问。试着改成连续访问,测一下性能变化。我保证你会惊讶。

可可精品资料库