多个线程访问同一CUDA流

Question

我正在编写一个混合cpu-gpu程序，需要多个cpu线程访问多个gpu。CUDA Stream是线程安全的吗？具体来说，我想知道以下做法是否正确：

// two threads concurrently enter cuda device 1 and 
// launch kernel on the same stream

std::thread t1([&](){
  cudaSetDevice(1);
  cudaEventRecord(begin_t1, stream);
  kernel<<<mygrid, myblock, 0, stream>>>(...);
  cudaEventRecord(end_t1, stream);
});

std::thread t2([&](){
  cudaSetDevice(1);
  cudaEventRecord(begin_t2, stream);
  kernel<<<mygrid, myblock, 0, stream>>>(...);
  cudaEventRecord(end_t2, stream);
});

- Jes

---

似乎表明CUDA是线程安全的：https://devblogs.nvidia.com/cuda-pro-tip-always-set-current-device-avoid-multithreading-bugs/ - Alan Birtles

在这种情况下，CUDA是线程安全的，您的代码应该能够顺利执行。然而，并没有任何东西来排序两个线程将工作放入单一流的顺序。因此，这两个线程之间发出的工作顺序可能是任意的。例如，线程2的事件可能捕获了两个内核启动，以选择一个示例。CUDA不会为您解决这个问题，也不会为发出工作到流中的线程提供某种锁。 - Robert Crovella

Answer 1

多个主机线程访问和使用同一流是合法的。

然而，CUDA没有任何保证不同线程的操作顺序。因此，对于这个流，下面的操作序列是可能的：

begin_t1,kernel,end_t1,begin_t2,kernel,end_t2

但这也是可能的：

begin_t1,begin_t2,kernel,kernel,end_t1,end_t2

如果您希望在线程之间强制排序，您需要使用您正在使用的线程系统提供的机制。

- Robert Crovella