首页 高科技实验园地 NVIDIA Spark
N

GPU, CUDA 与 LLM 训练原理互动实验室 费曼学习法

解密 NVIDIA 硬件架构、CUDA Tiling 内存优化与大模型算力/显存计算

Batch Size ($B$):
序列长度 ($S$):
隐藏层维度 ($H$):
Transformer 层数 ($L$):
计算精度 Precision
单卡预估总显存需求

费曼比喻:教授 vs 工厂万名流水线工人

CPU(中央处理器) 就像一个 诺贝尔奖数学教授。他智商极高,任何复杂的逻辑跳转(逻辑分支)都能瞬间搞定,但他只有一双手,一次只能一题一题算。
GPU(图形处理器) 则是由 10,000 个小学流水线工人(CUDA 核心) 组成的超大型工厂。工人单独拆出来连微积分都不会,但他们被编成一个个严格对齐的小组(Warp 线程束),在统一命令(SIMT)下同时干简单的加减乘除!

现代 CPU 架构 (延迟优化型) 8~64 个强力核心

追求单线程极速、分支预测与超大 L3 缓存,减少指令等待延迟。

巨型共享 L3 Cache & 逻辑预测单元 占用晶片 ~60% 面积
优势: 处理极复杂的条件分支代码,单核执行速度极快。
劣势: 面对巨型矩阵相乘,算力吞吐量(TFLOPS)严重不足。
NVIDIA GPU 架构 (吞吐量优化型) 10,000+ CUDA 核心 (100+ SM)

追求极致并行吞吐(SIMT 架构),成千上万线程并发执行。

高带宽显存 HBM3e / VRAM 高达 3.35 TB/s 超大吞吐带宽
优势: 巨量计算密集型(矩阵乘法)加速,海量核心并发。
劣势: 线程分支分化(Warp Divergence)会导致性能骤降。

流多处理器 SM (Streaming Multiprocessor) 深度拆解

GPU 的核心引擎块。旗舰级 H100 GPU 内部包含约 132 个 SM。

SM 结构布局 (划分为 4 个 Processing Blocks 处理块)
超高速共享内存 Shared Memory / L1 缓存 单个 SM 拥有 228 KB (超低延迟 1~2 纳秒)
~20 TB/s SM 汇总聚合带宽

SM 内部核心部件详解

  • FP32 / INT32 核心: 处理常规标量浮点与整数算术加减乘除。
  • Tensor Core(张量核心): 专为 AI 打造的专用硬件,一个时钟周期内直接完成矩阵乘加 $D = A \cdot B + C$。
  • Shared Memory(共享内存): 软件可调配的高速片上 SRAM。同一个 Block 块内的线程通过它极速共享数据,避免访问慢速显存。

Warp 线程束执行模型 (32 线程步调一致)

NVIDIA GPU 执行的基本硬件调度单位:32 个线程组成一个 Warp,同时执行同一条指令。

SIMT 单指令多线程
硬件指令流: FMA R1, R2, R3 (融合乘加运算) Clock Cycle #128
费曼提醒:注意“线程束分支分化 (Warp Divergence)”!

如果在 CUDA 代码里写了 if (threadIdx.x % 2 == 0) 这样的条件分支,16 个线程在做 IF 时,另外 16 个线程只能**傻傻干等待**!然后再交换。GPU 性能会直接砍半!因此 AI 算子代码必须极致规避逻辑分支。

Tensor Core: $16 \times 16$ 矩阵乘加 (MMA) 硬件引擎

底层指令 mma.sync 一口气对 $A$ 和 $B$ 矩阵进行乘法,并累加到 $C$ 矩阵中。

矩阵 A ($16 \times 16$)
FP16 / FP8 输入
×
矩阵 B ($16 \times 16$)
FP16 / FP8 输入
+ C = D
累加器 D ($16 \times 16$)
FP32 高精度累加输出

费曼比喻:CUDA 编程结构与“办公桌 vs 远处仓库”

