如何自动确定二维数组的 CUDA 块大小和网格大小？答案

【问题标题】：How to automatically determine the CUDA block size and grid size for a 2D array?如何自动确定二维数组的 CUDA 块大小和网格大小？
【发布时间】：2021-01-07 01:25:16
【问题描述】：

如何在 CUDA 中自动确定二维数组（例如图像处理）的块大小和网格大小？

CUDA 具有 cudaOccupancyMaxPotentialBlockSize() 函数，可以自动计算 cuda 内核函数的块大小。见here。在这种情况下，它适用于一维数组。

就我而言，我有一张 640x480 的图片。

如何确定块/网格大小？我用：

////image size: 640x480


int x_min_grid_size, x_grid_size, x_block_size;
int y_min_grid_size, y_grid_size, y_block_size;

cudaOccupancyMaxPotentialBlockSize
(
    &x_min_grid_size, &x_block_size,
    my_cuda_kernel,
    0, image.width()
);
cudaOccupancyMaxPotentialBlockSize
(
    &y_min_grid_size, &y_block_size,
    my_cuda_kernel,
    0, image.height()
);

x_grid_size = (image.width()  + x_block_size - 1) / x_block_size;
y_grid_size = (image.height() + y_block_size - 1) / y_block_size;

dim3 grid_dim(x_grid_size, y_grid_size);
dim3 block_dim(x_block_size, y_block_size);

my_cuda_kernel<<<grid_dim, block_dim>>>(<arguments...>)

////check cuda kernel function launch error
cudaError_t error = cudaGetLastError();
if(cudaSuccess != error)
{
    std::cout<<"CUDA Error! "<<cudaGetErrorString(error)<<std::endl;
    exit(1);
}
cudaDeviceSynchronize();

问题 1 我可以使用这种方法计算块/网格大小吗？

对于这段代码，内核函数启动后出现错误。

CUDA Error! invalid configuration arguments

如果我手动设置x_block_size = 32; y_block_size = 32，它可以工作并且没有错误。

我能问一下为什么 CUDA 得到invalid configuration arguments 错误消息吗？二维数组好像不能直接用cudaOccupancyMaxPotentialBlockSize()？

潜在解决方案 我对潜在的解决方案有了一个想法：

如果我先计算线程数，然后使用cudaOccupancyMaxPotentialBlockSize() 计算二维数组的块大小：

////total_thread_num = 640x480 = 307200
int total_thread_num = image.width * image.height;

////compute block/grid size
int min_grid_size, grid_size, block_size;
cudaOccupancyMaxPotentialBlockSize
(
    &min_grid_size, &block_size,
    my_cuda_kernel,
    0, total_thread_num
);

grid_size = (total_thread_num + block_size - 1) / block_size;

//launch CUDA kernel function
my_cuda_kernel<<<grid_size, block_size>>>(<arguments...>);

在my_cuda_kernel中，根据图片大小计算对应的索引：

__global__ void my_cuda_kernel()
{
    //compute 2D index based on 1D index;
    unsigned int idx = BlockIdx.x * blockDim.x + threadIdx.x;
    unsigned int row_idx = idx / image.width;
    unsigned int col_idx = idx % image_width;

    /*kernel function code*/

}

问题 2 如果问题1中的方法不可行，我可以使用上面的方法吗？

【问题讨论】：

标签： c++ image-processing cuda

【解决方案1】：

问题 1 我可以使用这种方法计算块/网格大小吗？

没有。

重要的是要记住，这些 API 调用提供的占用最大化每个块的线程数，而不是块尺寸。如果你在每个方向上运行 API 两次，当这两个值相结合时，你可能会得到一个非法的块大小。例如，如果内核的占用最大化线程数为 256，那么您最终可能会得到 256 x 256 的块大小，这远远大于每个块的 1024 个总线程，因此启动失败。

问题2 如果问题1中的方法不可行，我可以使用上面的方法吗？

原则上，这应该可行，尽管您会受到一点性能损失，因为整数模运算在 GPU 上并不是特别快。或者，您可以根据 API 返回的每个块的最大线程数来计算满足您需求的 2D 块大小。

例如，如果您只想要块维度中具有 32 个线程的块，您将映射到数据的主要顺序（用于内存合并），那么只需将线程数除以 32（注意 API 将始终每个块返回 32 个线程的整数倍，因为这是扭曲大小）。因此，例如，如果从 API 返回的每个块的线程数为 384，那么您的块大小将为 32 x 12。

如果您真的想要某种使用方形块的平铺方案，那么很容易计算出只有 64 (8 x 8)、256 (16 x 16)、576 (24 x 24) 和 1024 ( 32 x 32) 是可行的块大小，它既是 32 的平方数，又是整数的整数倍。在这种情况下，您可能希望选择小于或等于 API 返回的总线程数的较大块大小。

最终您选择如何执行此操作将取决于您的内核代码的要求。但当然可以设计一个与 CUDA 当前公开的块大小 API 兼容的 2D 块尺寸标注方案

【讨论】：