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)
パラメーター
p | ヴァル |
---|---|
セルへのポインタ | 細胞に加えられた |
備考
パフォーマンスは、原子操作数とメモリ空間に依存します。シリアル作業を行うと、gpuがSIMD配列であり、配列の各ユニットが同じタイプの作業をしなければ、他のユニットを待つため、ほとんど常にカーネルの実行が遅くなります。
原子追加機能
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実装に応じて、よりランダムまたは完全に値を並べることができます。
Modified text is an extract of the original Stack Overflow Documentation
ライセンスを受けた CC BY-SA 3.0
所属していない Stack Overflow