CUDA 基础概念 Q&A

Q1：CPU 与 GPU 的核心区别是什么？各自适合哪类任务？

CPU 和 GPU 的设计哲学从根本上就不同——CPU 追求的是把单条线程跑得尽可能快，GPU 追求的是同时把海量线程都跑起来。

CPU 擅长执行一条接一条的串行操作，可以并行运行的线程数通常只有几十个。GPU 则相反，它愿意牺牲单线程的执行速度，换取同时运行数万乃至数百万个线程的能力。这种设计差异直接体现在芯片的晶体管分配上：CPU 把大量晶体管用在数据缓存和流控制逻辑上，以减少单线程的延迟；GPU 则把更多晶体管堆在计算单元上，优先保证整体吞吐量。

正因如此，GPU 在相近的价格和功耗下，能提供远高于 CPU 的指令吞吐量和内存带宽。

核心对比：

维度 CPU GPU
设计目标 单线程串行执行速度 大规模并行吞吐量
并发线程数 数十个 数万至数百万
晶体管用途 缓存 + 流控制 计算单元（ALU）
单线程延迟 极低 较高
内存带宽 相对较低 极高

各自适合的任务：

  • CPU 擅长：逻辑复杂、分支密集、线程间依赖强的串行任务。比如操作系统调度、数据库事务处理、编译器、网络协议栈等。
  • GPU 擅长：能拆分为大量独立子任务的数据并行计算。比如深度学习训练与推理、图像/视频处理、科学仿真（流体动力学、分子动力学）、大规模矩阵运算等。

实际应用中，两者往往配合使用：CPU 负责主控逻辑，GPU 承担计算密集的并行部分，这也是 CUDA 的异构计算模型的核心思路。

Q2：什么是 CUDA？它与 OpenCL 的区别是什么？

CUDA 是什么？

CUDA（Compute Unified Device Architecture，统一计算设备架构）是 NVIDIA 于 2006 年推出的并行计算平台和编程模型。它的出现让 GPU 的计算能力彻底摆脱了图形 API 的束缚，任何计算任务都可以直接利用 GPU 的吞吐能力。

从那以后，CUDA 和 GPU 计算被用于加速几乎所有类型的计算任务——从流体动力学、能量传输等科学仿真，到数据库与商业分析，再到图像分类、扩散模型、大语言模型等 AI 应用。

CUDA 平台的核心组成：

  • CUDA C++：在标准 C++ 基础上扩展的 GPU 编程语言，是最常用的 CUDA 开发方式
  • NVCC：NVIDIA 的 CUDA 编译器，负责把包含 GPU 代码的源文件编译为可执行程序
  • CUDA Runtime API：提供内存分配、数据拷贝、内核启动等常用接口
  • CUDA Driver API：更底层的 GPU 控制接口
  • PTX 虚拟指令集：GPU 的中间表示层，实现跨架构的可移植性
  • 加速库：cuBLAS（线性代数）、cuFFT（傅里叶变换）、cuDNN（深度神经网络）、CUTLASS（高性能矩阵运算模板库）等

除了直接写 CUDA C++ 代码，还有更高层的方式使用 GPU，比如 NVIDIA Warp、OpenAI Triton 这类领域专用语言（DSL），它们最终也编译到 CUDA 平台上运行。

CUDA 与 OpenCL 的区别：

OpenCL（Open Computing Language）是由 Khronos Group 制定的开放标准，设计目标是跨厂商、跨硬件的通用并行计算接口，支持 NVIDIA/AMD/Intel GPU、多核 CPU、FPGA 等。

维度 CUDA OpenCL
开发者 NVIDIA（专有） Khronos Group（开放标准）
硬件支持 仅 NVIDIA GPU NVIDIA/AMD/Intel GPU、CPU、FPGA 等
生态系统 极其丰富（cuDNN、NCCL、Triton 等） 相对有限
编程体验 C++ 扩展，工具链完善，上手快 接近 C99 的内核语法，配置繁琐
性能 在 NVIDIA 硬件上通常更优 跨平台但通常有性能损耗
主要用途 AI/深度学习/HPC 的主流选择 跨厂商硬件可移植场景

