CUDA 线程/线程块之间的通信答案

【问题标题】：Communication between CUDA threads/thread blocksCUDA 线程/线程块之间的通信
【发布时间】：2012-09-05 22:21:21
【问题描述】：

我正在尝试将一些任务“映射”到 CUDA GPU。有 n 个任务要处理。（见伪代码）

malloc an boolean array flag[n] and initialize it as false.
for each work-group in parallel do
    while there are still unfinished tasks do
        Do something;
        for a few j_1, j_2, .. j_m (j_i<k) do
            Wait until task j_i is finished; [ while(flag[j_i]) ;  ]
            Do Something;
        end for
        Do something;
        Mark task k finished;  [  flag[k] = true;  ]
    end while
end for

出于某种原因，我将不得不在不同的线程块中使用线程。

问题是如何在CUDA中实现等待任务j_i完成；和标记任务k完成；。我的实现是使用布尔数组作为标志。然后在任务完成后设置标志，并读取标志以检查任务是否完成。

但它只适用于小案例，一个大案例，GPU 因未知原因而崩溃。有没有更好的方法来实现 CUDA 中的 Wait 和 Mark。

这基本上是CUDA上的线程间通信问题。

【问题讨论】：

您将不得不为此使用原子操作.. 这会花费很多。

标签： cuda synchronization gpu

【解决方案1】：

使用__syncthreads() 在线程块内同步很简单。但是线程块之间的同步比较棘手——编程模型方法是分成两个内核。

如果你仔细想想，这是有道理的。执行模型（对于 CUDA 和 OpenCL）适用于在处理单元上执行的一大堆块，但没有说明何时执行。这意味着一些块将被执行，但其他块不会（它们将等待）。因此，如果您有__syncblocks()，那么您将面临死锁的风险，因为那些已经执行的将停止，但那些未执行的将永远不会到达障碍。

您可以在块之间共享信息（例如，使用全局内存和原子），但不能共享全局同步。

根据您尝试执行的操作，通常有另一种解决或分解问题的方法。

【讨论】：

对不起，我想我没有把问题说清楚。我想要做的是要求一个线程等待来自另一个线程的数据。例如，线程#2 需要线程#1 中的结果。这意味着线程#2 的一部分必须在线程#1 完成后启动。
所以我想知道有没有像“sem_post”和“sem_wait”这样在两个线程之间传递按摩的函数？或者我将不得不写入（线程＃1）和读取（线程＃2）全局内存，即使它非常低效？ @Mark Ebersole
在一个块中，您可以使用共享内存在线程之间（使用 __syncthreads()）、不同块中的线程之间进行通信，请参阅我的回答：您不能同步块，因为这可能导致死锁。
我明白了。仍然在示例中，如果线程#2 需要来自线程#1 的数据。（线程#2 检查“flag”全局内存，线程#1 完成后写入“flag”全局内存）。那么线程#2 将一直检查内存并占用总线，因此线程#1 无法成功地将“标志”写入内存。如果它们在同一个线程块中，我们可以使用 __syncthreads() 确保写入完成，然后启动线程#2的其余部分？
感谢您的回答。我改写问题描述，现在看起来很清楚了。你能帮忙再看看这个问题吗？

【解决方案2】：

您所要求的并不容易完成，因为线程块可以按任何顺序安排，并且没有简单的方法可以在它们之间进行同步或通信。来自 CUDA 编程指南：

对于并行工作负载，在算法中由于某些线程需要同步以相互共享数据而破坏并行性的点，有两种情况：这些线程属于同一个块，在这种情况下它们应该使用 __syncthreads() 并在同一个内核调用中通过共享内存共享数据，或者它们属于不同的块，在这种情况下，它们必须使用两个单独的内核调用通过全局内存共享数据，一个用于写入，一个用于从全局读取记忆。第二种情况不太理想，因为它增加了额外内核调用和全局内存流量的开销。因此，应该通过将算法映射到 CUDA 编程模型来最小化它的发生，使需要线程间通信的计算尽可能在单个线程块中执行。

因此，如果您无法在一个线程块中满足您所需的所有通信，则需要进行多次内核调用才能完成您想要的。

我不认为与 OpenCL 有任何区别，但我也不在 OpenCL 中工作。

【讨论】：

而@Tom 比我快 4 分钟... :)
对不起，我想我没有把问题说清楚。我想要做的是让一个线程等待来自另一个线程的数据。例如，线程#2 需要线程#1 中的结果。这意味着线程#2 的一部分必须在线程#1 完成后启动。
所以我想知道有没有像“sem_post”和“sem_wait”这样在两个线程之间传递按摩的函数？或者我将不得不写（线程＃1）和读取（线程＃2）全局内存，即使它非常低效？
感谢您的回答。我改写问题描述，现在看起来很清楚了。你能帮忙再看看这个问题吗？

【解决方案3】：

这类问题最好通过稍微不同的方法来解决：

不要将固定任务分配给您的线程，从而强制您的线程等待它们的任务可用（这在 CUDA 中是不可能的，因为线程无法阻塞）。

相反，保留一个可用任务列表（使用原子操作）并让每个线程从该列表中获取一个任务。

这仍然很难实现并正确处理极端情况，但至少它是可能的。

【讨论】：

这样，我需要保留一个可用任务的列表，但是在这个应用程序中非常复杂。是否可以在CUDA中实现“等待”和“标记”？
不，线程无法等待它自己的块之外的任何东西。
您可以通过保留两个单独的可用任务列表来大大简化任务列表的簿记 - 一个用于处理任务，另一个用于在任务准备就绪时添加任务.处理完第一个列表中的所有任务后，交换列表并启动新内核。这会导致轻微的性能损失，但会使列表管理变得非常简单。

【解决方案4】：

我认为您不需要在 CUDA 中实现。每件事都可以在 CPU 上实现。您正在等待一项任务完成，然后随机执行另一项任务。如果你想在 CUDA 中实现，你不需要等待所有的标志都为真。您最初知道所有标志都是错误的。所以只需为所有线程并行实现Do something 并将标志更改为true。

如果你想在CUDA中实现，取int flag，在完成Do something之后继续加1，这样你就可以知道Do something前后flag的变化。

如果我的问题有误，请发表评论。我会努力改进答案。

【讨论】：