Re: [問題] cuda race condition

看板C_and_CPP (C/C++)作者littleshan (我要加入劍道社!)時間16年前 (2010/04/22 12:52)推噓8(8推 0噓 14→)

留言22則, 8人參與討論串2/2 (看更多)

※ 引述《evilned (君千殤)》之銘言： : 小弟目前在練習 cuda 程式 : 發現在多 thread 同時進行寫入下會有 race condition 問題 : ex : 計算整個 array 裡面某個值的數量 : 變的要用另一個相同大小 array 去判斷 : 這樣下來反而速度是拖慢的 : 請問有較快的解決方法嗎? : 說明一下我程式 : : 這個 array 是 1000 x 1000 大小 : 所以我開了 1000 x 1000 threads 下去判斷 : if (idx < 1000 x 1000 && array[idx] == 1) : sum++; : 就發生~ race condition 問題了~ 一次跑完1000x1000筆判斷 : thread 會搶寫入空間 , 導致sum錯誤關於這個問題，最標準的解決方法是 reduce 方法是這樣的 1. 開另外一個 1000x1000 的陣列，裡面存 array[i] 是否為 1 的比較結果是則存 1 不是則存 0 2. 把這個陣列內的值加總，結果即為 array 中為 1 的元素個數。加總的時候要注意到，為了能最佳化 GPU 的平行運算能力，通常會跑多個 pass，每個 pass 只把兩個元素相加具體來講，大概像這樣： __global__ void check(int* array, int* buf, int size) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if(idx >= size) return; if(data[idx] == 1) buf[idx] = 1; else buf[idx] = 0; } __global__ void reduce(int* buf, int range, int size) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if(idx >= range || idx+range >= size) return; buf[idx] += buf[idx+range]; } int main() { // ... int size = 1024 * 1024; int nblock = size / 64; int* buf; cudaMalloc((void**)&buf, sizeof(int)*size); check<<<nblock, 64>>>(array, buf, size); for(int range = size/2; range > 0; range /= 2){ nblock = (range-1) / 64 + 1; // get the correct block number reduce<<<nblock, 64>>>(buf, range, size); } // get the result int result; cudaMemcpy(&result, buf, sizeof(int), cudaMemcpyDeviceToHost); // ... } 另一個方法是 atomicAdd 我「感覺」你想要的好像是這個答案畢竟用 atomicAdd 你幾乎不用改你的程式碼： __global__ void check(int* array, int* result, int size) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if(idx < size && array[idx] == 1) atomicAdd(result, 1); } 不過很不幸的，atomic operation 的本質就是循序的、非平行化的，也就是慢儘管你會覺得用 atomicAdd 看起來很簡單，而用 reduce 很複雜但就我自己的實驗結果使用 atomicAdd 所花的時間超出 reduce 的十倍以上甚至你直接用 CPU 的迴圈都比它還快而且這還是在 array 中符合條件的元素僅占 3% 的情況下如果有更多元素為 1，atomicAdd 所花的時間會更久結論就是寫 CUDA 程式時要把平常寫循序程式的思維轉換到平行的思維下而不是直接把 CPU 程式中的迴圈拿來當 CUDA kernel 就完事 -- ※ 發信站: 批踢踢實業坊(ptt.cc) ◆ From: 140.112.29.108

推

VictorTom

04/22 13:05, , 1^F

04/22 13:05, 1^F

推

stupidbear

04/22 13:08, , 2^F

04/22 13:08, 2^F

→

stupidbear

04/22 13:08, , 3^F

04/22 13:08, 3^F