cudaStreamWaitEvent 似乎没有等待

Question

我正在尝试编写一个小型演示程序，该程序有两个 cuda 流正在进行，并且受事件控制，相互等待。到目前为止，这个程序看起来像这样：

// event.cu
#include <iostream>
#include <cstdio>
#include <cuda_runtime.h>
#include <cuda.h>

using namespace std;

__global__ void k_A1() { printf("\tHi! I am Kernel A1.\n"); }
__global__ void k_B1() { printf("\tHi! I am Kernel B1.\n"); }
__global__ void k_A2() { printf("\tHi! I am Kernel A2.\n"); }
__global__ void k_B2() { printf("\tHi! I am Kernel B2.\n"); }

int main()
{
  cudaStream_t streamA, streamB;
  cudaEvent_t halfA, halfB;
  cudaStreamCreate(&streamA);
  cudaStreamCreate(&streamB);
  cudaEventCreate(&halfA);
  cudaEventCreate(&halfB);

  cout << "Here is the plan:" << endl <<
    "Stream A: A1, launch 'HalfA', wait for 'HalfB', A2." << endl <<
    "Stream B: Wait for 'HalfA', B1, launch 'HalfB', B2." << endl <<
    "I would expect: A1,B1, (A2 and B2 running concurrently)." << endl;

  k_A1<<<1,1,0,streamA>>>(); // A1!
  cudaEventRecord(halfA,streamA); // StreamA triggers halfA!
  cudaStreamWaitEvent(streamA,halfB,0); // StreamA waits for halfB.
  k_A2<<<1,1,0,streamA>>>(); // A2!

  cudaStreamWaitEvent(streamB,halfA,0); // StreamB waits, for halfA.
  k_B1<<<1,1,0,streamB>>>(); // B1!
  cudaEventRecord(halfB,streamB); // StreamB triggers halfB!
  k_B2<<<1,1,0,streamB>>>(); // B2!

  cudaEventDestroy(halfB);
  cudaEventDestroy(halfA);
  cudaStreamDestroy(streamB);
  cudaStreamDestroy(streamA);

  cout << "All has been started. Synchronize!" << endl;
  cudaDeviceSynchronize();
  return 0;
}

我对 CUDA 流的掌握如下： 流是一种可以添加任务的列表。这些任务是按顺序处理的。因此，在我的程序中，我可以放心，streamA 会按顺序进行

1. 调用内核k_A1
2. 触发 halfA
3. 等待有人触发 halfB
4. 调用内核k_A2

而streamB会

1. 等待有人触发halfA
2. 调用内核k_B1
3. 触发器半B
4. 调用内核k_B2

通常两个流可能彼此异步运行。但是，我想阻止 streamB 直到 A1 完成，然后阻止 streamA 直到 B1 完成。

这似乎没有那么简单。在我的带有 Tesla M2090 (CC 2.0) 的 Ubuntu 上，

nvcc -arch=sm_20 event.cu && ./a.out

是

Here is the plan:
Stream A: A1, launch 'HalfA', wait for 'HalfB', A2.
Stream B: Wait for 'HalfA', B1, launch 'HalfB', B2.
I would expect: A1,B1, (A2 and B2 running concurrently).
All has been started. Synchronize!
        Hi! I am Kernel A1.
        Hi! I am Kernel A2.
        Hi! I am Kernel B1.
        Hi! I am Kernel B2.

我真的希望 B1 在 cudaEventRecord(halfB,streamB) 之前完成。尽管如此，流 A 显然不会等待 B1 的完成，因此也不会等待 halfB 的记录。

更重要的是：如果我完全删除 cudaEventRecord 命令，我希望程序锁定 cudaStreamWait 命令。但它不会产生相同的输出。我在这里俯瞰什么？

Answer 1

我认为这是因为在记录“halfB”之前调用了“cudaStreamWaitEvent(streamA,halfB,0);”(cudaEventRecord(halfB,streamB);)。很可能 cudaStreamWaitEvent 调用正在搜索它之前的封闭“halfB”之前；既然没有被发现，它就静静地往前走。请参阅以下文档：

流stream 将仅等待在event 上对cudaEventRecord() 的最近一次主机调用完成。此调用返回后，任何函数（包括cudaEventRecord() 和cudaEventDestroy()）都可能再次在event 上调用，后续调用不会对stream 产生任何影响。

如果您必须进行深度优先编码，我找不到解决方案；但是，以下代码可能会导致您想要的结果：

  k_A1<<<1,1,0,streamA>>>(d); // A1!
  cudaEventRecord(halfA,streamA); // StreamA triggers halfA!
  cudaStreamWaitEvent(streamB,halfA,0); // StreamB waits, for halfA.
  k_B1<<<1,1,0,streamB>>>(d); // B1!
  cudaEventRecord(halfB,streamB); // StreamB triggers halfB!
  cudaStreamWaitEvent(streamA,halfB,0); // StreamA waits for halfB.
  k_A2<<<1,1,0,streamA>>>(d); // A2!
  k_B2<<<1,1,0,streamB>>>(d); // B2!

已通过分析确认：

请注意，我更改了内核接口。