5

我正在尝试编写一个小型演示程序，该程序有两个 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. 触发半A
  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 命令。但它不会产生相同的输出。我在这里俯瞰什么？

4

1 回答 1

7

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

该流将仅等待对onstream的最新主机调用完成。一旦此调用返回，任何函数（包括和）都可能再次被调用，并且后续调用不会对 产生任何影响。cudaEventRecord()eventcudaEventRecord()cudaEventDestroy()eventstream

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

  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!

分析证实了这一点：

在此处输入图像描述

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

于 2013-03-23T22:17:49.280 回答