OpenCL：同じグローバルメモリアドレスに結果を保存する複数の作業項目

Question

特定の条件に応じて4つの異なる値を格納する必要がある場合に、縮減累積計算を実行しようとしています。私のカーネルは長い配列を入力として受け取り、入力上の各データポイントから得られた「グローバルな合計」である4つの値だけを格納する必要があります。例えば、ある条件を満たすすべてのデータ値の合計と、その条件を満たすデータ点の数を格納する必要があります。カーネルは以下のように定義しています。

__kernel void photometry(__global float* stamp, 
         __constant float* dark, 
         __global float* output) 
{ 
int x = get_global_id(0); 
int s = n * n; 

if(x < s){ 
    float2 curr_px = (float2)((x/n), (x % n)); 
    float2 center = (float2)(centerX, centerY); 
    int dist = (int)fast_distance(center, curr_px); 
    if(dist < aperture){ 
     output[0] += stamp[x]-dark[x]; 
     output[1]++; 
    }else if (dist > sky_inner && dist < sky_outer){ 
     output[2] += stamp[x]-dark[x]; 
     output[3]++; 
    } 
} 
} 

カーネルで宣言されていない値は、すべてマクロによって事前に定義されています。 sは、入力配列stampとdarkの長さであり、n x nの行列は1Dに平坦化されています。

結果は表示されますが、これは自分のCPUバージョンとは異なります。もちろん、私は不思議です：これは私がやろうとしていることを行うための正しい方法ですか？各ピクセルデータが1回だけ追加されていることを確認できますか？累積結果の値を保存する他の方法は考えられません。

Answer 1

あなたのケースでは、アトミック操作が必要です。そうしないと、データ競合によって結果が予測できなくなります。

問題はここにある：あなたが同じ波内のスレッドは、まだ同じステップをたどるかもしれないと想像できる

output[0] += stamp[x]-dark[x]; 
output[1]++; 

、したがって、それは同じ波の内部スレッドのOKかもしれません。グローバルロード命令（ブロードキャスト）を使用して同じ出力[0]値を読み取るためです。次に、計算が終了し、同じメモリアドレス（出力[0]）にデータを格納しようとすると、書き込み操作がシリアル化されます。この時点では、（同じ波の中の作業項目の）正しい結果が得られます。

しかし、プログラムが複数のウェーブを起動する可能性が高いため（ほとんどのアプリケーションでは、この場合が該当します）。異なる波が未知の順序で実行されることがあります。同じメモリアドレスにアクセスすると、その動作はより複雑になります。たとえば、wave0とwave1は、他の計算が行われる前に最初にoutput [0]にアクセスすることができます。つまり、output [0]から同じ値をフェッチします。計算を開始します。計算後、累積結果をoutput [0]に保存します。明らかに、後でメモリを書き込む人だけが実行されたかのように、波の1つの結果が別のものによって上書きされます。あなたが本当のアプリケーションではるかに多くの波を持っていると想像してください。間違った結果を出すのは奇妙ではありません。