我知道块同步是不可能的，唯一的办法是启动一个新内核。
但是，假设我启动了 X 个块，其中 X 对应于我 GPU 上 SM 的数量。我应该考虑调度程序将为每个 SM 分配一个块......对吗？如果 GPU 被用作辅助显卡（完全专用于 CUDA），这意味着理论上没有其他进程使用它……对吧？
我的想法是：隐式同步。
假设有时我只需要一个块，有时我需要所有 X 个块。好吧，在我只需要一个块的情况下，我可以配置我的代码，以便第一个块（或第一个 SM）将在“真实”数据上工作，而其他 X-1 块（或 SM）在某些“ dummy”数据，执行完全相同的指令，只是有一些其他偏移量。
这样所有这些都将继续同步，直到我再次需要它们为止。
在这种情况下调度器可靠吗？还是您永远无法确定？
【问题讨论】：
标签： cuda synchronization
你有几个问题，所以我会尝试分别解决。
每个 SM 一个区块
我在nVidia's own forums 上问过这个问题，因为我得到的结果表明这不是发生的情况。显然，如果块的数量等于 SM 的数量，则块调度器不会为每个 SM 分配一个块。
隐式同步
没有。首先，你不能保证每个区块都有自己的 SM（见上文）。其次，所有块不能同时访问全局存储。如果它们完全同步运行，它们将在第一次内存读/写时失去这种同步性。
块同步
现在有个好消息：是的，你可以。 CUDA C Programming Guide 的 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 = v iff *p == c 并返回*p 的旧值。这部分只是循环直到barrier 到达0，即直到所有块都穿过它。
请注意，您必须将此部分包装在对 __synchtreads() 的调用中，因为块中的线程不会以严格的锁步执行，您必须强制它们全部等待第零个线程。
请记住，如果您多次调用内核，则应将barrier 设置回N。
更新
在回复jHackTheRipper 的回答和Cicada 的评论时，我应该指出，您不应该尝试启动比GPU 上可以同时安排的更多的块！这受到许多因素的限制，您应该使用CUDA Occupancy Calculator 来查找内核和设备的最大块数。
不过，从最初的问题来看，只有与 SM 一样多的区块正在启动，所以这一点没有实际意义。
【讨论】：
__syncthread()，因为我每个块只有 32 个线程。如果您不愿相信我的话，您可能想查看“CUDA 示例：通用 GPU 编程简介”的附录 A，其中讨论了原子操作、互斥和块之间的同步。
@Pedro 肯定错了！
实现全局同步已成为最近几项研究工作的主题，最后是非 Kepler 架构（我还没有）。结论总是相同（或应该）：不可能在整个 GPU 上实现这样的全局同步。
原因很简单：CUDA 块不能被抢占，所以如果你完全占用了 GPU，等待屏障会合的线程将永远不会允许块终止。因此，它不会从 SM 中移除，并且会阻止剩余的块运行。
因此，您只会冻结永远无法摆脱这种死锁状态的 GPU。
-- 编辑以回答佩德罗的言论--
其他作者也注意到了这些缺点，例如： http://www.openclblog.com/2011/04/eureka.html
OpenCL 的作者在行动
-- 编辑回答佩德罗的第二句话--
@Jared Hoberock 在这篇 SO 帖子中得出了相同的结论： Inter-block barrier on CUDA
【讨论】：
while-loop 上旋转时，同一个 SM 上的其他块仍然可以运行，填充每个内存访问之间的插槽。对了，我还在等“几篇研究作品”。