企业官网建设流程全解析

所以，当我们聊AI芯片的时候，到底在聊什么？

如果只看厂商发布会上的TOPS、TFLOPS这些数字，很容易陷入一种"参数拜物教"——好像算力堆上去，一切问题就都解决了。但做过实际AI项目的人心里都清楚，真正决定一个模型训练快不快、推理顺不顺的，恰恰是那些发布会PPT上很少展开讲的硬件底层机制。而在这些底层机制里，最核心、也最容易被当成"理所当然"的，就是GPU的执行单元。

这篇文章我想把"GPU执行单元"这件事彻底聊透。不仅告诉你它是什么，更想拆开给你看：它怎么工作、为什么这样设计、以及AI计算和它之间那种近乎天作之合的匹配关系到底建立在什么原理之上。

不管你是刚入门CUDA编程的开发者，是在做模型训练/推理优化的算法工程师，还是纯粹对AI芯片底层架构好奇的读者，这篇文章都会给你一个从"能用GPU"到"懂GPU"的认知台阶。

1. 为什么CPU搞不定的活，要专门造一个GPU来干

1.1 执行单元的起源：从"计算"到"并行计算"的岔路口

在讨论执行单元之前，得先把一个根子上的问题说清楚：CPU里也有计算单元，为什么AI计算最后还是选择了GPU？

回到计算机体系结构的基本盘。CPU的设计哲学是"低延迟控制流优先"，它要应对的是指令高度分支、数据依赖复杂、逻辑判断密集的通用任务。为了做到这一点，CPU把大量的晶体管预算花在了分支预测器、乱序执行引擎、大容量缓存这些"辅助设施"上。真正干活的ALU（算术逻辑单元）反而在芯片面积里占比不高。

GPU则走了另一条路。它的假设是：如果我面对的工作负载本身就包含海量的、互不依赖的、可以同时执行的计算，那我为什么要把晶体管浪费在预测分支和乱序调度上？不如把每一分资源都堆到计算单元上，让几千个计算单元同时开工。

这就是执行单元在GPU上成规模出现的原因——不是为了替代CPU，而是为了承接CPU不适合做的"重计算、轻控制"的工作。

1.2 执行单元到底是"一块"还是"一堆"

必须先把概念理清。当你听到"GPU执行单元"这个词时，要意识到它是一个层级概念。

最宏观的层面，整颗GPU芯片可以被视作一个"由海量执行单元组成的计算阵列"。往下一层，GPU内部被划分为若干Graphics Processing Cluster（GPC），每个GPC里有若干Streaming Multiprocessor（SM）——在AMD平台上叫Compute Unit（CU），在Intel的Xe架构里叫Execution Unit或Xe Core。这些SM/CU/Execution Unit才是真正意义上的"执行单元核心"。再往下一层，每个SM内部还有数十到上百个更细粒度的处理元件，NVIDIA叫它们CUDA Core（或者说流处理器），AMD叫Stream Processor，Intel叫Xe Vector Engine。

所以"执行单元"这个话题，本质上是一个纵向贯穿整个GPU硬件体系的大话题。不是说SM是执行单元而CUDA Core不是，而是它们各自处在"执行"这个链条的不同层级上，各自承担不同的职责。

1.3 为什么AI芯片领域如此看重执行单元

答案不那么玄妙：因为AI计算就是建立在"大量简单算术操作并行执行"之上的。无论是卷积、全连接、矩阵乘、Attention里的QKV投影，还是Transformer里每一步非线性激活，拆到最底层，就是成千上万亿次的乘法和加法。这些东西没有复杂的控制依赖，没有诡异的分支跳转，只有一点——量大、管饱。

举个例子，一个7B参数的LLM做一次前向推理，参数量大约70亿，哪怕只处理一个token，也要把这70亿个参数全部读出来参与计算。对应的乘加操作数轻松突破百亿级别。这种规模的纯算术负载，如果靠CPU那十几个核心慢慢跑，根本不可能支撑起如今的大模型应用。

GPU执行单元就是为了"以最大吞吐量完成海量无关算术操作"而存在的。理解了这一点，后面所有的架构细节都变得顺理成章。

