GPU执行单元深度拆解：从CUDA Core到Tensor Core的AI算力真相

如果你最近在折腾 AI 芯片、大模型推理或者 GPU 编程，肯定绕不开一个词：执行单元。很多人跑 llama.cpp、微调大模型、装 PyTorch 的 GPU 版本时，总会在 benchmark 里看到算力数字，但真正问你“GPU 执行单元到底是什么、它怎么把活干完的”，可能又答不上来。这篇文章不聊空泛的架构宣传语，直接从执行单元的内部构造、任务派发、调度占用、AI 负载优化到实际查看 GPU 利用率这整个链路，把执行单元拆开揉碎讲清楚。

这篇文章适合三类人：第一类，正在用 PyTorch、Ollama、llama.cpp 跑模型，想知道为什么 GPU 没比 CPU 快多少；第二类，做 GPU 服务器运维或性能压测，需要理解 nvidia-smi 里的数字到底代表什么；第三类，纯粹想搞懂 AI 芯片中“执行单元”这个核心概念和背后原理的人。读完你至少能回答：“为什么 A100 算力这么高但我的模型这么慢”“GPU Util 100% 是不是就代表执行单元全部跑满了”这类高频问题。

1. 为什么说执行单元才是 GPU 的灵魂

1.1 从 CPU 和 GPU 的分工说起

CPU 和 GPU 最本质的区别，不在于谁的主频高，而在于芯片面积的用途。CPU 里大量晶体管被用在了乱序执行、分支预测、缓存一致性、大容量缓存这些“控制面”上。这样设计是为了让单条指令尽量快地跑完，应对逻辑跳跃、条件分支极其复杂的通用程序。真正执行算术运算的 ALU/FPU 数量其实很少，一个现代 CPU 物理核每周期能执行的乘加运算也就几条到十几条。

GPU 走了完全相反的路线。它把几十亿晶体管中的大部分都做成了执行单元——大量精简的 FP32、INT32 单元、专用的矩阵运算单元。控制逻辑被压缩得很小，缓存也不大，但数量多到可怕。一块消费级 RTX 4090 有 128 个 SM（Streaming Multiprocessor），每个 SM 里有 128 个 FP32 CUDA Core；一块 A100 有 108 个 SM，每个 SM 里有 64 个 FP32 CUDA Core 和 4 个第三代 Tensor Core。所以整张 GPU 加起来有上万甚至十几万个简单执行单元，它们并行工作，靠数量堆出极高的吞吐量。

这就是为什么 AI 芯片和大模型训练非 GPU 不可。神经网络里的卷积、矩阵乘法、Attention 这类计算，本质上全部由乘加运算组成，而且数据局部性强、并行度高。把一个大矩阵乘法拆成几千个子块，每个子块交给一个 SM 上的执行单元去算，GPU 比 CPU 能快一两个数量级。这里的“加速”不是单指令变快了，而是同时干活的执行单元变多了几百倍。

1.2 “执行单元”到底指什么

执行单元（Execution Unit）这个词在不同厂商的文档里有不同叫法。NVIDIA 官方经常用“SM 中的处理核心”来描述；AMD 的 RDNA/CDNA 架构里有“CU（Compute Unit）”“Stream Processor”；Intel 的 Xe 架构里有“Xe Core”“Vector Engine”；如果延伸到专用 AI 芯片，比如昇腾、Habana、Google TPU，又有“AI Core”“MXU”“脉动阵列”等名称。但无论叫什么，它们指的都是真正接收指令、完成算术运算、回写结果的那部分电路。

一个执行单元最基本的能力是执行指令。一条机器指令从取指、译码之后，最终会进入执行单元里完成具体操作。所以 GPU 算力的上限，直接取决于执行单元的数量、工作频率，以及每个周期能完成多少次浮点或整型运算。比如 A100 官方标称 FP32 算力约 19.5 TFLOPS，TFLOPS 的计算公式粗略看就是“执行单元数量 × 频率 × 每周期 FLOPs”。如果执行单元被闲置或者没有数据可算，这些书面上的算力就只是纸面数字。

