跳转到内容
Jason's Study

Lindholm 2008 Tesla — SM、warp、SIMT 这套词汇的官方出生证明

是什么

这是一篇 17 页的 IEEE Micro 论文，4 位 NVIDIA 工程师亲自把 G80 芯片（GeForce 8800 GTX 背后那块）拆给你看：每个零件叫什么名字、为什么这么摆、和上一代差在哪。

为什么把它单独拎出来读？因为今天所有讲 GPU 的文章——博客、教科书、Hopper 白皮书——都在用 SM / warp / SIMT 这套词汇。这些词的官方定义、缩写解释、画图惯例，全部出自这一篇。读后续任何一代 GPU 论文，省一半字典查询时间。

日常类比：英语里很多缩略词第一次出现在某个具体场合（比如 OK 来自 1839 年波士顿一份报纸的玩笑）。Lindholm 2008 就是 GPU 圈的”那份报纸”——不是它发明了硬件，但它把术语首次系统命名。

本站已有一篇 tesla-architecture-2008 讲 Tesla 是什么、为什么是 NVIDIA 的拐点。本篇专门做词汇表——把每个名词配一张脑内图像，读后续论文时能立刻接住。

为什么重要

不啃这一篇，下面这些词你只能模糊感觉：

  • 别人说”我这个 kernel 跑了 4 个 SM、occupancy 50%“——你能复述但画不出图
  • 看 flash-attention 论文里”shared memory 192KB”——不知道这层在哪一级、为什么是 KB 不是 GB
  • 读 hopper-architecture-2022 提到”warp specialization”——不清楚 warp 有多大、为什么必须是 32
  • 调 PyTorch profiler 看到 Achieved Occupancy 23%——不知道分子分母各是什么

核心要点

论文用 四层抽象 把 G80 描述清楚——你以后读任何 GPU 论文都用这四层：

  1. 芯片层：1 个 GPU = 8 个 TPC（Texture/Processing Cluster）。TPC 是物理打包单位，论文里之后基本不再提，因为程序员看不到它。

  2. SM 层（Streaming Multiprocessor）：每个 TPC 含 2 个 SM，全芯片共 16 个 SM。SM 是程序员能看到的最大独立计算单元——一个 block 整段挂在一个 SM 上跑，跨 SM 不能互通有无。

  3. SP 层（Streaming Processor，又叫 CUDA core）：每个 SM 内有 8 个 SP，全芯片 128 个。SP 跑标量浮点 / 整数指令。SP 的时钟（1.35 GHz）比芯片其他部分高，叫”shader clock”。

  4. warp 层：硬件强制把 32 个相邻线程编一组——这就是 warp。一个 warp 内 32 线程同一时刻执行同一条指令，硬件把这条指令在 8 个 SP 上跑 4 个时钟周期完成（8 × 4 = 32）。

SIMT 这个词的由来

论文第 42 页首次定义 SIMT：Single Instruction, Multiple Threads。它和 SIMD 的差别只在”程序员视角”：

  • SIMD：你必须写 vec4 a + vec4 b，向量是显式的
  • SIMT：你写 c[i] = a[i] + b[i] 像写普通线程，硬件偷偷把 32 个连续 i 捆成一束送进 SP

这句话以后会被引用上千次。所有”GPU 编程比 SIMD 友好”的说法都源于这一段。

存储分层（以后所有 GPU 论文都用这个表）

层级大小（G80）速度谁能访问
Register8K × 32-bit / SM1 周期单线程私有
Shared memory16 KB / SM~5 周期同 block 32 个线程共享
Local memory显存里划一块数百周期单线程私有但放显存
Global memory768 MB数百周期全 grid 共享
Constant64 KBcache 后 1 周期只读，全 grid
Texture显存 + 2D cachecache 命中快只读，带插值硬件

记住前 4 行就够了。Shared memory 这层是 G80 第一次引入的，后来 flash-attention 之类工作的命脉。

实践案例

案例 1：把”占用率”读懂

nvcc -Xptxas -v my_kernel.cu
> Used 24 registers, 4096 bytes smem

逐部分解释：

  • 24 寄存器/线程：SM 总共 8K 寄存器，所以 8192 ÷ 24 ≈ 341 线程能同时挂
  • 但 SM 上限是 768 线程（24 warp × 32），所以寄存器是当前瓶颈
  • 占用率 = 341 / 768 ≈ 44%
  • 想提升：降寄存器使用、或换更小的 block size

这个公式 2008 年论文里已经写清楚，今天 H100 改的只是数字（8K → 64K 寄存器、768 → 2048 线程），结构没变。

