Atomic max для поплавков в OpenCL

Question

Мне нужна функция атомного максимума для float в OpenCL. Это мой текущий наивным код с помощью atomic_xchg

float value = data[index]; 
if (value > *max_value) 
{ 
    atomic_xchg(max_value, value); 
} 

Этот код дает правильный результат при использовании процессора Intel, но не для GPU Nvidia. Правильно ли этот код, или кто-нибудь может мне помочь?

Answer 1

Вы можете сделать это следующим образом:

//Function to perform the atomic max 
inline void AtomicMax(volatile __global float *source, const float operand) { 
    union { 
     unsigned int intVal; 
     float floatVal; 
    } newVal; 
    union { 
     unsigned int intVal; 
     float floatVal; 
    } prevVal; 
    do { 
     prevVal.floatVal = *source; 
     newVal.floatVal = max(prevVal.floatVal,operand); 
    } while (atomic_cmpxchg((volatile __global unsigned int *)source, prevVal.intVal, newVal.intVal) != prevVal.intVal); 
} 

__kernel mykern(__global float *data, __global float *max_value){ 
    unsigned int index = get_global_id(0); 

    float value = data[index]; 
    AtomicMax(max_value, value); 
} 

Как указано в LINK.

Что он делает, это создать объединение float и int. Выполните математику на float, но сравните целые числа при выполнении атомарного xchg. Пока целые числа совпадают, операция завершается.

Однако снижение скорости из-за использования этих методов очень велико. Используйте их осторожно.