1.3 对 AI 工作负载有什么意义

AI 负载里最常见的算子就是深度神经网络中的卷积和全连接层，对应到底层是 GEMM（通用矩阵乘）。要跑 GPU 上的大模型推理，编译器或运行时会把 GEMM 切块，分配到各个执行单元上。这里的执行单元不仅包含普通 FP32 CUDA Core，更关键的是 Tensor Core 这类矩阵专用单元。Tensor Core 能在一条指令里完成一个 4×4×4 甚至更大的矩阵乘加操作，吞吐量是普通 CUDA Core 的好几倍。

所以当你去搜索“GPU 微调大模型”“llama.cpp 怎么跑 GPU”“yolov8 目标检测需不需要 GPU”时，真正决定体验的就是这些执行单元能否被正确调用。CPU 也能跑这些模型，但 CPU 的执行单元数量和专用矩阵单元规模有限，算力差了几个数量级。理解了这一点，你就知道为什么 GPU 驱动、CUDA 工具包、PyTorch GPU 版的安装如此重要——它们的任务，就是把你的上层算子顺利翻译成 GPU 执行单元能执行的指令。

2. 执行单元的组成：CUDA Core、Tensor Core、INT32、LD/ST 和 SFU

2.1 一个 SM 里住着哪些"工种"

如果细看一份 NVIDIA 架构图（比如 Ampere、Hopper、Ada Lovelace），你会发现每个 SM 内部并不是只有一种执行单元，而是包含多个不同类型的“工人”。它们各司其职，处理不同类型的指令。以典型的 Ampere A100 SM 为例，大致有这几种：

注意，这些数字不是固定的。图灵架构的 SM 是 64 个 FP32 + 64 个 INT32；到安培之后 FP32 和 INT32 是独立的，可以并行发射；而在消费级 Ada 架构里，FP32 单元数量有所增加，部分执行单元与 INT32 共享数据通路。但核心思想不变：不同类型的指令去不同的执行单元，能并行的类型就并行，不能并行的就按周期排队。

很多人会把 CUDA Core 直接等同于 CPU 的核，然后说“GPU 有几千个核”，这句话其实不太准确。CPU 的核是完整的独立计算核心，有完整的流水线和控制能力；而 CUDA Core 更像是 SM 内部的一个流水线部件，不是独立核心。你可以在编程模型层面把线程抽象成“每个线程跑在一个 CUDA Core 上”，但物理执行时一个 warp 的 32 个线程由多个执行单元协同完成。

2.2 FMA 与乘加指令为什么关键

AI 算力估算里最常用的指标是 FLOPS，即每秒浮点运算次数。而浮点运算里最划算的是 FMA（Fused Multiply-Add）——一条指令同时完成d = a * b + c。如果 CPU/GPU 支持 FMA，那么一次 FMA 会被计为两次 FLOPs（一次乘法、一次加法）。GPU 的 FP32 执行单元本质上是大量的 FMA 单元，一排一排地排列在 SM 内部。一个 FP32 CUDA Core 每周期能做一次 FMA，单精度算力就是 2 次 FLOPs。

拿 A100 来说，108 个 SM × 64 个 FP32 单元 × 1.41 GHz × 2（FMA 算两次）= 大约 19.5 TFLOPS。所以看到算力数字时，你可以反推出执行单元的大致频率和数量。而在 AI 模型里，卷积和全连接层的核心就是乘加操作，FMA 单元自然成了执行单元中的主力。

2.3 Tensor Core：专为 AI 打造的矩阵执行单元

普通 FP32 单元一次只能处理两个数相乘再加到一个数上。但在 GEMM 中，我们希望一次指令能处理一个矩阵块，比如 A 是 4×4 的 FP16，B 是 4×4 的 FP16，C 是 4×4 的 FP32 累加器。Tensor Core 就是在硬件层面把这种“小矩阵乘加”做成单条指令。A100 的第三代 Tensor Core 每个时钟周期可以执行 4×4×4 的矩阵乘加（不同代不同，Hopper 的 WGMMA 指令支持更大的 tile）。

