CUDA 块和网格大小效率

Question

在 cuda 中处理动态大小的数据集的建议方法是什么？

是“根据问题集设置块和网格大小”的情况，还是值得将块尺寸分配为 2 的因子并有一些内核逻辑来处理溢出？

我可以看出这对于块尺寸来说可能很重要，但是这对于网格尺寸有多大影响呢？据我了解，实际的硬件约束停止在块级别（即分配给具有一定数量的 SP 的 SM 的块，因此可以处理特定的扭曲大小）。

我仔细阅读了柯克的“编程大规模并行处理器”，但它并没有真正涉及这个领域。

Answer 1

通常是为了获得最佳性能而设置块大小，并根据工作总量设置网格大小。大多数内核在每 Mp 上都有一个“最佳点”数量的扭曲，在该点上它们工作得最好，您应该做一些基准测试/分析来看看它在哪里。您可能仍然需要内核中的溢出逻辑，因为问题大小很少是块大小的整数倍。

编辑：
给出一个具体示例，说明如何对简单内核执行此操作（在本例中，自定义 BLAS 1 级 dscal 类型操作作为压缩对称带矩阵 Cholesky 分解的一部分完成）：

// 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() );
}

也许这给出了一些关于如何设计一个“通用”方案来根据输入数据大小设置执行参数的提示。

Answer 2

好吧，我想我们在这里处理两个问题。

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

）分配网格大小（即块数）的好方法
这里有点棘手。因为我们可以而仅仅启动 10,000 个区块通常并不是最好的做法。将块移入或移出硬件的成本很高。需要考虑的两件事是每个块使用的共享内存和可用 SP 的总数，并求解最佳数量。

您可以从 推力。不过，可能需要一段时间才能弄清楚代码内部发生了什么。

Answer 3

我认为通常最好根据问题集设置块和网格大小，特别是出于优化目的。拥有不执行任何操作的额外线程实际上没有意义，并且可能会降低程序的性能。

Answer 4

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

这个网站有一些伟大的启发法。一些一般要点：

选择每个网格的块

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

选择每个块的线程

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

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