如何避免在OpenCL中回讀

我正在使用OpenCL實現算法。我將多次循環使用C++，並每次調用同一個OpenCL內核。內核將生成下一次迭代的輸入數據和這些數據的數量。目前，我在每個循環中讀回這個數字以用於兩種用法：如何避免在OpenCL中回讀

我使用這個數字來決定下一個循環需要多少個工作項目;和
我用這個數字來決定何時退出循環（當數字是0）。

我發現閱讀需要大部分時間的循環。有什麼辦法可以避免它？

一般來說，如果您需要反覆調用一個內核，並且退出條件依賴於內核生成的結果（不是固定數字循環），那麼您如何有效地執行該操作？有什麼像OpenGL中的遮擋查詢，你可以做一些查詢，而不是從GPU回讀？

來源

2012-07-20 redpearl

你有幾個選擇這裏：

如果可能的話，你可以簡單地移動循環和條件到內核？使用一種方案，其中額外的工作項目根據當前迭代的輸入而不執行任何操作。
如果1.是不可能的，我建議您將由「決定」內核生成的數據存儲在緩衝區中，並使用該緩衝區「引導」其他內核。

這兩個選項都允許您跳過回讀。

來源

2012-07-20 19:42:59 Ani

如果我啓動很多工作項目（比如1000），並且其中大部分（比如說990）都不做任何事情。例如。他們測試一個條件並立即返回，如果它是假的。劑量它有類似的速度，只是啓動10個工作項目？還是會慢得多？ – redpearl 2012-07-20 19:46:10

我不明白爲什麼這會是一件壞事。只要沒有回讀，它應該只要工作組中最長的工作項目一樣長，是正確的？ – Ani 2012-07-20 19:47:31

我的真實程序比較複雜。我在每個循環中調用許多內核。我不認爲有可能將整個循環的工作實現到內核中。 :( – redpearl 2012-07-20 19:54:43

從GPU讀取數字內核總是需要10s - 1000s微秒或更多。

如果控制編號總是減少，您可以保留在全局內存中，並根據全局標識進行測試，並確定每次迭代時內核是否工作。使用全局內存屏障來同步所有線程...

kernel void x(global int * the_number, constant int max_iterations, ...) 
{ 
    int index = get_global_id(0); 
    int count = 0; // stops an infinite loop 

    while(index < the_number[0] && count < max_iterations) 
    { 
     count++; 
     // loop code follows 

     .... 

     // Use one thread decide what to do next 
     if (index == 0) 
     { 
      the_number[0] = ... next value 
     } 

     barrier(CLK_GLOBAL_MEM_FENCE); // Barrier to sync threads 
    } 
}

來源

2012-07-20 20:16:12

我只是完成了一些研究，我們必須解決這個確切的問題！

我們發現了兩件事情：

使用兩個（或更多）的緩衝區！在內核的第一次迭代中，對b1中的數據進行操作，然後對b2中的下一個操作，然後再對b1中的數據進行操作。在每個內核調用之間的中，回讀其他緩衝區的結果並檢查是否該停止迭代。內核花費的時間比讀取時間長時效果最好。使用一個分析工具來確保你沒有等待讀取（如果你是，增加緩衝區的數量）。
過度拍攝！爲每個內核添加一個完成檢查，並在將數據複製回來之前將其稱爲幾（100s）次。如果你的內核是低成本的，這可以很好地工作。

來源

2012-07-23 02:23:08

如何避免在OpenCL中回讀

回答

相關問題