CUDA 线程、SMX、SP 和块，它们是如何工作的？

Posted 2023-04-15

【中文标题】CUDA 线程、SMX、SP 和块，它们是如何工作的？

【英文标题】：CUDA threads, SMX, SP and blocks, how do they work?

【发布时间】：2012-09-02 02:40:55

【问题描述】：

我对 CUDA 的工作原理有点困惑，线程是否执行每个相同的指令 (SIMT)，但使用通过不同索引访问的单个数据？或者它被认为是“不同的数据”（所以它也是 SIMD）？

SMX 是整个 GPU 芯片吗？一个 SMX 应该由多个 SP 组成，每个 SP 一次执行一个线程，是不是只分配给一个 SP 的线程块？

我现在有点迷茫

【参考方案1】：

网格启动是线程块的 1-3 维启动。线程块是 1-3 维的线程组。 CUDA 工作分配器将线程块分配给 SMX 单元。低端设备可能有 1 个 SMX 单元。一个高端设备可能有超过 10 个 SMX 单元。

SMX 单元将线程块分解为 32 个线程组，称为 warp。 SMX 单元一次最多可以分配 64 个 warp 或 16 个块。由于资源限制（块、warp、每个线程的寄存器、每个块的共享内存或屏障），数量可能会更少。

每个 SMX 单元都有 4 个 warp 调度程序，每个调度程序负责一个 warp 子集。在每个周期，warp 调度程序将选择一个合格的 warp 并发出 1 或 2 条指令。为了双重发布，两条指令必须是独立的并且使用不同的执行单元。例如，可以将一条指令分派到浮点单元，将第二条指令分派到加载存储单元。

除了双重发布之外，warp 调度程序还可以向 warp 发布背靠背独立指令。当检测到依赖关系，或者下一条指令的执行单元很忙，或者 warp 没有指令（等待获取）时，如果一个符合条件，warp 调度程序将选择不同的 warp。

每个线程都有自己的一组通用寄存器、条件代码、谓词代码和本地内存。每个线程都是线程块的成员。所有线程都可以访问线程块资源，包括共享内存和屏障。网格启动中的所有线程都可以访问网格资源，包括常量内存、纹理绑定和表面绑定。所有线程都可以访问全局内存。

【讨论】：

优秀的总结。那么障碍是一种有限的资源吗？这是为什么呢？

PTX bar 指令支持每个线程块 (CTA) 16 个屏障。为了提高效率，这些是有限的硬件资源。大多数计算程序使用 1 个屏障 (__syncthreads())。为 MAX_BLOCKS_PER_SM 支持 16 的好处非常小，使用 CUDA C/C++ 不太可能发生。