如何在CUDA应用程序中正确应用线程同步

Question

一般来说，我在应用程序中很少使用线程同步，因为我不经常需要这个功能。我不是真正的高级C/c++程序员，但我也不是初学者。我开始学习CUDA C，与CPU的能力相比，现在GPU的能力让我很兴奋，我意识到CUDA编程主要是关于并行线程执行的，有时适当的线程同步是必要的。事实上，我甚至不知道如何在C或c++中应用线程同步。我最后一次使用同步是在2年前，当时我正在用Java编写简单的应用程序，像这样:

synchronized returnType functionName(parameters)
{
    ...
}

什么允许'functionName'一次只由一个线程执行-即这个函数由不同的线程交替执行。现在回到CUDA C，如果我在一个块中有200个线程，在while循环中运行代码:

while(some_condition)
{
    ...
}

如何使线程<0 - 99>彼此同步，线程<100 - 199>彼此同步，但应用同步的方式线程<0 - 99>和<100 - 199>交替执行(即前100个线程运行'while'的内容，之后的100个线程运行'while'的内容，等等)?

Answer 1

我认为你可能只是需要更多地了解cuda。您可能会陷入这样一个陷阱:认为以前学过的编程范例应该在这里应用。我不确定是不是这样。

但是要回答你的问题，首先让我指出CUDA中的线程同步只能在线程块中实现。所以我的注释只适用于这里。

设备代码中的主要同步机制是__syncthreads()。为了大致按照您所描述的方式使用它，我可以编写如下代码:

__syncthreads();
if (threadIdx.x < 100){
   // code in this block will only be executed by threads 0-99, all others do nothing
  }
__syncthreads();
if ((threadIdx.x > 99) && (threadIdx.x < 200)){
  // code in this block will only be executed by threads 100-199, all others do nothing
  }
// all threads can begin executing at this point

注意，即使线程块中的线程也不是都是同步执行的。SM (CUDA GPU中的线程块处理单元)通常将线程块分成32个线程组，称为warp ，这些warp实际上(或多或少)是同步执行的。然而，如果您出于某种原因想要这样做，我上面列出的代码仍然具有我所描述的在线程组之间排序执行方面的效果。