c++ - 如何自动确定2D阵列的CUDA block 大小和网格大小？

Question

如何在CUDA中自动确定2D数组(例如图像处理)的块大小和网格大小？
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()用于2D数组？
潜在解决方案
我对潜在的解决方案有所了解:
如果我先计算线程号，然后使用cudaOccupancyMaxPotentialBlockSize()计算2D数组的块大小怎么办:

////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中的方法不可行，我可以使用上述方法吗？

Answer 1

Question 1 Can I calculate block/grid size using this method?

没有。
重要的是要记住，这些API调用提供的占用率是每个块的最大线程数，而不是块的最大大小。如果您在每个方向上两次运行API，则将两个值组合在一起时，可能会得到非法的块大小。例如，如果一个内核的最大占用线程数为256，则最终可能会出现256 x 256的块大小，这远远大于每个块的1024个总线程，因此启动失败。

Question 2 If the method in Question 1 is not feasible, can I use the method above?

从原理上讲，这应该起作用，尽管您会付出很小的性能损失，因为整数模运算在GPU上并不是特别快。或者，您可以根据API返回的每个块的最大线程数来计算满足您需求的2D块大小。
例如，如果您只希望块尺寸为32个线程的块映射到数据的主要顺序(以进行内存合并)，则只需将线程数除以32(请注意，API始终会返回一轮每个块32个线程的倍数，因为这是扭曲大小)。因此，举例来说，如果从API返回的每个块的线程数是384，则您的块大小将为32 x 12。
如果您真的想要某种使用正方形块的 slice 方案，那么很容易就能得出只有64(8 x 8)，256(16 x 16)，576(24 x 24)和1024(32 x 32) )是可行的块大小，它既是整数又是32的整数倍。在这种情况下，您可能希望选择更大的块大小，该大小小于或等于API返回的线程总数。
最终，您如何选择执行此操作将取决于内核代码的要求。但是，当然可以设计一种与CUDA当前公开的块尺寸调整API兼容的2D块尺寸调整方案。