2. 一块GPU芯片的计算骨架：从GPC到SM再到CUDA Core

2.1 NVIDIA Ampere架构下的芯片解剖

为了让讨论不那么飘在空中，我以目前生产环境中最常见的NVIDIA Ampere架构为例，把一颗典型的GPU芯片（比如GA102）做个自上而下的拆解。

一颗GA102芯片，总共有7个GPC。每个GPC内部有16个SM（部分GPC可能少一两个），所以整个芯片一共是84个SM左右。每个SM里有128个CUDA Core（FP32单元）、64个SFU（特殊函数单元）、4个Tensor Core（三代Tensor Core，每个Tensor Core内部实际上是一个四元组阵列）、4个Texture Unit、以及对应的LSU（载入/存储单元）、寄存器堆（Regster File）和共享内存（Shared Memory）。

这一层拆解下来，你会发现"执行单元"的含义已经分叉了：**如果你把SM看成执行单元，那它是"一组混合功能单元的组合体"；如果你把CUDA Core或Tensor Core看成执行单元，那它们是"能完成特定类型算术运算的最小硬件单元"。**这两种理解都对，关键看你讨论问题的粒度。

2.2 SM内部各功能单元的职责切分

在SM内部，各功能单元的分工是非常明确的。这里用表格来对比会更直观：

从这张表可以看出，GPU的SM不只是一群计算单元的简单堆叠，而是一个带有完整数据通路和存储体系的"微型处理器"。计算单元只是这个微型处理器里负责"算"的那部分，它能不能跑满，取决于数据能不能按时喂到嘴边——这就是为什么后面要专门讲访存和调度。

2.3 CUDA Core和Tensor Core的本质差异

很多初学者会把CUDA Core和Tensor Core混为一谈，这是执行单元话题里最容易出现的认知混淆。这两者的差异，说直白点就是：

CUDA Core是通用工具，什么都能干，但干矩阵乘这种大活效率一般。每个CUDA Core本质上是一个支持FMA（融合乘加）的浮点单元，一条指令完成d=a*b+c这个操作，吞吐是一拍一个。128个CUDA Core同时工作，一个SM一个时钟周期能做128次FMA。

Tensor Core是专用工具，只会干矩阵乘加，但干这个活的速度是CUDA Core的好几倍。它是专门为深度学习设计的硬件单元，执行的是D=A*B+C这种矩阵级别的操作。A、B、C、D都是矩阵，不是标量。以Ampere代的三代Tensor Core为例，一个Tensor Core一个时钟周期能完成一个4×4×4的矩阵乘加，等效于执行了64次FMA。所以一块SM里虽然只有4个Tensor Core，但它们加起来的矩阵计算吞吐远超那128个CUDA Core。

打个比方：CUDA Core是能做各种精细活计的工匠，Tensor Core是只做标准门窗的流水线工厂。做特殊造型（各种非矩阵计算），得靠工匠；批量盖楼（矩阵运算），工厂的效率碾压工匠。

3. 指令在SM里到底是怎么"跑"起来的

3.1 Warp：执行单元调度的时间切片

有了执行单元，下一步自然要知道指令怎么在这些单元上跑。NVIDIA的做法是把32个线程绑定为一个Warp，以Warp为单位进行调度和执行。

这是什么概念？就是你写CUDA代码的时候，虽然你定义的是成千上万个线程，但硬件执行时，SM里真正的最小调度单位是Warp。一个Warp里的32个线程执行同一条指令，只是各自处理不同的数据。这就是SIMT（单指令多线程）模型。

从SM内部看，一个SM通常有4个Instruction Scheduler（指令调度器），每个调度器每个时钟周期可以向相应的功能单元发出一条指令。但注意，这条指令是"一个Warp的指令"，也就是说，一条指令发出后，对应的功能单元要为这个Warp里的32个线程同时干活。

麻烦就出在这里。一个Warp有32个线程要做FMA，但每个SM只有128个CUDA Core。调度器一条指令要让32个线程同时执行FMA的话，就需要32个CUDA Core。那其余96个Core干什么？答案是：另外还有三个调度器在同时发出指令，各自服务一个Warp。四个调度器一拍发出四条指令，每条指令占32个线程的资源，总共占满128个CUDA Core。

