CUDA線程執行順序

我有一個CUDA程序下面的代碼：CUDA線程執行順序

#include <stdio.h> 

#define NUM_BLOCKS 4 
#define THREADS_PER_BLOCK 4 

__global__ void hello() 
{ 

    printf("Hello. I'm a thread %d in block %d\n", threadIdx.x, blockIdx.x); 

} 


int main(int argc,char **argv) 
{ 
    // launch the kernel 
    hello<<<NUM_BLOCKS, THREADS_PER_BLOCK>>>(); 

    // force the printf()s to flush 
    cudaDeviceSynchronize(); 

    return 0; 
}

，其中每個線程，將打印threadIdx.x和blockIdx.x。該計劃的一個可能的輸出是這樣的：

Hello. I'm a thread 0 in block 0 
Hello. I'm a thread 1 in block 0 
Hello. I'm a thread 2 in block 0 
Hello. I'm a thread 3 in block 0 
Hello. I'm a thread 0 in block 2 
Hello. I'm a thread 1 in block 2 
Hello. I'm a thread 2 in block 2 
Hello. I'm a thread 3 in block 2 
Hello. I'm a thread 0 in block 3 
Hello. I'm a thread 1 in block 3 
Hello. I'm a thread 2 in block 3 
Hello. I'm a thread 3 in block 3 
Hello. I'm a thread 0 in block 1 
Hello. I'm a thread 1 in block 1 
Hello. I'm a thread 2 in block 1 
Hello. I'm a thread 3 in block 1

運行程序幾次，我得到了類似的結果，除了塊順序是隨機的。例如，在上面的輸出中，我們有這個塊順序0,2,3,1。再次運行問題我得到1,2,3,0。這是預期的。但是，每個塊中的線程順序始終爲0,1,2,3。這是爲什麼發生？我認爲這也是隨機的。

我試圖改變我的代碼，強制每個塊中的線程0花費更長的時間執行。我這樣做是這樣的：

__global__ void hello() 
{ 

    if (threadIdx.x == 0) 
    { 
     int k = 0; 
     for (int i = 0; i < 1000000; i++) 
     { 
      k = k + 1; 
     } 
    } 

    printf("Hello. I'm a thread %d in block %d\n", threadIdx.x, blockIdx.x); 

}

我希望線程爲了1,2,3，0。然而，我得到了類似的結果，在一個我在上面，其中線順序總是0顯示，1 ，2，3.爲什麼會發生這種情況？

來源

2015-05-25 Blue

你確定這是C語言嗎？這應該被稱爲類似CUDA_的_C方言。這既不是C，也不是C++。 – ForceBru

好的，同意。我刪除了該標籤以避免混淆。謝謝。 – Blue

您的代碼嘗試在最後的內核中減慢線程0是完全沒有意義的。編譯器足夠聰明，可以優化整個循環。 – talonmies

但是，每個塊中的線程順序始終爲0,1,2,3。這是爲什麼發生？我認爲這將是隨機的太

由於每塊4個線程你只推出每塊一個經。 A warp是CUDA中的執行單元（以及調度和資源分配），而不是線程。目前，一個warp包含32個線程。

這意味着每個塊的所有4個線程（因爲在這種情況下沒有條件行爲）正在執行鎖步。當它們到達printf函數調用時，它們都在同一行代碼,中以鎖步執行對該函數的調用。

因此，在這種情況下，CUDA運行時如何調度這些「同時」函數調用？這個問題的答案沒有具體說明，但它不是「隨機的」。因此，對於一個warp內的操作的調度順序不會隨着運行而改變是合理的。

如果啓動足夠的線程爲每個塊創建多個warp，並且可能還包含一些其他代碼以分散或「隨機化」warp之間的行爲，則應該能夠看到printf操作是從「隨機順序。

來源

2015-05-25 18:28:15

謝謝，這是對我的問題的非常準確的答案。我將其標記爲答案，因爲我學到了很多東西。但是，如果我將代碼分散或「隨機化」，會發生什麼？我有興趣進一步瞭解鎖步部分。所以假設線程0進入一個if語句，並且線程1到31不（在同一個warp中）。假設線程1-31已準備好繼續，它們是否會等待線程0完成if語句並以鎖步方式繼續？再次感謝您的回答。 – Blue

是的，對於if-then-else構造，滿足if條件的warp中的所有線程將遵循if路徑，而其餘線程爲「inactive」。然後，滿足if條件的線程將變爲「不活動」，並且不滿足if條件的線程將被激活，並且這些線程將遵循else路徑。在完成這個序列時，warp中的所有線程都會隱式地重新同步，然後繼續以lockstep執行。（事實上，if路徑首先發生還是其他路徑首先發生的確切順序是未指定的。） –

完美，我明白了。謝謝。 – Blue

要回答您的問題的第二部分，當控制流在if聲明中發生分歧時，threadIdx.x != 0僅在if語句後的收斂點處等待的線程。它們不會繼續執行printf語句，直到線程0完成if塊爲止。

來源

2015-05-25 19:00:25

感謝您的評論，因爲它幫助我更快地理解了Robert Crovella的答案。 – Blue

CUDA線程執行順序

回答

相關問題