double原子操作

加速计算 CUDA
1

由于精度原因，把程序从单精度改为了双精度的。原子加操作就使用手册里的double原子加操作。在我屏蔽掉原子操作这一过程，程序出来的结果误差还能接受，但是有了原子操作这一步之后，程序迭代次数多了之后，结果误差就惨不忍睹了。各位大大，有什么办法解决不?

2

根据楼主的描述，楼主有2种实现：一个过程是不是用原子操作，结果还不错的；一个是使用了原子操作，结果惨不忍睹的。

那么我建议楼主使用前者。

3

此外，如果不回复你这个问题。而是就你的描述来说，敢问楼主是如何实现double原子加法的吗？能否给出实现过程？谢谢！

4

如您所说，您说不使用原子操作，精度可以接受，而使用原子操作，精度差很多。

您是如何评估的呢？
如果是同一个功能，那么推荐您果断使用精度较好的实现。
如果并不是同一个功能，那么似乎没有可比性。

以及，如果必须使用原子操作完成该任务，您还可以从您的数据特点等角度出发，分析看看是什么原因引起的。

大致如上，供LZ参考。

祝您编码愉快~

5

我这个程序是有两部分的，第二部分是原子操作的。我把原子操作的先给屏蔽了。然后和串行的代码比较结果的。我目前只是知道是原子操作这部分的误差。不知道是不是类似一个很大的数加一个比较小的数造成的。

6

1：根据您最新提供的信息，您否决了4#中我说的第一种情况。那么请您考虑，能否从算法上避免原子操作。

2：如果不可以，请参看并回答2#，3#的说法和问题，给出您实现的“double类型原子加”的具体实现。

3：是否是大数吃掉小数，这个和您的数据特点有关系，需要您根据实际情况分析，我们无法凭空猜测得到。

祝您编码愉快~

7

可能我没说清楚你误解了我的意思的。程序有：
·1、A过程
2、B过程，用原子操作去实现的这个过程。
3、其他过程
整个程序先执行1，然后2，然后3，只是我在测试的时候把2过程给屏蔽了。在这种情况下误差可以接受的。但是1,2,3,都执行的话，误差就随着迭代的次数增加而增大的。

__device__ double atomicAdd_double(double *address,double val)
{
   unsigned long long int *address_as_ull=(unsigned long long int *)address;
   unsigned long long int old=*address_as_ull, assumed;
   do{
   assumed = old;
   old=atomicCAS(address_as_ull,assumed,                                  __double_as_longlong(val+__longlong_as_double(assumed)));
   }while(assumed!=old);
   return __longlong_as_double(old);
}

这个我是直接利用了手册里的函数的。

8

这样的确是可以的。

误差可能存在于，大数吃掉了小数。请参考ICE在12月份的帖子，如何用kahan的加法公式来实现补偿。
我综合了一下ICE的帖子，以及你的需求。给出如下代码仅供参考。
device int lock;
device double sum;
device double compensation;
global void buddy_init()
{
lock = 0;
sum = 0.0;
compensation = 0.0;
}
device void buddy_add(double value)
{

while(atomicExch(&lock, 1) == 1) {}

double s0 = sum;
double c0 = compensation;

double v = value + c0;
double s1 = s0 + v;
double c1 = v + (s0 - s1);

sum = s1;
compensation = c1;

__threadfence();

atomicExch(&lock, 0);

}

这个是手写的，没经过验证。里面综合ice的锁的帖子，ice的kahan补偿加法公式帖子。供您参考。

EDIT: 请注意楼下，如果需要优化建议自己再动脑。

9

此外，这个原型实现是没有丝毫优化的。仅仅是示范哦。

10

非常感谢你细心的回答。只是在问一个问题，ice的kahan补偿的帖子，翻了12月份的帖子，都没看到ice的发帖。能否给个链接。谢谢！

11

这个。我也得翻翻。应该是某个人问，然后ICE回复帖子，不是单独的主题。我建议楼主仔细翻翻？论坛尚没有搜索功能，抱歉了。

12

谢谢横扫千军和ice热情的回答。Kahan summation algorithm在我程序中用到了，只是不知道为什么不起作用的。在理解你的回复内容之后，我知道我问题所在之处了。谢谢了~！

13

不客气不客气~帮助您是职责所在。

如果您已经明白了前面构造一个包含kahan算法的原子操作的实现方法，不妨先试试，看看精度如何？以及看看效率如何？回来反馈一下。

横扫斑竹在前已经暗示可能有效率更高的方法这个我会随便乱说么？

祝您编码愉快~

14

问题还是没解决的。我是对一个对称矩阵矢量乘的，只存储了上三角的数据，下面的操作是把下三角的值加入到上三角的乘积中。

__device__ void atomicAdd_double(double *address,double val,int j,double *d_eps,int *d_lock)
{

	bool finished = false;
	double sum;
	double eps;
	while(!finished){  
   if( atomicExch(&d_lock[j], 1) == 0){  
   sum=address[j];
   eps=d_eps[j];
   
   double y=val-eps;
		 double t=sum+y;
		 eps=(t-sum)-y;
   
   d_eps[j]=eps;
   address[j]=t;
   
   finished = true;  
   atomicExch(&d_lock[j], 0);  
   }  
   }  
}

数组address存储的是矩阵上三角矩阵矢量乘的结果，val是需要加入到address中的值。由于我是按照矩阵的行来并行的。所以可能会同时对同一个地址空间读写，所以就用了原子操作的。
我对每一行都设置了一个lock和kahan算法的补偿值。不知道为什么还是结果有问题的。

15

请重新阅读你的代码和我的代码。看看有啥区别。我们的代码不等价。

然后仔细考虑下。

此外，如果你修正了代码仍然不对。你还可以考虑其他方式累加，请ICE补充。

EDIT: 至少你不能去掉我的__threadfence…