CUDA 块和网格尺寸效率答案

【问题标题】：CUDA Block and Grid size efficienciesCUDA 块和网格尺寸效率
【发布时间】：2011-08-14 04:54:48
【问题描述】：

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

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

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

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

【问题讨论】：

标签： optimization cuda gpgpu

【解决方案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;
    }
}

为了启动这个内核，执行参数计算如下：

我们允许每个块最多 4 个 warp（因此 128 个线程）。通常，您会以最佳数量来解决此问题，但在这种情况下，内核通常会在非常小的向量上调用，因此具有可变块大小是有意义的。
然后我们根据总工作量计算块数，最多 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() );
}

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

【讨论】：

griddim 区域有什么想法吗？
没那么棒。我在发布代码之前对代码进行了一些清理，并在过程中的块大小计算中引入了一个错误。现在已修复，但希望你能明白这一点......
@talonmies，这很漂亮。我的一个查询是关于这个的。 "for(int i=imin; i
@Pavan：是的，这确实意味着某些块将比其他块更早完成，并且“最后一个”块将有一些扭曲发散。但总的来说，我仍然发现替代方案更好，比如在内核启动结束时拥有“一半”的 GPU 块。保持内核块常驻有助于分摊“设置”索引和平方根计算，从而降低它们对整体性能的影响。
@talonmies，我并不是说您需要启动更多区块，只是工作可能已经均匀地分布在各个区块中。就像如果你有 n = 1.5 * 步幅，而不是退出一半的块并为其余的执行另一个步骤，你可以尝试退出每个块的一半经纱并将工作分散到所有块上。只是提供一个想法，因为这就是我通常做事的方式。这里的这个（你的代码）有点新，对于一些应用程序来说可能是个好主意。我需要测试一下:)

【解决方案2】：

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

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

2) 分配网格大小的好方法（即块数）这里有点棘手。仅仅因为我们可以发布 10,000 个区块通常不是最好的做事方式。在硬件中切换块进出是昂贵的。要考虑的两件事是每个块使用的共享内存，以及可用的 SP 总数，并求解最佳数量。

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

【讨论】：

Pavan：你能指出这个计算在 Thrust 中的什么地方发生吗？
@Ashwin:thrust::detail::backend::cuda::detail::launch_closure 包含所有血淋淋的细节。

【解决方案3】：

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

【讨论】：

嗯，你说对了一部分。有 16 个线程（半经线）而不是说 14 是有道理的，而不是一直到 256。

【解决方案4】：

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

这个site 有一些很棒的启发式方法。一些一般亮点：

为每个网格选择块

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

选择每个块的线程

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

【讨论】：