CUDA: if语句中的__syncthreads()
我有一个关于CUDA同步的问题。特别是，我需要一些关于if语句中同步的说明。我的意思是，如果我把一个__syncthreads()放在一个if语句的作用域下，这个语句被块内的一小部分线程击中，会发生什么？我认为一些线程将“永远”等待其他不会达到同步点的线程。因此，我编写并执行了一些示例代码来检查：
__global__ void kernel(float* vett, int n)
{
int index = blockIdx.x*blockDim.x + threadIdx.x;
int gridSize = blockDim.x*gridDim.x;
while( index < n )
{
vett[index] = 2;
if(threadIdx.x < 10)
{
vett[index] = 100;
__syncthreads();
}
__syncthreads();
index += gridSize;
}
}令人惊讶的是，我观察到输出是相当“正常”的(64个元素，块大小32)：
100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2
100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2所以我稍微修改了一下代码，如下所示：
__global__ void kernel(float* vett, int n)
{
int index = blockIdx.x*blockDim.x + threadIdx.x;
int gridSize = blockDim.x*gridDim.x;
while( index < n )
{
vett[index] = 2;
if(threadIdx.x < 10)
{
vett[index] = 100;
__syncthreads();
}
__syncthreads();
vett[index] = 3;
__syncthreads();
index += gridSize;
}
}输出结果是：
3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3
3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 再一次，我错了:我认为if语句中的线程在修改完向量的元素后，将保持等待状态，永远不会离开if作用域。所以..。你能澄清一下发生了什么吗？在同步点之后获取的线程是否会解锁在屏障处等待的线程？如果你需要重现我的情况，我使用了CUDA Toolkit 5.0 RC和SDK 4.2。在此之前非常感谢。
回答 5
Stack Overflow用户
发布于 2012-09-21 04:28:03
简而言之，行为是未定义的。所以它有时可能会做你想做的，也可能不做，或者(很有可能)只是挂起或者使你的内核崩溃。
如果你真的很好奇事情是如何在内部工作的，你需要记住线程并不是独立执行的，而是一次执行一次warp ( 32个线程组)。
当然，这会产生条件分支的问题，条件分支在整个warp过程中不会统一计算。这个问题通过一个接一个地执行两个路径来解决，每个路径都禁用了那些不应该执行该路径的线程。IIRC在现有硬件上，首先执行分支，然后在未执行分支的位置执行路径，但此行为未定义，因此无法保证。
路径的这种单独执行一直持续到编译器可以确定它被两个单独执行路径的所有线程保证到达的某个点(“重新会聚点”或“同步点”)。当第一个代码路径的执行到达这一点时，它被停止，而第二个代码路径被执行。当第二个路径到达同步点时，将再次启用所有线程，并从那里均匀地继续执行。
如果在同步之前遇到另一个条件分支，情况会变得更加复杂。这个问题可以通过仍然需要执行的路径堆栈来解决(幸运的是，堆栈的增长是有限的，因为对于一个warp，我们最多可以有32个不同的代码路径)。
同步点的插入位置是未定义的，甚至在不同的体系结构之间略有不同，因此也不能保证。你会从Nvidia得到的唯一的(非官方的)评论是，编译器在寻找最佳同步点方面做得很好。然而，经常有一些微妙的问题可能会使最佳点比您预期的更低，特别是在线程提前退出的情况下。
现在要理解__syncthreads()指令的行为(它转换为PTX中的bar.sync指令)，重要的是要认识到，该指令不是针对每个线程执行的，而是一次针对整个warp执行的(不管是否禁用了任何线程)，因为只需要同步块的warp。warp的线程已经在同步执行，当尝试从不同的条件代码路径同步线程时，进一步的同步将不起作用(如果所有线程都已启用)或导致死锁。
您可以按照自己的方式从这个描述到您的特定代码的行为。但请记住，所有这些都是未定义的，没有保证，依赖于特定的行为可能会在任何时候破坏您的代码。
您可能希望查看PTX manual以了解更多细节，特别是__syncthreads()编译成的bar.sync指令。亨利·黄的"Demystifying GPU Microarchitecture through Microbenchmarking" paper，下面由艾哈迈德引用，也很值得一读。即使现在已经过时的架构和CUDA版本，关于条件分支和__syncthreads()的部分似乎仍然是有效的。
Stack Overflow用户
发布于 2012-09-21 04:27:18
CUDA模型是MIMD的，但目前的NVIDIA GPU在warp粒度而不是线程上实现__syncthreads()。这意味着，这些是同步的warps inside a thread-block，不一定是threads inside a thread-block。__syncthreds()等待线程块的所有“变形”命中屏障或退出程序。有关更多详细信息，请参阅Henry Wong's Demistifying paper。
Stack Overflow用户
发布于 2012-09-21 04:07:50
除非始终在一个线程块内的所有线程中都能到达该语句，否则不能使用__syncthreads()。从programming guide (B.6)：
在条件代码中允许
__syncthreads()，但仅当条件在整个线程块中计算相同时才允许，否则代码执行可能会挂起或产生意外的副作用。
基本上，您的代码不是格式良好的CUDA程序。
https://stackoverflow.com/questions/12519573
复制相似问题
腾讯云开发者
