max - OpenCL 中浮点数的原子最大值

Question

我需要一个用于 OpenCL 中浮点数的原子最大函数。这是我当前使用 atomic_xchg 的幼稚代码

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

此代码在使用 Intel CPU 时给出了正确的结果，但对于 Nvidia GPU 却没有。此代码是否正确，或者有人可以帮助我吗？

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中所述。

它所做的是创建一个浮点数和整数的联合。在浮点数上执行数学运算，但在执行原子 xchg 时比较整数。只要整数匹配，操作就完成。

然而，由于使用这些方法而导致的速度下降非常高。小心使用它们。