opencl
Атомные операции
Поиск…
Синтаксис
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)
параметры
| п | вал |
|---|---|
| указатель на ячейку | добавлен в ячейку |
замечания
Производительность зависит от количества атомных операций и объема памяти. Выполнение серийной работы почти всегда замедляет выполнение ядра из-за того, что gpu является SIMD-массивом, и каждый блок в массиве ждет других устройств, если они не выполняют такой же тип работы.
Функция добавления атома
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];
}
Вывод первых элементов:
иногда
192 193 194 195 196 197 198
иногда
0 1 2 3 4 5 6
иногда
128 129 130 131 132 133 134
для настройки с локальным диапазоном = 256.
Независимо от того, какой поток посещает fakeMalloc, он помещает свой собственный идентификатор локального потока в ячейку первого результата. Внутренняя структура SIMD и волнового фронта в примере gpu позволяет соседним 64 потокам помещать их значения в соседние ячейки результатов. Другие устройства могут ставить значения более случайным образом или полностью в порядке в зависимости от реализации этих устройств.