OpenCL内核编程：数据类型与设备内存详解

### OpenCL 内核编程：数据类型与设备内存详解 #### 1. OpenCL 数据处理基础 在 OpenCL 编程中，了解如何对内核进行编码和排队是很关键的，但如果没有数据可供处理，内核就毫无用处。OpenCL 提供了一种标准机制，通过内存对象（`cl_mem`）数据结构在主机和设备之间传输数据。 数据传输过程如下： 1. 从现有数据创建一个内存对象。 2. 调用 `clSetKernelArg` 函数，将该对象设置为内核参数。当内核执行时，就能像访问常规函数参数一样访问其数据。 3. 主机发送进一步命令后，设备可以将数据传输回主机，或者将数据复制到另一个缓冲区对象。 内存对象主要有两种类型： - 缓冲区对象：以一维形式存储通用数据。 - 图像对象：以二维或三维形式存储格式化的像素数据。 对于这两种类型，OpenCL 都提供了用于排队数据传输命令的函数。读写函数可在内存对象和主机之间传输数据，不过通常将内存对象的内存映射到主机内存可以提高性能。 数据分区对于要求高性能的 OpenCL 应用程序至关重要。基本的工作单元是工作项（work - item），它对应于传统 C/C++ 循环中执行的代码。每个工作项都会获得一个全局 ID，以便访问专门为其准备的数据。如果工作项需要同步，可以将它们分组到工作组（work - group）中，每个工作组在设备的单个计算单元上执行。 #### 2. 引入内核编码 现在我们来看看实际的 OpenCL 内核。以下是一个类似于 C 语言中常见的 “Hello World!” 函数的 OpenCL 内核示例： ```c __kernel void hello_kernel(__global char16 *msg) { *msg = (char16)('H', 'e', 'l', 'l', 'o', ' ', 'k', 'e', 'r', 'n', 'e', 'l', '!', '!', '!', '\0'); } ``` OpenCL 内核与常规 C 函数有一些主要区别： - 每个内核声明必须以 `__kernel` 开头。 - 每个内核函数必须返回 `void`。 - 有些平台在没有参数的情况下不会编译内核。 内核函数可以按值或按引用接受参数。按值传递时，提供实际数据，如 `char`、`int` 或 `float`，且内核函数不支持复合结构。按引用传递时，提供一个指向设备内存中数据的指针（通常是一个内存对象）。 所有传递给内核的指针都必须带有地址空间限定符，这告诉设备参数应存储在哪个地址空间。常见的限定符有 `__global`、`__constant`、`__local` 和 `__private`。在上述示例中，`msg` 参数应存储在设备的全局地址空间。 主机应用程序从内存对象创建内核参数的代码如下： ```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` 访问缓冲区数据： ```c clEnqueueReadBuffer(queue, msg_buffer, CL_TRUE, 0, sizeof(msg), &msg, 0, NULL, NULL); ``` 需要注意的是，主机将 `msg` 声明为 `char[16]`，而内核将其声明为 `char16`，但由于数据是按引用传递的，这对编译器没有影响。 #### 3. OpenCL 标量数据类型 在向量计算中，标量是指每个数据表示只包含单个值的数据类型。OpenCL 中的标量数据类型如下表所示： | 标量数据类型 | 用途 | | --- | --- | | `bool` | 布尔条件：真（1）或假（0） | | `char` | 有符号二进制补码 8 位整数 | | `unsigned char/uchar` | 无符号二进制补码 8 位整数 | | `short` | 有符号二进制补码 16 位整数 | | `unsigned short/ushort` | 无符号二进制补码 16 位整数 | | `int` | 有符号二进制补码 32 位整数 | | `unsigned int/uint` | 无符号二进制补码 32 位整数 | | `long` | 有符号二进制补码 64 位整数 | | `unsigned long/ulong` | 无符号二进制补码 64 位整数 | | `half` | 16 位浮点值，符合 IEEE - 754 - 2008 标准 | | `float` | 32 位浮点值，符合 IEEE - 754 标准 | | `intptr_t` | 可将 `void` 指针转换为的有符号整数 | | `uintptr_t` | 可将 `void` 指针转换为的无符号整数 | | `ptrdiff_t` | 指针相减产生的有符号整数 | | `size_t` | 由 `sizeof` 运算符产生的无符号整数 | | `void` | 无类型数据 | ##### 3.1 访问双精度数据类型 双精度（`double`）数据类型在目标设备支持 `cl_khr_fp64` 扩展时可用。从主机端，可以通过调用 `clGetDeviceInfo` 函数来确定该扩展是否可用。如果支持该扩展，可以在核中使用以下 `pragma` 语句启用其功能： ```c #pragma OPENCL EXTENSION cl_khr_fp64 : enable ``` 以下是一个示例代码，展示了如何根据设备是否支持 `double` 类型来进行不同的计算： ```c #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); *out = ```