OpenCL 主机端采样器对象

在主机端可以使用如下函数来创建主机端采样器对象：

cl_sampler clCreateSamplerWithProperties (
                      cl_context context,
                      const cl_sampler_properties ＊sampler_properties,
                      cl_int ＊errcode_ret)

• 参数context为一个有效的OpenCL上下文。

• 参数sampler_properties为采样器属性名称以及它们相应值的列表。每一个采样器名称后面紧跟着其对应的值，列表以0结束。采样器支持的属性在下表中列出。如果在sampler_properties中没有指定一个支持的属性，则使用其默认值。

  

• 参数errcode_ret为函数执行后的错误码。

上表中的采样器属性，可以分为如下三个操作。

规格化坐标

通常，内核从图像对象读取颜色值，先处理读取到的颜色值，把处理后的颜色值写回到其他图像对象（对于read_write属性的图像，可以写回到原图像对象）中。内核根据定义的坐标值从图像对象相应位置读取颜色值，坐标值可以有下列三种形式：

• 整数坐标：坐标范围在[0,image_dimension-1]。如果一个图像的维度为128×64，则中间值的颜色值可以通过坐标（64,32）来访问。

• 浮点坐标：坐标范围在[0.0,image_dimension-1]。如果一个图像的维度为128×64，则中间值的颜色值可以通过坐标（64.0,32.0）来访问。

• 规格化浮点坐标：坐标范围在[0.0,1.0]。如果一个图像的维度为128×64，则中间值的颜色值可以通过坐标（0.5,0.5）来访问。
  CL_SAMPLER_NORMALIZED_COORDS为CL_TRUE，则使用规格化浮点坐标，否则为非规格化坐标。当我们处理的图像不知道维度时，规格化坐标变得十分有用。

寻址模式

对于一张128×64的图片，它的有效像素坐标范围为（0,0）～（127,63）。如果内核访问的是坐标（150,200）或者（-5,-10）处的颜色值时，寻址模式决定返回什么颜色值。如果寻址模式在图像边界保持颜色值或者设置一个默认颜色值，我们称它为钳位（clamp）输出颜色。在寻址模式中，支持5种寻址模式。

• CL_ADDRESS_NONE：有效坐标范围外的颜色值为未定义。

• CL_ADDRESS_CLAMP：有效坐标范围外的颜色值为一个指定的边界颜色值，默认为黑色。

• CL_ADDRESS_CLAMP_TO_EDGE：有效坐标范围外的颜色值与图像边界颜色值一样。

• CL_ADDRESS_REPEAT：在范围内重复坐标。如果最大维度为N，坐标范围之外X处的颜色值与X%N处相同。（只在规格化坐标下有效。）

• CL_ADDRESS_MIRRORED_REPEAT：坐标范围外的颜色值设置为坐标范围内颜色值的反射。

上述5种寻址模式可以用下图来演示。

过滤模式

当使用整数坐标时，我们可以获得图像对象相应位置上的颜色值。但是，如果为浮点坐标，获得的颜色值为插值后的值。OpenCL支持两种插值方式：

• 最近邻插值（CL_FILTER_NEAREST）：采样器将使用最近邻插值来计算中间颜色值。

• 双线性插值（CL_FILTER_LINEAR）：对于2维图像，采样器取4个最近的颜色值，对它们求平均，这被称为双线性插值。对于3维图像，采样器将会从各个最接近的切片分别取4个颜色值，然后在这些平均值之间完成线性插值，这被称为三线性插值。

如下代码展示了在主机端创建采样器对象：

……
const cl_sampler_properties ＊sampler_properties[] =
{
    CL_SAMPLER_NORMALIZED_COORDS, CL_FASLE,
    CL_SAMPLER_ADDRESSING_MODE, CL_ADDRESS_CLAMP_TO_EDGE,
    CL_SAMPLER_FILTER_MODE, CL_FILTER_LINEAR,
    0
};
cl_sampler sampler = clCreateSamplerWithProperties(context,
sampler_properties, &err);

在创建采样器对象后，可以把采样器对象作为参数传给内核函数：

//内核函数如下
__kernel void ImageProcess(__global sampler_t sampler, ...)
{
    ……
}
//主机参数传入
clSetKernelArg(kernel, 0, sizeof(cl_sampler), &sampler);