OpenCL 2.0 中的 GPU-Quicksort: 嵌套并行性和工作组扫描函数

签署人: Robert M Ioffe 已发布: 09/29/2014 最后更新时间: 03/04/2015


简介

本教程展示了如何使用 OpenCL™ 2.0 的两个强大特性: enqueue_kernel 函数,支持您排列设备中的 kernel;work_group_scan_exclusive_addwork_group_scan_inclusive_add,两个添加至 OpenCL 2.0 的工作组函数,可方便扫描并减少对工作组工作项的操作。 本教程通过在 OpenCL 中实施我们自己的 GPU-Quicksort 来展示了这些特性,据我们所知,这是该算法在 OpenCL 中的首次实施。 本教程展示了一种重要的尺寸 1 NDRange 的排队 kernel 的设计模式,以执行之前为 CPU 保留的管理和调度操作。

在本教程中,您将了解如何使用面向 OpenCL 的英特尔® 工具。 面向 OpenCL™ 应用的英特尔® SDK 是专为英特尔® 架构上的 OpenCL 应用而提供的一个开发环境。 除 SDK 外,英特尔® VTune™ Amplifier XE 也能帮助分析和优化 CPU 和 GPU 上英特尔平台上的应用,并且还支持 OpenCL。 如欲了解面向 OpenCL 的英特尔工具的更多信息,请访问 http://intel.com/software/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 对我的坚定支持与理解。


Quicksort 简史

Quicksort 是 C.A.R. (“Tony”) Hoare 在 1960 年尚在莫斯科国立大学就读时发明的一套算法。 Tony Hoare 是概率论专业研究生,指导教授为 A.N. Kolmogorov。 在皇家海军服役期间,Hoare 学习了俄语(他的叔叔是皇家海军上校)。 在研究语言机器翻译的问题时,Hoare 曾尝试按升序来排列听到的每个俄语句子中的单词。 Quicksort 算法的基本理念是围绕主元要排序的序列(主元可通过不同的方式从序列中选择),所有小于主元的元素放到数组的左侧,所有大于主元的元素放到数组的右侧,所有等于主元的元素放到数组中间。 数组被分割后,序列的左侧和右侧将再次应用 Quicksort 算法。

Partitioning Sequence Quicksort
图 1. 使用 Quicksort 分割序列


GPU-Quicksort 简介

这是第一套面向 GPU 的 Quicksort 算法,由 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 的 Quicksort 改进。 该算法最初使用 CUDA* 编写以在 Nvidia* 独立显卡上运行,可轻松在 OpenCL 中实施以在任何支持 OpenCL API 的硬件上运行。 该算法专为充分利用高 GPU 带宽以及在 OpenCL 1.2(英特尔® 核芯显卡 4600)和 OpenCL 2.0 驱动程序(英特尔® 核芯显卡 4600 和英特尔® 核芯显卡 5300 或更高版本)上出色运行而设计。 如欲详细了解 GPU-Quicksort,请参阅 Cederman 和 Tsigas 的文章的第 3.1 节。

作为最初的 Quicksort,GPU-Quicksort 以递归方式围绕主元分割元素序列,直至整个序列变成有序序列。 由于该算法面向 GPU 而编写,它包含两个阶段,在第一阶段,多个工作组在同一序列的不同部分运行,直至结果变得足够小,能够在第二阶段中被每个工作组完全排序。

GPU-Quicksort 主要理念是将输入序列分割成区块,将它们分配至不同的工作组,每个工作组负责围绕主元通过双行程进程来分割输入序列:第一步,工作组中的每个工作项会计算低于和高于它在区块中发现的主元的元素数量。 工作组扫描添加内置功能用于计算低于和高于主元的元素的累积总数。 该算法会使用一个辅助数组为每个工作组分配空间。 第二步,数据会将低于和高于主元的元素写入分配的空间内。 最后,最后的工作组会使用主元值填补空缺。

Partitioning Sequence GPU Quicksort
图 2.  GPU-Quicksort分割序列


OpenCL 1.2 中的 GPU-Quicksort

OpenCL 1.2 中的 GPU-Quicksort 是 Cederman/Tsigas 文章所描述算法直接实施。 它包括三个部分,分别是 GPUQSort 函数实施的 CPU 部分, gqsort_kernel OpenCL kernel 函数实施的 GPU First Phase Kernel,以及 lqsort_kernel OpenCL kernel 函数实施的 GPU Second Phase Kernel。 GPUQSort 函数会反复启动 gqsort_kernel,直至初始序列被分割成足够小的区块,使每个区块能被一个工作组排序。 在这一点上,GPUQSort 会启动 lqsort_kernel 以完成排序任务。