所以，**一个SM能同时处于执行状态的Warp数量，理论上受限于它有多少个可用的功能单元。**这就是"吞吐"的本质。

3.2 延迟隐藏：为什么GPU不怕"慢"

CPU应对访存延迟的方式是缓存命中、乱序执行、分支预测，核心思路是"别让流水线停下来"。GPU的方法则完全是另一个次元：打不过就用车轮战。

GPU有一个术语叫Occupancy（占用率），指的是SM上同时活跃的Warp数量。为什么GPU要承载这么多Warp？就是为了延迟隐藏。当一个Warp的指令需要从全局内存取数据，可能需要几百个时钟周期才能拿到结果。这段时间里，如果调度器等这个Warp，那执行单元就闲置了，性能灾难。但如果SM上有足够多的Warp，调度器就可以立刻切到另一个已经准备好数据的Warp，让执行单元永远有活干。

这就是为什么你在编程时要关注"有没有足够多的线程/Block来填满SM"。线程太少，执行单元就会因为等待Data而挨饿；线程太多，寄存器不够用又会导致本地内存溢出，反而拖慢速度。找到这个平衡点，是每个CUDA性能优化工程师的基本功。

3.3 分支分化：执行单元最怕的程序结构

SIMT的执行方式带来一个经典问题：分支发散（或叫分支分化）。

想象一下一个Warp里的32个线程执行if (data[i] > 0) { do_a(); } else { do_b(); }。如果这32个线程的数据有的大于0有的小于0，那么执行单元没法同时执行do_a和do_b——它只能先让走true分支的线程执行do_a，此时走false分支的线程只能在一旁干等（被掩码屏蔽）；然后再反过来执行do_b。这个Warp一共执行了do_a和do_b两段代码，耗时翻倍，但有效计算量只有一半。

这直接导致了GPU执行单元的利用率下降。所以我们在写kernel时要想尽办法避免Warp内分支：可以把数据预先排序、把不同分支拆成独立的kernel、或者用算术方法替代分支。这些都是实际项目中能实打实省出几十毫秒的优化手段。

3.4 访存合并：让LSU吃饱的秘诀

除了分支，GPU执行单元最在意的事情就是访存模式。

LSU（载入/存储单元）负责把数据从显存搬进寄存器，或者把运算结果写回显存。当一个Warp的32个线程访问全局内存时，如果它们的地址是连续的（比如访问数组的相邻元素），硬件可以把这个访问合并成少数几次大的内存事务——这种模式叫合并访存（coalesced access）。反过来，如果32个线程各自访问毫无规律的地址，那内存控制器就要拆成几十次小事务，效率一落千丈。

这个特性和执行单元的关系太直接了：**访存效率低，再快的执行单元也得空转等数据。**很多人在写kernel时觉得自己的计算逻辑没问题，跑出来却慢得像爬，一查profile，大概率是访存没合并。

4. AI计算如何与执行单元互相成就

4.1 一个GEMM在SM层面是怎么被切分的

深度学习里最核心的运算GEMM（通用矩阵乘），是理解执行单元价值的最佳窗口。

假设要计算C[M,N] = A[M,K] * B[K,N]，M、N、K都是几千的规模。GPU的做法不是把一个巨大的矩阵整体丢给执行单元，那样数据放不下，计算也组织不起来。实际的策略是分块（tiling）：

1. 在Block层面，把最终结果矩阵C划分成若干个Block Tile，比如每个Block负责一个128×128的C子块。
2. 在Warp层面，一个Block内部又由多个Warp各自负责其中一部分，比如每个Warp处理一个64×64的子块。
3. 在指令层面，每个Warp的线程调用mma.sync指令（Tensor Core指令），一次完成m16n8k16之类的矩阵分块乘加。

这个分层切分策略，本质上就是为了让数据尽可能留在SM内部的共享内存/寄存器里，减少对全局内存的访问。每个Thread从全局内存中只读取自己负责的那一小块A和B数据，存入共享内存，然后反复复用——这正是"计算密度"的核心：执行单元在一个数据上做的计算次数越多，访存带宽压力就越小。

4.2 Tensor Core里的"乘法累加洪流"

