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，我只好厚着脸皮请教师兄。师兄给出的结果是出现了读后写的问题，劝我说这样
    写虽然效率高，但只能块同步而不能所有线程同步，让我试着分开来写，于是我把程序分
    成了三段以确保所有线程的同步，果然就对了。感谢师兄！