原子新增功能

    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 實現更隨機或完全按順序放置值。