max - OpenCL 中 float 的原子最大值

Question

我需要一个用于 OpenCL 中 float 的原子最大函数。这是我目前使用 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 中所述.

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

但是，由于使用这些方法导致的速度下降非常大。小心使用它们。