原子添加功能

    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 实现更随机或完全按顺序放置值。