GPU-Quicksort 实施需要使用 OpenCL 1.2 中的障碍函数和原子函数,能够运行 OpenCL 1.2 的所有现代化硬件支持。 具体来说,atomic_add、atomic_sub 和 atomic_dec 被用于实施 gqsort_kernel。 障碍函数被广泛用于实施 gqsort_kernel 和 lqsort_kernel。

使用 OpenCL 1.2 的弊端有:

  1. 缺少工作组扫描基元,这要求使用 Guy Blelloch 在一篇著名文章中描述的算法来实施前缀求和。
  2. gqsort kernel 启动和 lqsort kernel 最终启动需要在 CPU 和 GPU 之间频繁来回
  3. 在 CPU 上启动 kernel 之前需要执行管理操作,相关数据在 GPU 上生成和提供。

这三个弊端都可通过 OpenCL 2.0 转换来解决。 您会发现,通过解决这些弊端,该算法的性能已大幅提升。


将 GPU-Quicksort 转为 OpenCL 2.0

在本节中,我们将回顾将 OpenCL 1.2 实施转变为 OpenCL 2.0 实施需要做出的变动。

工作组函数可提升 OpenCL C 中代码的性能和可读性
将 GPU-Quicksort 转为 OpenCL 2.0 的第一步也是最简单的一步,是利用当前可用的工作组扫描函数:具体来说,是在 gqsort_kernellqsort_kernel 中使用 work_group_scan_inclusive_addwork_group_scan_exclusive_add 函数。 我们不仅能获得一些明显的性能优势(约 8%),还能提升代码简洁性、可维护性和清晰度,从而前缀求和计算相关的代码大小。

区块支持我们通过设备端排队利用嵌套并行化
接下来,我们把在每次 gqsort_kernel 运行后记录排序需要进一步细分、区块和上代记录gqsort_kernellqsort_kernel 的逻辑从 C++ OpenCL C。lqsort_kernel 将在 GPU-Quicksort 运行结束时仅启动一次。 在第一次实施中,该逻辑位于 gqsort_kernel 的末尾,这需要额外的全局变量和原子函数,并使 gqsort_kernel 变得相对较慢,尽管整个示例的性能相对于 OpenCL 1.2 版本有所提升。 我们将该逻辑移动到了一个单独的 relauncher_kernel,后者现在作为第一个 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)。 将该 relauncher 逻辑与 qsort_kernel 隔离会显著地简化 qsort_kernel,并具有提升算法性能的附加优势。

关于我们排列 relauncher_kernel 的方式的几个要点:我们使用了 OpenCL 2.0 中的一项新特性。 由于这些 kernel 不需要本地内存参数,我们使用了 void (^) (void) 类型的区块。 这些区块类似于 C++ lambdas,但拥有不同的句法。 这些区块会捕捉传递给它们的参数,因此可避免通过 clSetKernelArg 等价物为 kernel 设置参数的辛苦工作。 第二个值得一提的要点是 num_workgroups 在区块外计算然后传递给区块的方式。 请注意,在代码中,num_workgroups将捕捉,这正是我们所需要的。 如果我们直接使用了 get_num_groups(0) 而非 num_workgroups,它会在包含 relauncher_kernel 的区块被排列和执行后被调用,因此传递给 relauncher_kernel 的数值将会是 1,而不是 gqsort_kernel 工作组的数字。 如遇了解关于 OpenCL C 所使用的区块句法的更多信息,请参见 OpenCL C 规范 [3]。

relauncher_kernel 会根据当前是否有任务来针对要执行的 gqsort_kernel 启动 gqsort_kernellqsort_kernel (work_size != 0):

       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 后,后续的 kernel 启动都会通过 GPU 来执行。 在对所有输入数据进行完全排序后,控制权将回到 CPU 手中。

GPU-Quicksort kernel enqueueing sequence in OpenCL 2.0

图 3. OpenCL 2.0 中的 GPU-Quicksort kernel 排队序列


教程要求

构建和运行本教程需要有一台满足以下要求的 PC:

  • 英特尔® 酷睿™ 处理器系列(代号 Broadwell)处理器
  • Microsoft Windows* 8 或 8.1
  • 面向 OpenCL™ 应用的英特尔® SDK 2014 R2 或更高版本
  • Microsoft Visual Studio* 2012 或更高版本

运行教程

本教程是一个控制台应用,可生成由随机的无符号整数组成的“宽*高”尺寸数组。 之后它会使用 std::sort(常规的单线程 Quicksort 算法)对该数组的拷贝进行排序,再使用 OpenCL 2.0 中的 GPU-Quicksort 大量迭代。  本教程支持以下几个命令行选项:

