OpenCL - 双重原子操作 - 工作到极限
OpenCL - Atomic operation with double - works until limit
在这个 link 之后,我尝试实现一个原子函数来计算 double 数组的总和,所以我实现了自己的 atom_add
函数(对于 double)。
这是使用的内核代码:
#pragma OPENCL EXTENSION cl_khr_fp64: enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
void atom_add_double(__global double *val, double delta)
{
union {
double f;
ulong i;
} old, new;
do
{
old.f = *val;
new.f = old.f + delta;
}
while (atom_cmpxchg((volatile __global ulong *)val, old.i, new.i) != old.i);
}
__kernel void sumGPU ( __global const double *input,
__global double *finalSum
)
{
// Index of current workItem
uint gid = get_global_id(0);
// Init sum
*finalSum = 0.0;
// Compute final sum
atom_add_double(finalSum, input[gid]);
}
我的问题是内核代码会产生良好的结果,直到我达到 input
数组大小的大约 100000 个元素。
超过这个限制,计算不再有效(我可以很容易地检查结果,因为在我的测试用例中,我通过循环 for(i=0;i<sizeArray;i++) input[i]=i+1;
填充输入数组,所以总和等于sizeArray*(sizeArray+1)/2
).
我可以定义像atom_add_double
这样的函数并将其放入内核代码吗?
*finalSum = 0.0;
是所有 in-flight 线程的竞争条件。它使我的计算机的结果为零。删除它,从主机端初始化它。如果你的 gpu 非常好,in-flight 线程的数量可能高达 50000 甚至更多,并且每个线程在任何开始原子函数之前都达到 finalSum = 0.0 但是当你超过该限制时,第 50001 个(只是一个微不足道的数字) 将 re-initializes 线程化为零。
那么,所有元素的和不等于size*(size+1)/2 因为它是从0开始的(第0个元素为0)所以实际上是
(size-1)*(size)/2
并且在我从内核中删除 finalSum =0.0 时为我的计算机提供了正确的结果。
@huseyin 的回答是正确的,可以解决问题。
不过还是忍不住要说"Don't use atomics to reduce."
更糟糕的是锁定在 while 循环中并直接访问全局数据的原子。我们可能正在谈论至少 10 倍的性能损失。
如果可以,请使用proper automatic reduction (CL 2.0+)。
__kernel void sumGPU(__global const double *input, __global double *finalSum)
{
// Index of current workItem
uint gid = get_global_id(0);
// Sum locally without atomics
double sum = work_group_scan_inclusive_add(input[gid]);
// Compute final sum using atomics
// but it is even better if just store them in an array and do final sum in CPU
// Only add the last one, since it contains the total sum
if (get_local_id(0) == get_local_size(0) - 1) {
atom_add_double(finalSum, sum);
}
}
在这个 link 之后,我尝试实现一个原子函数来计算 double 数组的总和,所以我实现了自己的 atom_add
函数(对于 double)。
这是使用的内核代码:
#pragma OPENCL EXTENSION cl_khr_fp64: enable
#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
void atom_add_double(__global double *val, double delta)
{
union {
double f;
ulong i;
} old, new;
do
{
old.f = *val;
new.f = old.f + delta;
}
while (atom_cmpxchg((volatile __global ulong *)val, old.i, new.i) != old.i);
}
__kernel void sumGPU ( __global const double *input,
__global double *finalSum
)
{
// Index of current workItem
uint gid = get_global_id(0);
// Init sum
*finalSum = 0.0;
// Compute final sum
atom_add_double(finalSum, input[gid]);
}
我的问题是内核代码会产生良好的结果,直到我达到 input
数组大小的大约 100000 个元素。
超过这个限制,计算不再有效(我可以很容易地检查结果,因为在我的测试用例中,我通过循环 for(i=0;i<sizeArray;i++) input[i]=i+1;
填充输入数组,所以总和等于sizeArray*(sizeArray+1)/2
).
我可以定义像atom_add_double
这样的函数并将其放入内核代码吗?
*finalSum = 0.0;
是所有 in-flight 线程的竞争条件。它使我的计算机的结果为零。删除它,从主机端初始化它。如果你的 gpu 非常好,in-flight 线程的数量可能高达 50000 甚至更多,并且每个线程在任何开始原子函数之前都达到 finalSum = 0.0 但是当你超过该限制时,第 50001 个(只是一个微不足道的数字) 将 re-initializes 线程化为零。
那么,所有元素的和不等于size*(size+1)/2 因为它是从0开始的(第0个元素为0)所以实际上是
(size-1)*(size)/2
并且在我从内核中删除 finalSum =0.0 时为我的计算机提供了正确的结果。
@huseyin 的回答是正确的,可以解决问题。
不过还是忍不住要说"Don't use atomics to reduce."
更糟糕的是锁定在 while 循环中并直接访问全局数据的原子。我们可能正在谈论至少 10 倍的性能损失。
如果可以,请使用proper automatic reduction (CL 2.0+)。
__kernel void sumGPU(__global const double *input, __global double *finalSum)
{
// Index of current workItem
uint gid = get_global_id(0);
// Sum locally without atomics
double sum = work_group_scan_inclusive_add(input[gid]);
// Compute final sum using atomics
// but it is even better if just store them in an array and do final sum in CPU
// Only add the last one, since it contains the total sum
if (get_local_id(0) == get_local_size(0) - 1) {
atom_add_double(finalSum, sum);
}
}