CUDA問題—單執行緒塊 VS 多執行緒塊

我將我最近遇到的問題簡化成以下的問題:

const N = 32 // even number
int A[N/2] = {0};
int B[N] = {0};

// serial version
void func(int* d_A, int* d_B){
  for(int b_i=0; b_i<N; ++b_i){
    a_i = (b_i%2) == 0? b_i : (b_i - 1);   
    A[a_i] += B[b_i];
  }
}

// 版本一  
__global__ void func1(int* d_A, int* d_B){
  for(int b_i=0; b_i<N; ++b_i){
    a_i = (b_i%2) == 0? b_i : (b_i - 1);   
    A[a_i] += B[b_i];
  }
}
func1<<<1,1>>>(A, B);


// 版本二 
__global__ void func2(int* d_A, int* d_B){
  bid = blockIdx.x;
  a_i = (bid%2) == 0? bid : (bid - 1);   
  A[a_i] += B[bid];
}
func2<<<n,1>>>(A, B);

版本一和版本二的差別在於,一個使用單執行緒塊,另一個使用多執行緒塊。理論上來說版本二會有比較好的效率。
實驗後的結果版本二會有比較大的數值誤差(相對誤差大於0.01),想請問這種現象正常嗎?
PS 這邊我將問題簡化一個執行緒塊只用一個執行緒,在我的實際問題中,A會是一個array of matrix,因此一個執行緒塊會用多個執行緒。

a_i = (bid % 2) == 0 ? bid : (bid - 1);
可以改成 (當 bid 是整數型態的話)
a_i = bid / 2 * 2; 

代表會有兩個同時碰到 A[a_i] 從而導致 race condition
例如:
bid = 2, 3 都會用到 A[2]
最終應該要是 A[2] = A[2]_0 + B[2] + B[3]
不能保證順序,所以最終結果會不確定
用 E2 跟 E3 來表示執行 A[2] += B[2] 跟 A[2] += B[3]

  • A[2] + B[2]
    E2 讀 A[2]_0 跟 E3 讀 A[2]_0 (兩者皆在寫入前讀取資料)
    E3 寫入 A[2]_0 + B[3]
    E2 寫入 A[2]_0 + B[2] (覆蓋掉 E3 寫入的內容)
  • A[2] + B[3]
    跟上述依樣只是最後寫入順序不同,
  • A[2] + B[2] + B[3]
    E2/E3 讀取 A[2]_0 並寫入 A[2]_0 + B[2]/B[3]
    另一個 E3/E2 從 A[2] 讀取到 A[2]_0 + B[2]/B[3] ,加上 B[3]/B[2] 並寫入最後答案。

可能解決方案

  • 一個 block 直接處理 B[2*bid]B[2*bid+1]
  • 用 atomicAdd 確保同一個讀取寫入是 Atomic
  • 因為目前是不同的 block 分別處理,所以不能用 shuffle;但如果是在同一 block 中就可以用 shuffle 來加總,最後其一輸出答案至 global memory。

感謝回答!
想再追問,在這個例子可以事先預測哪些memory access會造成race condition,如果今天有一個情況是完全無法事先預測,並且由於資料型態非primitive type,atomicAdd也無法使用的話,有沒有什麼其他方法可以解決呢?比如可以lock不同threads block之間的操作。

比如說如果用上面的例子來說的話,大概像這樣

a_i = rand() % (N/2)
bid = blockIdx.x
A[a_i] += B[bid]

就是無法知道哪些block需要一起處理,然後atomicAdd也無法使用,這時候有沒有什麼方法可以避免“無法預期的”race condtion