opencl
Opérations atomiques
Recherche…
Syntaxe
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)
Paramètres
| p | val |
|---|---|
| pointeur vers la cellule | ajouté à la cellule |
Remarques
Les performances dépendent du nombre d'opérations atomiques et de l'espace mémoire. Faire du travail en série ralentit presque toujours l'exécution du noyau car gpu est un tableau SIMD et chaque unité d'un tableau attend d'autres unités si elles ne font pas le même type de travail.
Fonction d'ajout atomique
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];
}
Sortie des premiers éléments:
parfois
192 193 194 195 196 197 198
parfois
0 1 2 3 4 5 6
parfois
128 129 130 131 132 133 134
pour un paramètre avec plage locale = 256.
Quel que soit le thread qui visite d'abord fakeMalloc, il place son propre identifiant de thread local dans la première cellule de résultat. SIMD interne et la structure de front d'onde de l'exemple gpu permet à 64 threads voisins de mettre leurs valeurs aux cellules de résultats voisines. D'autres périphériques peuvent mettre des valeurs plus aléatoires ou totalement en ordre en fonction de l'implémentation opencl de ces périphériques.