OpenCL 将"cl_khr_fp64"双精度值求和为一个数字

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 中的规范，原子操作支持单精度浮点数，这是(AtomicAdd)将支持扩展到双精度数的可行方法还是我遗漏了什么？

当然，我正在测试的设备支持cl_khr_fp64当然。

Answer 1

您的AtomicAdd不正确。也就是说，这两个错误是：

1. 在union中，intVal必须是 64 位整数而不是 32 位整数。
2. 使用 64 位atom_cmpxchg函数，而不是 32 位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