optimization - CUDA block 和网格尺寸效率

Question

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

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

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

我已经仔细阅读过Kirk的“大规模并行处理器编程”，但是它并没有真正涉及到这一 Realm 。

Answer 1

通常是设置块大小以获得最佳性能，并根据工作总量设置网格大小的情况。大多数内核在每个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;
    }
}

为了启动该内核，执行参数的计算如下:

每个块最多允许4个经线(因此128个线程)。通常，您可以将其固定为最佳数目，但是在这种情况下，内核通常是在非常小的 vector 上调用的，因此具有可变的块大小在一定程度上是有意义的。

然后，我们根据工作总量(最多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() );
}

也许这为如何设计一种“通用”方案以针对输入数据大小设置执行参数提供了一些提示。