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 和 int 的联合。对 float 执行数学运算，但在执行原子 xchg 时比较整数。只要整数匹配，操作就完成。

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