第十六章 C++语言在深度学习中的应用
1. C++语言的优势
1.1 编译到底,中间没人抽成
C++ 程序从源码到运行,走一条直线。编译器把 .cpp 翻译成机器码,链接器拼成可执行文件,运行的时候 CPU 直接执行机器指令,中间没有第三个人。
┌─────────────────────────────────────────────────────────────┐ │ 从源码到执行 │ ├─────────────────────────────────────────────────────────────┤ │ 运行路径,一条直线对两道手续 │ │ │ │ ┌────────────────────────┐ ┌────────────────────────┐ │ │ │ C++ │ │ 解释型语言 │ │ │ │ │ │ │ │ │ │ 源码 │ │ 源码 │ │ │ │ v │ │ v │ │ │ │ 编译器 │ │ 解释器读一句 │ │ │ │ v │ │ v │ │ │ │ 机器码 │ │ 机器执行一句 │ │ │ │ v │ │ v │ │ │ │ CPU 直接执行 │ │ 每次都付读的钱 │ │ │ └────────────────────────┘ └────────────────────────┘ │ │ │ │ 循环一亿次,解释器读一亿次句子 │ └─────────────────────────────────────────────────────────────┘解释型语言多一道手续。每行代码运行时先过解释器,解释器读懂一句,机器执行一句。读句子的开销每次都付,循环一亿次就付一亿次。深度学习里一个矩阵乘法就是上亿次乘加,这道手续付不起。
C++ 的编译期干的活不止翻译。模板在编译期展开,inline 函数在编译期铺平,常量表达式在编译期算完。到了运行期,只剩下非算不可的部分。这是第一章说的编译型语言的底子,到本章才真正兑现它的价值。
1.2 内存布局自己说了算
C++ 的数组是一块连续内存,第几个元素在第几个字节,程序员心里能算出来。二维数组按行主序存,第一行摆完摆第二行,同一行的元素在内存里挨着。
┌─────────────────────────────────┐ │ 行主序,同一行挨着摆 │ ├─────────────────────────────────┤ │ 二维数组在内存里的真实摆法 │ │ │ │ 逻辑上 a[2][3] 是方阵 │ │ │ │ a00 a01 a02 │ │ a10 a11 a12 │ │ │ │ 内存里是一条直线,按行摆 │ │ │ │ [a00 a01 a02 a10 a11 a12] │ │ <第一行> <第二行> │ │ │ │ 顺着行读,地址连着走 │ │ 跳着列读,一步跨一整行 │ └─────────────────────────────────┘CPU 读内存有个脾气,一次不只读一个数,它把目标地址附近的一整段都搬进缓存,一段通常 64 字节。下一个数要是就在这段里,缓存直接给,快。不在,就得再去内存搬,慢一百倍。这就叫缓存命中和不命中。
回到师兄那两行代码。i-j-k 顺序里,最内层 k 循环挨个读 b[k][j],k 走一步跳一行,每个元素都隔着一整行内存,缓存次次扑空。i-k-j 顺序里,最内层 j 循环顺着 b[k][j] 读,j 走一步挪一个元素,全在缓存搬进来的那段里。算法没变,踩内存的姿势变了。
#include <iostream> #include <chrono> using namespace std; const int N = 512; int main() { static double a[N][N], b[N][N], c[N][N]; for (int i = 0; i < N; i++) { for (int j = 0; j < N; j++) { a[i][j] = (i + j) % 7 * 0.1; b[i][j] = (i * j) % 5 * 0.2; } } auto t1 = chrono::steady_clock::now(); for (int i = 0; i < N; i++) { for (int j = 0; j < N; j++) { double s = 0.0; for (int k = 0; k < N; k++) { s += a[i][k] * b[k][j]; } c[i][j] = s; } } auto t2 = chrono::steady_clock::now(); cout << "i-j-k 顺序,c[511][511] = " << c[511][511] << ",耗时 " << chrono::duration<double>(t2 - t1).count() << " 秒" << endl; for (int i = 0; i < N; i++) { for (int j = 0; j < N; j++) { c[i][j] = 0.0; } } t1 = chrono::steady_clock::now(); for (int i = 0; i < N; i++) { for (int k = 0; k < N; k++) { double r = a[i][k]; for (int j = 0; j < N; j++) { c[i][j] += r * b[k][j]; } } } t2 = chrono::steady_clock::now(); cout << "i-k-j 顺序,c[511][511] = " << c[511][511] << ",耗时 " << chrono::duration<double>(t2 - t1).count() << " 秒" << endl; return 0; }512 乘 512 的小矩阵,我机器上 i-j-k 跑 0.45 秒,i-k-j 跑 0.034 秒,差 13 倍,两种顺序算出的 c[511][511] 都是 61.18,结果一模一样。矩阵再大些,差距拉到几十倍,师兄那七个半小时对十七分钟就是这么来的。
┌──────────────────────────────────────────┐ │ 换个顺序,十三倍 │ ├──────────────────────────────────────────┤ │ 512 x 512 双精度矩阵,本机实测 │ │ │ │ +------------+------------+----------+ │ │ |循环顺序 |耗时 |结果 | │ │ +------------+------------+----------+ │ │ |i-j-k |0.45 秒 |61.18 | │ │ |i-k-j |0.034 秒 |61.18 | │ │ +------------+------------+----------+ │ │ │ │ 得数分毫不差,快慢差 13 倍 │ │ 跳列读的那个,缓存次次扑空 │ │ 顺行读的那个,缓存全程命中 │ └──────────────────────────────────────────┘这段代码还有两处 C++ 特有的讲究。数组用 static 摆进静态区,避开栈大小的限制,这是第六章指针那章聊过的地盘划分。计时用 chrono,标准库的时钟,不用引任何第三方。C++ 给你的不是现成的快,是让你够得着快的所有开关。
1.3 零开销抽象
有人担心 C++ 的抽象要付代价,类啊模板啊虚函数啊,听着就比 C 慢。C++ 的设计原则里有一条叫零开销抽象,用不上的特性一分钱不付,用得上的特性编译期能算就不留到运行期。
| 特性 | 抽象层 | 运行期代价 |
|---|---|---|
| 模板 | 编译期展开,一种类型一份代码 | 零 |
| inline 函数 | 编译期铺平,省掉调用 | 零 |
| RAII | 析构函数自动收尾 | 和手写释放一样 |
| 虚函数 | 查虚表跳转 | 一次间接寻址 |
| 异常 | 不抛时不走这条路 | 接近零 |
虚函数是少数要付钱的,第十二章聊过它查表的跳转,可正是这次跳转换来了多态。C++ 的账本公开,每一分钱花在哪,你能查到。这就是深度学习框架敢把底层交给它的原因。
2. C++在AI应用中的优势
2.1 训练和推理是两种活
AI 应用分两半。训练,把数据喂进去调权重,第十五章干的就是这件事的迷你版。推理,权重调好了不再动,来一个新样本算一个结果。
┌─────────────────────────────────────────────────────────┐ │ 两种活,两张账单 │ ├─────────────────────────────────────────────────────────┤ │ 一份模型,两种脾气 │ │ │ │ ┌──────────────────────┐ ┌──────────────────────┐ │ │ │ 训练 │ │ 推理 │ │ │ │ │ │ │ │ │ │ 数据灌进去 │ │ 权重冻结 │ │ │ │ 权重一轮轮调 │ │ 来一个算一个 │ │ │ │ 跑在机房 │ │ 跑在设备上 │ │ │ │ 重吞吐 │ │ 重延迟 │ │ │ └──────────────────────┘ └──────────────────────┘ │ │ │ │ 训练慢点能忍,推理慢半秒就砸招牌 │ └─────────────────────────────────────────────────────────┘两半的脾气完全不一样。
| 对比项 | 训练 | 推理 |
|---|---|---|
| 频率 | 偶尔跑一次,跑很久 | 每来一个请求跑一次 |
| 在乎什么 | 吞吐,单位时间吞多少数据 | 延迟,单个请求多快回 |
| 跑在哪 | 机房,显卡成排 | 手机、车、摄像头、服务器 |
| 权重 | 每轮更新 | 冻结不动 |
师兄的活是训练,算完一批就结束,慢点能忍。你手机里的人脸解锁是推理,按下去半秒没反应你就想摔手机,这种活对延迟和设备都挑剔。挑剔的地方,就是 C++ 出场的地方。
2.2 推理落地的地方,都是 C++ 的地盘
看推理都落在哪些设备上,再看这些设备的系统用什么写的。
| 落地场景 | 设备条件 | 为什么是 C++ |
|---|---|---|
| 手机 APP | 内存几个 G,要省电 | 直接调系统接口,无虚拟机开销 |
| 车载系统 | 实时性硬性要求 | 延迟可控,行为可预测 |
| 摄像头、无人机 | 嵌入式芯片,资源紧张 | 可裁剪,能跑在裸系统上 |
| 服务器推理引擎 | 高并发,机器贵 | 榨干每一个核,省钱 |
这些场景有个共同点,资源紧或者要求硬,没有多余算力养解释器和虚拟机。操作系统本身是 C 和 C++ 写的,C++ 程序和系统之间没有翻译官,调显卡驱动、调硬件加速指令,都是直达。
2.3 Python 和 C++ 的分工
那 Python 在 AI 里那么火,是不是矛盾。不矛盾,分工不同。研究和训练阶段,代码天天改,要的是写得快,Python 顺手。部署和推理阶段,代码定型了,要的是跑得快,C++ 顶上。
┌───────────────────────────────────┐ │ 写得快归 Python,跑得快归 C++ │ ├───────────────────────────────────┤ │ 一个 AI 项目的接力 │ │ │ │ 研究阶段 部署阶段 │ │ Python 写原型 C++ 扛上线 │ │ 代码天天改 代码定型 │ │ 要写得快 要跑得快 │ │ │ │ v 模型定型,交接 │ │ │ │ 壳子用 Python,引擎用 C++ │ └───────────────────────────────────┘真实的项目里这两半常常在同一个框架里接力。研究员用 Python 搭网络、调参数,那是前端壳子。壳子底下,真正算矩阵的那部分是 C++ 写的。第 4 节拆开这个壳子看。
3. 面向GPU的C++编程
3.1 GPU 为什么算得快
CPU 是几个很强的大核,每个核能独立干复杂的活,适合逻辑绕来绕去的程序。GPU 反过来,成百上千个小核,每个核只会干简单活,但架不住核多。
┌─────────────────────────────────────────────────────────────┐ │ 几个大核对上千小核 │ ├─────────────────────────────────────────────────────────────┤ │ 两种芯片,两种打法 │ │ │ │ ┌────────────────────┐ ┌────────────────────────────┐ │ │ │ CPU │ │ GPU │ │ │ │ │ │ │ │ │ │ [大核][大核] │ │ [小小小小小小小小] │ │ │ │ [大核][大核] │ │ [小小小小小小小小] │ │ │ │ │ │ │ │ │ │ 几个核 │ │ 上千个小核 │ │ │ │ 个个能打 │ │ 一人一份简单活 │ │ │ └────────────────────┘ └────────────────────────────┘ │ │ │ │ 矩阵乘法二十六万份简单活 │ │ 人多好干活,GPU 赢 │ └─────────────────────────────────────────────────────────────┘矩阵乘法恰好是 GPU 的菜。512 乘 512 的矩阵有二十六万个元素,每个元素的计算互不依赖,你算你的我算我的。这种活叫并行,二十六万份简单活分给上千个小核同时干,比几个大核排队干快得多。深度学习里九成计算都是矩阵乘加,GPU 就是为这口饭生的。
3.2 CUDA 是 C++ 的方言
NVIDIA 给 GPU 编程造的工具叫 CUDA,它不是一门新语言,是 C++ 的扩展。普通函数照写,想在 GPU 上跑的函数前面加global,调用时用三角括号告诉它开多少线程。
#include <iostream> using namespace std; // 本段是示意代码,需要 NVIDIA 的 nvcc 编译器,普通 g++ 编不了 __global__ void vecAdd(const float *a, const float *b, float *c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) { c[i] = a[i] + b[i]; } } int main() { int n = 1024; float *a, *b, *c; cudaMallocManaged(&a, n * sizeof(float)); cudaMallocManaged(&b, n * sizeof(float)); cudaMallocManaged(&c, n * sizeof(float)); for (int i = 0; i < n; i++) { a[i] = 1.0f; b[i] = 2.0f; } vecAdd<<<4, 256>>>(a, b, c, n); cudaDeviceSynchronize(); cout << "c[0] = " << c[0] << ",c[1023] = " << c[1023] << endl; cudaFree(a); cudaFree(b); cudaFree(c); return 0; }骨架全是老朋友。cudaMallocManaged 像 malloc,第六章的手艺。cudaFree 像 free,new 和 delete 的老规矩,借了就还。真正新的只有三样,global标记这个函数在 GPU 上跑,三角括号 <<<4, 256>>> 说开 4 个线程块每块 256 个线程,threadIdx 和 blockIdx 是每个线程问自己是谁。
┌──────────────────────────────────────┐ │ 一千零二十四个线程一起上 │ ├──────────────────────────────────────┤ │ vecAdd<<<4, 256>>> 的兵力部署 │ │ │ │ 线程块 0 线程 0 ... 线程 255 │ │ 线程块 1 线程 0 ... 线程 255 │ │ 线程块 2 线程 0 ... 线程 255 │ │ 线程块 3 线程 0 ... 线程 255 │ │ │ │ 4 块 x 256 = 1024 个线程 │ │ i = blockIdx.x*256 + threadIdx.x │ │ 一线程一份,互不等待 │ └──────────────────────────────────────┘1024 份加法,4 块乘 256 线程,正好 1024 个线程一人一份,同时做完。CPU 版本要写循环挨个加,GPU 版本每个线程只加自己那一份,谁也不用等谁。这就是 SIMT 的意思,同一条指令,成百上千个线程各算各的数据。
3.3 数据搬家比计算还贵
GPU 有自己的内存,叫显存,和 CPU 用的内存是两间仓库。数据要从内存搬到显存,GPU 才算得着,算完再搬回来。搬家的路是 PCIe 总线,比内存慢一个量级。
┌──────────────────────────────────────┐ │ 搬家费要先算进账 │ ├──────────────────────────────────────┤ │ 两间仓库,一条路 │ │ │ │ CPU 内存 GPU 显存 │ │ 数据在这 ==PCIe==> 算力在这 │ │ <==PCIe== 结果搬回 │ │ │ │ 搬家比计算慢一个量级 │ │ 一次搬定,千次摊薄 │ │ 来回搬,搬出亏本买卖 │ └──────────────────────────────────────┘这就有个账本要算。计算量小、数据量大,搬家费比算出来的价值还贵,这种活别上 GPU。计算量大、数据能留在显存里反复用,搬家一次摊薄到几千次计算上,GPU 才赚。深度学习的矩阵一进显存就反复参与成百上千次乘加,天然符合后面这种。师兄那批传感器数据也是,搬一次,算到底。
4. C++在深度学习框架底层设计中的应用
4.1 框架都是两层皮
PyTorch、TensorFlow 这些框架,上手都写 Python,可 Python 只是壳。掀开壳子,底下真正干活的引擎是 C++。
┌─────────────────────────────────┐ │ 壳子底下是引擎 │ ├─────────────────────────────────┤ │ 掀开 PyTorch 的壳 │ │ │ │ 上层 Python API │ │ 研究员在这层写网络 │ │ ──────────────────────── │ │ 中层 计算图与自动微分,C++ │ │ 张量调度在这层 │ │ ──────────────────────── │ │ 底层 算子库,C++ 与 CUDA │ │ 矩阵乘在这层真算 │ └─────────────────────────────────┘这么设计是两头都要。研究员要写得快,Python 壳子给。机器要跑得狠,C++ 引擎给。你在 Python 里写一句张量相乘,调用顺着壳子下去,最后落在 C++ 写的矩阵乘算子上,那个算子里就藏着本章第 1 节那种循环顺序的讲究。框架的作者把脏活累活干完了,用户只看见一句乘法。
4.2 张量、一个 struct 的学问
框架里到处飞的张量,说穿了是多维数组加个壳。一个最朴素的张量,三样东西就够,一维数据数组,形状 shape,步长 stride。
#include <iostream> #include <vector> using namespace std; struct Tensor { vector<float> data; vector<int> shape; vector<int> stride; Tensor(int rows, int cols) { shape = {rows, cols}; stride = {cols, 1}; data.assign(rows * cols, 0.0f); } float &at(int r, int c) { return data[r * stride[0] + c * stride[1]]; } void show() const { cout << "shape (" << shape[0] << ", " << shape[1] << ")"; cout << ",stride (" << stride[0] << ", " << stride[1] << ")" << endl; } }; int main() { Tensor t(2, 3); t.show(); float v = 1.0f; for (int r = 0; r < 2; r++) { for (int c = 0; c < 3; c++) { t.at(r, c) = v; v += 1.0f; } } for (int r = 0; r < 2; r++) { for (int c = 0; c < 3; c++) { cout << t.at(r, c) << " "; } cout << endl; } cout << "底层一维数组:"; for (float x : t.data) { cout << " " << x; } cout << endl; cout << "元素 (1, 2) 的地址偏移 = 1*" << t.stride[0] << " + 2*" << t.stride[1] << " = " << 1 * t.stride[0] + 2 * t.stride[1] << endl; return 0; }跑出来的样子。
shape (2, 3),stride (3, 1) 1 2 3 4 5 6 底层一维数组: 1 2 3 4 5 6 元素 (1, 2) 的地址偏移 = 1*3 + 2*1 = 5data 是一维的,2 行 3 列只是 data 的一种看法。stride 记着每个维度走一步要跳几个元素,第 0 维走一步跳 3 个,第 1 维走一步跳 1 个。元素 (1, 2) 的偏移就是 1 乘 3 加 2 乘 1,等于 5,数数一维数组第 5 个位置,正好是 6。
┌────────────────────────────────────────┐ │ 张量三件套 │ ├────────────────────────────────────────┤ │ data 一维,shape 定形,stride 定看法 │ │ │ │ data [1 2 3 4 5 6] │ │ shape (2, 3) │ │ stride (3, 1) │ │ │ │ 拿 (1, 2) 算一笔 │ │ = 1*3 + 2*1 = 5 │ │ │ │ 转置不搬数据 │ │ stride 换成 (1, 3),看法就变了 │ └────────────────────────────────────────┘stride 这个设计是框架省钱的关键。转置一个矩阵,不用搬数据,把 stride 里的 3 和 1 换个位置,看法变了,数据原地不动。切一行出来,也不用搬,记个起始偏移就行。几百万个元素的矩阵,换个看法零开销。第十章 struct 拼数据,第十三章模板让 Tensor 装任意类型,第十四章 vector 管内存,前十四章的招在这一个 struct 里会师。
4.3 计算图、把公式串成链
框架眼里,一个模型不是一段代码,是一张图。每个节点是一次运算或一份数据,节点之间的连线是数据流向。第十五章感知器那个式子,z 等于 w1 乘 x1 加 w2 乘 x2 加 b,画成图就一串节点。
#include <iostream> #include <vector> using namespace std; struct Node { double value; vector<Node *> inputs; char op; Node(double v) : value(v), op('c') {} Node(char o, Node *a, Node *b) : op(o) { inputs.push_back(a); inputs.push_back(b); } }; double eval(Node *n) { if (n->op == 'c') { return n->value; } double x = eval(n->inputs[0]); double y = eval(n->inputs[1]); if (n->op == '+') { n->value = x + y; } else { n->value = x * y; } return n->value; } int main() { Node x1(0.55), x2(0.61); Node w1(-0.095), w2(-0.076), b(0.1); Node m1('*', &x1, &w1); Node m2('*', &x2, &w2); Node s1('+', &m1, &m2); Node z('+', &s1, &b); double r = eval(&z); cout << "王五的 z = " << r << endl; cout << (r > 0 ? "判挂科" : "判不挂") << endl; return 0; }输出是王五的 z = 0.00139,判挂科。权重用的就是第十五章学出来那套,(-0.095, -0.076) 配偏置 0.1,王五的作业 0.55 月考 0.61 喂进去,结论和上一章手算的一样。只是这回公式不是写死的一行,是节点串出来的图。
┌──────────────────────────────────┐ │ 公式串成一张图 │ ├──────────────────────────────────┤ │ z = w1*x1 + w2*x2 + b 的图模样 │ │ │ │ x1 ─┐ x2 ─┐ │ │ [*] m1 [*] m2 │ │ w1 ─┘ w2 ─┘ │ │ └──>[+] s1 <──┘ │ │ │ │ │ v b ──┐ │ │ [+] z <──┘ │ │ │ │ │ v │ │ 输出落在挂科一侧 │ └──────────────────────────────────┘eval 是个递归,从输出节点往回要值,要不到的再往输入追,追到常量节点为止。框架里的前向传播就是这个递归的工业版,反向传播也走这张图的边,一条链把每个节点对结果的影响算出来。图的好处是可以优化,相邻的加法合并,没用的节点剪掉,相同的子图只算一次,这些优化全在 C++ 层做,Python 壳子碰不着。
4.4 框架剩下的活
一个真框架比我们这个玩具多三样大活。
算子库是其中之一。卷积、池化、注意力,每种运算一份高度优化的 C++ 实现,CPU 版一份,GPU 版一份,第 3 节的 CUDA 就藏在这里。再有自动微分,沿着计算图自动求导,训练不用手推公式。还有内存调度,大张量来来去去,谁复用谁的内存,哪块先释放,都是 C++ 层的精细活,第二章聊的内存分布在这里变成了真金白银。
| 框架的活 | 对应的 C++ 章节 | 本章位置 |
|---|---|---|
| 张量 struct | 第十章、第十三章、第十四章 | 4.2 节 |
| 内存调度 | 第二章、第六章 | 1.2 节 |
| GPU 算子 | CUDA 扩展 | 3.2 节 |
| 多态算子注册 | 第十二章 | 4.1 节分层 |
本章小结
本章把 C++ 放进深度学习的版图。语言层面,编译到底没有中间商,内存布局自己说了算,零开销抽象账本公开。应用层面,训练要吞吐,推理要延迟,挑剔的推理场景全是 C++ 的地盘。GPU 层面,CUDA 是 C++ 的方言,核多是本事,搬家是成本,数据留在显存反复用才划算。框架层面,Python 是壳,C++ 是芯,张量靠 stride 变看法,模型靠计算图串公式。
附录 A 速查卡
| 条目 | 一句话 |
|---|---|
| 编译型 | 源码直译机器码,运行期没中间商 |
| 行主序 | 二维数组按行摆,同行元素内存里挨着 |
| 缓存命中 | 要的数据在缓存里,快一百倍 |
| 零开销抽象 | 用不上的特性不付钱 |
| 训练 | 调权重,重吞吐,跑在机房 |
| 推理 | 用权重,重延迟,跑在设备上 |
| SIMT | 同一条指令,千个线程各算各的 |
| CUDA | C++ 的扩展,global加三角括号 |
| 显存 | GPU 自己的仓库,数据要搬过去才算得着 |
| 张量 | 一维数据加 shape 加 stride |
| stride | 每个维度走一步跳几个元素 |
| 计算图 | 节点是运算,连线是数据流向 |
附录 B 排错表
| 常见现象 | 病根 | 药方 |
|---|---|---|
| 矩阵乘慢得离谱 | 最内层循环跳着读内存 | 换循环顺序,让内层顺着行读 |
| 局部大数组一跑就崩 | 512 乘 512 的 double 数组栈上放不下 | static 或堆上申请 |
| 计时结果忽大忽小 | 没开优化,编译器各凭心情 | 加 -O2 再比 |
| GPU 程序比 CPU 还慢 | 数据在内存显存间来回搬 | 数据搬进显存就别来回,算完再搬 |
| kernel 跑了结果没变 | 没等 GPU 干完就读结果 | cudaDeviceSynchronize 等它 |
| 线程越界读垃圾 | 线程数大于数据量 | kernel 里加 if 判断 i 小于 n |
| 张量转置后数据乱了 | 以为转置要搬数据,手动瞎改 | 只换 stride,数据别动 |
| 递归 eval 栈溢出 | 图太深,递归层数超栈 | 改迭代或调大栈,工业框架都这么干 |