OpenCL 2.0 中的 GPU-Quicksort：嵌套并行性和工作组扫描函数
4.60/5 (3投票s)
本教程展示了如何使用 OpenCL™ 2.0 的两个强大功能：enqueue_kernel 函数，它允许您从设备 en-queue 内核，以及 work_group_scan_exclusive_add 和 work_group_scan_inclusive_add。
Intel® Developer Zone 为跨平台应用开发提供工具和操作指南、平台和技术信息、代码示例以及同行专业知识，帮助开发人员进行创新并取得成功。加入我们的社区，获取 Android、物联网、Intel® RealSense™ 技术和 Windows 的开发工具，下载工具，访问开发套件，与志同道合的开发人员交流想法，并参与黑客马拉松、竞赛、路演和本地活动。
- 引言
- 快速排序简史
- GPU-Quicksort 简介
- OpenCL 1.2 中的 GPU-Quicksort
- 将 GPU-Quicksort 转换为 OpenCL 2.0
- 教程要求
- 运行教程
- 结论
- 参考文献
引言
本教程展示了如何使用 OpenCL™ 2.0 的两个强大功能：enqueue_kernel 函数，它允许您从设备 en-queue 内核，以及 work_group_scan_exclusive_add 和 work_group_scan_inclusive_add，它们是 OpenCL 2.0 中新增的一组工作组函数，用于简化工作组内工作项之间的扫描和归约操作。本教程在 OpenCL 中的 GPU-Quicksort 实现上演示了这些功能，据我们所知，这是该算法在 OpenCL 中的首次实现。本教程展示了一种重要的设计模式，即 en-queue 大小为 1 的 NDRange 内核，以执行以前由 CPU 保留的后台操作和调度操作。
在本教程中，您将了解如何使用 Intel® OpenCL 工具。Intel® SDK for OpenCL™ Applications 是在 Intel® 架构上开发 OpenCL 应用程序的开发平台。除了 SDK，Intel® VTune™ Amplifier XE 还是分析和优化 Intel 平台上 CPU 和 GPU 应用程序的工具，并且支持 OpenCL。请访问 http://intel.com/software/opencl 了解更多关于 Intel OpenCL 工具的信息。
我要感谢 Deepti Joshi、Allen Hux、Aaron Kunze、Dillon Sharlet、Adam Lake、Michal Mrozek、Ben Ashbaugh、Ayal Zaks、Yuri Kulakov、Joseph Wells、Lynn Putnam、Jerry Baugh、Jerry Makare 和 Chris Davis 在开发、审查和发布本文、配套代码以及编辑视频方面的帮助。我还要感谢我的妻子 Ellen 和我的孩子 Michael 和 Celine 的坚定支持和理解。
快速排序简史
快速排序是由 C.A.R.（“Tony”）Hoare 于 1960 年在莫斯科国立大学期间发明的算法。Tony Hoare 是 A.N. Kolmogorov 教授门下攻读概率论的研究生。Hoare 在皇家海军服役期间学会了俄语，他的叔叔是一名海军上尉。Hoare 在从事语言机器翻译问题的研究时，试图将每句输入的俄语句子的单词按升序排序。快速排序算法的基本思想是将待排序序列围绕着一个可能从序列中以多种方式选出的枢轴元素进行分区，使得所有小于枢轴的元素都放在数组的左侧，所有大于枢轴的元素都放在数组的右侧，所有等于枢轴的元素都放在数组的中间。在对数组进行分区后，快速排序算法会重新应用于序列的左侧和右侧。
GPU-Quicksort 简介
首个已知的 GPU 快速排序算法由 Shubhabrata Sengupta、Mark Harris、Yao Zhang 和 John D. Owens 在他们的论文“Scan Primitives for GPU Computing”中开发和描述。我们在此介绍的 GPU-Quicksort 算法最初由 Daniel Cederman 和 Philippas Tsigas 在他们的论文“GPU-Quicksort: A practical Quicksort algorithm for graphics processors”中描述，并且是第一个 GPU 快速排序的改进。该算法最初是用 CUDA* 编写的，可在 Nvidia* 独立显卡上运行，但很容易在 OpenCL 中实现，可在支持 OpenCL API 的任何硬件上运行。该算法专门设计用于利用高 GPU 带宽，并且在 OpenCL 1.2（Intel® HD Graphics 4600）和 OpenCL 2.0 驱动程序（Intel® HD Graphics 4600 和 Intel® HD Graphics 5300 或更高版本）上都能很好地运行。有关 GPU-Quicksort 的精彩概述，请参阅 Cederman 和 Tsigas 论文的 3.1 节。
与原始的快速排序一样，GPU-Quicksort 递归地将元素序列围绕着枢轴进行分区，直到整个序列被排序。由于该算法是为 GPU 编写的，因此它包含两个阶段，在第一阶段中，多个工作组处理同一序列的不同部分，直到子序列足够小，可以在第二阶段中由每个工作组完全排序。
GPU-Quicksort 的主要思想是将输入序列分成块并分配给不同的工作组，每个工作组通过一个两阶段过程参与围绕枢轴对输入序列进行分区：首先，工作组中的每个工作项计算它在块中看到的小于和大于枢轴的元素数量。工作组扫描添加内置函数用于查找小于和大于枢轴的元素的累积总数。使用辅助数组为每个工作组分配空间。第二次遍历数据会将小于和大于枢轴的元素写入分配的空间。最后，最后一个工作组用枢轴值填充间隙。
OpenCL 1.2 中的 GPU-Quicksort
OpenCL 1.2 中的 GPU-Quicksort 是一个直接的实现，遵循 Cederman/Tsigas 论文中描述的算法。这三个部分是 GPUQSort 函数实现的 CPU 部分，gqsort_kernel OpenCL 内核函数实现的 GPU 第一阶段内核，以及 lqsort_kernel OpenCL 内核函数实现的 GPU 第二阶段内核。GPUQSort 函数会迭代地启动 gqsort_kernel，直到初始序列被划分为足够小的块，以便每个块可以由一个工作组进行排序。此时，GPUQSort 会启动 lqsort_kernel 来完成排序任务。
GPU-Quicksort 实现需要 OpenCL 1.2 中可用的障碍（barriers）和原子函数（atomic functions），并且几乎所有能够运行 OpenCL 1.2 的现代硬件都支持这些函数。具体来说，atomic_add、atomic_sub 和 atomic_dec 用于实现 gqsort_kernel。障碍在 gqsort_kernel 和 lqsort_kernel 中都得到了广泛使用。
使用 OpenCL 1.2 的缺点是：
- 缺乏工作组扫描原语，这需要使用 Guy Blelloch 的著名论文中描述的算法来实现前缀和。
- CPU 和 GPU 之间频繁的往返通信，用于启动 gqsort 内核和最后的 lqsort 内核启动。
- 在 CPU 上启动内核之前执行后台操作的必要性，而这些数据已在 GPU 上生成并立即可用。
这三个缺点都在我们的 OpenCL 2.0 转换中得到了解决。您会发现，通过解决这些不足之处，算法的性能得到了显著提升。
将 GPU-Quicksort 转换为 OpenCL 2.0
在本节中，我们将回顾为将 OpenCL 1.2 实现转换为利用新的设备端 en-queue 和工作组扫描功能的 OpenCL 2.0 实现所做的更改。
工作组函数可以提高 OpenCL C 代码的性能和可读性。
将 GPU-Quicksort 转换为 OpenCL 2.0 的第一步也是最简单的一步是利用现成的工作组扫描函数：具体来说，在 gqsort_kernel 和 lqsort_kernel 中使用 work_group_scan_inclusive_add 和 work_group_scan_exclusive_add 函数。我们不仅获得了一些可衡量的性能改进（约 8%），而且在代码简洁性、可维护性和清晰度方面也得到了提升，将与前缀和计算相关的代码量减少了近 3 倍。
块（Blocks）使我们能够通过设备端 en-queue 利用嵌套并行。
接下来，我们将对在每次 gqsort_kernel 运行后排序记录的逻辑进行转换，使其适用于 lqsort_kernel 处理的记录以及需要进一步细分的记录，计算块和父记录，并将 gqsort_kernel 或 lqsort_kernel 从 C++ en-queue 到 OpenCL C。lqsort_kernel 将仅在 GPU-Quicksort 运行的最后启动一次。在我们最初的实现中，该逻辑位于 gqsort_kernel 的末尾，这需要额外的全局变量和原子操作，并导致 gqsort_kernel 相对较慢，尽管整个样本的性能与 OpenCL 1.2 版本相比有所提高。我们将此逻辑移至一个单独的 relauncher_kernel，该内核现在作为第一个内核启动，然后在每次 gqsort_kernel 运行结束后启动。
// now let’s recompute and relaunch
if (get_global_id(0) == 0) {
uint num_workgroups = get_num_groups(0);
queue_t q = get_default_queue();
enqueue_kernel(q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL,
ndrange_1D(1),
^{ relauncher_kernel(d, dn,
blocks, parents, result, work, done, done_size,
MAXSEQ, num_workgroups);
});
}
请注意，relauncher_kernel 仅在一个工作项上启动（ndrange = 1）。将重新启动器逻辑与 qsort_kernel 分离，极大地简化了 qsort_kernel，并带来了提高算法性能的额外好处。
关于我们如何 en-queue relauncher_kernel 的几点说明：我们使用了 OpenCL 2.0 中引入的新功能——块（blocks）。由于内核不需要局部内存参数，因此我们使用 void (^)（void）类型的块。块类似于 C++ lambda 表达式，但语法不同。块会捕获传递给它的参数，从而避免了通过 clSetKernelArg 等效函数设置内核参数的繁琐工作。第二点值得一提的是 num_workgroups 的计算方式：它在块外部计算，然后传递给块。请注意，在上面的代码示例中，num_workgroups 将按值捕获，这正是我们想要的。如果我们直接使用 get_num_groups(0) 而不是 num_workgroups，它将在包含 relauncher_kernel 调用的块被 en-queue 和执行后调用，因此传递给 relauncher_kernel 的值将是 1，而不是 gqsort_kernel 的工作组数量。有关 OpenCL C 使用的块语法的更多信息，请参阅 OpenCL C 规范 [3]。
relauncher_kernel 根据是否仍有工作可供 gqsort_kernel 执行（work_size != 0）来启动 gqsort_kernel 或 lqsort_kernel。
if (work_size != 0) {
// Calculate block size, parents and blocks
...
enqueue_kernel(q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL,
ndrange_1D(GQSORT_LOCAL_WORKGROUP_SIZE * blocks_size,
GQSORT_LOCAL_WORKGROUP_SIZE),
^{ gqsort_kernel(d, dn,
blocks, parents, result, work, done,
done_size, MAXSEQ, 0); });
} else {
enqueue_kernel(q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL,
ndrange_1D(LQSORT_LOCAL_WORKGROUP_SIZE * done_size,
LQSORT_LOCAL_WORKGROUP_SIZE),
^{ lqsort_kernel(d, dn, done); });
}
一旦我们从 CPU 启动第一个 relauncher_kernel，后续的所有内核启动都将从 GPU 执行。在所有输入数据完全排序后，控制权将返回给 CPU。
教程要求
构建和运行本教程需要满足以下要求的 PC：
- 代号为 Broadwell 的 Intel® Core™ 处理器系列处理器
- Microsoft Windows* 8 或 8.1
- Intel® SDK for OpenCL™ Applications 2014 R2 或更高版本
- Microsoft Visual Studio* 2012 或更高版本
运行教程
本教程是一个控制台应用程序，它生成一个大小为 Width*Height 的随机无符号整数数组。然后，它使用 std::sort、常规单线程快速排序，以及随后多次使用 OpenCL 2.0 中的 GPU-Quicksort 进行排序来排序该数组的副本。该教程支持一些命令行选项：
| 选项 | 描述 |
|---|---|
-h, -? |
显示帮助文本并退出 |
[num test iterations] |
在同一数据上运行 GPU-Quicksort 的次数 |
[cpu|gpu] |
是使用 CPU 还是 GPU OpenCL 运行教程 |
[intel|amd|nvidia] |
选择要在其上运行 OpenCL 设备的供应商 |
[Width] |
输入的“宽度”——便于输入大数字 |
[Height] |
输入的“高度”——便于输入大数字，例如，我们可以提供 8192 8192 来代替 6700 万个元素 |
[show_CL|no_show_CL] |
是否显示详细的 OpenCL 设备信息 |
尝试使用参数 5 gpu intel 8192 8192 show_CL 运行教程。
结论
配备 Intel® HD Graphics 5300 或更高版本的 Intel 处理器是一个复杂的硬件，但当与支持 OpenCL 2.0 的驱动程序结合使用时，可以为您的 OpenCL 代码带来显著的性能改进。OpenCL 2.0 为 GPGPU 程序员提供了许多强大的功能。我们仅介绍了其中两个功能：enqueue_kernel 函数以及 work_group_scan_exclusive_add 和 work_group_scan_inclusive_add 函数，并展示了它们在付出相对较小的代价的同时带来了显著的性能改进，同时简化并提高了代码的可读性。请考虑在 Intel® SDK for OpenCL™ Applications Support 网站上阅读 Intel 关于其他 OpenCL 功能的其他教程。
参考文献
- Intel SDK for OpenCL Applications – Optimization Guide
- Khronos OpenCL 2.0 API Specification
- Khronos OpenCL 2.0 C Language Specification
- Wikipedia 关于 Quicksort 的文章
- Tony Hoare 著《My Early Days at Elliots》
- Shubhabrata Sengupta、Mark Harris、Yao Zhang 和 John D. Owens 著《Scan Primitives for GPU Computing》，Graphics Hardware 2007，第 97-106 页，2007 年 8 月。
- Daniel Cederman 和 Philippas Tsigas 著《GPU-Quicksort: A practical Quicksort algorithm for graphics processors》，Journal of Experimental Algorithmics，第 14 卷，2009 年，文章编号 4
- Guy Blelloch 著《Prefix Sums and Their Applications》
- enqueue_kernel 函数在线文档。
- work group scan inclusive 函数在线文档
- work group scan exclusive 函数在线文档
- Intel® SDK for OpenCL™ Applications
- Intel® SDK for OpenCL Applications 2013 R2 优化指南
- 优化 OpenCL 应用程序的完整清单
- 使用 Intel® OpenCL SDK 编写优化的 OpenCL 代码 - pdf 格式
关于作者
Robert Ioffe 是 Intel 软件与解决方案事业部的一名技术咨询工程师。他是 OpenCL 编程专家，擅长在 Intel Iris 和 Intel Iris Pro Graphics 上进行 OpenCL 工作负载优化，对 Intel 图形硬件有深入的了解。他深度参与了 Khronos 标准工作，专注于最新功能的原型开发，并确保它们能在 Intel 架构上良好运行。最近，他一直在从事 OpenCL 2.0 的嵌套并行（enqueue_kernel 函数）功能的原型开发，并编写了多个演示嵌套并行功能的示例，包括 OpenCL 2.0 的 GPU-Quicksort。他还录制并发布了两个“优化简单 OpenCL 内核”视频，并且正在录制第三个关于嵌套并行的视频。
您可能还对以下内容感兴趣：