对于OpenCL内建的标量或者矢量数据类型,我们可以直接通过赋值(=)语句来读等号右边的变量值,并把相应的值写入到等号左边的变量中。但是对于图像数据,OpenCL设备有专门的硬件来读/写图像。与其他通过赋值语句读/写不一样,OpenCL中有读写图像的内置函数,使得我们可以充分利用这个专用硬件。OpenCL设备对图像支持是可选的。我们可以使用clGetDeviceInfo()来查询CL_DEVICE_IMAGE_SUPPORT属性,来确定设备是否支持图像。
声明为read_only的图像对象,对图像进行写操作,则会产生编译错误;同样,如果声明为write_only的图像对象,对图像进行读操作,也会产生编译错误。只有声明为read_write的图像对象,才可以对图像进行读、写操作。
声明为read_only的图像,需要使用带有采样器的读函数来读取图像;而声明为read_write的图像,不能使用带有采样器的读函数来读取图像。
在内核中,我们可以查询图像的大小、深度、通道数据类型等信息。
内建图像读函数
我们来讲解声明为read_only的图像对象该如何读取。下表列出了内建图像读函数。
主机代码:
cl_image_format format;
cl_mem ImgMem;
……
format.image_channel_data_type = CL_UNORM_INT8;
format.image_channel_order = CL_RGBA;
ImgMem = clCreateImage2D(context, CL_MEM_READ_ONLY, &format,
width, height, 0, NULL, &err);
……
内核代码:
const sampler_t sampler = CLK_ADDRESS_CLAMP_TO_EDGE |
CLK_FILTER_LINEAR;
kernel void ImageTest(read_only image2d_t Src,
write_only image2d_t Dst)
{
int2 coord = (int2)(get_global_id(0), get_global_id(1));
float4 p00 = read_imagef(Src, sampler, coord);
……
}
例子中,设置图像数据类型为CL_UNORM_INT8。内核代码中,图像坐标为整数,故采样器过滤模式设置为CLK_FILTER_NEAREST,寻址模式设置为CLK_FILTER_LINEAR。