GPU memory hierarchy 是怎样的？
一个 block 最多多少 thread？

英文问法Walk me through the GPU memory hierarchy. · How is shared memory (SRAM) different from HBM? · What's the maximum number of threads per block? · What is occupancy?

考过面经：NVIDIA DevTech NG 2026（"GPU 内存架构 SRM HBM区别"、"最高多少thread"，1p3a 1161658）；国内：NVIDIA 中国 AI Developer 一面（"GPU上的存储, global memory, local memory, shared memory, register的大小, 带宽, 访存延迟大概是在什么级别"）

先弄清楚：每个概念是什么

SM 与 grid / block / warp: GPU 由很多 SM（Streaming Multiprocessor）组成。kernel 启动成一个 grid，grid 拆成 block，每个 block 整个放在一个 SM 上跑。block 里每 32 个 thread 是一个 warp，同时执行同一条指令（见 GPU vs CPU）。
Register: 每个 thread 私有，在 SM 上，最快。由 compiler 分配，用太多就 spill 到 local memory。
Shared memory 与 L1（SRAM）: 每个 SM 上有一块 SRAM（unified data cache），一部分划给 shared memory，剩下当 L1 cache。shared memory 是程序员自己管的 scratchpad，同一个 block 的 thread 共享。例：H100 每 SM 共 256 KB，其中 shared memory 最多 228 KB，一个 block 最多用 227 KB。
L2 与 HBM（global memory）: L2 在芯片上、所有 SM 共享。HBM 是片外的 GPU memory，所有 thread 和 host 都能访问，大但远。local memory 也在 HBM 上。
Occupancy: SM 上活跃的 warp 数 ÷ 最多能放的 warp 数。warp 越多，一个在等 memory 时越容易换另一个上来算，latency 就藏住了。例：每 block 768 thread，一个 SM 最多 2048 thread → 只放得下 2 个 block → 1536/2048 = 75%。

看图：H100 的存储一层套一层

图 1：H100 从里到外的存储。橙红色的都在 SM 上（register、shared memory、L1），蓝色的 HBM 在芯片外，所有 SM 共用 L2。一个 block 整个放在一个 SM 上，所以 block 里的 thread 能共用这块 shared memory。

图 2：横轴是访问一次要等多少个 cycle（H800 上的 microbenchmark）。SM 上的几层都在几十 cycle 以内，L2 约 263，HBM 约 479，差一个数量级。左边灰字是容量，越远越大也越慢。

30 秒回答

From fastest to slowest: per-thread registers, 64K per SM. Each SM has 256 KB of on-chip SRAM on H100, split between L1 and shared memory, a programmer-managed scratchpad shared by a thread block. L2 is on chip and shared by all SMs, 50 MB. Global memory is off-chip HBM, 80 GB at about 3.35 TB/s, with hundreds of cycles of latency; local memory for register spills lives there too. A block has at most 1024 threads and runs on one SM, threads run in 32-wide warps, and an SM holds at most 2048 threads. Fast kernels reuse data in SRAM and registers and keep enough warps resident to hide latency.

从快到慢

层	谁能访问	在哪	H100（CC 9.0）大小	latency（cycles）
register	1 个 thread	SM 上	64K × 32-bit / SM	≈ 0 额外
shared memory	同一 block	SM 上 SRAM	≤ 228 KB / SM	≈ 29
L1 cache	同一 SM	同一块 SRAM	256 KB 减去 shared	≈ 41
L2 cache	所有 SM	片上	50 MB	≈ 263
global（HBM）	所有 thread + host	片外	80 GB，3.35 TB/s	≈ 479
local	1 个 thread	片外（HBM）	register spill	同 global

latency 来自 H800（Hopper）的 microbenchmark，NVIDIA 文档对 global memory 只说 "hundreds of clock cycles"。记数量级：SRAM 几十 cycle，L2 几百，HBM 约 500。

Thread 上限（CC 9.0，和 8.0 / 10.0 相同）

限制	值
threads / block	1024（x、y ≤ 1024，z ≤ 64）
warp size	32
resident threads / SM	2048 = 64 warps（CC 8.6、8.9 是 1536）
resident blocks / SM	32
registers	64K / SM，≤ 255 / thread

\[ \text{occupancy} = \frac{\text{active warps / SM}}{\text{max warps / SM}} \]

受 thread 数、register、shared memory 三样里最紧的那个限制，算例见 occupancy。

必答点

从快到慢：register → shared memory / L1（SM 上 SRAM）→ L2（全卡共享）→ HBM（global、local）
SRAM vs HBM：片上、每 SM 几百 KB、几十 cycle、程序员管 vs 片外、几十 GB、几百 cycle
block ≤ 1024 thread；warp = 32；H100 每 SM ≤ 2048 thread、≤ 32 block
occupancy = active warps / max warps，被 thread、register、shared memory 限制
优化：数据搬进 shared memory / register 复用（GEMM tiling、FlashAttention），保证足够多 warp 藏 latency

追问

"最多多少 thread" 到底指什么？

分三个数。一个 block ≤ 1024，一个 SM 同时驻留 ≤ 2048（H100），整张卡同时驻留 = SM 数 × 2048。一次 launch 的 grid 可以远大于此（x 方向最多 \(2^{31}-1\) 个 block），没轮到的 block 排队等 SM 空出来。

带宽大概多少？

H800 实测：shared memory 每 SM 每 cycle 128 bytes，按 114 个 SM、1755 MHz 算全卡约 26 TB/s；L2 约是 HBM 的 4 倍；HBM 实测约 1.9 TB/s（H100 SXM 标称 3.35 TB/s）。

occupancy 越高越好吗？

不一定。occupancy 只是为了藏 latency，超过某个点就不再变快（见 occupancy）。

shared memory 和 L1 怎么分？怎么用？

同一块 SRAM，比例（carveout）可配。L1 由硬件自动 cache，shared memory 要自己读写，block 内用 __syncthreads() 同步。一个 block 超过 48 KB 要用 dynamic shared memory 显式 opt-in。32 个 bank 的访问冲突见 bank conflict。

来源CUDA Programming Guide §1.2.2–1.2.3.3、§2.3.3、§2.3.7、§5.1.3 Table 30–32 · CUDA C++ Best Practices Guide §10.2、§10.2.7、§11.1、§11.3、§12.2 · Hopper Tuning Guide §1.4.1.1、§1.4.2.1–1.4.2.4 · Luo et al. 2024, Dissecting Hopper Table III–V（pp.6–7） · NVIDIA H100 产品页（H100 SXM：80 GB、3.35 TB/s）

GPU memory hierarchy 是怎样的？一个 block 最多多少 thread？