C++ TBB 并行编程教程（四）

原文：C++ Parallel Programming With Threading Building Blocks

协议：CC BY-NC-SA 4.0

十八、使用异步节点增强流程图

早在 2005 年，Herb Sutter 写了“免费的午餐结束了” ¹ 的论文来警告我们多核时代的到来及其对软件开发的影响。在多核时代，关心性能的开发人员再也不能坐以待毙，懒洋洋地等待下一代处理器，以便高兴地看到他们的应用运行得更快。那些日子早已过去。Herb 的意思是，希望充分利用现代处理器的开发人员必须接受并行技术。在这本书的这一点上，我们当然知道这一点，那又怎么样呢？嗯，我们认为今天“午餐越来越贵了。”我们来详细阐述一下这个。

近年来，在能源限制的强烈推动下，更复杂的处理器已经出现。如今，不难发现包含一个或多个 GPU、FPGA 或 DSP 以及一个或多个多核 CPU 的异构系统。就像我们采用并行来充分利用所有 CPU 内核一样，现在将部分计算卸载到这些加速器上也是有意义的。但是，嘿，这太难了！是的，它是！如果顺序编程曾经是“免费的午餐”，那么今天的异构并行编程更像是三星级米其林餐厅的盛宴——我们必须付费，但它太棒了！

TBB 有助于节省一些晚餐的价格吗？当然可以！你怎么敢怀疑？在本书的这一章和下一章中，我们将介绍 TBB 图书馆最近整合的功能，以帮助人们再次负担得起午餐——我们将展示如何将计算卸载到异步设备，从而拥抱异构计算。在这一章中，我们将采用 TBB 流图接口，并用一种新的节点类型来加强它:async_node。在下一章中，我们将更进一步，将流图放在 OpenCL Steroids 上。

异步世界示例

让我们从使用async_node的最简单的例子开始。我们将说明为什么这个特定的流图节点是有用的，我们还将给出一个对下一章有用的更复杂的例子。

因为没有比“Hello World”代码片段更简单的了，所以我们提出了一个基于流图 API 的“Async World”替代方案，它在图中包含了一个async_node。如果您对 TBB 的流程图有疑问，您可能希望通读第三章以获得可靠的背景信息，并使用附录 B 中的“流程图”部分作为 API 的参考。我们在第一个例子中建立的流程图如图 18-1 所示。

图 18-1

“异步世界”示例的流程图

我们的目标是从一个source_node、in_node向一个asynchronous node、a_node发送一条消息，但是这个任务不是在a_node内部处理消息，而是被卸载到一个正在某个地方运行的异步活动(一个不同的 MPI 节点，一个 OpenCL 支持的 GPU，一个 FPGA，你能想到的)。一旦这个异步任务完成，流图引擎必须取回控制，读取异步活动的输出，并将消息传播到图中的后代节点。在我们非常简单的"Async World"例子中，in_node只是打印出"Async，并将a=10传递给a_node。a_node接收a=10作为input并将其转发给AsyncActivity。在这个例子中，AsyncActivity是一个类，它只是增加输入消息并输出“World!”。这两个动作在一个新的线程中执行，该线程在这里模拟一个异步操作或设备。只有当AsyncActivity设计用output=11响应时，out_node才会收到该值，程序结束。

图 18-2 中的代码包含了async_world()函数定义，我们在其中构建了由图 18-1 的三个节点组成的图g。

图 18-2

构建“Async World”示例的流程图

在附录 B 的图 B-37 的表格的第一个条目中描述了source_node接口。在我们的例子中，我们创建了source_node类型的in_node。lambda 的参数int& a实际上是将被发送到图中它的后继节点async_node的输出消息。当源节点在接近async_world()函数结束时被激活，通过使用in_node.activate()，lambda 将只被执行一次，因为它只为第一次调用返回true(最初n=false，n在 lambda 中被设置为true，只有在n=true时才返回 true)。在这个调用中，带有a=10的消息被发送到图中的下一个节点。in_node的最后一个参数是false，以便源节点在休眠模式下创建，并且仅在调用in_node.activate()后唤醒(否则，节点在输出边沿连接后立即开始发送消息)。

接下来是async_node定义。async_node接口所需的语法是

在我们的例子中，a_node是在这里构造的:

这在图g中创建了一个具有unlimited并发性的async_node<int, int>。通过使用unlimited，我们指示库在消息到达时立即生成一个任务，而不管已经生成了多少其他任务。如果我们只想同时调用a_node的4，我们可以将unlimited改为4。模板参数<int, int>指出类型int的消息进入a_node，类型int的消息离开a_node。a_node构造器中使用的 lambda 如下:

它通过引用捕获一个AsyncActivity对象asyncAct，并声明对于到达a_node的每个消息必须运行的仿函数。这个仿函数有两个参数，input和gateway，通过引用传递。但是等等，我们不是说过模板参数<int, int>意味着节点期望一个传入的整数并发出一个传出的整数吗？仿子的原型不应该是(const int& input) -> int吗？嗯，对于普通的function_node来说应该是这样，但是我们现在面对的是async_node和它的特殊性。这里，我们得到了预期的const int& input，但是还有第二个输入参数gateway_t& gateway，它作为一个接口将AsyncActivity的输出注入到图中。我们在讲解AsyncActivity类的时候会讲到这一招。现在，为了完成对这个节点的描述，让我们假设它基本上用asyncAct.run(input, gateway)调度AsyncActivity。

输出节点out_node是一个function_node，它在本例中被配置为不发送任何输出消息的端节点:

该节点接收来自AsyncActivity到gateway的整数，并完成打印“Bye!，后跟该整数的值。

在图 18-2 中Async World示例的最后几行，我们发现两个make_edge调用创建了图 18-1 中描述的连接，最后该图被in_node.activate()唤醒，立即等待，直到所有消息都被g.wait_for_all()处理完毕。

接下来是AsyncActivity类，它实现了我们例子中的异步计算，如图 18-3 所示。

图 18-3

异步活动的实现

公共成员函数“run”(在a_node的带有asyncAct.run的仿函数中调用)首先执行gateway.reserve_wait()，通知流程图工作已经提交给外部活动，因此在async_world()结束时g.wait_for_all()可以考虑到这一点。然后，产生一个异步线程来执行 lambda，它通过引用捕获gateway,通过值捕获input整数。通过值传递input很关键，因为否则引用的变量source_node中的a可能会在线程读取其值之前被破坏(如果source_node在asyncThread可以读取a的值之前结束)。

线程构造器中的 lambda 首先打印“World”消息，然后分配output=11 ( input+1，更准确地说)。这个输出通过调用成员函数gateway.try_put(output)传递回流程图。最后，通过gateway.release_wait()，我们通知流程图，就AsyncActivity而言，无需再等待。

注意

不需要为提交给外部活动的每个输入消息调用成员函数reserve_wait()。唯一的要求是每个对reserve_wait()的调用必须有一个对release_wait()的相应调用。请注意，当有一些reserve_wait()调用不匹配release_wait()时，wait_for_all()不会退出

结果代码的输出是

Async World! Input: 10
Bye! Received: 11

其中“Async由in_node写，“World! Input: 10由异步任务写，最后一行由out_node写。

为什么以及何时async_node？

现在，可能会有读者表现出自负的傻笑，并认为“我不需要一个async_node来实现它。”为什么我们不依靠好的 ol’ function_node？

例如，a_node可以如图 18-4 所示实现，这里我们使用一个function_node接收一个整数input，并返回另一个整数output。相应的 lambda 表达式生成一个线程asyncThread，它打印并生成output值，然后等待线程完成asyncThread.join()并愉快地返回output。

图 18-4

创建并等待异步线程的最简单的实现。有人说危险吗？

如果你以前不是那种傻笑的读者，那现在呢？因为，这个简单得多的实现有什么问题？为什么不依靠同样的方法将计算卸载到 GPU 或 FPGA，然后等待加速器完成它的任务呢？

要回答这些问题，我们必须回到 TBB 设计的一个基本标准，即可组合性要求。TBB 是一个可组合的库，因为如果开发人员决定或需要在其他并行模式中嵌套并行模式，无论嵌套了多少层，性能都不会受到影响。使 TBB 成为可组合的因素之一是，添加嵌套的并行级别不会增加工作线程的数量。这反过来又避免了超额认购及其相关的开销破坏我们的性能。为了充分利用硬件，TBB 通常被配置为运行与逻辑核心一样多的工作线程。各种 TBB 算法(嵌套或非嵌套)只添加足够的用户级轻量级任务来支持这些工作线程，从而利用内核。然而，正如我们在第五章中所警告的，在用户级任务中调用阻塞函数不仅会阻塞该任务，还会阻塞处理该任务的操作系统管理的工作线程。在这种不幸的情况下，如果我们每个内核都有一个工作线程，并且其中一个线程被阻塞，那么相应的内核可能会空闲。在这种情况下，我们将无法充分利用硬件！

在图 18-4 的简单例子中，asyncThread在运行流程图控制之外的任务时将使用空闲内核。但是把工作卸载到加速器(GPU/FPGA/DSP，随你挑！)，还等什么？如果一个 TBB 任务调用 OpenCL、CUDA 或 Thrust 代码(仅举几个例子)中的阻塞函数，运行这个任务的 TBB 工人将不可避免地阻塞。

在async_node出现在节点的流程图列表中之前，一个可能的，尽管不理想的解决方法是用一个额外的线程超额订阅系统。为了实现这一点(如第十一章中更详细的描述)，我们通常依赖于以下几行:

如果我们在代码中不需要流程图，只想将工作从parallel_invoke或parallel_pipeline的某个阶段转移到加速器，那么这个解决方案仍然是可行的。这里需要注意的是，我们应该知道，在等待加速器的大部分时间里，额外的线程都会被阻塞。然而，这种变通办法的缺点是，系统会在一段时间内超额订阅(在卸载操作之前和之后，或者甚至在加速器驱动程序决定阻止 ² 线程时)。

为了避免这个问题，async_node来拯救我们。当async_node任务(通常是它的 lambda)完成时，负责该任务的工作线程切换到流程图的其他未决任务上。这样，工作线程不会阻塞，留下一个空闲的内核。需要记住的关键是，在async_node任务完成之前，流程图应该被警告一个异步任务正在运行(使用gateway.reserve_wait())，并且在异步任务将其结果重新注入流程图之后(使用try_put())，我们应该通知异步任务已经在gateway.release_wait()完成。还傻笑？如果有，请告诉我们原因。

更现实的例子

众所周知的流基准测试 ³ 的三元组函数是一个基本的数组操作，也称为“链接三元组”，它主要计算C = A + α ∗B，其中A、B和C是 1D 数组。因此，它非常类似于实现A=A+ α ∗B的 BLAS 1 saxpy操作，但是将结果写入不同的向量。图示上，图 18-5 有助于理解该操作。

图 18-5