案例 2：warp 大小为什么是 32

__global__ void k() {
    int i = threadIdx.x;
    if (i % 2 == 0) A();  // 偶数线程
    else            B();  // 奇数线程
}

一个 warp 32 线程被切成 16 偶 16 奇。SIMT 硬件没有同时跑两条指令的能力——它先让 16 个偶线程跑 A（其余 16 个挂起），再让 16 个奇跑 B。总耗时 = A + B。

为什么是 32 不是 8 或 64？论文给出权衡：

  • 太小（如 8）：每个 warp 的调度开销摊销不开
  • 太大（如 64）：分支发散惩罚太重，且寄存器文件分组困难
  • 32 是 G80 内部数据通路宽度的整数倍——8 SP × 4 周期

案例 3：PTX 是为什么的

CUDA 源码先编译到 PTX（Parallel Thread Execution），再由驱动 JIT 到具体 GPU 的机器码：

.cu  →  nvcc  →  .ptx  →  driver JIT  →  SASS（具体硬件）

为什么多这一层？因为 NVIDIA 知道硬件每代会改（warp 大小可能变、寄存器组织会变）。PTX 是抽象 ISA，让 2006 年写的代码能在 2026 年的 Blackwell 上跑。这层设计在论文最后一节，是 NVIDIA 18 年生态护城河的工程根基。

踩过的坑

  1. 把 SP 想成 CPU 核心——错。SP 没有自己的取指 / 译码，它是 SM 共用一套指令通路上的”算术口子”。SM 才是”最像 CPU 核心”的单位。

  2. 以为 block 之间能同步——__syncthreads() 只在 block 内有效。block 之间想同步只能 kernel 退出后再启第二个 kernel。新人常写出”global barrier”然后发现根本不存在。

  3. shared memory 当 cache 用——shared memory 不是自动 cache，是程序员手动管的暂存板。要自己搬数据进来、自己同步、自己搬回去。论文里反复强调这点，但很多教程一笔带过。

  4. 以为 occupancy 越高越快——50% occupancy 经常比 100% 更快，因为高占用率意味着寄存器少、需要更多次访存。这条 2008 年论文已经埋了伏笔，到 Volta 时代被 Vasily Volkov 的 “Better performance at lower occupancy” 论文正式打出来。

适用 vs 不适用场景

这篇论文适合什么时候读：

  • 第一次接触 GPU 编程，需要建立词汇表
  • 读 Fermi/Volta/Hopper 白皮书前的”前置阅读”
  • 要给团队讲”GPU 为什么和 CPU 不一样”，需要权威引用源

不适合什么时候读：

  • 想学具体优化技巧 → 看 gpu-microbenchmarking-2010 或 NVIDIA 官方 Best Practices Guide
  • 想了解最新一代 GPU → 看 hopper-architecture-2022 直接对照
  • 想入门 CUDA 编程 → 跳过论文，直接 CUDA C++ Programming Guide 前 4 章

历史小故事（可跳过）

  • 2003 年：NVIDIA 内部代号 “Tesla” 启动，目标是把分散的 vertex / pixel / geometry shader 合并
  • 2006 年 11 月：G80 上市，Lindholm 等人写完芯片但论文没发
  • 2008 年 3-4 月：IEEE Micro 这一期发表，距离芯片上市晚了 16 个月——原因是 NVIDIA 想先让 CUDA 生态铺开，再把架构细节交给学界
  • 同期：AMD 的 R600 / Intel 的 Larrabee 都在做类似探索；Larrabee 失败、R600 思路接近但没配套软件，Tesla 一统江湖
  • 后续 18 年：每代 GPU（Fermi / Kepler / Maxwell / Pascal / Volta / Turing / Ampere / Hopper / Blackwell）的架构白皮书都按本论文的章节模板写

学到什么

  1. 词汇统一即生态：SM、warp、SIMT 这些词被 4 个工程师写进 17 页论文，之后所有人沿用。命名权是基础设施战争最便宜也最深的护城河
  2. 抽象层数刚刚好：grid / block / warp / thread 四层，少一层不够表达、多一层程序员记不住。论文里有句原话”每层抽象都对应一种硬件资源”——这是好抽象的判据
  3. PTX 是真正的护城河：硬件可以被对手追上，但 18 年向前兼容的 ISA + 工具链生态非一日之功
  4. 论文晚于产品 1.5 年：先让用户用上，再教学界怎么用——这种节奏适合正反馈强的平台型技术，不适合学术追求

延伸阅读

关联