[OpenCL] 主机编程：数据传输和分区(12)

3.6 数据分区

  如果您使用 OpenCL 实现算法，您可能需要处理大量数据。 这使得数据分区成为一个重要的优先事项——处理负载分配得越好，计算任务就会越早完成。

  您已经知道如何在多个设备之间划分数据，但您可以进一步划分数据。 大多数 OpenCL 设备包含多个处理元素，并且使用正确的代码，您可以控制每个处理元素接收的数据量。

  只有一个函数需要知道：clEnqueueNDRangeKernel。 这是 OpenCL API 中最重要的功能之一，也是最复杂的功能之一。 与 clEnqueueTask 一样，这会将 kernel 置于命令队列中以供执行。 但与 clEnqueueTask 不同，clEnqueueNDRangeKernel 允许您控制 kernel 执行在设备处理资源中的分布方式。 这通过它的签名来显示，如下所示：

clEnqueueNDRangeKernel(cl_command_queue queue, cl_kernel kernel,
	cl_uint work_dims, const size_t *global_work_offset,
	const size_t *global_work_size, const size_t *local_work_size,
	cl_uint num_events, const cl_event *wait_list, cl_event *event)

  这比 clEnqueueTask 要复杂得多。 这两个函数的区别在于 clEnqueueNDRangeKernel 接受四个附加参数：

• work_dims - 数据中的维数
• global_work_offset - 每个维度中的全局 ID 偏移量
• global_work_size - 每个维度中的工作项数
• local_work_size - 每个维度中工作组中的工作项数

  如果这些术语还没有意义，请不要担心。 本节的目的是解释它们的含义以及如何配置它们，以便您可以充分利用您的硬件。

3.6.1 循环和工作项

  当您有大量数据时，通常使用循环遍历数据。 如果您需要在常规 C/C++ 中处理多维数据，您可能会使用嵌套循环，例如：

for(i=0; i<Z; i++) {
	for(j=0; j<Y; j++) {
		for(k=0; k<X; k++) {
			process(point[i][j][k]);
		}
	}
}

  像这样的循环很常见但效率低下。 效率低下的原因是每次迭代都需要单独的比较和加法。 在最好的处理器上进行比较非常耗时，但在图形处理器单元 (GPU) 等专用数字处理器上比较慢。 GPU 擅长一遍又一遍地执行相同的操作，但它们不擅长做出决策。 如果 GPU 必须检查条件和分支，它可能需要数百个周期才能恢复全速处理数字。

  OpenCL 的一个令人着迷的方面是您不必在内核中配置这些循环。 相反，您的 kernel 只执行位于最内层循环内的代码。 我们称这个单独的 kernel 执行为工作项 (work-items) 。 在前面的示例循环中，工作项由单个函数调用组成：process(point[i][j][k])。

  了解 kernels 和工作项 (work-items) 之间的区别很重要。 内核标识要对数据执行的一组任务。 工作项是 kernels 在特定数据集上的单个实现。 对于每个 kernel ，可以有多个工作项。 在前面的示例中，内核可能由 process(point[i][j][k]) 表示。 此内核的特定实现，例如 process(point[1][2][3])，将是一个工作项。

  数组 {i, j, k} 称为工作项的全局 ID。 它唯一地标识工作项并允许它访问它应该处理的数据。 例如，以下内核代码访问项目 ID 的元素并处理一个点：

int i = get_global_id(0);
int j = get_global_id(1);
int k = get_global_id(2);
process(point[i][j][k]);

  一旦这个工作项被执行，一个新的工作项将以不同的全局 ID 执行。

  全局 ID 中的元素数量称为数据的维度。 您可以通过设置 clEnqueueNDRangeKernel 的 work_dims 参数来配置它。 最小维数为 1，最大维数取决于设备。 要查找最大维数，请使用 CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS 参数调用 clGetDeviceInfo。 在之前的 i-j-k 循环中，您将设置 work_dims 等于 3。

  我们之前在查看缓冲区对象和图像对象时检查了维度。 图像对象可以是二维或三维的，而缓冲区对象只能在一维中访问。 clEnqueueNDRangeKernel 不关心这种区别。 如果您正在处理图像对象，您可能应该将 work_dims 设置为 2 或 3。但是对于缓冲区对象，您可以设置您认为最好的任何维度。 对于包含二维矩阵的缓冲区对象，如图 3.5 所示，您可以将 work_dims 设置为等于 2。

3.6.2 工作尺寸和偏移量

  图 3.7 的左侧描绘了一个处理循环。 右侧呈现循环对应的索引空间。 索引空间包含所有可能的索引组合。 如果一个循环中有 N 个不同的索引，则对应的索引空间有 N 个维度。

  clEnqueueNDRangeKernel 的 global_work_sizes 参数标识每个维度需要处理多少工作项。 内循环从 k=3 开始，一直到 k=11，因此在 k 方向上有 9 个工作项要处理。 类似地，在 j 方向有 6 个要处理的工作项，在 i 方向有 4 个要处理的工作项。 因此，您需要将 global_work_sizes 设置为 {4, 6, 9}。

  当第一个工作项开始执行时，您希望它访问对应于索引三元组 (0, 2, 3) 的数据，因为这些是 i、j 和 k 的初始值。 换句话说，您希望第一个工作项的全局 ID 等于 {0, 2, 3}。 您可以通过将 clEnqueueNDRangeKernel 中的 global_work_offset 设置为 {0, 2, 3} 在代码中指定这一点。