计算C = A +α∗B(ci= ai+α∗bi,【∀】的三元向量运算

在我们的实现中，我们将假设数组大小由变量vsize决定，并且三个数组存储单精度浮点数。在这本书的这一点上，提出这种令人尴尬的并行算法的并行实现对我们来说还不够有挑战性。让我们来看一个异构实现。

好吧，那么你有一个集成的图形处理器？那没给我留下太多印象！ ⁴ 据报道，超过 95%的出货处理器都带有集成 GPU，与多核 CPU 共享芯片。在一个 CPU 内核上运行 triad 代码后，您会睡得很香吗？不完全是，对吗？CPU 核心不应该闲置。同理，GPU 核心也不应该闲置。在许多情况下，我们可以利用出色的 GPU 计算能力来进一步加快我们的一些应用程序。

在图 18-6 中，我们展示了三元组计算将在不同计算设备之间分配的方式。

图 18-6

三元组计算的异构实现

在我们的实现中，我们将依赖于offload_ratio变量，它控制卸载到 GPU 的迭代空间的一部分，而其余部分在 CPU 上并行处理。0 ≤ offload_ratio ≤ 1不言而喻。

代码将基于图 18-7 所示的流程图。第一个节点in_node是一个source_node，它向a_node和cpu_node发送相同的offload_ratio。前者是一个async_node，它将数组的相应子区域的计算卸载到支持 OpenCL 的 GPU 上。后者是一个常规的function_node,它嵌套了一个 TBB parallel_for,用于在可用的 CPU 内核之间分割分配给阵列的子区域。GPU 上的执行时间Gtime和 CPU 上的执行时间Ctime都被收集在相应的节点中，并被转换成join_node中的一个元组。最后，在out_node中，打印这些时间，并且将数组 C 的异构计算版本与三元组循环的普通串行执行获得的黄金版本进行比较。

图 18-7

实现异构三元组的流程图

注意

我们喜欢温和地引入新的概念，我们试图遵循这一点，尤其是当涉及到 TBB 内容时。然而，OpenCL 超出了本书的范围，所以我们不得不放弃我们自己的规则，仅仅简单地评论一下在下面的例子中使用的 OpenCL 结构。

为了简单起见，在本例中，我们将接受以下假设:

1. 为了利用零拷贝缓冲策略来减少设备间数据移动的开销，我们假设有一个 OpenCL 1.2 驱动程序可用，并且有一个 CPU 和 GPU 都可见的公共内存区域。这通常是集成 GPU 的情况。对于最近的异构芯片，OpenCL 2.0 也是可用的，在这种情况下，我们可以利用 SVM(共享虚拟内存)，我们也将在接下来说明。

2. 为了减少流图节点的参数数量，从而提高代码的可读性，指向三个数组A、B和C的 CPU 和 GPU 视图的指针是全局可见的。变量vsize也是全局的。

3. 为了跳过与 TBB 不太相关的方面，所有的 OpenCL 样板文件都被封装到一个函数opencl_initialize()中。该函数负责获取平台platform，选择 GPU 设备device，创建 GPU 上下文context和命令队列queue，读取 OpenCL 内核的源代码，编译它以创建内核，并初始化存储数组A、B和C的 GPU 视图的三个缓冲区。由于AsyncActivity也需要命令队列和程序处理程序，相应的变量queue和program也是全局变量。我们利用了 OpenCL C API 可用的 C++ 包装器。更准确地说，我们使用了可以在 https://github.com/KhronosGroup/OpenCL-CLHPP/ 上找到的cl2.hpp OpenCL C++ 头文件。

先说代码的主要功能；在图 18-8 中，我们只展示了两个第一节点的定义:in_node和cpu_node。

图 18-8

具有前两个节点的异构三元组计算的主要功能

我们首先读取程序参数并初始化调用opencl_initialize()的 OpenCL 样板文件。从这个函数中，我们只需要知道它初始化了一个 GPU 命令队列queue，和一个 OpenCL 程序program。线程数量的初始化以及初始化一个global_control对象的原因将在本节的最后进行说明。GPU 内核的源代码非常简单:

这实现了三元运算，C = A + α ∗B，假设α =0.5，并且浮点数组存储在全局内存中。在内核启动时，我们必须指定 GPU 将遍历的迭代范围，GPU 内部调度程序将使用指令i=get_global_id(0)从该空间中选取单次迭代。对于这些i中的每一个，计算C[i] = A[i] + alpha ∗ B[i]将在 GPU 的不同计算单元中并行进行。

在opencl_initialize()函数中，我们还分配了三个 OpenCL 缓冲区和从 CPU 端指向相同缓冲区的相应 CPU 指针(我们称之为数组的 CPU 视图)。假设我们有 OpenCL 1.2，对于输入数组 A，我们依靠 OpenCL cl::Buffer构造器来分配一个叫做Adevice的 GPU 可访问数组:

标志CL_MEM_ALLOC_HOST_PTR是利用零拷贝缓冲区 OpenCL 特性的关键，因为它强制分配主机可访问的内存。同样的调用用于数组的另外两个 GPU 视图，Bdevice和Cdevice。为了获得指向这些缓冲区的 CPU 视图的指针，OpenCL enqueueMapBuffer是可用的，其用法如下:

这为我们提供了一个浮点指针【the CPU 可以使用它在同一个内存区域中进行读写操作。指针Bhost和Chost也需要类似的调用。在具有集成 GPU 的现代处理器中，这种调用并不意味着数据复制开销，因此这种策略被称为零复制缓冲区。关于 OpenCL 还有其他一些微妙之处，比如clEnqueueUnmapMemObject()的含义和功能，以及在同一阵列的不同区域同时写入 CPU 和 GPU 所带来的潜在问题，但这些都超出了本书的范围。

注意

如果您的设备支持 OpenCL 2.0，实现起来会更容易，尤其是如果异构芯片实现了所谓的细粒度缓冲 SVM。在这种情况下，有可能分配一个不仅对 CPU 和 GPU 可见，而且可以同时更新并由底层硬件保持一致的内存区域。为了检查 OpenCL 2.0 和细粒度缓冲 SVM 是否可用，我们需要使用:device.getInfo<CL_DEVICE_SVM_CAPABILITIES>();

为了利用这个特性，在opencl_initialize()中，我们可以使用cl::SVMAllocator()并将其作为std::vector构造器的分配器模板参数传递。这将为我们提供一个std::vector A，即同时显示数据的 GPU 视图和 CPU 视图:

这就是，再也不需要Ahost和Adevice了。只是A。与任何共享数据一样，我们有责任避免数据竞争。在我们的示例中，这很容易，因为 GPU 在数组C的一个区域中写入，该区域与 CPU 写入的区域不重叠。如果这个条件不满足，在某些情况下，解决方案是求助于原子数组。这种解决方案通常被称为平台原子或系统原子，因为它们可以由 CPU 和 GPU 自动更新。这个特性是可选实现的，它要求我们用cl::SVMTraitAtomic<>实例化SVMAllocator。

图 18-8 中的下一件事是图g的声明和source_node、in_node的定义，它与图 18-2 中解释的非常相似，唯一的区别是它传递一个值为offload_ratio的消息。

我们示例中的下一个节点是一个function_node、cpu_node，它接收一个float(实际上是offload_ratio)并发送一个double(进行 CPU 计算所需的时间)。在cpu_node lambda 中，调用了一个parallel_for，它的第一个参数是一个阻塞范围，如下所示:

这意味着只有数组的上部会被遍历。这个parallel_for的 lambda 为不同的迭代块并行计算Chost[i] = Ahost[i] + alpha ∗ Bhost[i]，其中范围被自动划分。

我们可以继续图 18-9 中的下一个节点a_node，这是一个异步节点，它接收一个浮点值(同样是offload_ratio值)并发送 GPU 计算所需的时间。这是在a_node的 lambda 中异步完成的，其中AsyncActivity对象asyncAct的成员函数run被调用，类似于我们已经在图 18-2 中看到的。

图 18-9

具有最后三个节点定义的异构三元组计算的主要功能

join_node不值得我们在这里浪费时间，因为它已经在第三章中讨论过了。可以说，它将一个包含 GPU 时间和 CPU 时间的元组转发到下一个节点。

最后一个节点是一个function_node，out_node，它接收带有时间的元组。在打印它们之前，它检查产生的C数组是否部分在 CPU 上、部分在 GPU 上被正确计算。为此，分配C、CGold的黄金版本，然后使用 STL 算法transform进行串行计算。然后，如果Chost和CGold重合，我们就都定好了。STL 算法可以方便地实现这种比较。

图 18-10 通过节点连接完成main()功能，这得益于五个make_edge调用，随后是in_node激活以触发图的执行。我们用g.wait_for_all()等待完成。

图 18-10

三元组主函数的最后一部分，在这里连接节点并调度图形

最后，在图 18-11 中，我们展示了AsyncActivity类的实现，它的运行成员函数是从async_node调用的。

图 18-11

AsyncActivity实现，实际的 GPU 内核调用发生在这里

我们没有像在图 18-3 的AsyncActivity中那样生成一个线程，而是遵循一个更精细、更有效的替代方案。请记住，我们推迟了对为什么在图 18-8 中使用global_control对象的解释。在此图中，我们初始化了调度程序，如下所示:

如果您还记得第十一章中的内容，那么task_scheduler_init行将产生以下结果:

• 将创建一个带有nth槽的默认竞技场(其中一个槽是为主线程保留的)。

• 工作线程将被填充到全局线程池中，一旦该领域中有工作等待处理，全局线程池将占用该领域的工作线程槽。

但是后来，global_control对象，gc被构造，使得全局线程池中的实际工作线程数增加。这个额外的线程在默认的竞技场中没有空位，所以它将被休眠。

现在，AsyncActivity类，不是像我们以前那样产生一个新线程，而是唤醒休眠线程，这通常更快，特别是如果我们调用几次AsyncActivity。为此，该类的构造器初始化了一个新的 arena，a = tbb::task_arena{1,0}，它有一个工作线程槽，因为它为主线程保留了 0 个槽。当成员函数run()被调用时，一个新任务与a.enqueue()一起在这个竞技场中排队。这将导致休眠线程的分派，该线程将占据这个新竞技场的槽位并完成任务。

接下来，这个AsyncActivity中产生的任务按照通常的步骤将计算卸载到 GPU。首先，构造triad_kernel KernelFunctor，告知triad kernel有三个cl::Buffer参数。第二，调用triad_kernel通过NDRange，计算为ceil(vsize∗offload_ratio)，以及缓冲区的 GPU 视图Adevice、Bdevice、Cdevice。

在集成 GPU 的英特尔处理器上运行这段代码时，会生成以下两行代码:

Time cpu: 0.132203 sec.
Time gpu: 0.130705 sec.

其中vsize设置为 1 亿个元素，我们一直在玩offload_ratio,直到两个设备在计算分配给它们的数组子区域时消耗大约相同的时间。

摘要

在这一章中，我们首先介绍了async_node类，它增强了流程图的功能，可以处理脱离流程图控制的异步任务。在第一个简单的Async世界的例子中，我们展示了这个类和它的伙伴gateway接口的使用，这对于将来自异步任务的消息重新注入流图是有用的。然后，我们激发了这个扩展与 TBB 流图的相关性，如果我们认识到阻塞 TBB 任务会导致阻塞 TBB 工作线程，这就很容易理解了。async_node允许在流程图之外分派异步工作，但在等待异步工作完成时不会阻塞 TBB 工作线程。我们用一个更现实的例子结束了这一章，这个例子让async_node将parallel_for的一些迭代卸载到 GPU 上。我们希望我们已经提供了详细阐述更复杂的项目的基础，其中涉及到异步工作。然而，如果我们通常的目标是支持 OpenCL 的 GPU，我们有好消息:在下一章，我们将介绍 TBB 的opencl_node特性，它提供了一个更友好的界面来让 GPU 为我们工作！

更多信息

以下是我们推荐的一些与本章相关的额外阅读材料:

• 赫伯·萨特(Herb Sutter)，《免费的午餐结束了:软件并发性的根本转变》， www.gotw.ca/publications/concurrency-ddj.htm 。

• 约翰·麦卡尔平， www.cs.virginia.edu/stream/ref.html 。

• 大卫凯利，佩哈德米斯特里，达纳沙，张东平。使用 OpenCL 2.0 进行异构计算。摩根·考夫曼 2015。

开放存取本章根据知识共享署名-非商业-非专用 4.0 国际许可协议(http://Creative Commons . org/licenses/by-NC-nd/4.0/)的条款进行许可，该协议允许以任何媒体或格式进行任何非商业使用、共享、分发和复制，只要您适当注明原作者和来源，提供知识共享许可协议的链接，并指出您是否修改了许可材料。根据本许可证，您无权共享从本章或其部分内容派生的改编材料。

本章中的图像或其他第三方材料包含在本章的知识共享许可中，除非在材料的信用额度中另有说明。如果材料不包括在本章的知识共享许可中，并且您的预期使用不被法律法规允许或超出了允许的使用范围，您将需要直接从版权所有者处获得许可。

Footnotes 1

“免费的午餐结束了:软件并发性的根本转变，”Herb Sutter。 www.gotw.ca/publications/concurrency-ddj.htm 。

2

当线程使用阻塞调用将内核卸载到 GPU 时，驱动程序可能不会立即阻塞调用线程。例如，一些 GPU 驱动程序保持线程旋转，以便它可以更早地响应轻量级内核，并在一段时间后最终阻止线程，以避免在重量级内核完成时消耗资源。

3

约翰·麦卡尔平， www.cs.virginia.edu/stream/ref.html 。

4

仙妮娅·唐恩——那没给我留下太多印象，来看看 1997 年的专辑。

十九、Steroids 上的流程图：penCL 节点

async_node是否让你渴望更多？如果是这样，这就是你的篇章。在这里，我们将介绍一个高级流程图类opencl_node，它试图隐藏 OpenCL 设备的硬件细节和编程细节。为什么选择 OpenCL？原因有很多，仅举几个例子:OpenCL 是由一个大型联盟的成员贡献的开放标准，它被设计成一个独立于平台的 API，它的目标是灵活地发展以满足更新的需求。例如，OpenCL 一直是 C(不是 C++)的扩展，但是最新的 OpenCL 2.2 版本增加了对 C++14 子集的支持，包括类、lambda 表达式、模板等等。

这还不够吗？好吧，再来一个。对我们来说，所有原因中最突出的是可以使用 OpenCL 的平台的数量和种类。从笔记本电脑和台式机开始，这些系统中超过 95%的处理器都包含支持 OpenCL 的集成 GPU(通常来自英特尔或 AMD)。在移动领域，在大多数智能手机和平板电脑的核心，我们发现了一个片上系统，SoC，具有支持 OpenCL 的 GPU(是的，从 TBB 仓库我们也可以获得 Android 的 TBB 二进制文件)。这些例子看起来已经足够有说服力了，但还有更多！在嵌入式领域，多年来，我们已经能够购买和开发异构板，包括 OpenCL 可编程 FPGA(来自 Intel-Altera 和 Xilinx)。在服务器领域，在撰写这些文章时，英特尔的目标是采用 FPGA PCIe 卡和英特尔至强可扩展处理器 6138P 的数据中心，该处理器包括片上英特尔 Altera Arria 10 FPGA，当然，OpenCL 是受支持的编程模型之一。此外，OpenCL 代码还可以在许多 CPU 和其他类型的加速器上运行，如 Xeon Phi。

但是如果 OpenCL 不能满足您的需求，TBB 建筑师事务所也考虑了支持其他编程模型的可能性。他们将加速器编程模型的底层细节抽象成一个叫做工厂的模块。事实上，opencl_node是用特定工厂实例化一个名为streaming_node的通用类的结果。然后，工厂定义必要的方法来上传/下载数据到加速器和启动内核。也就是说，opencl_node是将streaming_node类与 OpenCL 工厂结合的结果。开发相应的工厂就可以支持更新的编程模型。

现在，这是一个相当长的章节，涵盖了几个概念(opencl_node、opencl_program、,、??)、OpenCL 内核的参数和范围、子缓冲区等。)因此意味着陡峭的学习曲线。但是我们将从简单的开始，逐步提升到更复杂的类和例子(就像我们一直试图做的那样)。正如我们在图 19-1 中所描绘的，我们将从一个简单的Hello World——类似于使用opencl_node的例子开始，随后实现与前一章相同的三元矢量计算，但是现在使用我们新的高级玩具。如果你想把最后的攀登留到顶峰，你可以在那里停止阅读。另一方面，如果你是一个有经验的攀岩者，在本章的最后，我们会先睹为快更高级的特性，比如微调 OpenCL NDRange和内核规范。

图 19-1

描绘本章的学习曲线

Hello OpenCL_Node 示例

这次让我们从结尾开始。这是我们第一个示例的输出:

Hello OpenCL_Node
Bye! Received from: OPENCL_NODE

这两行是运行图 19-2 所示的流程图的结果，其中的气泡标识了由图中三个节点中的每一个打印的字符串。

图 19-2

流程图以“Hello OpenCL_Node”为例

中间的节点gpu_node，是一个opencl_node，打印出OpenCL_Node\n。为此，它将被配置为运行存储在hello.cl文件中的以下 OpenCL 内核:

hello.cl文件包括cl_print()内核的定义，该内核将由流程图的一个特殊节点，一个opencl_node执行。如果我们仔细看看内核函数，它实际上打印了作为输入参数出现的任何字符数组。此外，为了产生明显的影响，内核还通过只大写小写字母来改变字符串。参数的char *str声明之前的global关键字声明字符数组应该存储在 OpenCL 全局内存中。对于这里的问题(即过于简化)，这意味着字符串存储在内存的一个区域中，可以由 CPU 和 GPU“以某种方式”读取和写入。在集成 GPU 的常见情况下，全局内存只是位于主内存中。这意味着opencl_node应该接收一个字符数组作为参数。在我们的例子中，这个字符数组包含的字符是"OpenCL_Node \n "。正如您可能已经猜到的，这个消息来自第一个节点in_node。对，指向字符串的指针(图 19-2 中的a)从in_node飞到gpu_node，在没有用户干预的情况下，在 CPU 上初始化的字符串最终到达 GPU。什么消息到达out_node？同样，指针a离开gpu_node并以名称m进入out_node。最后，图中的最后一个节点打印出了“Bye! Received from: OPENCL_NODE”，我们注意到了字符串的变化，也注意到了在 GPU 上处理的字符串已经可以被 CPU 访问了。现在，我们都渴望实际实现的细节，所以它们在图 19-3 中。

图 19-3

构建“Hello OpenCL_Node”示例的流程图

就这样！注意，GPU 节点配置只需要三行 C++ 代码。是不是很整洁？

放弃

在写这一章的时候，TBB 的最新版本是 2019 年。在这个版本中，opencl_node仍然是一个预览功能，这实质上意味着

• 它可能会发生变化。如果您依赖代码中的预览功能，请在更新到较新的 TBB 版本时仔细检查它是否继续工作。在最坏的情况下，预览功能甚至会消失！

• 它可能没有什么文档和支持。事实上，opencl_node和streaming_node文档在网络上并不丰富。有一些博客条目 ¹ 说明了这个特性，但是它们已经有 3 年的历史了，而且 API 的一部分也已经改变了。

• 必须明确启用它(即，默认情况下它是关闭的)。为了在我们的代码中使用opencl_node,我们必须添加这三行代码:

  

使用这个头文件的额外好处是你不需要手动包含tbb/flow_graph.h或者 OpenCL 头文件，因为它们已经包含在flow_graph_opencl_node.h中了。实际上，这个头文件和博客条目是目前我们关于这个特性提供的类和成员函数的最可靠的信息来源。这一章应该被认为是对包含在opencl_node头文件中的 1050 行代码的简单介绍。

好吧，我们一点一点来。如果你记得上一章的例子，第一个in_node看起来很熟悉。为了提醒我们，可以说:( 1)lambda(&的输入参数实际上是对将被发送到任何连接节点的消息的引用；(2)只有一条消息离开in_node,因为在第一次调用后它返回 false 以及(3) in_node.activate()实际上唤醒节点并触发该单个消息。但是等等，在这个节点中有一些新的东西是我们必须注意的！离开in_node的消息必须在 GPU 可访问的内存区域结束，这就是为什么参数a不仅仅是一个字符数组，而是对一个buffer_t的引用。就在定义in_node之前，我们看到buffer_t是 OpenCL chars ( cl_char)的一个opencl_buffer:

opencl_buffer是我们将在本章中看到的第一个opencl_node助手类，但是还有更多。它是一个模板类，抽象了强类型线性数组，封装了主机和加速器之间的内存事务逻辑。我们使用类的构造器来分配一个opencl_buffer<T>，就像我们的例子中的行a = buffer_t{sizeof(str)}，或者通过用

在这两种情况下，我们最终都会分配一个cl_char的opencl_buffer。我们现在使用的 OpenCL 工厂版本基于 OpenCL 1.2，并利用了零拷贝缓冲区方法。这意味着，在内部，当调用opencl_buffer构造器时，OpenCL 函数clCreateBuffer被调用，它的一个参数是CL_MEM_ALLOC_HOST_PTR。正如我们在前一章简单解释的那样，缓冲区是在 GPU 空间上分配的，但是 CPU 可访问的指针(缓冲区的 CPU 视图)可以使用映射函数(clEnqueueMapBuffer)获得。为了将缓冲区的控制权交还给 GPU，OpenCL 提供了一个unmap函数(clEnqueueUnmapMemObject)。在集成 GPU 的现代芯片上，map 和unmap函数很便宜，因为不需要实际的数据副本。对于这些情况，map和unmap函数负责保持 CPU 和 GPU 缓存与存储在全局内存(主内存)中的副本一致，这可能意味着也可能不意味着 CPU/GPU 缓存刷新。好消息是，所有这些低级的杂务都不关我们的事了！可以开发具有更好特性或支持其他加速器的新工厂，我们可以通过简单地重新编译我们的源代码来使用它们。考虑一下，如果明天公开一个 OpenCL 2.0 工厂，并且我们的加速器实现了细粒度的缓冲 SVM。仅仅通过使用新的 OpenCL 2.0 工厂而不是 1.2 工厂，我们将免费获得性能提升(因为现在map和unmap操作是不必要的，CPU 和 GPU 之间的缓存一致性由硬件自动保持)。

哎呀，抱歉让我们的思绪飘了一会儿。让我们回到正题。我们在图 19-3 中解释了我们例子中的source_node(是的，几段之前)。这个source_node、in_node，只是用字符串OpenCL_Node\n初始化一个chars、str的数组，分配适当大小的opencl_buffer、a，并使用std::copy_n STL 算法将字符串复制到那个缓冲区。就这样。当这个source_node的 lambda 结束时，引用opencl_buffer的消息将从in_node飞到gpu_node。

现在，记住配置gpu_node所需的行:

第一行使用了我们在本章中提到的第二个opencl_node助手类:opencl_program类。在这一行中，我们创建了program对象，并将文件名hello.cl传递给构造器，OpenCL 内核cl_print就存储在这里。如果我们想提供一个预编译的内核或者内核的 SPIR (OpenCL 中间表示)版本，还有其他的opencl_program构造器可用。为了不让人分心，并专注于我们的例子，我们将在后面讨论这些替代方法。

第二行创建了类型为opencl_node<tuple<buffer_t>>的gpu_node。这意味着gpu_node接收类型为buffer_t的消息，完成后，它发出类型为buffer_t的消息。对于单个参数/端口，我们真的需要一个元组吗？嗯，opencl_node被设计为从前面的节点接收几个消息，并向图中后面的节点发送几个消息，这些消息被打包到一个元组中。目前，接口中没有针对单个输入和输出的特殊情况，因此我们需要在这种情况下使用单个元素元组。关于opencl_node端口和内核参数之间的对应关系，默认情况下，opencl_node将第一个输入端口绑定到第一个内核参数，第二个输入端口绑定到第二个内核参数，依此类推。后面还会谈到其他的可能性。

我们真的需要为每个传入的消息发送一个传出的消息吗？嗯，opencl_node被设计成支持这种最大连接性(每个输入端口一个输出端口),如果输入少于输出，或者相反，我们总是可以保持相应的端口不连接。我们真的需要对输入和输出使用相同的数据类型吗？嗯，就目前的工厂来说，是的。如果输入端口 0 是类型T，输出端口 0 也是同样的T类型(指定参数类型的元组不区分输入和输出)。

注意

支持opencl_node实现决策的主要原因是每个opencl_node的端口都有可能被映射到每个 OpenCL 内核参数中。对于一个“输入-输出”参数，在输入和输出都有它当然是有意义的。对于一个“out”参数，我们仍然需要传入要写入的对象，因此需要一个输入来匹配输出——否则opencl_node将需要分配对象，但它没有。最后，对于一个“in”参数，让它在输出端可用可以让我们转发值，也就是说，不加修改地将它传递给下游节点。所以，最实际的事情就是把所有的论点都放进去。我们相信，如果我们将 OpenCL 节点的 tuple 视为一个参数列表，那么这是有意义的，我们可以将边连接到任何参数，以在执行之前/之后设置/获取值。对于“in”参数，相应的发出值不变。对于一个“out”参数，我们提供了要写入的内存，并在稍后获取值。对于“in-out”，我们发送值并接收修改后的值。

请记住，OpenCL 节点是一个预览功能。TBB 开发者渴望预览功能的输入——这就是为什么他们毕竟是预览功能。他们希望收集好的和坏的信息，这样他们就可以花时间完善图书馆中人们最关心的部分。这个 OpenCL 节点的预览版应该足够好，可以试用并提供反馈。如果我们对需要添加什么有强烈的意见，我们应该说出来！

现在，opencl_node的构造器包含流图对象g作为参数，以及应该包含在 OpenCL 程序文件中的内核函数的句柄。由于文件hello.cl包含内核函数cl_print，我们使用成员函数:program.get_kernel("cl_print")。

这意味着我们可以在同一个 OpenCL 源文件中有几个内核函数，并将每个函数分配给不同的opencl_nodes。我们真的必须用一个程序文件来解决吗？不完全是。如果我们将 OpenCL 内核分布在几个源文件中，我们可以实例化期望数量的opencl_program对象。

最后，配置gpu_node所需的第三行代码是gpu_node.set_range({{1}})。这个来自opencl_node的成员函数指定了 GPU 将要遍历的迭代空间。更正式地说，在 OpenCL 行话中，这个迭代空间被称为NDRange，但是我们现在不要详细讨论这些细节。现在，让我们大胆地相信,set_range({{1}})成员函数导致内核主体只被执行一次。

现在我们已经完成了source_node(in_node)opencl_node(gpu_node)，我们例子中的最后一个是一个名为out_node的常规function_node。对应的代码是

我们看到out_node收到了一条buffer_t类型的m消息。因为buffer_t实际上是一个opencl_buffer<cl_char>，所以调用m.begin()会产生一个 CPU 可见的指针，指向最初在in_node中设置的、后来被 GPU 内核修改的字符串。我们的最后一个节点只是打印这个字符串，然后死亡。

示例的其余部分是通常的流图粘合逻辑，它在节点之间形成边，唤醒源节点，并等待所有消息(在我们的示例中只有一条)通过节点。这里没什么新鲜的。

然而，在我们开始攀登我们的第一座山峰之前，我们将对我们刚刚解释的内容进行一次高级别的回顾，同时更深入地了解消息a发生了什么，该消息诞生在 CPU 上，发送到 GPU 并在那里进行修改，然后传递到最终节点，在那里我们可以看到 GPU 内核执行的效果。我们希望图 19-4 能在这方面很好地为我们服务。

图 19-4

包含消息操作细节的示例概述

图片假设 OpenCL 工厂是基于这个标准的 1.2 版本。在这种情况下，消息a作为opencl_buffer被分配在 GPU 内存空间中，但是如果我们首先使用a.begin()获得 CPU 可访问的迭代器，它也可以被写到 CPU 上。对a的引用是离开in_node并进入gpu_node的端口 0 的消息(这将总是导致消息-对a的引用-通过出发端口 0 离开)。gpu_node的端口 0 被绑定到具有兼容类型的内核函数的第一个参数(opencl_buffer<cl_char>可以被强制转换为char *)。内核可以安全地访问字符串，而不会出现缓存一致性问题，因为在启动内核之前，OpenCL 工厂会负责解除缓冲区的映射。最后，对缓冲区的引用到达out_node,，在这里字符串再次被映射，以便在 CPU 上访问和打印。

在继续之前，我们想在这里强调我们应该感到多么幸运，因为我们不必手动处理所有的 OpenCL 样板代码(平台、设备、上下文、命令队列、内核读取和编译、内核参数设置和启动、OpenCL 资源解除分配等)。).多亏了 OpenCL 工厂，所有这些现在都隐藏在引擎盖下。此外，正如我们所说的，新工厂可以使我们的代码更快，或者能够与其他加速器一起工作，只需对源代码进行很小的更改或不做任何更改。

我们在哪里运行我们的内核？

到目前为止一切顺利，对吧？但是说到 OpenCL 样板代码，控制我们在哪个设备上运行我们的opencl_nodes的旋钮在哪里呢？在我们之前的例子中，我们说过gpu_node正在 GPU 上运行指定的内核。还有哪里，对吗？但是如果我们在撒谎呢？令人不安，是吧？好的，让我们先看看我们的机器上是否有更多支持 OpenCL 的设备。希望只有一个单一的设备，它是一个 GPU，但我不会赌我的手指！我们将不得不嗅出它，但是我们在情感上还没有准备好编写旧式的普通 OpenCL 代码，不是吗？幸运的是，TBB OpenCL 工厂给了我们两个额外的有价值的助手类(现在已经有四个了)。这些是opencl_device和opencl_device_list助手类。让我们首先在流程图上下文之外使用它们，如图 19-5 所示。

图 19-5

查询 OpenCL 平台和可用设备的简单代码

首先，通过调用函数available_devices() .初始化一个opencl_device_list对象devices，该函数返回一个可迭代容器，其中包含第一平台中所有可用的 OpenCL 使能设备。是的，仅在第一个可用的平台中。 ² 然后，我们从列表中弹出第一个opencl_device、d，查询平台名称、概要文件、版本和厂商。平台中所有可用的设备将共享这些属性。

接下来，使用for(opencl_device d:devices)，我们遍历整个设备列表，获取并打印每个设备的名称、主版本和次版本以及设备类型。主版本和次版本信息已经由d.platform_version()提供，但是这个返回一个字符串，而d.major_version()和d.minor_version()都返回一个整数。在我们写这些代码的 MacBook 上运行这些代码的输出结果，以及我们运行之前例子的地方，可以在图 19-6 中看到。

注意

函数available_devices()实际上不是公共的，这就是我们必须使用这个错综复杂的名称空间链的原因:

tbb::flow::interface10::opencl_info::available_devices()

我们注意到，就在实现这个成员函数之前，在flow_graph_opencl_node.h内部有一个注释声明

// TODO: consider opencl_info namespace as public API

由于这是 TBB 的一个预览功能，界面还没有完全确定下来.考虑到这一点，以防这一考虑最终成为事实。

图 19-6

在 MacBook Pro 上运行图 19-5 的代码的结果

令人惊讶的是，一台笔记本电脑中可能有三个 OpenCL 设备！也就是说，一个英特尔 CPU 和两个 GPU，第一个集成在英特尔酷睿 i7 中，第二个是独立的 AMD GPU。请记住，OpenCL 是一种可移植的编程语言，也可以用来实现 CPU 代码。看，第一个支持 OpenCL 的设备不是 GPU，而是四核英特尔 CPU。现在，关于本章的第一个例子，内核在哪里运行？在第一点上，你是对的。默认情况下，OpenCL 工厂选择第一个可用的设备，不管它是 CPU 还是 GPU。所以…我们在撒谎！！！内核运行在伪装成 OpenCL 加速器的 CPU 上。如果我们在整本书里到处撒谎呢？想想看…那就更恐怖了(除非这是你正在读的第一章)。

好吧，我们来解决这个小麻烦。为了化险为夷，OpenCL 工厂提供了两个额外的特性:设备过滤器和设备选择器。设备过滤器用于用一组可用于内核执行的设备来初始化opencl_factory。所有过滤的设备必须属于同一个 OpenCL 平台。有一个默认的设备过滤器类default_device_filter，它自动从第一个 OpenCL 平台收集所有可用的设备，并返回一个包含这些设备的opencl_device_list。就其本身而言，设备选择器，顾名思义，选择那个opencl_device_list中的一个设备。不同的opencl_node实例可以使用不同的设备选择器。对于每个内核执行都要进行选择，所以对于不同的调用，在不同的设备上运行opencl_node也是可能的。默认选择器default_device_selector从设备过滤器构建的可用设备列表中选择并返回第一个设备。

让我们的gpu_node在真正的 GPU 上运行，而不是

我们应该使用

其中gpu_selector是我们自定义的对象class gpu_device_selector:

gpu_device_selector gpu_selector;

而这个类呈现在图 19-7 中。

图 19-7

我们的首款定制器件选择器

协议(更正式的说法是“概念”)是，opencl_node的第三个参数是一个函子(带有operator()成员函数的类的对象)，它返回一个设备。这样，我们可以在它的位置嵌入一个 lambda 表达式，而不是传递函子。operator()接收一个opencl_factory、f，并返回一个opencl_device。使用find_if STL 算法，我们在满足it->type()==CL_DEVICE_TYPE_GPU的容器devices()中返回第一个迭代器it。为了方便起见，我们声明了auto it并委托给编译器去发现it的类型实际上是

tbb::flow::opencl_device_list::const_iterator it = ...

考虑到找不到 GPU 设备的可能性，我们包括了一个返回第一个设备的回退(应该至少有一个！…没有任何设备的平台是没有意义的)。仿函数通过打印所选设备的名称并将其返回来结束。在我们的笔记本电脑中，输出将是:

注意，当该节点被激活时，新消息由gpu_node设备选择器仿函数打印出来。这是，首先in_node打印它的消息“Hello”并将消息传递给gpu_node，后者在启动内核之前选择设备(打印输出的粗体字)，然后运行内核。这是需要考虑的事情:流程图中的opencl_node通常会被激活几次，所以我们最好实现尽可能最轻的设备选择器。

比如std::find_if算法的 lambda 表达式不需要打印“找到 GPU！”消息，可以进一步简化:

现在，如果我们不喜欢必须显式添加gpu_device_selector类的源代码的样子，我们可以用 lambda 表达式代替仿函数。这有点棘手，因为这个类的operator()是一个模板化的函数，还记得吗？：

(据我们所知)实现 lambda 最简单的方法是依赖从 C++14 开始就有的多态 lambda。不要忘记用选项std=c++14编译图 19-8 中的代码。

图 19-8

使用 lambda 表达式而不是仿函数进行设备选择

注意 lambda 的(auto& f)参数，而不是我们在基于函子的替代方案中使用的(opencl_factory<DeviceFilter>& f)。这段代码遍历devices()容器，然后返回列表中的第二个设备，结果类似于

Available devices:
0.- Device: Intel(R) Core(TM) i7-6700HQ CPU @ 2.60GHz
1.- Device: Intel(R) HD Graphics 530
2.- Device: AMD Radeon Pro 450 Compute Engine
Running on Intel(R) HD Graphics 530

现在我们知道了我们的设备列表，并假设我们想要使用集成的 GPU，最好更改 lambda 以使其更快:

更快的替代方法是在我们第一次调用设备选择器时缓存opencl_device。例如，在图 19-9 中，我们绘制了图 19-7 中出现的gpu_device_selector类的修改草图。

图 19-9

第一次调用opencl_device时缓存它的设备选择器类

这个类现在有了一个opencl_device成员变量device。当第一次调用operator()时，遍历设备列表f.devices()，找到我们想要使用的设备(在本例中，是第二个可用的设备)。然后我们将它缓存到device变量中以备将来使用。请注意，如果可以从不同的线程同时调用该操作符，则需要进一步注意避免数据竞争。

我们希望你能保守我们对数字 19-8 和 19-9 的例子编码有多糟糕的秘密。在这些片段中，我们将设备硬编码为第二个设备，它可以在我们的测试机上工作，但在其他平台上可能会失败。实际上，如果有一个设备存储在f.devices()容器中，取消引用*(++f.devices().cbegin())将触发一个分段错误。这是便携性和性能之间权衡的又一个例子。如果我们不知道代码最终会在哪里运行，并且与 OpenCL 计算相比，设备选择时间可以忽略不计，那么我们最好使用图 19-7 (注释掉打印语句)的版本。

第十八章回到更现实的例子

你还记得我们在前一章介绍的三元组向量运算吗？这只是一个形式为C = A + α *B的基本数组操作，其中A、B和C是包含vsize浮点数的 1D 数组，α是一个标量，我们将其设置为 0.5(因为我们可以)。图 19-10 提醒我们三元组计算将根据变量offload_ratio在 GPU 和 CPU 之间分配的方式。

图 19-10

三元组计算的异构实现

重新实现这个例子的目的有两个。首先，通过重新访问我们的老熟人，但是现在从opencl_node的角度，我们将更好地欣赏 TBB 流图的这个更高层次特征的好处。其次，超越“Hello OpenCL_Node”将允许我们深入研究opencl_node类及其助手类的更高级的用法。在图 19-11 中，我们给出了我们将要实现的流程图的概述。

图 19-11

实现三元组的流程图，现在使用 OpenCL 节点

和我们前面的例子一样，source_node ( in_node)只是触发流程图的执行，在这个例子中，传递一条值为offload_ratio的消息。下游的下一个节点是multifunction_node (dispatch_node).,这种节点非常灵活，可以向图中的下一个节点发送消息。我们看到dispatch_node有五个输出端口，前四个针对gpu_node，最后一个连接到cpu_node。gpu_node是一个opencl_node，它将配置适当的三元组 GPU 内核，该内核期望数组A、B和C的“GPU 视图”作为输入参数(如前一章所述，它们被称为Adevice、Bdevice,和Cdevice)。然而，gpu_node有一个额外的端口来接收将要卸载的迭代次数，这取决于offload_ratio并且我们称之为NDRange来遵守 OpenCL 符号。cpu_node是一个常规的函数节点，它接收三个数组的“CPU 视图”以及offload_ratio,这样 CPU 就可以完成它的任务。cpu_node只有一个输入端口，所以dispatch_node必须将 CPU 所需的四个变量打包成一个元组。gpu_node和cpu_node都将它们自己的结果数组C的视图传递给join_node，后者又用两个视图构建一个元组，并将其转发给out_node。这个最终节点将验证计算是否正确，并打印出执行时间。事不宜迟，让我们从真正的实现开始，从图 19-12 中的数据类型定义和缓冲区分配开始。

图 19-12

三元组示例中的数据类型定义和缓冲区分配

从现在开始，buffer_f是cl_floats的opencl_buffer(OpenCL 中常规浮点数据类型的对应类型)。这样，我们将Adevice、Bdevice和Cdevice分配为我们三个数组的“GPU 视图”。opencl_buffer类还公开了data()成员函数，这是我们在这里第一次看到。该函数返回一个指向 GPU 缓冲区的 CPU 可访问指针，并负责映射缓冲区，以便 CPU 可以访问它。这允许我们初始化指针Ahost、Bhost和Chost。使用 STL generate算法，我们用 0 到 255 之间的随机数初始化数组A和B，使用 Mersenne Twister 生成器(正如我们在第五章中所做的)。

图的前两个节点in_node和dispatch_node在图 19-13 中定义。

图 19-13

三元组示例中的前两个节点in_node和dispatch_node

算法的这一部分非常简单。我们的老朋友in_node给offload_ratio=0.5发了一封短信给dispatch_node。dispatch_node属于以下类型:

这意味着它接收一个 float ( offload_ratio)并有五个输出端口发送对应于五个元组元素类型的消息。这个元组封装了这个多功能节点的五个输出端口的数据类型:三个数组的三个buffer_f(这里是opencl_buffers)、NDRange和一个为cpu_node打包所有信息的tuple_cpu。

定义dispatch_node主体的 lambda 表达式的两个输入参数是

在这里我们可以找到输入消息(offload_ratio)和一个句柄(ports)，它可以让我们访问五个输出端口中的每一个。现在，我们使用函数get<port_number>(ports).try_put(message)向相应的port_number发送消息。我们只需要对这个函数进行四次调用，就可以发送 GPU 正在等待的信息。请注意，这四个调用中的最后一个调用放置了一个只有一个元素等于ceil(vsize*offload_ratio)的 1D 数组，它对应于 GPU 上的迭代空间。使用get<4>(ports).try_put(cpu_vectors).，单个消息通过最后一个端口到达 CPU。之前，我们已经方便地将三个向量的 CPU 视图和向量分区信息(ceil(vsize*offload_ratio))打包在cpu_vectors元组中。

有什么问题吗？确定吗？我们不想落下任何读者。那好吧。让我们继续看下两个节点的实现，这是问题的核心，真正的计算发生在这里，如图 19-14 所示。

图 19-14

在三元组示例中真正肩负重任的节点:gpu_node和cpu_node

虽然cpu_node是图 19-14 中的第二个，但我们将首先介绍它，因为它不太需要澄清。模板参数<tuple_cpu, float*>指出节点接收到一个tuple_cpu并发送一个指向float的指针。lambda 输入参数cpu_vectors在主体中用于将指针解包为三个向量和变量start(获得已经在dispatch_node上计算的值ceil(vsize*offload_ratio))。利用该信息，a parallel_for在范围blocked_range<size_t>(start, vsize)中执行三元组计算，这对应于迭代空间的第二部分。

正如我们所说，GPU 负责这个迭代空间的第一部分，在这个上下文中称为NDRange=0, ceil(vsize*offload_ratio))。GPU 内核的源代码与我们在上一章中介绍的相同，它只是接收三个数组，并对NDRange中的每个i进行三元运算:

