opencl
アトミックオペレーション

構文

• int atomic_add（volatile __global int * p、int val）

• unsigned int atomic_add（volatile __global unsigned int * p、unsigned int val）

• int atomic_add（volatile __local int * p、int val）

• unsigned int atomic_add（volatile __local unsigned int * p、unsigned int val）

パラメーター

備考

原子追加機能

    int fakeMalloc(__local int * addrCounter,int size)
    {
        // lock addrCounter
        // adds size to addrCounter's pointed cell value
        // unlock
        // return old value of addrCounter's pointed cell

        // serial between all threads visiting -> slow
        return atomic_add(addrCounter,size);
    }

    __kernel void vecAdd(__global float* results )
    {
       int id = get_global_id(0);
       int lid=get_local_id(0);
       __local float stack[1024];
       __local int ctr;
       if(lid==0)
          ctr=0;
       barrier(CLK_LOCAL_MEM_FENCE);
       stack[lid]=0.0f;                    // parallel operation
       barrier(CLK_LOCAL_MEM_FENCE);
       int fakePointer=fakeMalloc(&ctr,1); // serial operation
       barrier(CLK_LOCAL_MEM_FENCE);
       stack[fakePointer]=lid;             // parallel operation
       barrier(CLK_GLOBAL_MEM_FENCE);
       results[id]=stack[lid];
    }

最初の要素の出力：

時々

192 193 194 195 196 197 198

時々

0 1 2 3 4 5 6

時々

128 129 130 131 132 133 134

ローカル範囲= 256の設定の場合。

最初にfakeMallocにアクセスしたスレッドは、最初の結果セルに自身のローカルスレッドIDを配置します。例のgpuの内部SIMDおよび波面構造は、隣接する64スレッドにそれらの値を隣接する結果セルに入れることを可能にする。他のデバイスは、そのデバイスのopencl実装に応じて、よりランダムまたは完全に値を並べることができます。