Tensor Core 对 AI 特别重要的原因是精度与吞吐的 trade-off。在 FP16 输入下，A100 Tensor Core 的算力大约 156 TFLOPS，是 FP32 CUDA Core 算力的 8 倍左右；如果再用 Sparsity（结构化稀疏，只计算非零元素），可以进一步提升到 312 TFLOPS。到了 H100，FP16 Tensor Core 算力接近 990 TFLOPS。对比它的 FP32 算力只有约 67 TFLOPS，你就知道 Tensor Core 才是大模型训练推理时代的主角。

所以当你在 PyTorch 里设置torch.float16或bfloat16跑模型，不只是降低显存占用，更重要的是让算子落到底层 Tensor Core 上，才能吃到这颗 GPU 相对于 CPU 最大的算力优势。很多初学者跑 AI 模型时发现 GPU 利用率上不去，其中一个常见原因就是算子还在用 FP32 执行，没有走 Tensor Core 路径，纸面上的混合精度算力根本没被用到。

2.4 为什么还需要 INT32、LD/ST 和 SFU

很多人会疑惑：AI 计算不都是浮点吗？为什么 SM 里还要放 INT32、LD/ST、SFU？因为它们承担的是一些“辅助但不可或缺”的工作。

• INT32 单元：指针运算、数组索引、整数循环计数、掩码计算。内核里写int idx = blockIdx.x * blockDim.x + threadIdx.x;这类代码时，地址计算走的就是 INT32 单元。访存地址算得快，数据才能更早送到浮点执行单元。
• LD/ST 单元：负责把数据从全局内存、共享内存加载到寄存器，或者把结果写回。一个 warp 执行一条ld.global.b32指令时，会通过 LD/ST 单元去访问缓存/显存。没有这些单元，任何执行单元都会因为没有“原料”而停工。
• SFU：求倒数、平方根、sin/cos、log/exp。神经网络里的归一化、激活函数有时会用到。虽然大部分算子现在被融合进了别的实现，但 SFU 依然存在，某些特殊算子里是性能关键。

理解这些部件的分工，对性能分析很有用。如果你用 Nsight Compute 去 profile 一个内核，发现Long Scoreboard或Wait这类 stall 原因很高，说明 LD/ST 单元和访存路径是瓶颈；如果发现Math Pipe Throttle太高，就说明数学执行单元（FP32/Tensor Core）已经饱和。这两种情况对应的优化方向完全不同。

3. 从线程到执行单元：SIMT、Warp 和调度器如何把活分下去

3.1 网格/线程块/线程是怎么映射到执行单元的

GPU 编程模型里有个经典层级：Grid（网格） → Block（线程块） → Thread（线程）。当你启动一个 CUDA kernel，比如<<<gridDim, blockDim>>>，GPU 的硬件调度器并不会把每个独立线程逐个分配给某个执行单元，而是以线程块为单位把任务分发到不同的 SM。一个 SM 可以在同一时刻驻留多个线程块，具体能驻留多少，取决于线程块的大小、寄存器用量、共享内存用量，以及硬件上限。

线程块内的线程会被继续细分成 Warp。NVIDIA GPU 中一个 Warp 固定包含 32 个线程，这是硬件调度和指令发射的基本单位。为什么是 32？这是从 G80 时代延续下来的传统，与执行单元的宽度和掩码设计有关。AMD 的 wavefront 是 64 个线程，Intel 的 EU 中则可能有不同宽度，但都是类似思路。

3.2 SIMT 模式一条指令伺候 32 个线程

执行单元干活时遵循 SIMT（单指令多线程）模式：一个 Warp 里的 32 个线程在同一个周期执行同一条指令，只是各自使用的数据不同。举个例子：c[i] = a[i] + b[i]，如果 i 从 0 到 31，这 32 个线程会统一执行加法指令，但分别读取自己那份 a[i] 和 b[i]，把结果写入自己的 c[i]。这种模式大幅减少了指令取指和译码的开销——本来需要 32 条指令完成的工作，现在只需要一条指令加一堆数据通路。

