2013-12-16 114 views
0

本書報價:無法理解__syncthreads()

在CUDA,一個__syncthreads()聲明,如果存在的話必須由塊中的所有線程執行。當__syncthreads()被置於if語句中時,塊中的所有線程都會執行包含__syncthreads()或其中沒有任何一個的路徑。對於if-then-else語句,如果每條路徑都有一條__syncthreads()語句,則塊中的所有線程都會在then路徑上執行__syncthreads()或者所有線程都執行else路徑。兩個__syncthreads()是不同的屏障同步點。如果一個塊中的線程執行then路徑,而另一個執行else路徑,則它們將在不同的障礙同步點處等待。他們最終會永遠等待對方。程序員有責任編寫代碼,以滿足這些要求。

沒有給出ifif-else-then個案的例子,所以我無法理解這個概念。請用簡單的話來解釋我的情況。 PS:我是並行編程和CUDA的初學者。

在此先感謝。

回答

6

讓我們假設你有一個啓動的內核,其中一個線程塊由32個線程組成。

kernel<<<1,32>>>()

有內核的代碼如下:

__global__ void kernel() 
{ 
    if (threadIdx.x < 16) 
    { 
    // do something 
    __syncthreads(); 
    } 
    else 
    { 
    // do something 
    __snycthreads(); 
    } 
} 

第16個的線程塊的線程將運行if語句。其他16個else語句。如果前16個線程中的每個線程都到達__syncthreads,那麼它們會阻塞,直到整個線程塊達到語句爲止。但是這種情況永遠不會出現,因爲其他16個線程在else分支中出現問題,並且會出現死鎖。

您應該避免在不同的if和else分支中使用__syncthreads,或者您必須確保整個線程塊在同一分支中運行!

+0

好吧,它是這樣的兩個等待點,第一個16到達一個,最後16個到達其他,但這兩個等待點不知道eac其他。我是這麼想的嗎 – AbKDs

+0

是的,你可以用這種方式表達。前16個無法如何達到其他'__syncthreads'來填補聲明。 – hubs

+0

感謝幫助哥們:) – AbKDs