[問題] cuda新手關於平行化的問題

[usshan (漢典)]

開發平台(Platform): (Ex: VC++, GCC, Linux, ...) 
windows 7  VC2012

額外使用到的函數庫(Library Used): (Ex: OpenGL, ...) 


問題(Question)：

原本的程式碼的部分
每條thread中會產生一個值及它所在的xy坐標,分別存進result,bx,by
每組陣列有 x1*y1 筆資料
這三組陣列會被複製回host端
並在host端找result陣列中的最小值 以及對應的xy坐標
這部分已經可以順利執行找到所要的答案

但因其他需求
我想在device端就將result[ x1 * y1 ]的最小值找到
並存在result[0],bx[0],by[0]複製回host

但存在result[0]的結果 檢查後大部分是最小值
小部分不正確的值  不是最小值 但他會很接近最小值


餵入的資料(Input)：
假設產生以下資料
result[1]=100
result[2]=15
result[3]=3
result[4]=99

預期的正確結果(Expected Output)：

result[0]=3

錯誤結果(Wrong Output)：

result[0]=15


程式碼(Code)：(請善用置底文網頁, 記得排版) 
__global__ static void test(略)
{
                const int x =blockIdx.x;
                const int y =threadIdx.x;
                const int xid = blockDim.x;

                中間過程code略過

                result[y+xid*x]=cost;
                bx[y+xid*x]=x;
                by[y+xid*x]=y;
               //以上是計算完  存好的資料
               //下面是找最小值的過程,也是我有問題的地方

                if(result[0]>cost )
                {
                        result[0]=cost;
                        bx[0]=x;
                        by[0]=y;

                }
                else if(result[0]<cost && (bx[0]==0&&by[0]==0) && result[0]==0)
                {
                        result[0]=cost;
                        bx[0]=x;
                        by[0]=y;
                }
}

for(略)
{
        for(略)
                test<<<(x1),y1, sizeof(int)  >>>(略)

}

補充說明(Supplement)：

略過的部分code 我是覺得不影響執行結果
為了方便閱讀 就省略掉

--
※ 發信站: 批踢踢實業坊(ptt.cc), 來自: 1.173.26.90
※ 文章網址: https://www.ptt.cc/bbs/C_and_CPP/M.1441385731.A.487.html
推 johnjohnlin: race condition, you need atomicMin
→ usshan: 那再請問 使用atomic的函式  要include哪些東西
→ usshan: 目前在VS2012會顯示 未定義
推 johnjohnlin: 看你的 cuda 是幾版
→ usshan: 目前查網路上的資料在屬性內的值改成compute_20,sm_20
→ usshan: 已經可以使用了   感謝J大提供的方法  謝謝
推 johnjohnlin: default 是 sm1.0 的 CUDA 已經很舊了...
推 nigue: 在vs2012使用atomic函式要在專案的屬性的CUDA C/C++中的
→ nigue: code generation的裡面改成compute_20,sm_20
推 nigue: atomic如此,使用 Thrust library亦為如此