GPU的一些概念-拓冰网站定制

GPU 的“心脏”是SM（流式多处理器），它负责将软件线程映射到硬件单元。为了理解它是如何驱动 AI 与高性能计算的，我们需要深入其内部执行单元（Warp、Tensor Core、TMA）的协同机制。

一、SM：GPU 的“计算堡垒”

SM 是 GPU 真正的计算核心，负责将软件线程映射到硬件单元。它内部包含四大关键模块：

模块	功能	角色定位
Warp Scheduler	指令调度	大脑。每个周期挑选就绪的 Warp 发射指令，通过零开销切换掩盖内存延迟。
Register File	寄存器堆	血液。存储线程的私有变量，容量巨大（A100 每个 SM 约 256KB）。
Execution Units	执行单元	肌肉。包含 CUDA Core（标量计算）、Tensor Core（矩阵计算）、LD/ST（内存读写）。
Shared Memory / L1	片上缓存	高速暂存区。由程序员显式管理，用于线程块内数据共享。

调度逻辑：SM 并不直接调度单个线程，而是将线程块（Block）切分为Warp（32 个线程一组）进行调度。当某个 Warp 因等待数据而“卡住”（Stall）时，调度器会立即切换到另一个就绪 Warp，保持计算单元忙碌。

二、Warp：SIMT 执行的基本单位

Warp 是 GPU 硬件调度的最小粒度，理解它是理解 GPU 并行性的关键。

SIMT 模型（单指令多线程）：一个 Warp 内的 32 个线程在同一周期执行同一条指令，但操作不同的数据。这就像 32 个士兵听同一个口令，同时做同样的动作。
分支发散（Divergence）：如果 Warp 内线程出现 if-else 分支（部分线程走 if，部分走 else），GPU 会串行化执行所有分支路径，未参与的线程被禁用。这是性能杀手。
状态机：Warp 在 SM 中拥有独立的状态（Active/ Eligible/ Stalled），调度器只挑选“Eligible”（就绪）的 Warp 发射。

三、Tensor Core：矩阵计算的“特种部队”

Tensor Core 是专门为矩阵乘累加（MMA）设计的专用单元，与通用 CUDA Core 有本质区别。

特性	CUDA Core	Tensor 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、Warp（32 线程）
软件栈	CUDA C/C++、Runtime、Driver
库	cuBLAS、cuDNN、NCCL、CUTLASS
工具	nvcc、Nsight、nvprof

2. CUDA 的线程层次（软件视角）

Grid └── Block └── Thread

Thread：一个线程，处理一个数据元素
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 的执行单位：Warp

32 个 thread 组成一个warp
同一个 warp 的线程同一条指令、不同数据（SIMT）
分支分歧会导致串行执行

七、TPC 是什么

TPC = Texture Processing Cluster（纹理处理簇）

是 NVIDIA GPU硬件架构中的中间层。

1. GPU 硬件层级

GPU └── GPC（Graphics Processing Cluster，图形处理簇） └── TPC（Texture Processing Cluster，纹理处理簇） └── SM（Streaming Multiprocessor，流多处理器） └── CUDA Core / Tensor Core

现代数据中心 GPU（A100/H100）：

1 个 TPC ≈2 个 SM
TPC 还包含：
- 纹理单元（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 深度优化的基础。

资讯详情

GPU的一些概念