opencl
Operaciones atómicas
Buscar..
Sintaxis
int atomic_add (volatile __global int * p, int val)
unsigned int atomic_add (__global volatile unsigned int * p, int int sin signo)
int atomic_add (volatile __local int * p, int val)
unsigned int atomic_add (volatile __local unsigned int * p, unsigned int val)
Parámetros
| pag | val |
|---|---|
| puntero a celda | añadido a la celda |
Observaciones
El rendimiento depende del número de operaciones atómicas y del espacio de memoria. El trabajo en serie casi siempre ralentiza la ejecución del kernel debido a que gpu es una matriz SIMD y cada unidad en una matriz espera otras unidades si no realizan el mismo tipo de trabajo.
Función de adición atómica
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];
}
Salida de primeros elementos:
algunas veces
192 193 194 195 196 197 198
algunas veces
0 1 2 3 4 5 6
algunas veces
128 129 130 131 132 133 134
para una configuración con rango local = 256.
Cualquiera que sea el hilo que visite fakeMalloc primero, coloca su propia ID de hilo local en la primera celda de resultados. El SIMD interno y la estructura de frente de onda del ejemplo gpu permiten que los 64 hilos vecinos coloquen sus valores en las celdas de resultados vecinas. Otros dispositivos pueden poner los valores de forma más aleatoria o totalmente en orden, dependiendo de la implementación abierta de dichos dispositivos.