CUDA，虚设/隐式块同步(cuda, dummy/implicit block synchroni

我知道块同步是不可能的，唯一的办法就是推出一个新的内核。

但是，让我们假设我启动X块，其中X对应于SM在我的GPU的数量。我要的方面是调度程序将分配给每个SM块......吧？并且如果GPU被用作次级图形卡（完全专用于CUDA），这意味着，从理论上讲，没有其他的过程中使用它是吗？

我的想法是这样的：隐含同步。

让我们假设有时我只需要一个块，有时我需要所有的X块。那么，在那些我需要的只是一个块的情况下，我可以配置我的代码，以使第一块（或第一SM）将在“真实”数据的工作，而另一个X-1块（或短信）上的一些“虚拟”的数据，执行完全相同的指令，只需用一些其他的偏移。

所以，他们都将继续同步，直到我将再次需要他们。

是调度此条件下可靠吗？或者，你可以永远确定？

Answer 1:

你真是一语中的几个问题，所以我会尽力单独解决这些问题。

每个SM一个块

我问这个了一段时间后上的nVidia自己的论坛，因为我得到的是表示，这是不会发生什么变化的结果。显然，如果块的数量等于短信号的块调度器将不分配每SM的块。

隐式同步

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

块同步

现在的好消息是：是的，可以。在第B.11所述的原子指令CUDA C编程指南可以被用来创建一个屏障。假设你有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 = v IFF *p == c并返回的旧值*p 。这部分只是循环，直到barrier到达0 ，即直到所有块已经越过它。

请注意，您必须包装这一部分中调用__synchtreads()如块不严格锁步执行线程，你必须强迫他们所有的等待线程零。

只要记住，如果你打电话给你的内核不止一次，你应该设置barrier回N 。

更新

在回答jHackTheRipper的答案，蝉的评论，我要指出，你不应该试图不是可以同时安排在GPU上展开更块！这是由许多因素的限制，你应该使用CUDA占有率计算器来寻找您的内核和设备的最大块数。

在原来的问题来看，虽然，只是因为有SM的许多块被启动，所以这点是没有实际意义。