CUDA，虚拟/隐式块同步

Question

CUDA，虚拟/隐式块同步

cudasynchronization

4

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

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

我的想法是：隐式同步。

假设有时候我只需要一个块，而有时候我需要所有的X个块。在那些只需要一个块的情况下，我可以配置我的代码使第一个块（或第一个SM）处理“真实”数据，而其他X-1个块（或SM）处理一些“虚拟”的数据，执行完全相同的指令，只是有一些其他偏移量。

这样，它们将继续同步，直到我再次需要它们全部。

在这种情况下，调度程序可靠吗？还是永远无法确定？

- elect

2个回答

-4

@Pedro绝对是错的！

最近几项研究工作都在探讨实现全局同步的问题，至少对于非Kepler架构（我还没有这样的设备），结论总是相同的（或者应该是）：无法在整个GPU上实现全局同步。

原因很简单：CUDA块无法被抢占，因此如果你完全占用了GPU，等待屏障会合的线程将永远不会允许块终止。因此，它将不会从SM中移除，并阻止其余块运行。

因此，你只会使GPU冻结，无法从这种死锁状态中逃脱。

-- 编辑以回答Pedro的评论 --

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

由OpenCL in action的作者撰写

-- 编辑以回答Pedro的第二条评论 --

@Jared Hoberock在这个SO帖子中得出了相同的结论: CUDA上的Inter-block屏障

- jopasserat

不，我并不是“绝对错误”的，否则这在我的代码中就行不通了。我已经添加了有关块的最大数量的注释，这解决了您对死锁的担忧。至于“几项研究工作”表明这是不可能的，您能指出其中一两项吗？ - Pedro

这不是同时安排的问题，而是同时运行的块。 - jopasserat

你如何定义并发运行？您可以在每个SM上调度最多八个块，这些块将交替运行。当一个块在while-循环上旋转时，同一SM上的其他块仍然可以运行，在每次内存访问之间填充插槽。顺便说一句，我仍在等待“几项研究工作”。 - Pedro

2

在使用Cuda一年多后，我认为理论只有到一定程度才有意义。我曾经通过大量的谷歌搜索得出块同步不可能的结论，但如果Pedro说他成功了，我不明白他为什么要撒谎。别误会，伙计们，我并不是说你们中的任何一个人是100%正确的，我只是想说我会尝试一下（只要我找到时间去实现它:D）。我相信游戏中有很多因素（硬件和软件），我们应该找出哪些因素起作用。无论如何，我会让你们两个都知道最新情况！ :) - elect

@elect：当你有时间的时候，请尝试使用不同的配置和不同的GPU。它应该可以在小配置下工作。顺便说一句，我也会尽快提供完整的基准测试。 - jopasserat

显示剩余9条评论

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Pedro · Accepted Answer

您提出了几个问题，因此我将尝试分别回答。

每个SM一个块

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

隐式同步

不可以。首先，您无法保证每个块都有自己的SM（请参见上文）。其次，所有块不能同时访问全局存储器。如果它们以任何形式同步运行，则将在第一次内存读/写时失去这种同步性。

块同步

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

__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 == c时，设置*p = v并返回*p的旧值。此部分只是循环直到barrier达到0，即直到所有块都穿过它。

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

请记住，如果您多次调用内核，则应将barrier重置为N。

更新

回复 jHackTheRipper 的答案和 Cicada 的评论，我应该指出，不应该尝试启动超过 GPU 可以同时调度的块数！这受到许多因素的限制，您应该使用 CUDA 计算器找到内核和设备的最大块数。

然而，根据原始问题，只有与 SM 数量相同的块正在启动，所以这一点没有意义。