OpenCL - Атомная операция с двойным действием до предела

Question

После этого link, я пытаюсь реализовать атомную функцию, которая вычисляет сумму массива double, поэтому я реализовал свою собственную функцию atom_add (для двойного).

Вот код ядра используется:

#pragma OPENCL EXTENSION cl_khr_fp64: enable 
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable 

void atom_add_double(__global double *val, double delta) 
{ 
    union { 
    double f; 
    ulong i; 
    } old, new; 

    do 
    { 
    old.f = *val; 
    new.f = old.f + delta; 
    } 
    while (atom_cmpxchg((volatile __global ulong *)val, old.i, new.i) != old.i); 

} 

__kernel void sumGPU (__global const double *input, 
       __global double *finalSum 
       ) 
{ 
    // Index of current workItem 
    uint gid = get_global_id(0); 

    // Init sum 
    *finalSum = 0.0; 

    // Compute final sum 
    atom_add_double(finalSum, input[gid]); 

}     

Моя проблема заключается в том, что коды ядра генерирует хорошие результаты, пока не достигнет примерно 100 000 элементов для размера input массива.

За этот предел, то вычисление не действует больше (я могу проверить результат легко, потому что в моем тесте я заполнить массив входной петлей for(i=0;i<sizeArray;i++) input[i]=i+1;, поэтому сумма равна sizeArray*(sizeArray+1)/2.

Любой уже получил такую ​​ошибку?

Могу ли я определить и поставить функцию как atom_add_double в коде ядра?

Любая помощь приветствуется, спасибо

Answer 1

*finalSum = 0.0; 

Состояние гонки для всех потоков в полете. Это приводит к нулю результат для моего компьютера. Удалите его, инициализируйте его со стороны хоста. Если ваш gpu очень хорош, количество потоков в полете может достигать 50000, может быть, даже больше, и каждый из них ударяет finalSum = 0.0 до того, как начнется атомная функция, но когда вы пройдете этот предел, 50001st (просто тривиальное число) повторно инициализирует его до нуля.

Тогда сумма всех элементов не равен размеру * (площадь + 1)/2, так как он начиная с нуля (нулевого элемента равна нулю), так что на самом деле

(size-1)*(size)/2 

и дает право результаты для моего компьютера, когда я удаляю finalSum = 0.0 из ядра.

Answer 2

Ответ @huseyin правильный, чтобы исправить проблему.

Однако я не могу удержаться, чтобы сказать «Не используйте атомику для уменьшения».

И еще хуже атомы, которые блокируют цикл while и напрямую получают доступ к глобальным данным. Вероятно, мы, по крайней мере, говорим о 10-кратной производительности.

Если вы можете, используйте proper automatic reduction (CL 2.0+).

__kernel void sumGPU(__global const double *input, __global double *finalSum) 
{ 
    // Index of current workItem 
    uint gid = get_global_id(0); 

    // Sum locally without atomics 
    double sum = work_group_scan_inclusive_add(input[gid]); 

    // Compute final sum using atomics 
    // but it is even better if just store them in an array and do final sum in CPU 
    // Only add the last one, since it contains the total sum 
    if (get_local_id(0) == get_local_size(0) - 1) { 
    atom_add_double(finalSum, sum); 
    } 
}