c++ - CUDA 扭曲和每个 block 的最佳线程数

标签 c++ cuda gpgpu nvidia

根据我对 Kepler GPU 和一般 CUDA 的了解,当单个 SMX 单元在一个 block 上工作时,它会启动由 32 个线程组成的 warp。现在这是我的问题:

1) 如果 SMX 单元可以在 64 个 warp 上工作,这意味着每个 SMX 单元有 32x64 = 2048 个线程的限制。但是 Kepler GPU 有 4 个 warp 调度器,那么这是否意味着在 GPU 内核中只能同时处理 4 个 warp?如果是这样,这是否意味着我真的应该寻找具有 128 个线程(假设线程没有分歧)的倍数的 block ,而不是推荐的 32 个?当然,这忽略了任何分歧,甚至忽略了全局内存访问之类的事情可能导致 warp 停滞并让调度程序切换到另一个的情况。

2) 如果以上是正确的,那么单个 SMX 单元同时在 128 个线程上工作的最佳结果是什么?对于具有 14 个 SMX 单元的 GTX Titan,总共有 128x14 = 1792 个线程?我在网上看到数字表明情况并非如此。 Titan 可以同时运行 14x64(每个 SMX 的最大扭曲数)x32(每个 SMX 的线程数)= 28,672。 SMX 单元怎么会启动 warp,并且只有 4 个 warp 调度程序?他们不能同时启动每个 SMX 的所有 2048 个线程?也许我对 GPU 可以同时启动的最大线程数的定义以及允许排队的线程数感到困惑?

我很感激对此的回答和澄清。

最佳答案

so does this mean that only 4 warps can be worked on simultaneously within a GPU kernel?

可以在开普勒 SMX 上在任何给定的时钟周期中安排多达 4 个 warp 的指令。然而,由于执行单元中的流水线,在任何给定的时钟周期,指令可能处于流水线执行的不同阶段,从当前驻留在 SMX 上的任何到所有扭曲。

And if so, does this mean I should really be looking for blocks that have multiples of 128 threads (assuming no divergence in threads) as opposed to the recommended 32?

我不确定您是如何从上一点跳到这一点的。由于指令混合可能因 warp 而异(因为不同的 warp 可能位于指令流中的不同点)并且指令混合在指令流中从一个地方到另一个地方不同,我看不到 4 个 warp 之间可调度的任何逻辑连接一个给定的时钟周期,以及任何需要有 4 个 warp 的组。给定的 warp 可能位于其指令高度可调度的点(可能在 SP FMA 序列中,需要 SP 内核,这些内核很多),而另外 3 个 warp 可能位于指令流中的另一点,它们的指令是“更难安排”(可能需要 SFU,而 SFU 较少)。因此,任意将 warp 分组为 4 组没有多大意义。请注意,我们不需要分歧来使扭曲彼此不同步。调度程序的自然行为加上执行资源的不同可用性可能会创建最初在一起的扭曲,使其位于指令流中的不同点。

对于你的第二个问题,我认为你的基本知识差距在于理解 GPU 如何隐藏延迟。假设 GPU 有一组 3 条指令要跨 warp 发出:

LD R0, a[idx]
LD R1, b[idx]
MPY R2, R0, R1

第一条指令是来自全局内存的 LD,它可以发出并且不会拖延 warp。同样可以发出第二条指令。然而,由于全局内存的延迟,扭曲将在第 3 条指令处停止。在正确填充 R0 和 R1 之前,无法调度乘法指令。主内存的延迟阻止了它。 GPU 通过(希望)准备好它可以转向的“其他工作”来处理这个问题,即处于未停滞状态的其他扭曲(即具有可以发出的指令)。促进此延迟隐藏过程的最佳方法是为 SMX 提供许多扭曲。这没有任何粒度(例如需要 4 个扭曲)。一般来说,网格中的线程/扭曲/ block 越多,GPU 隐藏延迟的机会就越大。

因此,GPU 确实无法在单个时钟周期内“启动”2048 个线程(即从 2048 个线程发出指令)。但是,当一个 warp 停滞时,它会被放入等待队列,直到解除停滞条件,并且在此之前,让其他 warp 在下一个时钟周期“准备就绪”是有帮助的。

GPU 延迟隐藏是一个经常被误解的话题。如果您搜索它们,有许多可用资源可以了解它。

关于c++ - CUDA 扭曲和每个 block 的最佳线程数,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/32855684/

相关文章:

c++ - Caffe SigmoidCrossEntropyLoss层多标签分类c++

c++ - 为进程添加权限

c++ - 如何使用 CUDA 内核更新 OpenGL VBO

cuda - 为什么 cuFFT 性能会因输入重叠而受到影响?

c++ - 如何在 Vulkan 中跨多个计算队列执行并行计算着色器?

multithreading - 在Metal中同步网格中的所有线程

c++ - 如何将宏生成的#foo 字符串传递给模板类?

c++ - 将文本文件中的数据存储到 C++ 中以便编辑的最佳方法是什么

cuda - 由 mexfunction 调用的内核中的矩阵行/列优先访问

c++ - 如何将类的非静态成员函数传递给 CUDA 内核函数(__global__ 函数)