擺脫異步執行期間的繁忙等待

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

沒有顯式的設備同步。