GPU结构与线程调度

1 GPU 的硬件组成

一个 GPU 芯片由多个 SM（Streaming Multiprocessor，流式多处理器）组成，并配备显存、缓存等硬件资源，以支持大规模并行计算和高效的数据传输。
每个 SM 内部包含 Warp 调度器、多个 CUDA Core、多个 Tensor Core（具体数量取决于 GPU 架构）、寄存器文件、共享内存等硬件资源，用于调度和执行并行计算任务。

如下图所示：

2 GPU 是如何组织线程的？

CUDA 编程模型中，线程的组织层级如下：

注意：

• 一个 Block 只能分配到一个 SM 上执行，不能拆分到多个 SM。
• 一个 SM 可以同时驻留多个 Block，前提是硬件资源足够。
• Warp 是 SM 内部进行线程指令调度的重要单位。

3 线程究竟负责哪部分数据？

GPU 硬件负责调度线程，而线程处理哪份数据，通常由程序员编写的索引公式决定。

假设向量有 1024 个元素：

我们要执行计算，让向量中的每个元素都+1。

3.1 开发自定义线程组织方式

add_one<<<4, 256>>>(x);

表示：

• Grid 包含 4 个 Block。
• 每个 Block 有 256 个线程。
• 总计启动 1024 个线程。

3.2 每个线程计算自己的数据下标

int i = blockIdx.x * blockDim.x + threadIdx.x;
x[i] = x[i] + 1;

其中：

变量	含义
blockIdx.x	当前 Block 的编号
blockDim.x	每个 Block 的线程数
threadIdx.x	当前线程在 Block 内的编号
i	当前线程负责的全局数据下标

示例如下：

Block 编号	线程编号	负责的数据
0	0	x[0]
0	1	x[1]
0	255	x[255]
1	0	x[256]
1	255	x[511]
2	0	x[512]
2	255	x[767]
3	0	x[768]
3	255	x[1023]

这样，要启动的 1024 个线程就分别知道自己要处理哪个元素。

4 GPU 如何真正执行这些线程？

完整过程可以理解为：

1. 启动 Kernel
   1. 程序员指定 Grid 和 Block 的大小。
2. 分配 Block
   1. GPU 将 Block 分配到具有足够资源的 SM。
3. 组织 Warp
   1. Block 内的线程按 Warp 分组，NVIDIA GPU 通常每组 32 个线程。
4. 调度 Warp
   1. SM 的 Warp 调度器选择可执行的 Warp，向相应的执行管线发射指令。
5. 执行计算
   1. 线程根据自己的索引定位数据，并执行 x[i] + 1。
6. 写入结果
   1. 将结果写入目标内存，相关计算完成后，整个向量的操作完成。

需要注意，线程数量不等于同时执行的线程数量。GPU 可以分批执行 Block，并在多个 Warp 之间切换调度，以提高硬件利用率。

5 和 CPU 线程的区别

对比	CPU	GPU（CUDA）
线程	OS Thread	CUDA Thread
主要调度机制	OS 调度器	GPU 硬件调度器
调度特点	通常围绕单个软件线程调度	通常以 Warp 为单位发射指令
并行组织	CPU 核心执行线程	Grid → Block → Warp → Thread
数据由谁分工	软件逻辑、线程任务划分	Kernel 中的索引逻辑通常决定