7

我知道“每个经线都包含连续的、递增的线程 ID，第一个经线包含线程 0”，因此前 32 个线程应该在第一个经线中。另外我知道一个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] += data [tid + s]; 
  __syncthreads (); 
 } 

 if ( tid < 32 )
 { 
  data [tid] += data [tid + 32];
  __syncthreads (); 
  data [tid] += data [tid + 16];
  __syncthreads (); 
  data [tid] += data [tid + 8];
  __syncthreads (); 
  data [tid] += data [tid + 4];
  __syncthreads (); 
  data [tid] += data [tid + 2];
  __syncthreads (); 
  data [tid] += data [tid + 1];
  __syncthreads (); 
 }
 if ( tid == 0 )
  outData [blockIdx.x] = data [0];
}

void main()
{
...
 reduce<<<dim3(128), dim3(128)>>>(dev_data, dev_res);
...
}

PS我用的是GT560Ti

4

1 回答 1

7

您应该将共享内存变量声明为 volatile：

__shared__ volatile int data [BLOCK_SIZE];

您看到的问题是 Fermi 架构和编译器优化的产物。Fermi 架构缺乏直接操作共享内存的指令（它们存在于 G80/90/GT200 系列中）。因此，所有内容都被加载以注册、操作并存储回共享内存。但是编译器可以自由地推断，如果在寄存器中进行一系列操作，而不需要从共享内存进行中间加载和存储，那么代码可以变得更快。这非常好，除非您依赖于同一扭曲中的线程的隐式同步来操作共享内存，就像在这种缩减代码中一样。

通过将共享内存缓冲区声明为易失性，您将强制编译器在缩减的每个阶段之后强制执行共享内存写入，并恢复 warp 中线程之间的隐式数据同步。

这个问题在 CUDA 工具包随附（或可能随附）的 Fermi 编程说明中进行了讨论。

于 2012-12-21T17:07:00.537 回答