opencl
Operazioni atomiche
Ricerca…
Sintassi
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 (volatile __local unsigned int * p, unsigned int val)
Parametri
| p | val |
|---|---|
| puntatore alla cella | aggiunto alla cella |
Osservazioni
Le prestazioni dipendono dal numero di operazioni atomiche e dallo spazio di memoria. Il lavoro seriale rallenta quasi sempre l'esecuzione del kernel a causa del fatto che gpu è un array SIMD e ogni unità in un array attende altre unità se non eseguono lo stesso tipo di lavoro.
Funzione di aggiunta atomica
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];
}
Uscita dei primi elementi:
qualche volta
192 193 194 195 196 197 198
qualche volta
0 1 2 3 4 5 6
qualche volta
128 129 130 131 132 133 134
per un'impostazione con intervallo locale = 256.
Qualunque thread faccia prima visita a fakeMalloc, inserisce il proprio id del thread locale nella prima cella dei risultati. La SIMD interna e la struttura del fronte d'onda del gpu di esempio consentono a 64 thread vicini di mettere i loro valori nelle celle dei risultati adiacenti. Altri dispositivi possono mettere i valori in ordine casuale o totalmente a seconda dell'implementazione opencl di tali dispositivi.