mio*_*ura 11 d image-processing opencl
我正在开发一个处理大图像的OpenCL 1.2应用程序.目前,我正在测试的图像是16507x21244像素.我的内核在一个循环中运行,该循环对图像块进行操作.内核需要32bpp(rgba)的图像块,并将float4像素块传递出去.
让我们将(正方形)块的一侧(以像素为单位)定义为块大小.也就是说,8192x8192像素正方形的块大小为8192.当然,在右侧和底侧,如果图像不能被块大小干净地分割,则我们有较小的矩形块.我的代码处理这个问题,但对于本文的其余部分,让我们为了简单起见忽略它.
我试图确定我在循环的每次迭代中可以操作的最大块大小,以及最佳块大小(可能不是最大块大小).
作为参考,这是我的机器上的clinfo实用程序报告的信息.我Geforce GTX 560 Ti使用他们专有的Linux驱动程序在Nvidia平台上运行我的内核.
我最初的天真假设是我可以操作最大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
解决我们得到的块大小
chunk size = sqrt(1GiB/20)
插入OpenCL报告的内存量(略小于1GiB - 1023 MiB)并对结果进行分层,我们得到:
floor(sqrt(1072889856/20)) = 7324
但是,仍然会产生7324的块大小CL_MEM_OBJECT_ALLOCATION_FAILURE.
我的下一个猜测是我们无法传递大于最大分配大小的图像,OpenCL报告为我的卡的268222464字节.因为我的输出图像具有更大的像素宽度,它将决定我的块大小.
floor(sqrt(268222464/16)) = 4094
嘿,这确实有效!现在如果我们试图变大呢？令我惊讶的是,它并没有失败.通过反复试验,我将6784缩小为实际的最大块大小.在6785,它开始抱怨CL_MEM_OBJECT_ALLOCATION_FAILURE.我不知道为什么max似乎是6784,我不知道这是否可重复或者值是否波动(例如视频内存中存在的其他状态会影响它能保持多少.)我也发现运行时块大小为6784比基于最大分配的大小运行慢几秒.我想知道这是否是因为OpenCL需要在引擎盖下执行多个(昂贵的)分配？我还注意到OpenCL能够报告的"内核参数的最大大小"(CL_DEVICE_MAX_PARAMETER_SIZE).但是,这个价值似乎是假的.如果我只能传入4096个字节,那么我将限制为16x16像素!
所以我留下了两个基本问题:
作为一个额外的问题,我是否有任何好的资源可以用于关于低级OpenCL硬件交互的这种性质的未来问题？
最后,我将为同行评审提供一些代码片段; 我会非常感谢任何建设性的批评!
一如既往,提前感谢您的帮助!
回答您的直接问题：
1)要确定可用于单个内核操作的绝对最大块大小，必须知道“块大小”所指的是什么。例如，OpenCL 内存结构中有五种定义的内存模型。其中之一是主机内存，我们将忽略它。其他四个是全局的、恒定的、本地的和私有的。
要获取有关您的硬件及其支持功能的任何信息，我强烈建议您访问底部记录的 Khronos API 文档。您可以收集大量有关您的设备的元数据。例如，查询设备可以支持的 2D 和/或 3D 图像的最大高度和最大宽度。我还建议查看 CL_DEVICE_LOCAL_MEM_SIZE 和 CL_DEVICE_MAX_COMPUTE_UNITS 来定义您的工作组。甚至允许 CL_DEVICE_MAX_MEM_ALLOC_SIZE 查询。
指出您对性能的担忧是，为您提供的用于工作的内存大小是工作组或项目的最佳最大大小（取决于查询）。可能发生的情况是内存溢出到全局空间。这需要在不同的工作线程之间分配更多的内存，从而导致性能下降。不能 100% 确定该声明，但当您超过建议的缓冲区大小时，它很可能是问题的一部分。
2)确定最快的块大小不需要反复试验。在 Addison-Wesley 出版的《OpenCL 编程指南》一书中，有一节介绍如何在主机应用程序中使用事件进行分析。有一些函数可以被分析。这些功能如下：
要启用此分析，当创建队列时，需要设置标志 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); 可用于提取计时数据。然后，您可以让主机应用程序根据需要处理这些数据，例如：
通过使用此分析信息，您可以通过软件或分析方式确定最快的“块大小”，然后全面使用该块大小的常量。
额外问题一些很好的资源是 Addison Wesley 出版的“OpenCL 编程指南”，作者是 Aaftab Munshi、Benedict R. Gaster、Timothy G. Mattson、James Fung 和 Dan Ginsburg。我还想说Khronos 文档有很多信息。
作为旁注，您正在主机代码中的双重嵌套循环内运行此内核...这打破了使用并行编程的全部原因。尤其是在图像上。我建议重构您的代码并研究 GPU 操作的并行编程模型。还对 OpenCL 中设置和使用内存屏障进行一些研究。英特尔和英伟达在这方面有一些很棒的论文和例子。 最后，API 文档始终可用
| 归档时间： |
|
| 查看次数： |
902 次 |
| 最近记录： |