CUDA块和网格尺寸效率

Question

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

是"基于问题集设置块和网格大小"的情况还是值得将块维度分配为2的因子并且有一些内核逻辑来处理溢出？

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

我已经仔细阅读过Kirk的"大规模并行编程器编程",但它并没有触及这个领域.

Answer 1

通常情况下,设置块大小以获得最佳性能,并根据工作总量设置网格大小.大多数内核在每个Mp上都有一个"最佳位置"的warp数,它们效果最好,你应该做一些基准测试/分析,看看它在哪里.您可能仍需要内核中的溢出逻辑,因为问题大小很少是块大小的倍数.

编辑:给出一个具体的例子,说明如何对一个简单的内核进行这种操作(在这种情况下,自定义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个warp(所以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() );
}

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