4、GPU计算单元详解：SM/CU架构、CUDA核心/流处理器、张量核心（Tensor Core）、RT核心，以及专用单元的作用

GPU之所以能在并行计算、图形渲染和AI推理训练领域展现出远超CPU的吞吐能力，其根本原因在于内部计算单元的规模化组织方式。无论是NVIDIA的SM（Streaming Multiprocessor，流式多处理器）、AMD的CU（Compute Unit，计算单元），还是Intel的Xe核显架构，其本质都是将数千个简单计算核心按照“线程组—调度器—执行单元”的层级进行封装。理解SM/CU的内部结构、各类核心（CUDA Core、Tensor Core、RT Core）的职责边界，是进行GPU选型、内核调优和性能瓶颈分析的前提。本章将从硬件微架构视角，拆解这些计算单元的组成原理、工作流程与适用场景。

GPU计算单元层级架构与专用单元协同关系 GPU 芯片（GA102 / Navi 21 / Ponte Vecchio） GPC（图形处理簇）/ Shader Engine（着色器引擎）— 包含光栅化、裁剪、纹理单元 SM（流式多处理器） / CU（计算单元）— 每个SM包含以下子单元 CUDA Core / 流处理器 FP32/INT32 标量运算 MUFU（特殊函数） 张量核心 Tensor Core FP16/BF16/INT8 矩阵乘加 稀疏化加速 RT Core 光线追踪核心 BVH遍历/光线求交 三角形/包围盒测试 专用单元 LSU（加载/存储） SFU / 纹理单元 / 调度器 共享内存（Shared Memory） + L1 Cache + 寄存器堆（Register File） 每个SM独享，低延迟高带宽，线程间数据交换的关键 L2 Cache（芯片级共享，所有SM/CU共用） 显存（HBM2e / GDDR6X）— 容量大、延迟高 调度流程 1. Warp/Wavefront 2. 指令分发 3. 执行单元映射 4. 写回/同步 图：GPU计算单元层级架构与专用单元协同关系（教学简化图）

4.1 SM/CU架构：GPU并行计算的组织核心

核心原理：SM（NVIDIA）和CU（AMD）是GPU内部最基本的并行计算单元，相当于一个“微型多核处理器”。每个SM/CU内部包含多个处理块（Processing Block），每个处理块又由若干CUDA Core（或流处理器）组成。以NVIDIA Ampere架构的GA102芯片为例，一个SM包含128个CUDA Core、4个Tensor Core、1个RT Core（第二代）、4个纹理单元、32个LSU（Load/Store Unit）以及256KB的寄存器堆和128KB的共享内存/L1缓存组合。SM通过Warp Scheduler（线程束调度器）将32个线程组成一个Warp进行指令分发，每个调度器每个时钟周期可以发射一条指令到执行单元。

判断方法：在评估GPU计算能力时，不能只看“核心数”这一表面参数。需要关注以下关键指标：

  • SM数量与频率：SM数量乘以每个SM的CUDA Core数，再乘以2（FMA操作算两次浮点运算），再乘以核心频率，即可得到理论FP32算力（FLOPS）。
  • Occupancy（占用率）：指活跃Warp数量与最大Warp容量之比。高占用率有助于隐藏内存延迟，但并非越高越好，需结合寄存器压力和共享内存使用量综合判断。
  • 共享内存容量：决定线程块内数据复用能力。例如A100的SM拥有164KB共享内存，而消费级RTX 3080只有100KB，这直接影响某些矩阵转置、卷积内核的性能。
实战场景：在优化一个矩阵乘法内核时，我们观察到RTX 3090上的占用率仅为37%（理论最大为50%）。通过分析CUDA Occupancy Calculator发现，每个线程使用了40个寄存器，导致每个SM只能容纳8个线程块（每个块256线程）。将内核拆分为两个阶段，减少每个线程的寄存器压力至24个后，占用率提升至62%，实际性能提升了1.8倍。这说明SM的资源分配（寄存器、共享内存）比单纯增加线程数更重要。

4.2 CUDA Core与流处理器：标量计算的主力军