不过，SIMT 也带来了一个著名问题：分支分歧（branch divergence）。如果一个 Warp 里的线程因为数据不同走到了 if 和 else 两个分支，硬件无法让这 32 个线程并行执行两条不同路径，只能先执行 if 路径（此时走 else 的线程被掩蔽），再执行 else 路径（此时走 if 的线程被掩蔽）。两条路径串行完成，整个 Warp 的执行时间就是两段之和，执行单元里总有一半席位在空转。AI 算子中这类分支相对少，但游戏中或通用计算中很常见。

3.3 调度器如何让执行单元保持忙碌

SM 内部有多个“调度器（Scheduler）”或“处理分区（Processing Block）”，每个分区有自己的执行单元集合、寄存器文件子集和指令缓存。以 A100 SM 为例，它被分成 4 个处理分区，每个分区包含 16 个 FP32 CUDA Core、16 个 INT32 单元、4 个 Tensor Core，以及一个 Warp Scheduler。

Warp Scheduler 每个时钟周期需要从一组可运行的 Warp 中选择一个发射指令。当一个 Warp 在执行一条长延迟指令（比如读取全局内存）时，这个 Warp 在结果返回之前无法继续执行后续依赖指令。如果只有一个 Warp，执行单元就得干等着。但 GPU 的做法是：每个 SM 至少驻留成千上万个线程（即几十个 Warp），当一个 Warp 等待访存时，调度器立刻切换到其他 Warp 发射指令，把执行单元的每一个周期都填满。

这就是 GPU “靠隐藏延迟换取吞吐”的思路。CPU 为了隐藏访存延迟，用大缓存、乱序执行、多级分支预测；GPU 则用海量 Warp 轮流上阵，把访存延迟“盖”在别的 Warp 的计算时间里。这也是为什么执行单元利用率对 Warp 数量极其敏感：驻留 Warp 太少，盖不住延迟，执行单元就会频繁空闲。

3.4 指令发射到运算完成的全过程

如果你跟踪一个 Warp 的一条 FMA 指令，大致流程是这样的：

1. 取指：Instruction Cache（指令缓存）中按程序计数器取出当前 Warp 的指令。
2. 译码：解析出操作码、目标寄存器、源寄存器、立即数。
3. 操作数收集：从寄存器文件读取源操作数。如果操作数来自共享内存或全局内存，需要先通过加载指令把数据搬到寄存器。
4. 发射：Warp Scheduler 把指令分配给对应的执行单元端口（FP32/Tensor Core/LD-ST）。
5. 执行：执行单元的流水线开始计算。FP32 FMA 通常有几个周期的流水线延迟。
6. 写回：结果写回目标寄存器，然后这个 Warp 可以被再次调度执行后续指令。

每一步都可能成为瓶颈。比如操作数不在寄存器里而在全局内存里，就会产生长等待；执行端口本身被其他 Warp 占满，这条指令就要在发射队列里排队。这也是为什么优化 GPU 程序时总在强调：减少全局内存访问、提高计算访存比、让寄存器里的数据尽量被重用。

4. 让执行单元真正忙起来：Occupancy、访存与算法设计

4.1 占用率（Occupancy）不是越大越好

你打开 Nsight Compute 或 Nsight Systems，能看到一个指标叫 Achieved Occupancy，即“实际达到的占用率”。它表示平均每个周期驻留在 SM 上的 Warp 数量占理论最大 Warp 数量的比例。如果最大能驻留 64 个 Warp，实际平均只有 32 个，占用率就是 50%。

一般来说，占用率高有利于隐藏访存延迟，但如果占用率已经能盖住延迟，继续提高占用率可能因为缓存/寄存器文件竞争反而降低性能。因为每个 SM 的寄存器文件、共享内存是固定的，比如 A100 SM 有 64K 个 32 位寄存器（256KB）。如果每个线程用 32 个寄存器，理论上一个 SM 能驻留 2048 个线程（64 个 Warp）；但如果每个线程用 255 个寄存器，那么最多只能驻留 256 个线程（8 个 Warp），占用率会掉到 12.5%。

