cuda - 如何在给定 block 中的线程之间共享公共(public)值？

Question

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

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

__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

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

constant memory的优点是它经过专用缓存，所以不会污染L1，如果你每 block 访问的数据量比较小，在每个 block 内第一次访问之后，你在每个线程 block 中的初始强制未命中之后，应该始终命中缓存。

您也不需要使用任何共享内存或从全局内存传输到共享内存。