归约求和问题,求不全!!!!!!!!

用到一个归约求和,参考了书上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

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

!楼主注意这里!你不能尝试让一个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}),才能得到正确的结果。

搬个板凳来围观