OpenCL编程指南-5.1工作项函数-整数函数-公共函数

工作项函数

应用程序使用clEnqueueNDRangeKernel和 clEnqueueTask API将OpenCL中的数据并行和任务并行内核排队。对于一个数据并行内核（使用clEnqueueNDRangeKernel排队等待执行)，应用程序会指定全局工作大小，即可以并行执行这个内核的工作项总数，还会指定局部工作大小，即可以归组到一个工作组中的工作项数目。

如下描述了OpenCL内核可以调用的一些内置函数，以便得到有关工作项和工作组的信息，如工作项的全局ID和局部ID，以及全局工作大小和局部工作大小。

uint get_work_dim()                            返回使用的维度数目。这是clEnqueueNDRangeKernel中指定的work_dim参数值对于clEnqueueTask，这个函数返回1size_t get_global_size(uint dimindx)           返回为维度dimindx指定的全局工作项数目。这个值由clEnqueueNDRangeKernel的global_work_size参数给定。dimindx的合法取值范围为0 ~ get_work_dim() - 1。对于其他dimindx值，get_global_size()返回1对于clEnqueueTask，这个函数总是返回1size_t get_global_id(uint dimindx)             返回维度dimindx的唯一全局工作项ID值。全局工作项ID根据为执行内核指定的全局工作项数目来指定工作项ID。dimindx的合法取值范围为0 ~ get _work_dim() - 1。对于其他dimindx值，get_global_id()返回0对于clEnqueueTask，这个函数总是返回0size_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，这个函数总是返回1size_t get_local_id(uint dimindx)              返回维度dimindx的唯一局部工作项ID值（即一个特定工作组中的一个工作项)。dimindx的合法取值范围为0 ~ get_work_dim()-1。对于其他dimindx值，get_local_id()返回0对于clEnqueueTask，这个函数总是返回0size_t get_num_groups(uint dimindx)            返回为维度dimindx指定的将执行内核的工作组数目。dimindx的合法取值范围为0 ~ get_work_dim()-1。对于其他dimindx值，get_num_groups()返回1对于clEnqueueTask，这个函数总是返回1size_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，这个函数总是返回0size_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 否则返回xgentype 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) | lolong 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)*radiansgentype 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)*degreesgentype 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作为参数类型。