使用 OpenCL™ 2.0 读写图片

致谢

非常感谢 Javier Martinez、Kevin Patel 和 Tejas Budukh 在审核本文和相关示例过程中所提供的帮助。

简介

OpenCL™ 2.0 问世之前,人们无法在同一 kernel 中读写图像。 图像通常声明为 “CL_MEM_READ_WRITE”,但图像传递至 kernel 后,必须是 “__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 );

代码 1. 图像缓冲区可通过 CL_MEM_READ_WRITE 创建

__kernel void Alpha( __read_write image2d_t inputImage1,
__read_only image2d_t inputImage2, 
uint width, 
uint height, 
float alpha, 
float beta, 
int gamma )

代码 2. OpenCL 2.0 推出了在 kernel 中读写图像的功能

新增功能尽管非常直观,但需要注意几点(下一节讨论)。

Read-Write 图像的价值

尽管图像卷积不如新的读写图像功能有效,但需要就地完成的所有图像处理技巧都可从 Read-Write 图像中获益。 例如,可有效使用的流程是图像合成。

在 OpenCL 1.2 及早期版本中,图像通过 “__read_only” 和 “__write_only” 限定符来限定。 在 OpenCL 2.0 中,图像可通过 “__read_write” 限定符限定,并将输出复制到输入缓冲区中。 这样可减少所需的资源量。

自 OpenCL 1.2 起,图像是 read_only 或 write_image。 执行就地图像调整需要将图像视作缓冲区,并在缓冲区上操作(见cl_khr_image2d_from_buffer: https://software.intel.com/zh-cn/articles/using-image2d-from-buffer-extension)。

当前的解决方案将图像视作缓冲区,并操纵缓冲区。 将 2d 图像视作缓冲区也许不是自由操作,也会阻碍 read_image 中锁定和过滤功能的发挥。 因此,用户更希望使用 read_write 限定的图像。

示例概要

示例提取两张窗口位图图像 “input1.bmp” 和 “input2.bmp”,并将它们放在图像缓冲区中。 然后,这两张图像根据 α 值(已计算像素方程的权重因素,可作为选项传递)进行合成。

Using Alpha value 0.84089642

图 1. 使用 α 值 0.84089642

图像必须为 24/32 位图像。 输出为 24 位图像。 图像必须大小相同。 图像也可以是 ARGB 格式,因此加载时需将其考虑在内。

Using Alpha value of 0.32453

图 2. 使用 α 值 0.32453

ARGB 格式转换成 RGBA 格式。 改变 β 值会导致输出发生重大变化。

使用示例 SDK

SDK 展示了如何将图像合成用于读写图像。 使用以下命令行选项控制该示例:

选项

描述

-h, --help

显示文本并退出

-p, --platform number-or-string

选择设备所使用的平台

-t, --type all | cpu | gpu | acc | default | <OpenCL constant for device type>

按照类型选择执行 OpenCL Kernel 的设备

-d, --device number-or-string

选择执行所有材料的设备

-i, --infile 24/32-bit .bmp file

首个待读取 .bmp 文件的基本名称。 默认为 input1.bmp

-j, --infile 24/32-bit .bmp file

第二个待读取 .bmp 文件的基本名称。默认为 input2.bmp

-o, --outfile 24/32-bit .bmp file

待写入输出的基本名称。 默认为 output.bmp (面向 OCL1.2)和 20_output.bmp(面向 OCL2.0)

-a, --alpha floating point value between 0 and 1

非零正值,确定两个图像的合成程度。 默认 α 值等于 0.84089642。 默认 β 值等于 0.15950358。

示例 SDK 包含许多默认值,支持应用在不进行任何用户输入的情况下运行。 用户将能够使用 input .bmp 文件。 该文件也必须为 24/32 bmp 文件。 α 值可用来确定图像 1 高于图像 2 的凸显度,比如:

calculatedPixel = ((currentPixelImage1 * alpha) + (currentPixeImage2 * beta) + gamma);

用 1 减去 α 值就可得出 β 值。

浮动 β = 1- α;

这两个值可以确定图像 1 到图像 2 的权重分布。

γ 值可用于提高各像素。 默认值为 0。 但用户可以提高合成后整张图像的亮度。 

程序运行示例

Read Write Image Sample Program running on OCL2.0 Device

图 3. 在 OpenCL 2.0 设备上运行的程序

Read-Write 图像的局限性

障碍不能用于需要不同工作组同步的图像。 图像卷积需要同步所有线程。 有关图像的卷积通常会涉及两个矩阵进行数学运算,从而生成第三个矩阵。 使用高斯模糊就是图像卷积的例子。 其他例子包括图像锐化、边缘检测和浮雕。

我们来看看如何使用高斯模糊。 高斯滤波器是一种低通滤波器,可以清除高频率值。 这样会降低清晰度,进而导致类似模糊的效果。 采用高斯模糊与采用高斯函数(通常称为 “mask”)进行图像卷积相同。 为有效展示 Read-Write 图像的功能,需要进行横向模糊和纵向模糊。

这需要 OpenCL 1.2 通过两个步骤来执行。 一个 kernel 专门用于横向模糊,另一个用于纵向模糊。 这样,其中一种模糊将用作输入,而另一种取决于第一种模糊的执行结果。

__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); 
}

代码 3. OpenCL 1.2 中的高斯模糊 kernel

使用 OpenCL 2.0 是希望将两个 kernel 整合成一个。 使用障碍迫使横向模糊或纵向模糊在下一个模糊开始前完成。

__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); 
}

代码 4. OpenCL 2.0 中的高斯模糊 kernel

我们发现,障碍的效率非常低。 如果不首先执行横向模糊,使用障碍将无法确保横向模糊在纵向模糊开始之前完成。 这会产生不一致,进而导致多次运行。 障碍可用来同步一个组内的线程。 出现这种问题的原因是,边缘像素从多个工作组读取,且无法同步多个工作组。 我们最初假设可以使用 read_write 图像实施单种高斯模糊,但证明是错误的,因为工作组内的数据相关性无法在 OpenCL 中同步。

参考资料

关于作者

Oludemilade Raji 是英特尔公司视觉与并行计算事业部的一名显卡驱动程序工程师。 他拥有四年的工作经验,致力于 OpenCL 编程语言和开发英特尔高清显卡驱动程序,包括开发 OpenCL 2.0。

 

Robert Ioffe 是英特尔软件及解决方案事业部的技术咨询工程师。 他是在英特尔锐炬和英特尔锐炬 Pro 显卡上进行 OpenCL 编程和 OpenCL 工作负载优化的专家,拥有丰富的英特尔显卡硬件经验。 他积极参与 Khronos 标准工作,关注于构建最新的特性原型并确保它们在英特尔架构上出色运行。 最近,他一直从事于构建 OpenCL 2.0 的嵌套并行化 (enqueue_kernel functions) 函数的原型,并编写了大量示例来演示嵌套并行化功能,包括面向 OpenCL 2.0 的 GPU-Quicksort。 他还录制并发布了两段 “优化简单的 OpenCL Kernel” 视频和 “OpenCL 2.0 中的 GPU-Quicksort 和 Sierpinski Carpet” 视频。

 

您可能还会对以下内容感兴趣:

优化简单的 OpenCL Kernel: 调节 Kernel 优化

优化简单的 OpenCL Kernel: Sobel Kernel 优化

OpenCL 2.0 中的 GPU-Quicksort: 嵌套并行性和工作组扫描函数

OpenCL 2.0 中的 Sierpiński Carpet

下载

有关编译器优化的更完整信息,请参阅优化通知