第十六章 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::durationdouble(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::durationdouble(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; } vecAdd4, 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 像 freenew 和 delete 的老规矩借了就还。真正新的只有三样global标记这个函数在 GPU 上跑三角括号 4, 256 说开 4 个线程块每块 256 个线程threadIdx 和 blockIdx 是每个线程问自己是谁。┌──────────────────────────────────────┐ │ 一千零二十四个线程一起上 │ ├──────────────────────────────────────┤ │ vecAdd4, 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 { vectorfloat data; vectorint shape; vectorint 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; vectorNode * 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同一条指令千个线程各算各的CUDAC 的扩展global加三角括号显存GPU 自己的仓库数据要搬过去才算得着张量一维数据加 shape 加 stridestride每个维度走一步跳几个元素计算图节点是运算连线是数据流向附录 B 排错表常见现象病根药方矩阵乘慢得离谱最内层循环跳着读内存换循环顺序让内层顺着行读局部大数组一跑就崩512 乘 512 的 double 数组栈上放不下static 或堆上申请计时结果忽大忽小没开优化编译器各凭心情加 -O2 再比GPU 程序比 CPU 还慢数据在内存显存间来回搬数据搬进显存就别来回算完再搬kernel 跑了结果没变没等 GPU 干完就读结果cudaDeviceSynchronize 等它线程越界读垃圾线程数大于数据量kernel 里加 if 判断 i 小于 n张量转置后数据乱了以为转置要搬数据手动瞎改只换 stride数据别动递归 eval 栈溢出图太深递归层数超栈改迭代或调大栈工业框架都这么干