Tensor Core执行GEMM的微观机制非常有意思。以A100的第三代Tensor Core为例，它执行的是D = A * B + C，其中A、B、C、D都是4×4的矩阵。也就是说，一次Tensor Core指令可以完成：

• 4行4列的A矩阵（16个FP16元素）
• 4行4列的B矩阵（16个FP16元素）
• 4行4列的C矩阵（16个FP32元素）
• 输出4行4列的D矩阵（16个FP32元素）

等效于64次FMA运算。而且Tensor Core支持在一个时钟周期内完成这个操作。对比一下，一个CUDA Core一个周期只能完成1次FMA。虽然Tensor Core占用的晶体管面积比CUDA Core大不少，但就矩阵运算而言，它的能效比是碾压级别的。

后续代际的Tensor Core规格继续膨胀：H100的第四代Tensor Core直接支持m16n8k4到m16n8k16等多种形状，还有FP8支持；到了H200和B200时代，Tensor Core的矩阵形状进一步扩大，单指令的FLOPs还在攀升。可以说，大模型的每一次规模跃迁，背后都是Tensor Core这个"专用执行单元"的算力密度升级。

4.3 训练和推理对执行单元的需求为什么不同

聊执行单元，必须区分训练和推理两种场景，因为它们对硬件的压力点完全不一样。

训练阶段的特点是海量计算、数据精度要求高、需要反向传播求梯度。此时执行单元的压力是"持续满载"，GPU需要以最高吞吐执行前向和反向的GEMM，而且通常使用FP16/BF16精度配合混合精度训练，因为梯度更新和损失计算需要FP32精度。训练时，执行单元最容易被算力瓶颈卡住，因为计算量实在太大，每个时钟周期都要最大程度地压榨Tensor Core。

推理阶段则复杂很多。在线推理通常有延迟要求，Batch Size往往很小（甚至为1）。此时执行单元可能吃不饱，因为单次请求的计算量太小，内存带宽反而成了瓶颈。这就是为什么很多推理优化工作会做Continuous Batching（连续批处理）、PagedAttention，本质上都是想办法把不同的请求拼成更大的矩阵，让执行单元不至于因为"活太少"而空转。另外，推理还常常用量化（INT8/FP8），把执行单元的单位算力进一步转化成吞吐优势。

4.4 Roofline模型：判断程序瓶颈的一把尺

理解了执行单元和访存系统，你就有了一把分析程序瓶颈的尺子：Roofline模型。

Roofline的核心思想是：一个kernel能达到的实际性能，不大于"算力上限"和"带宽上限"两者中的较小者。如果一个程序的计算密度（每字节访存对应的计算量）很低，那么它大概率是访存瓶颈；反之，如果计算密度很高，但执行单元已经被压满，那就是算力瓶颈。

这张"屋顶"图像是在如何在GEMM优化、卷积算子开发时做决策的关键依据：如果一个op是访存瓶颈，那你去调Tensor Core的指令形状一点用都没有，应该去优化共享内存的复用、改善访存合并；如果一个op是算力瓶颈，那就要考虑能否用Tensor Core替代CUDA Core、能否用更低精度的数据格式增加单指令的FLOPs。

5. 谁在"喂"执行单元：存储与调度体系

5.1 寄存器堆、共享内存与全局内存的分工

执行单元本身跑得快，但数据链路跟不上，一切白搭。GPU里存在一个金字塔形的存储体系：

• 寄存器堆：每个SM里有256KB的寄存器堆（Ampere），分割给所有活跃线程。这是速度最快的存储，执行单元做任何计算，数据基本都要先在寄存器里。但寄存器是隐私的，一个线程的寄存器，其它线程访问不了。
• 共享内存：每个SM有最多164KB（可配置）的共享内存，整个Block的所有线程都能访问。这是执行单元之间协作的最佳场所，也是tiling策略发挥作用的地方。共享内存速度远快于全局内存，相当于一个用软件管理的"二级缓存"。
• 全局内存：就是显存（HBM/GDDR），所有Block都能访问，容量大但延迟高。执行单元从全局内存取数据是最贵的操作，要尽量避免。

