4

我知道块同步是不可能的，唯一的方法是启动一个新内核。

但是，假设我启动了 X 个块，其中 X 对应于我 GPU 上 SM 的数量。我应该考虑调度程序将为每个 SM 分配一个块......对吗？如果 GPU 被用作辅助图形卡（完全专用于 CUDA），这意味着理论上没有其他进程使用它......对吗？

我的想法如下：隐式同步。

假设有时我只需要一个块，有时我需要所有 X 个块。好吧，在我只需要一个块的情况下，我可以配置我的代码，以便第一个块（或第一个 SM）将在“真实”数据上工作，而其他 X-1 块（或 SM）在某些“ dummy”数据，执行完全相同的指令，只是有一些其他偏移量。

这样所有这些都将继续同步，直到我再次需要它们。

在这种情况下调度器可靠吗？或者你永远不能确定？

4

2 回答 2

3

您有几个问题合二为一，所以我将尝试分别解决。

每个 SM 一个块

我在nVidia 自己的论坛上问过这个问题，因为我得到的结果表明这不是发生的情况。显然，如果块的数量等于 SM 的数量，则块调度器不会为每个 SM 分配一个块。

隐式同步

不，首先，你不能保证每个区块都有自己的 SM（见上文）。其次，所有块不能同时访问全局存储。如果它们完全同步运行，它们将在第一次内存读/写时失去这种同步性。

块同步

现在有个好消息：是的，你可以。CUDA C 编程指南第 B.11 节中描述的原子指令可用于创建屏障。假设您有N块在 GPU 上同时执行。

__device__ int barrier = N;

__global__ void mykernel ( ) {

    /* Do whatever it is that this block does. */
    ...

    /* Make sure all threads in this block are actually here. */
    __syncthreads();

    /* Once we're done, decrease the value of the barrier. */
    if ( threadIdx.x == 0 )
        atomicSub( &barrier , 1 );

    /* Now wait for the barrier to be zero. */
    if ( threadIdx.x == 0 )
        while ( atomicCAS( &barrier , 0 , 0 ) != 0 );

    /* Make sure everybody has waited for the barrier. */
    __syncthreads();

    /* Carry on with whatever else you wanted to do. */
    ...

    }

该指令以原子方式atomicSub(p,i)计算*p -= i并且仅由块中的第零个线程调用，即我们只想递减barrier一次。该指令atomicCAS(p,c,v)设置*p = viff*p == c并返回 的旧值*p。这部分只是循环直到barrier到达0，即直到所有块都穿过它。

请注意，您必须将此部分包装在调用中，__synchtreads()因为块中的线程不会以严格的锁步执行，您必须强制它们全部等待第零个线程。

请记住，如果您多次调用内核，则应设置barrier回N.

更新

在回复jHackTheRipper的回答和Cicada的评论时，我应该指出，您不应该尝试启动比 GPU 上可以同时安排的更多的块！这受到许多因素的限制，您应该使用CUDA 占用计算器来查找内核和设备的最大块数。

不过，从最初的问题来看，只有与 SM 一样多的区块正在启动，所以这一点没有实际意义。

于 2012-07-04T10:58:40.670 回答
-4

@Pedro 绝对是错误的！

实现全球同步一直是最近几项研究工作的主题，最后是非开普勒架构（我还没有）。结论总是相同（或应该是）：不可能在整个 GPU 上实现这样的全局同步。

原因很简单：CUDA 块不能被抢占，所以如果你完全占用了 GPU，等待屏障 Rendez-vous 的线程将永远不会允许块终止。因此，它不会从 SM 中删除，并且会阻止剩余的块运行。

因此，您只会冻结永远无法摆脱这种死锁状态的 GPU。

-- 编辑回答佩德罗的评论 --

其他作者也注意到了这些缺点，例如： http ://www.openclblog.com/2011/04/eureka.html

由 OpenCL 的作者在行动

-- 编辑回答佩德罗的第二句话 --

@Jared Hoberock 在这篇 SO 帖子中得出了相同的结论： CUDA 上的块间障碍

于 2012-07-04T12:42:45.590 回答