如何选择？ 如果目标硬件是 NVIDIA GPU，CUDA 几乎是唯一正确答案——生态、工具链、性能都碾压其他方案。只有当你需要在 AMD/Intel GPU 上运行，或者客户硬件不确定时，才需要考虑 OpenCL 或 SYCL。

Q3：GPU 的硬件架构是怎样的？SM、SP、Warp 分别是什么？

GPU 的整体架构

从 CUDA 编程的视角来看，GPU 是若干个流式多处理器（SM，Streaming Multiprocessor）的集合，多个 SM 组成一个图形处理簇（GPC，Graphics Processing Cluster），若干 GPC 再加上共享的 L2 缓存和全局显存，构成整个 GPU。

1
2
3
4
5
6
7
8
9
10
11
GPU
├── GPC（图形处理簇）
│   ├── SM（流式多处理器）
│   │   ├── 寄存器文件（Register File）
│   │   ├── 统一数据缓存（= L1 Cache + Shared Memory，可配比）
│   │   ├── 功能单元（CUDA Core / SP，负责实际计算）
│   │   └── Warp 调度器（Warp Scheduler）
│   └── SM × N
└── GPC × N
├── L2 Cache（所有 SM 共享）
└── 全局显存 DRAM（Global Memory）

下图展示了 GPU 与 CPU 的整体架构对比：

GPU 通过 PCIe 或 NVLink 与 CPU 相连。CPU 负责主控逻辑，GPU 负责并行计算。

SM（Streaming Multiprocessor，流式多处理器）

SM 是 GPU 的基本计算单元。每个 SM 拥有独立的寄存器文件、共享内存/L1 缓存，以及一套完整的 Warp 调度器和功能单元。CUDA 中的一个线程块（Thread Block）只会被分配到某一个 SM 上执行，这保证了同一线程块内的线程可以高效地通过共享内存通信和同步。

一个 GPU 可以拥有数十乃至上百个 SM，不同架构代际的 SM 数量和每个 SM 的资源规格差异较大。

SP（Streaming Processor / CUDA Core，流处理器）

SP 是 SM 内部的单个计算核心，通常也称为 CUDA Core，负责执行具体的浮点或整数运算指令。一个 SM 包含若干个 SP，不同 GPU 每个 SM 的 SP 数量不同（例如 Ampere 架构的 SM 包含 128 个 FP32 CUDA Core）。从编程角度，开发者不需要直接管理 SP，只需管理线程——硬件自动将线程映射到 SP 上执行。

Warp

Warp 是 SM 调度和执行的基本单位，由 32 个线程组成。SM 不是逐线程调度指令，而是以 Warp 为粒度——每个调度周期，Warp 调度器挑选一个就绪的 Warp，向其中所有线程同时发射同一条指令。Warp 这个名字来自纺织业中的"经线"，代表第一种并行线程技术。

三者的关系总结：

概念 层级 作用
SM 硬件执行单元 承载线程块，包含计算资源和调度器
SP（CUDA Core） SM 内部的计算核心 执行具体的算术指令
Warp SM 的调度单位 32 个线程打包，统一发射指令

Q4：什么是 Warp？一个 Warp 有多少个线程？

Warp 的定义

在线程块内部，线程被按线程 ID 顺序每 32 个分为一组，这一组就叫一个 Warp。Warp 是 GPU 硬件调度和执行的最小单位。

一个 Warp 固定包含 32 个线程，这在所有当前支持的 NVIDIA GPU 架构中都是不变的常量（warpSize = 32）。

Warp 的执行机制

