有关ReduceSum中线程同步的问题

otakuxiang

    我在使用CUDA对核函数中每个线程产生的6x6矩阵做ReduceSum时，发现程序结果与中间添加的__syncthread()个数相关，很奇怪，代码如下

1. __syncthreads();
   
2.         for(int ii = 0 ; ii < 6 ; ii ++) for(int jj = 0 ; jj < 6 ; jj ++){
   
3.     dim_shared1[tid] = new_H(ii,jj);
   
4.     __syncthreads();
   
5.     for(int s = 128;s>0;s>>=1){
   
6.         if(tid < s){
   
7.             dim_shared1[tid] += dim_shared1[tid+s];
   
8.         }
   
9.         __syncthreads();
   
10.         __syncthreads();
    
11.         __syncthreads();
    
12.         __syncthreads();
    
13.         __syncthreads();
    
14.         __syncthreads();
    
15.         __syncthreads();
    
16.         __syncthreads();
    
17.     }
    
18.     __syncthreads();
    
19.     __syncthreads();
    
20.     __syncthreads();
    
21.     if(tid == 0){
    
22.         atomicAdd(&acc_H->at(ii,jj),dim_shared1[0]);

复制代码

如果中间一串__syncthreads();变成一个__syncthreads();程序结果就会出错，请问这是什么问题？