使用 OpenCL™ 2.0 读写图像
5.00/5 (1投票)
使用 OpenCL™ 2.0 读写图像
英特尔® 开发人员专区 提供跨平台应用程序开发的工具和操作指南、平台和技术信息、代码示例以及同行专业知识，旨在帮助开发人员创新并取得成功。加入我们的 Android、物联网、英特尔® 实感™ 技术 和 Windows 社区，以下载工具、访问开发套件、与志同道合的开发人员分享想法，并参与编程马拉松、竞赛、路演和本地活动。
致谢
我们感谢 Javier Martinez、Kevin Patel 和 Tejas Budukh 在审阅本文及相关示例时提供的帮助。
引言
在 OpenCL™ 2.0 之前，无法在同一个内核中对图像进行读写。图像总是可以声明为 "CL_MEM_READ_WRITE"，但一旦图像传递给内核，它必须是 "__read_only" 或 "__write_only"。
input1 = clCreateImage(
oclobjects.context,
CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,
&format,
&desc,
&input_data1[0],
&err );
SAMPLE_CHECK_ERRORS( err );
__kernel void Alpha( __read_write image2d_t inputImage1,
__read_only image2d_t
inputImage2,
uint width,
uint height,
float alpha,
float beta,
int gamma )
这项直观的附加功能附带了一些注意事项，将在下一节中讨论。
读写图像的价值
虽然图像卷积对于新的读写图像功能而言效果不佳，但任何需要就地完成的图像处理技术都可能受益于读写图像。一个可以有效利用该过程的例子是图像合成。
在 OpenCL 1.2 及更早版本中，图像使用 "__read_only" 和 "__write_only" 修饰符进行限定。在 OpenCL 2.0 中，图像可以使用 "__read_write" 修饰符进行限定，并将输出复制到输入缓冲区。这减少了所需的资源数量。
由于 OpenCL 1.2 图像是只读或只写图像，对图像进行就地修改需要将图像视为缓冲区并对缓冲区进行操作（参见 cl_khr_image2d_from_buffer：https://software.intel.com/en-us/articles/using-image2d-from-buffer-extension）。
当前的解决方案是将图像视为缓冲区，并对缓冲区进行操作。将二维图像视为缓冲区可能不是免费操作，并且会阻止使用 read_images 中可用的钳位和过滤功能。因此，使用 read_write 限定的图像可能更可取。
示例概述
该示例读取两个 Windows 位图图像“input1.bmp”和“input2.bmp”，并将它们放入图像缓冲区中。然后根据 alpha 值（即计算像素方程中的权重因子，可以作为选项传入）对这些图像进行合成。
图像必须是 24 位/32 位图像。输出是 24 位图像。图像必须具有相同的大小。图像的格式也是 ARGB，因此在加载时考虑了这一事实。
ARGB 被转换为 RGBA。改变 beta 值会显著改变输出。
使用示例 SDK
SDK 演示了如何使用读写图像进行图像合成。使用以下命令行选项控制此示例
|
选项 |
描述 |
|
-h, --help |
显示此文本并退出 |
|
-p, --platform number-or-string |
选择平台，其设备将被使用 |
|
-t, --type all | cpu | gpu | acc | default | <OpenCL constant for device type> |
按类型选择执行 OpenCL 内核的设备 |
|
-d, --device number-or-string |
选择执行所有操作的设备 |
|
-i, --infile 24/32 位 .bmp 文件 |
要读取的第一个 .bmp 文件的基本名称。默认为 input1.bmp |
|
-j, --infile 24/32 位 .bmp 文件 |
要读取的第二个 .bmp 文件的基本名称。默认为 input2.bmp |
|
-o, --outfile 24/32 位 .bmp 文件 |
要写入的输出文件的基本名称。对于 OCL1.2，默认为 output.bmp；对于 OCL2.0，默认为 20_output.bmp |
|
-a, --alpha 介于 0 到 1 之间的浮点值 |
非零正值，决定两幅图像在合成中混合的程度。默认 alpha 值为 0.84089642。默认 beta 值为 0.15950358。 |
该示例 SDK 具有许多默认值，允许应用程序在无需任何用户输入的情况下运行。用户将能够使用他们自己的输入 .bmp 文件。这些文件也必须是 24/32 位 bmp 文件。alpha 值用于确定图像 1 相对于图像 2 的显著程度，例如
calculatedPixel = ((currentPixelImage1 * alpha) + (currentPixeImage2 * beta) + gamma);
beta 值通过用 1 减去 alpha 值来确定。
float beta = 1 – alpha;
这两个值决定了图像 1 和图像 2 的加权分布。
伽马值可用于提亮每个像素。默认值为 0。但用户可以提亮整体合成图像。
程序运行示例
读写图像的局限性
屏障不能与需要跨不同工作组同步的图像一起使用。图像卷积需要同步所有线程。图像卷积通常涉及对两个矩阵进行数学运算，从而生成第三个矩阵。图像卷积的一个例子是使用高斯模糊。其他例子包括图像锐化、边缘检测和浮雕。
让我们以高斯模糊为例。高斯滤波器是一个低通滤波器，用于去除高频值。这意味着减少细节并最终产生模糊效果。应用高斯模糊与将图像与高斯函数（通常称为掩码）进行卷积是相同的。为了有效展示读写图像的功能，必须进行水平和垂直模糊处理。
在 OpenCL 1.2 中，这必须分两步完成。一个内核专门用于水平模糊，另一个内核进行垂直模糊。其中一个模糊的结果将用作下一个模糊的输入，具体取决于哪个先完成。
__kernel void GaussianBlurHorizontalPass( __read_only image2d_t inputImage, __write_only image2d_t outputImage, __constant float* mask, int maskSize)
{
int2 currentPosition = (int2)(get_global_id(0), get_global_id(1));
float4 currentPixel = (float4)(0,0,0,0);
float4 calculatedPixel = (float4)(0,0,0,0);
for(int maskIndex = -maskSize; maskIndex < maskSize+1; ++maskIndex)
{
currentPixel = read_imagef(inputImage, imageSampler, currentPosition + (int2)(maskIndex, 0));
calculatedPixel += currentPixel * mask[maskSize + maskIndex];
}
write_imagef(outputImage, currentPosition, calculatedPixel);
}
__kernel void GaussianBlurVerticalPass( __read_only image2d_t inputImage, __write_only image2d_t outputImage, __constant float* mask, int maskSize)
{
int2 currentPosition = (int2)(get_global_id(0), get_global_id(1));
float4 currentPixel = (float4)(0,0,0,0);
float4 calculatedPixel = (float4)(0,0,0,0);
for(int maskIndex = -maskSize; maskIndex < maskSize+1; ++maskIndex)
{
currentPixel = read_imagef(inputImage, imageSampler, currentPosition + (int2)(0, maskIndex));
calculatedPixel += currentPixel * mask[maskSize + maskIndex];
}
write_imagef(outputImage, currentPosition, calculatedPixel);
}
OpenCL 2.0 的想法是将这两个内核合并为一个。使用屏障强制每个水平或垂直模糊完成后再开始下一个模糊。
__kernel void GaussianBlurDualPass( __read_only image2d_t inputImage, __read_write image2d_t tempRW, __write_only image2d_t outputImage, __constant float* mask, int maskSize)
{
int2 currentPosition = (int2)(get_global_id(0), get_global_id(1));
float4 currentPixel = (float4)(0,0,0,0);
float4 calculatedPixel = (float4)(0,0,0,0)
currentPixel = read_imagef(inputImage, currentPosition);
for(int maskIndex = -maskSize; maskIndex < maskSize+1; ++maskIndex)
{
currentPixel = read_imagef(inputImage, currentPosition + (int2)(maskIndex, 0));
calculatedPixel += currentPixel * mask[maskSize + maskIndex];
}
write_imagef(tempRW, currentPosition, calculatedPixel);
barrier(CLK_GLOBAL_MEM_FENCE);
for(int maskIndex = -maskSize; maskIndex < maskSize+1; ++maskIndex)
{
currentPixel = read_imagef(tempRW, currentPosition + (int2)(0, maskIndex));
calculatedPixel += currentPixel * mask[maskSize + maskIndex];
}
write_imagef(outputImage, currentPosition, calculatedPixel);
}
发现屏障无效。使用屏障不能保证在垂直模糊开始之前完成水平模糊，假设您首先进行了水平模糊。这意味着多次运行会产生不一致的结果。屏障可用于同步组内的线程。问题发生的原因是边缘像素从多个工作组读取，并且无法同步多个工作组。我们最初假设可以使用读写图像实现单一高斯模糊被证明是不正确的，因为 OpenCL 中无法同步工作组间的数据依赖性。
参考文献
- OpenCL 2.0 规范：https://www.khronos.org/registry/cl/specs/opencl-2.0.pdf
- OpenCL 编程指南 – A. Munshi 等人 (Pearson 2012)
- OpenCL 编程示例 – Banfer, Ravishekhar
- Shapiro, L. G. & Stockman, G. C: "Computer Vision", 第 137, 150 页。Prentice Hall, 2001
- https://en.wikipedia.org/wiki/Alpha_compositing
- http://mathnathan.com/2010/07/5opencv/
关于作者
Oludemilade Raji 是英特尔视觉与并行计算部门的图形驱动工程师。他从事 OpenCL 编程语言四年，并为英特尔高清显卡驱动程序的开发做出了贡献，包括 OpenCL 2.0 的开发。
Robert Ioffe 是英特尔软件与解决方案部门的技术咨询工程师。他是 OpenCL 编程和在 Intel Iris 和 Intel Iris Pro Graphics 上 OpenCL 工作负载优化的专家，对 Intel 图形硬件有深入的了解。他曾大量参与 Khronos 标准工作，专注于最新功能的原型设计，并确保它们能在 Intel 架构上良好运行。最近，他一直在为 OpenCL 2.0 的嵌套并行性（enqueue_kernel 函数）功能进行原型设计，并编写了许多演示嵌套并行性功能的示例，包括 OpenCL 2.0 的 GPU-Quicksort。他还录制并发布了两段优化简单 OpenCL 内核的视频，以及 OpenCL 2.0 中的 GPU-Quicksort 和 Sierpinski 地毯视频。
您可能还会对以下内容感兴趣