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

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

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

希望，在CUDA 4.0，将有可能：CUDA 4.0 RC news

编辑：我在CUDA 4.0 RC测试，使用开放的熔点。我为每个cuda流创建了一个主机线程。它开始工作。

Answer 2

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

编辑：BTW流是得心应手内核和内存传输的并发执行。为什么要通过等待当前流完成序列化执行？

Answer 3

相反的cudaStreamQuery的，你想cudaStreamSynchronize

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 4

您需要复制数据块并在该数据块上执行不同的内核以获得循环。这会更有效率。

这样的：

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); 
} 

这样的内存拷贝不必等待前面的流，反之亦然内核执行。

Answer 5

真正的答案是使用的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 is asychronously！主机继续运行，并在kernel1完成之前启动kernel2！然而，在内核1完成之后，kernel2将不会执行到，因为它们都在流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 /内核，然后不要使用的流。同样，如果你想内核在一个特定的顺序来执行，那么不要使用的流。

也就是说，记住，任何东西放到一个流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; 
} 

没有显式的设备同步。