cuda - 流式多处理器、块和线程 (CUDA)

Question

CUDA 内核、流式多处理器和块和线程的 CUDA 模型之间有什么关系？

什么被映射到什么，什么被并行化以及如何？什么更有效，最大化块数或线程数？

---

我目前的理解是每个多处理器有 8 个 cuda 内核。并且每个 cuda 核心将能够一次执行一个 cuda 块。并且该块中的所有线程都在该特定核心中连续执行。

它是否正确？

Answer 1

CUDA 编程指南中详细描述了线程/块布局。特别是，第 4 章指出：

CUDA 架构围绕可扩展的多线程流式多处理器 (SM) 阵列构建。当主机 CPU 上的 CUDA 程序调用内核网格时，网格的块被枚举并分发到具有可用执行能力的多处理器。一个线程块的线程在一个多处理器上并发执行，多个线程块可以在一个多处理器上并发执行。当线程块终止时，新块在空出的多处理器上启动。

每个 SM 包含 8 个 CUDA 内核，并且在任何时候它们都在执行 32 个线程的单个 warp - 因此为整个 warp 发出一条指令需要 4 个时钟周期。您可以假设任何给定 warp 中的线程以锁步执行，但要跨 warp 同步，您需要使用__syncthreads().

Answer 2

对于 GTX 970，有 13 个流式多处理器 (SM)，每个具有 128 个 Cuda 内核。Cuda 核心也称为流处理器 (SP)。

您可以定义将块映射到 GPU 的网格。

您可以定义将线程映射到流处理器的块（每个 SM 有 128 个 Cuda 核心）。

一个warp总是由32个线程组成，一个warp的所有线程同时执行。

要使用 GPU 的全部可能功能，每个 SM 需要比 SM 拥有的 SP 更多的线程。对于每个计算能力，有一定数量的线程可以一次驻留在一个 SM 中。您定义的所有块都排队等待 SM 拥有资源（可用的 SP 数量），然后加载。SM 开始执行 Warps。由于一个 Warp 只有 32 个线程，而一个 SM 有例如 128 个 SP，一个 SM 可以在给定时间执行 4 个 Warp。问题是如果线程进行内存访问，线程将阻塞，直到它的内存请求得到满足。在数字上：SP 上的算术计算有 18-22 个周期的延迟，而非缓存的全局内存访问可能需要 300-400 个周期。这意味着如果一个 warp 的线程正在等待数据，则只有 128 个 SP 的一个子集可以工作。因此，如果可用，调度程序会切换到执行另一个扭曲。如果这个经纱阻塞，它会执行下一个，依此类推。这个概念称为延迟隐藏。warp 的数量和块大小决定了占用率（从 SM 可以选择执行多少 warp）。如果占用率很高，则 SP 更不可能没有工作。

您关于每个 cuda 核心一次执行一个块的说法是错误的。如果您谈论流式多处理器，它们可以从驻留在 SM 中的所有线程执行扭曲。如果一个块的大小为 256 个线程，并且您的 GPU 允许每个 SM 驻留 2048 个线程，则每个 SM 将有 8 个块驻留，SM 可以从中选择要执行的 warp。执行的warp的所有线程都是并行执行的。

您可以在此处找到不同计算能力和 GPU 架构的数字： https ://en.wikipedia.org/wiki/CUDA#Limitations

您可以从 Nvidia Occupancy Calculation sheet (by Nvidia)下载占用计算表。

Answer 3

只有当 SM 有足够的资源用于线程块（共享内存、warp、寄存器、屏障等）时，计算工作分配器才会在 SM 上调度线程块 (CTA)。分配诸如共享内存之类的线程块级资源。分配为线程块中的所有线程创建了足够的扭曲。资源管理器使用循环法将扭曲分配给 SM 子分区。每个 SM 子分区都包含一个 warp 调度程序、寄存器文件和执行单元。一旦一个 warp 被分配给一个子分区，它将保留在子分区上，直到它完成或被上下文切换（Pascal 架构）抢占。在上下文切换恢复时，warp 将恢复到相同的 SM 相同的 warp-id。

当 warp 中的所有线程都完成时，warp 调度程序等待 warp 发出的所有未完成的指令完成，然后资源管理器释放 warp 级别的资源，包括 warp-id 和 register 文件。

当线程块中的所有 warp 完成时，块级资源被释放，并且 SM 通知 Compute Work Distributor 该块已完成。

一旦一个 warp 被分配给一个子分区并且所有的资源都被分配，这个 warp 被认为是活跃的，这意味着 warp 调度程序正在积极地跟踪这个 warp 的状态。在每个周期，warp 调度器确定哪些活动的 warp 被停止，哪些有资格发出指令。warp 调度程序选择最高优先级的合格 warp 并从 warp 发出 1-2 个连续指令。双重问题的规则特定于每个架构。如果 warp 发出内存加载，它可以继续执行独立指令，直到它到达依赖指令。然后，warp 将报告停止，直到加载完成。对于从属数学指令也是如此。SM 架构旨在通过在 warp 之间切换每个周期来隐藏 ALU 和内存延迟。

这个答案没有使用术语 CUDA 核心，因为这引入了一个不正确的心智模型。CUDA 内核是流水线单精度浮点/整数执行单元。问题率和依赖延迟因每个架构而异。每个 SM 子分区和 SM 都有其他执行单元，包括加载/存储单元、双精度浮点单元、半精度浮点单元、分支单元等。

为了最大限度地提高性能，开发人员必须了解块与扭曲与寄存器/线程的权衡。

术语占用率是 SM 上的活动扭曲与最大扭曲的比率。Kepler - Pascal 架构（GP100 除外）每个 SM 有 4 个 warp 调度程序。每个 SM 的最小 warp 数应至少等于 warp 调度程序的数量。如果架构具有 6 个周期（Maxwell 和 Pascal）的相关执行延迟，那么每个调度程序至少需要 6 个 warp，即每个 SM 24 个（24 / 64 = 37.5% 占用率）来覆盖延迟。如果线程具有指令级并行性，那么这可以减少。几乎所有内核都会发出可变延迟指令，例如可能需要 80-1000 个周期的内存加载。这需要每个 warp 调度程序有更多的活动 warp 来隐藏延迟。对于每个内核，warp 的数量和其他资源（如共享内存或寄存器）之间存在一个权衡点，因此不建议优化 100% 占用，因为可能会做出一些其他牺牲。CUDA 分析器可以帮助识别指令发出率、占用率和停顿原因，以帮助开发人员确定平衡。

线程块的大小会影响性能。如果内核有大块并使用同步屏障，那么屏障停顿可能是一个停顿的原因。这可以通过减少每个线程块的扭曲来缓解。

Answer 4

一台设备上有多个流式多处理器。
一个 SM 可能包含多个块。每个块可能包含多个线程。
一个 SM 有多个 CUDA 内核（作为开发人员，你不应该关心这个，因为它是被 warp 抽象出来的），它将在线程上工作。SM 总是处理线程的扭曲（总是 32）。扭曲只会在同一块的线程上工作。
SM 和 block 都对线程数、寄存器数和共享内存有限制。