Re: [問題] cuda race condition

[littleshan (我要加入劍道社!)]

※ 引述《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:推l大:)
推 stupidbear:void裡面為什麼會有return阿?通常像這種累加的動作
→ stupidbear:我都會丟給cpu去算,除非資料量真的很大,不然gpu可能
→ stupidbear:還比較慢@@
推 VictorTom:void func裡也可以return啊, 代表func執行到return就不
→ VictorTom:執行下去結束了啊, return不是非回傳值不可的@_@"
推 stupidbear:所以void中的return就是交回執行權摟?不回傳值也可以
→ littleshan:看資料量的大小，以 1024x1024 這樣的資料量來說
→ littleshan:我用 GPU 的速度約是 CPU 的 12 倍
→ littleshan:大量資料的累加本來就很適合平行化的
→ blackwindy:現在感覺CUDA比較多人在討論了...去年我專題弄的要死..
→ blackwindy:給原PO，『用空間換時間』
→ blackwindy:然後請去想想演算法，單累加應該不需要用到這麼複雜
→ blackwindy:喔重看了一下，這樣確實是正確的寫法，眼殘了抱歉
推 BIGHEADXYZ:這個不就是tree reduction嗎?
推 evilned:感謝~~ 就是想知道atomicAdd ~ 非常感謝
→ sunneo:atomic真的會很慢的
→ littleshan:我開始後悔發這篇文了，atomicAdd真的不該用在這個case
推 VictorTom:l大別這麼說嘛, 這篇文很有幫助啊, 不然幫atomicAdd加
→ VictorTom:個警訊好了XD
推 xxxx9659:誰知道 if(idx >= range || idx+range >= size)
→ xxxx9659:|| 後面的 idx+range >= size 在做什麼 看不懂＠＠