[OpenCL] 内核编程：数据类型和设备内存(13)

本章涵盖

介绍一个简单的 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_64double 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 如何处理浮点值。