opencl
Atoomoperaties
Zoeken…
Syntaxis
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 (vluchtig __local unsigned int * p, unsigned int val)
parameters
| p | val |
|---|---|
| aanwijzer naar cel | toegevoegd aan cel |
Opmerkingen
De prestaties zijn afhankelijk van het aantal atomaire bewerkingen en geheugenruimte. Serieel werk vertraagt bijna altijd de uitvoering van de kernel omdat gpu een SIMD-array is en elke eenheid in een array wacht op andere eenheden als ze niet hetzelfde soort werk doen.
Atomic Add Function
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 van eerste elementen:
soms
192 193 194 195 196 196 197 198
soms
0 1 2 3 4 5 6
soms
128 129 130 131 132 133 134
voor een instelling met lokaal bereik = 256.
Welke thread eerst fakeMalloc bezoekt, het plaatst zijn eigen lokale thread-ID in de eerste resultaatcel. Interne SIMD en golffrontstructuur van het voorbeeld gpu laat aangrenzende 64 threads toe om hun waarden aan aangrenzende resultaatcellen te geven. Andere apparaten kunnen waarden meer willekeurig of volledig in volgorde plaatsen, afhankelijk van de opencl-implementatie van die apparaten.