问题遇到的现象和发生背景
问题现象：在使用CUDA并行计算时，有一个步骤需要叠加数组内元素的和，直接相加必然会出现线程抢夺资源导致的计算错误，所以要使用原子操作功能。但是我发现当每个线程计算次数超过一定数时，计算结果就会出现错误，结果变成一个特别大的数，仿佛是超过表示限制了。
发生背景：使用cuda计算点云平均点距，需要对图像上像素对应的每个点的距离进行总和再除以点数，这就需要使用原子加和操作atomicAdd()。操作的数据类型是双精度浮点数double。但是我使用的是计算能力为3.5的GT720显卡（mutiProcessorCount = 1，maxThreadsPerBlock = 1024），不支持double类型原子操作，需要对原子操作进行重载，按照官网的重载代码完成了对double类型的计算功能实现。
但是实际实验发现，我在核函数当中想要用for循环计算所有点的点距和，一共有1920*1200个点，我调用了共2048个线程，需要每个线程执行for循环计算1125次，但是这样计算结果就会出错，变成一个非常大的数字(-6.27744e+066)，但是当我只计算一两万个点的时候，计算结果就没有错误。
为了方便查看测试结果是否正确，我假设每个点的点距都是1，这样有多少点正确的数就是多少。
问题相关代码
int iThreadNum = 256;
int iBlockNum = 8;
__global__ add(const int iImageW, const int iImageH,int const iThreadNum ,const int iBlockNum ,,double* result)
{
const int iImageSize = iImageW * iImageH;
const int tid = threadIdx.x;
const int bid = blockIdx.x;
for (int t = bid *iThreadNum + tid ; t < iImageSize ; t += iBlockNum * iThreadNum)
{
atomicAdd(&(result[0]), 1.0);
}
}
把iImageSize 改成 2048 后计算结果正常，
经过测试iImageSize一旦大于368640就会出错。
int iThreadNum = 256;
int iBlockNum = 8;
__global__ add(const int iImageW, const int iImageH,int const iThreadNum ,const int iBlockNum ,,double* result)
{
const int iImageSize = 2048;
const int tid = threadIdx.x;
const int bid = blockIdx.x;
for (int t = bid *iThreadNum + tid ; t < iImageSize ; t += iBlockNum * iThreadNum)
{
atomicAdd(&(result[0]), 1.0);
}
}
运行结果及报错内容
第一个核函数结果：
result[0] = -6.27744e+066
第二个核函数结果
result[0] = 2048
我想要达到的结果
想要搞清楚问题出在哪，在英伟达显卡的官网cuda指导书上没有发现类似问题的描述，以及如何才能正确计算大量数的双精度浮点数的和。请您能够指点一下万分感谢。
