【笔记】理解 CUDA 底层概念

1. 示意图

image

2. 详解

2.1 基础概念

（1）一个流（CUDA Stream）代表一系列的操作序列，这些操作包括：核函数启动、内存拷贝、事件同步等。

（2）一个 Grid 是一个核函数在 GPU 上执行的所有线程的集合。一个核函数等于一个 Grid。

（3）一个 Grid 可以拆分为多个 block，多个 block 可以分散到不同的 SM 并行执行，block 之间没有依赖关系，不能直接通信。

（4）一个 block 是指多个线程组成的线程块，一个 warp 是指 32 个线程组成的线程束，SM 调度以 warp 为基本单位。

（5）一个 SM 包含多个 warp，warp 是执行单元，SM 内有多个 warp 调度器，调度器负责将 warp 里的指令发射到具体的计算硬件（如：Tensor Core、CUDA Core、FP32 Core、FP64 Core）。多个 warp 调度器可以同时调度，且同一类计算硬件有多套，不同的计算硬件可以同时使用，所以一个 SM 内多个 warp 可以并行执行。一个 SM 最多包含多少个 warp 由硬件层面限制，这些硬件包括：warp 上下文存储、寄存器数量、共享内存容量。

（6）block、warp、SM 的关系：一个 block 只能被一个 SM 执行，而一个 SM 可以执行多个 block。一个 SM 内包含多个 warp，因一个 block 只能属于一个 SM，因此一个 block 最大时，是刚好能把 SM 的所有 warp 都用上。另外，block 内线程数还受限于硬件，譬如：有些硬件规定一个 block 最多 1024 个线程，同时每个 block 的线程数还需要能给 warp_size（即 32）整除。

（7）GPC 是 SM 的硬件载体，一个 GPU 有多个 GPC，一个 GPC 承载多个 SM。但 GPC 不参与调度，GPU 的全局调度器（如 NVIDIA 的 GigaThread 调度器）直接将一个核函数调度到 SM 上。

2.2 block（线程块）

（1）一个 block 内的多个 warp 有共同的：

  • 共享内存，这部分属于线程间通信
  • 线程块的全局 ID 和局部 ID，每个 block 都有唯一的 blockIdx，每个线程在 block 内有私有的 threadIdx
  • 同步屏障（Barrier Synchronization），同一个 block 内的所有 warp 共享，确保 block 内的所有线程完成指定操作后继续执行，当一个 warp 执行到 __syncthreads() 时，会等待 Block 内其他所有 warp 的线程到达统一屏障点，才能继续执行后续指令
  • 常量内存（Constant Memory），GPU 上的只读全局资源，在 block 内共享，硬件会通过缓存机制（常量缓存）为整个 block 的线程共享该数据，减少重复访问开销
  • 纹理内存（Texture Memory），一种优化图像/数组访问的只读资源。

（2）block 存在的意义：是为了将一批 warp 放到一起，制定它们资源共享和同步的边界。

（3）一个 block 有多少个 warp 需要考虑多个因素，不能太多也不能太少：

  • 硬件限制

    • GPU 硬件限制一个 SM 最多有 N 个 warp 槽位，譬如 A100 限制为 64 个 warp 槽位，而 block 需要被 SM 加载，因此一个 block 最大只能有 64 个 warp。通常达不到 64 个 warp，因为还有其他硬件：共享内存、寄存器的限制。
    • 多个 block 时，即使总寄存器需求超过了一个 SM 能提供的总量，但若每个 block 的需求小于 SM 总量，SM 可以分批调度（例如先调度 2 个 block，完成后再调度剩余的），避免因单个 block “过大” 而无法调度。

  • 性能要求

    • 一个 block 所包含的 warp 所执行的代码通常是相同的，如果一个 block 独占了一个 SM，那么这个 SM 激活的 warp 更可能会因为相同的全局内存访问而挂起，导致这个 SM 的硬件占用率不高。因此尽量让一个 SM 可以装下多个 block。注：一个 block 所包含的 warp 如果有差异逻辑，则需要复杂的控制流和数据结构，会使代码难以维护。
    • SM 为每个 block 分配的共享内存，有最小限制，譬如：8KB，如果一个 block 只有一个 warp，会存在共享内存浪费。
    • 一个 block 里的线程共享 SM 为 block 分配的寄存器配额，如果用超了，则会导致寄存器溢出到全局内存，性能下降。
    • 单个 block 若因线程数不是 “寄存器分配粒度的整数倍”，可能导致部分寄存器被闲置（例如，SM 按 256 寄存器块分配，若单个 block 需求为 300 寄存器，则实际分配 512，浪费 212）。

2.3 warp（线程束）

（1）SIMT

一个 warp 内的多有线程执行相同的指令，但操作不同的数据（单指令多线程，SIMT）。如果 warp 内的线程执行不同的分支，则 warp 会串行执行所有分支，导致性能下降。SIMT 使用一套指令控制多个计算单元，是 GPU 对比 CPU 的最大不同点，如果 SIMT 里每个线程都可以并行执行不同的指令，则变成了 CPU。

（2）线程

