如何自动确定2D阵列的CUDA块大小和网格大小？

问题描述

如何在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,image.width()
);
cudaOccupancyMaxPotentialBlockSize
(
    &y_min_grid_size,&y_block_size,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,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中的方法不可行，我可以使用上面的方法吗？

解决方法

问题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块大小调整方案。