CUDA规约前缀求和问题

CUDA规约前缀求和问题

前缀求和问题算是个比较常见的问题了，这周项目中遇到了个这样的问题，查找数组a中所有值为M的元素，
用数组b将依次存储这些元素的下标。
这也算是比较经典的前缀求和的问题了。

1.问题的求解思路

并发求解这个数组的思路是这样的，判断所求值与当前数组值是否相等，使用临时变量temp存储，如果
相等设为1，否则为零，然后对所有线程中的temp进行前缀求和，我们通过求解前缀和结果，可以
发现，当前线程对应的前缀和值-1，即为按次序数组b中所求对应线程(数组a的下标)。

2.解决代码

__global__ void scan(int *a,int *b,int equal_value, int N)
{
extern    __share__ int share_sum[];
int tid=thread.x+blockIdx.x*blockDim.x;
int temp1,temp=0;
int i=0;
int t_temp;
int laneid=thread.x&0x1f,warpid=thread.x/warp_size;
if((tid<N)&&(a[tid]==equal_value))
{
temp=1;
}
temp1=temp;//作为标记，用来标记是否写入
/**************这里用来进行前缀求和，将一个线程束中的值先进行求和，然后再对所有线程求和***************/
for(i=1;i<warp_size;i*=2)
{
t_temp = __shfl_up(temp,i,warp_size);
if(laneid>=i)
{
temp+= t_temp;
}
}
//这里得出每个线程束的前缀和，且最后一个为最大
if(laneid==(warp_size-1))
{
share_sum[warpid]=temp;
}
__sychthread();
if(!tid)
{
for(i=1;i<(N+blockDim.x-1)/warp_size;i++)
{
share_sum[i]=share_sum[i]+share_sum[i-1];
}
}
__sychthread();
if((laneid!=(warp_size-1))&&(warpid>0))
{
temp+=share_sum[warpid-1];
}
__sychthread();
if(temp1)
{
b[temp-1]=tid;
}
}

这段代码在调试的时候一直出问题，从第一个block 中读取的数据是没有错的，但此后
的都是0，我只好厚着脸皮请教师兄。师兄给出的结果是出现了读后写的问题，劝我说这样
写虽然效率高，但只能块同步而不能所有线程同步，让我试着分开来写，于是我把程序分
成了三段以确保所有线程的同步，果然就对了。感谢师兄！