求和问题

加速计算 CUDA
1

斑竹您好：
采用如下方式求和，代码如下

 __global__ void Mean_680(float *p,float *d_mean,int Col_num,int Row_num)
 {
   int y_id=blockIdx.y*blockDim.y+threadIdx.y;
   int x_id=blockIdx.x*blockDim.x+threadIdx.x;
//简单低效
   //if (x_id==0&&y_id<75)
   //{
   // for (int i=1;i<56644;i++)
   // {
   //         p[y_id*56644]+=p[i+y_id*56644];
   // } 
   //}
//规约         

   int a=Col_num/2;
   int b=Col_num%2;
   int c;
   while (a!=0)
   {
   if (x_id<a&&y_id<Row_num)
   {
   p[x_id+y_id*Col_num]=p[x_id+y_id*Col_num]+p[x_id+a+b+y_id*Col_num];
   }
   c=a+b;
   a=c/2;
   b=c%2;
   __syncthreads();
   }
}

其中：threaddim（512,1）；blockdim（111,75），col_num=56644,row_num=75,p中为75行56644列的的float数据，代码实现列方向上的和，最终结果为p的第一列为所有列之和，采用代码中注释代码是正确的，现采用规约的方式，结果不正确，辛苦斑竹帮忙看一下，谢谢

2

楼主您好：

所有参与的线程不是一个block, 更不是一个warp。
他们无法在完成 p[x_id+y_idCol_num]+=p[x_id+a+b+y_idCol_num]; 后，集体同步，然后再执行下一次循环的。自然您的结果无法正确。

3

如斑竹所说，我在第27行添加同步，可是结果和没有添加时一样，仍然不对，不知该如何解决这个问题呢？

4

楼主您好，您无法添加全局同步的。

所有您的代码无法让grid同步进行规约的。

5

斑竹您好：
我查了下材料，同意您的观点，那么对于我上文提到的该类型问题，大量数据的求和，该怎么处理呢？

6

您可以考虑多个block, (例如100个), 每个block计算总数据的1/100.
最近将这100个block的和再次求和（第二次kernel启动, 可以用1个block).

您觉得呢？

7

恩，谢谢斑竹，

8

感谢您莅临CUDAZone China,
欢迎下次来访。