核心原理：CUDA Core（NVIDIA）和流处理器（AMD，Stream Processor）本质上是支持FP32和INT32运算的标量执行单元。每个CUDA Core在一个时钟周期内可以执行一次FMA（Fused Multiply-Add，融合乘加）操作，即计算 a = a * b + c。在Volta架构之前，CUDA Core仅支持FP32；从Turing架构开始，引入了独立的INT32执行单元，使得整数运算和浮点运算可以并行执行，提升了地址计算、索引操作等混合负载的效率。

操作流程：一个典型的CUDA Core执行流程如下：

  1. 取指（Fetch）：Warp Scheduler从指令缓存中取出当前Warp的下一条指令。
  2. 译码（Decode）：指令被解码为操作类型（FP32、INT32、分支跳转等）。
  3. 发射（Issue）：调度器将指令发送到对应的执行端口（如FP32单元、INT32单元、SFU特殊函数单元）。
  4. 执行（Execute）：CUDA Core完成算术运算，结果写入寄存器堆。
  5. 写回（Writeback）：结果被写回寄存器文件，供后续指令使用。
架构代际 CUDA Core/流处理器特性 FP32算力（RTX 3090） INT32算力 适用负载
Ampere（GA102） FP32+INT32独立单元，可并行 35.6 TFLOPS 17.8 TOPS 通用计算、图形渲染、物理模拟
Ada Lovelace（AD102） FP32+INT32独立，新增FP32吞吐提升 82.6 TFLOPS 41.3 TOPS 高精度科学计算、AI推理
RDNA 2（Navi 21） 统一着色器，FP32/INT32共享 23.0 TFLOPS 11.5 TOPS 游戏渲染、通用计算
常见误区：很多开发者认为CUDA Core越多性能一定越好。实际上，当内核存在大量分支发散（divergence）时，同一Warp内的线程会串行执行不同路径，导致CUDA Core利用率大幅下降。例如，一个包含 if (threadIdx.x % 2 == 0) 的内核，理论上会浪费一半的CUDA Core周期。此时，增加核心数并不能线性提升性能，优化分支结构或使用Warp级别编程（如Warp Shuffle）才是关键。

4.3 张量核心（Tensor Core）：AI与HPC的加速引擎

核心原理：Tensor Core是NVIDIA从Volta架构开始引入的专用矩阵乘加单元，AMD的CDNA架构也提供了类似的Matrix Core。Tensor Core的核心优势在于执行4×4矩阵乘加操作：D = A × B + C，其中A、B、C、D均为4×4矩阵。在A100上，每个SM包含4个第三代Tensor Core，每个时钟周期可执行1024次FP16 FMA操作（即2048 FLOPs）。Tensor Core支持FP16、BF16、INT8、INT4、TF32等多种精度，其中TF32（19位精度）专为AI训练设计，无需修改模型即可获得接近FP32的精度。

判断方法：在AI推理场景中，判断是否使用Tensor Core的关键指标是矩阵乘法的维度。如果模型中的全连接层、卷积层可以转化为批量矩阵乘法（Batch GEMM），且维度是8的倍数（FP16）或4的倍数（INT8），则Tensor Core可以发挥最大效率。使用cuBLAS、cuDNN或PyTorch的 torch.matmul 时，框架会自动调用Tensor Core，但需要确保输入数据的内存对齐（通常要求16字节对齐）。

案例拆解：某AI推理服务使用RTX 3080部署BERT-base模型，初始延迟为12ms/请求。通过NVIDIA Nsight Compute分析发现，FP32 CUDA Core利用率仅为23%，而Tensor Core利用率为0%。将模型从FP32转换为FP16（使用PyTorch的 model.half()），并确保输入张量调用 contiguous() 保证内存连续后，延迟降至4.2ms，吞吐量提升2.8倍。进一步使用INT8量化（通过TensorRT），延迟降至2.1ms，但精度损失了0.3%的F1分数，在可接受范围内。此案例说明Tensor Core的启用需要数据精度、内存布局和框架支持三方面配合。

4.4 RT Core与专用单元：图形与特殊计算任务

