CUDA 基础概念测验

Q1. 请描述 GPU 的内存层级结构。 从 Global Memory → Shared Memory → Register，数据是如何流动的？各层的容量、带宽、延迟大致是什么量级？

内存层级结构：Global Mem，L2 cache，L1 / Shared Mem，Register。

数据流动：Global Mem 到 L1 / Shared Mem（通过 L2 cache），计算的时候从 L1 / Shared Mem 取出。

越靠近计算单元，容量越小，带宽越高，延迟越低。

（以H200为例）
层级 容量 带宽 延迟
Global Mem 141G 4.815TB/s（官方数据）

3627.52GB/s （实测）		 622 cycles
L2 60M 5704.74 GB/s 270 cycles
L1 / Shared Mem 每 SM 228KB 21772.26 GB/s 30 cycles
Register 每 SM 65536 个 32-bit 寄存器 （256KB）

每个 thread 至多 255 个 32-bit 寄存器		 29.33 TFLOP/s

(计算****吞吐量 而非 带宽）		 1 cycle

Q2. 什么是合并访问（Coalesced Access）？ 为什么它对带宽利用率至关重要？

当一个 warp 的 32 个线程同时发起全局内存访问时，内存控制器会检查这些地址的分布模式。如果地址落在同一个缓存行/内存段内，硬件会将多个请求合并为一次（或少数几次）内存事务，而不是发起 32 次独立请求。

合并访问可以让访问接近理论带宽，非合并性能会退化。

Q3. 什么是向量化加载？ 对于 float32 数据，应该使用什么指令来保证单指令周期搬运更多数据？

单条指令加载多个元素，减少指令开销。

LDG128指令，加载128位的数据，代码层面可以是float4，int4，double2。

Q4. 什么是 Bank Conflict（存储体冲突）****？

  • Shared Memory 有多少个 bank？每个 bank 的宽度是多少？

32（和 warp 的线程数量一致），每个 bank 的宽度是 4 字节。

  • 如果一个 Warp 内的 32 个线程同时访问同一个 bank 的不同地址，会发生什么？

（如果是完全相同的地址那么是 broadcast，不算冲突）

不同地址会发生串行，性能变为 1/32。

  • 如何通过 Padding 避免 Bank Conflict？加多少 Padding？

让访问的 index%32 尽可能不相同。

Padding 一般多加一个就可以让跨行同列的映射到不同的 bank。

另外好像发现 LDG128 对地址有特殊要求，好像4路冲突是不可避免的。(无论是padding还是swizzle）

补充：swizzle 大概是 (row,col) 映射到 (row,col^(row&M-1))，在 a[N][M] 的 M>=32 的时候可以使用，当 M 较小的时候感觉还是 padding 比较对。

Cutlass swizzle 实践

针对 4096*4096 的矩阵做的如下实验：
Kernel Shared Mem Write Read Bank Conflict
copy (baseline) tile[32][32] tile[ty][tx] tile[ty][tx] 无（同行同列）
transposeNaive tile[32][32] tile[ty][tx] tile[tx][ty] 32-way
transposePadded tile[32][32+1] tile[ty][tx] tile[tx][ty] 无
transposeSwizzle tile[32][32] tile[ty][tx ^ ty] tile[tx][ty ^ tx] 无
td {white-space:nowrap;border:0.5pt solid #dee0e3;font-size:10pt;font-style:normal;font-weight:normal;vertical-align:middle;word-break:normal;word-wrap:normal;}
Kernel ld conflicts st conflicts
transposeNaive 16,421,810 24,364
transposePadded 43,209 24,758
transposeSwizzle 44,608 24,034
transposeSwizzleCutlass 41,651 23,068 
copy (baseline)                 0.036 ms   3696.84 GB/s  PASS
transpose naive                 0.092 ms   1459.12 GB/s  PASS
transpose padded                0.039 ms   3475.59 GB/s  PASS
transpose swizzle               0.040 ms   3370.90 GB/s  PASS
transpose swizzle (CUTLASS)     0.040 ms   3373.39 GB/s  PASS

Q5. 什么是 Memory Bound 和 Compute Bound？ 如果一个 Kernel profiling 显示计算单元利用率低但显存带宽打满，这属于哪种 Bound？应该怎么优化？

Mem Bound 是数据搬运速度受限，计算资源空余。

Compute Bound 是数据就绪但是计算资源受限。

Mem Bound

  1. 从输入输出角度（看能否调整算法降低访存量）

kernel 合并 / 算子融合、提高计算访存比、减少冗余加载、低精度（FP8/FP16/INT4）、Flash Attention、重计算、稀疏存储、在线压缩。

  1. 从 cache 角度，在访存中间是否可以加cache（指任何比 GMEM 快的存储）

合并访问、128B 对齐、float4 向量化、避免 bank conflict、shared memory tile 复用、Cluster + cluster-local shared、__ldg 只读缓存、L2 友好策略、寄存器缓存。

  1. 从并行角度，overlap 访存耗时

Async MemCopy + 流并发、cp.async、多级流水线（double/triple buffering）、Grid/stream 级流水线、固定内存 + 预取、CUDA Graph + 流水线

Q6. 请****描述 CUDA 的线程组织层级。 Grid、Block、Thread、Warp 之间是什么关系？硬件调度的最小单位是什么？一个 Warp 包含多少个线程？

Grid 包含全部的 Block，Block 里面有 N 个 Thread，每 32 个 Thread 构成一个 Warp。

硬件调度的最小单位是 Warp。

一个 Warp 包含 32 个线程。

每个 SM 最大 Warp 数量：64。

（cudaDeviceProp 中查询 maxThreadsPerMultiProcessor: 2048，warpSize: 32）

每个 SM 最大线程块数：32。

Q7. 什么是 Warp Divergence（线程束分歧）？ 它为什么会降低性能？请举一个会导致 Warp Divergence 的代码例子。

同一 Warp 内的线程遇到条件分支走了不同的执行路径，SIMT 架构要求 Warp 内所有线程执行相同指令，不同的指令会被 mask 掉， ALU 空转但是占着位置，最坏情况执行时间是所有分支时间的总和（而非最长分支）。

例如按照 threadidx 奇偶分支，奇偶走不同的路径，利用率降低到 50%。

可以把分支粒度提升到 warp 级，比如用 threadIdx.x / 32 做判断，保证每个 Warp 的每个线程走相同分支。

Q8. 什么是 Occupancy（占用率）？ 它受哪些因素限制？（关键词：寄存器数量、Shared Memory 大小、Block 内线程数） Occupancy 越高性能一定越好吗？为什么？

Occupancy 是 SM 上实际活跃的 Warp 数量 / 支持的最大 Warp 数量。

  • 寄存器数量

单个 SM 内寄存器总量固定，如果每个线程用足够多个寄存器，那么线程数量就有一个上限，影响占用率。

  • Shared Mem

单个 SM 内 Shared Mem 有限，如果 Block 用的多，那么同时驻留的 Block 数少，线程数就少。

  • Block 内线程数

每个 SM 内最多驻留的 Block 数数量有限，Block 内线程数会影响单个 SM 内的线程数量。

不一定。

高 Occupancy 的好处：有足够的 warp 可以切换，非空闲来隐藏内存延迟。

有时低 Occupancy 会更快：矩阵乘法使用更大的 Tile （使用了更多的 Shared Mem），Occupancy 低但是性能最优。

Q9. 什么是寄存器溢出（Register Spill）？ 溢出到哪里？对性能有什么影响？

线程所需寄存器超过分配给他的数量。

会溢出到 Local Mem（实际上是 Global Mem）。

性能严重退化：增加内存带宽压力，增加指令数。

Q10. 什么是 Roofline Model？什么是算术强度（Arithmetic Intensity）？

分析 kernel 性能瓶颈的模型，用于判断是 Mem Bound 还是 Compute Bound。

算术强度是 浮点计算次数 / 内存搬运字节数量。

Q11. 如何检测和控制寄存器压力？

实践记录：

简单的代码一般编译器都会优化，以矩阵乘法为例。

开发机的环境没办法直接使用 ncu（ncu 版本比驱动版本低），使用一个 docker 镜像。

sudo docker run --rm --gpus all --privileged -it -v $(pwd):/workspace -w /workspace nvcr.io/nvidia/pytorch:25.06-py3 bash

nvcc --ptxas-options=-v -arch=sm_90 --maxrregcount=24 -o reg_test_24 reg_test.cu
nvcc --ptxas-options=-v -arch=sm_90 --maxrregcount=32 -o reg_test_32 reg_test.cu
nvcc --ptxas-options=-v -arch=sm_90 --maxrregcount=40 -o reg_test_40 reg_test.cu

ncu --metrics launch__registers_per_thread ./reg_test_24
ncu --metrics launch__registers_per_thread ./reg_test_32
ncu --metrics launch__registers_per_thread ./reg_test_40

// 不限制寄存器
nvcc --ptxas-options=-v -arch=sm_90 -o reg_test reg_test.cu
ncu --metrics launch__registers_per_thread ./reg_test

实验结果


    296 bytes stack frame, 560 bytes spill stores, 560 bytes spill loads
ptxas info    : Used 24 registers, used 1 barriers, 296 bytes cumulative stack size, 8192 bytes smem

    256 bytes stack frame, 492 bytes spill stores, 492 bytes spill loads
ptxas info    : Used 32 registers, used 1 barriers, 256 bytes cumulative stack size, 8192 bytes smem

    216 bytes stack frame, 396 bytes spill stores, 396 bytes spill loads
ptxas info    : Used 40 registers, used 1 barriers, 216 bytes cumulative stack size, 8192 bytes smem

    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 90 registers, used 1 barriers, 8192 bytes smem


==PROF== Connected to process 1439 (/workspace/reg_test_24)
==PROF== Profiling "MatMulKernel" - 0: 0%....50%....100% - 1 pass
GPU time: 2038.68 ms, Throughput: 0.00105337 TFLOP/s

==PROF== Connected to process 1457 (/workspace/reg_test_32)
==PROF== Profiling "MatMulKernel" - 0: 0%....50%....100% - 1 pass
GPU time: 2897.33 ms, Throughput: 0.000741195 TFLOP/s

==PROF== Connected to process 1475 (/workspace/reg_test_40)
==PROF== Profiling "MatMulKernel" - 0: 0%....50%....100% - 1 pass
GPU time: 1160.4 ms, Throughput: 0.00185065 TFLOP/s

==PROF== Connected to process 1526 (/workspace/reg_test)
==PROF== Profiling "MatMulKernel" - 0: 0%....50%....100% - 1 pass
GPU time: 1210.95 ms, Throughput: 0.00177338 TFLOP/s


// 经过测试好像是 48 最优。
==PROF== Connected to process 1628 (/workspace/reg_test_48)
==PROF== Profiling "MatMulKernel" - 0: 0%....50%....100% - 1 pass
GPU time: 933.311 ms, Throughput: 0.00230093 TFLOP/s

并不一定是 0 spill / 寄存器数量最少 最好，大概要 trade-off 一下。

Q12. 什么是 Tensor Core？它和普通 CUDA Core 的区别是什么？

Tensor Core 每个时钟周期执行一个小矩阵乘加（数百次 FMA / Tensor Core。）

CUDA Core 每个时钟周期执行一次 FMA（Fused Multiply-Add）（一次 FMA / Core）

Tensor Core 使用的编程接口是 WMMA API / MMA PTX / cuBLAS/cuDNN（CUDA Core 是普通的 * +），但是支持的精度是 FP16, BF16, TF32, FP8, INT8, INT4 （而 CUDA Core 是 FP64, FP32, INT32）。

Q13. 简述 Volta → Ampere → Hopper → Blackwell 架构的关键变化。

AIGC。

NVIDIA的数据中心GPU架构从Volta到Blackwell，其核心演进逻辑是不断突破“内存墙”瓶颈，并针对Transformer类大模型进行极致优化。以下将简述各架构的关键变化，并整理详细参数对比。

  • Volta（伏特，2017）- AI****的起点：首次引入第一代Tensor Core，开创了混合精度计算时代，专门为深度学习训练加速。

  • Ampere（安培，2020）- 多精度与弹性：Tensor Core演进至第三代，支持创新的TF32和BF16精度，并引入了结构化稀疏与MIG（多实例GPU） 功能，大幅提升通用性与利用率。

  • Hopper（赫柏，2022）- Transformer引擎：为大模型而生，新增Transformer引擎，是首款支持FP8精度的架构，并引入了TMA（张量内存加速器）来优化数据搬运。

  • Blackwell（布莱克韦尔，2024）- 万亿参数时代：引入第五代Tensor Core及FP4/FP6精度，支持更庞大的模型。通过NVLink-C2C技术将两颗芯片互联，并配备了速度更快、容量更大的HBM3e内存。

以下是针对四代架构旗舰数据中心GPU的参数汇总。由于官方数据会因配置（如SXM vs PCIe）不同而有细微差异，下表取峰值理论性能数据。
指标维度 Volta (V100) Ampere (A100) Hopper (H100) Blackwell (B200)
SM（流式多处理器）数量 80 108 132 208 (基于B200)
显存大小 32 GB HBM2 80 GB HBM2e 80 GB HBM3 / 141 GB HBM3e (H200) 192 GB HBM3e
Tensor Core 算力 (FP16/FP8) 125 TFLOPS (FP16) 312 TFLOPS (FP16)
支持稀疏性		 1,979 TFLOPS (FP8)
约 1,000 TFLOPS (FP16)		 20,000 TFLOPS (FP4)
10,000 TFLOPS (FP8)
显存带宽 900 GB/s 2.0 TB/s (2,000 GB/s) 3.35 TB/s (H100)
4.8 TB/s (H200)		 8.0 TB/s

关于Tensor Core算力的说明：各代架构支持的精算精度不同，为了直观体现算力代际，上表选用了该架构主打的精度进行对比。例如，Volta主打FP16，而Hopper和Blackwell分别主打FP8和FP4精度，性能呈现出跨越式增长。

posted @ 2026-04-03 03:41  EternalEpic  阅读(6)  评论(0)    收藏  举报