图 19-14 开始。定制型tuple_gpu包三个buffer_f和NDRange。据此，我们将gpu_node声明为

它选择程序文件的内核triad，并指定我们最喜欢的设备选择器gpu_selector。

现在出现了一个有趣的配置细节。四条消息到达gpu_node，我们之前提到过“opencl_node将第一个输入端口绑定到第一个内核参数，将第二个输入端口绑定到第二个内核参数，依此类推。”但是等等！内核只有三个参数！我们又说谎了！！？？好吧，这次不会。我们还说过这是默认行为，可以修改。以下是方法。

使用gpu_node.set_args(port_ref<0,2>)，我们声明到达端口 0、1 和 2 的消息应该绑定到内核的三个输入参数(A、B和C)。那NDRange呢？在图 19-3 中的第一个例子Hello OpenCL_Node中，我们只是使用gpu_node.set_range({{1}})来指定可能最小的NDRange常量值 1。但是在第二个更详细的例子中，NDRange是可变的，来自dispatch_node。我们可以绑定节点的第三个端口，它用set_range()函数接收NDRange，就像我们对行gpu_node.set_range(port_ref<3>)所做的那样。这意味着我们可以通过端口向set_range()传递一个常量或变量NDRange。成员函数set_args()应该支持同样的灵活性吧？我们知道如何将内核参数绑定到opencl_node端口，但是通常内核参数只需要设置一次，而不是每次调用都设置。