• Grid (网格) = 整个大工厂 (GPU 全局)。
• Block (线程块) = 工厂里的某个工作小组 (分配给 1 个 SM)。
• Thread (线程) = 工人组里的某个工人 (CUDA 核心)。
• 寄存器与共享内存 = 你办公桌上的草稿纸(几纳秒拿取)。而 全局显存 (VRAM) = 距离你 50 公里外的大仓库(拿趟数据要耗费 100 倍的时间!)。

CUDA 编程模型层次 $\rightarrow$ 物理硬件映射表

1. CUDA 网格 (Grid) Kernel 级
整个 GPU 设备 (所有 SMs)

一次内核调用启动的全局海量并发任务。

2. 线程块 (Thread Block) 协作组
调度到单个 SM 上

共享 SRAM,可通过 __syncthreads() 屏障同步。

3. 线程束 (Warp) 硬件执行
32 线程步调一致

由 SM 硬件 Warp Scheduler 调度器管理。

4. 线程 (Thread) 标量单元
运行于 1 个 CUDA 核心

拥有专属私有寄存器与 threadIdx.x。

GPU 内存金字塔带宽与访问延迟光谱

1. 寄存器 Registers (线程私有) 最快存储,紧贴核心算术单元
~0.5 ns (1 时钟周期) 30~40 TB/s 聚合总带宽
2. 共享内存 Shared Memory / L1 (Block 级) 软件可调配的高速片上 SRAM 缓存
~1.5~2 ns (15~20 周期) 15~20 TB/s 聚合带宽
3. L2 缓存 L2 Cache (GPU 全局) 所有 SM 共享的高速二级缓存 (约 50MB~500MB)
~15~20 ns (150 周期) 6~12 TB/s 带宽
4. 全局显存 Global Memory (HBM3e / VRAM) 片外巨型显存 (80GB~141GB),大模型权重存放处
~200~300 ns (400+ 周期!) 3.35 TB/s 极限带宽

当前选中内存层级说明

费曼核心:用共享内存 SRAM 做分块缓存,是把 GPU 算力提升 10 倍的突破口!

CUDA 共享内存分块算法 (Shared Memory Tiling) 模拟器

演示如何通过 SRAM 缓存切块,减少对慢速 VRAM 显存重复读取。

矩阵 A (VRAM 显存)
+
矩阵 B (VRAM 显存)
=
片上 SRAM 共享内存

VRAM 显存访问次数直接降低 16 倍!

费曼比喻:LLM 计算的两大瓶颈——计算密集型 vs 内存带宽瓶颈

1. 计算吞吐瓶颈 (Compute-Bound): 比如巨型矩阵相乘 ($Q \times K^T$, FFN 线性层)。这种任务全是密集数学题,GPU Tensor Core 核心可以全力开火 100% 跑满。
2. 显存带宽瓶颈 (Memory-Bound): 比如 Softmax, LayerNorm, 旋转位置编码 (RoPE)。数学公式极简单,但需要把整个大矩阵从 VRAM 读一遍再写一遍!GPU 绝大多数时间在**傻等显存数据送达**。FlashAttention 算法就是为了消除这种等待而发明的!

Transformer 层前向传播 (Forward Pass) 算子流

分步查看单个 Attention & Feed-Forward 层内部的数据计算流程。

单步估算浮点运算量 (FLOPs)
主要硬件执行单元:

LLM 训练显存占用 (VRAM Breakdown)

1. 模型静态权重 ( 参数, )
2. 梯度数据 ($\nabla L$)
3. Adam 优化器状态 (一阶矩 $m$ 与二阶矩 $v$ FP32)
4. 前向传播激活值 (Activation)
单 GPU 必备最小显存:

训练一轮 (Forward + Backward) 算力总量

# 1. 单 Step Token 吞吐量
# 2. 算力消耗估算 (每 Token 约 6P 次浮点运算)
FLOPs ≈ 6 × 参数量($P$) × Token数
费曼提醒: 发现了吗?Adam 优化器占用的显存竟然是权重本身的 2~4 倍!这就是为什么分布式训练中必须引入 DeepSpeed ZeRO-3 将优化器切碎分摊到不同 GPU 上。