GPU 的“心脏”是SM流式多处理器它负责将软件线程映射到硬件单元。为了理解它是如何驱动 AI 与高性能计算的我们需要深入其内部执行单元Warp、Tensor Core、TMA的协同机制。一、SMGPU 的“计算堡垒”SM 是 GPU 真正的计算核心负责将软件线程映射到硬件单元。它内部包含四大关键模块模块功能角色定位Warp Scheduler​指令调度大脑。每个周期挑选就绪的 Warp 发射指令通过零开销切换掩盖内存延迟。Register File​寄存器堆血液。存储线程的私有变量容量巨大A100 每个 SM 约 256KB。Execution Units​执行单元肌肉。包含 CUDA Core标量计算、Tensor Core矩阵计算、LD/ST内存读写。Shared Memory / L1​片上缓存高速暂存区。由程序员显式管理用于线程块内数据共享。调度逻辑SM 并不直接调度单个线程而是将线程块Block切分为Warp32 个线程一组进行调度。当某个 Warp 因等待数据而“卡住”Stall时调度器会立即切换到另一个就绪 Warp保持计算单元忙碌。二、WarpSIMT 执行的基本单位Warp 是 GPU 硬件调度的最小粒度理解它是理解 GPU 并行性的关键。SIMT 模型单指令多线程一个 Warp 内的 32 个线程在同一周期执行同一条指令但操作不同的数据。这就像 32 个士兵听同一个口令同时做同样的动作。分支发散Divergence如果 Warp 内线程出现 if-else 分支部分线程走 if部分走 elseGPU 会串行化执行所有分支路径未参与的线程被禁用。这是性能杀手。状态机Warp 在 SM 中拥有独立的状态Active/ Eligible/ Stalled调度器只挑选“Eligible”就绪的 Warp 发射。三、Tensor Core矩阵计算的“特种部队”Tensor Core 是专门为矩阵乘累加MMA设计的专用单元与通用 CUDA Core 有本质区别。特性CUDA CoreTensor Core计算粒度​标量一次算一个数矩阵块如 4x4, 8x8, 16x16操作​单精度浮点 (FP32) 等混合精度FP16/BF16/FP8 输入FP32 累加吞吐量​高极高专为深度学习 GEMM 优化执行机制Tensor Core 指令通常由 Warp 发起。在 Hopper 架构中它支持更灵活的Warp Group​ 模式允许跨 Warp 协作进行更大规模的矩阵计算同时引入了TF32​ 和FP8​ 格式以加速 AI 训练。四、TMA数据搬运的“异步引擎”Tensor Memory Accelerator (TMA)​ 是 Hopper 及后续架构引入的专用 DMA 单元旨在解决“喂饱 Tensor Core”的内存瓶颈。异步搬运TMA 负责在全局显存Global Memory和共享内存Shared Memory之间搬运大块数据。它独立于计算单元运行实现了“计算与数据搬运”的流水线并行。地址计算卸载传统模式下线程需要计算复杂的内存地址如矩阵分块消耗寄存器。TMA 通过TMA Descriptor​ 描述张量形状由硬件自动计算地址极大减轻了 CUDA Core 的负担。多播Multicast在Thread Block Cluster线程块集群模式下TMA 可将一份数据从全局内存一次性搬运到多个 SM 的共享内存中减少冗余读取。五、四者协同从指令到结果的全流程以 Hopper 架构执行一次大矩阵乘法GEMM为例SM 内部的协同流程如下线程映射Grid 中的 Thread Block 被分配到 SM 上并划分为 Warp。TMA 预加载某个 Warp 中的线程配置 TMA Descriptor发起异步拷贝cp.async.bulk将数据从全局内存搬至共享内存。此时 CUDA Core 无需等待可立即执行其他计算。流水线重叠当 TMA 在后台搬运数据时Warp 调度器切换到其他就绪的 Warp利用 Tensor Core 计算上一批数据。屏障同步计算 Warp 通过mbarrier异步内存屏障等待 TMA 搬运完成随后从共享内存加载数据到寄存器交由 Tensor Core 进行 MMA 运算。结果写回计算结果通过 TMA 或普通存储指令写回全局内存。关键进化在 Hopper 之前数据搬运依赖 Warp 的 LD/ST 指令会占用计算资源。TMA 的引入实现了真正的异步数据流让 Tensor Core 能持续处于“饱腹”状态这是现代大模型训练性能飞跃的关键。CUDA 是编程模型 / 软件平台六、CUDA 是什么CUDA Compute Unified Device Architecture​是 NVIDIA 的通用 GPU 并行计算平台和编程模型。一句话CUDA 让你用 C/C/Python通过 PyTorch/TensorFlow 底层写程序把计算放到 GPU 上并行执行。1. CUDA 包含什么层级内容编程模型Grid / Block / Thread执行模型Kernel、Warp32 线程软件栈CUDA C/C、Runtime、Driver库cuBLAS、cuDNN、NCCL、CUTLASS工具nvcc、Nsight、nvprof2. CUDA 的线程层次软件视角Grid └── Block └── ThreadThread一个线程处理一个数据元素Block一组线程跑在同一个 SM 上Grid一次 kernel 启动的所有 block示例__global__ void add(int *a, int *b, int *c) { int i blockIdx.x * block_size threadIdx.x; c[i] a[i] b[i]; }3. CUDA 的执行单位Warp32 个 thread 组成一个warp同一个 warp 的线程同一条指令、不同数据SIMT分支分歧会导致串行执行七、TPC 是什么TPC Texture Processing Cluster纹理处理簇​是 NVIDIA GPU硬件架构中的中间层。1. GPU 硬件层级GPU └── GPCGraphics Processing Cluster图形处理簇 └── TPCTexture Processing Cluster纹理处理簇 └── SMStreaming Multiprocessor流多处理器 └── CUDA Core / Tensor Core现代数据中心 GPUA100/H1001 个 TPC ≈2 个 SMTPC 还包含纹理单元Texture Unit纹理缓存指令缓存PolyMorph Engine图形相关2. TPC 的作用TPC不是 CUDA 编程模型里的概念而是硬件资源分组单位调度 / 功耗 / 图形处理的中间层帮助 GigaThread Engine 把 block 分给 SM简化理解TPC “SM 的小区”GPC “城区”SM “楼”CUDA Core “工人”总结GPU 的高性能源于精细的分工SM​ 作为容器管理资源Warp​ 作为调度单位隐藏延迟Tensor Core​ 作为算力引擎加速矩阵TMA​ 作为后勤部队保障数据供给。理解这四者的层级关系与协作机制是进行 CUDA 深度优化的基础。