如果你最近在折腾 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 个 SMStreaming 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 架构里有“CUCompute Unit”“Stream Processor”Intel 的 Xe 架构里有“Xe Core”“Vector Engine”如果延伸到专用 AI 芯片比如昇腾、Habana、Google TPU又有“AI Core”“MXU”“脉动阵列”等名称。但无论叫什么它们指的都是真正接收指令、完成算术运算、回写结果的那部分电路。一个执行单元最基本的能力是执行指令。一条机器指令从取指、译码之后最终会进入执行单元里完成具体操作。所以 GPU 算力的上限直接取决于执行单元的数量、工作频率以及每个周期能完成多少次浮点或整型运算。比如 A100 官方标称 FP32 算力约 19.5 TFLOPSTFLOPS 的计算公式粗略看就是“执行单元数量 × 频率 × 每周期 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 和 SFU2.1 一个 SM 里住着哪些工种如果细看一份 NVIDIA 架构图比如 Ampere、Hopper、Ada Lovelace你会发现每个 SM 内部并不是只有一种执行单元而是包含多个不同类型的“工人”。它们各司其职处理不同类型的指令。以典型的 Ampere A100 SM 为例大致有这几种执行单元类型数量A100 SM主要职责FP32 CUDA Core64单精度浮点加减乘、FMA融合乘加INT32 单元64整型运算、地址计算、逻辑操作Tensor Core第三代4矩阵乘加、混合精度FP16/BF16/FP32 累加等LD/ST 单元32通常每分区若干加载/存储访存指令SFU 特殊函数单元16求倒数、平方根、三角函数、超越函数等注意这些数字不是固定的。图灵架构的 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即每秒浮点运算次数。而浮点运算里最划算的是 FMAFused 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 × 2FMA 算两次 大约 19.5 TFLOPS。所以看到算力数字时你可以反推出执行单元的大致频率和数量。而在 AI 模型里卷积和全连接层的核心就是乘加操作FMA 单元自然成了执行单元中的主力。2.3 Tensor Core专为 AI 打造的矩阵执行单元普通 FP32 单元一次只能处理两个数相乘再加到一个数上。但在 GEMM 中我们希望一次指令能处理一个矩阵块比如 A 是 4×4 的 FP16B 是 4×4 的 FP16C 是 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。到了 H100FP16 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, blockDimGPU 的硬件调度器并不会把每个独立线程逐个分配给某个执行单元而是以线程块为单位把任务分发到不同的 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 指令大致流程是这样的取指Instruction Cache指令缓存中按程序计数器取出当前 Warp 的指令。译码解析出操作码、目标寄存器、源寄存器、立即数。操作数收集从寄存器文件读取源操作数。如果操作数来自共享内存或全局内存需要先通过加载指令把数据搬到寄存器。发射Warp Scheduler 把指令分配给对应的执行单元端口FP32/Tensor Core/LD-ST。执行执行单元的流水线开始计算。FP32 FMA 通常有几个周期的流水线延迟。写回结果写回目标寄存器然后这个 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/sH100 约 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 增大到一定程度矩阵乘法变成了大 GEMMTensor Core 能充分发挥显存访存也被摊薄。反过来 batch1 时很多时间花在访存和延迟暴露上执行单元喂不满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 ComputeNVIDIA 官方 profiling 工具能逐 kernel 地分析执行单元的活动。比较关键的指标包括指标含义瓶颈指向SM BusySM 有活动指令的时间占比整体执行单元忙碌程度Issue ActiveWarp Scheduler 发射指令的周期占比指令供给是否充足Executed Ipc每周期实际执行指令数是否达到执行单元带宽上限Stall Long ScoreboardWarp 等待全局内存访问访存延迟Stall WaitWarp 等待固定硬件资源执行端口/固定单元冲突Tensor ActiveTensor Core 活跃周期Tensor Core 是否被利用FMA ActiveFP32/FP64 FMA 单元活跃周期CUDA Core 是否被利用通常判断逻辑是如果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 CoreSM 同样忙碌但 FLOPS 和吞吐会高很多。对于不想写代码的人最简单的方法是直接跑一个已优化好的 AI benchmark比如 llama.cpp 的-t多线程与-nglGPU 层数对比或者用 vLLM 跑一个 batch 推理压测。观察相同显存占用下GPU Util 和执行单元是否随 batch 增大而上升。这能帮你理解“执行单元饱和”和“访存饱和”的区别。6.4 GPU 虚拟化、实例化与执行单元隔离热词里有人问“GPU 实例化到底减少的是什么具体原理是什么”。这里指的是 MIGMulti-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 上。