OpenCL 将“cl_khr_fp64”双精度值求和为单个数字
OpenCL sum `cl_khr_fp64` double values into a single number
来自 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
˙当然
您的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
来自
/* 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
˙当然
您的AtomicAdd
不正确。即,2 个错误是:
- 在
union
中,intVal
必须是 64 位整数而不是 32 位整数。 - 使用 64 位
atom_cmpxchg
function and not the 32-bitatomic_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