opencl
Atomoperationer

Syntax

• int atomic_add (flyktig __global int * p, int val)

• unsigned int atomic_add (flyktig __global unsigned int * p, osignerad int val)

• int atomic_add (flyktig __local int * p, int val)

• unsigned int atomic_add (flyktig __local unsigned int * p, unsigned int val)

parametrar

Anmärkningar

Atomic Add-funktion

    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];
    }

Output från första element:

ibland

192 193 194 195 196 197 198

ibland

0 1 2 3 4 5 6

ibland

128 129 130 131 132 133 134

för en inställning med lokalt intervall = 256.

Oavsett vilken tråd som besöker fakeMalloc först, sätter den sin egen lokala tråd-id i första resultatcellen. Intern SIMD- och vågfrontstruktur i exemplet gpu låter angränsande 64 trådar sätta sina värden till angränsande resultatceller. Andra enheter kan sätta värden mer slumpmässigt eller helt i ordning beroende på den öppna implementeringen av enheterna.