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

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

然而,由于使用这些方法而导致的速度降低非常高.仔细使用它们.