比方说，我们的内核接收α的值，它现在是一个用户定义的参数(不像以前那样硬连接到 0.5):

然后我们可以编写如下代码:gpu_node.set_args(port_ref<0,2>, 0.5f)，它将前三个内核参数绑定到到达端口 0、1 和 2 的数据，并将第四个参数绑定到… 0.5(哦不！又硬连线了！更严重的是，没有什么可以阻止我们传递一个变量alpha，这个变量之前被设置为…0.5).

现在，让我们来看看最后两个节点，node_join和out_node，它们在图 19-15 中有详细描述。

图 19-15

异源三元组向量运算的最后两个节点node_join和out_node

如粗体所示，node_join接收一个buffer_f(来自gpu_node)和一个指向float(来自cpu_node)的指针。创建这个节点只是为了将这两条消息连接成一个元组，该元组将被转发到下一个节点。说到这里，下一个节点是out_node，一个function_node，接收join_t::output_type类型的消息，不发送任何输出消息。注意join_t是node_join的类型，所以join_t::output_type是tuple<buffer_f, float*>.的别名实际上，lambda 的输入参数m就有这种类型。解包元组m的一种便捷方式是执行std::tie(Cdevice, Chost) = m，这完全等同于

Cdevice = std::get<0>(m);
Chost = std::get<1>(m);

