在异步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等待先前的所有启动完成，使用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被异步启动！主机继续运行，然后在kernel1完成之前启动kernel2！但是，kernel2直到kernel1完成后才会执行，因为它们都在流0（默认流）上启动。考虑以下替代方法：

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

绝对不需要这样做，因为设备已经同步了在同一流上启动的内核。

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

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 / kernel，则不要使用流。同样，如果您希望内核按特定顺序执行，则不要使用流。

也就是说，请记住，放入流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;
 }

无需显式的设备同步。