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

Posted 2023-03-23

【中文标题】如何在 CUDA 应用程序中正确应用线程同步？

【英文标题】：How to properly apply thread synchronization in CUDA app?

【发布时间】：2013-05-19 10:24:42

【问题描述】：

通常我在我的应用程序中偶尔使用线程同步，因为我并不经常需要这个功能。我不是真正的高级 C/C++ 程序员，但我也不是初学者。与 CPU 的强大功能相比，我开始学习 CUDA C，因为现在 GPU 的强大功能让我兴奋不已，我意识到 CUDA 编程主要是关于并行线程执行，有时需要适当的线程同步。事实上，我什至还不知道如何在 C 或 C++ 中应用线程同步。我上一次使用同步是大约 2 年前，当时我正在用 Java 编写这样的简单应用程序：

synchronized returnType functionName(parameters)

    ...

什么允许在一个 tmie 上仅由一个线程执行 'functionName' - 也就是说，该函数由不同的线程交替执行。现在回到 CUDA C，如果我有例如一个块中有 200 个线程，在 while 循环中运行代码：

while(some_condition)

    ...

如何使线程 彼此同步，线程 也彼此同步，但以线程 和 交替执行的方式应用同步（即前 100 个线程运行 'while' 的内容，然后接下来的 100 个线程运行 'while' 的内容，依此类推）？

【参考方案1】：

我认为您可能只需要learn more about cuda。您可能陷入了一个陷阱，认为您以前学习的编程范式应该在这里应用。我不确定是不是这样。

但要回答您的问题，首先让我指出 CUDA 中的线程同步只能在线程块内实现。所以我的 cmets 只适用于那里。

设备代码中的主要同步机制是__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，这些线程实际上（或多或少）以锁步方式执行。但是，我上面列出的代码仍然具有我所描述的效果，就线程组之间的顺序执行而言，如果您出于某种原因想要这样做的话。