CUDA 块和网格大小的效率

Question

如何处理CUDA中大小不固定的数据集？

是根据问题集设置块和网格大小，还是值得分配2的因子作为块维度，并有一些内核逻辑来处理溢出?

我可以看出这对块维度很重要，但这对网格维度有多重要呢？据我所知，实际硬件约束仅限于块级别（即分配给具有一组SP的SM的块，因此可以处理特定的warp大小）。

我已经浏览了Kirk的“大规模并行处理器编程”，但它并没有真正涉及这个领域。

- Bolster

Answer 1

通常情况下，为了达到最佳性能，需要设置块大小并根据总工作量确定网格大小。大多数内核都有一定数量的线程束（warps）每个MP运行时表现最佳，您应该进行一些基准测试/性能分析以确定该点。由于问题大小很少是块大小的整数倍，所以您可能仍需要在内核中使用溢出逻辑。

编辑：为了给出一个具体的例子，说明如何对简单内核执行此操作（在这种情况下，它是作为打包对称带状矩阵的Cholesky分解的一部分完成的自定义BLAS第1级dscal类型操作）：

// Fused square root and dscal operation
__global__ 
void cdivkernel(const int n, double *a)
{
    __shared__ double oneondiagv;

    int imin = threadIdx.x + blockDim.x * blockIdx.x;
    int istride = blockDim.x * gridDim.x;

    if (threadIdx.x == 0) {
        oneondiagv = rsqrt( a[0] );
    }
    __syncthreads();

    for(int i=imin; i<n; i+=istride) {
        a[i] *= oneondiagv;
    }
}

启动此内核，执行参数计算如下:

1. 我们允许每块最多4个线程束（即128个线程）。通常情况下，您会将其固定为最佳数量，但在这种情况下，内核经常针对非常小的向量调用，因此具有可变块大小是有意义的。
2. 然后，根据总工作量计算块计数，最多112个总块，这相当于14 MP Fermi Telsa上每个MP 8个块。如果工作量超过网格大小，则内核会迭代。

包含执行参数计算和内核启动的结果包装器函数如下:

// Fused the diagonal element root and dscal operation into
// a single "cdiv" operation
void fusedDscal(const int n, double *a)
{
    // The semibandwidth (column length) determines
    // how many warps are required per column of the 
    // matrix.
    const int warpSize = 32;
    const int maxGridSize = 112; // this is 8 blocks per MP for a Telsa C2050

    int warpCount = (n / warpSize) + (((n % warpSize) == 0) ? 0 : 1);
    int warpPerBlock = max(1, min(4, warpCount));

    // For the cdiv kernel, the block size is allowed to grow to
    // four warps per block, and the block count becomes the warp count over four
    // or the GPU "fill" whichever is smaller
    int threadCount = warpSize * warpPerBlock;
    int blockCount = min( maxGridSize, max(1, warpCount/warpPerBlock) );
    dim3 BlockDim = dim3(threadCount, 1, 1);
    dim3 GridDim  = dim3(blockCount, 1, 1);

    cdivkernel<<< GridDim,BlockDim >>>(n,a);
    errchk( cudaPeekAtLastError() );
}

也许这能提供一些关于如何设计针对输入数据大小设置执行参数的"通用"方案的线索。

- talonmies

---

在griddim领域有什么想法吗？ - Bolster

@talonmies，这太棒了。我只有一个问题，关于这个。 "for(int i=imin; i<n; i+=istride)" 看起来这会使得工作在块之间有点不平衡。特别是当 n = 1.5 * stride 时。 - Pavan Yalamanchili

@Pavan：是的，这意味着有些块会比其他块更早地完成，而“最后”一个块将有一些warp分歧。但总体而言，我仍然发现它比其他选择更好，比如在内核启动的末尾拥有“半个”GPU的块。保持内核块常驻有助于分摊“设置”索引和平方根计算的影响，从而降低它们对整体性能的影响。 - talonmies

@talonmies，我并不是说你需要启动更多的块，只是工作可能会均匀地分布在这些块中。比如，如果你有n = 1.5 * stride，那么你可以尝试每个块退役一半的线程束，并将工作分配到所有块中，而不是让一半的块退役并为其余的块执行另一个步骤。只是提供一个想法，因为这是我通常做事情的方式。这里（你的代码）稍微有点新颖，可能对一些应用程序来说是一个好主意。我需要测试一下 :) - Pavan Yalamanchili

在这里使用max()的目的是什么： max(1, min(4, warpCount))？ - syntagma

Answer 2

好的，我想我们需要处理两个问题。

1）分配块大小（即线程数）的好方法 这通常取决于您正在处理的数据类型。您正在处理向量吗？您正在处理矩阵吗？建议的方法是保持线程数为32的倍数。因此，在处理向量时，启动256 x 1，512 x 1块可能很好。同样，在处理矩阵时，选择32 x 8、32 x 16等。

2）分配网格大小（即块数）的好方法 这有些棘手。仅仅因为我们可以启动10000个块并不是通常做事情的最佳方式。切换块进入和退出硬件是昂贵的。要考虑的两个因素是每个块使用的共享内存和可用的总SP数量，并求解最优数量。

您可以从thrust找到一个非常好的实现方法。不过可能需要一段时间才能弄清楚代码内部发生了什么。

- Pavan Yalamanchili

---

Pavan：您能指出在Thrust中这个计算发生的位置吗？ - Ashwin Nanjappa

@Ashwin：thrust::detail::backend::cuda::detail::launch_closure 包含了所有的细节。 - talonmies

Answer 3

我认为根据问题集设置块和网格大小通常是最好的选择，特别是为了优化目的。拥有无用的额外线程真的没有意义，可能会使程序的性能变差。

- Adam27X

---

你部分正确。拥有16个线程（半warp）而不是14个是有意义的，而不是一路提高到256个。 - Pavan Yalamanchili

Answer 4

如果您具有动态大小的数据集，那么您可能会遇到一些延迟问题，因为某些线程和块正在等待其他线程完成。

这个网站有一些很好的启发。一些一般性的亮点：

选择每个网格的块数

• 每个网格的块数应该大于或等于多处理器的数量。
• 在内核中更多使用__syncthreads()，块就越多（这样一个块可以在另一个块等待同步时运行）

选择每个块的线程数

• 线程应该是warp大小的倍数（即通常为32）

• 通常最好选择线程数，使每个块的最大线程数（基于硬件）是线程数的倍数。例如，使用256个线程每块的最大线程数为768，比使用512个线程要好，因为多个线程可以同时在一个块上运行。

- Michael Ohlrogge