理解了这套体系，你就知道为什么优化CUDA kernel的核心思路永远是：别让执行单元频繁去全局内存取数，把数据尽量往共享内存和寄存器里搬。

5.2 调度器如何决定"下一个执行谁"

前面提过SM里有多个Instruction Scheduler，那调度器到底按什么规则选Warp？

NVIDIA的Warp调度策略在不同代际略有差异，大体上是一个循环式的优先级调度——调度器维护一个Warp列表，每次选一个"准备好了的（即依赖数据已到齐、且没有被阻塞）"Warp发出下一条指令。这个选择过程每个时钟周期都在进行，目的是最大化功能单元的利用率。

有一个值得了解的概念是两个Warp的"发射间隔"（issue interval）。如果两条指令之间没有数据依赖，那它们可以流水线式地连续发射；如果有依赖，就必须等前一条指令的结果写回。这要求开发者对指令级并行有意识，在写代码时尽量让计算链不要过长。当然，对绝大多数应用开发者来说，这个问题已经被编译器处理掉了大半，但知道底层逻辑还是有帮助的。

5.3 为什么说"SM的占用率不是越高越好"

这是个反直觉但极其重要的优化细节。

Occupancy高，意味着SM上有更多Warp可以切换，延迟隐藏效果更好。但代价是每个Warp分到的资源变少——尤其是寄存器。如果你的kernel因为占用率太高一拍，每个线程只分到少量的寄存器，被迫把局部变量"溢出"到本地内存（其实落在L1/全局内存里），那性能反而会暴跌。

在实际调优中，我经常用__launch_bounds__来限制一个Block里的线程数，从而给每个线程争取更多寄存器，换来更高的单线程指令执行效率，有时性能反而比强行拉满Occupancy更好。这就是执行单元话题里最有意思的平衡术：你要的是"效率"而不是"占用率数字"。

6. 从CUDA代码到执行单元：一次性能优化的实战视角

6.1 一个矩阵乘kernel的优化路径

纸上谈兵够了，用一个实战视角来看执行单元是如何被优化的。假设我们要实现一个FP16 GEMM的CUTLASS式kernel，优化的路径大致是：

1. 朴素版本：每个线程直接算C中的一行一列，访存完全不合并，执行单元利用率低得可怜。
2. 共享内存分块：把A和B矩阵分块载入共享内存，每块被多个Warp复用，大幅减少全局内存访问。这个阶段执行单元开始能吃上"热饭"。
3. 寄存器优化：每个线程用寄存器数组手动缓存A、B的切片，减少重复读共享内存。
4. 向量化访存：用float4/half2这类向量类型一次读多个float，增加访存吞吐。
5. Tensor Core：把核心计算换成mma.sync指令，将执行单元从CUDA Core切换到Tensor Core，算力瞬间翻几倍。
6. Pipeline：用CUDA Graph或异步拷贝（cp.async）实现共享内存的prefetch，把计算和访存重叠起来，让执行单元永远有数据可算。

每一步的改进，本质上都是在调整"执行单元与数据之间的距离"。距离越短，性能越高。

6.2 用Profile工具看到执行单元的"饥饿"

理论说了一大堆，真正要定位性能问题，最终还是得靠工具。我推荐从NVIDIA Nsight Compute入手，因为它是可以直接观测SM内部执行单元状态的神器。

Nsight Compute里的几个关键指标直接对执行单元负责：

• SM Busy：SM内计算单元活跃占比。如果很低，说明执行单元在等数据，优先排查访存。
• Issue Active：调度器发出指令的活跃度。如果低，说明Warp处于阻塞状态（等待数据或同步）。
• Executed Ipc Active：平均每周期执行的指令数。接近4（Ampere SM有4个调度器）说明执行单元喂得很饱，远低于4则说明有阻塞。
• Memory Throughput：全局内存吞吐。如果这个指标先打满，那就别折腾执行单元了，它在等着吃饭。

我遇到过很多次"代码写得像优化过但实际还是慢"的kernel，一测Executed Ipc Active只有0.3。这意味着执行单元97%时间在发呆。这种问题你用多少算力纸面数据都救不回来，必须从数据流入手。

6.3 大模型时代的执行单元真相：算力规格之外更值得关注的指标