同一 Warp 内的 32 个线程同时执行同一条指令，但每个线程有自己的寄存器状态，可以操作不同的数据。这就是 SIMT（单指令多线程）模型的体现。

线程块被分配到 SM 后，SM 按以下规则将其拆分为 Warp：

  • 第 0 个 Warp 包含线程 ID 0~31
  • 第 1 个 Warp 包含线程 ID 32~63
  • 以此类推，连续线程 ID 构成一个 Warp

一个线程块包含的 Warp 数量 = ⌈线程总数 / 32⌉。

零开销的 Warp 切换

每个 Warp 的执行上下文（程序计数器、寄存器等）在其整个生命周期内都保存在片上，不需要像 CPU 那样做上下文保存/恢复。因此，Warp 之间的切换是零开销的。

这正是 GPU 隐藏内存延迟的核心手段：当一个 Warp 在等待全局内存数据返回时，调度器立刻切换到另一个就绪的 Warp 继续执行，让计算单元始终保持忙碌。

Warp 分歧（Warp Divergence）

如果 Warp 内的线程因为条件判断走了不同的代码路径，就会发生 Warp 分歧。此时 GPU 必须串行执行每条分支路径，并将不在该路径上的线程暂时屏蔽（mask off）。所有分支都执行完后，Warp 才重新汇合。

下图展示了这一情况——只有偶数 ID 线程执行 if 体，奇数 ID 线程被屏蔽等待：

分歧只会发生在同一个 Warp 内部；不同 Warp 之间的执行路径完全独立，互不影响。

最佳实践

  • 线程块的线程总数应设为 32 的整数倍。若不是，最后一个 Warp 中会有若干 lane 始终空闲，造成计算资源浪费。
  • 尽量让同一 Warp 内的线程走相同的代码路径，减少分歧，最大化 GPU 利用率。

Q5：CUDA 的线程层次结构是什么？Thread、Block、Grid 的关系？

启动一个 CUDA kernel 时，往往需要数百万个线程同时执行。为了管理这些线程，CUDA 将它们组织成三级层次：

1
2
3
Grid（网格）
└── Block（线程块）× 多个
    └── Thread（线程）× 多个（最多 1024）

Thread（线程）

线程是 CUDA 并行执行的最小逻辑单元。每个线程运行相同的 kernel 代码，但通过各自的索引（threadIdx、blockIdx）知道自己该处理哪部分数据，从而实现数据并行。

每个线程拥有独立的寄存器，在 CC 7.0 及以上的 GPU 上还有独立的程序计数器，可以在 Warp 内走不同的控制流路径。

Block（线程块）

线程被分组为线程块，同一线程块内的线程保证在同一个 SM 上并发执行，因此它们可以：

  • 通过**共享内存（Shared Memory）**高速交换数据
  • 使用 __syncthreads() 进行块内同步

线程块可以是 1、2 或 3 维的，每个块最多包含 1024 个线程。使用多维线程块只是为了方便映射数据结构（如 2D 图像、3D 体素），对性能没有影响。

Grid（网格）

所有线程块合在一起构成 Grid。一次 kernel 启动对应一个 Grid。Grid 同样可以是 1、2 或 3 维的，其中的线程块总数可以达到数百万。

不同线程块之间没有执行顺序保证，调度器可以以任意顺序、串行或并行地将它们分配到各个 SM。因此，正确的 CUDA 程序不能依赖跨块的执行顺序。这一约束反过来使 CUDA 程序天然具备可扩展性——同样的代码在 1 个 SM 和 1000 个 SM 的 GPU 上都能正确运行。

下图展示了 Grid 与 Block 的结构关系：

下图展示了线程块如何被分配到 SM：

关系总结：

层级 通信范围 硬件映射 数量限制
Thread 寄存器（私有） SM 内某个 lane —
Block 共享内存（块内） 整块调度到同一 SM 最多 1024 线程/块
Grid 全局内存（全局） 分散到所有 SM 理论可达数亿线程块

