opencl
Operacje atomowe
Szukaj…
Składnia
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)
Parametry
| p | val |
|---|---|
| wskaźnik do komórki | dodano do komórki |
Uwagi
Wydajność zależy od liczby operacji atomowych i przestrzeni pamięci. Wykonywanie pracy szeregowej prawie zawsze spowalnia wykonywanie jądra, ponieważ GPU jest tablicą SIMD, a każda jednostka w tablicy czeka na inne jednostki, jeśli nie wykonują tego samego rodzaju pracy.
Funkcja dodawania atomowego
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];
}
Wyjście pierwszych elementów:
czasami
192 193 194 195 196 197 198
czasami
0 1 2 3 4 5 6
czasami
128 129 130 131 132 133 134
dla ustawienia z zakresem lokalnym = 256.
Niezależnie od tego, który wątek najpierw odwiedza fakeMalloc, umieszcza swój własny identyfikator wątku lokalnego w pierwszej komórce wyniku. Wewnętrzna struktura SIMD i wavefront przykładowego gpu pozwala sąsiednim 64 wątkom na umieszczenie ich wartości w sąsiednich komórkach wynikowych. Inne urządzenia mogą uporządkować wartości bardziej losowo lub całkowicie w kolejności, w zależności od implementacji opencl tych urządzeń.