摆脱异步 cuda 流执行期间的繁忙等待

Question

我正在寻找一种方法来摆脱闲置代码中主机线程中的忙碌等待（不要复制该代码，它仅显示了我的问题的一个想法，它有许多基本错误）：

cudaStream_t steams[S_N];
for (int i = 0; i < S_N; i++) {
    cudaStreamCreate(streams[i]);
}
int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
     while (true) {
         if (cudaStreamQuery(streams[sid])) == cudaSuccess) { //BUSY WAITING !!!!
             cudaMemcpyAssync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
             kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
             break;
         }
         sid = ++sid % S_N;
     }

}

有没有办法让主机线程空闲并以某种方式等待某个流完成，然后准备并运行另一个流？

编辑：我在代码中添加了 while(true)，以强调忙等待。现在我执行所有流，并检查它们中的哪一个完成了运行另一个新流。 cudaStreamSynchronize 等待特定流完成，但我想等待任何作为第一个完成工作的流。

EDIT2：我以休闲的方式摆脱了忙碌等待：

cudaStream_t steams[S_N];
for (int i = 0; i < S_N; i++) {
    cudaStreamCreate(streams[i]);
}
int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
    cudaMemcpyAssync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
    kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
    sid = ++sid % S_N;
}
for (int i = 0; i < S_N; i++) {
    cudaStreamSynchronize(streams[i]);
    cudaStreamDestroy(streams[i]);
}

但它似乎比在主机线程上忙等待的版本慢一点。我认为这是因为，现在我在流上静态分配作业，所以当一个流完成工作时，它处于空闲状态，直到每个流完成工作。之前的版本动态将工作分配到第一个空闲流，效率更高，但是宿主线程有忙等待。

Answer 1

真正的答案是使用 cudaThreadSynchronize 等待 all 之前的启动完成，cudaStreamSynchronize 等待某个流中的所有启动完成，cudaEventSynchronize 只等待某个流上的某个事件被记录。

但是，您需要先了解流和同步的工作原理，然后才能在代码中使用它们。

---

如果您根本不使用流会发生什么？考虑以下代码：

kernel <<< gridDim, blockDim >>> (d_data, DATA_STEP);
host_func1();
cudaThreadSynchronize();
host_func2();

内核启动，主机继续同时执行 host_func1 和内核。然后，主机和设备同步，即主机等待内核完成，然后继续执行 host_func2()。

现在，如果您有两个不同的内核怎么办？

kernel1 <<<gridDim, blockDim >>> (d_data + d1, DATA_STEP);
kernel2 <<<gridDim, blockDim >>> (d_data + d2, DATA_STEP);

kernel1 异步启动！主机继续运行，并且 kernel2 在 kernel1 完成之前启动！然而，kernel2 直到 在 kernel1 完成后才会执行，因为它们都在流 0（默认流）上启动。考虑以下替代方案：

kernel1 <<<gridDim, blockDim>>> (d_data + d1, DATA_STEP);
cudaThreadSynchronize();
kernel2 <<<gridDim, blockDim>>> (d_data + d2, DATA_STEP);

完全没有必要这样做，因为设备已经同步了在同一流上启动的内核。

所以，我认为您正在寻找的功能已经存在......因为内核总是在启动之前等待同一流中的先前启动完成（即使主机经过)。也就是说，如果您想等待 any 之前的启动完成，那么只需 不要 使用流。这段代码可以正常工作：

for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
    cudaMemcpyAsync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, 0);
    kernel<<<gridDim, blockDim, smSize, 0>>>(d_data, DATA_STEP);
 }

---

现在，开始直播。您可以使用流来管理并发设备执行。

将流视为队列。您可以将不同的 memcpy 调用和内核启动放入不同的队列中。然后，流 1 中的内核和流 2 中的启动是异步的！它们可以同时执行，也可以以任何顺序执行。如果您想确保设备上一次只执行一个 memcpy/内核，那么不要使用流。同样，如果您希望内核按特定顺序执行，则不要使用流。

也就是说，请记住，放入流 1 中的任何内容都是按顺序执行的，因此不必费心进行同步。同步是用于同步主机和设备调用，而不是两个不同的设备调用。因此，如果您想同时执行多个内核，因为它们使用不同的设备内存并且彼此没有影响，那么请使用流。比如……

cudaStream_t steams[S_N];
for (int i = 0; i < S_N; i++) {
    cudaStreamCreate(streams[i]);
}

int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
     cudaMemcpyAsync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
     kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
     sid = ++sid % S_N;
 }

无需显式设备同步。

Answer 2

我解决这个问题的想法是每个流都有一个主机线程。该主机线程将调用 cudaStreamSynchronize 以等待流命令完成。 不幸的是，在 CUDA 3.2 中这是不可能的，因为它只允许一个主机线程处理一个 CUDA 上下文，这意味着每个启用 CUDA 的 GPU 一个主机线程。

希望在 CUDA 4.0 中可以实现：CUDA 4.0 RC news

编辑：我在 CUDA 4.0 RC 中使用 open mp 进行了测试。我为每个 cuda 流创建了一个主机线程。它开始起作用了。

Answer 3

有：cudaEventRecord(event, stream) 和 cudaEventSynchronize(event)。参考手册http://developer.download.nvidia.com/compute/cuda/3_2/toolkit/docs/CUDA_Toolkit_Reference_Manual.pdf 包含所有详细信息。

编辑：BTW 流对于并发执行内核和内存传输很方便。为什么要通过等待当前流完成来序列化执行？

Answer 4

你想要cudaStreamSynchronize而不是cudaStreamQuery

int sid = 0;
for (int d = 0; d < DATA_SIZE; d+=DATA_STEP) {
     cudaStreamSynchronize(streams[sid]);
     cudaMemcpyAssync(d_data, h_data + d, DATA_STEP, cudaMemcpyHostToDevice, streams[sid]);
     kernel<<<gridDim, blockDim, smSize streams[sid]>>>(d_data, DATA_STEP);
     sid = ++sid % S_N;
}

（您还可以使用 cudaThreadSynchronize 等待所有流的启动，以及使用 cudaEventSynchronize 的事件以进行更高级的主机/设备同步。）

您可以进一步控制使用这些同步功能发生的等待类型。查看 cudaDeviceBlockingSync 标志和其他标志的参考手册。不过，默认值可能是您想要的。

Answer 5

您需要复制数据块并在不同的 for 循环中对该数据块执行内核。这样会更有效率。

像这样：

size = N*sizeof(float)/nStreams;

for (i=0; i<nStreams; i++){
    offset = i*N/nStreams;
    cudaMemcpyAsync(a_d+offset, a_h+offset, size, cudaMemcpyHostToDevice, stream[i]);
}


for (i=0; i<nStreams; i++){
    offset = i*N/nStreams;
    kernel<<<N(nThreads*nStreams), nThreads, 0, stream[i]>>> (a_d+offset);
}

通过这种方式，内存副本不必等待内核执行前一个流，反之亦然。