OpenCL 学习笔记 - OpenCL C built-in functions - OpenCL C的内建函数

work-item funciton - 工作项函数

应用程序使用函数clEnqueueNDRangeKernel和clEnqueueTask将OpenCL中的数据并行和任务并行kernel入队。使用clEnqueueNDRangeKernel将数据并行的kernel入队并执行，应用程序指明全局的工作量(global work size，即并行执行这个kernel的工作项(work item)的个数)，局部的工作量(local work size，即一个工作组(work-group)中工作项的个数)。表5.1列出了为了查询work-item和work-group相关信息，OpengCL kernel可以调用的内建函数，比如获得work-item全局的和局部的ID，或者全局的和局部的工作量。

图5.1的例子说明了设备上执行的kernel如何获得由函数clEnqueueNDRangeKernel指定的全局和局部工作量。这里例子里面，kernel的全局工作量是16，工作组的大小是8(即一个工作组有8个工作项)。

OpenCL没有说明全局和局部ID是如何映射到工作项和工作组的。应用程序不能假定：组ID(group ID)为0的工作组将包含全局ID为0...get_local_size(0) - 1的工作项。这个映射是由OpenCL实现和执行kernel的设备所决定的。

math-function - 数学函数

OpenCL实现了C99规范里面描述的数学函数。要使用这些函数，必须包含头文件math.h，而OpenCL kernel是可以直接使用这些函数的，就是kernel可以不用包含math.h。

表格5.2和5.3列出的数学函数中的gentype为通用类型，包括float, float2, float3, float4, float8, float16, 如果支持双精度数的扩展，那么还包括double, double2, double3, double4, double8, double16。gentypei表示int, int2, int3, int4, int8, 或者 int16。gentypef表示float, float2, float3, float4, float8或float16。gentyped表示double, double2, double3, double4, double8或double16。

除了表5.2列出的数学函数之外，OpenCL C为单精度浮点数标量和变量实现了常用的数学函数的两个变体(variant)。这些附加的函数(见表5.3)为了性能而牺牲了精度，使得开发者的选择更加灵活。这些数学函数分为：

• 表5.2中的部分函数加上half_的前缀。这些函数最小的精度是10位，也就是说，ulp <= 8292ulp。
• 表5.2中的部分函数加上native_的前缀。比起没有native_前缀或者前缀为half_的函数，这类函数有最高的性能，而具体的精度是由实现决定的。
• 前缀为half_, native_的、处理除法和倒数运算的函数。

floating-point pragmas - 浮点编译提示

OpenCL C所支持的唯一的编译提示是FP_CONTRACT编译提示。这个编译提示可以用来禁止收缩表达式：

#pragma OPENCL FP_CONTRACT on-off-switch

on-off-switch的值为：ON, OFF或者DEFAULT。DEFAULT的值是ON。

FP_CONTRACT编译提示用来打开或关闭收缩表达式(contract expression)。如果FP_CONTRACT设为ON，浮点表达式可能被收缩，也就是像原子操作那样计算值。比如，表达式a*b+c能够被一个FMA(fused multiply-add)指令代替。

每个FP_CONTRACT编译提示可以在外部申明，也可以在一个复合语句中。在外部申明时，这个编译提示作用域一直到下一个编译提示出现，或者在翻译单元的结束。在一个复合语句中使用时，这个编译提示的作用域一直到下一个编译提示的出现（包括在内嵌的符合语句中），或者到这个符合语句的结尾。在复合语句的结尾，编译提示的状态恢复到这个复合语句之前的状态。

floating-pointing constants - 浮点常量

表5.4列出了可用的常量。以_F为后缀的常量的类型为float，有float类型数据的精度。没有_F后缀的常量的类型是double，有double的精度，只有在OpenCL的实现支持双精度数扩展的时候可以使用。

relative error as ulps - ulps相对误差

