CUDA.如何展开前32个线程，以便并行执行

CUDA. How to unroll first 32 threads so they will be executed in parallel?

本文关键字：并行执行 线程 32个 何展开 CUDA      更新时间：2023-10-16

我知道"每个经线包含连续的线程，增加第一个经线包含线程0的线程id "所以前32个线程应该在第一个经线。我还知道，在一个warp中的所有线程在任何可用的流多处理器上同时执行。

据我所知，因为没有必要在线程同步，如果只有一个经纱正在执行。但是如果我在倒数第二个if块中删除任何__syncthreads()，下面的代码会产生错误的答案。我试图找出原因，但却一无所获。我非常希望得到你的帮助，你可以告诉我这个代码有什么问题吗?为什么我不能只留下最后一个__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);
...
}

注:我用的是GT560Ti

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

__shared__ volatile int data [BLOCK_SIZE];

你看到的问题是费米架构和编译器优化的产物。费米架构缺乏直接操作共享内存的指令(它们存在于G80/90/GT200系列中)。因此，所有内容都被加载到注册、操作并存储回共享内存中。但是编译器可以自由地推断，如果将一系列操作暂存到寄存器中，而不需要从共享内存中加载和存储，则代码可以编写得更快。这是非常好的，除了，当您依赖于同一warp内操作共享内存的线程的隐式同步时，就像在这种缩减代码中一样。

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

这个问题在Fermi的编程说明中讨论，它随CUDA工具包一起发货(或可能发货)。

相关文章：