小编Вит*_*ров的帖子

我知道"每个warp包含连续增加的线程ID的线程,第一个warp包含线程0",因此前32个线程应该在第一个warp中.另外我知道一个warp中的所有线程都在任何可用的Streaming Multiprocessor上同时执行.

据我所知,因为如果只执行一个warp,就不需要线程同步.但是如果我删除__syncthreads()倒数第二个if块中的任何一个,下面的代码会产生错误的答案.我试图找到原因,但最终没有任何结果.我真的希望得到你的帮助,所以你可以告诉我这段代码有什么问题？为什么我不能只留下最后__syncthreads()得到正确答案？

#define BLOCK_SIZE 128

__global__ void reduce ( int * inData, int * outData )
{
 __shared__ int data [BLOCK_SIZE]; 
 int tid = threadIdx.x; 
 int i   = blockIdx.x * blockDim.x + threadIdx.x; 

 data [tid] = inData [i] + inData [i + blockDim.x / 2 ];
 __syncthreads ();

 for ( int s = blockDim.x / 4; s > 32; s >>= 1 ) 
 {
  if ( tid < s ) 
   data [tid] += …