out_node正文的下几行检查异构计算是否正确，首先串行计算三元数组运算的黄金版本CGold，然后使用std::equal算法与Chost进行比较。由于Chost , Cdevice.data() ,和Cdevice.begin()实际上都指向同一个缓冲区，所以这三个比较是等价的:

std::equal (Chost, Chost+vsize, CGold.begin())
std::equal (Cdevice.begin(), Cdevice.end(), CGold.begin())
std::equal (Cdevice.data(), Cdevice.data()+vsize, CGold.begin())

是时候结束我们的代码了。在图 19-16 中，我们添加了make_edge调用并触发流程图的执行。

图 19-16

三元组主函数的最后一部分，在这里连接节点并调度图形

注意，虽然gpu_node的四个输入端口连接到前面的dispatch_node，但是只有gpu_node的 2 号端口连接到node_join。这个端口承载产生的Cdevice缓冲区，所以它是我们唯一关心的端口。其他三个被忽视的端口不会觉得被冒犯。

我们花了一段时间来解释整个例子，但我们仍然需要添加一个东西。它与我们在前一章介绍的async_node版本相比如何？我们的async_node版本包含了 OpenCL 样板文件，它隐藏在OpenCL_Initialize()函数中，但却是必需的，因为它让我们可以访问上下文、命令队列和内核处理程序。如果我们使用cl.h OpenCL 头文件，这个async_node版本有 287 行代码(不包括注释和空行),或者使用cl.h头文件的 cl.hpp C++ 包装器有 193 行代码。这个基于opencl_node特性的新版本进一步将源文件的大小减少到只有 144 行代码。

细节决定成败

我们这些以前开发过 OpenCL 代码的人知道，如果我们直接使用原始 OpenCL 库，我们可以“享受”相当大的自由度。乍看之下，这种灵活性并没有体现在opencl_node中。怎样才能定义一个多维的NDRange？除了NDRange的全局尺寸，我们如何指定局部尺寸？我们如何提供一个预编译的内核来代替 OpenCL 源代码呢？也许问题是我们还没有涵盖所有可用的配置旋钮。让我们开始回答这些问题。

启动内核所需的主要 OpenCL 函数是clSetKernelArg(如果我们使用 OpenCL 2.x 共享虚拟内存指针，则为clSetKernelArgSVMPointer)和clEnqueueNDRangeKernel。这些函数在 OpenCL 工厂中被内部调用，我们可以控制将哪些参数传递给它们。为了说明opencl_node成员函数和助手函数如何被转换成原始 OpenCL 调用，我们放大了图 19-17 中的opencl_node。

图 19-17

opencl_node函数和本地 OpenCL 调用之间的内部和对应关系

在这个图中，我们使用前面三元组示例中的gpu_node，其中我们配置了一个opencl_node来接收三个opencl_buffers和NDRange(总共四个进出节点的端口)。正如我们在几页前解释的那样，由于gpu_node.set_args(port_ref<0,2>, alpha)，我们清楚地说明了携带A、B和C向量的前三个输入端口(0、1 和 2)应该绑定到内核的前三个参数，内核的最后一个参数(乘法因子α)静态绑定到变量alpha，该变量不来自图的前面的节点。现在，我们已经获得了进行图 19-17 中所示的四个clSetKernelArg()调用所需的所有信息，这四个调用依次发挥它们的魔力，使这四个参数作为输入出现在kernel void triad(...) OpenCL 函数中。

现在，让我们看看如何适当地配置clEnqueueNDRangeKernel调用。这是最复杂的 OpenCL 调用之一；这需要我们在图 19-18 中列出的九个参数。然而，这不是一本 OpenCL 初级读本，对于本章来说，只讨论第二到第六个参数就足够了。用变量“kernel”标识的一个将在后面讨论，为了理解其他四个，我们必须更深入地研究 OpenCL 的基本概念之一:NDRange。

图 19-18

OpenCL clEnqueueNDRangeKernel调用的签名

NDRange概念

一个NDRange定义了一个独立工作项的迭代空间。这个空间可能是三维的，但也可能是 2D 或 1D。在我们的三元组示例中,NDRange是 1D。图 19-17 和 19-18 中clEnqueueNDrangeKernel调用中的参数dim应相应地包含 1、2 或 3，并将由gpu_node.set_range()调用正确设置。在图 19-17 的例子中，这个set_range()调用指出NDRange信息从图的前一个节点到达gpu_node的端口 3。NDRange信息应该在一个或者可选的两个容器中，这两个容器提供了begin()和end()成员函数。许多标准 C++ 类型都提供了这些成员函数，例如std::initializer_list、std::vector、std::array和std::list。如果我们只指定一个容器，opencl_node只设置clEnqueueNDRangeKernel()函数的global_work_size参数(在图 19-17 和 19-18 中用变量global标识)。否则，我们也指定第二个容器，opencl_node也设置local_work_size参数(图 19-17 和 19-18 中的local)。

注意

正如我们所说的，NDRange global_work_size定义了将由加速器执行的并行迭代空间。使用 OpenCL 俚语，这个空间中的每个点都被称为一个工作项(如果您熟悉 CUDA，它相当于一个 CUDA 线程)。因此，工作项目可以在不同的加速器计算单元 CUs 上并行处理，相应的计算由内核代码定义，也就是说，如果我们的内核函数包括C[i]=A[i]+B[i],，这是将应用于该 1D 迭代空间的每个工作项目i的表达式。

现在，工作项被分组为所谓的工作组(或者使用 CUDA 符号的块)。由于架构实现的细节，属于同一个工作组的工作项之间的联系更加紧密。例如，在 GPU 上，可以保证在单个 GPU 计算单元上调度一个工作组。这意味着我们可以用 OpenCL barrier 同步单个工作组的工作项，这些工作项共享一个称为“本地内存”的每 CU 内存空间，它比全局内存快。

参数local_work_size指定了工作组的规模。如果没有提供，OpenCL 驱动程序可以自动计算推荐的local_work_size。然而，如果我们想要强制一个特定的工作组规模，我们必须设置local_work_size参数。

这里的一些例子将使它变得非常清楚。假设我们有维度为h x w的 2D 数组A、B和C，我们想计算矩阵运算 C=A+B。虽然矩阵是二维的，但在 OpenCL 中，它们是作为指向行为主的线性化 1Dcl_mem缓冲区的指针传递给内核的。这并不妨碍我们从 2D 指数计算 1D 指数，所以内核看起来像这样

尽管表达相同内容的奇特方式使用了int2类型，读作

  int2 gId = (int2)(get_global_id(0),   get_global_id(1));
  C[gId.y*w+gId.x] = A[gId.y*w+gId.x] + B[gId.y*w+gId.x];

为了获得内核执行期间每个工作项的更多信息，我们将打印出一些附加信息，如图 19-19 所示。

图 19-19

添加两个矩阵并打印出相关工作项信息的内核示例

前三个变量gId、lId和grId分别在维度x和y中存储每个工作项的全局 ID、本地 ID 和组 ID。接下来的三个变量gSize、lSize和numGrp被设置为全局大小、局部大小和工作组数量。第一个 if 条件仅由具有全局ID (0,0).的工作项满足，因此只有该工作项打印出不同大小和数量的组，这对于所有工作项都是相同的。第二个printf语句由每个工作项执行，并打印该工作项的全局、局部和组 id。当与dim = 2、global = {4,4}和local = {2,2}.一起排队时，这将产生如图 19-20 所示的输出

图 19-20

图 19-19 配置dim=2、global={4,4}、local={2,2} --set_range({{4, 4}, {2, 2}})--时的内核输出

在这个图中，我们用一个彩色的方框描述了每个工作项。有 16 个工作项排列在一个 4×4 的网格中，我们用四种不同的颜色来标识每个工作组。由于局部尺寸是{2,2}，每个工作组是一个 2×2 的子空间。难怪组的数量是 4，但是为了给这一章提供一些形式主义，我们在这里添加了一些我们可以很容易证明的不变量:

numGrp.x = gSize.x/lSize.x
0 <= gId.x < gSize
0 <= lId.x < lSize
gId.x = grId * lSize.x + lId.x

同样，对于.y坐标(或者甚至是 3D 空间中的.z)

现在，我们如何指定一个opencl_node的全局和局部大小？到目前为止，我们只是在本章前面的例子中使用了gpu_node.set_range({{<num>}})。这将转化为dim=1、global={<num>}和local=NULL，这导致 1D NDRange的本地大小由 OpenCL 驱动程序决定。

在一般情况下，我们可能需要global={gx, gy, gz}和local={lx, ly, lz}.实现这一点最简单的方法是使用

    gpu_node.set_range({{gx, gy, gz},{lx, ly, lz}});

然而，正如我们所说的，任何可以用begin()成员函数迭代的容器也将满足我们的需求。例如，一种更复杂的表达方式是

结果范围的维度与容器中的元素数量一样多，每个维度的大小都设置为相应的元素值。这里的警告是为全局和局部容器指定相同的维度。

为了让事情变得有趣，我们必须添加可以启动图 19-19 的内核的 TBB 驱动程序代码。我们所知道的最简洁的方法是建立一个只有一个opencl_node的图，如图 19-21 所示。

图 19-21

opencl_node孤立地练习

看到了吗？只需几行代码，我们就可以开始运行添加两个矩阵 A 和 b 的 OpenCL 代码。请注意，opencl_node、gpu_node只有一个端口port<0>，它绑定到内核的第三个参数 matrix C，它携带内核中执行的计算的结果。使用set_args成员函数直接传递输入矩阵A和B以及矩阵宽度w。还要注意的是，opencl_node必须至少有一个端口，并且只有当一个消息到达这个入口端口时，它才被激活。实施gpu_node的替代方案如下:

其中gpu_node在port<0>,上接收Cdevice，在port<1>上接收NDRange，其余的内核参数由set_range()成员函数指定。到达和离开gpu_node的port<1>的消息类型是tbb::flow::opencl_range(到目前为止是第无数个opencl_node助手类！)，我们依靠try_put()来传递一个用两个容器初始化的opencl_range对象。

玩弄偏移

我们留下了clEnqueueNDRangeKernel函数的另外两个参数(见图 19-18 )。一个是 offset 参数，可以用来跳过迭代空间开始处的一些第一个工作项。在 OpenCL 工厂的当前实现中，这个偏移量是硬连线到{0,0,0}.的，没什么大不了的。有两种可能的解决方法来克服这个限制。

第一种方法是将偏移量传递给内核，并在索引数组之前将其添加到全局 ID 中。例如，对于一维的C=A+B操作，我们可以这样写

当然，我们可以修改NDRange来避免数组溢出。虽然实用，但不是一个超级优雅的解决方案。那么哪个是超级优雅的解决方案呢？嗯，我们可以使用opencl_subbuffer类来实现相同的结果。例如，如果我们只想添加向量A和B的一个子区域，我们可以保留一个简单版本的向量添加内核:

但是将以下参数传递给set_args()成员函数:

  Adevice.subbuffer(offset, size)

同样，对于Bdevice和Cdevice。创建Cdevice子缓冲区的另一种方法是调用

  tbb::flow::opencl_subbuffer<cl_float>(Cdevice, offset, size)

指定 OpenCL 内核

最后，我们必须花些时间来讨论kernel的论点(见图 19-18 )。到目前为止，我们使用 OpenCL 源文件来提供我们的内核。在图 19-21 的最后一个例子中，我们再次使用了opencl_program类:

这相当于更显式的构造器:

这是提供内核函数的常用方法，一方面，它需要在运行时编译源代码，另一方面，它提供了可移植性，因为源代码将为所有可用的设备编译(在opencl_program构造时只编译一次)。在内部，OpenCL 工厂依赖于 OpenCL 函数clCreateProgramWithSource和clBuildProgram。

如果我们确信不需要将我们的代码移植到任何其他平台，和/或如果对于生产版本，我们需要最后一点性能，我们也可以预编译内核。例如，借助英特尔 OpenCL 工具链，我们可以运行

ioc64 -cmd=build -input=my_kernel.cl -ir=my_kernel.clbin
      -bo="-cl-std=CL2.0" -device=gpu

它生成预编译文件my_kernel.clbin。现在，我们可以使用

当将这种类型的文件传递给opencl_program构造器时，工厂内部使用clCreateProgramWithBinary来代替。另一种可能性是使用opencl_program_type::SPIR提供内核的 SPIR 中间表示。要生成 SPIR 版本，我们可以使用

ioc64 -cmd=build -input=my_kernel.cl -spir64=my_kernel.spir
      -bo="-cl-std=CL1.2"

在这两种情况下，ioc64编译器都会提供一些有用的信息。最后一次运行的输出如下所示

Using build options: -cl-std=CL1.2
OpenCL Intel(R) Graphics device was found!
Device name: Intel(R) HD Graphics
Device version: OpenCL 2.0
Device vendor: Intel(R) Corporation
Device profile: FULL_PROFILE
fcl build 1 succeeded.
bcl build succeeded.
my_kernel info:
       Maximum work-group size: 256
       Compiler work-group size: (0, 0, 0)
       Local memory size: 0
       Preferred multiple of work-group size: 32
       Minimum amount of private memory: 0
Build succeeded!

这个输出告诉我们关于这个特定内核的最大工作组大小 256，以及工作组大小的首选倍数 32。

更多关于设备选择的信息

在上一节中，我们意识到我们用来进行实验的笔记本电脑包括两个 GPU。让我们看一个简单的例子，在这个例子中，我们在同一个流程图中使用了它们。在图 19-22 中，我们链接了两个opencl_nodes以便第一个计算C=A+B并将C发送给下一个执行C = C – B的。当两个节点都完成时，我们检查常规function_node中的C == A。数组尺寸为rows × cols。

图 19-22

两个opencl_node的例子，每个配置使用不同的 GPU

