opencl
Atomoperationen
Suche…
Syntax
int atomic_add (volatile __global int * p, int val)
unsigned int atomic_add (volatile __global unsigned int * p, vorzeichenloser Int-Wert)
int atomic_add (volatile __local int * p, int val)
unsigned int atomic_add (volatile __local unsigned int * p, unsigned int val)
Parameter
| p | val |
|---|---|
| Zeiger auf Zelle | zur Zelle hinzugefügt |
Bemerkungen
Die Leistung hängt von der atomaren Operationsnummer und dem Speicherplatz ab. Bei seriellen Arbeiten wird die Kernel-Ausführung fast immer verlangsamt, da gpu ein SIMD-Array ist und jede Unit in einem Array auf andere Units wartet, wenn sie nicht dieselbe Art von Arbeit ausführen.
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];
}
Ausgabe der ersten Elemente:
manchmal
192 193 194 195 196 198 198
manchmal
0 1 2 3 4 5 6
manchmal
128 129 130 131 132 133 134
für eine Einstellung mit Ortsbereich = 256.
Was auch immer FakeMalloc zuerst besucht, es wird seine eigene lokale Thread-ID in die erste Ergebniszelle eingefügt. Durch die interne SIMD- und Wellenfrontstruktur des Beispiels gpu können benachbarte 64 Threads ihre Werte in benachbarte Ergebniszellen setzen. Andere Geräte können abhängig von der Opencl-Implementierung dieser Geräte Werte zufällig oder völlig in der richtigen Reihenfolge festlegen.