OpenCL - метод для выполнения СОКРАЩЕНИЯ

Question

Из following post, я пытаюсь осуществить сокращение суммы массива с этим ядром кода:

#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable 

__kernel void sumGPU (__global const long *input, 
       __global long *finalSum 
       ) 
{ 
    uint local_id = get_local_id(0); 
    uint group_size = get_local_size(0); 

    // Temporary local value 
    local long tempInput; 

    tempInput = input[local_id]; 

    // Variable for final sum 
    local long totalSumIntegerPart[1]; 

    // Initialize sums 
    if (local_id==0) 
    totalSumIntegerPart[0] = 0; 

    // Compute atom_add into each workGroup 
    barrier(CLK_LOCAL_MEM_FENCE); 

    atom_add(&totalSumIntegerPart[0], tempInput); 

    barrier(CLK_LOCAL_MEM_FENCE); 

    // Perform sum of each workGroup sum 
    if (local_id==(get_local_size(0)-1)) 
    atom_add(finalSum, totalSumIntegerPart[0]); 

}     

Но значение finalSum не ожидаемое значение (я имею сначала установите input массив:

for (i=0; i<nWorkItems; i++) 
    input[i] = i+1; 

Итак, я ожидаю с nWorkItems = 1024: finalSum = nWorkItems*(nWorkItems+1)/2=524800

И на самом деле, я получаю finalSum = 16384.

Получаю этот результат, принимая sizeWorkGroup = 16 и nWorkItems = 1024.

Как ни странно, с sizeWorkGroup = 32 и nWorkItems = 1024, я получаю еще одно значение: finalSum = 32768

Я не понимаю последнюю команду (которая, как предполагается, чтобы вычислить сумму каждой частичной суммы, то есть для каждой рабочей группы):

// Perform sum of each workGroup sum 
    if (local_id==(get_local_size(0)-1)) 
    atom_add(finalSum, totalSumIntegerPart[0]); 

Действительно, я бы подумал, что инструкция atom_add(finalSum, totalSumIntegerPart[0]); не зависит от local_idif condition.

Важнейшей является эта инструкция, которая должна быть выполнена «number of workGroups» раз (предполагается, что finalSum является общим значением между всеми рабочими группами, не так ли?).

Так я думал, что я мог бы заменить:

// Perform sum of each workGroup sum 
    if (local_id==(get_local_size(0)-1)) 
    atom_add(finalSum, totalSumIntegerPart[0]); 

по

// Perform sum of each workGroup sum 
     if (local_id==0) 
     atom_add(finalSum, totalSumIntegerPart[0]); 

Любой может помочь найти правильное значение с моими параметрами (sizeWorkGroup = 16 и nWorkItems = 1024), т.е. finalSum, равных 524800?

или exlain мне, почему эта окончательная сумма не очень хорошо выполнена?

UPDATE:

Вот код ядра на following link (она немного отличается от моего, потому что atom_add здесь только шаг 1 для каждого WorkItem):

kernel void AtomicSum(global int* sum) 

{ 
local int tmpSum[1]; 
if(get_local_id(0)==0){ 
tmpSum[0]=0;} 

barrier(CLK_LOCAL_MEM_FENCE);   
atomic_add(&tmpSum[0],1);   
barrier(CLK_LOCAL_MEM_FENCE);  

if(get_local_id(0)==(get_local_size(0)-1)){ 
    atomic_add(sum,tmpSum[0]); 
} 

} 

Является ли это правильный код ядра , Я имею в виду, что дает хорошие результаты?

Может быть, решение могло бы положить на начало моего кода ядра:

unsigned int tid = get_local_id(0); 
unsigned int gid = get_global_id(0); 
unsigned int localSize = get_local_size(0); 
// load one tile into local memory 
int idx = i * localSize + tid; 
localInput[tid] = input[idx]; 

Я собираюсь проверить его и держать вас в курсе.

Благодаря

Answer 1

Эта линия является неправильным:

tempInput = input[local_id]; 

Должно быть:

tempInput = input[get_global_id(0)]; 

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

16*16*64 = 16384 
32*32*32 = 32768 

Кроме того, ваш код может быть немного упрощен:

uint local_id = get_local_id(0); 

    // Variable for final sum 
    local long totalSumIntegerPart; 

    // Initialize sums 
    if (local_id==0) 
    totalSumIntegerPart = 0; 

    // Compute atom_add into each workGroup 
    barrier(CLK_LOCAL_MEM_FENCE);  
    atom_add(&totalSumIntegerPart, input[get_global_id(0)]);  
    barrier(CLK_LOCAL_MEM_FENCE); 

    // Perform sum of each workGroup sum 
    if (local_id==0) 
    atom_add(finalSum, totalSumIntegerPart); 

И я бы не злоупотреблять, как вы делаете из Атомикса, потому что они не являются наиболее эффективным способом делать сокращение. Вероятно, вы сможете получить 10-кратную скорость с помощью надлежащих методов сокращения. Тем не менее, это нормально, как PoC или для изучения локальной памяти и CL.