__local atomic в opencl - PullRequest
Новая
       36

__local atomic в opencl

4 голосов
/ 14 марта 2012

Об атомарном доступе __local переменных:

Я знаю, что глобальные операции выполняются медленно по сравнению с локальными. В этом смысле я хотел бы сделать атомарный доступ к некоторым переменным.

Я знаю, что могу делать атомарные операции в OpenCL: 

// Program A:
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
__kernel void test(global int * num)
{
    atom_inc(&num[0]);
}

Как я могу делиться атомарными данными между рабочими в данной рабочей группе?

Например: я хотел бы сделать что-то вроде этого: 

// Program B: (it doesn't work, just to show how I'd like it to be)
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
__kernel void test(global int * num, const int numOperations)
{
    __local int num;
    if (get_global_id(0) < numOperations) {
        atom_inc(&num);
    }
}

В конце должно возвращаться значение num: numOperations - 1;

Разве это не возможно? Если нет, то как я мог это сделать?

1 Ответ

5 голосов
/ 14 марта 2012

Как правило, у вас есть один поток, который инициализирует общий (локальный) атом, за которым следует некоторый барьер.Т.е. ваше ядро ​​запускается так: 

__local int sharedNum;
if (get_local_id (0) == 0) {
    sharedNum = 0;
}
barrier (CLK_LOCAL_MEM_FENCE);

// Now, you can use sharedNum
while (is_work_left ()) {
    atomic_inc (&sharedNum);
}

В этом нет особой магии - все элементы в рабочей группе могут видеть одни и те же локальные переменные, поэтому вы можете просто обращаться к ним как обычно.

...