所以在写 CUDA 或优化算子库时，常常需要在“每个线程多干活（多用寄存器做缓存）”和“多驻留 Warp（隐藏延迟）”之间平衡。经典的做法是让 occupancy 不低于某个阈值（比如 50%），然后再通过增加块内数据复用、向量化等来提升单线程效率。盲目调大块大小不一定有用，因为块大小受限于每个 SM 最大线程数、最大块数、寄存器/共享内存限制。

4.2 执行单元喜欢“数据一直在身边”

执行单元本身只是一个运算器，它的输入端是寄存器。寄存器里的数据从哪里来？从共享内存、L1/L2 缓存、全局内存（显存）一层层搬上来。GPU 的显存带宽虽然比 CPU 内存高不少（A100 约 2TB/s，H100 约 3.35TB/s），但相比动不动几百 TFLOPS 的算力来说，还是远远不够。你要是让每个执行单元每做一次乘加都要从显存读两个数，那么算力再高也跑不过内存带宽。

比如 A100 的 FP32 算力约 19.5 TFLOPS，意味着每秒需要约 19.5G 次 FMA 操作，如果每次 FMA 需要 8 字节（两个 FP32 操作数）的新鲜数据，那么理想带宽需求是 156GB/s，看似不高。但 Tensor Core 的 FP16 算力 156 TFLOPS，对应 78G 次 FMA 每秒，假如每次需要从显存读取两个 FP16 操作数（4 字节），带宽需求就是 312GB/s 以上，已经迫近带宽上限。如果做矩阵乘法时不把某个小矩阵块重复加载到共享内存/寄存器，而是每个乘积都从显存里搬，真实程序会慢几十倍。

这就是为什么矩阵乘法优化（也叫 GEMM 优化）里面一定要做“分块”：把大矩阵 C = A×B 分成多个 tile，每个 tile 由同一个线程块负责。先把 A 的 tile 和 B 的 tile 从显存加载到共享内存，再从共享内存读到寄存器，让数据在共享内存和寄存器之间反复重用。这样执行单元每次计算消耗的显存带宽极小，计算访存比大幅提升。CUDA 之所以提供shared共享内存，就是为了配合执行单元的数据复用。

4.3 算子融合与 Flash Attention 的本质

Transformer 大模型里最耗时的算子之一是 Attention。经典实现需要将 S = Q × K^T 写回显存，再做 softmax 后与 V 相乘。每一步中间结果都很大，如果全部在显存里晃一圈，执行单元会被迫等待数据搬运。Flash Attention 的思路就是不把中间矩阵写回显存，而是在一个 kernel 内部、通过分块循环计算 attention，把中间结果留在寄存器/共享内存里，减少全局内存读写。

执行单元本质上喜欢“一个 kernel 里连续做很多次计算”，而不是频繁地“写一个中间结果 → 启动另一个 kernel 读它”。这也是为什么 PyTorch 里的 torch.compile，或者更底层的 CUDA Graph 能带来加速：它们把原来多次 launch 的计算图合并，减少 kernel 间的启动和内存读开销，让执行单元在前一个任务刚结束时马上进入下一个任务，保持高利用率。

对于跑大模型的人来说，理解这一点非常有价值。你让 llama.cpp 或用 vLLM 跑推理时，为什么批处理（batch）越大吞吐越高？因为当 batch 增大到一定程度，矩阵乘法变成了大 GEMM，Tensor Core 能充分发挥，显存访存也被摊薄。反过来 batch=1 时，很多时间花在访存和延迟暴露上，执行单元喂不满，GPU 利用率就低。

5. 实操视角：执行单元在 AI 框架和推理引擎里如何被“喂饱”

5.1 PyTorch 训练/推理时发生了什么

当你写一行loss = model(x)，PyTorch 会经过如下链路：Python 层调度 → 算子分发（dispatch） → CUDA 实现（ATen 库） → cuDNN/cuBLAS/cutlass 等底层库 → 生成 CUDA kernel → 驱动把 kernel 加载到 GPU → SM 调度器把线程块分配给 SM → 执行单元执行。