表5.5中为定义为ulp(units in the last place)的、对双精度和单精度数的基本运算和函数的、最大相对误差(maximum relative error)。ulp定义为(注：ulp的定义遵从Jean-Michel Muller的、关于0的行为的申明，参见ftp://ftp.inria.fr/INRIA/publication/publi-pdf/RR/RR-5504.pdf.)：

如果x是两个连续的浮点数a, b之间的一个实数，且不等于a或b，那么ulp(x)=|b-a|，否则ulp(x)是最靠近x的两个不等有限浮点数之间的距离。ulp(NaN)为NaN

下面列出了ulp值和截断模式的附加申明：

• 在完整的版本(full profile)中最近截断(round-to-nearest)是默认的阶段模式。在嵌入的版本(embedded profile)中，默认的截断模式可以是向0截断，也可能是最近阶段。如果在CL_DEVICE_SINGLE_FP_CONFIG支持CL_FP_ROUND_TO_NEAREST(参见OpengCL 1.1说明文档的表4.3)，那么嵌入版本默认是最近截断；否则默认为向零阶段。
• 0 ulp用于没有截断的数学函数
• 数学函数lgama和lgama_r的ulp值是未定义的。

integer funcitons - 整数函数

表5.6列出了OpenCL C中的内建的整数函数。

gentype指明函数可以用char(1, 2, 3, 4, 8, 16), uchar, short, ushort, int, uint, long, ulong作为参数。用ugentype指明无符号的gentype。比如，如果gentype是char4，那么ugentype是uchar4。sgentype说明函数以标量作为参数，也就是char, uchar, short, ushort, int, uint, long, ulong。这些以gentype和sgentype作为参数的内建的整数函数，gentype类型的参数必须是sgentype类型的向量或者标量。比如，如果sgentype是uchar，gentype必须是uchar, uchar2, uchar3, uchar4, uchar8或者uchar16。

下面是一些关于整数的宏。它们的值是一些常量表达式，可以用#if处理指令来控制。

#define CHAR_BIT 8
#define CHAR_MAX SCHAR_MAX
#define CHAR_MIN SCHAR_MIN
#define INT_MAX 2147483647
#define INT_MIN (-2147483647 – 1)
#define LONG_MAX 0x7fffffffffffffffL
#define LONG_MIN (-0x7fffffffffffffffL – 1)
#define SCHAR_MAX 127
#define SCHAR_MIN (-127 – 1)
#define SHRT_MAX 32767
#define SHRT_MIN (-32767 – 1)
#define UCHAR_MAX 255
#define USHRT_MAX 65535
#define UINT_MAX 0xffffffff
#define ULONG_MAX 0xffffffffffffffffUL

common functions - 通用函数

表5.7列出了可用的通用函数。

gentype表示float, float2, float3, float4, float8, float16, 如果支持双精度数的扩展，那么还包括double, double2, double3, double4, double8, double16。gentypef表示float, float2, float3, float4, float8或float16。gentyped表示double, double2, double3, double4, double8或double16。

geometric funtions - 几何函数

表5.8为可用的几何函数。gentypef表示float, float2, float3, float4, float8或float16。gentyped表示double, double2, double3, double4, double8或double16。

下面是关于这些通用几何函数的一些说明：

• 几何函数可以通过缩略(contraction, 比如mad, fma)来实现。
• fast_使得开发者可以在精度之前让性能优先。
• distance, length和normalize函数不会产生上溢或者由下溢产生的不正常的精度丢失。

relational function - 关系函数

表5.9列出了OpenCL C中的关系函数。

函数isequal, isnotequal, isgreater, isgreaterequal, isless, islessequal, islessgreater, isfinite, isinf, isnan, isnormal, isordered, isunordered,signbit，参数为标量时，关系为false，这些函数返回0，关系为true，则返回1。如果参数为向量，则关系为true时返回-1，关系为false时，返回0。

函数isequal, isgreater, isgreaterequal, isless, islessequal islessgreater，有参数为NaN时，返回0。如果参数为标量，且一个或者两个参数为NaN时，isnotequal返回1；如果参数为向量，且一个或两个参数为NaN时，返回-1。

表5.10中为附加的关系函数。gentype表示char, char2, char3, char4, char8, char16, uchar, uchar2, uchar3, uchar4, uchar8, uchar16, short, short2, short3, short4, short8, short16, ushort, ushort2, ushort3, ushort4, ushort8, ushort16, int, int2, int3, int4, int8, int16, uint, uint2, uint3, uint4, uint8, uint16, long, long2, long3, long4, long8, long16, ulong, ulong2, ulong3, ulong4, ulong8, ulong16, float, float2, float3, float4, float8, float16，如果支持双精度扩展，还表示double, double2, double3, double4, double8, double16。

sgentype表示char, char2, char3, char4, char8, char16, short, short2, short3, short4, short8, short16, int, int2, int3, int4, int8, int16, long, long2, long3, long4, long8, or long16。

ugentype表示uchar, uchar2, uchar3, uchar4, uchar8, uchar16, ushort, ushort2, ushort3, ushort4, ushort8, ushort16, uint, uint2, uint3, uint4, uint8, uint16, ulong, ulong2, ulong3, ulong4, ulong8, ulong16.

Vector Data Load and Store Functions - 向量数据装载和存储函数

表5.11描述了用来从一个指针读取和写入向量类型的函数。gentype表示标量内建数据类型char, uchar, short, ushort, int, uint, long, ulong, float, double。使用gentypen来说明n个元素的、gentype类型的向量。后缀n同样也使用在函数中，比如vloadn, vstoren，其中n=2，3,4,8,16。

synchronization functions - 同步函数

OpenCL C实现一个名为barrier的同步函数，该函数用来保证供一个工作组里面、各个工作项之间的内存一致性。

||一个工作组在一个计算单元上执行kernel，其中的所有工作项必须在任意一个工作项继续执行后面的内容之前，执行这个函数。一个工作组里的所有工作项必须运行到这个函数。

如果barrier函数在一个条件语句内，且有一个工作项进入了条件语句并执行了这个barrier，那么所有工作项必须进入这个条件。

如果barrier函数在一个循环内，任何一个工作项想要执行后面的barrier函数，必须在所有的工作项在一次迭代中执行这个barrier函数之后。

barrier函数同样会建立一个存储的围墙（读与写），用来保证local global存储操作的正确顺序。

该函数的参数指明存储地址空间，可以是一下值的混合：

• CLK_LOCAL_MEM_FENCE：barrier函数要么刷新local存储中的所有变量，要么建立一个围墙来保证正确的、local存储的操作顺序。
• CLK_GLOBAL_MEM_FENCE：barrier函数要么刷新global存储中的所有变量，要么建立一个围墙来保证正确的、global存储的操作顺序。当工作组中的工作项要写到一个global的缓冲对象然后又要读这个更新之后的数据时，必须使用这个参数。

如果工作组中所有的工作项都没有运行到barrier，结果是未定义的。在一些设备上，特别是GPU，这很有可能产生一个硬件的死锁。下面例子中的barrier用法是错误的。

kernel void
read(global int *g, local int *shared)
{
if (get_global_id(0) < 5)
barrier(CLK_GLOBAL_MEM_FENCE); ← illegal since not all workitems
encounter barrier.
else
k = array[0];
}

注意，这里的内存一致性是一个工作组中，不同工作项之间的，而不是工作组之间的。下面的例子说明了这一点：

kernel void
smooth(global float *io)
{
float temp;
int id = get_global_id(0);
temp = (io[id – 1] + id[id] + id[id + 1]) / 3.0f;
barrier(CLK_GLOBAL_MEM_FENCE);
io[id] = temp;
}

如果kernel smooth执行的全局工作大小为16，在2个大小为8的工作组上执行，那么值被存储在io[7]还是io[8]是不确定的。因为在两个工作组中的工作项使用io[7]和io[8]来计算temp。工作组0使用它来计算io[7]的temp，工作组1使用它来计算io[8]的temp。因为不能确定工作组执行的时间先后，以及在哪个计算单元执行，barrier又只能保证同一个工作组里的工作项之间的内存一致性，这里我们就不能保证io[7]和io[8]的值。

async copy and prefetch functions - 同步复制和预读取函数

OpenCL C通过表5.13所列的函数，为global和local存储之间复制以及从global存储中预读取，提供了一个方便高效方法。在全局和局部之间复制的函数称为同步复制。

下面的例子中，使用async_work_group_strided_copy进行跳跃的复制，先从全局到局部，在从局部到全局。有多个元素组成的缓冲区，其中的每个元素表示三维几何体的一个顶点。每个顶点是一个存储了位置，法向量，纹理坐标和其他信息的一个结构体。一个OpenCL kernel想要读取顶点的位置坐标，并进行一些计算，再存储这些更新之后的值。这就需要跳跃地、从全局存储复制到局部存储中，然后进行计算，在通过将局部的数据复制到全局存储中，来更新顶点位置。

typedef struct {
float4 position;
float3 normal;
float2 texcoord;
...
} vertex_t;
kernel void
update_position_kernel(global vertex_t *vertices,
local float4 *pos_array)
{
event_t evt = async_work_group_strided_copy(
(local float *)pos_array,
(global float *)vertices,
4, sizeof(vertex_t)/sizeof(float),
NULL);
wait_group_events(evt);
// do computations
. . .
evt = async_work_group_strided_copy((global float *)vertices,
(local float *)pos_array,
4, sizeof(vertex_t)/sizeof(float),
NULL);
wait_group_events(evt);
}

kernel运行结束之前必须等待同步复制完成，这是通过内建函数wait_group_events实现的。如果等待，结果是不确定的。

atomic function - 原子函数

OpenCL C通过表5.14中的内建函数，提供了在局部和全局存储中的32位有符号整数、无符号整数和单精度数的原子操作。

注意：atom_xchg是唯一一个使用浮点数作为参数的原子函数。

miscellaneous vector functions - 各种向量函数

OpenCL实现了表5.25中的一些内建向量函数。

下面的例子说明如何使用函数shuffle和shuffle2。

uint mask = (uint4)(3, 2, 1, 0);
float4 a;
float4 r = shuffle(a, mask); // r.s0123 = a.wzyx
uint8 mask = (uint8)(0, 1, 2, 3, 4, 5, 6, 7);
float4 a, b;
float8 r = shuffle2(a, b, mask); // r.s0123 = a.xyzw,
// r.s4567 = b.xyzw

下面的例子会导致编译错误。

uint8 mask;
short16 a;
short8 b;
b = shuffle(a, mask); // not valid

在需要做交换操作时，使用shuffle和shuffle2函数比自己编写交换的代码更好，因为编译器可以很好将内建的函数映射到对应的硬件指令集架构上面。

image read and write functions - 图像读写函数

这部分，我们将介绍一些用来函数用来读取图像、写入图像和查询一些图像信息的函数。

OpenCL GPU设备有专门的硬件用来读写图像。OpenCL C的图像读写函数允许开发者利用者的专用的硬件。OpenCL的图像支持是可选的。用CL_DEVICE_IMAGE_SUPPORT调用clGetDeviceInfo函数可以查看设备是否支持图像。

reading from an image - 从图像中读取数据

表5.16和5.17中的函数分别可以用来从二维和三维的图像中读取数据。

read_imagef, read_imagei和read_imageui分别返回一个float4, int4和uint4类型的颜色值。表5.18列出了图像中没有的分量的值。

sampler - 采样器

limitation - 限制

determining the border color - 确定边界颜色

writing to an image - 写入图像

querying image information - 查询图像信息