归约求和问题，求不全！！！！！！！！

加速计算 CUDA
1

用到一个归约求和，参考了书上180页的reduce2的算法，但是结果与单个累加的结果相差一点点，感觉是漏掉了哪些数据没有求和进去一样！
相关代码如下：
#define DataNum 1024 //数据量
device int d_data_cost[DataNum];
……
global void ****()
{
int tid_in_grid = blockDim.x * blockIdx.x + threadIdx.x;
int tid_in_block= threadIdx.x;
int Tcost;
……
__syncthreads();
for(int s=DataNum/2;s>0;s>>=1)
{
if(tid_in_grid<s)
{
d_data_cost[tid_in_grid]+=d_data_cost[tid_in_grid+s];
}
__syncthreads();
}
//__syncthreads();
if(tid_in_block==0)Tcost=d_data_cost[0];
__syncthreads();

结果：
单个累加值：25821
归约求和值：18392

求高手指点！！！！！万分感谢！

2

！楼主注意这里！你不能尝试让一个grid通过__syncthreads()来同步。__syncthreads只对一个block里面的线程有效。
if(tid_in_grid<s) //如果你只有一个block的话，可以考虑使用你的上面的那个tid_in_block.
{
d_data_cost[tid_in_grid]+=d_data_cost[tid_in_grid+s];
}
__syncthreads(); //所以这里并不是你想象的所有线程都同步，而只是一个block中的都同步。

所以，你的上述代码，只在仅有一个block的情况下（也就是你的block形状是{N,1,1}, grid形状是{1,1,1})，才能得到正确的结果。

3

搬个板凳来围观