核心原理：RT Core（Ray Tracing Core）是NVIDIA Turing架构引入的光线追踪专用硬件，负责执行光线与三角形/包围盒的求交测试以及BVH（Bounding Volume Hierarchy，层次包围体）树的遍历。在Ampere架构中，RT Core支持光线追踪、着色器重排序（Shader Execution Reordering）和透明度映射。AMD的RDNA 2也提供了类似的光线加速器（Ray Accelerator），但实现方式不同。RT Core的算力通常用“Giga Rays/s”（每秒十亿条光线）衡量，RTX 3090可达35 GRays/s。

专用单元的作用：除了RT Core，SM内部还包含以下专用单元：

  • SFU（Special Function Unit）：执行超越函数（sin、cos、log、exp、rsqrt等），每个SFU每时钟周期处理4个线程的运算。SFU的吞吐量远低于CUDA Core，因此应避免在热循环中频繁调用这些函数。
  • LSU（Load/Store Unit）：负责内存访问指令，包括全局内存、共享内存和本地内存的读写。LSU的带宽决定了数据搬运效率，通常每个SM包含16-32个LSU。
  • 纹理单元（Texture Unit）：提供纹理过滤、采样和缓存功能，在图形渲染中用于贴图映射，在计算中可用于图像处理（如双线性插值）。
  • Warp Scheduler：每个SM包含4个（Ampere）或8个（Ada）调度器，每个调度器管理一组Warp，负责指令发射和线程调度。
专用单元 主要功能 典型吞吐量（RTX 3090） 优化建议
SFU 超越函数计算 16 ops/clock/SM 用多项式近似替代，或使用CUDA内置快速数学函数（__sinf）
LSU 内存访问 32 bytes/clock/SM 合并内存访问（Coalescing），避免未对齐访问
RT Core 光线求交与BVH遍历 35 GRays/s 使用OptiX或DXR API，减少无效光线发射
Tensor Core 矩阵乘加 1024 FP16 FMA/clock/SM 使用cuBLAS Lazy Loading，避免小矩阵（<64×64）
风险提示：在混合精度训练中，Tensor Core虽然能大幅提升吞吐，但FP16的有限动态范围可能导致梯度下溢（underflow）。建议使用损失缩放（Loss Scaling）技术，或采用BF16（Brain Floating Point，与FP32有相同指数位）来避免精度问题。此外，RT Core在非图形计算中（如物理模拟中的碰撞检测）也能发挥作用，但需要将场景数据转换为BVH结构，这本身有较高的构建成本，需评估是否值得。

4.5 总结与性能调优策略

GPU计算单元的设计哲学是“以空间换时间”——通过大量简单核心并行执行，掩盖内存延迟，提升吞吐量。在实际调优中，应遵循以下原则：

  1. 先分析瓶颈类型：使用NVIDIA Nsight Compute或AMD ROCm Profiler判断是计算密集型（Compute-bound）还是访存密集型（Memory-bound）。计算密集型优先优化Tensor Core或CUDA Core利用率；访存密集型则需优化数据布局和缓存命中率。
  2. 匹配数据精度与硬件能力：AI推理优先使用INT8/FP16，科学计算使用FP32/FP64，图形渲染使用FP32。不要盲目追求低精度，需验证精度损失。
  3. 关注Warp级效率：避免分支发散，使用Warp Shuffle减少共享内存访问，利用异步拷贝（Async Copy）隐藏数据搬运延迟。
  4. 合理配置资源：通过 __launch_bounds__ 限制每个线程块的线程数和寄存器使用量，以提高占用率。但需注意，过高的占用率可能导致缓存抖动，需通过实验找到平衡点。
经验总结：在GPU性能调优中，“先测量，后优化”是铁律。不要凭经验猜测瓶颈，而是使用Profiler获取SM利用率、内存带宽、Warp Stall原因等数据。例如，如果发现“Long Scoreboard”事件占比高，说明内存延迟未隐藏，应增加并行线程数或使用预取指令；如果“Wait”事件多，则可能是同步开销过大，需减少 __syncthreads() 调用或使用细粒度同步。
← 上一章 📖 返回目录 下一章 →