OpenCL снижение поплавка сумма

Question

Я хотел бы применить уменьшить на эту часть моего кода ядра (1 размерные данные):

__local float sum = 0; 
int i; 
for(i = 0; i < length; i++) 
    sum += //some operation depending on i here; 

Вместо того, чтобы только 1 поток, который выполняет эту операцию, я хотел бы иметь n потоков (с n = длина), а в конце имеет 1 поток, чтобы сделать общую сумму.

В псевдокоде, я хотел бы иметь возможность написать что-то вроде этого:

int i = get_global_id(0); 
__local float sum = 0; 
sum += //some operation depending on i here; 
barrier(CLK_LOCAL_MEM_FENCE); 
if(i == 0) 
    res = sum; 

Есть ли способ?

У меня есть состояние гонки на сумму.

Answer 1

Чтобы начать работу, вы можете сделать что-то вроде приведенного ниже примера (see Scarpino). Здесь мы также используем векторную обработку с использованием типа данных OpenCL float4.

Имейте в виду, что ядро ​​ниже возвращает несколько частичных сумм: по одной для каждой локальной рабочей группы, обратно на хост. Это означает, что вам нужно будет выполнить окончательную сумму, добавив все частичные суммы обратно на хост. Это связано с тем, что (по крайней мере с OpenCL 1.2) нет барьерной функции, которая синхронизирует рабочие элементы в разных рабочих группах.

Если суммирование частичных сумм на хосте нежелательно, вы можете обойти это, запустив несколько ядер. Это вводит некоторые служебные данные ядра, но в некоторых приложениях дополнительный штраф является приемлемым или несущественным. Чтобы сделать это в следующем примере, вам нужно будет изменить свой хост-код, чтобы повторно вызвать ядро, а затем включить логику, чтобы остановить выполнение ядра после того, как количество выходных векторов опустится ниже локального размера (подробности остались вам или проверьте Scarpino reference) ,

EDIT: Добавлен дополнительный аргумент ядра для вывода. Добавлен точечный продукт для суммирования по векторам float 4.

__kernel void reduction_vector(__global float4* data,__local float4* partial_sums, __global float* output) 
{ 
    int lid = get_local_id(0); 
    int group_size = get_local_size(0); 
    partial_sums[lid] = data[get_global_id(0)]; 
    barrier(CLK_LOCAL_MEM_FENCE); 

    for(int i = group_size/2; i>0; i >>= 1) { 
     if(lid < i) { 
      partial_sums[lid] += partial_sums[lid + i]; 
     } 
     barrier(CLK_LOCAL_MEM_FENCE); 
    } 

    if(lid == 0) { 
     output[get_group_id(0)] = dot(partial_sums[0], (float4)(1.0f)); 
    } 
} 

Answer 2

Простым и быстрым способом сокращения данных является многократное складывание верхней половины данных в нижнюю половину.

Например, используйте следующий смехотворно простой код CL:

__kernel void foldKernel(__global float *arVal, int offset) { 
    int gid = get_global_id(0); 
    arVal[gid] = arVal[gid]+arVal[gid+offset]; 
} 

С помощью следующего кода Java/JOCL хоста (или портировать его на C++ и т.д.):

int t = totalDataSize; 
    while (t > 1) { 
     int m = t/2; 
     int n = (t + 1)/2; 
     clSetKernelArg(kernelFold, 0, Sizeof.cl_mem, Pointer.to(arVal)); 
     clSetKernelArg(kernelFold, 1, Sizeof.cl_int, Pointer.to(new int[]{n})); 
     cl_event evFold = new cl_event(); 
     clEnqueueNDRangeKernel(commandQueue, kernelFold, 1, null, new long[]{m}, null, 0, null, evFold); 
     clWaitForEvents(1, new cl_event[]{evFold}); 
     t = n; 
    } 

Хост кода циклов log2 (n) раз, поэтому он быстро заканчивается даже с огромными массивами. Скрипка с «m» и «n» предназначена для обработки массивов без питания двух.

• Легко для OpenCL хорошо распараллелить для любого графического процессора платформы (т.е. быстро).
• Низкая память, потому что он работает на месте
• эффективно работает с неэнергетических-оф-двух данных размеров
• Flexible, например, Вы можете изменить ядро, чтобы сделать «мин» вместо «+»

Answer 3

Вы можете использовать новую функцию work_group_reduce_add() для уменьшения суммы внутри одной рабочей группы, если у вас есть поддержка OpenCL C 2.0 features