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 - 浮点常量
relative error as ulps - ulps相对误差
下面列出了ulp值和截断模式的附加申明:如果x是两个连续的浮点数a, b之间的一个实数,且不等于a或b,那么ulp(x)=|b-a|,否则ulp(x)是最靠近x的两个不等有限浮点数之间的距离。ulp(NaN)为NaN
- 在完整的版本(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 - 整数函数
#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 - 通用函数
geometric funtions - 几何函数
- 几何函数可以通过缩略(contraction, 比如mad, fma)来实现。
- fast_使得开发者可以在精度之前让性能优先。
- distance, length和normalize函数不会产生上溢或者由下溢产生的不正常的精度丢失。
relational function - 关系函数
Vector Data Load and Store Functions - 向量数据装载和存储函数
synchronization functions - 同步函数
- CLK_LOCAL_MEM_FENCE:barrier函数要么刷新local存储中的所有变量,要么建立一个围墙来保证正确的、local存储的操作顺序。
- CLK_GLOBAL_MEM_FENCE:barrier函数要么刷新global存储中的所有变量,要么建立一个围墙来保证正确的、global存储的操作顺序。当工作组中的工作项要写到一个global的缓冲对象然后又要读这个更新之后的数据时,必须使用这个参数。
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 - 同步复制和预读取函数
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 - 原子函数
miscellaneous vector functions - 各种向量函数
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函数比自己编写交换的代码更好,因为编译器可以很好将内建的函数映射到对应的硬件指令集架构上面。