OpenCL编程指南-4.1OpenCL C编程

使用OpenCL C编写数据并行内核

OpenCL中的数据并行性表述为一个N维计算域，其中N=1、2或3。N-D域定义了可以并行执行的工作项的总数。下面通过一个简单的例子来了解如何用OpenCL C编写一个数据并行内核，将两个浮点数数组相加。这个代码的串行版本求和时需要通过一个for循环将两个数组中的各个元素相加：

void scalar_add (int n, const float *a, const float *b, float *result)
{int i;for (i = 0; i < n; i++){result[i] = a[i] + b[i];}
}

使用OpenCL C的数据并行代码如下所示。

kernel void scalar_add(global const float *a, global const float *b, global float *result)
{int id = get_global_id(0);result[id] = a[id] + b[id];
}

scalar_add函数声明使用kernel限定符指示这是一个OpenCL C内核。需要说明的是，scalar_add内核只包括计算单个元素求和的代码，也就是内循环。N-D域是设置为n的1维域。对于n个工作项，要为每个工作项分别执行内核来生成数组a和b的和。为此，每个要执行内核的工作项需要知道要对数组a和 b中的哪个元素求和。对于各个工作项来说，这必须是唯一的值，要由N-D域（将内核入队等待执行时指定）得出。get_global_id(0)返回各个工作项的1维全局ID。

图4-1显示了如何使用get_global_id唯一标识执行内核的工作项列表中的一个工作项。

OpenCL C编程语言用来创建描述数据并行内核和任务的程序，这些内核和任务可以在一个或多个异构设备上执行，如 CPU、GPU和另外一些称为加速器的处理器（如DSP和 Cell Broad-band Engine (B.E.）处理器)。OpenCL程序类似于一个动态库，OpenCL 内核则类似于动态库的一个导出函数。应用程序可以直接从代码调用由动态库导出的函数。不过，应用程序不能直接调用OpenCL内核，只能将内核的执行放在一个为设备创建的命令队列中排队。内核与宿主机CPU上运行的应用代码异步执行。

OpenCL C基于ISO/IEC 9899：1999 C语言规范（简称为C99)，并针对并行性对语言做了一些限制和特定扩展。

OpenCL C还为C99增加了以下特性：
1）矢量数据类型 很多OpenCL设备（如Intel SSE、面向POWER和Cell的AltiVec，以及ARM NEON）都支持矢量指令集。这个矢量指令集在C/C++代码中通过内置函数访问(其中一些可能特定于设备)，或者利用设备特定的汇编指令访问。类似于C语言中使用标量类型，在OpenCL C中可以采用同样的方式使用矢量数据类型。基于这一点，开发人员可以更容易地编写矢量代码，因为可以对矢量和标量数据类型使用类似的操作符。这样还便于编写可移植的矢量代码，因为现在的OpenCL编译器会负责将OpenCL C中的矢量操作映射到设备上适当的矢量ISA。基于常规的内存访问以及更好地结合这些内存访问，矢量代码还有助于提高内存带宽。
2）地址空间限定符 诸如GPU等OpenCL设备实现了一个内存层次结构。地址空间限定符用来标识这个层次结构中的一个特定内存区域。
3）对语言的并行性补充 这包括对工作项和工作组的支持，还包括对工作组中工作项之间的同步提供支持。
4）图像 OpenCLC增加了图像和采样器数据类型，还增加了读、写图像的内置函数。
5）庞大的内置函数集合 如数学函数、整数函数、几何函数和关系函数。

标量数据类型

OpenCL C支持的C99标量数据类型如下1。与C不同，OpenCL C指定了整数和浮点数据类型的大小，也就是具体的位数。

bool                       这是一个条件数据类型，可以为true或false值true可以扩展为整数常量1，值false扩展为整数常量0char                       有符号8位整数，2的补值
unsigned char uchar        无符号8位整数
short                      有符号16位整数，2的补值
unsigned short ushort      无符号16位整数
int                        有符号32位整数，2的补值
unsigned int uint          无符号32位整数
long                       有符号64位整数，2的补值
unsigned long ulong        无符号64位整数float                      32位浮点数float数据类型必须符合IEEE 754单精度存储格式double                     64位浮点数double 数据类型必须符合IEEE 754双精度存储格式这是一个可选的格式，只有当设备支持双精度扩展(cl_khr_fp64)时才可用half                       16位浮点数half数据类型必须符合IEEE 754-2008半精度存储格式size_t                     无符号整数类型，这是 sizeof操作符结果的类型如果设备的地址空间为32位，这就是一个32位无符号整数;如果设备的地址空间是64位，这就是个64位无符号整数ptrdiff_t                  有符号整数类型，这是两个指针相减结果的类型如果设备的地址空间为32位，这就是一个32位有符号整数;如果设备的地址空间是64位，这就是一个64位有符号整数intptr_t                   有符号整数类型，它有一个性质，任何指向void的合法指针都可以转换为这个类型，然后还可以再转换回指向void的指针，其结果与原指针比较是相等的uintptr_t                  无符号整数类型，它有一个性质，任何指向void 的合法指针都可以转换为这个类型，然后还可以再转换回指向void的指针，其结果与原指针比较是相等的void                       void类型构成值的一个空集这是一个不完整的类型，而且不能补全

half数据类型

half数据类型必须符合IEEE 754—2008。Half数有一个符号位、5个指数位和10个尾数位。符号、指数和尾数的解释与IEEE754浮点数的相应解释类似。指数偏差为15。half数据类型必须能表示有限标准数、非规格化数、无穷大和NaN（非数字)。half数据类型的非规格化数可能在使用内置函数vstore_half将一个float转换为half时生成，也可能在使用内置函数vload_half将一个half转换为float时生成，这些非规格化数不能刷新为0。

从float转换为half时会适当地将尾数舍入为11位精度。Half到float的转换则是无损的，所有half数都可以准确地表示为float值。

half数据类型只能用来声明指针指向一个包含half值的缓冲区。下面给出几个合法使用half数据类型的例子：

void bar(global half *p)
{....
}void foo(global half *pg, local half *pl)
{global half *ptr;int offset;ptr = pg + offset;bar(ptr);
}

下面是一个非法使用half类型的例子：

half a;
half a[100];half *p;
a = *p; //not allowed. must use vload_half function

可以使用vload_half、vload_halfn、vloada_halfn 以及 vstore_half、vstore_halfn和vstorea_halfn函数分别完成对half指针的加载和存储。加载（load）函数从内存读取标量或矢量half值，将其转换为一个标量或矢量浮点值。存储（store）函数将一个标量或矢量浮点值作为输入，将它转换为一个half标量或矢量值（采用适当的舍入模式)，并把这个half标量或矢量值写入内存。