很多项目里我都见过同一个场面：有人把一段三层循环挪到 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 程序性能，你更应该看的是计算单元数量、显存带宽、吞吐量和占用率。

这也是为什么 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，大约能得到这样一组示意数据（不同硬件会有差异，但数量级和趋势是可靠的）：

N 从 512 涨到 1024，计算量翻 8 倍，时间也近似翻 8 倍，这正是 O(N^3) 复杂度在单线程下的直白体现。看到 2048 要等 16 秒，任何有点耐心的人都会开始琢磨"有没有更快的方法"。

4.2 多线程 CPU：简单的行切片就能提速

CPU 多线程最直观的方案是把行切开：0 到 N/4 行交给第一个线程，N/4 到 N/2 行交给第二个线程，以此类推。每个线程内部依然走原来的串行循环，只不过各算各的行，之间完全没有通信。

还是同样的环境，假设机器上有 8 个物理核心，多线程版本的大致表现如下：

8 核跑出约 4 倍加速，并没有到 8 倍，原因很常见：内存带宽有限、超线程收益没想象中高、编译器自动向量化也有开销。但不管怎样，N=2048 可以从 16.8 秒压到 4 秒左右，体验已经好了很多。

4.3 GPU kernel：小矩阵无优势，大矩阵通杀

GPU 版本使用前面的二维线程映射 kernel，线程块大小设为 16×16，网格按需展开。同样的矩阵规模，加上必要的内存拷贝时间之后，示意数据如下：

N=256 时 GPU 只比 CPU 快 2 倍，甚至可能不如优化很好的 CPU 多线程；N=2048 时，GPU 已经把单线程 CPU 甩开近 50 倍。加速比会随矩阵规模继续拉大，前提是你内核写得别太离谱。

4.4 为什么小矩阵完全体现不出 GPU 优势

很多人第一次跑 GPU 程序会沮丧：明明号称几千个核心，算个 128×128 矩阵居然比 CPU 还慢，是不是哪里搞错了？

没搞错。GPU 不是一个"拿到即用"的独立设备，它和 CPU 之间有数据通道。每一步并行计算都伴随着固定开销：把矩阵从内存拷到显存、启动内核、等所有线程同步完成、再把结果拷回来。矩阵只有 64KB 时，计算本身一瞬间就结束了，但固定开销可能就要几毫秒，总耗时自然难看。

所以在实际项目里，把数据留在 GPU 上反复做多步运算，远比"算一步拷一次"划算。这个"降低固定开销占比"的思路，是 GPU 优化的第一课。

5. 性能解读：让 GPU "快不起来"的四个陷阱

5.1 数据搬运和启动开销才是真正的隐形杀手

假设 N=2048，单个 float 矩阵占 16MB，A、B、C三个矩阵加起来 48MB。常见的设备间传输带宽大约每秒几 GB 到几十 GB，来回搬运 48MB 至少需要毫秒级时间；而 GPU 上计算这组矩阵乘法本身可能只需要 0.35 秒的很小一部分毫秒数。你会发现"搬运时间接近甚至超过计算时间"是 GPU 编程里最反直觉的常态。

正确态度是：尽量让数据在显存里多活一会儿。把多个 kernel 串成一条流水线，只在真正需要把结果送回主存时才做一次拷贝。很多初学者喜欢每一步都 printf 结果验证，这种习惯在 GPU 开发里会严重拖慢性能。

5.2 线程块配置直接影响占用率

线程块的尺寸不是随便填的。块开太小，比如 4×4=16 个线程，会导致调度单元无法有效隐藏访存延迟；块开太大，比如 64×64=4096 个线程一块，又可能超过硬件的线程数量上限，反而被频繁调度卡住。

常见的经验区间是块内放 128 到 512 个线程，像 16×16=256 就是入门阶段很稳的选择。网格大小可以粗略取ceil(N / 16)，保证覆盖整个矩阵即可。

我见过不少初学者把 block 拿到后就用默认值，结果内核性能差了几倍还找不到原因。花两分钟手动调一下块大小，有时变化大到能惊掉下巴，这正是"占用率"在起作用：块越是能把计算单元塞满，越能藏住访存的等待时间。

5.3 访存模式决定数据能不能被"喂饱"

第 3.3 节埋了个伏笔：朴素 GPU kernel 里读 B 矩阵的列时，访存是低效的。这里解释底层逻辑。

GPU 访存时，会一次性搬运一段连续的数据，然后把这批数据分给相邻的线程。理想情况是：线程 0 读地址 X，线程 1 读地址 X+4，线程 2 读地址 X+8……这样硬件只需要一次事务就能把所有线程的数据拿齐，也就是所谓的"合并访问"。

回到我们的 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 代码一旦出错，调试信息又少又难读，但有对照版本在，你就能逐步二分定位问题。这个项目里我靠这个习惯省了不知道多少时间。矩阵乘法只是起点，但把它的"为什么快、为什么慢、瓶颈在哪"想明白，后面遇到再复杂的并行问题，你都不会再一头雾水了。