大模型火起来之后，各家芯片厂商的算力数字一个比一个吓人。但我要泼一盆冷水：在执行单元的话题里，峰值算力只是入场券，持续稳定的有效算力才是胜负手。

有效算力=峰值算力×实际利用率。实际利用率取决于你的算子库有没有针对具体硬件做深度优化（FlashAttention就是一个绝佳案例），取决于你的显存带宽能不能喂饱执行单元，取决于你的矩阵形状填充度够不够高（比如seq_len=3的短序列推理，Tensor Core根本跑不满）……

所以，做AI基础设施选型时，别只看TFLOPS，要看基于你真实模型benchmark出来的有效吞吐。这也是为什么H100、A100这种经过市场反复打磨的芯片在生产环境里地位稳固——不只是因为纸面强，是因为在真实负载下，它们的执行单元确实能被高比例地喂饱。

7. 执行单元设计的边界与未来：功耗墙、存储墙与不断扩张的AI需求

7.1 功耗墙：芯片能跑多快，其实由散热决定

GPU执行单元的发展，视觉上看起来是"核心越来越多、频率越来越高"，但物理上其实一直被功耗墙卡着。

制造工艺不断缩小，晶体管密度不断提高，但单位功耗的散热量并没有跟上传导速度的提升。一块700W功耗的GPU把几乎所有功耗都喂给了执行单元，然而这带来的热量极其恐怖。于是我们看到一个趋势：算力增长不再主要依靠堆核心数量，而是依靠每瓦性能的提升——例如引入MXFP4格式、用更小的数据位宽换更快速度、或者用更微妙的激活稀疏化。

执行单元的设计越来越受到能效比约束。谁能在同等功耗下跑出更高的有效算力，谁就在AI芯片竞争中占据优势。

7.2 存储墙：算得再快，数据喂不进来也是白搭

与功耗墙并行的还有存储墙。HBM的带宽这两年提升很快，但相比执行单元的算力提升还是慢了不少。

以H100为例，FP16稠密算力接近1000 TFLOPS，而HBM3的带宽只有3.35TB/s。你可以算一笔账：每个字节的数据，执行单元可以做约300次FP16运算。如果你的算法不注重数据复用，那么即便执行单元全速运转，也会被内存带宽卡死。这也是为什么AI推理和训练都要强调"计算密度"和"kernel融合"——本质都是为了减少对存储带宽的需求，让执行单元在"数据复用"上多做文章。

7.3 下一代执行单元往哪里走

展望未来，GPU执行单元的设计方向其实已经非常清晰：

• 更宽的Tensor Core指令：单指令覆盖更大的矩阵形状，减少指令发射开销。
• 更低精度的数据格式：FP8、MXFP4、甚至更低位的三值/二值量化，让执行单元每个周期能做更多次有效乘法。
• 更为异构的SM：CPU、GPU、张量单元、稀疏加速器组合共存，各司其职。英伟达的Grace Hopper以及更近的Blackwell架构，已经在走向"把整个计算系统的各个部分都做成专用执行器"的路线。
• chiplet化设计：把执行单元按区块拆分到不同小芯片上，再用高速互连连接，突破单芯片光罩面积的限制。

执行单元这个概念的边界，会从"GPU内部的硬件单元"逐渐扩展到"整个计算系统的执行能力"。但对从事AI开发的人而言，最核心的那句话依然没变：让执行单元吃饱、吃好、别乱跑。

我在实际项目中看到过太多类似的案例：一个模型推理慢得离谱，开发者第一反应是"GPU太弱了要换卡"，结果一分析发现是kernel写得稀烂，执行单元利用率不到20%。换卡之前先做一次Profile，通常能分分钟找出问题——无论是warp分化、访存不合并，还是Tensor Core压根没用上。这是我在这个领域反复体验到的真理：硬件在进步，但能否吃透这些执行单元的脾性，才是决定你是"NVIDIA显卡用户"还是"GPU性能优化工程师"的分水岭。

等你真正把一个kernel从"能跑"优化到"跑满"，看着Nsight Compute里Executed Ipc Active从0.3爬到3.7的时候，那种对整个计算体系的理解，比任何参数表上的数字都有说服力。