线程代表一段独立的指令流和数据上下文。这点和 CPU 的线程很像，但 GPU 线程的寄存器在 warp 因内存访问被挂起的时候，寄存器不会让出。这使得 warp 切换时的消耗非常低。当寄存器不足时，编译器会将部分变量存储到全局内存（成为 spilling），导致性能下降。

（3）warp 划分规则

CUDA 硬件将每个 Block 内的线程自动划分为固定大小的 warp（32 线程/warp），划分规则仅基于线程 ID。

（4）warp调度器

warp 调度器是 SM 粒度的，block 中的线程以 warp 为单位被 SM 中的 warp 调度器调度执行。不同架构的 SM 中的 warp 调度器数量不同，譬如：Fermi 架构的 SM 包含 2 个 warp 调度器、Hopper 架构的 SM 包含 4 个调度器。 调度是异步的，将指令发射到计算单元之后，不需要等计算完成就可以继续下一次发射。

（5）挂起和就绪

如果当前 warp 如果无法继续（譬如：等待内存数据）则会挂起，warp 调度器会从就绪队列里取出 ready 的 warp 做执行，当前 warp 数据获取完成后会放到就绪队列等待执行。

2.4 存储的访问速度

寄存器 > 共享内存/L1缓存 > L2缓存 > 全局内存/常量内存/纹理内存 > 主机内存（CPU内存）。

L1 缓存在 SM 上，一般比较小，譬如：48KB。

L2 缓存在 GPU 上，所有 SM 共享，譬如：4MB。

2.5 流、grid 和代码的关联

TL;DR: 一个核函数的启动对应一个 Grid 的执行。一个流可以顺序地承载多个 Grid 的执行。不同的流可以并发地执行各自的 Grid。

（1）代码和 grid

vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, VECTOR_SIZE);

这个语法是 启动了一个 CUDA 核函数，它启动了一个 Grid。blocksPerGrid 指定了这个 Grid 包含的块（Blocks）的数量，每个块中又有 threadsPerBlock 数量的线程（Threads）。

（2）默认流

Stream（流） 在 CUDA 中是一个独立的执行队列。它允许你在 GPU 上进行并发操作。默认情况下，如果你不明确指定流，所有的 CUDA 操作（包括核函数启动、内存拷贝等）都会在 默认流 (null stream) 中执行。默认流是一个隐式的、同步的流，意味着其中的操作是按顺序执行的，并且在某个点上可能会隐式同步（例如，所有在默认流中的操作在任何显式同步命令之前都必须完成）。

如果你想在一个非默认流中启动核函数，你需要这样写：

cudaStream_t stream1;
cudaStreamCreate(&stream1); // 创建一个流

// ... 内存拷贝等操作 ...

vectorAdd<<<blocksPerGrid, threadsPerBlock, 0, stream1>>>(d_a, d_b, d_c, VECTOR_SIZE); // 在 stream1 中启动核函数

// ... 其他操作或销毁流 ...
cudaStreamDestroy(stream1);

在这个例子中：

  • 0 是动态共享内存的字节数，这里设置为 0，表示不使用。
  • stream1 是你明确指定的 CUDA 流。

2.6 获取设备信息的代码

2.6.1 一个 SM 最多可以驻留多少个 warp

一个 SM 最多可以驻留多少个 warp

import pycuda.driver as cuda
import pycuda.autoinit

dev = cuda.Device(0)  # 第一个GPU
max_warps_per_sm = dev.get_attribute(cuda.device_attribute.MAX_THREADS_PER_MULTIPROCESSOR) // 32
print(f"每个SM最多驻留 {max_warps_per_sm} 个warp")

# L20：48 个 warp
# H20: 64 个 warp

2.6.2 一个 GPU 有多少个 SM

一个 GPU 最多有多少个 SM

import pycuda.driver as cuda
import pycuda.autoinit

try:
    dev = cuda.Device(0) # 假设L20是你的第一个GPU
    total_sm_count = dev.get_attribute(cuda.device_attribute.MULTIPROCESSOR_COUNT)
    print(f"你的 GPU ({dev.name()}) 共有 {total_sm_count} 个 SMs。")
except Exception as e:
    print(f"An error occurred: {e}")

# L20 92 个
# H20 78 个

2.6.3 获取 CUDA GPU 计算能力

计算能力（CUDA GPU Compute Capability）的官方定义：Compute capability (CC) defines the hardware features and supported instructions for each NVIDIA GPU architecture.

获取 CUDA GPU Compute Capability

import pycuda.driver as cuda
import pycuda.autoinit

dev = cuda.Device(0)
cc_major = dev.get_attribute(cuda.device_attribute.COMPUTE_CAPABILITY_MAJOR)
cc_minor = dev.get_attribute(cuda.device_attribute.COMPUTE_CAPABILITY_MINOR)

print(f"计算能力: {cc_major}.{cc_minor}")

# L20 8.9
# H20 9.0

官方文档：https://developer.nvidia.com/cuda-gpus

posted on 2026-09-16 18:26  -银光-  阅读(23)  评论(0)    收藏  举报