注意 global_work_offset 在本书的示例代码中总是设置为 NULL。

3.6.3 一个简单的一维示例

  了解 clEnqueueNDRangeKernel 的一个好方法是查看它是如何在代码中使用的。 在第 1 章中，我尽我所能通过展示完整 OpenCL 应用程序的代码来吓唬你。 此应用程序将向量乘以矩阵并生成向量。 图 3.8 显示了计算。 它还描述了包含矩阵数据的缓冲区对象以及矩阵数据在四个工作项之间的划分方式。

  矩阵向量乘法由四个点积组成，我选择使用四个工作项来执行乘法。 这是通过以下代码完成的：

work_items_per_kernel = 4;
clEnqueueNDRangeKernel(queue, kernel, 1, NULL,
	&work_items_per_kernel, NULL, 0, NULL, NULL);

  这告诉 OpenCL，要分区的数据只有一个维度，并且应该生成四个工作项来执行内核。 全局偏移设置为 0。

  在内核方面，每个工作项检查其全局 ID 并访问矩阵的一行。 它使用 dot 函数将此行 (1×4) 与向量 (4×1) 相乘，并将结果 (1×1) 放在由其 ID 确定的数组位置中。 这显示在以下代码中：

int i = get_global_id(0);
result[i] = dot(matrix[i], vector[0]);

  看？ 没有循环。 这四个工作项并行运行，没有与 for 语句或类似结构相关的延迟。

3.6.4 工作组和计算单元

  工作组是访问相同处理资源的工作项的组合。 在编程方面，工作组提供了两个主要优势：

• 工作组中的工作项可以访问称为本地内存的同一块高速内存
• 工作组中的工作项可以使用栅栏和障碍进行同步。

  第 4 章解释了 OpenCL 设备中不同类型的内存，第 7 章讨论了工作项同步。 目前，我的目标是足够深入地解释工作组，以便您完全理解 clEnqueueNDRangeKernel 的论点。

  除了全局 ID，每个工作项都有一个本地 ID，将其与工作组中的所有其他工作项区分开来。 工作组中工作项的数量通过 clEnqueueNDRangeKernel 的 local_work_size 参数设置。 此数组中的元素标识每个维度中工作组中可以容纳多少工作项。

  例如，让我们从图 3.7 中描述的索引空间中的二维切片创建工作组。 有四个切片，所以我们将有四个工作组。 每组在 j 方向包含 6 个工作项，在 k 方向包含 9 个工作项。 因此，我们将 clEnqueueNDRangeKernel 中的 local_work_size 设置为 {0, 6, 9}。

  在 OpenCL 中，能够支持工作组的处理资源称为计算单元。 每个工作组在单个计算单元上执行，每个计算单元一次只执行一个工作组。 图 3.9 以图形方式显示了这种关系。

  您不必创建工作组。 如果您将 local_work_size 设置为 NULL，OpenCL 将决定如何最好地在设备的处理元素之间分配工作项。

注意 如果您对工作组、工作组 ID、计算单元和本地 ID 的主题仍然不满意，请不要担心。 本书将讨论这些抽象但重要的概念

3.7 小结

  知道如何对内核进行编码和入队至关重要，但是如果没有要处理的数据，内核就毫无用处。 本章的目标是展示 OpenCL 如何打包和分区数据以供设备处理。

  OpenCL 提供内存对象 (cl_mem) 数据结构作为在主机和设备之间传输数据的标准机制。 传输内存对象的过程很简单：从现有数据创建一个内存对象，然后调用 clSetKernelArg 将该对象变成内核参数。 当内核执行时，内核将能够以常规函数参数的形式访问其数据。 然后，随着主机发送进一步的命令，设备可以将数据传输到主机或将数据复制到另一个缓冲区对象。

  有两种类型的内存对象。 缓冲区对象以一维存储一般数据，图像对象以二维或三维存储格式化的像素数据。 对于这两种类型，OpenCL 都提供了将数据传输命令排入队列的功能。 读/写函数在内存对象和主机之间传输数据，但您通常可以通过将内存对象的内存映射到主机内存来提高性能。

  本章的最后一部分讨论了数据分区，这对于任何需要高性能的 OpenCL 应用程序都至关重要。 工作的基本单元是工作项，它对应于在传统 C/C++ 循环中执行的代码。 每个工作项都会收到一个全局 ID，允许它访问专门用于它的数据。 如果工作项需要同步，可以将它们放入工作组中。 每个工作组在设备上的单个计算单元上执行。

  本章与第 2 章一起，解释了您需要了解的有关主机应用程序的几乎所有内容。 唯一需要讨论的主题是同步、事件处理和线程，它们将在第 7 章中讨论。在下一章中，我们将离开主机编程并开始我们对内核开发的讨论。