在我们的笔记本电脑上，我们已经知道设备列表f.devices()包括三个设备，第二个和第三个是两个 GPU。这样，我们可以安全地使用f.devices().begin() +1和+2来获得指向每个 GPU 的迭代器，正如我们在图 19-22 的两个opencl_node定义的装箱语句中看到的。除了针对不同的 GPU，每个opencl_node都被配置为运行程序的两个不同内核fig_19_23.cl: cl_add和cl_sub。从gpu_node1流向gpu_node2的信息就是opencl_buffer Cdevice。在 OpenCL 工厂内部，数据移动被最小化，例如，如果一个opencl_buffer必须由映射到同一 GPU 的两个连续的opencl_nodes访问，则在图形的第一个 CPU 节点尝试访问相应的缓冲区(通过使用opencl_buffer.begin()或opencl_buffer.data()成员函数)之前，分配在 GPU 上的数据不会被移动到 CPU。

在图 19-23 中，我们展示了程序fig_19_23.cl，包括前面代码中引用的两个内核。注意，我们没有将行宽作为第四个参数传递，而是使用包含相同值的gSz.x。

图 19-23

fig_19_23.cl的内容，我们看到两个内核，每个内核都是从不同的opencl_node调用的

在我们的笔记本电脑上运行图 19-22 的代码产生的输出如下:

Running gpu_node1 on Intel(R) HD Graphics 530
Running gpu_node2 on AMD Radeon Pro 450 Compute Engine
gSz.x=4, gSz.y=4
gSz.x=4, gSz.y=4

也可以用一个opencl_node来改变 OpenCL 设备，每次调用节点时工作都被卸载到这个设备上。图 19-24 的例子显示了一个被调用三次的opencl_node，对于每一次调用，不同的设备被用于运行一个简单的内核。

图 19-24

一个opencl_node就可以改变每次调用的目标加速器

代码使用初始化为0的原子变量device_num。对gpu_node的每次调用返回不同的设备，循环遍历所有设备(在我们的平台中有三个)。以及以下内核:

产生的输出是

Iteration: 0
Iteration: 1
Iteration: 2
Running on Intel(R) Core(TM) i7-6700HQ CPU @ 2.60GHz
Running on Intel(R) HD Graphics 530
Running on AMD Radeon Pro 450 Compute Engine
A[0]=1
A[0]=2
A[0]=3

其中我们可以证实数组Adevice的元素在gpu_node的三次连续调用中已经增加了三次，并且相应的内核已经在三个不同的 OpenCL 设备上执行。

关于订单的警告是适当的！

我们应该注意的最后一个警告是，当从几个节点提供服务时，消息到达一个opencl_node的顺序。例如，在图 19-25 中，我们展示了一个流程图g，它包括一个由两个功能节点filler0和filler1提供的gpu_node。每个“填充器”发送 1000 个缓冲区，b，每个缓冲区有 10 个整数，形式为{i,i,i,…,i},，范围从 1 到 1000。接收方gpu_node接收两个消息作为b1和b2，并调用一个 OpenCL 内核，就像这样简单:

正如我们看到的，它基本上是乘以b1[i]=b1[i]*b2[i]。如果b1和b2相等(等于{1,1,1,…}，或者{2,2,2,…}，等等。)，我们应该在输出端得到 1000 个平方输出的缓冲器({1,1,1,…}，然后是{4,4,4,…}，等等)。正确确定吗？我们不想说谎，所以为了以防万一，让我们在图的最后一个节点checker中仔细检查一下，它验证了我们的假设。

图 19-25

两个功能节点向opencl_node提供缓冲区，这些缓冲区将在 GPU 上相乘

图 19-26 中列出了实现上图的代码。我们同意乔治·萧伯纳的观点:“说谎者的惩罚丝毫不在于他不被人相信，而在于他不能相信任何人。”作为骗子鉴赏家，我们在代码中使用了一个专门用来捕捉骗子的 try-catch 结构。

图 19-26

与图 19-25 中描绘的图形相对应的源代码

我们首先将buffer_i定义为整数的opencl_buffer。两个“填充器”接收一个整数i，并用 10 个i填充一个buffer_i，然后发送到gpu_node。用于配置opencl_node的三行代码对我们来说太基础了，不需要进一步阐述。最后一个节点是检查器，如果在 GPU 上处理的缓冲区中接收的任何值不是平方整数，它将抛出异常。在制作边缘之后，1000 次迭代循环使两个填充器工作。现在，关键时刻到了，结果是

Liar!!: 42 is not a square of any integer number

好吧，我们被抓了！显然，6*7是在 GPU 上计算的，而不是在6*6或7*7上。为什么呢？答案是我们没有采取足够的措施来确保到达gpu_node的消息被正确配对。记住“填充符”的主体是由任务执行的，我们不能假定任务执行的任何特定顺序。

幸运的是，opencl_node带有一个方便的特定类型的键匹配特性，这将扭转局面。我们在图 19-27 中使用了这个特性。

图 19-27

修正图 19-26 的代码

基本上，现在的buffer_i是一个继承自opencl_buffer<cl_int>的新类，增加了一个int my_key成员变量和一个返回该键的key()成员函数。现在填充器必须使用不同的构造器(buffer_i b{N,i})，但更重要的是，opencl_node接收第二个模板参数(key_matching<int>)。这将自动指示opencl_node调用key()函数，并等待具有相同键值的消息被传递到所有输入端口。搞定了。如果我们用这些小的修改来运行我们的代码，我们将会看到现在我们已经被宣判伪证罪不成立了！

摘要

在这一章中，我们介绍了 TBB 流图的opencl_node特征。我们从一个简单的Hello OpenCL_Node例子开始，它代表了对opencl_node的初步了解，涵盖了这个类的基础知识。然后我们开始深入研究一些助手类，比如opencl_device对象的容器opencl_device_list，以及设备过滤器和设备选择器实体。为了说明其他助手类并给出一个更复杂的例子，我们还使用一个opencl_node实现了三元向量运算，以处理部分计算，而其余部分在 CPU 内核上同时处理。在那里，我们更好地介绍了opencl_buffer助手类和opencl_node类的set_range和set_args成员函数。NDRange概念以及如何设置全局和局部 OpenCL 大小几乎需要一个章节，在这里我们还解释了如何使用opencl_subbuffer类和其他变体来提供内核程序(预编译或 SPIR 中间表示)。接下来，我们介绍了两个例子，说明了如何将流程图的不同opencl_node映射到不同的设备上，或者甚至如何在每次调用时更改opencl_node卸载计算的设备。最后，我们描述了当一个opencl_node来自不同的节点时如何避免排序问题。

最后一个免责声明。也许最后我们真的在撒谎。在写这一章的时候，opencl_node仍然是一个预览功能，所以它可能会被修改。经过 3 年的发展，我们不期望有大的变化，但我们不能承诺这一点。如果这样的变化在未来的版本中结束，我们保证会写这一章的更新版本！你相信我们吗？

更多信息

以下是我们推荐的一些与本章相关的额外阅读材料:

• Alexei Katranov，Opencl_node 概述。英特尔开发人员专区博客中的系列文章: https://software.intel.com/en-us/blogs/2015/12/09/opencl-node-overview 。

• 大卫凯利，佩哈德米斯特里，达纳沙，张东平。使用 OpenCL 2.0 进行异构计算。摩根·考夫曼 2015。

图中的徒步图标 19-1 由来自 www.flaticon.com 的 Scott de Jonge 制作。

