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

我正试图将几个任务“映射”到CUDA GPU。有n个任务需要处理。（看伪代码）CUDA线程/线程块之间的通信

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

由于某些原因，我将不得不使用不同线程块中的线程。

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

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

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

来源

2012-09-12 wwnigel

你将不得不使用原子操作这..这花费了很多。 – angainor

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

如果要在CUDA中实现，请采用int flag并在完成Do something后继续添加1，以便您可以在执行Do something之前和之后知道标志的更改。

如果我的问题不对，请评论。我会尽力改进答案。

来源

2012-09-14 18:43:50 Fr34K

使用__syncthreads()直接在线程块内同步。然而，线程块之间的同步更棘手 - 编程模型方法是分成两个内核。

如果你仔细想想，这是有道理的。执行模型（针对CUDA和OpenCL）适用于处理单元上执行的大量模块，但没有提及什么时候执行。这意味着一些块会执行，但其他块不会（他们会等待）。所以如果你有一个__syncblocks()，那么你会冒死锁的风险，因为那些已经执行的将会停止，但那些没有执行的将永远不会到达屏障。

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

根据你想要做的事情，经常会有另一种解决或解决问题的方法。

来源

2012-09-12 20:54:07 Tom

对不起，我想我没有把问题弄清楚。我想要做的是**请求一个线程等待来自另一个线程**的数据。例如，线程＃2需要线程＃1中的结果。这意味着线程＃2的一部分必须在线程＃1完成后启动。
所以我想知道是否有像“sem_post”和“sem_wait”，通过2线程之间的按摩功能？或者我将不得不编写（线程＃1）和读取（线程＃2）全局内存，尽管它效率很低？ @Mark Ebersole – wwnigel

在一个块中，您可以使用共享内存来进行线程间的通信（使用__syncthreads（）），在不同块的线程之间进行通信，请参阅我的答案：由于可能导致死锁，因此无法同步块。 – Tom

我明白了。仍然在这个例子中，如果线程＃2需要来自线程＃1的数据。（线程＃2检查“标志”全局内存，线程＃1在完成时写入“标志”全局内存）。然后，线程＃2将总是检查内存并占用总线，因此线程＃1不能成功地将“标志”写入内存。如果它们在同一个线程块中，我们可以使用__syncthreads（）来确保写入完成，然后启动线程＃2的其余部分？ – wwnigel

由于线程块可以按任意顺序进行调度，并且在它们之间没有简单的同步或通信方法，所以您要求的操作并不容易。从CUDA编程指南：

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

所以，如果你不能满足你在一个线程块中需要的所有通信，你将需要多个内核调用才能完成你想要的。

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

来源

2012-09-12 20:58:10

@Tom在4分钟内击败了我...... :) –

对不起，我想我没有把问题弄清楚。我试图做的是让一个线程等待来自另一个线程的数据。例如，线程＃2需要线程＃1中的结果。这意味着线程＃2的一部分必须在线程＃1完成后启动。
所以我想知道是否有像“sem_post”和“sem_wait”，通过2线程之间的按摩功能？或者我将不得不编写（线程＃1）和读取（线程＃2）全局内存，尽管它效率很低？ – wwnigel

谢谢你的回答。我重新描述了问题描述，现在看起来很清楚。你可以再次看看这个问题吗？ – wwnigel

这种问题最好由一个稍微不同的方法解决：

不分配固定的任务，你的线程，迫使你的线程等待，直到他们的任务变得可用（这是不可能的，因为CUDA线程不能阻塞）。

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

这仍然是棘手的实施和获得正确的案件，但至少是可能的。

来源

2012-09-12 22:59:25 tera

这样，我需要保留一个可用任务的列表，但在这个应用程序中它非常复杂。是否有可能在CUDA中实现“等待”和“标记”？ – wwnigel

不，线程无法等待任何超出其自身块的内容。 – tera

您可以通过保留两个单独的可用任务列表来简化任务列表的簿记工作 - 其中一个任务从中被处理，另一个任务在准备就绪时添加。一旦处理完第一个列表中的所有任务，就交换列表并启动一个新的内核。这会导致轻微的性能损失，但是使列表管理非常简单。 – tera

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

回答

相关问题