CUDA：__syncthreads（）里面if语句

我有一个关于CUDA同步的问题。特别是，我需要澄清if语句中的同步问题。我的意思是，如果我将一个__syncthreads（）放在if语句的作用域内，那么会发生什么？我认为有些线程会“永远”等待其他不会触及同步点的线程。所以，我写和执行的一些示例代码来检查：CUDA：__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个元素，BLOCKSIZE 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语句中的线程在修改vector的元素之后将保持等待状态，并且永远不会超出if范围。所以......请你说清楚发生了什么？同步点之后获得的线程是否会阻止等待屏障的线程？如果您需要重现我的情况，我使用SDK 4.2的CUDA Toolkit 5.0 RC。提前致谢。

来源

2012-09-20 biagiop1986

给予回答您问题的人的选中标记（接受的答案）。 – Yakk

除非总是在一个线程块内的所有线程中都达到语句，否则不得使用__syncthreads()。从programming guide（B.6）：

__syncthreads()被允许在条件代码，但只有当条件估值相同整个线程块，否则代码执行将有可能挂起或产生不期望的副作用。

基本上，您的代码不是一个格式良好的CUDA程序。

来源

2012-09-20 20:07:50

当然不是！但我只是为了检查它的行为而写它。 – biagiop1986

@ biagiop1986：恩...你有一段库代码和硬件，附带一份说明“你不能做X”的文档。现在你问*我们*，公众，如果你这样做会发生什么X - 我们*应该怎么知道？询问供应商！知道该计划会不合格是不是足够了？ –

这取决于......说我应该避免在我的程序中那样的代码，因为它不合格（我发誓，我会的！），但我对“如何”很好奇。此外，我经常在这里发现关于问题的说明，比供应商说明更好。所以，我会回到这里问你，而不是每个人，因为我将来会遇到每一个编码问题。 Stackoverflow是最好的！顺便说一句，谢谢大家。 – biagiop1986

CUDA模型是MIMD，但目前的NVIDIA GPU以warp粒度而不是线程实现__syncthreads()。这意味着，这些是warps inside a thread-block谁同步不一定threads inside a thread-block。 __syncthreds()等待线程块的所有'warps'击中屏障或退出程序。有关更多详细信息，请参阅Henry Wong's Demistifying paper。

来源

2012-09-20 20:27:17 ahmad

那篇论文确实是一个很好的参考。我忘记了它也涵盖了条件分支。 – tera

谢谢你，很好的资源。 – biagiop1986

总之，行为是undefined。所以它有时可能会做你想做的事情，或者它可能不会，或者（很可能）会挂起或崩溃你的内核。

如果您真的很好奇内部是如何工作的，您需要记住线程不会独立执行，而是一次执行一个warp（32个线程组）。

这当然会造成条件分支的问题，条件分支不会在整个变形中统一计算。通过执行两个路径，一个接一个地执行，每个都禁用那些不应该执行该路径的线程来解决问题。 IIRC在现有硬件上首先采用分支，然后在未采用分支的位置执行路径，但此行为是undefined，因此无法保证。

这种单独的路径执行继续到某个点，编译器可以确定它确保两条独立执行路径（“重合点”或“同步点”）的所有线程都能够达到该路径。当第一条代码路径的执行到达这一点时，它停止，代替执行第二条代码路径。当第二条路径到达同步点时，所有线程都会再次启用，并从此处继续执行。

如果在同步之前遇到另一个条件分支，情况会变得更加复杂。这个问题是通过一堆仍然需要执行的路径来解决的（幸运的是堆栈的增长是有限的，因为我们对一个warp最多可以有32个不同的代码路径）。

在插入同步点的地方是undefined甚至在不同架构之间也会有所不同，所以再次没有保证。您将从Nvidia获得的唯一（非官方）评论是编译器非常擅长寻找最佳同步点。然而，通常存在一些微妙的问题，可能会使最佳点进一步下降，特别是线程早退的时候。

现在，为了理解__syncthreads（）指令的行为（在PTX中转换成bar.sync指令），意识到这条指令不是每个线程都执行，而是一次执行整个warp是否禁用任何线程），因为只有块的变形需要同步。一个warp的线程已经同步执行，并且进一步的同步将不会起作用（如果所有的线程都被启用），或者当试图从不同的条件代码路径同步线程时导致死锁。

你可以按照你的方式从这个描述到你的特定代码行为如何。但请记住，所有这些都是undefined，没有任何保证，依靠特定的行为可能会随时破坏您的代码。

有关更多详细信息，请参阅PTX manual，特别是编号为__syncthreads()的bar.sync指令。亨利王的"Demystifying GPU Microarchitecture through Microbenchmarking" paper，由ahmad引用，也很值得一读。尽管对于现在过时的架构和CUDA版本，关于条件分支和__syncthreads()的部分似乎仍然是一般有效的。

来源

2012-09-20 20:28:03 tera

谢谢，非常清楚的说明。 – biagiop1986

__syncthreads（）用于同步块内的线程。这意味着在继续之前，块中的所有线程都会等待所有线程完成。

考虑一个块中有一些线程的情况，它会进入if语句而有些线程不会。那些等待的线程将被阻止;永远在等待。

通常，在if条件语句中放置同步并不是一种好的方式。最好避免它，并重新设计你的代码，如果你有它。同步的目的是为了确保所有线程一起进行，为什么你首先使用if语句将它们过滤掉？

要添加，如果需要跨块的同步。重新启动内核是必需的。

来源

2012-09-21 07:28:33

CUDA：__syncthreads（）里面if语句

回答

相关问题