1. CUDA 执行模型解决什么问题CUDA 编程模型告诉我们如何使用 Thread、Block 和 Grid 描述一个并行任务。CUDA 执行模型进一步回答GPU 如何把这些线程组织起来并实际执行它们一次 Kernel 启动后GPU 大致会经历以下过程Grid ↓ Block 被分配到 SM ↓ Block 中的 Thread 被划分为 Warp ↓ Warp 以 SIMT 方式执行 ↓ 线程访问寄存器、共享内存或全局内存2. Thread、Block 和 Grid 的实际含义2.1 Thread逻辑执行单位Thread 是 CUDA 中最小的逻辑执行单位。例如计算两个数组相加C[i] A[i] B[i]可以让每个线程负责一个元素Thread 0计算 C[0] Thread 1计算 C[1] Thread 2计算 C[2]每个线程通常拥有自己的线程索引寄存器局部变量执行状态不过Thread 是程序员看到的逻辑概念并不等于 GPU 中一个独立的物理核心。2.2 Block线程协作和调度单位Block 是一组线程的集合。例如Block 0 ├─ Thread 0 ├─ Thread 1 ├─ Thread 2 └─ ...Block 有两个重要作用作为 GPU 调度的基本单位允许其中的线程进行协作和同步同一个 Block 中的线程可以通过共享内存交换数据并使用线程同步机制协调执行。需要注意Block 是调度单位但不是一个固定的硬件核心。一个 Block 可能包含多个 Warp并由一个 SM 负责执行。2.3 Grid一次 Kernel 启动的全部线程块一次 Kernel 启动产生的所有 Block 组成一个 GridGrid ├─ Block 0 ├─ Block 1 ├─ Block 2 └─ Block 3Grid 只表示一次 Kernel 启动包含的全部计算任务不表示这些 Block 会同时执行。实际运行时GPU 可能同时执行部分 Block其他 Block 等待资源空闲后再执行Block 的执行顺序通常不应由程序假设3. SM 与线程块调度3.1 SM 是什么SM 是Streaming Multiprocessor的缩写可以理解为 GPU 中负责执行线程块的处理单元。一个 SM 通常包含用于执行指令的计算资源寄存器共享内存Warp 调度相关资源GPU 中通常有多个 SMGPU ├─ SM 0 ├─ SM 1 ├─ SM 2 └─ ...3.2 Block 如何分配到 SM当 Kernel 启动后GPU 会把 Grid 中的 Block 分配给可用的 SMGrid ├─ Block 0 ── SM 0 ├─ Block 1 ── SM 1 ├─ Block 2 ── SM 0 └─ Block 3 ── 等待资源一个 SM 可以同时驻留多个 Block但能驻留多少 Block取决于资源限制例如SM 中可用的寄存器数量共享内存使用量线程数量限制Block 数量限制因此Block 不一定会立即执行也不一定全部同时执行。3.3 Block 为什么应该相互独立CUDA 通常按照“Block 之间相互独立”的方式设计程序。原因是不同 Block 可能被分配到不同 SM不同 Block 的执行速度可能不同不同 Block 的执行顺序通常不固定某些 Block 可能要等资源释放后才能启动因此下面这种假设是不安全的Block 0 一定先于 Block 1 执行 Block 0 执行完后 Block 1 才开始 所有 Block 一定同时执行更可靠的设计方式是让每个 Block 尽量独立完成自己的工作。如果不同阶段之间存在严格依赖通常可以拆分为多个 KernelKernel 1 完成 ↓ Kernel 2 开始4. Warp 与 SIMT 执行4.1 Warp 是什么Warp 是 GPU 实际调度线程的重要执行单位。在 NVIDIA CUDA 中一个 Warp 通常包含 32 个线程。例如一个 Block 有 128 个线程通常会被划分为Warp 0Thread 0 ~ Thread 31 Warp 1Thread 32 ~ Thread 63 Warp 2Thread 64 ~ Thread 95 Warp 3Thread 96 ~ Thread 127Block 是调度单位而 Warp 是 SM 中实际进行指令调度的重要单位。4.2 什么是 SIMTSIMT 是Single Instruction, Multiple Threads的缩写意思是单条指令同时作用于多个线程。例如下面的 Kernel__global__voidadd(int*data){inti/* 当前线程索引 */;data[i]data[i]1;}当一个 Warp 中的线程执行这段代码时通常会执行相同的加法指令但每个线程使用自己的i因此操作不同的数据。可以理解为同一条指令 data[i] data[i] 1 Thread 0 使用 i 0 Thread 1 使用 i 1 Thread 2 使用 i 2 ...这就是“相同指令不同线程不同数据”。4.3 Warp 中的线程不一定完全独立虽然程序员把每个 Thread 看成独立执行但硬件通常以 Warp 为单位调度。如果一个 Warp 中的线程走不同分支if(condition){// 路径 A}else{// 路径 B}而同一个 Warp 中部分线程满足 condition 部分线程不满足 condition就会出现分支分化。4.4 什么是分支分化分支分化是指同一个 Warp 中的线程需要执行不同的分支路径。GPU 通常需要分阶段执行这些路径先执行路径 A - 满足条件的线程有效 - 其他线程暂时等待 再执行路径 B - 不满足条件的线程有效 - 其他线程暂时等待因此虽然线程逻辑上走了不同路径但硬件可能需要依次执行这些路径。示例if(threadIdx.x%20){doA();}else{doB();}同一个 Warp 中通常既有偶数线程也有奇数线程因此可能发生分支分化。需要注意分支分化不一定导致程序错误但可能降低执行效率。5. 一个完整的执行示例假设启动一个 Kernelkernel2,64();它表示Grid 中有 2 个 Block每个 Block 有 64 个 Thread总线程数为 128执行过程可以理解为Grid ├─ Block 064 个线程 │ ├─ Warp 032 个线程 │ └─ Warp 132 个线程 │ └─ Block 164 个线程 ├─ Warp 232 个线程 └─ Warp 332 个线程随后GPU 将 Block 分配给可用的 SM每个 Block 中的线程被划分成 WarpSM 调度 Warp 执行指令Warp 中的线程以 SIMT 方式执行每个线程根据自己的索引访问对应数据需要注意Block 0 不一定先于 Block 1 执行Warp 0 不一定一直连续执行到结束GPU 可能在多个 Warp 之间切换程序不应依赖这些未保证的顺序6. 线程执行与内存访问线程执行不仅包括计算指令也包括访问数据。CUDA 中常见的内存层次如下内存类型主要使用者特点寄存器单个线程速度快每个线程独有共享内存同一个 Block 的线程速度较快可用于线程协作全局内存Grid 中的所有线程容量较大但访问延迟较高常量内存等只读空间多个线程适合读取不会改变的数据6.1 寄存器每个线程通常有自己的寄存器用于保存临时变量计算中间结果线程索引一个线程不能直接使用另一个线程的寄存器。6.2 共享内存同一个 Block 中的线程可以使用共享内存交换数据。例如Block 内的多个线程 ↓ 共同使用共享内存共享内存适合用于缓存重复访问的数据线程之间交换中间结果减少对全局内存的访问但使用共享内存时通常需要考虑线程同步避免一个线程读取到尚未准备好的数据。6.3 全局内存全局内存可以被多个 Block 访问容量通常较大但访问延迟也较高。如果一个 Warp 中的线程访问相邻地址通常更有利于内存访问效率。相反如果线程访问地址非常分散可能降低效率。本次只需记住GPU 程序的执行效率不仅取决于计算还取决于线程如何访问内存。7. 同步与执行边界7.1 Block 内同步同一个 Block 中的线程可以使用同步机制__syncthreads();它的作用是让同一个 Block 中的线程在某个位置等待直到该 Block 中满足条件的线程都到达这里。典型场景线程将数据写入共享内存 ↓ 等待其他线程写入完成 ↓ 线程读取共享内存如果没有正确同步可能出现线程 A 还没有写完数据 线程 B 已经开始读取数据从而导致结果错误。7.2 Block 之间不能随意同步普通 CUDA Kernel 中不同 Block 之间通常不能直接使用普通的 Block 内同步机制。原因是不同 Block 可能运行在不同 SM不同 Block 的执行时间不同某些 Block 可能还没有开始执行如果互相等待可能造成死锁因此跨 Block 的阶段依赖通常使用多个 Kernel 表示Kernel 1 不同 Block 完成第一阶段 ↓ Kernel 1 结束 ↓ Kernel 2 不同 Block 开始第二阶段7.3 Kernel 是重要的执行边界不同 Kernel 之间通常形成明确的阶段边界Kernel 1 ↓ Kernel 2如果 Kernel 2 依赖 Kernel 1 的结果就应确保 Kernel 1 已经完成或者使用适当的 CUDA 同步机制。所以Kernel 不只是一个函数也是一个 GPU 任务阶段一个全局依赖边界一个程序组织并行计算的基本单位8. 常见误解误解一一个线程对应一个物理核心不是。Thread 是逻辑执行单位。GPU 通常以 Warp 为单位调度线程实际硬件资源会被多个线程和 Warp 共同使用。误解二一个 Block 就是一个 Warp不一定。Block 可以包含多个 Warp。例如Block 有 256 个线程 256 ÷ 32 8 个 Warp误解三所有线程都严格同时执行不完全是。一个 Warp 中的线程通常执行相同指令但不同 Warp 之间可能交替执行Block 之间也可能按不同顺序执行。误解四Block 一定按照编号顺序执行不是。Block 的执行顺序通常不应由程序假设。误解五__syncthreads()可以同步整个 Grid不是。它主要用于同步同一个 Block 中的线程不能直接作为普通的 Grid 级同步机制。误解六分支分化会让程序一定出错不一定。分支分化通常首先影响性能。只有在程序错误地依赖线程执行顺序或同步关系时才可能导致结果错误。9. 编程模型与执行模型的联系两者可以按照下面的方式理解层次关注内容编程模型程序员如何组织 Host、Device、Kernel、Grid、Block、Thread执行模型GPU 如何将 Block 分配到 SM如何划分 Warp 并执行硬件实现SM 中的计算单元、寄存器、缓存和调度资源整体关系如下程序员编写 Kernel ↓ 使用 Grid、Block、Thread 描述任务 ↓ GPU 将 Block 分配到 SM ↓ Thread 被划分为 Warp ↓ Warp 按 SIMT 方式执行 ↓ 线程访问不同层次的内存编程模型是程序员使用的抽象执行模型是 GPU 实际运行这些抽象的方式。10. 总结CUDA 执行模型的核心可以概括为Grid 中的 Block 被分配到 SMBlock 中的 Thread 被划分为 WarpWarp 以 SIMT 方式执行指令。需要掌握的关键关系是Grid一次 Kernel 启动的全部 Block BlockGPU 调度和线程协作的基本单位 Thread最小逻辑执行单位 Warp通常包含 32 个线程是重要的硬件调度单位 SM负责执行一个或多个 Block 的 GPU 处理单元 SIMT多个线程执行相同指令但处理各自的数据还需要记住三个边界不要假设 Block 的执行顺序不要把普通 Thread 等同于物理核心不要把 Block 内同步机制当作整个 Grid 的同步机制