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. 触发一半
3. 等待某人触发HalfB
4. 调用内核k_A2

和streamB会

1. 等待某人触发一半
2. 调用内核k_B1
3. 触发一半
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”（cudaEventRecord（halfB，streamB）;）之前调用了“ cudaStreamWaitEvent（streamA，halfB，0）;”。可能cudaStreamWaitEvent调用在其之前搜索了已关闭的“ halfB” 。由于找不到它，它只是悄悄地前进。请参阅以下文档：

流stream只会等待对cudaEventRecord()on的最新主机调用完成event。返回此调用后，可以再次调用任何函数（包括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!

由配置文件确认：

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