CUDA race condition

綜合討論
1

這是回答 PTT 上的一篇 CUDA 問題 [問題] CUDA 程式 - 看板 C_and_CPP - 批踢踢實業坊
下面簡略地敘述

// dim3 gridsize(1, 1, 1);
// dim3 blocksize(600, 1, 1);
for (int j = 0; j < 800; j+= 1){
  sum[(blockDim.x*bdx + tdx)] += data[600*j + (blockDim.x*bdx +tdx)];
}

// dim3 gridsize(40, 1, 1);
// dim3 blocksize(600, 1, 1);
for (int j = 0; j < 800; j+= 40){
  sum[tdx] += data[600*(j + bdx) +tdx];
}

第二個就會執行錯誤
猜測是: 例如, sum[1]無法同時處理40筆資料

答案會錯的原因是 race condition
現在有 40 threads 可能會同時處理 sum += data → (讀 sum, 算 sum+data 並把他寫回去)
但這個並非 atomic ，也就是同一條讀和寫中間有可能會被其他插入
例如: sum 一開始是 0，A: sum += 10, B: sum += 3
有可能是

  1. A 讀 sum = 0
  2. B 讀 sum = 0
  3. B 寫入 sum = 0 + 3 = 3
  4. A 寫入 sum = 0 (因為一開始讀到是 0) + 10 = 10

最後答案就變成 10 而非 13
根據執行順序也有可能是 3 或 13

可以用 atomicAdd (但可能因為衝突太多導致很慢)
或者不同切法，例如用 (記得處理邊界) sum[threadIdx.x + blockIdx.x * blockDim.x] 然後平行後面加總 data ，記得可以先放入 register 最後再放入 global memory 會比較好。