Option Description
-h, -? 显示帮助文本并退出
[num test iterations] 在同一数据上运行 GPU-Quicksort 的次数
[cpu|gpu] 是否使用 CPU 或 GPU OpenCL 运行本教程
[intel|amd|nvidia] 选择您想在其 OpenCL 设备上运行的厂商
[Width] 输入的“宽度” – 支持更轻松地输入较大的数字
[Height] 输入的“高度” – 支持更轻松地输入较大的数字,例如,我们可以提供 8192 8192 而不是 67M 元素
[show_CL|no_show_CL] 是否显示详细的 OpenCL 设备信息

尝试使用参数 5 gpu intel 8192 8192 show_CL运行本教程。


总结

带有英特尔® 核芯显卡 5300 或更高版本的英特尔处理器是一个复杂的硬件,但与支持 OpenCL 2.0 的驱动程序一起使用,可显著提升您 OpenCL 代码的性能。 OpenCL 2.0 为 GPGPU 程序员提供了众多强大特性。 我们只讲到了其中两个特性: enqueue_kernel 函数及 work_group_scan_exclusive_addwork_group_scan_inclusive_add 函数,并表明了它们能以相对较低的成本带来显著的性能提升,同时简化和提升您代码的可读性。 建议在面向 OpenCL™ 应用的英特尔® SDK 支持网站阅读关于其他 OpenCL 特性的其他英特尔教程。


参考资料

  1. 面向 OpenCL 应用的英特尔 SDK – 优化指南
  2. Khronos OpenCL 2.0 API 规范
  3. Khronos OpenCL 2.0 API 语言规范
  4. 维基百科关于 Quicksort 的文章
  5. 《My Early Days at Elliots》,作者 Tony Hoare
  6. 《Scan Primitives for GPU Computing》,作者 Shubhabrata Sengupta、Mark Harris、Yao Zhang 和 John D. Owens, Graphics Hardware 2007,第 97-106 页,2007 年 8 月。
  7. 《GPU-Quicksort: A practical Quicksort algorithm for graphics processors》,作者 Daniel Cederman 和 Philippas Tsigas,Journal of Experimental Algorithmics,第 14 卷,2009 年,第 14 篇文章
  8. 《Prefix Sums and Their Applications》,作者 Guy Blelloch
  9. enqueue_kernel 函数在线文档。
  10. work_group_scan_inclusive 函数在线文档
  11. work_group_scan_exclusive 函数在线文档
  12. 面向 OpenCL™ 应用的英特尔® SDK
  13. 面向 OpenCL™ 应用的英特尔® SDK R2 优化指南
  14. 面向优化的 OpenCL 应用的完整核对清单
  15. 使用英特尔® OpenCL SDK 编写最佳 OpenCL 代码 - pdf 格式

关于作者

Robert Ioffe 是英特尔软件及解决方案事业部的技术咨询工程师。 他是在英特尔锐炬和英特尔锐炬 Pro 显卡上进行 OpenCL 编程和 OpenCL 工作负载优化的专家,拥有丰富的英特尔显卡硬件经验。 他积极参与 Khronos 标准工作,关注于构建最新的特性原型并确保它们在英特尔架构上出色运行。 最近,他一直于构建 OpenCL 2.0 的嵌套并行化 (enqueue_kernel functions) 函数的原型,并编写了大量示例来演示嵌套并行化功能,包括面向 OpenCL 2.0 的 GPU-Quicksort。 他还记录和发布了两段优化简单的 OpenCL Kernel 的视频,目前正在关于嵌套并行化的第三段视频。

您可能还会对以下内容感兴趣:

OpenCL 2.0 中的 Sierpiński Carpet

优化简单的 OpenCL Kernel: 调节 Kernel 优化

优化简单的 OpenCL Kernel: Sobel Kernel 优化

 

下载代码

产品和性能信息

1

英特尔的编译器针对非英特尔微处理器的优化程度可能与英特尔微处理器相同(或不同)。这些优化包括 SSE2、SSE3 和 SSSE3 指令集和其他优化。对于在非英特尔制造的微处理器上进行的优化,英特尔不对相应的可用性、功能或有效性提供担保。该产品中依赖于微处理器的优化仅适用于英特尔微处理器。某些非特定于英特尔微架构的优化保留用于英特尔微处理器。关于此通知涵盖的特定指令集的更多信息,请参阅适用产品的用户指南和参考指南。

通知版本 #20110804