CUDA線程發散和分支，示例

Question

我有一些例子給了我一些奇怪的問題： 我產生了一個線程分歧，但我無法確定哪個分支或哪些語句是首先計算的？

第一個例子：
我有以下的內核，這是我在1塊開始由2個線程。 [0] = 0，1 = 0。

__global__ void branchTest_kernel(float* a){ 

    int tx = threadIdx.x; 

    if(tx==0){     // or tx==1 
    a[1] = a[0] + 1; (a) 
    }else if(tx==1){    // or tx==0 
    a[0] = a[1] + 1;;   (b) 
    } 
} 

輸出

a[0] = 1 
a[1] = 1 

我assum因爲兩個線程都在一個經紗，它們在鎖步執行，和（a）和（b）在同一時間同時讀取一個[ 0]和1。

第二個例子：
完全一樣的第一，但現在除去否則，如果部分：

__global__ void branchTest_kernel(float* a){ 

    int tx = threadIdx.x; 

    if(tx==0){ 
    a[1] = a[0] + 1; (a) 
    }else{ 
    a[0] = a[1] + 1; (b) 
    } 


} 

輸出

a[0] = 1 
a[1] = 2 

現在是什麼原因導致這種行爲突然（b）是第一個，（a）第二個......（最可能是內部分支） 有人可以解釋優先rul ES是分支機構？或者在哪裏可以找到這些信息？行爲是不確定的 - Gauss Seidel See Figure 3, (a) diagonal block

Answer 1

有在CUDA Warp中的分支執行順序沒有優先級規則：

我高斯 - 塞德爾求解器的實現過程中遇到的這個例子。編譯器，彙編器和JIT運行時可以自由地按照他們認爲合適的順序對指令進行重新排序，並且絕對不能試圖依賴經驗推導的任何順序，因爲它可能會發生變化（就像您發現的那樣）。在這種情況下強制形式正確性的唯一方法是使用原子內存訪問操作，該操作將強制序列化。更好的是，尋找另一種算法。

在Gauss-Seidel案例中，正統方法是對矩陣或計算網格的圖分解中的每種顏色使用單獨的內核啓動。