CC 9.0+ 的 Cluster（线程块簇）

从 Hopper 架构（CC 9.0）起，CUDA 引入了可选的第四个层级——Cluster，介于 Block 和 Grid 之间。同一 Cluster 内的线程块保证调度在同一个 GPC 上，线程之间可以访问彼此的共享内存（称为分布式共享内存），并进行簇级别的同步。这为需要更大线程组协作的算法提供了新的优化空间。

Q6：blockDim、gridDim、threadIdx、blockIdx 分别代表什么？

这四个是在 kernel 代码内部可以直接使用的内置变量，让每个线程知道"我是谁、我在哪、我该做什么"。

四个内置变量：

变量 含义 类型
threadIdx 当前线程在其所在 Block 内的索引 dim3（.x/.y/.z）
blockDim 线程块的维度（每个维度的线程数） dim3（.x/.y/.z）
blockIdx 当前线程块在 Grid 内的索引 dim3（.x/.y/.z）
gridDim Grid 的维度（每个维度的块数） dim3（.x/.y/.z）

所有变量均为三维向量（.x、.y、.z），未指定的维度默认为 1，索引均从 0 开始。threadIdx.x 的取值范围是 [0, blockDim.x - 1]，blockIdx.x 的取值范围是 [0, gridDim.x - 1]。

最常用的模式——1D 全局线程索引：

1
2
3
4
5
__global__ void vecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    C[i] = A[i] + B[i];
}

每个线程通过 threadIdx.x + blockDim.x * blockIdx.x 计算出自己在整个数组中的唯一位置，从而知道要处理哪个元素。这是 CUDA 1D 并行化的标准套路。

2D 示例——矩阵加法：

1
2
3
4
dim3 grid(16, 16);   // 16×16 个线程块
dim3 block(8, 8);    // 每块 8×8 个线程，共 64 个线程/块

MatAdd<<<grid, block>>>(A, B, C);

在 kernel 内部：

1
2
3
4
5
6
__global__ void MatAdd(float A[N][M], float B[N][M], float C[N][M])
{
    int row = threadIdx.y + blockDim.y * blockIdx.y;
    int col = threadIdx.x + blockDim.x * blockIdx.x;
    C[row][col] = A[row][col] + B[row][col];
}

多维线程块的线性化顺序

多维线程块在映射到 Warp 时，线程 ID 的线性化顺序是：x 变化最快，然后是 y，最后是 z。即连续的 threadIdx.x 对应连续的线程 ID，threadIdx.y 的步长为 blockDim.x，threadIdx.z 的步长为 blockDim.x × blockDim.y。这个顺序直接决定了哪些线程被分到同一个 Warp，对内存访问的合并（coalescing）有直接影响。

Q7：什么是 SIMT（单指令多线程）执行模型？

SIMT 的定义

SIMT（Single Instruction, Multiple Threads，单指令多线程）是 CUDA GPU 的核心执行模型。每个 SM 以 Warp 为单位调度执行：同一 Warp 内的 32 个线程在同一时刻执行同一条指令，但每个线程有自己独立的寄存器状态，可以操作不同的数据。

用一句话概括：程序逻辑是"每人一份"，指令发射是"集体行动"。

SIMT 与 SIMD 的区别

SIMT 常被拿来和 CPU 中的 SIMD（单指令多数据，如 x86 的 AVX、SSE）比较，两者都是"一条指令驱动多个执行单元"，但有本质区别：

维度 SIMD SIMT
编程视角 向量操作，程序员需显式管理向量宽度 标量线程，程序员按单线程思维编写
数据宽度 固定（如 256 位、512 位） 无固定宽度
分支处理 需手动管理掩码或向量化 硬件自动处理线程屏蔽
编程复杂度 较高 较低（对程序员透明）

