image-processing - 如何确定我可以安全地传入/传出 OpenCL 内核的最大图像大小？

Question

我正在开发 OpenCL 1.2 处理大型图像的应用程序。
目前，我正在测试的图像是 16507x21244 像素。我的 内核在一个循环中运行，该循环对图像的 block 进行操作.内核将 32bpp (rgba) 的图像 block 放入，并将 float4-pixel block 传递出去。

让我们以像素为单位将(方形) block 的一侧定义为 block 大小。也就是说，一个 8192x8192 像素正方形的 block 大小为 8192 .当然，如果图像不能完全被 block 大小整除，则在右侧和底部我们有更小的矩形 block 。我的代码处理了这个问题，但对于本文的其余部分，为了简单起见，让我们忽略它。

我是试图确定最大块大小 我可以在循环的每次迭代中进行操作，以及最佳 block 大小(可能不是最大块大小)。

供引用，here is the information reported by the clinfo utility on my machine .我在 Geforce GTX 560 Ti 上运行我的内核Nvidia 平台使用其专有的 linux 驱动程序。

我最初的天真假设是我可以在最大 2d 图像大小上进行操作。
但是，这会导致 clEnqueueNDRangeKernel返回错误代码 -4 ( CL_MEM_OBJECT_ALLOCATION_FAILURE )。

仔细想想，这对我来说很有意义。对于 1 GiB 的视频内存，人们期望能够保存单个 16384x16384 像素纹理 (32bpp) 或 8192x8192 像素纹理 (float4)。如果在内核运行时两者都需要缓存在卡上，我们可以预期使用以下内存量:

   4 bytes-per-pixel * chunk size^2 (input image) 
+ 16 bytes-per-pixel * chunk size^2 (output image) 
= 1 GiB total video memory

求解我们得到的 block 大小

chunk size = sqrt(1GiB/20)

插入 OpenCL 报告的内存量(略小于 1GiB - 1023 MiB)并计算结果，我们得到:

floor(sqrt(1072889856/20)) = 7324

但是，7324 的 block 大小仍然会导致 CL_MEM_OBJECT_ALLOCATION_FAILURE .

我的下一个猜测是我们不能传递大于 的图像。最大分配大小 ，对于我的卡，OpenCL 报告为 268222464 字节。因为我的输出图像具有较大的像素宽度，所以它将决定我的 block 大小。

floor(sqrt(268222464/16)) = 4094

嘿，这确实有效! 现在，如果我们尝试做大怎么办？令我惊讶的是，它并没有失败 .通过反复试验，我缩小了 6784 作为实际的最大块大小。在 6785，它开始提示 CL_MEM_OBJECT_ALLOCATION_FAILURE .我不知道为什么最大值似乎是 6784，我不知道这是否可重复或值是否波动(例如视频内存中存在的其他状态会影响它可以容纳多少。)我还发现运行6784 的 block 大小比使用基于最大分配的大小运行慢几秒钟。我想知道这是否是因为 OpenCL 需要在后台执行多个(昂贵的)分配？我还注意到 OpenCL 能够报告的“内核参数的最大大小”( CL_DEVICE_MAX_PARAMETER_SIZE )。然而，这个值(value)似乎是虚假的。如果我只能传入 4096 字节，那将限制我为 16x16 像素!

所以我有两个基本问题:

如何确定绝对最大块大小？

如何确定最快的 block 大小？ (除了试错还有其他方法吗？)

作为奖励问题，有没有 好资源我可以求助于关于 的此类性质的 future 问题。低级 OpenCL 硬件交互？

最后，我将提供一些代码片段供同行评审；我将非常感谢任何 build 性的批评!

The host code module (dlang)

Wrapper that sets up the OpenCL program (dlang)

The kernel code

一如既往， 提前感谢您的帮助!

Answer 1

要回答您的直接问题:

1) 要确定可用于单个内核操作的绝对最大块大小，必须知道“ block 大小”指的是什么。例如，OpenCL 内存结构中有五个已定义的内存模型。其中之一是主机内存，我们将忽略它。其他四个是全局的、常量的、本地的和私有(private)的。

要获取有关您的硬件可以支持什么的任何信息，我强烈建议您访问底部记录的 Khronos API 文档。您可以收集大量有关您设备的元数据。例如，查询设备可以支持的 2D 和/或 3D 图像的最大高度和最大宽度。我还建议查看 CL_DEVICE_LOCAL_MEM_SIZE 和 CL_DEVICE_MAX_COMPUTE_UNITS 来定义您的工作组。甚至还有一个允许的 CL_DEVICE_MAX_MEM_ALLOC_SIZE 查询。

指出您对性能的关注是因为为您提供的内存大小是工作组或项目的最佳最大大小(取决于查询)。可能发生的事情是内存溢出到全局空间。这需要在不同的工作人员之间分配更多的内存，从而导致性能下降。该声明不是 100% 确定的，但是当您超过建议的缓冲区大小时，它很可能是问题的一部分。

2) 确定最快的 block 大小不需要反复试验。在 Addison-Wesley 出版的“OpenCL Programming Guide”一书中，有一节介绍了在主机应用程序中使用事件进行分析。有一组功能可以被分析。这些功能如下:

clEnqueue{读取|写入|映射}缓冲区

clEnqueue{读|写}BufferRect

clEnqueue{读|写| map }图片

clEnqueueUnmapMemObject

clEnqueueCopyBuffer

clEnqueueCopyBufferRect

clEnqueueCopyImage

clEnqueueCopyImageToBuffer

clEnqueueCopyBufferToImage

clEnqueueNDRange内核

clEnqueueTask

clEnqueueNativeKernel

clEnqueueAcquireGLObjects

clEnqueueReleaseGLObject

要启用此分析，一旦创建队列，需要设置标志 CL_QUEUE_PROFILING_ENABLE。然后函数 clGetEventProfilingInfo(cl_event event, cl_profiling_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret);可用于提取时序数据。然后，您可以让主机应用程序随意处理这些数据，例如:

写入分析日志

写入输出缓冲区

负载均衡

通过使用此分析信息，您可以通过您的软件或分析确定最快的“ block 大小”，然后全面使用该 block 大小的常数。

奖金问题 一些很好的资源将是 Addison Wesley 出版的“OpenCL Programming Guide”，作者是 Aaftab Munshi、Benedict R. Gaster、Timothy G. Mattson、James Fung 和 Dan Ginsburg。我还想说 Khronos docs有很多信息。

作为旁注您正在主机代码中的双重嵌套循环内运行此内核......这种破坏了使用并行编程的全部原因。尤其是在图像上。我建议重构您的代码并研究 GPU 操作的并行编程模型。还对在 OpenCL 中设置和使用内存屏障进行一些研究。英特尔和英伟达在这方面有一些很棒的论文和例子。 Finally, the API docs are always available