如果这中间任何一层没配对，最后执行单元就吃不到任务。常见问题是：装了 CPU 版 PyTorch，或者 CUDA 工具包版本与显卡驱动不匹配，程序完全跑在 CPU 上，执行单元自然零使用。比如torch.cuda.is_available()返回 False、nvidia-smi 里看不到进程，大概率就是安装问题。另外，CUDA 不一定等于 CuDNN：跑卷积网络时如果 CuDNN 没装好，某些算子会退回慢速路径，执行单元利用率也会受影响。

5.2 llama.cpp / Ollama 的 GPU 执行单元利用

llama.cpp 最早以纯 CPU 推理出名，但后来也支持 CUDA/OpenCL/Metal 后端。它把 Transformer 的算子拆成一个个小 GEMM 或 GEMV 内核。当你用./llama-cli -m model.gguf -ngl 999时，-ngl 参数控制多少层放到 GPU 上。放得少，执行单元只负责部分算子；放得多，GPU 才真正成为主力。Ollama 底层也是类似逻辑，只要正确安装了 NVIDIA 容器工具包或驱动，模型会优先加载进显存由 GPU 执行。

很多 Windows 用户遇到“Ollama 未使用 GPU”或“llama.cpp 没跑 GPU”的问题，原因通常是驱动版本太旧、CUDA 环境没配好，或者环境变量CUDA_VISIBLE_DEVICES把 GPU 屏蔽了。从执行单元视角看，这相当于你把一整套工厂建好了，但工人（执行单元）全被锁在车间外面，只能靠 CPU 这个“外包团队”干慢活。

5.3 GPU 显存和执行单元到底谁更重要

这是热词搜索里一个高频问题：“GPU 显存容量是测算推理还是训练用的？”答案其实要分开看。显存容量决定你能不能放下模型和中间激活值，属于“仓库容量”；执行单元的算力决定你计算快不快，属于“加工速度”。两者缺一不可。

推理场景中，显存需求主要由模型权重、KV Cache、输入激活决定。如果显存不够，要么换小模型，要么量化到 INT8/INT4，要么做 KV Cache 分页（vLLM 等推理框架的做法）。训练场景中，显存要额外放下优化器状态和梯度，通常比推理大好几倍。但执行单元的利用率同样关键：有人用一张 4090 双卡跑大模型微调，结果发现 GPU Util 一直在 30% 以下，这往往不是因为显存不够，而是因为数据加载、CPU 预处理、梯度同步、算子实现低效导致执行单元喂不饱。这种情况下加显存没有意义，应该先修数据管线和算子融合。

5.4 微调大模型时的执行单元调度观察

用 LoRA 微调大模型时，参与计算的有两类执行单元：一部分是基座模型的 weight matrix × input 的 GEMM，另一部分是 LoRA 分支的低秩矩阵乘法。LoRA 分支参数量小，如果不仔细 Pin 内存、不把多次小 GEMM 融合，很容易出现 kernel launch 生开销大于 GPU 计算开销的情况。这也是为什么业界喜欢用 Unsloth、FlashAttention、torch.compile 这类工具：它们通过融合、算子改写，把那些“碎片化”的小算子合并成大 kernel，让执行单元能保持较长时间的高密度运算。

6. 用 nvidia-smi 和 Nsight 看清执行单元的真实状态

6.1 nvidia-smi 里的 GPU-Util 到底代表什么

先做一个很反直觉的说明：nvidia-smi里显示的 GPU-Util 并不是“所有执行单元的平均利用率”。它的含义是在过去一段时间窗口内，至少有一个计算内核（kernel）在 GPU 上运行的百分比。换句话说，只要任何一个 SM 在跑，GPU-Util 就可能显示 100%，哪怕只用了一个 SM 的计算能力。反过来，如果 GPU-Util 显示 50%，不等于一半执行单元空闲，可能是整个 GPU 只有一半时间在跑 kernel，另一半完全空闲。

