关于树约简中共享内存的问题

NVIDIA第22页 Optimizing Parallel Reduction in CUDA ,

__device__ void warpReduce(volatile int* sdata, int tid) {
sdata[tid] += sdata[tid + 32];
sdata[tid] += sdata[tid + 16];
sdata[tid] += sdata[tid + 8];
sdata[tid] += sdata[tid + 4];
sdata[tid] += sdata[tid + 2];
sdata[tid] += sdata[tid + 1];
}

// later…
for (unsigned int s=blockDim.x/2; s>32; s>>=1) {
if (tid < s)
    sdata[tid] += sdata[tid + s];
__syncthreads();
}
if (tid < 32) warpReduce(sdata, tid);

我不确定CUDA如何处理函数最后一行中的操作 warpReduce(...) : sdata[tid] += sdata[tid + 1]; . 如果扭曲具有线程ID tid 从0到31,是否 每日三次 此时0和1发生冲突?我的意思是,CUDA将如何选择 sdata[0] += sdata[0 + 1]; 和 sdata[1] += sdata[1 + 1]; ? 如果它运行 sdata[0]+=sdata[0+1]; 首先,结果是正确的。或者,CUDA采取措施防止这种混淆/冲突?