SIMT 让程序员可以按照普通标量线程的方式思考和编写代码，硬件负责把 32 个线程"打包"成一个 Warp 高效执行——编程友好性大幅提升。

Warp 分歧与串行执行

SIMT 的代价是分歧问题。当 Warp 内线程因条件分支走向不同路径时，GPU 必须串行执行每条路径，非活跃线程被屏蔽掉，直到所有路径执行完毕 Warp 才重新汇合。极端情况下，如果 32 个线程走了 32 条不同路径，性能下降至 1/32。

分歧只影响 Warp 内部；不同 Warp 之间完全独立，即使它们走不同代码路径也不会互相拖累。

独立线程调度（CC ≥ 7.0，Volta 架构起）

在 CC 7.0 之前，Warp 内所有线程共享一个程序计数器，线程必须严格锁步执行，这导致一些需要细粒度线程间通信的算法（如依赖锁或互斥量的算法）可能发生死锁。

从 Volta 架构（CC 7.0）起，GPU 为每个线程维护独立的程序计数器和调用栈，可以在子 Warp 粒度上进行调度，让线程在分歧后以更细的粒度重聚，提供了更大的灵活性。不过这也意味着：之前依赖"Warp 内线程一定锁步"这一隐式假设写的代码，在 CC 7.0+ 的 GPU 上可能行为改变，需要显式加 __syncwarp() 来保证正确性。

实践建议

  • 从正确性角度，可以完全按标量线程的方式写 CUDA 代码，SIMT 对你透明。
  • 从性能角度，要尽量让同一 Warp 内的线程走相同的代码路径，避免分歧。这就像 CPU 缓存行——不了解它程序也能跑，但了解它才能写出高性能代码。

Q8：GPU 的 Occupancy（占用率）是什么？如何计算？

占用率的定义

Occupancy（占用率）是衡量 SM 利用率的核心指标，定义为：

Occupancy=SM 上当前活跃的 Warp 数该 SM 支持的最大活跃 Warp 数\text{Occupancy} = \frac{\text{SM 上当前活跃的 Warp 数}}{\text{该 SM 支持的最大活跃 Warp 数}}

占用率越高，SM 上等待执行的 Warp 越多，调度器在某个 Warp 等待内存时就有更多选择可以切换，从而更好地隐藏内存访问延迟，保持计算单元持续忙碌。

限制占用率的三类资源

每个 SM 的资源是有限的，线程块上 SM 需要消耗这些资源：

资源 说明
寄存器 每个线程用的寄存器越多，SM 能同时容纳的线程（Warp）就越少
共享内存 每个线程块申请的共享内存越多，SM 能同时容纳的线程块就越少
线程块数上限 SM 对同时驻留的线程块数量有硬件上限

通过 cudaGetDeviceProperties() 可以查询 SM 的各项资源限制：

1
2
3
4
5
6
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
// prop.maxThreadsPerMultiProcessor  -- SM 最大线程数
// prop.maxBlocksPerMultiprocessor   -- SM 最大线程块数
// prop.regsPerMultiprocessor        -- SM 寄存器总数
// prop.sharedMemPerMultiprocessor   -- SM 共享内存总量

计算示例

以 CC 10.0 的 GPU 为例，SM 资源参数如下：

资源 值
SM 最大线程数（maxThreadsPerMultiProcessor） 2048
SM 最大线程块数（maxBlocksPerMultiProcessor） 32
SM 共享内存（sharedMemPerMultiprocessor） 228 KB

场景 1：768 线程/块

  • ⌊2048 / 768⌋ = 2，每个 SM 最多放 2 个块
  • 活跃线程数 = 768 × 2 = 1536
  • Occupancy = 1536 / 2048 = 75%

场景 2：32 线程/块

  • 按线程数算可放 2048 / 32 = 64 个块，但 maxBlocksPerMultiProcessor = 32，被此限制卡住
  • 活跃线程数 = 32 × 32 = 1024
  • Occupancy = 1024 / 2048 = 50%（受线程块数上限约束）