所以用 GPU-Util 判断“执行单元是否跑满”是不严谨的。更细的指标需要它旁边的Memory也算不上计算利用率，而是一个采样点上的显存读写活动。要看执行单元本身的状态，推荐nvidia-smi dmon或nvidia-smi -q -d UTILIZATION，其中SM字段是 SM 忙碌时间占比，更接近执行单元被调用的情况。不过 SM 忙碌也包含访存等待时间，真正区分“算数执行单元忙碌”和“访存等待”还得用 Nsight Compute。

6.2 用 Nsight Compute 定位执行单元瓶颈

如果你在做 AI 算子开发或性能调优，Nsight Compute（NVIDIA 官方 profiling 工具）能逐 kernel 地分析执行单元的活动。比较关键的指标包括：

通常判断逻辑是：如果Long Scoreboard占比很高，说明执行单元在等待数据，需要优化访存，比如用共享内存、向量化、提高缓存命中。如果Executed Ipc接近理论峰值但整体性能还行，说明执行单元已经被喂满，想提速只能从降低计算量（量化、稀疏）入手。如果 Tensor Active 很低，但模型用了 FP16，那你就要检查算子是否真的走到 Tensor Core 路径了。

6.3 一个快速验证执行单元的 mini 实验

你可以用一段非常简单的 CUDA 代码验证执行单元的行为：分别跑一个涉及大量整数地址计算的 kernel 和只做连续 FMA 计算的 kernel，用nvidia-smi dmon观察 SM 和 Memory 的差别。如果你没有 CUDA 编程环境，也可以用 PyTorch 做类似实验：创建两个很大的随机矩阵，分别进行 FP32 和 FP16 的矩阵乘法，然后看nvtop或 Nsight 中 SM 利用率变化。FP16 乘法如果能调用 Tensor Core，SM 同样忙碌，但 FLOPS 和吞吐会高很多。

对于不想写代码的人，最简单的方法是直接跑一个已优化好的 AI benchmark，比如 llama.cpp 的-t多线程与-nglGPU 层数对比，或者用 vLLM 跑一个 batch 推理压测。观察相同显存占用下，GPU Util 和执行单元是否随 batch 增大而上升。这能帮你理解“执行单元饱和”和“访存饱和”的区别。

6.4 GPU 虚拟化、实例化与执行单元隔离

热词里有人问“GPU 实例化到底减少的是什么？具体原理是什么”。这里指的是 MIG（Multi-Instance GPU）或 vGPU 这类技术。MIG 把一块物理 GPU 切分成多个独立实例，每个实例拥有独立的 SM 子集、L2 缓存切片和显存控制器。例如 A100 可以切成 1 个 7 个 compute instance 或 2 个 2g.10gb 等多个配置。切分后，不同实例的执行单元彼此隔离，一个实例的任务不会抢占另一个实例的 SM。这在大模型训练里很有用：如果你用 N 个租来的 GPU 虚机，可能底层就是一台物理 GPU 用 MIG 拆出来的，每个虚机拿到的“GPU 算力”只是整卡执行单元的一部分。

GPU 服务器运维中常说的“GPU 调度”，本质上就是决定哪个容器/进程使用哪张卡、哪个 MIG 实例，以及在抢占场景下怎么隔离。这和 CPU 的 cgroup 调度不同，GPU 调度器往往由驱动、容器运行时和编排平台（Kubernetes + Device Plugin）共同完成。如果调度不当，多个容器挤在同一张卡上，各自抢执行单元，性能就会急剧下降。

最后，我想分享一个我个人的习惯。每次接手一个新环境或者部署新的 AI 推理服务，我都会先跑一分钟 nvidia-smi dmon 和一次简单的矩阵乘 benchmark，确认执行单元确实被调用，再开始跑真正的模型。因为这些问题如果不先排除，后面会浪费非常多的时间在模型参数和缓存设置上瞎折腾。GPU 执行单元的知识，说到底是帮助你建立一个“算力是怎么被用掉的”心智模型：显存是仓库，执行单元是工人，调度器是车间主任，算法是生产流程。只有把这四者对齐了，你的 AI 任务才能丝般顺滑地跑在 GPU 上。