如何在给定块中的线程之间共享公共值？

Question

我有一个内核,对于给定块中的每个线程,它计算具有不同迭代次数的for循环.我使用大小为N_BLOCKS的缓冲区来存储每个块所需的迭代次数.因此,给定块中的每个线程必须知道其块的特定迭代次数.

但是,我不确定哪种方式最好(性能说话)来读取值并将其分发给所有其他线程.我只看到一个好方法(请告诉我是否有更好的东西):将值存储在共享内存中并让每个线程读取它.例如:

__global__ void foo( int* nIterBuf )
{
   __shared__ int nIter;

   if( threadIdx.x == 0 )
      nIter = nIterBuf[blockIdx.x];

   __syncthreads();

   for( int i=0; i < nIter; i++ )
      ...
} 

其他更好的解决方案？我的应用程序将使用大量数据,因此我希望获得最佳性能.

谢谢!

Answer 1

在块中的所有线程上统一的只读值可能最好存储在__constant__数组中.在某些CUDA架构(如Fermi(SM 2.x))上,如果使用C++ const关键字声明数组或指针参数并在块内统一访问它(即索引仅取决于blockIdx,而不是threadIdx),则编译器可能会自动促进对常量记忆的引用.

常量内存的优点是它通过专用缓存,因此它不会污染L1,如果每个块访问的数据量相对较小,在每个块中第一次访问后,你应该总是命中在每个线程块中的初始强制未命中后的缓存中.

您也不需要使用任何共享内存或从全局内存转移到共享内存.