本章涵盖
- 介绍一个简单的 OpenCL 内核
- 使用 OpenCL 的标量和矢量数据类型
- 了解 OpenCL 设备模型
在本章中,我们将抛开创建和部署内核的脚手架,开始编写内核本身。 我们将检查 OpenCL 内核中可用的数据类型,这意味着我们最终将讨论向量。 当您使用向量处理数据时,您可以抛开 char、float 和 int 等枯燥的、几十年前的数据类型,而使用新的、令人兴奋的数据类型,例如 char16、int3 和 float4。
直到大学毕业后我才学习矢量编程,但从那以后我一直很喜欢它。 无论是英特尔的 Streaming SIMD Extensions (SSE)、摩托罗拉的 AltiVec,还是 IBM 设计用于在 Cell 处理器上对协同处理器单元 (SPU) 进行编程的奇怪语言,都无关紧要。 我只是觉得用一个命令来处理几个数字很令人满意,当我同时在几个核心上处理数字时,我的乐趣就会增加。 还有什么可以问的?
在检查了不同类型的数据之后,我们将看看这些数据的存储方式和位置。 OpenCL 有一个包含四个不同地址空间的设备模型。 本章的最后几节将讨论这些空间以及如何在代码中配置数据存储。
但在深入了解数据和内存存储的细节之前,了解内核函数的基本结构很重要。 我们将首先讨论这个问题。
4.1 介绍内核编码
第 2 章解释了主机应用程序如何将内核发送到设备,第 3 章解释了如何为内核设置参数。 现在,终于,我们准备好查看一个实际的内核了。 下面的清单展示了古老的 Hello World 的 OpenCL 等价物! 在 C 编程文献中如此常见的函数。
清单 4.1 一个基本内核:hello_kernel.cl
__kernel void hello_kernel(__global char16 *msg) {
*msg = (char16)('H', 'e', 'l', 'l', 'o', ' ',
'k', 'e', 'r', 'n', 'e', 'l', '!', '!', '!', '\0');
}
如果你看一下这个函数的整体结构,你会发现它类似于一个普通的 C 函数:一个函数名、括号中的参数和大括号中的可执行语句。 但是 OpenCL kernel 和常规 C 函数之间存在三个主要区别:
- 每个内核声明都必须以 __kernel 开头。
- 每个内核函数都必须返回void。
- 一些平台不会编译没有参数的内核。
本书中的每个示例项目都将内核函数存储在 *.cl 文件中,但这个后缀不是必需的。 事实上,内核根本不必存储在单独的文件中。 但是每个内核函数都必须以 __kernel 关键字开头。 如果 __kernel 存在,编译器将知道该函数旨在在设备上运行,而不是在主机上运行。
clSetKernelArg 函数为内核设置参数,但没有访问内核返回值的函数。 这是因为内核没有返回值——每个内核函数都返回 void。 因此,本书中的每个内核都具有相同的基本结构:
__kernel void func_name(args) {
...
}
… 部分是最难的部分,讨论这个需要很多章。 现在,让我们看看论据。 内核函数只能通过它的参数访问和返回数据,如果你试图编译一个没有参数的内核,一些编译器会给你一个错误。
与常规 C 函数一样,内核函数按值或按引用接受参数。 当您按值传递数据时,您提供的是实际数据,例如 char、int 或 float。 内核函数不支持复合结构。 如果通过引用传递数据,则提供一个引用设备内存(通常是内存对象)中数据的指针。 在清单 4.1 中,msg 参数引用了一个 16 字节的缓冲区对象,主机应用程序将在内核执行后读取该对象。
现在我们来到重要的一点:传递给内核的所有指针都必须以地址空间限定符开头。 这告诉设备参数应该存储在哪个地址空间。第 4.5 节深入讨论了这个主题,但现在,请记住有四个可能的限定符:__global、__constant、__local 和 __private。 在清单 4.1 中,函数声明声明 msg 参数应该存储在设备的全局地址空间中。
在继续之前,让我们回顾一下主机应用程序如何从内存对象创建内核参数。 在 hello_kernel.c 中,这是通过以下代码行完成的:
char msg[16];
cl_mem msg_buffer;
msg_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
sizeof(msg), NULL, &err);
clSetKernelArg(kernel, 0, sizeof(cl_mem), &msg_buffer);
在内核入队并且设备执行函数后,主机应用程序使用 clEnqueueReadBuffer 访问缓冲区数据。 这显示在这里:
clEnqueueReadBuffer(queue, msg_buffer, CL_TRUE, 0,
sizeof(msg), &msg, 0, NULL, NULL);
请注意,主机将 msg 声明为 char[16],内核将 msg 声明为 char16。 这些是不同的数据类型,但是因为数据是通过引用传递给内核的,所以对编译器没有任何影响。
char16 数据类型是 OpenCL 的向量数据类型之一,4.3 节将详细讨论这些类型。 本书中的内核代码将尽可能依赖向量,但在我们研究向量之前,我们需要检查 OpenCL 对传统数据类型(如整数和浮点数)的支持。 与向量类型相比,它们被称为标量数据类型,它们将在下一节中讨论。
4.2 标量数据类型
术语标量和向量具有不同的含义,具体取决于您与数学家、科学家还是程序员交谈。 在向量计算中,标量是一种数据类型,其中每个数据表示都包含一个值。 在 OpenCL 中,标量是表 4.1 中列出的任何数据类型。
这些数据类型简单明了,功能与它们的 C/C++ 对应物一样。 但是当我第一次阅读这个列表时,我想到了一个突出的问题:double 在哪里? 对于非图形应用程序,我更喜欢使用 64 位浮点值。 OpenCL 中是否提供 double ? 答案是也许。
4.2.1 访问双精度数据类型
如果目标设备支持 cl_khr_fp64 扩展,则可以访问双精度数据类型。 在主机上,您可以通过调用 clGetDeviceInfo(第 2 章中介绍的函数)来确定此扩展是否可用。如果支持该扩展,您可以使用以下 pragma 语句在内核中启用其功能:
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
当它存在时,您可以声明 double 变量并正常操作它们。 如果要启用每个受支持的扩展,请将 cl_khr_fp64 替换为 all。 要禁用扩展,请将 enable 替换为 disable。
在 Ch4/double_test 项目中,如果支持,内核使用 double 类型,如果不支持,则使用 float 类型。 这显示在以下清单中。
清单 4.2 检查双精度数据类型:double_test.cl
// 启用扩展(如果可用)
#ifdef FP_64
#pragma OPENCL EXTENSION cl_khr_fp64: enable
#endif
__kernel void double_test(__global float* a,
__global float* b,
__global float* out) {
#ifdef FP_64
double c = (double)(*a / *b); // 如果可用,用double计算
*out = (float)c;
#else
*out = *a * *b;
#endif
}
宿主应用程序调用 clGetDeviceInfo 来获取设备支持的扩展。 如果 cl_khr_fp64 是其中之一,则主机将选项 -DFP_64 添加到 clBuildProgram。 如清单 4.2 所示,该选项告诉内核启用 cl_khr_fp64 扩展。 一旦启用此扩展,内核就可以声明双精度值并对其进行操作。
主机代码还检查目标设备的地址宽度。 如果您在位级别处理 size_t 和 ptrdiff_t 类型,这将变得很重要。 size_t 和 ptrdiff_t 类型在 64 位系统上为 64 位宽,在 32 位系统上为 32 位宽。
4.2.2 字节顺序
表 4.1 告诉您数据类型中有多少字节,但它没有说明字节的顺序。 OpenCL 标准也没有。 原因是不同的设备和操作系统对字节的顺序不同。
因此,如果您要执行涉及字节顺序的操作,例如使用指针访问数据,则需要确定目标设备的字节序。 这告诉您随着内存地址从低到高运行,字节是否变得或多或少重要。 图 4.1 以图形方式描述了这一点。
我发现通过记住 big-endianness 对我来说更直观(我宁愿口袋里有 43.21 美元而不是 12.34 美元),很容易区分这两者。 但是 little-endianness 更为普遍,因为 x86 设备是 little-endian。 最常见的大端处理器是 IBM 的 POWER 和 PowerPC 架构。
有两种方法可以确定设备是 little-endian 还是 big-endian。 您可以从主机调用 clGetDeviceInfo,并将 CL_DEVICE_ENDIAN_LITTLE 作为参数。 如果返回 CL_TRUE,则设备是 little-endian。 如果它返回 CL_FALSE,则设备为大端。
在内核中,您可以使用#ifdef 来确定是否定义了__ENDIAN_LITTLE__ 宏。 如果定义了此宏,则设备为 littleendian。 如果不是,则设备为大端
当我们查看向量时,我们将进一步讨论字节顺序。 但在我们离开标量主题之前,我们需要检查 OpenCL 如何处理浮点值。