流式多处理器，块和线程（CUDA）

Question

CUDA核心，流媒体多处理器与块和线程的CUDA模型之间有什么关系？

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

---

我目前的理解是，每个多处理器有8个CUDA核心。并且每个CUDA核心都可以一次执行一个CUDA块。该块中的所有线程均在该特定核心中串行执行。

它是否正确？

Answer 1

线程 /块布局在 CUDA编程指南. 。特别是第4章指出：

CUDA体系结构围绕多线程流多处理器（SMS）的可扩展阵列构建。当主机CPU上的CUDA程序调用内核网格时，将枚举网格的块并分配给具有可用执行能力的多处理器。线程块的线程在一个多处理器上同时执行，并且多个线程块可以在一个多处理器上同时执行。随着螺纹块终止，在空置的多处理器上启动了新块。

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

Answer 2

对于GTX 970，有13个流媒体多处理器（SM），每个都有128个CUDA核。 CUDA核心也称为流处理器（SP）。

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

您可以定义哪些映射线程到流处理器（每SM的128个CUDA内核）。

一个经纱总是由32个线形成，并且经纱的所有线程被模拟执行。

要使用GPU的全部功能，您每SM所需的线程比SM SPS需要更多的线程。对于每个计算功能，都有一定数量的线程，一次可以在一个SM中驻留。您定义的所有块均已排队，并等待SM具有资源（免费的SPS数量），然后加载它。 SM开始执行扭曲。由于一个扭曲只有32个线程，并且一个SM具有128 sps，因此SM可以在给定时间执行4个扭曲。问题是，如果线程确实访问内存，则该线程将阻止线程，直到满足其内存请求为止。数量：SP上的算术计算的潜伏期为18-22个周期，而非平衡的全局内存访问最多可能需要300-400个周期。这意味着，如果一个翘曲的线程仅等待数据，则只有128个SP的子集可以工作。因此，调度程序切换如果可用，可以执行另一个纱。如果此翘曲阻塞将执行下一个，依此类推。这个概念称为延迟隐藏。扭曲的数量和块大小决定了占用率（从SM可以选择执行多少WARPS）。如果占用率很高，那么SPS的不可能更不可能。

您的声明说，每个CUDA核心一次都会执行一个块是错误的。如果您谈论流媒体多处理器，他们可以从SM中的所有线程中执行WARPS。如果一个块的尺寸为256个线程，并且您的GPU允许2048个线程每SM居民每张SM的居民将有8个块，则SM可以从中选择执行Warps。已执行的WARPS的所有线程并行执行。

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

您可以从NVIDIA下载占用计算表 占用计算表（由NVIDIA）.

Answer 3

仅当SM有足够的资源来为线程块（共享存储器，扭曲，寄存器，屏障，...）时，计算工作分销商才能安排SM上的线程块（CTA）。线程块级别资源分配了共享内存。分配为线程块中的所有线程创建足够的扭曲。资源经理使用 罗宾 到SM子分区。每个SM子分区都包含一个经线调度程序，注册文件和执行单元。一旦将经纱分配给一个子分区，它将保留在该子分区上，直到它完成或被上下文开关（Pascal Architecture）抢占。在上下文开关上，还原纱将恢复到同一SM相同的翘曲ID。

经纱中的所有线程都完成了经线调度程序，等待经纱发出的所有未偿还说明完成，然后资源经理释放包括翘曲ID和注册文件在内的经线级别资源。

当螺纹块中的所有扭曲完成后，会发布块级别资源，并将SM通知计算工作分销商该块已完成。

一旦将经纱分配给一个子分区并分配了所有资源，经纱就被认为是主动的，这意味着经线调度程序正在积极跟踪翘曲状态。在每个循环中，经纱调度程序确定哪个主动扭曲被停滞不前，哪些有资格发布指令。经线调度程序选择了最高优先级的符合条件的纱，并从经纱中发出了1-2连续的说明。双标准规则针对每个体系结构。如果翘曲发出内存负载，它可以继续执行独立说明，直到达到依赖指令为止。然后，翘曲将报告停滞直到负载完成。依赖的数学说明也是如此。 SM体系结构旨在通过在WARPS之间切换每个周期来隐藏ALU和内存延迟。

该答案不使用术语CUDA核心，因为这引入了不正确的心理模型。 CUDA内核是管道的单精度浮点/整数执行单元。问题速率和依赖性延迟是每个体系结构的特异性。每个SM子分区和SM都有其他执行单元，包括负载/存储单元，双精度浮点单元，半精度浮点单元，分支单元等。

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

该术语占用率是主动扭曲与最大扭曲的比率。开普勒 - Pascal架构（GP100除外）每SM有4个经纱调度程序。每SM的扭曲数量最少，至少应等于经纱调度程序的数量。如果体系结构的执行延迟为6个周期（Maxwell和Pascal），则每个调度程序至少需要6个WARPS，每SM（24 /64 = 37.5％的占用率）才能覆盖延迟。如果线程具有并行性的指令级别，则可以减少这一点。几乎所有内核都会发出可变的延迟说明，例如可以进行80-1000个周期的内存负载。这需要每个扭曲调度程序的更多活跃的扭曲才能隐藏延迟。对于每个内核，不建议您对100％占用率进行优化的扭曲数量和其他资源之间的权衡点，因为可能会做出其他牺牲。 CUDA PROFILER可以帮助确定指令问题，占用率和失速原因，以帮助开发人员确定这种平衡。

螺纹块的大小会影响性能。如果内核具有较大的块并使用同步屏障，那么屏障摊位可能是一个失速的原因。可以通过减少每个螺纹块的扭曲来缓解这一点。

Answer 4

一台设备上有多个流媒体多处理器。
SM可能包含多个块。每个块可能包含几个线程。
SM具有多个CUDA核心（作为开发人员，您不应该关心它，因为它是由Warp抽象的），这将在线程上起作用。 SM总是在线程上工作（总是32）。翘曲只能从同一块上处理线程。
SM和Block都对线程数，寄存器数量和共享内存的数量都有限制。