Параллельные обновления (x += a) глобальной памяти в OpenCL

Я делаю следующее в ядре OpenCL (упрощенный пример):

__kernel void step(const uint count, __global int *map, __global float *sum)
{
    const uint i = get_global_id(0);
    if(i < count) {
        sum[map[i]] += 12.34;
    }
}

Здесь sum — это некоторая величина, которую я хочу вычислить (ранее установленная в ноль в другом ядре), а map — это отображение целых чисел i в целые числа j, так что несколько i могут отображаться в одно и то же j.

(map может быть в постоянной памяти, а не в глобальной, но кажется, что объем постоянной памяти на моем графическом процессоре невероятно ограничен)

Будет ли это работать? Является ли "+=" реализованным атомарным способом или есть вероятность того, что параллельные операции перезапишут друг друга?


opencl concurrentmodification
person Michael Clerx    schedule 04.03.2015    source источник

Ответы (1)


arrow_upward
3
arrow_downward

Будет ли это работать? Является ли "+=" реализованным атомарным способом или есть вероятность того, что параллельные операции перезапишут друг друга?

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

Что-то типа:

__kernel void step(const uint count, __global int *map, __global double *sum)
{
    const uint i = get_global_id(0);
    if(i < count) {
        atomic_add(&sum[map[i]], 1234);
    }
}
person user703016    schedule 04.03.2015
comment
Спасибо! Есть ли что-нибудь подобное для операций с плавающей запятой? - person Michael Clerx; 04.03.2015
comment
Для OpenCL ›= 2.0 да. До этого нет. Но вы можете реализовать свои собственные atomic_add для чисел с плавающей запятой на основе существующих атомарных операций за счет производительности. - person user703016; 04.03.2015