stream - cudaStreamWaitEvent似乎没有等待

我正在尝试编写一个小型演示程序，该程序具有两个正在运行的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会按顺序执行

调用内核k_A1

触发一半

等待某人触发HalfB

调用内核k_A2

和streamB会

等待某人触发一半

调用内核k_B1

触发halfB

调用内核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命令。但它不会产生相同的输出。我在这里俯瞰什么？

最佳答案

我认为这是因为在记录“halfB”(cudaEventRecord(halfB，streamB);)之前调用了“cudaStreamWaitEvent(streamA，halfB，0);”。可能cudaStreamWaitEvent调用正在搜索封闭的“halfB” ，而不是；由于找不到它，它只是悄悄地前进了。请参阅以下文档:



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

  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!

由配置文件确认:

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

关于stream - cudaStreamWaitEvent似乎没有等待，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/15501699/