开放存取本章根据知识共享署名-非商业-非专用 4.0 国际许可协议(http://Creative Commons . org/licenses/by-NC-nd/4.0/)的条款进行许可，该协议允许以任何媒体或格式进行任何非商业使用、共享、分发和复制，只要您适当注明原作者和来源，提供知识共享许可协议的链接，并指出您是否修改了许可材料。根据本许可证，您无权共享从本章或其部分内容派生的改编材料。

本章中的图像或其他第三方材料包含在本章的知识共享许可中，除非在材料的信用额度中另有说明。如果材料不包括在本章的知识共享许可中，并且您的预期使用不被法律法规允许或超出了允许的使用范围，您将需要直接从版权所有者处获得许可。

Footnotes 1

https://software.intel.com/en-us/blogs/2015/12/09/opencl-node-overview 。

2

请再次记住，这是一个预览功能。如果您在这方面需要更多的灵活性，我们将非常感谢您提出请求，让英特尔知道您认为 OpenCL 节点很有用，但仍有一些限制需要解决。

二十、NUMA 架构上的 TBB

关心性能的高级程序员知道利用本地性是最重要的。谈到局部性，缓存局部性是第一个想到的，但是在许多情况下，对于运行在大型共享内存架构上的重型应用程序，还应该考虑非统一内存访问(NUMA)局部性。众所周知，NUMA 传达了这样一个信息:内存组织在不同的存储体中，一些内核对一些“近”存储体的访问速度要快于对“远”存储体的访问速度。更正式地说， NUMA 节点是内核、高速缓存和本地存储器的分组，其中所有内核共享对本地共享高速缓存和存储器的相同访问时间。从一个 NUMA 节点到另一个节点的访问时间可能要长得多。出现了一些问题，例如程序数据结构如何在不同的 NUMA 节点上分配，以及处理这些数据结构的线程在哪里运行(它们是靠近还是远离数据？).在本章中，我们将解决这些问题，但更重要的是，在 TBB 并行应用程序中，如何利用 NUMA 局部性。

针对 NUMA 系统的性能调优归结为四项活动:(1)发现您的平台拓扑结构，(2)了解从系统的不同节点访问内存的相关成本，(3)控制数据的存储位置(数据放置)，以及(4)控制工作的执行位置(处理器关联性)。

为了防止你进一步失望(也就是现在就让你失望！)，我们应该提前声明:目前，TBB 不提供利用 NUMA 本地性的高级特性。或者换句话说，在前面列出的四个活动中，TBB 只在第四个活动中提供了一些帮助，在第四个活动中，我们可以依靠 TBB task_arena(参见第十二章)和本地task_sheduler_observer(参见第十三章)类来识别应该限制在 NUMA 节点中的线程。对于所有其他活动，甚至对于将线程实际固定到 NUMA 节点(这是第四个活动的基本部分)，我们需要使用低级的依赖于操作系统的系统调用或高级别的第三方库和工具。这意味着，即使这是一本 TBB 的书，这最后一章也不完全是关于 TBB 的。我们的目标是详细阐述如何实现利用 NUMA 局部性的 TBB 代码，即使大多数必需的活动与 TBB 没有直接关系。

既然我们已经提醒了读者，让我们把这一章分解成几个部分。我们基本上按顺序遵循前面列出的四项活动。第一部分展示了一些工具，它们可以用来发现我们平台的拓扑结构，并检查有多少 NUMA 节点可用。如果有多个 NUMA 节点，我们可以继续下一部分。在这里，我们使用一个基准来了解在我们的特定平台上利用 NUMA 本地性时潜在的加速效果。如果预期的收益令人信服，我们应该开始考虑在我们自己的代码中利用 NUMA 局部性(不仅仅是在一个简单的基准中)。如果我们认识到我们自己的问题可以受益于 NUMA 局部性，我们就可以进入问题的核心，即掌握数据放置和处理器关联性。有了这些知识，在 TBB task_arena和task_scheduler_observer类的帮助下，我们实现了第一个简单的 TBB 应用程序，该应用程序利用了 NUMA 局部性，并评估了相对于基线实现所获得的加速。整个过程总结在图 20-1 中。我们结束这一章，概述可以考虑用于更复杂应用的更高级和更通用的替代方案。

图 20-1

开发 NUMA 地区所需的活动

注意

如果你想知道为什么在当前版本的 TBB 中没有高级别的支持，这里有一些原因。首先，这是一个棘手的问题，高度依赖于必须并行化的特定应用程序及其运行的架构。因为没有一个放之四海而皆准的解决方案，所以由开发人员来决定最适合当前应用的特定数据放置和处理器关联性替代方案。其次，TBB 的架构师和开发人员总是试图避免在 TBB 库中使用特定于硬件的解决方案，因为它们可能会损害代码的可移植性和 TBB 的可组合性。该库不仅仅是为了执行 HPC 应用程序而开发的，在 HPC 应用程序中，我们通常可以独占访问整个高性能平台(或它的一个分区)。TBB 还应该在其他应用程序和进程也在运行的共享环境中尽力而为。在许多情况下，将线程绑定到内核并将内存绑定到 NUMA 节点会导致底层架构的利用不尽人意。在任何具有动态特性的应用程序或系统中，手动锁定被反复证明是一个坏主意。我们强烈建议不要采用这种方法，除非您确信您将在您的特定并行平台上提高您的特定应用程序的性能，并且您不关心可移植性(或者付出额外的努力来实现可移植的 NUMA 感知应用程序)。

考虑到 TBB 并行算法基于任务的特性和支持并行执行的工作窃取调度程序，让任务在接近本地内存的内核中运行似乎具有挑战性。但这不会阻止像我们这样勇敢无畏的程序员。让我们去吧！

发现您的平台拓扑

“知己知彼，百战不殆.”—孙子兵法。这句千年名言告诉我们，在解决问题之前，首先要努力仔细理解我们所面临的问题。有一些工具可以方便地理解底层的 NUMA 架构。在本章中，我们将使用hwloc和likwid ¹ 来收集关于架构和代码执行的信息。hwloc是一个软件包，它提供了一种便捷的方式来查询关于系统拓扑的信息，以及应用一些 NUMA 控制，如数据放置和处理器关联性。likwid是另一个软件包，它告知硬件拓扑结构，可用于收集硬件性能计数器，还提供一组有用的微基准，可用于描述系统特征。我们还可以使用 VTune 来分析代码的性能。虽然likwid只适用于 Linux，但是hwloc和 VTune 也可以很容易地安装在 Windows 和 MacOS 上。然而，由于用于说明我们代码的共享内存平台运行 Linux，除非另有说明，否则这将是我们假定的操作系统。

因为针对 NUMA 的调优需要对所使用的平台有深入的理解，所以我们将从描述两台机器的特征开始，这两台机器将贯穿本章。我们接下来介绍的两台机器被称为yuca(来自丝兰工厂)和aloe(来自芦荟工厂)。首先，我们可以收集这些机器的基本信息。在 Linux 上，可以使用命令“lscpu”获得这些信息，如图 20-2 所示。

图 20-2

尤卡和芦荟的lscpu产量

乍一看，我们看到 yuca 有 64 个逻辑内核，编号从 0 到 63，每个物理内核有两个逻辑内核(超线程又称 SMT 或同步多线程，可用)，每个插槽有八个物理内核，四个插槽也是四个 NUMA 节点或 NUMA 域。就其本身而言，aloe 有 32 个禁用超线程的物理内核(每个内核只有一个线程)，每个插槽有 16 个物理内核，还有两个插槽(NUMA 节点)。在lscpu输出的最后，我们可以看到 NUMA 节点和每个节点中包含的逻辑核心的 id，但是如果我们使用来自hwloc库的lstopo实用程序，画面会变得更加清晰。在图 20-3 中，我们包括了执行lstopo --no-io yuca.pdf命令时在 yuca 上生成的 PDF 文件(参数--no-io不考虑 I/O 设备拓扑)。

图 20-3

在 yuca 上执行lstopo的结果

从这个图中，我们可以清楚地看到尤卡的 NUMA 组织。四个 NUMA 节点包括八个物理核心，操作系统将其视为 16 个逻辑核心(也称为硬件线程)。请注意，逻辑内核 id 取决于架构、固件(电脑上的 BIOS 配置)和操作系统版本，因此我们不能从编号中做出任何假设。对于 yuca 的特定配置，逻辑核心 0 和 32 共享同一个物理核心。现在我们更好地理解了 yuca 上lscpu最后四行的意思:

NUMA node0 CPU(s):     0-7,32-39
NUMA node1 CPU(s):     8-15,40-47
NUMA node2 CPU(s):     16-23,48-55
NUMA node3 CPU(s):     24-31,56-63

在 yuca 上，每个 NUMA 节点有 63 GB 的本地内存，总共 252 GB。类似地，aloe 也具有 252 GB 的容量，但仅组织在两个 NUMA 节点中。在图 20-4 中，我们看到了芦荟上lstopo输出的编辑版本。

图 20-4

对芦荟执行lstopo的结果

我们看到，在 aloe 上，每个物理核心都包含一个逻辑核心，在第一个域中编号为 0-15，在第二个域中编号为 16-31。

了解访问内存的成本

现在我们知道了平台的拓扑结构，假设我们已经控制了处理器关联性和数据放置，让我们量化由于非本地访问而产生的开销。实际上，我们确实在已经可用的基准上控制这两个方面，比如在likwid工具中可用的likwid -bench。使用这个基准，我们可以使用一个命令行运行流三元组代码(参见前两章):

likwid-bench -t stream -i 1 -w S0:12GB:16-0:S0,1:S0,2:S0

它运行用-w参数配置的流基准的单次迭代(-i 1)

• S0:线程被固定到 NUMA 节点 0。

• 12 GB:三个三元组阵列占用 12 GB(每个阵列 4 GB)。

• 16: 16 个线程将共享计算，每个线程处理 31，250，000 个 double 的数据块(即 40 亿字节/每个 double/16 个线程 8 个字节)。

• 0:S0,1:S0,2:S0:三个数组分配在 NUMA 节点 0 上。

在 yuca 上，该命令的结果报告了 8219 MB/s 的带宽。但是，更改三个数组的数据放置是很容易的，例如，更改到 NUMA 节点 1(使用0:S1,1:S1,2:S1)将 16 个线程的计算限制在 NUMA 节点 0 中。毫不奇怪，我们现在得到的带宽只有 5110 MB/s，这意味着我们损失了 38%的带宽，这是我们在利用 NUMA 本地性时测量的。对于计算本地数据的其他配置(数据放置在线程固定的内核上)和不利用本地性的配置(数据放置在没有线程关联性的内核上)，我们得到了类似的结果。在 yuca 上，所有非本地配置都会导致相同的带宽冲击，但是在其他 NUMA 拓扑上，我们会根据数据放置的位置和线程运行的位置而付出不同的代价。

在芦荟上，我们只有两个 NUMA 节点 0 和 1。将数据和计算放在同一个域中可以获得 38671 MB/s 的速度，而沿着错误的路径只能获得 20489 MB/s 的速度(几乎是一半，整整少了 47%的带宽)。我们确信，像您这样渴望阅读和学习性能编程主题的读者，现在正积极地在您自己的项目中利用 NUMA 本地性！

我们的基线示例

图 20-5 显示了我们最近一直在使用的三元组示例的并行版本，只有一个parallel_for算法。

图 20-5

对基线算法进行评估和改进

这段代码的最后两行报告了执行时间和获得的带宽，它还没有针对 NUMA 进行优化。对于后者，访问的总字节数计算为每个数组元素的vsize × 8 字节/double × 3 次访问(两次加载和一次存储)，然后除以执行时间和一百万(转换为每秒兆字节)。在 yuca 上，当使用 32 个线程和一个千兆元素的数组运行时，会产生以下输出:

./fig_20_05 32 1000000000
Time: 2.23835 seconds; Bandwidth: 10722.2MB/s

关于芦荟:

./fig_20_05 32 1000000000
Time: 0.621695 seconds; Bandwidth: 38604.2MB/s

请注意，我们的 triad 实现获得的带宽不应与之前由likwid-bench报告的带宽进行比较。现在，我们使用 32 个线程(而不是 16 个)，根据操作系统调度程序，这些线程可以在每个内核上自由运行(而不是局限于单个 NUMA 节点)。类似地，阵列现在由操作系统按照自己的数据放置策略来放置。在 Linux 中，默认的策略 ² 是“本地分配”,其中执行分配的线程决定数据的位置:如果有足够的空间，则在本地内存中，否则在远程。这种策略有时被称为“首次接触”，因为数据放置不是在分配时完成的，而是在首次接触时完成的。这意味着一个线程可以分配一个区域，但是首先访问这个区域的线程是引发页面错误的线程，并且实际上是将内存中的页面分配给该线程。在我们的图 20-5 的例子中，相同的线程分配并初始化数组，这意味着在相同的 NUMA 节点上运行的parallel_for工作线程将具有更快的访问速度。最后一个区别是likwid-bench用汇编语言实现三元组计算，这阻止了进一步的编译器优化。

掌握数据放置和处理器关联性

绑定数据和计算一点也不简单。主要是因为它依赖于操作系统，每个操作系统都有自己的系统调用。在 Linux 中，低级接口由libnuma ³ 提供，其包括控制在 Linux 内核中实现的数据放置和处理器亲缘关系策略的功能。一个更高级的替代命令是numactl ⁴ 命令，它解决了同样的问题，但是灵活性较差。

然而，破坏我们的 TBB 应用程序与依赖于操作系统的 NUMA 库的可移植性并不是最好的主意。已经提到的hwloc库是一个可移植且广泛使用的替代方案。目前，TBB 没有提供自己的 API 来处理 NUMA 本地数据，但是正如我们将在后面看到的，我们可以采取一些措施来让我们的 TBB 任务在可能的时候访问本地数据。在撰写本文时，必须通过第三方库来手动控制数据放置和处理器关联性，不失一般性，我们将求助于本章中的hwloc。这个库可以在 Windows、MacOS 和 Linux 中使用(实际上，在 Linux 中hwloc使用下面的numactl/libnuma)。

在图 20-6 中，我们展示了一个例子，它查询 NUMA 节点的数量，然后在每个节点上分配一些数据，稍后为每个节点创建一个线程，并将其绑定到相应的域。我们在下面使用的是hwloc 2.0.1。

图 20-6

使用hwloc为每个 NUMA 节点分配内存和绑定线程

所有hwloc函数的一个反复出现的参数是对象拓扑，在我们的例子中是topo。这个对象首先被初始化，然后加载平台的可用信息。之后，我们准备从topo数据结构中获取信息，正如我们对hwloc_get_nbobjs_by_type所做的那样，当第二个参数是HWLOC_OBJ_NUMANODE时，它返回 NUMA 节点的数量(其他几种类型也是可用的，如HWLOC_OBJ_CORE or HWLOC_OBJ_PU–逻辑核心或处理单元)。NUMA 节点的数量存储在变量num_nodes中。

该示例继续创建一个指向 doubles 的指针数组num_nodes，该数组将在函数alloc_mem_per_node中初始化。对alloc_thr_per_node的函数调用创建了num_nodes个线程，每个线程都被固定到相应的 NUMA 节点。这两个功能分别在图 20-7 和 20-8 中描述。这个例子通过释放分配的内存和topo数据结构来结束。

图 20-7

为每个 numa 节点分配双精度数组的函数

图 20-7 显示了功能alloc_mem_per_node的实现。关键操作是hwloc_get_obj_by_type，当第二个和第三个参数分别为HWLOC_OBJ_NUMANODE和i时，它返回一个句柄给i th NUMA 节点对象numa_node。这个numa_node有几个属性，如numa_node->cpuset(标识节点中包含的逻辑内核的位掩码)和numa_node->nodeset(标识节点的类似位掩码)。函数hwloc_bitmap_asprintf可以方便地将这些集合转换成字符串，我们将在程序的输出中看到后面的内容。使用nodeset位掩码，我们可以在带有hwloc_alloc_membind的节点中分配内存。

当运行代码直到alloc_mem_per_node返回到主函数时，我们在 yuca 上得到的输出是

There are 4 NUMA node(s)
NUMA node 0 has cpu bitmask: 0x000000ff,0x000000ff
Allocate data on node 0 with node bitmask 0x00000001
NUMA node 1 has cpu bitmask: 0x0000ff00,0x0000ff00
Allocate data on node 1 with node bitmask 0x00000002
NUMA node 2 has cpu bitmask: 0x00ff0000,0x00ff0000
Allocate data on node 2 with node bitmask 0x00000004
NUMA node 3 has cpu bitmask: 0xff000000,0xff000000
Allocate data on node 3 with node bitmask 0x00000008

这里我们看到每个 NUMA 节点的cpuset和nodeset。如果我们再次刷新我们的记忆，查看图 20-3 ，我们会看到在节点 0 中我们有 8 个内核和 16 个逻辑内核，编号从 0 到 7 和从 32 到 39，在hwloc中用位掩码0x000000ff,0x000000ff表示。请注意，“”分隔了共享八个物理内核的两组逻辑内核。与禁用超线程的平台相比，这是 aloe 上的相应输出:

There are 2 NUMA node(s)
NUMA node 0 has cpu bitmask: 0x0000ffff
Allocate data on node 0 with node bitmask 0x00000001
NUMA node 1 has cpu bitmask: 0xffff0000
Allocate data on node 1 with node bitmask 0x00000002

在图 20-8 中，我们列出了为每个 NUMA 节点生成一个线程的函数alloc_thr_per_node，然后使用cpuset属性绑定它。

图 20-8

为每个 NUMA 节点创建并固定一个线程的函数

这个函数还查询 NUMA 节点的数量num_nodes，以便稍后在创建线程的循环中迭代这个次数。在每个线程执行的 lambda 表达式中，我们使用hwloc_set_cpubind将线程绑定到每个特定的 NUMA 节点，现在依赖于numa_node->cpuset。为了验证锁定，我们打印线程 id(使用std::this_thread::get_id)和运行线程的逻辑内核的 id(使用sched_getcpu)。接下来是 yuca 上的结果，也如图 20-9 所示。

图 20-9

描绘了由于固定到 yuca 上的 NUMA 节点而导致的线程移动

Before: Thread 0 with tid 873342720 on core 33
After: Thread 0 with tid 873342720 on core 33
Before: Thread 1 with tid 864950016 on core 2
After: Thread 1 with tid 864950016 on core 8
Before: Thread 2 with tid 856557312 on core 33
After: Thread 2 with tid 856557312 on core 16
Before: Thread 3 with tid 848164608 on core 5
After: Thread 3 with tid 848164608 on core 24

这里有两件事值得一提。首先，线程最初由操作系统分配到同一个 NUMA 节点中的逻辑核心上，因为它假设它们会协作。线程 0 和 2 甚至被分配在同一个逻辑核心上。其次，线程不是固定在单个内核上，而是固定在属于同一个 NUMA 节点的整个内核集上。如果操作系统认为将一个线程移动到同一个节点的不同内核会更好，这就留有余地。为了完整起见，下面是芦荟的等效输出:

Before: Thread: 0 with tid 140117643171584 on core 3
After: Thread: 0 with tid 140117643171584 on core 3
Before: Thread: 1 with tid 140117634778880 on core 3
After: Thread: 1 with tid 140117634778880 on core 16

有兴趣的读者可以从各自的文档和在线教程中了解到hwloc和likwid的更多特性。然而，我们在本节中所介绍的内容足以让我们继续前进，卷起袖子，使用 TBB 实现一个 NUMA 意识版本的 triad 算法。

和 TBB 一起工作

显然，首要目标是最大限度地减少非本地访问的数量，这意味着在离存储数据的内存最近的内核上进行计算。一种非常简单的方法是在 NUMA 节点上手动划分数据，并将处理这些数据的线程限制在相同的节点上。出于教育目的，我们将首先描述这个解决方案，并在下一节简要阐述更高级的替代方案。

我们可以依靠hwloc API 来完成数据放置和处理器关联任务，但是我们想要一个 NUMA 感知的 triad 基准的 TBB 实现。在这种情况下，管理线程的是 TBB 调度程序。从第十一章中，我们知道在tbb::task_scheduler_init函数中创建了许多线程。此外，这个 TBB 函数创建了一个默认的竞技场，它有足够的工作线程槽来允许线程参与执行任务。在我们 triad 的基线实现中(见图 20-5),parallel_for负责将迭代空间划分为不同的任务。所有线程将协作处理这些任务，而不管每个任务处理的迭代块以及线程运行的内核。但我们不希望它出现在 NUMA 的平台上，对吗？

我们最简单的基线三元组实施替代方案将通过执行以下三个步骤来增强实施:

• 它将在不同的 NUMA 节点上划分和分配三元组算法的三个向量 A、B 和 C。作为最简单的解决方案，静态块分区现在就可以了。在 yuca 上，这意味着 A、B 和 C 这四大块将被分配到四个节点中的每一个上。

• 它将在每个 NUMA 节点上创建一个主线程。每个主线程将创建自己的任务竞技场和自己的本地task_scheduler_observer。然后，每个主线程执行自己的tbb::parallel_for算法来处理对应于这个 NUMA 节点的 A、B 和 C 的分数。

• 它会自动将连接每个竞技场的线程固定到相应的 NUMA 节点。我们为每个竞技场创建的本地task_scheduler_observer将会负责此事。

让我们来看看所描述的每一个要点的实现。对于主函数，我们稍微修改了我们为图 20-6 的hwloc示例提供的函数。在图 20-10 中，我们列出了这个新示例所需的新行，在没有变化的行上使用省略号(…)。

图 20-10

NUMA 意识的主要功能是实现三和弦

程序参数thds_per_node允许我们在每个 NUMA 节点上使用不同数量的线程。如图 20-6 所示，num_nodes是我们使用hwloc API 获得的 NUMA 节点数。因此，我们传递给 TBB 调度器构造器(thds_per_node-1)*(num_nodes)而不是thds_per_node*num_nodes，因为我们将在alloc_thr_per_node中显式创建额外的num_nodes主线程。

函数alloc_mem_per_node本质上与图 20-7 中列出的函数相同，但现在用不同的大小参数调用它:doubles_per_node = vsize*3/num_nodes，其中vsize是三个向量的大小，所以 doubles 的总量乘以 3，但除以节点数来实现块划分。为了简洁起见，我们假设vsize是num_nodes的倍数。alloc_mem_per_node完成后，data[i]指向i th NUMA 节点上分配的数据。

如图 20-11 所示，alloc_thr_per_node功能的改编版本还有其他不同之处。它现在接收数据的句柄，每个节点将要遍历的本地向量的大小，lsize，以及用户设置的每个节点的线程数，thds_per_node。

图 20-11

该函数为每个节点创建一个线程，以计算每个 NUMA 节点上的三元组计算

注意，在图 20-11 呈现的代码片段中，在遍历num_nodes的i-循环内部，有三个嵌套的 lambda 表达式:(1)对于线程对象；(2)进行task_arena::execute会员功能；和(3)用于parallel_for算法。在外层，我们首先将线程固定到相应的 NUMA 节点i。

第二步是初始化在data[i]数组中分配的指向数组A、B和C的指针。在图 20-10 中，我们调用alloc_thr_per_node作为第三个参数vsize/num_nodes，因为在每个节点上，我们只遍历三个数组的块分布中的一个块。因此，函数的参数lsize = vsize/num_nodes，在初始化数组A和B的循环中使用，并作为计算C的parallel_for的参数。

接下来，我们初始化一个每个 NUMA 节点的 arena，numa_arena，它随后作为参数传递给一个task_scheduler_observer对象，p，并用于调用一个局限于这个 arena 的parallel_for(使用numa_arena.execute)。这就是我们 NUMA 感知的 triad 实现的关键。

parallel_for将创建遍历三个向量的局部分区块的任务。这些任务将由运行在同一个 NUMA 节点内核上的线程执行。但是到目前为止，我们只有thds_per_node*num_nodes个线程，其中num_nodes已经被明确衍生为主线程，并被固定到不同的 NUMA 节点，但是其余的仍然可以在任何地方自由运行。全局线程池中可用的线程将各自加入一个num_nodes竞技场。方便的是，每个numa_arena都已经用thds_per_node槽初始化，一个槽已经被主线程占用，其余的可供工作线程使用。我们现在的目标是将进入每个numa_arena的第一个thds_per_node-1线程固定到相应的 NUMA 节点。为此，我们创建了一个PinningObserver类(从task_scheduler_observer派生而来)并构造了一个对象p，向构造器传递了四个参数:PinningObserver p{numa_arena, topo, i, thds_per_node}。记住这里，i是主线程i的 NUMA 节点的 id。

在图 20-12 中，我们看到了PinningObserver类的实现。

图 20-12

实现本地task_scheduler_observer为三元组

在第十三章中介绍了task_scheduler_observer类。它有一个预览功能，允许我们在每个任务竞技场都有一个观察者——也称为本地task_scheduler_observer。这种观察者通过引用 arena 来初始化，就像我们在使用task_scheduler_observer{arena}的PinningObserver构造器的初始化列表中所做的那样。这导致进入这个特定领域的每个线程的成员函数on_scheduler_entry的执行。该类的构造器还设置了 NUMA 节点的数量，num_nodes和numa_node对象，它们将为我们提供对numa_node->cpuset位掩码的访问。构造器最后调用成员函数observe(true)开始跟踪任务是否进场。

函数on_scheduler_entry跟踪已经固定到原子变量thds_per_node中的numa_node的线程数量。这个变量在构造器的初始化列表中被设置为每个节点的线程数，用户将它作为程序的第一个参数传递。对于每个进入 arena 的线程，该变量递减，只有当值大于 0 时，该变量才会被固定到节点。由于每个numa_arena都是用thds_per_node槽初始化的，并且创建竞技场的已经被钉住的主线程占据了其中一个槽，所以首先加入竞技场的thds_per_node - 1线程将被钉住到节点，并处理由该竞技场正在执行的parallel_for生成的任务。

注意

我们的 PinningObserver 类的实现并不完全正确。一个线程可能离开竞技场并重新进入同一个竞技场，被钉住两次，但数量会减少thds_per_node。一个更正确的实现是检查进入竞技场的线程是否是一个还没有被固定到这个竞技场的新线程。为了避免这个例子变得复杂，我们把这个修正留给读者作为练习。

我们现在可以在 yuca 和 aloe 上评估这种 NUMA 优化版 triad 算法的带宽(以每秒兆字节为单位)。为了与图 20-5 中的基线实现进行比较，我们将向量大小设置为 10⁹double，并设置每个 NUMA 节点的线程数量，这样我们最终总共有 32 个线程。例如，在 yuca 中，我们将可执行文件称为:

baseline:           ./fig_20_05 32 1000000000
NUMA conscious:     ./fig_20_10  8 1000000000

图 20-13 的表格中显示的结果是十次运行的平均值，其中 yuca 和 aloe 有一个用户专门使用该平台进行实验。

图 20-13

由于 NUMA 意识的实施而加速

这在 yuca 上快了 74%,在 aloe 上快了 54%!您会忽略我们通过一些额外的实现工作从 NUMA 架构中挤出的额外性能吗？

为了进一步研究这种改进，我们可以利用能够读出硬件性能计数器的likwid-perfctr应用程序。通过调用likwid-perctr -a，我们可以得到一个事件组列表，只需使用组名就可以指定这些事件组。在 aloe 中，likwid提供了一个NUMA组，它收集关于本地和远程内存访问的信息。要在我们的基线和 NUMA 感知实现上测量该组中的事件，我们可以调用以下两个命令:

likwid-perfctr -g NUMA ./fig_20_05 32 1000000000
likwid-perfctr -g NUMA ./fig_20_10 16 1000000000

这将报告关于所有核心上的一些性能计数器的值的大量信息。被统计的事件包括

OFFCORE_RESPONSE_0_LOCAL_DRAM
OFFCORE_RESPONSE_1_REMOTE_DRAM

它为我们提供了本地内存和远程内存中被访问数据量的大致信息(因为是基于事件的采样)。对于基线 triad 实现，本地数据与远程数据的比率仅为 3.25，但在 NUMA 优化的 triad-numa 版本中，该比率高达 25.5。这证实了，对于这个内存受限的应用程序，我们利用 NUMA 局部性的努力在本地访问数量和执行带宽方面都取得了回报。

更高级的替代方案

对于常规的三元组代码，我们实现的简单解决方案是可以的，但是 TBB 的偷工减料调度器仅限于独立地平衡每个 NUMA 节点上的负载。在 yuca 上，将有四个parallel_for算法在运行，每个算法运行在一个 NUMA 节点上，该节点有八个线程，由八个物理内核提供服务。这种简单方法的缺点是，四个竞技场配置了八个插槽，这对于执行的稳态部分来说是没问题的，但是如果 NUMA 节点之间的负载没有完全平衡，就会限制 TBB 的灵活性。

例如，如果其中一个parallel_for算法首先结束，那么八个线程就变成空闲。他们回到全局线程池，但不能加入其他三个繁忙的舞台，因为所有的位置都已被填满。一个简单的解决方案是增加竞技场的插槽数量，同时保持固定线程的数量为thds_per_node。在这种情况下，如果一个parallel_for首先完成，那么返回全局池的八个线程可以在其他三个竞技场的空闲槽中重新分配。请注意，这些线程仍然被固定到原始节点，尽管它们现在将在不同节点的不同领域中工作，因此内存访问将是远程的。

当进入扩展竞技场的线程占用其空闲槽时，我们可以将它们固定到相应的 NUMA 节点(即使它们之前被固定到不同的 NUMA 节点)。现在这些帮助线程也将访问本地内存。但是，该节点可能会超额预订，这通常会影响性能(否则，您应该从一开始就超额预订每个 NUMA 节点)。对于每个特定的应用程序和体系结构，都应该进行彻底的实验，以决定是将线程迁移到 NUMA 节点有利，还是从原始节点远程访问数据有利。对于简单和规则的三元组算法，这些讨论的方法都没有显著提高性能，但是在更复杂和不规则的应用中，它们可能会提高性能。不仅远程访问有开销，而且线程从一个领域到另一个领域的迁移，以及再次锁定线程，都代表了必须通过更好的工作负载平衡来分摊的开销。

我们可以选择进行的另一场战斗与数据分区有关。在我们简单的三元组实现中，我们使用了三个数组的基本块分布，但是我们当然知道对于更不规则的应用程序有更好的数据分布。例如，代替在 NUMA 节点之间预先划分迭代空间，我们可以遵循引导调度方法。在每个 NUMA 节点上领导计算的每个主线程可以在计算开始时获得更大的迭代空间块，并且随着我们接近空间的末端而变小。这里需要注意的是，要保证数据块有足够的粒度，可以在每个 NUMA 节点的内核之间再次重新划分。

一个更复杂的替代方案是以一种分层的方式来概括工作窃取框架。为了允许竞技场之间和每个竞技场内部的偷工减料，可以实现竞技场的层次结构。Chen 和 Guo 为 Cilk 实现了类似的思想(参见“更多信息”部分),他们提出了一个三级工作窃取调度器，对于内存受限的应用程序，与更传统的工作窃取替代方案相比，性能提高了 54%。请注意，与受 CPU 限制的应用程序相比，受内存限制的应用程序将从 NUMA 局部性利用中获益更多。对于后者，内存访问开销通常被 CPU 密集型计算所隐藏。实际上，对于 CPU 受限的应用程序，为了利用 NUMA 局部性而增加调度程序的复杂性会导致额外的开销，最终得不偿失。

摘要

在这一章中，我们探索了一些利用 NUMA 局部性的替代方法，结合了 TBB 和第三方库，有助于控制数据放置和处理器关联性。我们从研究我们想要打败的敌人开始:NUMA 建筑。为此，我们引入了一些盟友库，hwloc和likwid。有了它们，我们不仅可以查询 NUMA 拓扑的底层细节，还可以控制数据放置和处理器关联性。我们展示了使用一些hwloc函数来分配每个节点的内存，为每个 NUMA 节点创建一个线程，并将线程固定到节点的内核。有了这个模板，我们重新实现了 triad 算法的基线版本，现在注意 NUMA 局部性。最简单的解决方案是将三个三元组数组分布在块中，在不同的 NUMA 节点中分配和遍历这些块。库hwloc是分配和固定线程的关键，TBB task_arena和task_scheduler_observer类有助于识别进入特定 NUMA 节点的线程。对于像 triad 基准测试这样常规的代码来说，这个初始解决方案已经足够好了，在两个不同的 NUMA 平台上，分别报告了 74%和 54%的性能提升(相对于基线 triad 实现)。对于更不规则和复杂的应用，本章的最后一节概述了更高级的替代方案。

更多信息

以下是我们推荐的一些与本章相关的额外阅读材料:

• Christoph Lameter，NUMA(非统一内存访问):概述，ACMqueue，第 11 卷，第 7 期，2013 年。

• Ulrich Drepper，每个程序员都应该知道的内存知识， www.akkadia.org/drepper/cpumemory.pdf ，2017。

• 、郭敏毅、关海兵，LAWS:面向多插槽多核架构的局部感知工作窃取，国际超级计算大会，ICS，2014。

开放存取本章根据知识共享署名-非商业-非专用 4.0 国际许可协议(http://Creative Commons . org/licenses/by-NC-nd/4.0/)的条款进行许可，该协议允许以任何媒体或格式进行任何非商业使用、共享、分发和复制，只要您适当注明原作者和来源，提供知识共享许可协议的链接，并指出您是否修改了许可材料。根据本许可证，您无权共享从本章或其部分内容派生的改编材料。

本章中的图像或其他第三方材料包含在本章的知识共享许可中，除非在材料的信用额度中另有说明。如果材料不包括在本章的知识共享许可中，并且您的预期使用不被法律法规允许或超出了允许的使用范围，您将需要直接从版权所有者处获得许可。

Footnotes 1

www.open-mpi.org/projects/hwloc 和 https://github.com/RRZE-HPC/likwid 。

2

我们可以使用 numatl-show 查询强制执行的 NUMA 策略。

3

http://man7.org/linux/man-pages/man3/numa.3.html 。

4

http://man7.org/linux/man-pages/man8/numactl.8.html 。

第一部分

第二部分