OpenCL - 制限付きまでのdouble型のアトミック演算

Question

この後、linkの配列の合計を計算するアトミック関数を実装しようとしました。したがって、自分自身のatom_add関数（double型）を実装しました。ここで

は使用カーネルのコードです：

#pragma OPENCL EXTENSION cl_khr_fp64: enable 
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable 

void atom_add_double(__global double *val, double delta) 
{ 
    union { 
    double f; 
    ulong i; 
    } old, new; 

    do 
    { 
    old.f = *val; 
    new.f = old.f + delta; 
    } 
    while (atom_cmpxchg((volatile __global ulong *)val, old.i, new.i) != old.i); 

} 

__kernel void sumGPU (__global const double *input, 
       __global double *finalSum 
       ) 
{ 
    // Index of current workItem 
    uint gid = get_global_id(0); 

    // Init sum 
    *finalSum = 0.0; 

    // Compute final sum 
    atom_add_double(finalSum, input[gid]); 

}     

私の問題は、私はinput配列のサイズのためにおよそ100000要素に到達するまで、カーネルコードが良い結果を生成していることです。

この制限を超える、計算はもはや有効ではない私のテストケースでは、私はループfor(i=0;i<sizeArray;i++) input[i]=i+1;により入力配列を埋めるため、合計はsizeArray*(sizeArray+1)/2に等しいので、（私は簡単に結果を確認することができます。

誰もがバグのこの種をすでに持っている？

は、私が定義し、カーネルコードにatom_add_doubleのような機能を入れることはできますか？

をすべてのヘルプは大歓迎です、ありがとう

Answer 1

*finalSum = 0.0; 

は、すべての実行中スレッドの競合状態です。私のコンピュータでは結果がゼロになっています。削除し、ホスト側から初期化します。あなたのgpuが非常に良い場合は、飛んでいるスレッドの数はおそらく50000以上になる可能性があります。さらに、atomicSize = 0.0となる前にfinalSum = 0.0を打ちますが、その制限を超えると50001st（ちょっとしたスレッド）ゼロに再初期化します。

それはゼロから開始されるので、次に、すべての要素の合計サイズ*（サイズ+ 1）/ 2に等しくない（ゼロ番目の要素がゼロである）ので、実際に

(size-1)*(size)/2 

、右与えていますカーネルからfinalSum = 0.0を削除したときの私のコンピュータの結果。

Answer 2

@huseyin答えは問題を解決するために正しいです。

しかし、私はと言うことに抵抗することはできません "削減するためにアトミックを使用しないでください。

また、whileループをロックしてグローバルデータに直接アクセスするアトミック性がさらに悪化します。おそらく、少なくとも10倍のパフォーマンスヒットを話しているでしょう。

可能であれば、proper automatic reduction (CL 2.0+)を使用してください。私はすでにあなたに言ったよう

__kernel void sumGPU(__global const double *input, __global double *finalSum) 
{ 
    // Index of current workItem 
    uint gid = get_global_id(0); 

    // Sum locally without atomics 
    double sum = work_group_scan_inclusive_add(input[gid]); 

    // Compute final sum using atomics 
    // but it is even better if just store them in an array and do final sum in CPU 
    // Only add the last one, since it contains the total sum 
    if (get_local_id(0) == get_local_size(0) - 1) { 
    atom_add_double(finalSum, sum); 
    } 
}