使用图像对象的理由
虽然对于图像也可以把它的像素数据当做一般的缓存数据来处理,但是如果把它当做图像来处理有如下好处:
- 在GPU中,图像数据是保存在特殊的全局内存中,即纹理内存,它和一般的全局内存不相同,它是被缓存的,用于高速访问处理
- 只要OpenCL支持该图像格式,那么就可以不用考虑图像格式的前提下使用读写图像数据的函数
- 可以使用采样器来配置读取图像中数据的的方式
- OpenCL提供函数来获取图像相关信息,比如宽度等
主机与内核的命名
图像相关应用主要涉及两个数据类型,即图像对象和采样器。图像对线用来保存主机和设备上的图像对象,而采样器则在设备接收数据时,说明如何读取这些颜色值
| |图像对象|采样器|
|---|---|---|
|主机|cl_mem (和一般缓存对象一样)|cl_sampler|
|设备|image2d_t或者image3d_t|sampler_t|
主机编程的主要接口
创建图像对象
cl_mem clCreateImage2D()
cl_mem clCreateImage3D()
图像格式
cl_image_format
创建采样器
cl_sampler clCreateSampler()
设备上的接口
图像数据
由于很多设备都是将图像对象保存在特定的内存中,因此image2d_t以及image3d_t数据的前面通常会有read_only或者write_only等修饰符。图像不能即可写又可读。
图像数据作为内核的参数时,和普通的缓存对象不同,它不是指针参数,因为图像数据并不希望直接通过内存操作来访问数据。
采样器
可以通过设置参数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等图像信息