OpenCL 将“cl_khr_fp64”双精度值求和为单个数字

Question

来自 and this question 我设法在 OpenCL 1.2 中编译了一个将矢量求和为单个双精度的最小示例。

    /* https://suhorukov.blogspot.com/2011/12/opencl-11-atomic-operations-on-floating.html */
    inline void AtomicAdd(volatile __global double *source, const double operand) {
      union { unsigned int intVal; double floatVal; } prevVal, newVal;
      do {
        prevVal.floatVal = *source;
        newVal.floatVal = prevVal.floatVal + operand;
      } while( atomic_cmpxchg((volatile __global unsigned int *)source, prevVal.intVal, newVal.intVal) != prevVal.intVal );
    }

    void kernel cost_function(__constant double* inputs, __global double* outputs){
      int index = get_global_id(0);

      if(0 == error_index){ outputs[0] = 0.0; }
      barrier(CLK_GLOBAL_MEM_FENCE);

      AtomicAdd(&outputs[0], inputs[index]); /* (1) */
      //AtomicAdd(&outputs[0], 5.0); /* (2) */

    }

实际上这个解决方案是不正确的，因为访问缓冲区时结果始终为 0。这可能有什么问题？

/* (1) */ 处的代码不起作用，/* (2) */ 处的代码也不起作用，它仅用于测试独立于任何输入的逻辑。

此处是否正确使用 barrier(CLK_GLOBAL_MEM_FENCE); 以在对它进行任何计算之前重置输出？

根据 OpenCL 1.2 中的 the specs 原子操作支持单精度浮点数，这是（AtomicAdd）将支持扩展到双精度数的可行方法还是我错过了什么东西？

我测试的设备当然支持cl_khr_fp64˙当然

Answer 1

您的AtomicAdd不正确。即，2 个错误是：

在 union 中，intVal 必须是 64 位整数而不是 32 位整数。
使用 64 位 atom_cmpxchg function and not the 32-bit atomic_cmpxchg 函数。

正确的实现是：

#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
inline void AtomicAdd(volatile __global double *source, const double operand) {
    union { unsigned ulong u64; double f64; } prevVal, newVal;
    do {
        prevVal.f64 = *source;
        newVal.f64 = prevVal.f64 + operand;
    } while(atom_cmpxchg((volatile __global ulong*)source, prevVal.u64, newVal.u64) != prevVal.u64);
}

barrier(CLK_GLOBAL_MEM_FENCE); 在这里使用正确。请注意，barrier 不得位于 if- 或 else- 分支中。

更新：根据 STREAMHPC，您使用的原始实现不能保证产生正确的结果。有一个改进的实现：

void __attribute__((always_inline)) atomic_add_f(volatile global float* addr, const float val) {
    union {
        uint  u32;
        float f32;
    } next, expected, current;
    current.f32 = *addr;
    do {
        next.f32 = (expected.f32=current.f32)+val; // ...*val for atomic_mul_f()
        current.u32 = atomic_cmpxchg((volatile global uint*)addr, expected.u32, next.u32);
    } while(current.u32!=expected.u32);
}

#ifdef cl_khr_int64_base_atomics
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
void __attribute__((always_inline)) atomic_add_d(volatile global double* addr, const double val) {
    union {
        ulong  u64;
        double f64;
    } next, expected, current;
    current.f64 = *addr;
    do {
        next.f64 = (expected.f64=current.f64)+val; // ...*val for atomic_mul_d()
        current.u64 = atom_cmpxchg((volatile global ulong*)addr, expected.u64, next.u64);
    } while(current.u64!=expected.u64);
}
#endif

OpenCL 将“cl_khr_fp64”双精度值求和为单个数字

OpenCL sum `cl_khr_fp64` double values into a single number

opencl

reduction