工作项函数
应用程序使用clEnqueueNDRangeKernel和 clEnqueueTask API将OpenCL中的数据并行和任务并行内核排队。对于一个数据并行内核（使用clEnqueueNDRangeKernel排队等待执行)，应用程序会指定全局工作大小，即可以并行执行这个内核的工作项总数，还会指定局部工作大小，即可以归组到一个工作组中的工作项数目。
如下描述了OpenCL内核可以调用的一些内置函数，以便得到有关工作项和工作组的信息，如工作项的全局ID和局部ID，以及全局工作大小和局部工作大小。
uint get_work_dim() 返回使用的维度数目。这是clEnqueueNDRangeKernel中指定的work_dim参数值
对于clEnqueueTask，这个函数返回1
size_t get_global_size(uint dimindx) 返回为维度dimindx指定的全局工作项数目。这个值由clEnqueueNDRangeKernel的global_work_size参数给定。
dimindx的合法取值范围为0 ~ get_work_dim() - 1。
对于其他dimindx值，get_global_size()返回1
对于clEnqueueTask，这个函数总是返回1
size_t get_global_id(uint dimindx) 返回维度dimindx的唯一全局工作项ID值。全局工作项ID根据为执行内核指定的全局工作项数目来指定工作项ID。
dimindx的合法取值范围为0 ~ get _work_dim() - 1。
对于其他dimindx值，get_global_id()返回0
对于clEnqueueTask，这个函数总是返回0
size_t get_local_size(uint dimindx) 返回为维度dimindx指定的局部工作项数目。
如果local_work_size不为NULL，则这个值由clEnqueueNDRangeKernel的local_work_size参数给定;
否则，OpenCL实现会选择一个适当的local_work_size值。dimindx的合法取值范围为0 ~ get_work_dim() -1。
对于其他dimindx值，get_local_size()返回1
对于clEnqueueTask，这个函数总是返回1
size_t get_local_id(uint dimindx) 返回维度dimindx的唯一局部工作项ID值（即一个特定工作组中的一个工作项)。
dimindx的合法取值范围为0 ~ get_work_dim()-1。
对于其他dimindx值，get_local_id()返回0
对于clEnqueueTask，这个函数总是返回0
size_t get_num_groups(uint dimindx) 返回为维度dimindx指定的将执行内核的工作组数目。
dimindx的合法取值范围为0 ~ get_work_dim()-1。
对于其他dimindx值，get_num_groups()返回1
对于clEnqueueTask，这个函数总是返回1
size_t get_group_id(uint dimindx) 返回工作组 ID，这是介于0 ~ get_num_groups(dimindx)-1之间的一个数。
dimindx的合法取值范围为0 ~ get_work_dim() - 1。
对于其他dimindx值，get_group_id()返回0
对于clEnqueueTask，这个函数总是返回0
size_t get_global_offset(uint dimindx) 返回clEnqueueNDRangeKernel的global_work_offset参数中指定的偏移值，dimindx的合法取值范围为0~get_work_dim()-1。
对于其他dimindx值,get_global_offset()返回0
对于clEnqueueTask，这个函数总是返回0
图5-1给出了一个例子，展示设备上执行的内核可以访问clEnqueueNDRangeKernel中指定的全局工作大小和局部工作大小。在这个例子中，执行内核的全局工作大小为16个工作项，工作组大小为每组8个工作项。
OpenCL没有描述全局ID和局部ID如何映射到工作项和工作组。例如，一个应用程序不会假设组ID为0的工作组就包含全局ID为0 ~ get_local_size(0)-1的工作项。这个映射由OpenCL实现以及执行内核的设备来确定。
整数函数
ugentype abs(gentype x) 返回|x|
ugentype abs_diff(gentype x, gentype y) 返回|x-y|,无模上溢出
gentype add_sat(gentype x, gentype y) 返回x+y, 且使结果饱和
gentype hadd(gentype x, gentype y) 返回(x+y)>>1 中间的和无模上溢出
gentype rhadd(gentype x, gentype y) 返回(x+y+1)>>1 中间的和无模上溢出
gentype clamp(gentype x, gentype minval, gentype maxval) 返回min(max(x,minval),maxval)
gentype clamp(gentype x, sgentype minval, sgentype maxval) 如果minval>maxval 则结果未定义
gentype clz(gentype x) 返回x中前导0的位数，从最高有效位开始
gentype mad_hi(gentype a, gentype b, gentype c) 返回mul_hi(a,b)+c
gentype mad_sat(gentype a, gentype b, gentype c) 返回a*b+c 且使结果饱和
gentype max(gentype x, gentype y) 如果x<y 返回y 否则返回x
gentype max(gentype x, sgentype y)
gentype min(gentype x, gentype y) 如果x<y 返回y 否则返回x
gentype min(gentype x, gentype y)
gentype mul_hi(gentype x, gentype y) 计算x*y 返回y 否则返回x
gentype rotate(gentype v, gentype i) 对于v中的各个元素，使各位按i中相应元素给定的位数左移。
元素左边移出的位再从右边移入
gentype sub_sat(gentype x, gentype y) 返回x-y 并使结果饱和
short upsample(char hi, uchar lo) 如果hi和lo是标量
ushort upsample(uchar hi, uchar lo) result=((short)hi << 8) | lo
shortn upsample(charn hi, ucharn lo) result=((ushort)hi << 8) | lo
ushortn upsample(ucharn hi, ucharn lo)
int upsample(short hi, ushort lo) result = ((int)hi << 16) | lo
uint upsample(ushort hi, ushort lo) result = ((uint)hi << 16) | lo
long upsample(int hi, uint lo) result = ((long)hi << 32) | lo
ulong upsample(uint hi, uint lo) result = ((ulong)hi << 32) | lo
longn upsample(intn hi, uintn lo) 如果hi和lo是标量，则对于矢量的各个元素：
ulongn upsample(uintn hi, uintn lo) result[i] = ((short)hi[i] << 8) | lo[i]
result[i] = ((ushort)hi[i] << 8) | lo[i]
result[i] = ((int)hi[i] << 16) | lo[i]
result[i] = ((uint)hi[i] << 16) | lo[i]
result[i] = ((long)hi[i] << 32) | lo[i]
result[i] = ((ulong)hi[i] << 32) | lo[i]
gentype mad24(gentype x, gentype y, gentype z) 使用mul24将两个24位整数值x和y相乘
并把32位整数结果与32位整数z相加
gentype mul24(gentype x, gentype y) 将两个24位整数值x和y相乘。
尽管×和y是32位整数，但是只使用其低24位完成乘法运算。
只有当x和y中的值在[-2^23,2^23 -1]范围内（如果x和y是有符号整数）
或者在[0, 224-1]范围内（如果×和y是无符号整数)时，才使用mu124。
如果x和y不在这个范围内，乘法结果由具体实现定义
使用泛型类型名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作为参数类型。
使用泛型类型名ugentype指示无符号的gentype。例如，如果 gentype是 char4，ugentype就是uchar4。另外使用泛型类型名sgentype指示有符号的gentype。
使用泛型类型名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_BTT 8
#define CHAR__MAX SCHAR__MAX
#define CHAR_MIN SCHAR_MIN
#define INT_MAX 2147483647
#define INT_MIN (-2147483647-1)
#define LONG_MAX Ox7ffffffffffffffL
#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 Oxffffffff
#define ULONG_MAX OxffffffffffffffffUL
公共函数
gentype clamp(gentype x, gentype minval, gentype maxval) 返回fmin(fmax(x,minval),maxval)
gentype clamp(gentypef x, float minval, float maxval) 如果minval>maxval 结果未定义
gentype clamp(gentyped x, double minval, double maxval)
gentype degrees(gentype radians) 将弧度转换为度，即(180/p)*radians
gentype max(gentype x, gentype y) 如果x<y 则返回y; 否则返回x。
gentypef max(gentypef x, float y) 这与fmax类似, 只不过如果x或y为infinite或NaN时返回值未定义
gentyped max(gentyped x, double y)
gentype min(gentype x, gentype y) 如果y<x则返回y; 否则，返回x。
gentypef min(gentypef x, float y) 这与fmin类似，只不过如果x或y为infinite或NaN时返回值未定义
gentyped min(gentyped x, double y)
gentype mix(gentype x,gentype y.gentype a) 返回x和y的线性混合，实现为x + (y - x)* a
gentypef mix(gentypef x,float y,gentype a) 如果a不在此范围内，返回值未定义
gentyped mix(gentyped x, double y,gentype a) a必须是范围0.0...1.0中的一个值。
gentype radians(gentype degrees) 将degrees转换为弧度，即(p/180)*degrees
gentype step(gentype edge, gentype x) 如果x<edge返回0.0; 否则返回1.0
gentypef step(float edge, gentypef x) step函数可以用于在任意点创建一个不连续的跳步
gentyped step(double edge, gentyped x)
gentype smoothstep(gentype edge0, gentype edgel, gentype x) 如果x≤edge0返回0.0，如果x>=edge1 则返回1.0;
gentypef smoothstep(float edge0, float edgel, gentypef x) edge0<x<edge1时，在0~1之间完成一个平滑hermite插值。
gentyped smoothstep(double edge0, double edgel, gentyped x) 这在需要一个平滑过渡的阈值函数时很有用
t与x类型相同时，等价于:
t = clamp((x-edgeO) / (edgel-edge0),0,1);return t * t * (3 - 2 * t)
如果edge0≥ edgel或者x、edge0或edge1 为NaN 时，结果未定义
gentype sign(gentype x) 如果x>0，返回1.0
如果x=-0.0,返回-0.0;
如果x=+0.0，返回+0.0;
如果x<0，则返回-1.0;
如果x是NaN,则返回0.0
使用泛型类型名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作为参数类型。
