OpenCL编程指南-5.5图像读、写函数

news2024/10/6 20:26:25

读图像

OpenCL GPU设备有专用硬件来读、写图像。OpenCL C图像读、写函数允许开发人员充分利用这个专用硬件。OpenCL 中的图像支持是可选的。要了解一个设备是否支持图像，可以使用clGetDeviceInfo API查询CL_DEVICE_IMAGE_SUPPORT属性。
在这里插入图片描述

需要说明的是，read_imagef、read _imagei和read_imageui分别返回一个float4、int4或uint4颜色值。这是因为颜色值最多有4个分量。表5-18列出了图像中缺失分量所用的值。

采样器

图像读函数取一个采样器参数。采样器指定了如何从图像采样像素。可以使用clsetKernelArg API把采样器作为参数传递到一个内核，或者采样器可以是程序源代码中声明的一个类型为sampler_t的恒定变量。

作为参数传递或者在程序源代码中声明的采样器变量都必须是 sampler_t类型。sampler_t类型是一个32位无符号整数常量，可以解释为一个位域。采样器描述了以下信息：
1）规格化坐标：指定coord. xy或coord. xyz值是规格化还是非规格化值，可以设置为CLK_NORMALIZED_COORDS_TRUE或CLK_NORMALIZED_COORDS_FALSE。
2）寻址模式：指定coord.xy或coord.xyz图像坐标如何映射到图像中适当的像素位置，越界的图像坐标如何处理。表5-19描述了目前支持的寻址模式。
3）过滤模式：指定使用的过滤模式，可以设置为CLK_FILTER_NEAREST（即最近过滤器）或CLK_FILTER_LINEAR（即双线性过滤器）。
在这里插入图片描述
下面的例子中，采样器作为参数传递到一个内核

kernel void
my_kernel(read_only image2d_t imgA, sampler_t sampler, write_only image2d imgB)
{
      int2 coord = (int2)(get_global_id(0), get_global_id(1));
      float4 clr = read_imagef(imgA, sampler, coord);
      write_imagef(imagB, coord, color);
}

下面的例子中，采样器在程序源代码中声明：

const sampler_t samplerA = CLK_NORMALIZED_COORDS_FALSE |
                           CLK_ADDRESS_CLAMP |
                           CLK_FILTER_LINEAR;

kernel void
my_kernel(read_only image2d_t imgA, read_only image2d_t imgB, write_only image2d imgB)
{
      int2 coord = (int2)(get_global_id(0), get_global_id(1));
      float4 clr = read_imagef(imgA, samplerA, coord);
      clr *= read_imagef(imgA,
                          (CLK_NORMALIZED_COORDS_FALSE |
                           CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST),
                           imgB);
}

内核中可以使用的最大采样器数目可以通过使用clGetDeviceInfo API查询CL_DEVICE_MAX_SAMPLERS属性来得到。

限制
从相同图像读取时，为read_imagef、read_imagei或read_imageui 指定的采样器必须使用相同的规格化坐标值。下面的例子可以说明这一点（这里突出显示了采样器使用了不同的规格化坐标值）。这会导致未定义的行为，也就是说，返回的颜色可能不正确。

const sampler_t samplerA = CLK_NORMALIZED_COORDS_FALSE |
                           CLK_ADDRESS_CLAMP |
                           CLK_FILTER_LINEAR;

kernel void
my_kernel(read_only image2d_t imgA, write_only image2d imgB)
{
    float4 clr;
    int2 coord = (int2)(get_global_id(0), get_global_id(1));
    float2 normalized_coords;
    float w = get_image_width(imgA);
    float h = get_image_height(imgA);
    
    clr = read_imagef(imgA, samplerA, coord);
    normalized_coords = convert_float2(coord) * (float2)(1.0f / w, 1.0f / h);
    clr *= read_imagef(imgA,
                          (CLK_NORMALIZED_COORDS_FALSE |
                           CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST),
                           normalized_coords);
}

另外，采样器不能声明为数组或指针，也不能作为函数中的局部变量或程序中定义的函数的返回值。函数的采样器参数不能修改。下面的例子给出了一些非法使用采样器的情况，这些会导致一个编译时错误：

sampler_t <- error. return type cannot be sampler_t
internal_proc(read_only image2d_t imgA, write_only image2d imgB)
{
    ...
}

kernel void
my_kernel(read_only image2d_t imgA, sampler_t sampler, write_only image2d imgB)
{
     sampler_t *ptr_sampler; <- error. pointer to sampler not allowed
     my_func(imgA, &sampler); <- error. passing a pointer to a sampler
     ...
}

确定边界颜色

如果采样器寻址模式为CLK_ADDRESs_CLAMP，则越界的图像坐标会返回边界颜色。返回的边界颜色取决于图像通道次序，见表5-20。
在这里插入图片描述

写图像

表5-21和表5-22分别描述了写2D图像和3D图像的内置函数。
在这里插入图片描述

如果×坐标不在范围（0 …图像宽度–1）内，或者y坐标不在范围（0 …图像高度-1)内，对2D图像的write_imagef、write_imagei或write_imageui行为未定义。

如果×坐标不在范围(0 …图像宽度–1）内，或者y坐标不在范围（0 …图像高度-1)内，或者z坐标不在范围（0 …图像深度–1)内，对3D图像的write_imagef、write_imagei或write_imageui行为未定义。

查询图像信息

表5-23描述了图像查询函数。get_image_channel_data_type和get_image_channelorder返回的值使用了一个CLK_前缀。有 CLK_前缀的值与clcreateImage2D和clCreateImage3D的cl_image_format参数中image_channel_order和image_channel_data_type域指定的相应CL_前缀存在一一映射。
在这里插入图片描述