CUDA __syncthreads（）和遞歸

我想用__syncthreads（）的遞歸像CUDA __syncthreads（）和遞歸

__device__ void foo(int k) { 
    if (some_condition) { 
    for (int i=0;i<8;i++) { 
     foo(i+k); // foo might take longer with some inputs 
     __syncthreads(); 
    } 
    } 
}

請問這個__syncthreads（）現在申請？我知道它只適用於一個區塊。據我瞭解，這適用於所有本地線程獨立的遞歸深度？但是如果我想確保這個__syncthreads（）到某個遞歸深度呢？這甚至有可能嗎？我可以檢查遞歸深度，但我相信這也不起作用。

有沒有可能的選擇？

我見過的，有用於CUDA設備> = 2.0

int __syncthreads_count(int predicate); 
int __syncthreads_and(int predicate); 
int __syncthreads_or(int predicate);

3個syncthread擴展版，但我不認爲他們會幫助，因爲他們看起來像一個原子計數器。

來源

2011-07-19 Pascal

我沒有確切的答案，因爲我自己從來沒有這樣做過，但只是爲了檢查您是否知道，您放入代碼中的some_condition必須對同一個塊中的所有線程計算相同的值，或者它會陷入僵局。 – jmsu

是的，這也是我所害怕的。 – Pascal

請你澄清你的問題嗎？我真的不明白你在這裏問什麼。 – Tom

如您所知，__syncthreads()只有在塊內的所有線程都達到屏障的情況下才是安全的。這意味着如果您在條件內調用__syncthreads()，則條件必須在塊內的所有線程上評估爲相同。

對於遞歸中的__syncthreads()，這意味着塊中的所有線程都必須執行遞歸到相同的深度，否則不是所有線程都會到達同一個障礙。

來源

2011-07-20 11:02:10 Tom

你的推理是有道理的，但我可以想象，因爲遞歸需要一個費米GPU，所以在遞歸深度上並不重要，只是在代碼中。它可能可以看看堆棧深度，是的，但是爲什麼呢，這會帶來很多潛在的問題（比如死鎖）。我試圖找到關於此的更多信息。這是在什麼地方定義的？最簡單也是最可能的解決方案是：**只是不要在遞歸中使用** – Pascal

我會更強烈地說：除非必須，否則不要在CUDA中使用遞歸。每個線程必須維護自己的堆棧，導致大量額外的片外內存訪問，如果您可以用迭代來替換遞歸，那麼這將不是必需的。如果你不能，那麼你可以在共享內存中維護一個更簡單的堆棧。或者您可能能夠在共享內存或寄存器中保留堆棧的前幾個級別，從而減少總的片外訪問（通常用於GPU光線跟蹤）。至於syncthreads（），在任何非發散代碼，遞歸或其他方面都是安全的。 – harrism

好的，謝謝澄清。我認爲cuda堆棧可能更有效率，那是一個糟糕的自制堆棧。因爲我沒有太多的接觸cuda，所以這是表達它最簡單的方法。我目前正在對上述代碼進行重寫（無遞歸和無堆棧），但我對遞歸中的__syncthreads行爲非常感興趣。因爲我使用__syncthreads（）啓用了此代碼，並且它的工作原理並令人驚訝地沒有發生死鎖。 – Pascal

有沒有可能的選擇？

是，不要使用遞歸的範式來表達你的功能邏輯

來源

2011-07-21 21:17:58 fabrizioM

當然，你說的對__syncthreads（）是真實的，它僅適用於本地線程塊中，因此你無法控制其他區塊正在發生什麼。減少的最佳方法是先對整個數組進行縮減，這個數組的大小通常等於塊的大小。然後，不要將數組複製回主機，而是調用另一個約減，它將有1個塊和線程類似於前一個調用中的塊數，然後將大小爲1的陣列從設備複製到主機。但要確保在兩次調用之間使用cudaThreadSynchronize（），除非生成第一個縮減，否則可以進行縮減。這是兩步減少，但它適用於我。

乾杯！ saif

來源

2011-07-22 07:11:01 saif

CUDA __syncthreads（）和遞歸

回答

相關問題