很多项目里我都见过同一个场面有人把一段三层循环挪到 GPU 上跑立刻收获几十倍的速度提升但你问它为什么快答案往往只有一句因为 GPU 核心多。这不算错可也太粗糙了。真要把并行计算这个东西讲透找个具体例子比空谈理论有用得多而矩阵乘法就是最合适的那把钥匙。这篇内容会从矩阵乘法出发把 GPU 单挑 CPU 这件事拆开看两者的架构到底差在哪一个串行循环怎么变成成千上万个并行任务以及为什么有些时候 GPU 不但没快多少、反而更慢。适合刚接触高性能计算、想搞明白并行编程到底在做什么的开发者也适合那些已经会用某个 GPU 接口但一直没想通底层逻辑的人。1. 为什么偏选矩阵乘法来约架1.1 先认清朴素实现长什么样矩阵乘法的数学定义很简单C[i][j] 等于 A 的第 i 行与 B 的第 j 列做点积。落到代码上绝大多数程序员第一反应就是三层 for 循环// C语言风格朴素实现 void matmul_naive(float* A, float* B, float* C, int N) { for (int i 0; i N; i) { for (int j 0; j N; j) { float sum 0.0f; for (int k 0; k N; k) { sum A[i * N k] * B[k * N j]; } C[i * N j] sum; } } }这段代码的计算量是 N^3 次乘加也就是 2N^3 次浮点运算。N 取 1024 时总共要做大约 21 亿次乘加浮点运算。单线程 CPU 每秒能做的浮点运算通常以亿为单位计算完这样一个矩阵乐观估计也要一秒以上N 到 2048 时计算量直接翻八倍单线程能把你等得失去耐心。为什么偏偏拿它开刀因为矩阵乘法是计算密集的典型代表内存里一圈数据反复读写利用了很多次算术运算占比极高。这种形态的大型计算任务最能展现并行加速的价值也最能暴露并行优化的难点。1.2 它天生就是为并行而生的矩阵乘法最迷人的地方在于C 矩阵里的每个元素在数学上完全独立。计算 C[i][j] 只需要 A 的第 i 行和 B 的第 j 列跟其它任何一个 C 元素没有任何依赖关系。你可以把 C 想象成一张 N×N 的答题卡每个格子是一道独立的题目考生之间不许互相商量题目与题目之间也没有先后顺序。那么最简单的并行策略就出现了找几千个计算工人每人分几个格子各算各的最后拼成整张答题卡。这跟很多串行算法有本质区别。比如计算斐波那契数列的每一项依赖前两项拆并行还需要动很多脑筋又比如求解线性方程组时消元步骤天然有先后顺序。而矩阵乘法所有格子之间零通信、零依赖是教科书级的尴尬并行问题。1.3 它的江湖地位比想象中高得多矩阵乘法不只是考试题。现代深度学习里的全连接层、卷积层底层都能被改写成大规模矩阵乘法图像处理和信号处理里的各种变换很多也依赖矩阵乘法物理仿真、流体力学、推荐系统里更是处处是它的影子。你甚至可以说矩阵乘法是现阶段计算机世界里最硬核的基础算子。把这一块啃明白后面理解深度学习训练为什么吃 GPU、图像处理怎么加速、科学计算怎么写高性能代码都会顺很多。2. CPU 和 GPU 不是跑得快与慢的区别2.1 CPU少数精英单任务开荒高手CPU 的设计哲学是把单个任务做到极致快。一颗现代 CPU 核心里有非常复杂的控制单元、乱序执行逻辑、分支预测器、好几层容量不小的缓存。这些东西都是为了一个目的把一条指令流尽可能快地执行完。你可以把 CPU 内核想象成一个带着全套工具箱的专家。碰到一个复杂任务专家能根据经验判断下一步做什么遇到分支能提前猜路径甚至能同时准备好几件事的操作顺序。但专家的培养成本高一个芯片上放不下太多这样的专家——通常就是几个到几十个核心。所以在 CPU 上做并行本质上是你只有一小组高能力专家让他们各自负责一块区域。提升很可观但天花板取决于你的专家数量。2.2 GPU大量工人靠规模碾压GPU 的设计哲学完全不同。它的核心目标是把吞吐量做上去比如每秒处理多少像素、多少矩阵元素。GPU 芯片上放着上千个小计算单元每个单元的控制逻辑极简、缓存极小几乎不会做分支预测和乱序执行这些聪明事。有人喜欢把 GPU 比作一个大型车间里面几百上千个只会做基础计算的操作工每人拿着一个简易计算器你只需要把任务拆得非常碎、分给每一个人让他们同步干活。单个工人的效率远不如一个 CPU 专家但架不住人多整体吞吐量可以把 CPU 甩开一两个数量级。所以 CPU 和 GPU 的差别不是一个快一个慢而是两种完全不同的取舍CPU 追求把一个复杂任务尽快做完GPU 追求在一秒内把海量简单任务全部做完。2.3 架构取向决定了你看什么指标这对实际开发有一个很直接的影响评估 CPU 程序性能你习惯看主频、单核跑分、缓存命中率评估 GPU 程序性能你更应该看的是计算单元数量、显存带宽、吞吐量和占用率。维度CPU 取向GPU 取向核心目标降低延迟提高吞吐核心数量少而强多而简控制逻辑复杂支持乱序、分支预测简单依赖大规模并行掩盖延迟缓存容量大层级深容量小但配合大带宽典型场景操作系统、串行业务逻辑图像渲染、深度学习、科学计算这也是为什么 GPU 跑矩阵乘法这种计算密集任务有天然优势计算单元多内存带宽大每个单元干的活又都是类似的算术操作。3. 把三重循环翻译成并行任务3.1 粒度选择每个元素一个任务还是一行一个任务有了可并行这个大前提下一步是决定拆成多少块、每块多大。初学者最容易陷入的误区是觉得拆得越细越好其实粒度选择要看任务数和任务本身的开销。一个线程负责计算一个 C[i][j]任务总数是 N^2。这个粒度最细、负载最均衡而且方便映射到 GPU 的二维线程坐标。代价是每个线程都要自己遍历一整行 A 和一整列 B读入的数据有大量重复访存效率不高。一个线程负责计算一整行 C[i]任务总数只有 N。每个线程的循环依然串行并行度大幅下降。在 N 只有几百时GPU 上可能连计算单元都填不满根本体现不出并行的好处。入门阶段我建议直接采用一个线程算一个输出元素的方案理由很简单可直观、易实现、不容易出错。真正做性能优化时再进入分块思路让每个线程负责一个小块通过复用数据来减少访存那是第二步的事。3.2 线程坐标与矩阵坐标的映射GPU 生态里线程不是一维排列的而是以线程块 网格这种层次结构组织。你通常在三维索引里申请一定数量的线程块然后在每个块内部再组织成二维或三维的线程矩阵。对应到矩阵乘法最自然的是开一个二维线程网格// 以某主流GPU厂商C风格接口为例示意线程索引映射 __global__ void matmul_kernel(float* A, float* B, float* C, int N) { int row blockIdx.y * blockDim.y threadIdx.y; int col blockIdx.x * blockDim.x threadIdx.x; if (row N || col N) return; float sum 0.0f; for (int k 0; k N; k) { sum A[row * N k] * B[k * N col]; } C[row * N col] sum; }这里 blockIdx 表示线程块在网格中的位置threadIdx 表示线程在块内的位置blockDim 表示块的大小。三者的关系可以类比成一栋多层写字楼blockIdx 是楼层号threadIdx 是房间号blockDim 是每层有多少房间。整个网格覆盖了 N×N 个房间正好对应 C 矩阵的每个格子。边界判断那行if (row N || col N)是必须的因为 N 不一定能被块大小整除超出矩阵范围的线程要提前退出。3.3 行主序存储下的访存陷阱代码能跑通只是入门第一步真正决定性能的是访存方式。矩阵在内存里以行主序连续存放一行 N 个 float 紧紧挨着下一行紧接其后。看上面的 kernel 循环。内存访问分成两路读取A[row * N k]时k 变化地址永远是连续的一行数据顺序往后读非常友好读取B[k * N col]时col 不变而 k 在变每次要跳过一整行 N 个 float 才能读到下一个元素这种访问会让缓存几乎失效每次都要重新从显存拉数据速度直接掉一大截。为什么调用现成的优化库比自己写快那么多有很大一部分原因就是它们对这种访存模式做了分块和复用而朴素 GPU 实现只把串行代码机械地搬到了 GPU 上算术快了内存却成了新瓶颈。这个问题在第 5 章还会详细展开。4. 三种实现方式的实测对比4.1 单线程 CPU串行基准线我们先用单线程版本当基准。在通用 PC 环境下用编译器开满优化N 从 256 涨到 2048大约能得到这样一组示意数据不同硬件会有差异但数量级和趋势是可靠的矩阵规模 NCPU 单线程耗时2560.03 s5120.25 s10242.10 s204816.80 sN 从 512 涨到 1024计算量翻 8 倍时间也近似翻 8 倍这正是 O(N^3) 复杂度在单线程下的直白体现。看到 2048 要等 16 秒任何有点耐心的人都会开始琢磨有没有更快的方法。4.2 多线程 CPU简单的行切片就能提速CPU 多线程最直观的方案是把行切开0 到 N/4 行交给第一个线程N/4 到 N/2 行交给第二个线程以此类推。每个线程内部依然走原来的串行循环只不过各算各的行之间完全没有通信。还是同样的环境假设机器上有 8 个物理核心多线程版本的大致表现如下矩阵规模 NCPU 多线程耗时对比单线程2560.008 s约 4 倍5120.060 s约 4 倍10240.500 s约 4 倍20484.300 s约 4 倍8 核跑出约 4 倍加速并没有到 8 倍原因很常见内存带宽有限、超线程收益没想象中高、编译器自动向量化也有开销。但不管怎样N2048 可以从 16.8 秒压到 4 秒左右体验已经好了很多。4.3 GPU kernel小矩阵无优势大矩阵通杀GPU 版本使用前面的二维线程映射 kernel线程块大小设为 16×16网格按需展开。同样的矩阵规模加上必要的内存拷贝时间之后示意数据如下矩阵规模 NGPU kernel 耗时GPU 相对 CPU 单线程加速2560.015 s约 2 倍5120.020 s约 12 倍10240.060 s约 35 倍20480.350 s约 48 倍N256 时 GPU 只比 CPU 快 2 倍甚至可能不如优化很好的 CPU 多线程N2048 时GPU 已经把单线程 CPU 甩开近 50 倍。加速比会随矩阵规模继续拉大前提是你内核写得别太离谱。4.4 为什么小矩阵完全体现不出 GPU 优势很多人第一次跑 GPU 程序会沮丧明明号称几千个核心算个 128×128 矩阵居然比 CPU 还慢是不是哪里搞错了没搞错。GPU 不是一个拿到即用的独立设备它和 CPU 之间有数据通道。每一步并行计算都伴随着固定开销把矩阵从内存拷到显存、启动内核、等所有线程同步完成、再把结果拷回来。矩阵只有 64KB 时计算本身一瞬间就结束了但固定开销可能就要几毫秒总耗时自然难看。所以在实际项目里把数据留在 GPU 上反复做多步运算远比算一步拷一次划算。这个降低固定开销占比的思路是 GPU 优化的第一课。5. 性能解读让 GPU 快不起来的四个陷阱5.1 数据搬运和启动开销才是真正的隐形杀手假设 N2048单个 float 矩阵占 16MBA、B、C三个矩阵加起来 48MB。常见的设备间传输带宽大约每秒几 GB 到几十 GB来回搬运 48MB 至少需要毫秒级时间而 GPU 上计算这组矩阵乘法本身可能只需要 0.35 秒的很小一部分毫秒数。你会发现搬运时间接近甚至超过计算时间是 GPU 编程里最反直觉的常态。正确态度是尽量让数据在显存里多活一会儿。把多个 kernel 串成一条流水线只在真正需要把结果送回主存时才做一次拷贝。很多初学者喜欢每一步都 printf 结果验证这种习惯在 GPU 开发里会严重拖慢性能。5.2 线程块配置直接影响占用率线程块的尺寸不是随便填的。块开太小比如 4×416 个线程会导致调度单元无法有效隐藏访存延迟块开太大比如 64×644096 个线程一块又可能超过硬件的线程数量上限反而被频繁调度卡住。常见的经验区间是块内放 128 到 512 个线程像 16×16256 就是入门阶段很稳的选择。网格大小可以粗略取ceil(N / 16)保证覆盖整个矩阵即可。我见过不少初学者把 block 拿到后就用默认值结果内核性能差了几倍还找不到原因。花两分钟手动调一下块大小有时变化大到能惊掉下巴这正是占用率在起作用块越是能把计算单元塞满越能藏住访存的等待时间。5.3 访存模式决定数据能不能被喂饱第 3.3 节埋了个伏笔朴素 GPU kernel 里读 B 矩阵的列时访存是低效的。这里解释底层逻辑。GPU 访存时会一次性搬运一段连续的数据然后把这批数据分给相邻的线程。理想情况是线程 0 读地址 X线程 1 读地址 X4线程 2 读地址 X8……这样硬件只需要一次事务就能把所有线程的数据拿齐也就是所谓的合并访问。回到我们的 kernel假设线程块内 col 变化得快那么线程 0、1、2 同时读B[k*N col]col 分别相差 1对应地址相差了一个 N 的长度。它们根本不是连续地址硬件被迫发起很多次小事务访存带宽被白白浪费。最简单的改造方案有三种思路。一是转置 B 矩阵让原来按列访问变成按行访问c 代码上改动很小性能立刻上一个台阶。二是用共享内存做分块让一个线程块先把 A 和 B 的小块搬到片内高速内存里再反复复用这是正规军做法。三是让每个线程一次负责计算一个小输出块而非单个元素数据复用率提高访存次数下降。这里不粘贴完整的分块代码了但记住矩阵乘法优化的半壁江山都在访存上。很多号称我用了 GPU 怎么还是很慢的问题最后查下来几乎全是访存没优化好而不是算力不够。5.4 分支分歧和负载不均并行不是全自动的GPU 指令执行有一个特性它在硬件层面把线程分组每次以组为单位执行同一条指令。如果组内出现 if-else 分支一半线程走 if、一半线程走 else那这一组线程就不得不先把 if 路径跑完、再去跑 else 路径。结果就是一半线程在等待流水线空转整体效率打折。矩阵乘法里的边界判断if (row N || col N) return;恰好会产生这种分歧。当 N 不能被块整除时处在矩阵边缘的线程块里一部分线程干活、另一部分直接退出。工作量不大时看不出问题N 很大时边缘块占全部块的比例很低影响有限但如果你把这种分支放在大循环内部性能会显著恶化。负载不均是另一个隐形问题。矩阵乘法划分天然均匀所有线程工作量几乎一样但很多真实任务不是这样比如稀疏矩阵、不规则图计算某些线程要做大量计算另一些线程瞬间跑完最终耗时由最慢的那个线程决定。这也是为什么并不是所有算法都适合简单粗暴地扔给 GPU。6. 从矩阵乘法迁移到真实业务并行思维怎么通用6.1 拿到一个任务先问三个问题学完矩阵乘法你掌握的其实不是这段代码而是一套判断框架。任何一个计算任务到了手里先依次自问可并行度任务能不能切成互相独立的子任务切完后需要多少通信和同步计算强度每个子任务里算术运算多还是访存和搬运多如果计算强度低GPU 优势会被内存瓶颈吞掉。数据规模任务总量够不够大如果总计算量连固定开销都盖不住并行是负优化。用这套框架去套实际场景会有很直观的结论图像滤波这种每个像素点独立、计算量均匀的任务非常适合 GPU迭代式算法或者需要频繁全局同步的任务就要看同步开销能否接受把一个小数组倒序排个序老老实实用 CPU 就好。6.2 实际业务里哪里最容易看到矩阵乘法的影子深度学习训练应该是最典型的一个模型里全连接层对输入做线性变换本质是矩阵乘法卷积运算也能通过 im2col 之类方式转成矩阵乘法再调用优化内核这也是为什么 GPU 在 AI 时代翻红。图像处理中很多操作比如卷积模糊、缩放、色彩空间转换也天然是逐像素独立任务和矩阵乘法的并行结构几乎一样。物理仿真里的格子方法、电磁场计算、流体求解里的大规模线性代数运算同样离不开矩阵乘法这个底层操作。还有一个反常识的应用方向是数据库。很多大规模数据筛选和连接运算可以改写成批量矩阵运算的形式在 GPU 上做粗筛再用 CPU 做精细处理两者各取所长。6.3 从入门到进阶的一条实在路径如果你看完这篇准备继续往深走我给一条自己验证过比较顺的路线第一步把朴素 CPU 矩阵乘法写好确认 O(N^3) 规模感第二步用线程库做 CPU 多线程版本感受并行度带来的提升第三步写一个最朴素的 GPU kernel哪怕性能一般关键是建立线程坐标映射的概念第四步做共享内存分块优化这一大步走完你对访存优化的理解会超过大多数人最后再去看现代优化库和深度学习框架怎么调度、怎么自动调参这个时候很多内核优化技巧就不会再看不懂了。我自己刚接触 GPU 时犯过一个经典错误一看现成库算矩阵又快又省事就再也不肯碰底层实现结果讲到访存、共用内存的时候全是凭感觉。后来花了整整一个周末把手写朴素 kernel 改成共享内存分块版本测出的加速比变化让很多原理一下子扎进脑子里。矩阵乘法这东西捏着鼻子手写一遍比刷十篇论文都管用。最后再分享一个小习惯做并行优化的时候永远准备一个串行正确版本当对照。GPU 代码一旦出错调试信息又少又难读但有对照版本在你就能逐步二分定位问题。这个项目里我靠这个习惯省了不知道多少时间。矩阵乘法只是起点但把它的为什么快、为什么慢、瓶颈在哪想明白后面遇到再复杂的并行问题你都不会再一头雾水了。