OpenCL编程指南-4.1OpenCL C编程

news2024/11/17 17:28:35

使用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扩展为整数常量0

char                       有符号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标量或矢量值写入内存。