OpenCL：图像处理基础note

使用图像对象的理由

虽然对于图像也可以把它的像素数据当做一般的缓存数据来处理，但是如果把它当做图像来处理有如下好处：

1. 在GPU中，图像数据是保存在特殊的全局内存中，即 纹理内存 ，它和一般的全局内存不相同，它是被缓存的，用于高速访问处理。GPU中有专门 支持图像读写的硬件，使用内置读写函数可以充分发挥这个优势 。
2. 只要OpenCL支持该图像格式，那么就可以不用考虑图像格式的前提下使用读写图像数据的函数
3. 可以使用 采样器 来配置读取图像中数据的的方式
4. OpenCL提供函数来获取图像相关信息，比如宽度等

主机与内核的命名

图像相关应用主要涉及两个数据类型，即 图像对象 和 采样器 。图像对线用来保存主机和设备上的图像对象，而采样器则在设备接收数据时，说明如何读取这些颜色值

主机编程的主要接口

创建图像对象

cl_mem clCreateImage2D()

cl_mem clCreateImage3D()

图像格式

cl_image_format

创建采样器

cl_sampler clCreateSampler()

设备上的接口

图像数据

由于很多设备都是将图像对象保存在特定的内存中，因此image2d_t以及image3d_t数据的前面通常会有read_only或者write_only等修饰符。 图像不能即可写又可读 。

图像数据作为内核的参数时，和普通的缓存对象不同，它不是指针参数，因为图像数据并不希望直接通过内存操作来访问数据。

采样器

根据khronos官网的说明，采样器其实就是一个uint类型（可以直接给内置cl读写函数传递标识位的组合来替代sampler_t，例如： float4 clr = read_imagef(img, CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST, (float2)(x, y)); ），通过bit-mode来说明 坐标模式，超出边界的处理模式以及插值模式 。在官方的sample中采用的是constant申明的模式，虽然也说了可以采用global声明，但这也许说明constant更好。

可以通过设置参数 clSetKernelArg 的方式来给内核传递采样器。但是 更为方便 的方式是在cl文件中，在内核函数之前创建一个sampler_t对象，如下case：

__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |

CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;

__kernel void simple_image(read_only image2d_t src_image,

write_only image2d_t dst_image) {

/* Compute value to be subtracted from each pixel */

uint offset = get_global_id(1) * 0x4000 + get_global_id(0) * 0x1000;

/* Read pixel value */

int2 coord = (int2)(get_global_id(0), get_global_id(1));

uint4 pixel = read_imageui(src_image, sampler, coord);

/* Subtract offset from pixel */

pixel.x -= offset;

/* Write new pixel value to output */

write_imageui(dst_image, coord, pixel);

}

内置图像函数

OpenCL的内置图像相关的函数主要包括三类：read_imageT读取类，write_imageT写入类以及get_image_X读取图像信息类。其中T表示数据类型，比如f，i，ui等，而X表示width，dim等图像信息

read_image

• read_image函数返回的都是四分量的数据，如果实际图片没有四个通道的数据，那么那些不含有的的通道对应的分量会被置为0，具体请看 官网
• 很多时候都会看到代码中用 float4 read_imagef() 去读图像数据为int或者uint类型的image，可能你会困惑 为什么整形要用float去读 ？这是因为，这些图片在创建的时候就是以正则NORM化的数据类型，这个类型在 cl_image_format 的 image_channel_data_type 中指定。以int8为例， image_channel_data_type 可取的值如下表： 参考

• 使用 read_imagei 和 read_imageui 时，采样器的 filter_mode 必须设置为 CLK_FILTER_NEAREST
• 使用 read_imagef 时，如果坐标是整数，那么 filte_mode 必须设置为 CLK_FILTER_NEAREST
• 使用 read_imageX 时， X 表示**，如果坐标是整数，那么采样器的坐标正则化必须设置为 CLK_NORMALIZED_COORDS_FALSE ，且 address mode 必须设置为 CLK_ADDRESS_CLAMP_TO_EDGE, CLK_ADDRESS_CLAMP 或者 CLK_ADDRESS_NONE

Linux公社的RSS地址： https://www.linuxidc.com/rssFeed.aspx

本文永久更新链接地址： https://www.linuxidc.com/Linux/2018-09/154274.htm