场景 3：每块使用 100 KB 共享内存

  • ⌊228 / 100⌋ = 2，SM 最多放 2 个块（第 3 块需要 300 KB，超出）
  • Occupancy 受共享内存约束

如何优化占用率

  1. 用 nvcc --resource-usage 编译时查看 kernel 的寄存器和共享内存用量
  2. 调整线程块大小（blockDim）找到最优配置
  3. 用 --maxrregcount 限制每线程最大寄存器数（会触发寄存器溢出到本地内存，需权衡）
  4. 用 cudaOccupancyMaxActiveBlocksPerMultiprocessor() API 自动计算最优块大小

占用率不是越高越好

高占用率是手段而非目的。有时 75% 占用率的 kernel 比 100% 的更快——比如每个线程使用大量寄存器、缓存命中率高的 kernel，强行提高占用率反而会因寄存器溢出而损害性能。应以实际 profiling 数据为准，而非盲目追求 100% 占用率。

Q9：什么是 Compute Capability？不同版本有何差异？

Compute Capability 是什么

每款 NVIDIA GPU 都有一个 **Compute Capability（CC，计算能力）**版本号，格式为 X.Y（主版本号.次版本号）。它有两个作用：

  1. 标识该 GPU 支持哪些硬件特性
  2. 规定该 GPU 的具体硬件参数（SM 资源规格等）

CC 直接对应 SM 的版本。例如 CC 12.0 的 GPU，其 SM 版本为 sm_120，编译时通过 -arch=sm_120 或 -arch=compute_120 来指定目标。

各主要版本的关键特性

CC 版本 代表架构（GPU 系列） 关键新特性
7.5 Turing（RTX 20 系列） 每 SM 最多 32 个活跃 Warp，1024 个线程
8.0 Ampere（A100） 每 SM 最多 64 个 Warp，2048 个线程；硬件加速异步内存拷贝和屏障；FP64/BF16 Tensor Core
8.6 Ampere（RTX 30 系列） 每 SM 最多 48 个 Warp，1536 个线程；FP8/TF32 Tensor Core
9.0 Hopper（H100） Thread Block Cluster；张量内存加速器（TMA）；FP8 Tensor Core；异步事务屏障
10.0 Blackwell（B100/B200） FP6/FP4 Tensor Core；更大共享内存和 Warp 上限
12.x 最新架构 持续迭代，详见附录

各版本 SM 资源对比

CC 最大 Warp/SM 最大线程/SM SM 最大共享内存
7.5 32 1024 64 KB
8.0 64 2048 164 KB
8.6 48 1536 100 KB
9.0 64 2048 228 KB
10.x — — 228 KB
12.x — — 100 KB

二进制兼容性规则

  • 同主版本内（如 8.0 → 8.6）：为低版本编译的 cubin 可以在高版本 GPU 上运行（向前兼容），但无法使用新架构特性。
  • 跨主版本（如 8.x → 9.0）：不兼容，必须重新编译。针对 CC 8.6 编译的 cubin 无法在 CC 9.0 的 GPU 上加载。
  • PTX 代码：PTX 是 GPU 的虚拟指令集中间层，可以在运行时被 JIT 编译为目标 GPU 的 cubin，是实现"一次编译、未来 GPU 兼容"的机制。Fatbin（胖二进制）打包了多个 cubin 和 PTX，运行时自动选择最合适的版本。

查询当前 GPU 的 Compute Capability

1
2
3
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);  // 0 为设备编号
printf("Compute Capability: %d.%d\n", prop.major, prop.minor);

也可以直接用命令行工具：

1
nvidia-smi --query-gpu=compute_cap --format=csv

编写需要特定架构特性的代码时，应在运行时检查 CC 版本，以保证在不支持的 GPU 上给出清晰的错误提示，而不是产生未定义行为。