C---数据并行教程-一-C++ 数据并行教程（一）一、介绍本章通过涵盖核心概念(包括术语)奠定了基础，当我们学习如

C++ 数据并行教程（一）

原文：Data Parallel C++

协议：CC BY-NC-SA 4.0

一、介绍

本章通过涵盖核心概念(包括术语)奠定了基础，当我们学习如何使用数据并行性来加速 C++ 程序时，这些核心概念对于保持头脑中的新鲜感至关重要。

img/489625_1_En_1_Figa_HTML.png

This chapter lays the foundation by covering core concepts, including terminology, that are critical to have fresh in our minds as we learn how to accelerate C++ programs using data parallelism.

C++ 中的数据并行性支持在现代异构系统中访问并行资源。单个 C++ 应用程序可以使用任何设备组合，包括 GPU、CPU、FPGAs 和人工智能专用集成电路(ASICs)，这些设备组合都适用于当前的问题。

这本书教授使用 C++ 和 SYCL 的数据并行编程。

SYCL(读作镰刀)是一个行业驱动的 Khronos 标准，它为异构系统的 C++ 增加了数据并行性。SYCL 程序在与支持 SYCL 的 C++ 编译器(如本书中使用的开源数据并行 C++ (DPC++ ))编译器)配合使用时性能最佳。SYCL 不是首字母缩写词；SYCL 只是一个名字。

DPC++ 是一个开源编译器项目，最初由英特尔员工创建，致力于在 C++ 中大力支持数据并行。DPC++ 编译器基于 SYCL、一些扩展、 ¹ 和广泛的异构支持，包括 GPU、CPU 和 FPGA 设备。除了开源版本的 DPC++ 之外，英特尔 oneAPI 工具包中还提供了商业版本。

基于 SYCL 实现的特性受到 DPC++ 编译器的开源和商业版本的支持。本书中的所有例子都可以用或者版本的 DPC++ 编译器编译和工作，而且几乎所有例子都可以用最新的 SYCL 编译器编译。在发布时，我们会仔细注意哪些地方使用了特定于 DPC++ 的扩展。

读这本书，而不是说明书

没有人想被告知“去读说明书吧！”规范很难读懂，SYCL 规范也不例外。像每一个伟大的语言规范一样，它在动机、用法和教学方面都非常精确和简洁。这本书是教授 SYCL 和使用 DPC++ 编译器的“学习指南”。

正如序言中提到的，这本书无法一次性解释所有的事情。因此，这一章做了其他章节都不会做的事情:代码示例包含的编程结构在以后的章节中才会解释。我们应该试着不要完全理解第一章中的代码示例，相信每一章都会变得更好。

SYCL 1.2.1 与 SYCL 2020 和 DPC++ 的对比

在本书付印之际，SYCL 2020 临时规范已经公开征求意见。随着时间的推移，将会出现当前 SYCL 1.2.1 标准的继任者。这个预期的继任者被非正式地称为 SYCL 2020。虽然说这本书教授 SYCL 2020 很好，但这是不可能的，因为该标准尚不存在。

这本书教授 SYCL 扩展，以估计 SYCL 在未来的位置。这些扩展是在 DPC++ 编译器项目中实现的。几乎所有在 DPC++ 中实现的扩展都是临时 SYCL 2020 规范中的新特性。DPC++ 支持的值得注意的新特性是 USM、子组、C++17 支持的语法简化(称为 class 类模板参数演绎)，以及无需命名即可使用匿名 lambdas 的能力。

在发布时，没有任何 SYCL 编译器(包括 DPC++)实现了 SYCL 2020 临时规范中的所有功能。

本书中使用的一些特性是特定于 DPC++ 编译器的。其中许多特性最初是英特尔对 SYCL 的扩展，后来被纳入 SYCL 2020 临时规范，在某些情况下，它们的语法在标准化过程中略有变化。其他功能仍在开发或讨论中，可能会包含在未来的 SYCL 标准中，它们的语法也可能类似地被修改。在语言开发过程中，这样的语法变化实际上是非常可取的，因为我们希望特性不断发展和改进，以满足更广泛的开发人员群体的需求和各种设备的功能。本书中的所有代码示例都使用 DPC++ 语法来确保与 DPC++ 编译器的兼容性。

在努力接近 SYCL 的发展方向的同时，随着标准的发展，几乎肯定需要对本书中的信息进行调整，以与标准保持一致。更新信息的重要资源包括 GitHub 一书和勘误表，可从该书的网页(www . a press . com/9781484255735)以及在线 oneAPI DPC++ 语言参考( tinyurl. com/ dpcppref )中找到。

获得 DPC++ 编译器

DPC++ 可以从 GitHub 资源库( github. com/ intel/ llvm )获得。可以在Intel . GitHub . io/llvm-docs/GetStartedGuide . html找到 DPC++ 入门指南，包括如何使用 GitHub 的克隆版本构建开源编译器。

还有 DPC++ 编译器的捆绑版本，增加了用于 DPC++ 编程和支持的其他工具和库，作为更大的 oneAPI 项目的一部分提供。该项目带来了对异构系统的广泛支持，包括库、调试器和其他工具，称为 oneAPI。包括 DPC++ 在内的 oneAPI 工具都是免费提供的(oneAPI . com/implementations)。官方 oneAPI DPC++ 编译器文档，包括扩展列表，可以在Intel . github . io/llvm-docs找到。

这本书的在线伴侣，oneAPI DPC++ 语言参考 online ，是一个很好的资源，可以在这本书的基础上获得更多正式的细节。

GitHub 图书

很快我们会遇到图 1-1 中的代码。如果我们想避免全部键入，我们可以很容易地从 GitHub 存储库中下载本书中的所有示例(www . a press . com/9781484255735—寻找本书的服务:源代码)。该存储库包括带有构建文件的完整代码，因为大多数代码清单省略了重复的或不必要的细节。存储库中有这些例子的最新版本，如果有更新的话，这是很方便的。

img/489625_1_En_1_Fig1_HTML.png

图 1-1

你好数据并行编程

你好，世界！和 SYCL 程序剖析

图 1-1 显示了一个样本 SYCL 程序。使用 DPC++ 编译器编译并运行它，会打印出以下内容:

你好，世界！(还有一些额外的文本留给运行它的人去体验)

在第四章结束时，我们会完全理解这个特殊的例子。在此之前，我们可以观察定义所有 SYCL 构造所需的<CL/sycl.hpp>(第 1 行)的单个 include。所有 sycl 构造都存在于一个名为 SYCL 的名称空间中:

第 3 行让我们避免一遍又一遍地写sycl::。
第 11 行为指向特定设备的工作请求建立了一个队列(第二章)。
第 13 行为与设备共享的数据创建一个分配(第三章)。
第 16 行将工作排入设备队列(第章第 4 )。
第 17 行是将在设备上运行的唯一一行代码。所有其他代码都在主机(CPU)上运行。

第 17 行是我们希望在设备上运行的内核代码。内核代码减少一个字符。借助于parallel_for()的能力，内核在我们的秘密字符串中的每个字符上运行，以便将它解码成result字符串。所需的工作没有顺序，一旦parallel_for将工作排队，它实际上相对于主程序异步运行。在查看结果之前有一个等待(第 18 行)是很关键的，以确保内核已经完成，因为在这个特定的例子中，我们使用了一个方便的特性(统一共享内存，第六章)。如果没有等待，输出可能会在所有字符被解密之前发生。还有更多要讨论的，但那是后面章节的工作。

队列和操作

第二章将讨论队列和动作，但是我们现在可以从一个简单的解释开始。队列是唯一允许应用程序在设备上直接完成工作的连接。有两种类型的操作可以放入队列中:(a)要执行的代码和(b)内存操作。要执行的代码通过single_task、parallel_for(用于图 1-1 )或parallel_for_work_group来表示。内存操作执行主机和设备之间的复制操作或填充操作来初始化内存。如果我们寻求比自动为我们做的更多的控制，我们只需要使用内存操作。这些都将在本书后面从第二章开始讨论。现在，我们应该意识到，队列是允许我们命令设备的连接，我们有一组可用于放入队列的操作来执行代码和移动数据。理解被请求的动作被放入队列而不等待也是非常重要的。在将动作提交到队列中之后，主机继续执行程序，而设备将最终异步地执行通过队列请求的动作。

队列将我们与设备联系起来。

我们将动作提交到这些队列中，请求计算工作和数据移动。

动作异步发生。

这完全是关于并行性

因为用 C++ 进行数据并行编程完全是关于并行性的，所以让我们从这个关键概念开始。并行编程的目标是更快地计算一些东西。事实证明这有两个方面:增加吞吐量和减少延迟。

生产能力

当我们在一定的时间内完成更多的工作时，程序的吞吐量就会增加。像流水线这样的技术实际上可能会延长完成一项工作所需的时间，以允许工作重叠，从而导致单位时间内完成更多的工作。人类在一起工作时经常会遇到这种情况。分担工作的行为本身就包含了协调的开销，这通常会拖慢做一件事情的时间。然而，多人的力量导致更多的吞吐量。计算机也不例外——将工作分散到更多的处理核心会增加每个工作单元的开销，这可能会导致一些延迟，但目标是完成更多的总工作，因为我们有更多的处理核心一起工作。

潜伏

如果我们想更快地完成一件事——例如，分析一个语音命令并制定一个响应，该怎么办？如果我们只关心吞吐量，响应时间可能会变得难以忍受。减少延迟的概念要求我们将一项工作分解成可以并行处理的部分。对于吞吐量，图像处理可能会将整个图像分配给不同的处理单元，在这种情况下，我们的目标可能是优化每秒图像数。对于延迟，图像处理可能会将图像中的每个像素分配给不同的处理核心，在这种情况下，我们的目标可能是最大化单幅图像每秒的像素。

平行思考

成功的并行程序员在他们的编程中使用这两种技术。这是我们寻求平行思考的开始。

我们希望调整我们的思维，首先考虑在我们的算法和应用程序中哪里可以找到并行性。我们还会思考表达并行性的不同方式如何影响我们最终实现的性能。那是一次要接受的很多东西。寻求思考并行成为并行程序员一生的旅程。我们可以在这里学到一些技巧。

阿姆达尔和古斯塔夫森

阿姆达尔定律是由超级计算机先驱吉恩·阿姆达尔在 1967 年提出的，是一个预测使用多个处理器时理论上最大加速的公式。Amdahl 哀叹道，并行性的最大收益受限于(1/(1-p))，其中p是并行运行的程序的一部分。如果我们只并行运行程序的三分之二，那么程序最多可以加速 3 倍。我们绝对需要这个概念深入人心！这是因为无论我们让程序的三分之二运行得多快，另外三分之一仍然需要同样的时间来完成。即使我们添加 100 个 GPU，我们也只能获得 3 倍的性能提升。

多年来，一些人认为这证明了并行计算不会有成效。1988 年，约翰·古斯塔夫森发表了一篇题为“重新评估阿姆达尔定律”的文章。他观察到并行性不是用来加速固定工作负载的，而是用来支持工作的扩展。人类也经历同样的事情。一个送货人不可能在更多人和卡车的帮助下更快地运送一个包裹。然而，一百个人和一辆卡车可以比一个司机驾驶一辆卡车更快地运送一百个包裹。多个驱动程序肯定会增加吞吐量，通常还会减少包裹交付的延迟。阿姆达尔定律告诉我们，一个司机不可能通过增加九十九个司机自己的卡车来更快地运送一个包裹。古斯塔夫森注意到，有了这些额外的司机和卡车，就有机会更快地运送 100 个包裹。

缩放比例

“缩放”一词出现在我们之前的讨论中。缩放是一种衡量当额外计算可用时程序加速多少(简称为“加速”)的方法。如果 100 个包裹与一个包裹在同一时间交付，只要有 100 辆卡车和司机，而不是一辆卡车和司机，就可以实现完美的加速。当然，事实并非如此。在某种程度上，存在一个限制速度提升的瓶颈。配送中心可能没有一百个地方供卡车停靠。在计算机程序中，瓶颈通常涉及到将数据移动到需要处理的地方。向一百辆卡车分发数据类似于向一百个处理核心分发数据。分发的行为不是瞬间的。第三章将开始我们探索如何在异构系统中将数据分布到需要的地方的旅程。我们必须知道数据分发是有成本的，而这种成本会影响我们对应用程序的可伸缩性的预期。

异构系统

短语“异构系统”偷偷溜进了前一段。出于我们的目的，异构系统是任何包含多种类型的计算设备的系统。例如，具有中央处理单元(CPU)和图形处理单元(GPU)的系统是异构系统。CPU 通常只是被称为处理器，尽管当我们把异构系统中的所有处理单元都称为计算处理器时，这可能会令人混淆。为了避免混淆，SYCL 将处理单元称为设备。第二章将开始讨论如何将工作(计算)导向异构系统中的特定设备。

GPU 已经发展成为高性能计算设备，因此有时被称为通用 GPU 或 GPGPUs。出于异构编程的目的，我们可以简单地假设我们正在编写这样强大的 GPGPUs，并将它们称为 GPU。

今天，异构系统中的设备集合可以包括 CPU、GPU、FPGAs(现场可编程门阵列)、DSP(数字信号处理器)、ASICs(专用集成电路)和 AI 芯片(图形、神经形态等)。).

这种设备的设计通常包括复制计算处理器(多处理器)和增加与存储器等数据源的连接(增加带宽)。第一种，多重处理，对于提高吞吐量特别有用。在我们的类比中，这是通过增加额外的司机和卡车来完成的。后者，更高的数据带宽，对于减少延迟特别有用。在我们的类比中，这是通过更多的装载码头来实现的，以使卡车能够平行满载。

拥有多种类型的设备，每种设备具有不同的架构，因此具有不同的特性，这导致每种设备具有不同的编程和优化需求。这成为 SYCL、DPC++ 编译器以及本书大部分内容的动机。

SYCL 的创建是为了应对异构系统的 C++ 数据并行编程的挑战。

数据并行编程

从这本书的标题开始,“数据并行编程”这个短语就一直没有得到解释。数据并行编程侧重于并行性，可以将并行性想象为一组并行操作的数据。这种重心的转移就像古斯塔夫森对阿姆达尔。我们需要交付 100 个包(实际上是大量数据)，以便在 100 辆卡车和司机之间分配工作。关键概念归结为我们应该划分什么。我们应该处理整个图像还是在更小的图块中处理它们还是逐个像素地处理它们？我们应该将一组对象作为一个单独的集合来分析，还是作为一组更小的对象组来分析，还是逐个对象地分析？

任何使用 SYCL 和 DPC++ 的并行程序员都有责任选择正确的工作分工，并将工作有效地映射到计算资源上。第四章开始了这一讨论，并贯穿全书的其余部分。

DPC++ 和 SYCL 的关键属性

每个 DPC++(或 SYCL)程序也是一个 C++ 程序。SYCL 和 DPC++ 都不依赖于 C++ 的任何语言变化。两者都可以用模板和 lambda 函数完全实现。

SYCL 编译器 ² 存在的原因是以一种依赖于 SYCL 规范的内置知识的方式来优化代码。缺乏任何 SYCL 内置知识的标准 C++ 编译器无法获得与支持 SYCL 的编译器相同的性能水平。

接下来，我们将检查 DPC++ 和 SYCL 的关键属性:单源样式、主机、设备、内核代码和异步任务图。

单源

程序可以是单源的，这意味着同一个翻译单元 ³ 既包含定义要在设备上执行的计算内核的代码，也包含协调这些计算内核的执行的主机代码。第二章从更详细地了解这种能力开始。如果我们愿意，我们仍然可以将我们的程序源分成不同的文件和主机和设备代码的翻译单元，但关键是我们不必这样做！

圣体

每个程序都是从在主机上运行开始的，程序中的大部分行代码通常是给主机的。迄今为止，主机一直是 CPU。标准对此没有要求，所以我们小心翼翼地将其描述为主机。这似乎不太可能是 CPU 以外的任何东西，因为主机需要完全支持 C++17 才能支持所有的 DPC++ 和 SYCL 程序。我们很快就会看到，设备不需要支持所有的 C++17。

设备

在一个程序中使用多个设备是异构编程的原因。这就是为什么自从几页前对异构系统的解释以来，设备这个词一直在本章中反复出现。我们已经了解到，异构系统中的设备集合可以包括 GPU、FPGAs、DSP、ASICs、CPU 和 AI 芯片，但不限于任何固定列表。

设备是 SYCL 承诺的加速卸载的目标。卸载计算的想法通常是将工作转移到可以加速工作完成的设备。我们不得不担心弥补移动数据所损失的时间，这是一个需要我们不断思考的话题。

共享设备

在一个有设备的系统上，比如一个 GPU，我们可以想象两个或者更多的程序正在运行并且想要使用一个设备。它们不必是使用 SYCL 或 DPC++ 的程序。如果另一个程序正在使用该设备，则该设备在处理程序时可能会遇到延迟。这与 C++ 程序中通常用于 CPU 的原理是一样的。如果我们在 CPU 上运行太多活动程序(邮件、浏览器、病毒扫描、视频编辑、照片编辑等)，任何系统都可能过载。)一下子。

在超级计算机上，当节点(CPUs 所有连接的设备)被专门授予单个应用程序时，共享通常不是一个问题。在非超级计算机系统上，我们可以注意到，如果有多个应用程序同时使用相同的设备，数据并行 C++ 程序的性能可能会受到影响。

一切仍然工作，没有我们需要做不同的编程。

内核代码

设备的代码被指定为内核。这不是 SYCL 或 DPC++ 独有的概念:它是其他卸载加速语言(包括 OpenCL 和 CUDA)的核心概念。

内核代码有一定的限制，以允许更广泛的设备支持和大规模并行。内核代码中不支持的特性列表包括动态多态、动态内存分配(因此没有使用 new 或 delete 操作符的对象管理)、静态变量、函数指针、运行时类型信息(RTTI)和异常处理。不允许从内核代码中调用虚拟成员函数和变量函数。内核代码中不允许递归。

第三章将描述如何在内核被调用之前和之后进行内存分配，从而确保内核专注于大规模并行计算。第五章将描述与设备相关的异常处理。

C++ 的其余部分是内核中的公平游戏，包括 lambdas、操作符重载、模板、类和静态多态。我们还可以与主机共享数据(参见第三章)并共享(非全局)主机变量的只读值(通过 lambda 捕获)。

内核:向量加法(DAXPY)

任何从事计算复杂代码工作的程序员都应该对内核很熟悉。考虑实现 DAXPY，它代表“双精度 A 乘以 X 加 Y”，这是几十年来的经典。图 1-2 显示了用现代 Fortran、C/C++ 和 SYCL 实现的 DAXPY。令人惊讶的是，计算行(第 3 行)实际上是相同的。第 4 和 10 章将详细解释内核。图 1-2 应该有助于消除对内核难以理解的任何担忧——即使术语对我们来说是新的，它们也应该感觉熟悉。

img/489625_1_En_1_Fig2_HTML.png

图 1-2

Fortran、C++ 和 SYCL 中的 DAXPY 计算

异步任务图

使用 SYCL/DPC++ 编程的异步特性必须而不是被忽略。理解异步编程是至关重要的，原因有两个:(1)正确的使用会给我们带来更好的性能(更好的伸缩性),以及(2)错误会导致并行编程错误(通常是竞争条件),使我们的应用程序不可靠。

异步的本质是因为工作是通过请求动作的“队列”转移到设备上的。宿主程序将请求的动作提交到一个队列中，程序继续运行，不等待任何结果。这个无等待很重要，这样我们就可以努力让计算资源(设备和主机)一直保持忙碌。如果我们必须等待，那将会束缚主机，而不是让主机做有用的工作。它还会在设备完成时产生串行瓶颈，直到我们排队等待新的工作。如前所述，阿姆达尔定律惩罚我们没有平行工作的时间。我们需要构建我们的程序，以便在设备繁忙时将数据移入和移出设备，并在工作可用时保持设备和主机的所有计算能力繁忙。如果做不到这一点，将会给我们带来阿姆达尔法则的诅咒。

第四章将开始讨论把我们的程序想成一个异步任务图，第八章大大扩展了这个概念。

我们犯错时的竞争条件

在我们的第一个代码示例(图 1-1 )中，我们特别在第 18 行做了一个“等待”，以防止第 20 行在result的值可用之前将其写出。我们必须记住这种异步行为。在同一个代码示例中还做了另一件微妙的事情——第 14 行使用std::memcpy来加载输入。因为std::memcpy在主机上运行，所以第 16 行和之后的代码直到第 15 行完成后才执行。在阅读完第三章之后，我们可能会尝试将其改为使用myQ.memcpy(使用 SYCL)。我们已经在第 8 行的图 1-3 中完成了。因为这是一个队列提交，所以不能保证它会在第 10 行之前完成。这就产生了一个竞争条件，这是一种并行编程错误。当程序的两个部分不协调地访问相同的数据时，就存在争用情况。因为我们希望使用第 8 行写入数据，然后在第 10 行读取数据，所以我们不希望出现第 17 行在第 8 行完成之前执行的竞争！这样的竞争条件会使我们的程序不可预测——我们的程序可能在不同的运行和不同的系统上得到不同的结果。解决这个问题的方法是通过在第 8 行末尾添加.wait()来明确等待myQ.memcpy完成后再继续。这不是最好的解决办法。我们可以使用事件依赖来解决这个问题(第八章)。将队列创建为有序队列还会在memcpy和parallel_for.之间添加一个隐含的依赖关系。作为替代，在第七章中，我们将看到如何使用缓冲区和访问器编程风格来让 SYCL 管理依赖关系并自动等待我们。

img/489625_1_En_1_Fig3_HTML.png

图 1-3

添加一个竞争条件来说明关于异步的一点

添加一个wait()强制在memcpy和内核之间进行主机同步，这与之前让设备一直忙碌的建议背道而驰。本书的大部分内容涵盖了不同的选项和权衡，平衡了程序的简单性和系统的有效使用。

为了帮助检测程序(包括内核)中的数据争用情况，Intel Inspector(在“获取 DPC++ 编译器”中提到的 oneAPI 工具中提供)等工具可能会有所帮助。这些工具使用的有些复杂的方法通常不能在所有设备上工作。检测竞争条件最好的方法是让所有的内核都在一个 CPU 上运行，这可以作为开发工作中的一种调试技术来完成。这个调试技巧在第二章中作为方法#2 讨论。

第四章会告诉我们“lambdas 不被认为是有害的。”为了更好地使用 DPC++、SYCL 和现代 C++，我们应该熟悉 lambda 函数。

C++ Lambda 函数

并行编程技术大量使用的现代 C++ 的一个特性是 lambda 函数。内核(在设备上运行的代码)可以有多种表达方式，最常见的是 lambda 函数。第十章讨论了内核可以采取的各种形式，包括 lambda 函数。在这里，我们复习了 C++ lambda 函数以及一些关于定义内核的注意事项。第十章在我们在中间章节中学习了更多关于 SYCL 的知识后，将详细阐述内核方面。

图 1-3 中的代码具有 lambda 函数。我们可以看到它，因为它从非常确定的[=]开始。在 C++ 中，lambda 以方括号开始，右方括号前的信息表示如何捕获在 lambda 中使用的变量，但这些变量没有作为参数显式传递给它。对于内核，捕获必须是值为的*，这由括号内包含的等号表示。*

对 lambda 表达式的支持是在 C++11 中引入的。它们用于创建匿名函数对象(尽管我们可以将它们赋给命名变量)，这些对象可以从封闭范围中捕获变量。C++ lambda 表达式的基本语法是

img/489625_1_En_1_Figb_HTML.png

在哪里

捕获列表是一个逗号分隔的捕获列表。我们通过在捕获列表中列出变量名来按值捕获变量。我们通过引用捕获一个变量，在它前面加上一个&符号，例如，&v **。**还有适用于所有作用域内自动变量的简写:[=]用于通过值和通过引用捕获主体中使用的所有自动变量，[&]用于通过引用捕获主体和当前对象中使用的所有自动变量，[]什么都不捕获。在 SYCL 中，[=]几乎总是被使用，因为在内核中不允许通过引用来捕获变量。根据 C++ 标准，全局变量在 lambda 中不是被捕获的。非全局静态变量可以在内核中使用，但是只能在const中使用。
params是函数参数列表，就像命名函数一样。SYCL 提供了参数来标识内核被调用来处理的元素:这可以是唯一的 id(一维的)或 2D 或 3D id。这些将在第四章中讨论。
ret是返回类型。如果未指定->ret，则从返回语句中推断出来。缺少 return 语句，或者 return 没有值，意味着 return 类型为void。SYCL 内核必须总是有一个返回类型void，所以我们不应该用这个语法来指定内核的返回类型。
body是函数体。对于 SYCL 内核，这个内核的内容有一些限制(参见本章前面的“内核代码”一节)。

图 1-4 显示了一个 C++ lambda 表达式，它通过值捕捉一个变量i，通过引用捕捉另一个变量j。它还有一个参数k0和另一个通过引用接收的参数l0。运行该示例将产生如图 1-5 所示的输出。

img/489625_1_En_1_Fig5_HTML.png

图 1-5

图 1-4 中 lambda 函数演示代码的输出

img/489625_1_En_1_Fig4_HTML.png

图 1-4

C++ 代码中的 Lambda 函数

我们可以把 lambda 表达式看作一个函数对象的实例，但是编译器为我们创建了类定义。例如，我们在前面的例子中使用的 lambda 表达式类似于图 1-6 中所示的类的实例。无论我们在哪里使用 C++ lambda 表达式，我们都可以用一个函数对象的实例来代替它，如图 1-6 所示。

img/489625_1_En_1_Fig6_HTML.png

图 1-6

函数对象而不是 lambda(在第十章中有更多关于这方面的内容)

每当我们定义一个函数对象时，我们都需要给它赋一个名字(图 1-6 中的函子)。内嵌表达的 Lambdas(如图 1-4 所示)是匿名的，因为它们不需要名字。

可移植性和直接编程

可移植性是 SYCL 和 DPC++ 的一个关键目标；但是，两者都不能保证。一门语言和编译器所能做的就是当我们想在应用程序中实现可移植性时，让它变得更容易一些。

可移植性是一个复杂的话题，包括功能可移植性和性能可移植性的概念。有了功能上的可移植性，我们希望我们的程序可以在各种各样的平台上同等地编译和运行。有了性能可移植性，我们希望我们的程序能在各种平台上获得合理的性能。虽然这是一个相当软的定义，但反过来可能更清楚——我们不希望编写一个在一个平台上运行超快的程序，却发现它在另一个平台上慢得不合理。事实上，我们更希望它能充分利用运行它的任何平台。考虑到异构系统中各种各样的设备，性能可移植性需要我们作为程序员付出巨大的努力。

幸运的是，SYCL 定义了一种可以提高性能可移植性的编码方式。首先，通用内核可以在任何地方运行。在有限的情况下，这可能就足够了。更常见的是，可以为不同类型的设备编写几个版本的重要内核。具体来说，一个内核可能有一个通用的 GPU 版本和一个通用的 CPU 版本。有时候，我们可能想为特定的设备(比如特定的 GPU)专门化我们的内核。当这种情况发生时，我们可以编写多个版本，并针对不同的 GPU 模型进行专门化。或者我们可以参数化一个版本，使用 GPU 的属性来修改我们的 GPU 内核如何运行，以适应现有的 GPU。

当我们作为程序员自己负责设计一个有效的性能移植计划时，SYCL 定义了允许我们实现计划的结构。如前所述，功能可以分层，首先为所有设备提供一个内核，然后根据需要逐渐引入更多、更专业的内核版本。这听起来很棒，但是程序的整体流程也会产生深远的影响，因为数据移动和整体算法选择很重要。了解了这一点，就能理解为什么没有人会声称 SYCL(或其他直接编程解决方案)解决了性能可移植性。然而，它是我们工具箱中帮助我们应对这些挑战的工具。

并发性与并行性

术语并发和平行是不等价的，尽管它们有时会被误解。重要的是要知道，并发性所需的任何编程考虑对于并行性也很重要。

术语并发指的是可以前进但不一定在同一时刻的代码。在我们的计算机上，如果我们有一个打开的Mail程序和一个Web Browser，那么它们是并发运行的。在只有一个处理器的系统上，通过时间分片过程(在运行每个程序之间快速来回切换)，可以发生并发。

Tip

并发性所需的任何编程考虑对于并行性也很重要。

术语并行是指代码可以在同一时刻前进。并行性要求系统实际上一次可以做多件事情。异构系统总是可以并行地做事情，这是由它至少具有两个计算设备的本质决定的。当然，SYCL 程序不需要异构系统，因为它可以在只有主机的系统上运行。今天，任何主机系统都不可能不具备并行执行的能力。

代码的并发执行通常面临与代码的并行执行相同的问题，因为任何特定的代码序列都不能假定它是改变世界(数据位置、I/O 等)的唯一代码。).

摘要

本章提供了 SYCL 和 DPC++ 所需的术语，并提供了对 SYCL 和 DPC++ 至关重要的并行编程和 C++ 的关键方面的更新。第 2 、 3 和 4 章详细阐述了 SYCL 编程的三个关键:需要给设备分配工作(发送代码以在其上运行)、提供数据(发送数据以在其上使用)以及拥有编写代码的方法(内核)。

Creative Commons

开放存取本章根据知识共享署名 4.0 国际许可证(http://Creative Commons . org/licenses/by/4.0/)的条款获得许可，该许可证允许以任何媒体或格式使用、共享、改编、分发和复制，只要您适当注明原作者和来源，提供知识共享许可证的链接并指明是否进行了更改。

Footnotes 1

DPC++ 团队很快指出，他们希望在将来的某个时候，SYCL 标准会考虑并接受他们所有的扩展。

称之为支持 SYCL 的 C++ 编译器可能更正确。

我们可以只说“文件”，但这在这里并不完全正确。翻译单元是编译器的实际输入，由 C 预处理器处理后的源文件构成，用于内联头文件和扩展宏。

二、代码执行的地方

img/489625_1_En_2_Figa_HTML.gif

并行编程并不是真的在快车道上行驶。这实际上是在 ?? 所有的车道上开快车。这一章是关于让我们能够把我们的代码放在我们能放的任何地方。只要有意义，我们就会选择启用异构系统中的所有计算资源。因此，我们需要知道这些计算资源藏在哪里(找到它们)，并让它们发挥作用(在其上执行我们的代码)。

我们可以控制代码在哪里执行——换句话说，我们可以控制哪些设备用于哪些内核。SYCL 为异构编程提供了一个框架，其中代码可以在主机 CPU 和设备上混合执行。决定代码在哪里执行的机制对于我们理解和使用非常重要。

本章描述了代码可以在哪里执行，何时执行，以及用于控制执行位置的机制。第三章将描述如何管理数据，以便它到达我们执行代码的地方，然后第四章回到代码本身，讨论内核的编写。

单源

SYCL 程序可以是单源的，这意味着同一个翻译单元(通常是一个源文件及其头文件)既包含定义要在 SYCL 设备上执行的计算内核的代码，也包含协调这些内核执行的主机代码。图 2-1 以图形方式显示了这两条代码路径，图 2-2 提供了一个标记了主机和设备代码区域的示例应用。

将设备和宿主代码组合到一个源文件(或翻译单元)中，可以使理解和维护异构应用程序变得更加容易。这种组合还提供了改进的语言类型安全性，并能使我们的代码得到更多的编译器优化。

img/489625_1_En_2_Fig2_HTML.png

图 2-2

简单 SYCL 程序

img/489625_1_En_2_Fig1_HTML.png

图 2-1

单源代码包含主机代码(运行在 CPU 上)和设备代码(运行在 SYCL 设备上)

主机代码

应用程序包含 C++ 宿主代码，由操作系统启动应用程序的 CPU 执行。宿主代码是应用程序的主干，它定义和控制向可用设备分配工作。它也是一个接口，通过它我们可以定义应该由运行时管理的数据和依赖关系。

宿主代码是标准的 C++，增加了特定于 SYCL 的构造和类，这些构造和类被设计成可作为 C++ 库来实现。这使得推断宿主代码中允许的内容(C++ 中允许的任何内容)变得更加容易，并且可以简化与构建系统的集成。

SYCL 应用程序是标准的 C++，增加了可以作为 C++ 库实现的结构。

SYCL 编译器可以通过“理解”这些结构为程序提供更高的性能。

应用程序中的主机代码协调数据移动和设备的计算卸载，但也可以自己执行计算密集型工作，并可以像任何 C++ 应用程序一样使用库。

设备码

设备对应于加速器或处理器，它们在概念上独立于执行主机代码的 CPU。如本章后面所述，实现必须将主机处理器也作为设备公开，但是主机处理器和设备应该被认为是逻辑上相互独立的。主机处理器运行本机 C++ 代码，而设备运行设备代码。

队列是一种机制，通过它可以将工作提交给设备以供将来执行。需要理解设备代码的三个重要属性:

它与主机代码异步执行。主机程序将设备代码提交给设备，只有当所有的执行依赖关系都满足时，运行时才会跟踪并启动该工作(更多信息请参见第三章)。主机程序执行在提交的工作在设备上开始之前进行，提供了设备上的执行与主机程序执行异步的属性，除非我们明确地将二者联系在一起。
为了能够在加速器设备上编译和实现性能，对设备代码有一些限制。例如，设备代码中不支持动态内存分配和运行时类型信息(RTTI ),因为它们会导致许多加速器的性能下降。第十章详细介绍了设备代码限制。
由 SYCL 定义的一些函数和查询只在设备代码中可用，因为它们只在那里有意义，例如，允许设备代码的执行实例在更大的数据并行范围内查询其位置的工作项标识符查询(在第四章中描述)。

一般来说，我们将包括提交到队列的设备代码的工作称为动作。在第三章中，我们将了解到动作不仅仅包括要执行的设备代码；动作还包括内存移动命令。在这一章中，由于我们关心的是动作的设备代码方面，我们将在大部分时间里特别提到设备代码。

选择设备

为了探索让我们控制设备代码将在何处执行的机制，我们将查看五个用例:

方法#1:在某个地方运行设备代码*，而我们并不关心使用的是哪个设备。这通常是开发的第一步，因为这是最简单的。*

** 方法#2:在主机设备上显式运行设备代码，这通常用于调试。保证主机设备在任何系统上都始终可用。

*   方法#3:将设备代码分派给 GPU 或另一个加速器。

*   方法#4:将设备代码分派给一组不同的设备，比如 GPU 和 FPGA。

*   方法#5:从更一般的器件类别中选择特定的器件，例如从一组可用的 FPGA 类型中选择特定类型的 FPGA。*

*开发人员通常会尽可能多地使用方法 2 来调试他们的代码，并且只有当代码已经用方法 2 尽可能多地进行了测试时，才转移到方法 3 到 5。

方法 1:在任何类型的设备上运行

当我们不关心我们的设备代码将在哪里运行时，很容易让运行时为我们选择。这种自动选择是为了在我们还不关心选择什么设备时，使开始编写和运行代码变得容易。这个设备选择没有考虑要运行的代码，所以应该被认为是一个任意的选择，可能不是最佳的。

在讨论设备的选择之前，即使是实现为我们选择的设备，我们也应该首先了解程序与设备交互的机制:队列。

行列

一个queue是一个抽象，动作被提交给它以便在单个设备上执行。图 2-3 和 2-4 中给出了queue等级的简化定义。动作通常是数据并行计算的启动，尽管其他命令也是可用的，例如当我们需要比运行时提供的自动移动更多的控制时，可以手动控制数据移动。提交给queue的工作可以在运行时跟踪的先决条件满足后执行，比如输入数据的可用性。这些先决条件包含在第 3 和 8 章中。

img/489625_1_En_2_Fig4_HTML.png

图 2-4

queue类中关键成员函数的简化定义

img/489625_1_En_2_Fig3_HTML.png

图 2-3

queue类的构造器的简化定义

一个queue被绑定到一个单独的device，这个绑定发生在队列的构造上。理解提交给队列的工作是在该队列所绑定的单个设备上执行的是很重要的。不能将队列映射到设备集合，因为这将导致哪个设备应该执行工作不明确。类似地，队列不能将提交给它的工作分散到多个设备上。相反，在一个队列和提交给该队列的工作将在其上执行的设备之间有一个明确的映射，如图 2-5 所示。

img/489625_1_En_2_Fig5_HTML.png

图 2-5

一个队列绑定到一个设备。提交到队列的工作在该设备上执行

一个程序中可以创建多个队列，按照我们对应用程序架构或编程风格所期望的任何方式。例如，可以创建多个队列，每个队列与不同的设备绑定，或者由主机程序中的不同线程使用。多个不同的队列可以绑定到单个设备，如 GPU，提交到这些不同的队列将导致在设备上执行组合工作。这方面的一个例子如图 2-6 所示。相反，正如我们前面提到的，一个队列不能绑定到多个设备，因为在请求执行动作的位置上不能有任何模糊性。例如，如果我们想要一个跨多个设备负载平衡的队列，那么我们可以在代码中创建这个抽象。

img/489625_1_En_2_Fig6_HTML.png

图 2-6

多个队列可以绑定到一个设备

因为队列被绑定到一个特定的设备，所以队列构造是代码中最常见的选择设备的方式，提交到队列的操作将在该设备上执行。构建队列时设备的选择是通过设备选择器抽象和相关的device_selector类实现的。

将队列绑定到设备，任何设备都可以

图 2-7 是一个没有指定队列应该绑定的设备的例子。不带任何参数的普通队列构造器(如图 2-7 所示)只是在幕后选择一些可用的设备。SYCL 保证至少有一个设备始终可用，即主机设备。主机设备可以运行内核代码，并且是主机程序在其上执行的处理器的抽象，因此总是存在。

img/489625_1_En_2_Fig7_HTML.png

图 2-7

通过简单的队列构造实现隐式默认设备选择器

使用简单的队列构造器是开始应用程序开发和启动并运行设备代码的简单方法。当它变得与我们的应用程序相关时，可以添加对绑定到队列的设备的选择的更多控制。

方法#2:使用主机设备进行开发和调试

主机设备可以被认为是使主机 CPU 能够像一个独立的设备一样工作，允许我们的设备代码执行，而不管系统中可用的加速器。我们总是有一些处理器运行主机程序，因此主机设备对我们的应用程序总是可用的。主机设备保证设备代码可以一直运行(不依赖于加速器硬件),并且有几个主要用途:

在没有任何加速器的低性能系统上开发设备代码:一个常见的用途是在本地系统上开发和测试设备代码，然后部署到 HPC 集群进行性能测试和优化。
使用非加速器工具调试设备代码:加速器通常通过较低级别的 API 公开，这些 API 可能没有主机 CPU 可用的高级调试工具。考虑到这一点，主机设备应该支持使用 CPU 开发人员熟悉的标准工具进行调试。
备份如果没有其他设备可用，保证设备代码可以功能性地执行:主机设备实现可能不会将性能作为主要目标，因此应被视为功能性备份，以确保设备代码可以始终在任何应用中执行，但不一定是性能的途径。

主机设备在功能上类似于硬件加速器设备，因为队列可以绑定到它，并且它可以执行设备代码。图 2-8 显示了主机设备如何成为系统中其他可用加速器的对等设备。它可以执行设备代码，就像 CPU、GPU 或 FPGA 一样，并且可以构建一个或多个绑定到它的队列。

img/489625_1_En_2_Fig8_HTML.png

图 2-8

始终可用的主机设备可以像任何加速器一样执行设备代码

应用程序可以选择创建一个绑定到主机设备的队列，方法是将host_selector显式传递给队列构造器，如图 2-9 所示。

img/489625_1_En_2_Fig9_HTML.png

图 2-9

使用host_selector类选择主机设备

即使没有特别请求(例如使用host_selector)，默认选择器也可能会选择主机设备，如图 2-7 中的输出所示。

定义了设备选择器类的几个变体，以便于我们定位设备类型。host_selector是这些选择器类的一个例子，我们将在接下来的章节中讨论其他的。

方法 3:使用 GPU(或其他加速器)

下一个例子展示了 GPU，但是任何类型的加速器都同样适用。为了更容易找到常见的加速器类别，设备被分成几大类，SYCL 为它们提供了内置的选择器类别。要从广泛的设备类型类别中进行选择，如“系统中可用的任何 GPU”，相应的代码非常简短，如本节所述。

设备类型

队列可以绑定到两大类设备:

已经描述过的主机设备。
加速器设备，如 GPU、FPGA 或 CPU 设备，用于加速我们应用程序中的工作负载。

加速器设备

有几大类促进剂类型:

CPU 设备
GPU 设备
加速器，它捕获既不是 CPU 设备也不是 GPU 设备的设备。这包括 FPGA 和 DSP 器件。

这些类别中的任何一个设备都可以很容易地使用内置的选择器类绑定到队列，这些选择器类可以传递给队列(和其他一些类)构造器。

设备选择器

必须绑定到特定设备的类，比如queue类，有可以接受从device_selector派生的类的构造器。例如，队列构造器是


queue( const device_selector &deviceSelector,
  const property_list &propList = {});

有五个内置的选择器用于各种常见设备:

DPC++ 中包含的一个附加选择器(SYCL 中没有)可以通过包含头"CL/sycl/intel/fpga_extensions.hpp":来获得

| `INTEL::fpga_selector` | 选择将自身标识为 FPGA 的设备。 |

可以使用内置选择器之一来构造队列，例如

img/489625_1_En_2_Figc_HTML.gif

图 2-10 显示了使用cpu_selector的完整示例，图 2-11 显示了队列与可用 CPU 设备的对应绑定。

图 2-12 显示了一个使用各种内置选择器类的例子，也展示了设备选择器与另一个接受构造上的device_selector的类(device)的使用。

img/489625_1_En_2_Fig12_HTML.png

图 2-12

来自各种设备选择器类的示例设备标识输出，以及设备选择器不仅可用于构建队列(在这种情况下，构建设备类实例)的演示

img/489625_1_En_2_Fig11_HTML.png

图 2-11

绑定到应用程序可用的 CPU 设备的队列

img/489625_1_En_2_Fig10_HTML.png

图 2-10

CPU 设备选择器示例

当设备选择失败时

如果在创建一个对象(比如队列)时使用了一个gpu_selector，并且没有可供运行时使用的 GPU 设备，那么选择器就会抛出一个runtime_error异常。对于所有的设备选择器类都是如此，因为如果没有所需类的设备可用，那么就会抛出一个runtime_error异常。对于复杂的应用程序来说，捕捉该错误并获取不太理想的(对于应用程序/算法)设备类作为替代是合理的。异常和错误处理将在第五章中详细讨论。

方法 4:使用多种设备

如图 2-5 和 2-6 所示，我们可以在一个应用中构建多个队列。我们可以将这些队列绑定到单个设备(队列的总工作量集中到单个设备中)、多个设备，或者这些设备的某种组合。图 2-13 提供了一个创建一个绑定到 GPU 的队列和另一个绑定到 FPGA 的队列的例子。相应的映射如图 2-14 所示。

img/489625_1_En_2_Fig14_HTML.png

图 2-14

GPU + FPGA 设备选择器示例:一个队列绑定到 GPU，另一个绑定到 FPGA

img/489625_1_En_2_Fig13_HTML.png

图 2-13

为 GPU 和 FPGA 设备创建队列

方法 5:定制(非常具体的)设备选择

我们现在来看看如何编写一个自定义选择器。除了本章中的示例，第十二章中还显示了一些示例。内置的设备选择器旨在让我们快速启动并运行代码。实际应用通常需要专门选择设备，例如从系统中可用的一组 GPU 类型中选择所需的 GPU。设备选择机制很容易扩展到任意复杂的逻辑，因此我们可以编写任何需要的代码来选择我们喜欢的设备。

`device_selector`基础类

所有的设备选择器都从抽象的device_selector基类派生，并在派生类中定义函数调用操作符:

img/489625_1_En_2_Figb_HTML.png

在从device_selector派生的类中定义这个操作符是定义任何复杂的选择逻辑所需要的，一旦我们知道了三件事:

对于运行时发现应用程序可以访问的每个设备，包括主机设备，都会自动调用一次函数调用运算符。
该运算符每次被调用时都返回一个整数。所有可用设备中得分最高的是选择器选择的设备。
函数调用操作符返回的负整数意味着不能选择所考虑的设备。

对设备进行评分的机制

我们有许多选项来创建对应于特定设备的整数分数，例如:

为特定设备类别返回正值。
设备名称和/或设备供应商字符串匹配。
基于设备或平台查询，我们在代码中可以想象的任何导致整数值的东西。

例如，选择英特尔 Arria 系列 FPGA 器件的一种可能方法如图 2-15 所示。

img/489625_1_En_2_Fig15_HTML.png

图 2-15

面向英特尔 Arria FPGA 设备的定制选择器

第十二章有更多关于器件选择的讨论和示例(图 12-2 和 12-3 )并更深入地讨论get_info方法。

在 CPU 上执行设备代码的三种途径

一个潜在的混淆来源是多种机制，通过这些机制，CPU 可以执行代码，如图 2-16 所示。

img/489625_1_En_2_Fig16_HTML.png

图 2-16

在 CPU 上执行的 SYCL 机制

CPU 执行的第一个也是最明显的路径是宿主代码，它或者是单源应用程序(宿主代码区域)的一部分，或者是链接到宿主代码并从宿主代码中调用，如库函数。

另外两条可用路径执行设备代码。设备代码的第一个 CPU 路径是通过主机设备，这在本章前面已经描述过了。它总是可用的，并被期望在执行主机代码的同一 CPU 上执行设备代码。

在 SYCL 中，在 CPU 上执行设备代码的第二条路径是可选的，它是一个针对性能进行了优化的 CPU 加速器设备。该设备通常由 OpenCL 等较低级别的运行时实现，因此其可用性可能取决于系统上安装的驱动程序和其他运行时。SYCL 描述了这一原理，其中主机设备旨在可使用本机 CPU 工具进行调试，而 CPU 设备可以构建在针对性能优化的实现上，而本机 CPU 调试器不可用。

虽然我们在本书中没有涉及到，但是当任务图中的先决条件得到满足时，有一种机制可以将常规 CPU 代码排队(图 2-16 的顶部)。这项高级功能可用于在任务图中执行常规 CPU 代码和设备代码，称为主机任务。

在设备上创建作品

应用程序通常包含主机代码和设备代码的组合。有几个类成员允许我们提交设备代码以供执行，因为这些工作分派构造是提交设备代码的唯一方式，它们允许我们容易地将设备代码与主机代码区分开。

本章的剩余部分介绍了一些工作分派结构，目的是帮助我们理解和识别设备代码和在主机处理器上本地执行的主机代码之间的区别。

任务图简介

SYCL 执行模型中的一个基本概念是节点图。该图中的每个节点(工作单元)都包含一个要在设备上执行的操作，最常见的操作是数据并行设备内核调用。图 2-17 显示了一个有四个节点的示例图，其中每个节点都可以被认为是一个设备内核调用。

图 2-17 中的节点具有依赖边，定义了何时开始执行节点的工作是合法的。依赖边通常是从数据依赖关系自动生成的，尽管我们可以在需要时手动添加额外的自定义依赖关系。例如，图中的节点 B 具有与节点 A 的依赖边。该边意味着在节点 B 的动作开始之前，节点 A 必须完成执行，并且最有可能(取决于依赖关系的细节)使生成的数据在节点 B 将执行的设备上可用。运行时完全与宿主程序的执行异步地控制依赖关系的解析和节点执行的触发。定义应用程序的节点图在本书中将被称为任务图，在第三章中会有更详细的介绍。

img/489625_1_En_2_Fig18_HTML.png

图 2-18

提交设备代码

img/489625_1_En_2_Fig17_HTML.png

图 2-17

任务图定义了要在一个或多个设备上执行的动作(与主机程序异步),还定义了确定何时执行动作是安全的依赖关系

设备代码在哪里？

有多种机制可用于定义将在设备上执行的代码，但一个简单的示例显示了如何识别此类代码。即使示例中的模式初看起来很复杂，但该模式在所有设备代码定义中保持不变，很快就成为第二天性。

作为最后一个参数传递给parallel_for的代码，在图 2-18 中定义为λ，是要在设备上执行的设备代码。本例中的parallel_for是让我们区分设备代码和主机代码的结构。parallel_for是一小组设备调度机制中的一个，所有成员都是handler类的成员，它们定义了要在设备上执行的代码。图 2-19 给出了handler等级的简化定义。

img/489625_1_En_2_Fig19_HTML.png

图 2-19

handler类中成员函数的简化定义

除了调用handler类的成员提交设备代码，还有queue类的成员允许提交工作。图 2-20 中显示的queue类成员是简化某些模式的快捷方式，我们将在以后的章节中看到这些快捷方式的使用。

img/489625_1_En_2_Fig20_HTML.png

图 2-20

queue类中成员函数的简化定义，作为handler类中等价函数的简写符号

行动

图 2-18 中的代码包含一个parallel_for，它定义了要在设备上执行的工作。parallel_for位于提交给queue的命令组(CG)内，queue定义了将要执行工作的设备。在命令组中，有两类代码:

恰好一个对动作的调用，该动作或者将设备代码排队等待执行，或者执行手动内存操作，例如copy。
建立依赖关系的宿主代码，定义运行时何时开始执行(1)中定义的工作是安全的，例如创建缓冲区的访问器(在第三章中描述)。

handler 类包含一小组成员函数，这些函数定义了执行任务图节点时要执行的操作。图 2-21 总结了这些动作。

img/489625_1_En_2_Fig21_HTML.png

图 2-21

调用设备代码或执行显式内存操作的操作

在一个命令组中只能调用图 2-21 中的一个动作(调用多个是错误的)，并且每个submit调用只能提交一个命令组到一个队列中。这样做的结果是，图 2-21 中的单个操作存在于每个任务图节点中，当满足节点依赖性并且运行时确定可以安全执行时，该操作将被执行。

一个命令组中必须有一个动作，例如内核启动或显式内存操作。

将来异步执行代码的想法是作为主机程序的一部分在 CPU 上运行的代码和将来在满足依赖性时运行的设备代码之间的关键区别。命令组通常包含每个类别的代码，定义依赖关系的代码作为宿主程序的一部分运行(以便运行时知道依赖关系是什么),设备代码在依赖关系得到满足后运行。

图 2-22 中有三类代码:

img/489625_1_En_2_Fig22_HTML.png

图 2-22

提交设备代码

宿主代码:驱动应用程序，包括创建和管理数据缓冲区，以及将工作提交到队列中，以在任务图中形成新的节点来进行异步执行。
命令组中的主机代码:该代码运行在主机代码正在执行的处理器上，并在submit调用返回之前立即执行。例如，这段代码通过创建访问器来设置节点依赖关系。任何任意的 CPU 代码都可以在这里执行，但是最佳实践是将其限制为配置节点依赖关系的代码。
一个动作:图 2-21 中列出的任何动作都可以包含在一个命令组中，它定义了未来满足节点需求时异步执行的工作(由(2)设置)。

要了解应用程序中的代码何时运行，请注意传递给图 2-21 中列出的启动设备代码执行的动作的任何东西，或者图 2-21 中列出的显式内存操作，将在满足 DAG 节点依赖关系后异步执行*。所有其他代码作为宿主程序的一部分立即运行，正如典型的 C++ 代码所预期的那样。*

撤退

通常一个命令组是在我们提交给它的命令队列中执行的。然而，可能存在命令组未能提交到队列的情况(例如，当所请求的工作大小对于设备的限制来说太大时)，或者当成功提交的操作不能开始执行时(例如，当硬件设备发生故障时)。为了处理这种情况，可以为要执行的命令组指定一个后备队列。作者不推荐这种错误管理技术，因为它提供的控制很少，相反，我们建议捕捉和管理初始错误，如第五章所述。我们在这里简单介绍一下回退队列，因为有些人更喜欢这种风格，它是 SYCL 中众所周知的一部分。

这种回退方式适用于机器上存在的设备的失败队列提交。这不是解决加速器不存在问题的后备机制。在没有 GPU 设备的系统上，图 2-23 中的程序会在Q声明(试图构造)中抛出一个错误，表明“没有请求类型的设备可用”

img/489625_1_En_2_Fig23_HTML.png

图 2-23

回退队列示例

基于现有设备的回退主题将在第十二章中讨论。

图 2-23 显示了由于所要求的工作组规模，将无法在某些 GPU 上开始执行的代码。我们可以指定一个辅助队列作为 submit 函数的参数，如果命令组无法加入主队列，就使用这个辅助队列(在本例中是主机设备)。

通过将辅助队列传递给submit调用来启用回退队列。作者建议捕捉初始错误并处理它，如第五章所述，而不是使用提供较少控制的回退队列机制。

摘要

在本章中，我们提供了队列的概述，选择与队列相关的设备，以及如何创建自定义设备选择器。我们还概述了当满足依赖性时在设备上异步执行的代码和作为 C++ 应用程序宿主代码的一部分执行的代码。第三章描述了如何控制数据移动。

Creative Commons

三、数据管理

img/489625_1_En_3_Figa_HTML.gif

超级计算机架构师经常感叹我们需要“喂野兽”。短语“喂养野兽”指的是当我们使用大量并行时，我们创建的计算机的“野兽”,向它提供数据成为需要解决的关键挑战。

在异构机器上输入数据并行 C++ 程序需要注意确保数据在需要的时候出现在需要的地方。在大型程序中，这可能需要大量的工作。在一个预先存在的 C++ 程序中，仅仅是整理如何管理所有需要的数据移动就可能是一场噩梦。

我们将仔细解释管理数据的两种方法:统一共享内存(USM)和缓冲区。USM 是基于指针的，这是 C++ 程序员所熟悉的。缓冲区提供了更高层次的抽象。选择是好事。

我们需要控制数据的移动，这一章涵盖了实现这一点的选项。

在第二章中，我们学习了如何控制代码在哪里执行。我们的代码需要数据作为输入，并产生数据作为输出。因为我们的代码可能在多个设备上运行，而这些设备不一定共享内存，所以我们需要管理数据移动。即使数据是共享的，例如 USM，同步和一致性也是我们需要理解和管理的概念。

一个合乎逻辑的问题可能是“为什么编译器不自动为我们做所有的事情？”虽然我们可以自动处理很多事情，但是如果我们不坚持自己的程序员身份，性能通常是次优的。实际上，为了获得最佳性能，在编写异构程序时，我们需要关注代码放置(第二章)和数据移动(本章)。

本章概述了数据管理，包括控制数据使用的顺序。它是对前一章的补充，前一章向我们展示了如何控制代码在哪里运行。本章帮助我们有效地使数据出现在我们要求代码运行的地方，这不仅对于正确执行我们的应用程序很重要，而且对于最小化执行时间和功耗也很重要。

介绍

没有数据，计算什么都不是。加速计算的目的是为了更快地得出答案。这意味着数据并行计算最重要的方面之一是它们如何访问数据，在机器中引入加速器设备会使情况进一步复杂化。在传统的基于单插槽 CPU 的系统中，我们只有一个内存。加速器设备通常有自己的附加存储器，不能从主机直接访问。因此，支持分立设备的并行编程模型必须提供管理这些多个存储器并在它们之间移动数据的机制。

在这一章中，我们将概述各种数据管理机制。我们介绍了用于数据管理的统一共享内存和缓冲区抽象，并描述了内核执行和数据移动之间的关系。

数据管理问题

从历史上看，并行编程的共享内存模型的优势之一是它们提供了一个单一的共享内存视图。拥有这种单一的内存视图简化了生活。我们不需要做任何特殊的事情来从并行任务中访问内存(除了正确的同步以避免数据竞争)。虽然某些类型的加速器设备(如集成 GPU)与主机 CPU 共享内存，但许多分立加速器拥有自己的独立于 CPU 的本地内存，如图 3-1 所示。

img/489625_1_En_3_Fig1_HTML.png

图 3-1

多个离散存储器

设备本地与设备远程

当使用直接连接到设备的内存而不是远程内存来读写数据时，设备上运行的程序性能会更好。我们将对直接连接的存储器的访问称为本地访问。对另一个设备内存的访问是远程访问。远程访问往往比本地访问慢，因为它们必须通过带宽较低和/或延迟较高的数据链路传输。这意味着将计算和它将使用的数据放在一起通常是有利的。为了做到这一点，我们必须设法确保数据在不同的内存之间复制或迁移，以便将数据移动到离计算发生地更近的地方。

img/489625_1_En_3_Fig2_HTML.png

图 3-2

数据移动和内核执行

管理多个存储器

大体上说，管理多个内存可以通过两种方式来完成:通过我们的程序显式地管理内存，或者由运行时隐式地管理内存。每种方法都有其优点和缺点，我们可以根据情况或个人喜好选择其中之一。

显式数据移动

管理多个存储器的一种选择是在不同的存储器之间显式地复制数据。图 3-2 显示了一个带有独立加速器的系统，我们必须首先将内核需要的任何数据从主机内存复制到 GPU 内存。在内核计算结果之后，我们必须将这些结果复制回 CPU，然后主机程序才能使用这些数据。

显式数据移动的主要优势在于，我们可以完全控制数据在不同内存之间传输的时间。这一点很重要，因为在某些硬件上，重叠计算和数据传输对于获得最佳性能至关重要。

显式数据移动的缺点是，指定所有数据移动可能会很繁琐且容易出错。传输不正确的数据量，或者在内核开始计算之前不确保所有数据都已传输，都可能导致不正确的结果。从一开始就让所有数据正确移动可能是一项非常耗时的任务。

隐式数据移动

程序控制的显式数据移动的替代方法是由并行运行时或驱动程序控制的隐式数据移动。在这种情况下，并行运行时不需要在不同的内存之间进行显式复制，而是负责确保数据在使用之前被传输到适当的内存。

隐式数据移动的优势在于，让应用程序利用直接连接到设备的更快的内存需要更少的努力。所有繁重的工作都由运行时自动完成。这也减少了将错误引入程序的机会，因为运行时将自动识别何时必须执行数据传输以及必须传输多少数据。

隐式数据移动的缺点是我们对运行时隐式机制的行为控制很少或没有控制。运行时将提供功能正确性，但可能不会以确保计算与数据传输最大重叠的最佳方式移动数据，这可能会对程序性能产生负面影响。

选择正确的策略

为一个项目选择最佳策略取决于许多不同的因素。不同的策略可能适用于程序开发的不同阶段。我们甚至可以决定最好的解决方案是混合和匹配程序不同部分的显式和隐式方法。我们可能会选择开始使用隐式数据移动来简化将应用程序移植到新设备的过程。当我们开始调整应用程序的性能时，我们可能会在代码的性能关键部分用显式数据移动来代替隐式数据移动。未来的章节将会介绍数据传输如何与计算重叠以优化性能。

USM、缓冲区和图像

管理内存有三个抽象概念:统一共享内存(USM)、缓冲区和映像。USM 是一种基于指针的方法，应该为 C/C++ 程序员所熟悉。USM 的一个优势是更容易与现有的操作指针的 C++ 代码集成。由buffer模板类表示的缓冲区描述了一维、二维或三维数组。它们提供了可以在主机或设备上访问的内存的抽象视图。程序不直接访问缓冲区，而是通过accessor对象来使用。图像充当一种特殊类型的缓冲区，提供特定于图像处理的额外功能。这个功能包括支持特殊的图像格式，使用 sampler 对象读取图像，等等。缓冲区和映像是解决许多问题的强大抽象，但是重写现有代码中的所有接口以接受缓冲区或访问器可能非常耗时。由于缓冲区和图像的接口基本相同，本章的其余部分将只关注 USM 和缓冲区。

统一共享内存

USM 是我们可以使用的一种数据管理工具。USM 是一种基于指针的方法，使用malloc或new分配数据的 C 和 C++ 程序员应该很熟悉。当移植大量使用指针的现有 C/C++ 代码时，USM 简化了工作。支持 USM 的设备支持统一的虚拟地址空间。拥有统一的虚拟地址空间意味着主机上的 USM 分配例程返回的任何指针值都将是设备上的有效指针值。我们不必手动转换主机指针来获得“设备版本”，我们在主机和设备上都看到相同的指针值。

USM 的更详细描述可在第六章中找到。

通过指针访问内存

由于当系统包含主机内存和一定数量的设备连接本地内存时，并非所有内存都是相同的，USM 定义了三种不同类型的分配:device、host和shared。所有类型的分配都在主机上执行。图 3-3 总结了每种分配类型的特点。

img/489625_1_En_3_Fig3_HTML.png

图 3-3

USM 分配类型

在设备连接内存中进行device分配。这种分配可以在设备上读取和写入，但不能从主机直接访问。我们必须使用显式复制操作在主机内存的常规分配和device分配之间移动数据。

host分配发生在可在主机和设备上访问的主机内存中。这意味着相同的指针值在主机代码和设备内核中都有效。然而，当访问这样的指针时，数据总是来自主机内存。如果在设备上访问，数据不会从主机迁移到设备本地内存。相反，数据通常通过将设备连接到主机的总线(例如 PCI-Express (PCI-E ))发送。

在主机和设备上都可以访问shared分配。在这方面，它与主机分配非常相似，但不同之处在于数据现在可以在主机内存和设备本地内存之间迁移。这意味着在迁移发生后，对设备的访问是从速度更快的设备本地内存进行的，而不是通过更高延迟的连接远程访问主机内存。通常，这是通过运行时内部的机制和对我们隐藏的低级驱动程序来实现的。

USM 和数据移动

USM 支持显式和隐式数据移动策略，不同的分配类型映射到不同的策略。设备分配要求我们在主机和设备之间显式移动数据，而主机和共享分配提供隐式数据移动。

USM 中的显式数据移动

使用 USM 的显式数据移动是通过device分配以及队列和处理程序类中的特殊memcpy()来完成的。我们对memcpy()操作(动作)进行排队，以便将数据从主机传输到设备，或者从设备传输到主机。

图 3-4 包含一个操作设备分配的内核。在内核执行之前和之后，使用memcpy()操作在hostArray和deviceArray之间复制数据。对队列上的wait()的调用确保在内核执行之前对设备的复制已经完成，并且确保在数据复制回主机之前内核已经完成。我们将在本章的后面学习如何消除这些调用。

img/489625_1_En_3_Fig4_HTML.png

图 3-4

USM 显式数据移动

USM 中的隐式数据移动

使用 USM 的隐式数据移动是通过host和shared分配完成的。使用这些类型的分配，我们不需要显式插入拷贝操作来在主机和设备之间移动数据。相反，我们只需访问内核中的指针，任何所需的数据移动都会自动执行，无需程序员干预(只要您的设备支持这些分配)。这极大地简化了现有代码的移植:只需用适当的 USM 分配函数替换任何 malloc 或 new(以及对free释放内存的调用)，一切都将正常工作。

img/489625_1_En_3_Fig5_HTML.png

图 3-5

USM 隐式数据移动

在图 3-5 中，我们创建了两个数组hostArray和sharedArray，它们分别是主机和共享分配。虽然主机和共享分配都可以在主机代码中直接访问，但是我们在这里只初始化hostArray。类似地，可以在内核内部直接访问它，执行数据的远程读取。运行时确保sharedArray在内核访问它之前在设备上可用，并且当它稍后被主机代码读取时被移回，所有这些都不需要程序员的干预。

缓冲

为数据管理提供的另一个抽象是缓冲区对象。缓冲区是一种数据抽象，表示一个或多个给定 C++ 类型的对象。缓冲区对象的元素可以是标量数据类型(如int、float或double)、矢量数据类型(第十一章)或用户定义的类或结构。缓冲区中的数据结构必须是 C++ 普通可复制的，这意味着一个对象可以被安全地逐字节复制，而不需要调用复制构造器。

虽然缓冲区本身是单个对象，但缓冲区封装的 C++ 类型可以是包含多个对象的数组。缓冲区代表的是数据对象而不是具体的内存地址，所以不能像普通的 C++ 数组一样直接访问。实际上，出于性能原因，缓冲区对象可能会映射到几个不同设备上的多个不同内存位置，甚至是同一设备上的多个不同内存位置。相反，我们使用访问器对象来读写缓冲区。

第七章对缓冲器进行了更详细的描述。

创建缓冲区

可以通过多种方式创建缓冲区。最简单的方法是简单地用指定缓冲区大小的范围构造一个新的缓冲区。然而，以这种方式创建缓冲区并不初始化其数据，这意味着我们必须先通过其他方式初始化缓冲区，然后再尝试从中读取有用的数据。

也可以从主机上的现有数据创建缓冲区。这是通过调用几个构造器中的一个来完成的，这些构造器要么接受一个指向现有主机分配的指针，一组InputIterators，要么接受一个具有特定属性的容器。在缓冲区构造期间，数据从现有的主机分配中复制到缓冲区对象的主机内存中。如果我们使用 OpenCL 的 SYCL 互操作性特性，也可以从现有的cl_mem对象创建一个缓冲区。

访问缓冲区

主机和设备可能无法直接访问缓冲区(除非通过此处未描述的高级且不常用的机制)。相反，我们必须创建访问器来读写缓冲区。访问器为运行时提供关于我们计划如何使用缓冲区中的数据的信息，允许它正确地调度数据移动。七

img/489625_1_En_3_Fig7_HTML.png

图 3-7

缓冲区访问模式

img/489625_1_En_3_Fig6_HTML.png

图 3-6

缓冲区和存取器

访问模式

当创建访问器时，我们可以通知运行时我们将如何使用它来提供更多的优化信息。我们通过指定一个访问模式来做到这一点。访问模式在图 3-7 中描述的access::mode enum中定义。在图 3-6 所示的代码示例中，访问器myAccessor是用默认的访问模式access::mode::read_write创建的。这让运行时知道我们打算通过myAccessor读写缓冲区。访问模式是运行库优化隐式数据移动的方式。例如，access::mode::read告诉运行时，在内核开始执行之前，数据需要在设备上可用。如果内核只通过一个访问器读取数据，那么在内核完成后就没有必要将数据复制回主机，因为我们没有修改它。同样，access::mode::write让运行时知道我们将修改缓冲区的内容，并且可能需要在计算结束后将结果复制回来。

用适当的模式创建访问器给了运行时更多关于我们如何在程序中使用数据的信息。运行时使用访问器对数据的使用进行排序，但它也可以使用这些数据来优化内核和数据移动的调度。访问模式和优化标签在第七章中有更详细的描述。

数据使用的排序

内核可以被看作是提交执行的异步任务。这些任务必须提交到一个队列中，在那里它们被安排在一个设备上执行。在许多情况下，内核必须按照特定的顺序执行，这样才能计算出正确的结果。如果获得正确的结果需要任务A在任务B之前执行，我们说任务A和B之间存在依赖 ¹ 。

然而，内核并不是必须被调度的任务的唯一形式。在内核开始执行之前，内核访问的任何数据都需要在设备上可用。这些数据依赖性会以从一个设备到另一个设备的数据传输的形式产生额外的任务。数据传输任务可以是显式编码的复制操作，也可以是运行时执行的更常见的隐式数据移动。

如果我们把一个程序中的所有任务以及它们之间存在的依赖关系都拿出来，我们就可以用这个来把信息可视化为一个图形。该任务图具体是有向无环图(DAG ),其中节点是任务，边是依赖关系。该图是定向的，因为依赖关系是单向的:任务A必须发生在任务B之前。这个图是非循环的，因为它不包含任何从一个节点回到自身的循环或路径。

在图 3-8 中，任务A必须在任务B和C之前执行。同样，B和C必须在任务D之前执行。由于B和C彼此之间没有依赖关系，只要任务A已经执行，运行时就可以自由地以任何顺序(甚至并行)执行它们。因此，该图可能的法律顺序是A``B``C``D``A``C``B``D，如果B和C可以并发执行，甚至是A``{B,C}``D。

img/489625_1_En_3_Fig8_HTML.png

图 3-8

简单任务图

任务可能依赖于所有任务的子集。在这些情况下，我们只想指定关系到正确性的依赖关系。这种灵活性为运行时优化任务图的执行顺序提供了空间。在图 3-9 中，我们从图 3-8 扩展了之前的任务图，增加了任务E和F，其中E必须在F之前执行。然而，任务E和F与节点A、B、C和D没有依赖关系。这允许运行时从许多可能的合法顺序中选择来执行所有任务。

img/489625_1_En_3_Fig9_HTML.png

图 3-9

具有不相交依赖关系的任务图

有两种不同的方法来模拟任务在队列中的执行，比如内核的启动:队列可以按照提交的顺序执行任务，也可以按照我们定义的任何依赖关系按照任何顺序执行任务。我们有几种机制来定义正确排序所需的依赖关系。

有序队列

对任务进行排序的最简单的选择是将它们提交给有序的queue对象。有序队列按照任务提交的顺序执行任务，如图 3-10 所示。尽管有序队列的直观任务排序在简单性方面提供了优势，但它也提供了一个缺点，即即使独立任务之间不存在依赖性，任务的执行也会串行化。有序队列在启动应用程序时非常有用，因为它们简单、直观、确定执行顺序，并且适用于许多代码。

img/489625_1_En_3_Fig10_HTML.png

图 3-10

有序队列使用

无序(OoO)队列

由于queue对象是无序队列(除非用in-order queue 属性创建)，它们必须提供对提交给它们的任务进行排序的方法。队列通过让我们通知运行时它们之间的依赖关系来排序任务。这些依赖性可以使用命令组明确或隐含地指定。

命令组是指定任务及其依赖性的对象。命令组通常被写成 C++ lambdas，作为参数传递给队列对象的submit()方法。这个 lambda 唯一的参数是对一个handler对象的引用。handler 对象在命令组中用于指定操作、创建访问器和指定依赖关系。

事件的显式依赖性

任务之间的显式依赖看起来就像我们已经看到的例子(图 3-8 )，其中任务 A 必须在任务 b 之前执行。以这种方式表达依赖侧重于基于发生的计算的显式排序，而不是基于计算访问的数据。请注意，表达计算之间的依赖性主要与使用 USM 的代码相关，因为使用缓冲区的代码通过访问器表达大多数依赖性。在图 3-4 和 3-5 中，我们只是告诉队列等待所有之前提交的任务完成，然后再继续。相反，我们可以通过事件对象来表达任务依赖性。向队列提交命令组时，submit()方法返回一个事件对象。然后，这些事件可以以两种方式使用。

首先，我们可以通过在事件上显式调用wait()方法来通过主机进行同步。这将强制运行时等待生成事件的任务完成执行，然后宿主程序才能继续执行。显式等待事件对于调试应用程序非常有用，但是wait()会过度限制任务的异步执行，因为它会停止主机线程上的所有执行。类似地，也可以在队列对象上调用wait()，这将阻塞主机上的执行，直到所有排队的任务完成。如果我们不想跟踪排队任务返回的所有事件，这可能是一个有用的工具。

这就把我们带到了使用事件的第二种方式。handler 类包含一个名为depends_on()的方法。此方法接受单个事件或事件向量，并通知运行时正在提交的命令组要求在命令组中的操作可以执行之前完成指定的事件。图 3-11 显示了如何使用depends_on()来订购任务的示例。

img/489625_1_En_3_Fig11_HTML.png

图 3-11

使用事件和depends_on

访问器的隐式依赖

任务之间的隐式依赖关系是从数据依赖关系创建的。任务之间的数据依赖有三种形式，如图 3-12 所示。

img/489625_1_En_3_Fig12_HTML.png

图 3-12

三种形式的数据相关性

数据依赖以两种方式向运行时表达:访问器和程序顺序。运行时必须使用这两者来正确计算数据依赖关系。如图 3-13 和 3-14 所示。

img/489625_1_En_3_Fig14_HTML.png

图 3-14

原始任务图

img/489625_1_En_3_Fig13_HTML.png

图 3-13

写后读

在图 3-13 和 3-14 中，我们执行三个内核— computeB、readA和computeC—and，然后在主机上读回最终结果。内核computeB的命令组创建了两个访问器，accA和accB。这些访问器使用访问标签read_only和write_only进行优化，指定我们不使用默认的访问模式access::mode::read_write。我们将在第七章中了解更多关于访问标签的内容。内核computeB读取缓冲区A并写入缓冲区B。内核开始执行之前，必须将缓冲区A从主机复制到设备。

内核readA也为缓冲区A创建一个只读访问器。由于内核readA是在内核computeB之后提交的，这就产生了一个读后读(RAR)的场景。然而，rar 对运行时没有额外的限制，内核可以以任何顺序自由执行。事实上，运行时可能更喜欢在内核computeB之前执行内核readA，或者甚至同时执行两者。两者都要求缓冲区A被复制到设备，但是内核computeB也要求缓冲区B被复制，以防任何现有的值没有被computeB覆盖并且可能被后来的内核使用。这意味着当缓冲区B的数据传输发生时，运行时可以执行内核readA,这也表明即使内核只写入缓冲区，缓冲区的原始内容仍可能被移动到设备，因为不能保证缓冲区中的所有值都会被内核写入(参见第七章，了解在这些情况下让我们进行优化的标签)。

内核computeC读取缓冲区B，这是我们在内核computeB中计算的。因为我们在提交内核computeB之后提交了内核computeC，这意味着内核computeC对缓冲区B有原始数据依赖。原始相关性也被称为真实相关性或流相关性，因为数据需要从一个计算流到另一个计算，以便计算正确的结果。最后，我们还在内核computeC和主机之间创建了一个对缓冲区C的原始依赖，因为主机希望在内核完成后再调用read C。这迫使运行时将缓冲区C复制回主机。由于没有写入设备上的缓冲区A，运行时不需要将该缓冲区复制回主机，因为主机已经有了最新的副本。

img/489625_1_En_3_Fig16_HTML.png

图 3-16

战争和战时任务图

img/489625_1_En_3_Fig15_HTML.png

图 3-15

读后写和写后写

在图 3-15 和 3-16 中，我们再次执行三个内核:computeB、rewriteA和rewriteB。内核computeB再次读取缓冲区A并写入缓冲区B，内核rewriteA写入缓冲区A，内核rewriteB写入缓冲区B。内核rewriteA理论上可以早于内核computeB执行，因为在内核准备好之前需要传输的数据较少，但是它必须等到内核computeB完成之后，因为存在对缓冲区A的 WAR 依赖。

在这个例子中，内核computeB需要来自主机的 A 的原始值，如果内核rewriteA在内核computeB之前执行，它将读取错误的值。战争依赖也被称为反依赖。原始依赖关系确保数据以正确的方向正确流动，而 WAR 依赖关系确保现有值在被读取之前不会被覆盖。内核重写中发现的 WAW 对缓冲区B的依赖类似地起作用。如果在内核computeB和rewriteB之间提交了对缓冲区B的任何读取，它们将导致原始和 WAR 依赖，这将正确排序任务。然而，在这个例子中，内核rewriteB和主机之间有一个隐含的依赖关系，因为最终的数据必须写回主机。我们将在第七章中了解导致这种写回的更多原因。WAW 依赖性，也称为输出依赖性，确保最终输出在主机上是正确的。

选择数据管理策略

为我们的应用程序选择正确的数据管理策略在很大程度上取决于个人偏好。事实上，我们可能从一种策略开始，随着程序的成熟，我们会切换到另一种策略。然而，有一些有用的指导方针可以帮助我们选择满足我们需求的策略。

要做的第一个决定是我们想要使用显式还是隐式数据移动，因为这极大地影响了我们需要对程序做什么。隐式数据移动通常是一个更容易开始的地方，因为所有的数据移动都是为我们处理的，让我们专注于计算的表达式。

如果我们决定从一开始就完全控制所有数据移动，那么使用 USM 设备分配的显式数据移动就是我们想要开始的地方。我们只需要确保在主机和设备之间添加所有必要的副本！

在选择隐式数据移动策略时，我们仍然可以选择是使用缓冲区还是 USM 主机或共享指针。同样，这个选择是个人偏好的问题，但是有几个问题可以帮助我们选择其中一个。如果我们正在移植一个使用指针的现有 C/C++ 程序，USM 可能是一个更容易的途径，因为大多数代码不需要改变。如果数据表示没有引导我们选择，我们可以问的另一个问题是我们希望如何表达内核之间的依赖关系。如果我们更喜欢考虑内核之间的数据依赖，选择缓冲区。如果我们更愿意将依赖关系理解为在执行一个计算之前执行另一个计算，并希望使用有序队列、显式事件或内核间等待来表达，请选择 USM。

当使用 USM 指针时(通过显式或隐式数据移动)，我们可以选择使用哪种类型的队列。有序队列简单而直观，但是它们约束了运行时间，并且可能会限制性能。无序队列更复杂，但是它们给了运行时更多的重新排序和重叠执行的自由。如果我们的程序在内核之间有复杂的依赖关系，无序队列类是正确的选择。如果我们的程序只是一个接一个地运行许多内核，那么有序队列将是我们更好的选择。

处理程序类:关键成员

我们已经展示了许多使用handler类的方法。图 3-17 和 3-18 提供了这个非常重要的类别的关键成员的更详细的解释。我们还没有使用所有这些成员，但是在本书的后面会用到它们。这是摆放它们的最佳地点。

一个密切相关的类，即queue类，在第二章的结尾有类似的解释。在线 oneAPI DPC++ 语言参考对这两个类提供了更详细的解释。

img/489625_1_En_3_Fig18_HTML.png

图 3-18

处理程序类的访问器成员的简化定义

img/489625_1_En_3_Fig17_HTML.png

图 3-17

处理程序类的非访问器成员的简化定义

摘要

在这一章中，我们介绍了解决数据管理问题的机制以及如何安排数据的使用。使用加速器设备时，管理对不同内存的访问是一个关键挑战，我们有不同的选项来满足我们的需求。

我们概述了数据使用之间可能存在的不同类型的依赖关系，并描述了如何向队列提供关于这些依赖关系的信息，以便它们正确地对任务进行排序。

本章概述了统一共享内存和缓冲区。我们将在第六章中更详细地探讨 USM 的所有模式和行为。第七章将更深入地探索缓冲区，包括创建缓冲区和控制其行为的所有不同方法。第八章将再次讨论控制内核执行和数据移动顺序的队列调度机制。

Creative Commons

Footnotes 1

请注意，你可能会在其他文本中看到“dependency”和“dependencies”有时拼写为“dependency”和“dependencies”。它们的意思是一样的，但是我们倾向于在几篇关于数据流分析的重要论文中使用的拼写。

参见dl . ACM . org/doi/pdf/10 . 1145/75277 . 75280和dl . ACM . org/doi/pdf/10 . 1145/113446 . 113449。

四、表达并行

img/489625_1_En_4_Figa_HTML.gif

现在我们可以把第一批拼图拼在一起了。我们已经知道如何在设备上放置代码(第二章)和数据(第三章)——我们现在必须做的就是决定如何处理它们。为此，我们现在来补充一些我们到目前为止方便地忽略或掩饰的东西。本章标志着从简单的教学示例到真实世界并行代码的过渡，并扩展了我们在前面章节中随意展示的代码示例的细节。

用一种新的并行语言编写我们的第一个程序似乎是一项艰巨的任务，尤其是如果我们是并行编程的新手。语言规范不是为应用程序开发人员编写的，通常假设他们熟悉一些术语；它们不包含以下问题的答案:

为什么排比的表达方式不止一种？
我应该用哪种表达排比的方法？
我真的需要了解执行模型多少？

本章试图解决这些问题以及更多的问题。我们介绍了数据并行内核的概念，使用工作代码示例讨论了不同内核形式的优缺点，并强调了内核执行模型最重要的方面。

内核内部的并行性

近年来，并行内核作为一种表达数据并行性的强大手段出现了。基于内核的方法的主要设计目标是跨多种设备的可移植性和高程序员生产率。这样，内核通常不被硬编码以与特定数量或配置的硬件资源(例如，内核、硬件线程、SIMD[单指令、多数据]指令)一起工作。相反，内核根据抽象概念来描述并行性，实现(即编译器和运行时的组合)然后可以映射到特定目标设备上可用的硬件并行性。尽管这种映射是由实现定义的，但是我们可以(也应该)相信实现会选择一种合理的、能够有效利用硬件并行性的映射。

以独立于硬件的方式展示大量并行性可确保应用程序可以扩展(或缩小)以适应不同平台的功能，但是…

保证功能的可移植性并不等同于保证高性能！

支持的器件有很大的多样性，我们必须记住，不同的架构是针对不同的使用情况而设计和优化的。每当我们希望在特定设备上实现最高水平的性能时，我们应该总是期望需要一些额外的手动优化工作——不管我们使用的是什么编程语言！这种特定于设备的优化的示例包括针对特定高速缓存大小的分块、选择分摊调度开销的粒度、利用专门的指令或硬件单元，以及最重要的是，选择适当的算法。其中一些例子将在第 15 、 16 和 17 章中再次出现。

在应用程序开发过程中，在性能、可移植性和生产力之间取得恰当的平衡是我们都必须面对的挑战——也是本书无法完全解决的挑战。然而，我们希望表明，DPC++ 提供了使用一种高级编程语言来维护通用可移植代码和优化的特定于目标的代码所需的所有工具。剩下的留给读者作为练习！

多维核

许多其他语言的并行构造是一维的，将工作直接映射到相应的一维硬件资源(例如，硬件线程的数量)。并行内核是比这更高级的概念，它们的维度更能反映我们的代码通常试图解决的问题(在一维、二维或三维空间中)。

然而，我们必须记住，由并行内核提供的多维索引是在底层一维空间之上实现的，方便了程序员。理解这种映射的行为是某些优化的重要部分(例如，调整内存访问模式)。

一个重要的考虑是哪个维度是连续的或单位步长(即，多维空间中的哪些位置在一维空间中彼此相邻)。SYCL 中与并行性相关的所有多维量都使用相同的约定:维度从 0 到 N-1 编号，其中维度 N-1 对应于连续维度。在多维数量被写成列表(例如，在构造器中)或者一个类支持多个下标操作符的地方，这种编号从左到右应用。这个约定与标准 C++ 中多维数组的行为一致。

图 4-1 显示了一个使用 SYCL 约定将二维空间映射到线性索引的例子。我们当然可以打破这一惯例，采用自己的指数线性化方法，但必须谨慎行事，因为打破 SYCL 惯例可能会对受益于 stride-1 访问的器件产生负面性能影响。

img/489625_1_En_4_Fig1_HTML.png

图 4-1

映射到线性索引的二维大小范围(2，8)

如果一个应用程序需要三个以上的维度，我们必须负责使用模运算手动映射多维和线性索引。

循环与内核

迭代循环是一种内在的串行结构:循环的每次迭代都是按顺序执行的(即，按次序)。优化编译器可能能够确定循环的一些或所有迭代可以并行执行，但它必须是保守的——如果编译器不够智能或没有足够的信息来证明并行执行总是安全的，它必须保持循环的顺序语义的正确性。

img/489625_1_En_4_Fig2_HTML.png

图 4-2

将向量加法表示为串行循环

考虑图 4-2 中的循环，它描述了一个简单的向量加法。即使在这种简单的情况下，证明循环可以并行执行也不是小事:只有当c不与a或b重叠时，并行执行才是安全的，在一般情况下，没有运行时检查就无法证明这一点！为了解决这种情况，语言增加了一些功能，使我们能够为编译器提供额外的信息，这些信息可以简化分析(例如，断言指针不与restrict重叠)或完全覆盖所有分析(例如，声明循环的所有迭代都是独立的，或准确定义如何将循环调度到并行资源)。

并行循环的确切含义有些模糊——由于不同的并行编程语言对该术语的重载——但是许多常见的并行循环结构表示应用于顺序循环的编译器转换。这种编程模型使我们能够编写连续的循环，并在以后提供关于如何安全并行执行不同迭代的信息。这些模型非常强大，可以与其他最新的编译器优化很好地集成，并极大地简化了并行编程，但并不总是鼓励我们在开发的早期阶段考虑并行性。

并行内核不是一个循环，没有迭代。更确切地说，一个内核描述了一个单一的操作，它可以被实例化多次并应用于不同的输入数据；当内核并行启动时，该操作的多个实例同时执行。

img/489625_1_En_4_Fig3_HTML.png

图 4-3

将循环重写(用伪代码)为并行内核

图 4-3 显示了我们使用伪代码重写为内核的简单循环示例。这个内核中并行性的机会是清楚而明确的:内核可以由任意数量的实例并行执行，并且每个实例独立地应用于单独的数据。通过将该操作编写为内核，我们断言并行运行是安全的(理想情况下应该是安全的)。

简而言之，基于内核的编程不是一种将并行性改进到现有顺序代码中的方法，而是一种用于编写显式并行应用程序的方法。

我们越早将我们的思维从并行循环转移到内核，使用数据并行 C++ 编写有效的并行程序就越容易。

语言功能概述

一旦我们决定编写一个并行内核，我们必须决定我们想要启动什么类型的内核，以及如何在我们的程序中表示它。表达并行内核的方式有很多种，如果我们想掌握这门语言，我们需要熟悉每一种方式。

从主机代码中分离内核

我们有几种分离主机和设备代码的替代方法，我们可以在应用程序中混合和匹配它们:C++ lambda 表达式或函数对象(函子)、OpenCL C 源代码字符串或二进制文件。这些选项中的一些已经在第二章中介绍过了，所有这些选项都将在第十章中详细介绍。

所有这些选项都共享表达并行性的基本概念。为了一致和简洁，本章中的所有代码示例都使用 C++ lambdas 表示内核。

Lambdas Not Considered Harmful

为了开始使用 DPC++，不需要完全理解 C++ 规范中关于 lambda 的所有内容——我们只需要知道 lambda 的主体表示内核，并且捕获的变量(通过值)将作为参数传递给内核。

使用 lambdas 而不是更详细的机制来定义内核不会对性能产生影响。DPC++ 编译器总是能理解 lambda 何时代表并行内核的主体，并能相应地针对并行执行进行优化。

关于 C++ lambda 函数的复习，以及它们在 SYCL 中的用法，请参见第一章。关于使用 lambdas 定义内核的更多细节，请参见第十章。

不同形式的并行内核

有三种不同的内核形式，支持不同的执行模型和语法。使用任何内核形式编写可移植的内核都是可能的，并且以任何形式编写的内核都可以进行调整，以在各种设备类型上实现高性能。然而，有时我们可能希望使用特定的形式来使特定的并行算法更容易表达，或者利用否则无法访问的语言功能。

第一种形式用于基本的数据并行内核，并为编写内核提供了最温和的介绍。对于基本内核，我们牺牲了对调度等底层特性的控制，以使内核的表达尽可能简单。单个内核实例如何映射到硬件资源完全由实现来控制，因此随着基本内核复杂性的增加，推断它们的性能变得越来越困难。

第二种形式扩展了基本内核，以提供对低级性能调优特性的访问。出于历史原因，这第二种形式被称为 ND-range (N 维范围)数据并行，要记住的最重要的事情是，它使某些内核实例能够被分组在一起，允许我们对数据局部性以及各个内核实例和将用于执行它们的硬件资源之间的映射进行一些控制。

第三种形式提供了另一种语法，使用嵌套的内核结构来简化 ND 范围内核的表达式。这第三种形式被称为分层数据并行，指的是出现在用户源代码中的嵌套内核结构的层次结构。

一旦我们更详细地讨论了它们的特性，我们将在本章的最后再次讨论如何在不同的内核形式之间进行选择。

基本数据并行内核

并行内核的最基本形式适用于令人尴尬的并行操作(例如，可以完全独立地以任何顺序应用于每一段数据的操作)。通过使用这种形式，我们可以让实现完全控制工作的调度。因此，这是一个描述性编程结构的例子——我们描述操作是令人尴尬的并行操作，所有的调度决策都由实现做出。

基本的数据并行内核是以单个程序、多数据(SPMD)风格编写的——单个“程序”(内核)应用于多个数据片段。注意，这种编程模型仍然允许内核的每个实例在代码中采用不同的路径，这是数据相关分支的结果。

SPMD 编程模型的最大优势之一是，它允许同一个“程序”映射到多个并行级别和类型，而无需我们给出任何明确的指示。同一个程序的实例可以流水线化，打包在一起用 SIMD 指令执行，分布在多个线程上，或者三者兼而有之。

理解基本数据-并行内核

一个基本并行内核的执行空间称为其执行范围，内核的每个实例称为一个项。这在图 4-4 中有图解表示。

img/489625_1_En_4_Fig4_HTML.png

图 4-4

基本并行内核的执行空间，显示了 64 个项目的 2D 范围

基本数据并行内核的执行模型非常简单:它允许完全并行执行，但不保证或需要它。项目可以按任何顺序执行，包括在单个硬件线程上按顺序执行(即没有任何并行性)！假设所有项目将被并行执行(例如，通过尝试同步项目)的内核因此非常容易导致程序在一些设备上挂起。

然而，为了保证正确性，我们必须总是在假设它们可以被并行执行的情况下编写我们的内核。例如，我们有责任确保对内存的并发访问被原子内存操作适当地保护(见第十九章)，以防止竞争情况。

编写基本数据-并行内核

基本的数据并行内核使用parallel_for函数来表示。图 4-5 展示了如何使用这个函数来表达一个向量加法，这是我们对“你好，世界！”用于并行加速器编程。

img/489625_1_En_4_Fig5_HTML.png

图 4-5

用parallel_for表示向量加法核

该函数只接受两个参数:第一个参数是指定在每个维度中启动的项目数量的range,第二个参数是为该范围中的每个索引执行的内核函数。有几个不同的类可以作为内核函数的参数，应该使用哪个取决于哪个类公开了所需的功能——我们将在后面再讨论这个问题。

图 4-6 显示了该函数的一个非常类似的用法来表示矩阵加法，除了二维数据之外，它(在数学上)与向量加法相同。这反映在内核中——两个代码片段之间的唯一区别是所使用的range和id类的维度！这样写代码是可能的，因为一个 SYCL accessor可以被一个多维id索引。虽然看起来很奇怪，但这可能非常强大，使我们能够根据数据的维度编写内核模板。

img/489625_1_En_4_Fig6_HTML.png

图 4-6

用parallel_for表示矩阵加法核

在 C/C++ 中更常见的是使用多个索引和多个下标操作符来索引多维数据结构，并且这种显式索引也受到访问器的支持。当内核同时对不同维度的数据进行操作时，或者当内核的内存访问模式比直接使用项目的id更复杂时，以这种方式使用多个索引可以提高可读性。

例如，图 4-7 中的矩阵乘法内核必须提取索引的两个独立分量，以便能够描述两个矩阵的行和列之间的点积。在我们看来，一致地使用多个下标操作符(如[j][k])比混合多种索引模式和构造二维id对象(如id(j,k))更具可读性，但这只是个人喜好问题。

本章剩余部分的例子都使用了多个下标操作符，以确保被访问的缓冲区的维数没有歧义。

img/489625_1_En_4_Fig8_HTML.png

图 4-8

将矩阵乘法工作映射到执行范围内的项目

img/489625_1_En_4_Fig7_HTML.png

图 4-7

用parallel_for表示方阵的简单矩阵乘法核

图 4-8 中的图表显示了矩阵乘法内核中的工作是如何映射到单个项目的。注意，项目的数量来自于输出范围的大小，并且相同的输入值可以由多个项目读取:每个项目通过顺序迭代 A 矩阵的(连续)行和 B 矩阵的(非连续)列来计算 C 矩阵的单个值。

基本数据的细节-并行内核

基本数据并行内核的功能通过三个 C++ 类公开:range、id和item。我们已经在前面的章节中见过几次range和id类，但是我们在这里以不同的焦点重新审视它们。

`range`类

range代表一维、二维或三维范围。range的维度是一个模板参数，因此必须在编译时知道，但是它在每个维度上的大小是动态的，在运行时传递给构造器。range类的实例用于描述并行结构的执行范围和缓冲区的大小。

图 4-9 显示了range类的简化定义，显示了构造器和查询其范围的各种方法。

img/489625_1_En_4_Fig9_HTML.png

图 4-9

range类的简化定义

`id`类

id表示一维、二维或三维范围的索引。id的定义在许多方面与range相似:它的维数在编译时也必须是已知的，并且它可以用于索引并行结构中内核的单个实例或缓冲区中的偏移量。

如图 4-10 中id类的简化定义所示，id在概念上只不过是一个、两个或三个整数的容器。我们可用的操作也非常简单:我们可以在每个维度中查询索引的组成部分，并且我们可以执行简单的运算来计算新的索引。

虽然我们可以构造一个id来表示任意的索引，但是为了获得与特定内核实例相关联的id，我们必须接受它(或者包含它的item)作为内核函数的参数。这个id(或者由它的成员函数返回的值)必须被转发到我们想要在其中查询索引的任何函数——目前没有任何免费的函数可以在程序中的任意点查询索引，但是这可能会在 DPC++ 的未来版本中解决。

接受id的内核的每个实例只知道它被分配计算的范围中的索引，而对范围本身一无所知。如果我们希望我们的内核实例知道它们自己的索引和范围，我们需要使用item类来代替。

img/489625_1_En_4_Fig10_HTML.png

图 4-10

id类的简化定义

`item`类

一个item代表一个内核函数的单个实例，封装了内核的执行范围和该范围内实例的索引(分别使用一个range和一个id)。像range和id一样，它的维数必须在编译时已知。

图 4-11 给出了item等级的简化定义。item和id的主要区别在于item公开了额外的函数来查询执行范围的属性(例如，大小、偏移量)以及计算线性化索引的便利函数。与id一样，获得与特定内核实例相关联的item的唯一方式是接受它作为内核函数的参数。

img/489625_1_En_4_Fig11_HTML.png

图 4-11

item类的简化定义

显式 ND 范围核

第二种形式的并行内核用一个项目属于组的执行范围代替了基本数据并行内核的平面执行范围，并且适用于我们希望在内核中表达一些局部性概念的情况。为不同类型的组定义和保证不同的行为，使我们能够更深入地了解和/或控制如何将工作映射到特定的硬件平台。

因此，这些显式的 ND-range 内核是一个更加规定的并行构造的例子——我们规定将工作映射到每种类型的组，并且实现必须服从该映射。然而，它并不是完全规定的，因为组本身可以以任何顺序执行，并且实现在如何将每种类型的组映射到硬件资源上保留了一些自由。这种说明性和描述性编程的结合使我们能够针对局部性设计和调优我们的内核，而不影响它们的可移植性。

像基本的数据并行内核一样，ND-range 内核以 SPMD 风格编写，其中所有工作项执行应用于多条数据的相同内核“程序”。关键的区别在于，每个程序实例可以查询它在包含它的组中的位置，并且可以访问特定于每种类型的组的附加功能。

理解显式 ND 范围并行核

ND-range 内核的执行范围被分为工作组、子组和工作项目。ND-range 表示总的执行范围，该范围被划分成统一大小的工作组(即，工作组大小必须在每个维度上精确地划分 ND-range 大小)。每个工作组可以通过实现进一步划分为子组。理解为工作项和每种类型的组定义的执行模型是编写正确的可移植程序的重要部分。

图 4-12 显示了一个 ND 尺寸范围(8，8，8)的示例，该范围分为 8 个尺寸工作组(4，4，4)。每个工作组包含 4 个工作项的 16 个一维子组。请仔细注意维度的编号:子组始终是一维的，因此 nd 范围和工作组的维度 2 成为子组的维度 0。

img/489625_1_En_4_Fig12_HTML.png

图 4-12

分为工作组、子组和工作项目的三维 ND-range

从每种类型的组到硬件资源的精确映射是实现定义的，正是这种灵活性使得程序能够在各种各样的硬件上执行。例如，工作项目可以完全顺序执行，由硬件线程和/或 SIMD 指令并行执行，或者甚至由为特定内核专门配置的硬件流水线执行。

在这一章中，我们只关注 ND-range 执行模型在通用目标平台方面的语义保证，我们不会涉及它到任何一个平台的映射。分别参见第 15 、 16 和 17 章了解 GPU、CPU 和 FPGAs 的硬件映射和性能建议的详细信息。

工作项目

工作项代表一个内核函数的单个实例。在没有其他分组的情况下，工作项目可以以任何顺序执行，并且不能相互通信或同步，除非通过对全局内存的原子内存操作(见第十九章)。

工作组

ND 范围中的工作项被组织成工作组。工作组可以以任何顺序执行，不同工作组中的工作项目不能相互通信，除非通过对全局内存的原子内存操作(见第十九章)。然而，当使用某些结构时，一个工作组中的工作项具有并发调度保证，并且这种局部性提供了一些额外的能力:

一个工作组中的工作项目可以访问工作组本地存储器，它可以映射到一些设备上的专用快速存储器(参见第九章)。
一个工作组中的工作项可以使用工作组屏障来同步，并使用工作组内存屏障来保证内存一致性(参见第九章)。
工作组中的工作项目可以访问组功能，提供通用通信例程(参见第九章)和通用并行模式(如缩减和扫描)的实现(参见第十四章)。

通常在运行时为每个内核配置工作组中的工作项目数量，因为最佳分组将取决于可用的并行性数量(即 nd 范围的大小)和目标设备的属性。我们可以使用device类的查询函数来确定特定设备支持的每个工作组的最大工作项目数(参见第十二章),我们有责任确保每个内核请求的工作组大小是有效的。

工作组执行模型中有一些微妙之处值得强调。

首先，尽管工作组中的工作项目被调度到单个计算单元，但是在工作组的数量和计算单元的数量之间不需要任何关系。事实上，ND-range 中的工作组数量可能比给定设备可以并发执行的工作组数量大很多倍！我们可能会尝试通过依赖非常聪明的特定于设备的调度来编写跨工作组同步的内核，但我们强烈建议不要这样做——这样的内核今天可能看起来可以工作，但不能保证它们可以在未来的实现中工作，并且在移动到不同的设备时很可能会崩溃。

第二，尽管一个工作组中的工作项是并发调度的，但不能保证它们独立地向前进展——在一个工作组内，在障碍和集合之间顺序地执行工作项是一种有效的实现。只有在使用提供的屏障和集合函数执行时，才能保证同一工作组中工作项之间的通信和同步是安全的，并且手工编码的同步例程可能会死锁。

Thinking in Work-Groups

工作组在许多方面类似于其他编程模型中的任务概念(例如，线程构建块):任务可以以任何顺序执行(由调度器控制)；让一台机器超额预定任务是可能的(甚至是可取的);试图在一组任务之间实现屏障通常不是一个好主意(因为它可能非常昂贵或者与调度器不兼容)。如果我们已经熟悉了基于任务的编程模型，我们可能会发现将工作组想象成数据并行任务是很有用的。

子群体

在许多现代硬件平台上，工作组中被称为子组的工作项目子集在额外的调度保证下执行。例如，作为编译器向量化的结果，子组中的工作项目可以同时执行，和/或子组本身可以在向前进度保证下执行，因为它们被映射到独立的硬件线程。

当使用单一平台时，很容易将关于这些执行模型的假设融入到我们的代码中，但这使得它们本质上不安全且不可移植——当在不同编译器之间移动时，甚至当在同一供应商的不同代硬件之间移动时，它们可能会中断！

将子组定义为语言的核心部分，为我们提供了一个安全的替代方案，来做出后来可能被证明是特定于设备的假设。利用子组功能还允许我们在较低的级别(即，接近硬件)推理工作项目的执行，并且是在许多平台上实现非常高的性能水平的关键。

与工作组一样，子组中的工作项可以同步，保证内存一致性，或者通过组函数执行常见的并行模式。但是，对于子组，没有等效的工作组本地内存(即，没有子组本地内存)。相反，子组中的工作项可以使用混洗操作(第九章)直接交换数据，而不需要显式的内存操作。

子组的某些方面是实现定义的，不在我们的控制之内。然而，对于设备、内核和 ND-range 的给定组合，一个子组有一个固定的(一维)大小，我们可以使用kernel类的查询函数来查询这个大小(参见第十章)。默认情况下，每个子组的工作项数量也是由实现选择的——我们可以通过在编译时请求特定的子组大小来覆盖这种行为，但是必须确保我们请求的子组大小与设备兼容。

像工作组一样，子组中的工作项只能保证并发执行——实现可以自由地顺序执行子组中的每个工作项，并且只在遇到子组集合函数时在工作项之间切换。子组的特殊之处在于，一些设备保证它们独立地向前进展——在一些设备上，一个工作组内的所有子组都保证最终执行(取得进展)，这是几个生产者-消费者模式的基石。可以使用设备查询来确定这种独立的前向进度保证是否成立。

Thinking in Sub-Groups

如果我们来自一个要求我们考虑显式矢量化的编程模型，那么将每个子组视为一组打包到 SIMD 寄存器中的工作项可能是有用的，其中子组中的每个工作项对应于一个 SIMD 通道。当多个子组同时运行，并且设备保证它们将向前推进时，这种心理模型扩展到将每个子组视为并行执行的独立矢量指令流。

img/489625_1_En_4_Fig13_HTML.png

图 4-13

用 ND-range 表示一个简单的矩阵乘法核parallel_for

编写显式 ND 范围数据并行内核

图 4-13 重新实现了我们之前看到的使用 ND-range 并行内核语法的矩阵乘法内核，图 4-14 中的图表显示了该内核中的工作如何映射到每个工作组中的工作项目。以这种方式对我们的工作项进行分组确保了访问的局部性，并且有望提高缓存命中率:例如，图 4-14 中的工作组具有(4，4)的局部范围，并且包含 16 个工作项，但是访问的数据是单个工作项的四倍——换句话说，我们从内存中加载的每个值都可以重用四次。

img/489625_1_En_4_Fig14_HTML.png

图 4-14

将矩阵乘法映射到工作组和工作项

到目前为止，我们的矩阵乘法示例依赖于硬件缓存来优化来自同一工作组中的工作项对 A 和 B 矩阵的重复访问。这种硬件高速缓存在传统 CPU 架构上很常见，并且在 GPU 架构上变得越来越常见，但是还有其他具有显式管理的“便笺式”存储器的架构(例如，上一代 GPU、FPGAs)。ND-range 内核可以使用本地访问器来描述应该放在工作组本地内存中的分配，然后实现可以自由地将这些分配映射到特殊内存(如果存在的话)。该工作组本地存储器的使用将在第九章中介绍。

显式 ND 范围数据并行内核的详细信息

与基本数据并行内核相比，ND-range 数据并行内核使用不同的类:range由nd_range ,代替，item由nd_item代替。还有两个新的类，代表一个工作项可能属于的不同类型的组:绑定到工作组的功能封装在group类中，绑定到子组的功能封装在sub_group类中。

`nd_range`类

一个nd_range使用两个range类的实例表示一个分组的执行范围:一个表示全局执行范围，另一个表示每个工作组的局部执行范围。图 4-15 给出了nd_range等级的简化定义。

可能有点奇怪的是，nd_range类根本没有提到子组:子组范围在构造时没有指定，无法查询。这一遗漏有两个原因。首先，子组是底层的实现细节，对于许多内核来说可以忽略。其次，有几个设备正好支持一个有效的子组大小，在任何地方指定这个大小都是不必要的冗长。所有与子组相关的功能都封装在一个专门的类中，稍后将讨论这个类。

img/489625_1_En_4_Fig15_HTML.png

图 4-15

nd_range类的简化定义

`nd_item`类

一个nd_item是一个item的 ND-range 形式，同样封装了内核的执行范围和该范围内的项目索引。nd_item与item的不同之处在于其在范围中的位置是如何查询和表示的，如图 4-16 中简化的类定义所示。例如，我们可以使用get_global_id()函数查询(全局)ND 范围中的项目索引，或者使用get_local_id()函数查询(本地)父工作组中的项目索引。

nd_item类还提供了获取描述项目所属的组和子组的类的句柄的函数。这些类为查询 ND 范围内的项目索引提供了另一种接口。我们强烈建议使用这些类来编写内核，而不是直接依赖于nd_item:使用group和sub_group类通常更干净，更清楚地传达意图，并且更符合 DPC++ 的未来方向。

img/489625_1_En_4_Fig16_HTML.png

图 4-16

nd_item类的简化定义

`group`类

group类封装了所有与工作组相关的功能，简化的定义如图 4-17 所示。

img/489625_1_En_4_Fig17_HTML.png

图 4-17

group类的简化定义

group类提供的许多函数在nd_item类中都有等价的函数:例如，调用group.get_id()相当于调用item.get_group_id(),，调用group.get_local_range()相当于调用item.get_local_range().如果我们没有使用该类公开的任何工作组函数，我们还应该使用它吗？直接使用nd_item中的函数，而不是创建一个中间的group对象，不是更简单吗？这里有一个折衷:使用group需要我们编写稍微多一点的代码，但是这些代码可能更容易阅读。例如，考虑图 4-18 中的代码片段:很明显body期望被group中的所有工作项调用，很明显parallel_for体中的get_local_range()返回的range是group的范围。同样的代码可以很容易地只用nd_item来编写，但是读者可能很难理解。

img/489625_1_En_4_Fig18_HTML.png

图 4-18

使用group类提高可读性

`sub_group`类

sub_group类封装了与子组相关的所有功能，简化的定义如图 4-19 所示。与工作组不同，sub_group类是访问子组功能的唯一方式；它的功能在nd_item中没有任何重复。sub_group类中的查询都是相对于调用工作项来解释的:例如，get_local_id()返回子组中调用工作项的本地索引。

img/489625_1_En_4_Fig19_HTML.png

图 4-19

sub_group类的简化定义

注意，有单独的函数用于查询当前子组中的工作条目的数量以及工作组内任何子组中的工作条目的最大数量。这些是否不同以及如何不同取决于子组对于特定设备是如何实现的，但是目的是反映编译器所针对的子组大小和运行时子组大小之间的任何差异。例如，非常小的工作组可能包含比编译时子组大小更少的工作项，或者不同大小的子组可用于处理不能被子组大小整除的工作组。

分层并行核

分层数据并行内核提供了一种实验性的替代语法，用于根据工作组和工作项来表达内核，其中使用嵌套调用parallel_for函数来编程分层的每一层。这种自顶向下的编程风格旨在类似于编写并行循环，可能比其他两种内核形式使用的自底向上的编程风格更熟悉。

分层内核的一个复杂性是对parallel_for的每次嵌套调用都会创建一个单独的 SPMD 环境；每个作用域定义了一个新的“程序”，该程序应该由与该作用域相关联的所有并行工作器执行。这种复杂性要求编译器执行额外的分析，并且会使某些设备的代码生成变得复杂；某些平台上的分层并行内核的编译器技术仍然相对不成熟，性能将与特定编译器实现的质量紧密相关。

由于分层数据并行内核与为特定设备生成的代码之间的关系依赖于编译器，因此分层内核应被视为比显式 ND-range 内核更具描述性的构造。然而，由于分级内核保留了控制工作到工作项和工作组的映射的能力，它们比基本内核更具规定性。

理解分层数据-并行内核

分层数据并行内核的底层执行模型与显式 ND 范围数据并行内核的执行模型相同。工作项、子组和工作组具有相同的语义和执行保证。

然而，编译器将分层内核的不同范围映射到不同的执行资源:外部范围为每个工作组执行一次(如同由单个工作项目执行)，而内部范围由工作组内的工作项目并行执行。不同的作用域还控制着不同变量在内存中的分配位置，作用域的打开和关闭意味着工作组的障碍(以加强内存的一致性)。

尽管一个工作组中的工作项仍然被分成子组，但是目前不能从一个分层的并行内核中访问sub_group类；将子组的概念结合到 SYCL 层次并行中需要比引入一个新类更大的改变，这方面的工作正在进行中。

编写分层数据并行内核

在分层内核中，parallel_for函数被parallel_for_work_group和parallel_for_work_item函数所取代，它们分别对应于工作组和工作项并行性。在parallel_for_work_group范围内的任何代码对于每个工作组只执行一次，在parallel_for_work_group范围内分配的变量对于所有工作项都是可见的(也就是说，它们被分配在工作组本地内存中)。parallel_for_work_item范围内的任何代码都由工作组的工作项并行执行，分配在parallel_for_work_item范围内的变量对单个工作项可见(即，它们被分配在工作项私有内存中)。

如图 4-20 所示，使用层次并行表示的内核与 ND-range 内核非常相似。因此，我们应该将层次并行主要视为一种生产力特征；它不会暴露任何尚未通过 ND-range 内核暴露的功能，但它可能会提高我们代码的可读性和/或减少我们必须编写的代码量。

img/489625_1_En_4_Fig20_HTML.png

图 4-20

用层次并行表达一个简单的矩阵乘法核

值得注意的是，传递给parallel_for_work_group函数的范围指定了组的数量和可选的组大小，没有指定工作项目的总数和组大小，就像 ND-range parallel_for的情况一样。内核函数接受一个group类的实例，反映出外部作用域与工作组相关联，而不是与单个工作项相关联。

parallel_for_work_item是group类的成员函数，只能在parallel_for_work_group范围内调用。在其最简单的形式中，它唯一的参数是一个接受h_item类实例的函数，该函数执行的次数等于每个工作组请求的工作项的数量；该功能按物理工作项目执行一次。parallel_for_work_item的一个额外的生产力特性是它能够支持一个逻辑范围，这个范围作为一个额外的参数传递给函数。当指定了逻辑范围时，每个物理工作项目执行零个或多个函数实例，并且逻辑范围的逻辑项目被分配给物理工作项目的循环。

图 4-21 显示了由 11 个逻辑工作项组成的逻辑范围和由 8 个物理工作项组成的底层物理范围之间的映射示例。前三个工作项被分配了两个函数实例，所有其他工作项只被分配了一个。

img/489625_1_En_4_Fig21_HTML.png

图 4-21

将大小为 11 的逻辑范围映射到大小为 8 的物理范围

如图 4-22 所示，将可选的组大小parallel_for_work_group与逻辑范围parallel_for_work_item结合起来，实现可以自由选择工作组大小，而不会牺牲我们使用嵌套并行结构方便地描述执行范围的能力。请注意，每组完成的工作量与图 4-20 中的相同，但是工作量已经从实际的工作组规模中分离出来。

img/489625_1_En_4_Fig22_HTML.png

图 4-22

用分层并行性和逻辑范围表达一个简单的矩阵乘法核心

分层数据并行内核的详细信息

分层数据并行内核重用了 ND-range 数据并行内核中的group类，但是用h_item替换了nd_item。引入了一个新的private_memory类来对parallel_for_work_group范围内的分配提供更严格的控制。

`h_item`类

一个h_item是一个item的变体，它只在一个parallel_for_work_item范围内可用。如图 4-23 所示，它提供了一个与nd_item类似的接口，有一个显著的区别:该项的索引可以相对于一个工作组的物理执行范围(用get_physical_local_id())或者一个parallel_for_work_item构造的逻辑执行范围(用get_logical_local_id())来查询。

img/489625_1_En_4_Fig23_HTML.png

图 4-23

h_item类的简化定义

`private_memory`类

private_memory类提供了一种机制来声明每个工作项私有的变量，但是这些变量可以通过嵌套在同一个parallel_for_work_group范围内的多个parallel_for_work_item构造来访问。

这个类是必要的，因为在不同的层次并行作用域中声明的变量的行为方式:如果编译器可以证明这样做是安全的，在外部作用域中声明的变量才是私有的，而在内部作用域中声明的变量是逻辑工作项而不是物理工作项的私有变量。对于我们来说，仅仅使用作用域来表达一个变量对于每个物理工作项来说是私有的是不可能的。

为了了解为什么这是一个问题，让我们回到图 4-22 中的矩阵乘法内核。ib和jb变量是在parallel_for_work_group范围内声明的，默认情况下应该分配在工作组本地内存中！一个优化的编译器很有可能不会犯这个错误，因为变量是只读的，它们的值足够简单，可以在每个工作项上进行冗余计算，但是语言没有这样的保证。如果我们想确定变量是在工作项私有内存中声明的，我们必须将变量声明包装在private_memory类的实例中，如图 4-24 所示。

img/489625_1_En_4_Fig24_HTML.png

图 4-24

private_memory类的简化定义

例如，如果我们要使用private_memory类重写矩阵乘法内核，我们会将变量定义为private_memory<int> ib(grp)，并且对这些变量的每次访问都会变成ib[item]。在这种情况下，使用private_memory类会导致代码更难阅读，而在parallel_for_work_item范围内声明值会更清晰。

我们的建议是，如果一个工作项私有变量在同一个parallel_for_work_group内的多个parallel_for_work_item范围内使用，重复计算的代价太大，或者它的计算有副作用，阻止它被冗余地计算，那么只使用private_memory类。否则，我们应该默认依赖现代优化编译器的能力，只有在分析失败时才在parallel_for_work_item范围内声明变量(记住还要向编译器供应商报告这个问题)。

将计算映射到工作项

到目前为止，大多数代码示例都假设一个内核函数的每个实例对应于对一段数据的单个操作。这是一种编写内核的简单方法，但是这种一对一的映射并不是由 DPC++ 或任何内核形式决定的——我们总是能够完全控制数据(和计算)到单个工作项的分配，并且使这种分配参数化是提高性能可移植性的好方法。

一对一映射

当我们编写内核时，工作与工作项之间存在一对一的映射，这些内核必须总是以大小与需要完成的工作量完全匹配的range或nd_range启动。这是编写内核最显而易见的方式，在许多情况下，它工作得非常好——我们可以相信一个实现可以有效地将工作项映射到硬件。

但是，在针对系统和实现的特定组合进行性能调优时，可能有必要更加关注底层调度行为。计算资源的工作组调度是由实现定义的，并且可能是动态的(即，当计算资源完成一个工作组时，它执行的下一个工作组可能来自共享队列)。动态调度对性能的影响不是固定的，并且其重要性取决于包括内核功能的每个实例的执行时间以及调度是在软件(例如，在 CPU 上)还是硬件(例如，在 GPU 上)中实现的因素。

多对一映射

另一种方法是编写工作到工作项的多对一映射的内核。在这种情况下，范围的含义发生了微妙的变化:范围不再描述要完成的工作量，而是要使用的工人数量。通过改变工人的数量和分配给每个工人的工作量，我们可以微调工作分配以最大化性能。

编写这种形式的内核需要做两处修改:

内核必须接受一个描述工作总量的参数。
内核必须包含一个将工作分配给工作项的循环。

图 4-25 给出了这种内核的一个简单例子。注意内核内部的循环有一个稍微不寻常的形式——起始索引是工作项在全局范围内的索引，步距是工作项的总数。数据到工作项的这种循环调度确保循环的所有N迭代将由一个工作项执行，而且线性工作项访问连续的内存位置(以改善缓存局部性和矢量化行为)。工作可以类似地跨组分布，或者将工作项分布在单个组中，以进一步提高局部性。

img/489625_1_En_4_Fig25_HTML.png

图 4-25

具有独立数据和执行范围的内核

这些工作分配模式很常见，当使用具有逻辑范围的分层并行时，可以非常简洁地表达它们。我们期望 DPC++ 的未来版本将引入语法糖来简化 ND-range 内核中工作分配的表达。

选择内核形式

在不同的内核形式之间进行选择很大程度上是个人偏好的问题，并且受到其他并行编程模型和语言的经验的严重影响。

选择特定内核形式的另一个主要原因是，它是公开内核所需的某些功能的唯一形式。不幸的是，在开发开始之前很难确定哪些功能是必需的——尤其是当我们还不熟悉不同的内核形式以及它们与各种类的交互时。

为了帮助我们驾驭这个复杂的空间，我们根据自己的经验编写了两本指南。这些指南应该被认为是经验法则，绝对不是要取代我们自己的实验——在不同的内核形式之间进行选择的最佳方式总是花一些时间来编写每一种形式，以便了解哪种形式最适合我们的应用程序和开发风格。

第一个指南是图 4-26 中的流程图，它选择一个内核表单基于

img/489625_1_En_4_Fig26_HTML.png

图 4-26

帮助我们为内核选择正确的形式

我们是否有并行编程的经验
无论我们是从头开始编写新代码，还是移植用不同语言编写的现有并行程序
我们的内核是令人尴尬的并行，已经包含嵌套并行，还是在内核函数的不同实例之间重用数据
无论我们是在 SYCL 中编写一个新的内核来最大化性能还是提高代码的可移植性，还是因为它提供了一种比低级语言更有效的表达并行性的方式

第二个指南是图 4-27 中的表格，它总结了每种内核形式的功能。值得注意的是，该表反映了本书出版时 DPC++ 的状态，并且每个内核形式可用的特性应该会随着语言的发展而变化。然而，我们预计基本趋势将保持不变:基本数据并行内核将不会公开位置感知功能，显式 ND-range 内核将公开所有性能支持功能，而分层内核在公开功能方面将落后于显式 ND-range 内核，但它们对这些功能的表达将使用更高级别的抽象。

img/489625_1_En_4_Fig27_HTML.png

图 4-27

每种内核形式可用的特性

摘要

本章介绍了在 DPC++ 中表达并行性的基础，并讨论了编写数据并行内核的每种方法的优缺点。

DPC++ 和 SYCL 支持多种形式的并行性，我们希望我们已经提供了足够的信息，让读者可以开始编写代码了！

我们只是触及了表面，对本章中介绍的许多概念和类的更深入的探究即将到来:本地内存、屏障和通信例程的使用将在第九章中讨论；除了使用 lambda 表达式，定义内核的不同方法将在第十章中讨论；第 15 、 16 和 17 章将探讨 ND-range 执行模型到具体硬件的详细映射；使用 DPC++ 表达通用并行模式的最佳实践将在第十四章中介绍。

Creative Commons

五、错误处理

img/489625_1_En_5_Figa_HTML.gif

阿加莎·克里斯蒂在 1969 年写道:“只要计算机努力，人为错误就不算什么。”作为程序员，我们要收拾残局，这并不奇怪。错误处理机制可以捕捉其他人可能犯的程序员错误。由于我们不打算自己犯错，我们可以专注于使用错误处理来处理现实世界中由于和其他原因而可能发生的情况。

检测和处理意外情况和错误在应用程序开发过程中可能是有帮助的(想想:在项目中工作的另一个程序员确实犯了错误)，但更重要的是在稳定和安全的生产应用程序和库方面起着关键作用。我们用这一章来描述 SYCL 中可用的错误处理机制，这样我们就可以了解我们的选择，以及如果我们关心错误的检测和管理，如何构建应用程序。

本章概述了 SYCL 中的同步和异步错误，描述了如果我们在代码中不做任何事情来处理错误时应用程序的行为，并深入探讨了 SYCL 特有的允许我们处理异步错误的机制。

安全第一

C++ 错误处理的一个核心方面是，如果我们对已经检测到(抛出)的错误不做任何处理，那么应用程序将会终止并指示出错。这种行为允许我们在编写应用程序时不必关注错误管理，并且仍然相信错误会以某种方式通知开发人员或用户。当然，我们并不是建议我们应该忽略错误处理！生产应用程序应该将错误管理作为架构的核心部分来编写，但是应用程序通常在开始开发时没有这样的关注点。C++ 的目标是让不处理错误的代码仍然能够观察到错误，即使它们没有被显式处理。

由于 SYCL 是数据并行 C++，同样的原理也成立:如果我们在代码中不做任何事情来管理错误，并且检测到错误，程序将发生异常终止，让我们知道发生了不好的事情。生产应用程序当然应该将错误管理视为软件架构的核心部分，不仅仅是报告错误，还经常从错误状态中恢复。

如果我们不添加任何错误管理代码，当错误发生时，我们仍然会看到一个异常的程序终止，这是一个深入挖掘的指示。

错误类型

C++ 通过其异常机制提供了一个通知和处理错误的框架。除此之外，异构编程还需要额外级别的错误管理，因为有些错误发生在设备上，或者在尝试在设备上启动工作时。这些错误通常在时间上与宿主程序的执行分离，因此它们不能与经典的 C++ 异常处理机制完全集成。为了解决这个问题，有额外的机制使异步错误像常规 C++ 异常一样易于管理和控制。

图 5-1 显示了一个典型应用的两个组成部分:(1)顺序运行的主机代码，并将工作提交给任务图以备将来执行;( 2)任务图，它与主机程序异步运行，并在必要的依赖关系满足时在设备上执行内核或其他动作。该示例显示了作为任务图的一部分异步执行的操作parallel_for，但是其他操作也是可能的，如第 3 、 4 和 8 章中所讨论的。

img/489625_1_En_5_Fig1_HTML.png

图 5-1

主机程序和任务图执行的分离

图 5-1 的左侧和右侧(主机和任务图)的区别是理解同步和异步错误之间差异的关键。

同步当主机程序执行某项操作(如 API 调用或对象构造器)时检测到错误条件时，就会发生错误。它们可以在图左侧的指令完成之前被检测到，并且错误可以由导致错误的操作立即抛出。我们可以在图的左侧用一个try-catch结构包装特定的指令，期望在 try 块结束之前检测到由于try内的操作而产生的错误(并因此被捕获)。C++ 异常机制就是为处理这些类型的错误而设计的。

或者，异步错误出现在图 5-1 右侧的部分，只有当执行任务图中的操作时才会检测到错误。当异步错误作为任务图执行的一部分被检测到时，主机程序通常已经继续执行了，所以没有代码可以用try-catch构造来捕获这些错误。取而代之的是一个异步异常处理框架来处理这些相对于主机程序执行看似随机发生的错误。

让我们制造一些错误！

作为本章剩余部分的例子，并允许我们进行实验，我们将在下面的部分创建同步和异步错误。

img/489625_1_En_5_Fig2_HTML.png

图 5-2

创建同步错误

同步误差

在图 5-2 中，从一个缓冲区创建了一个子缓冲区，但其大小非法(大于原始缓冲区)。子缓冲区的构造器检测到这个错误，并在构造器执行完成之前抛出异常。这是一个同步错误，因为它作为宿主程序执行的一部分(与之同步)发生。在构造器返回之前，错误是可以检测到的*，因此可以在宿主程序中的错误起源点或检测点立即处理错误。*

我们的代码示例不做任何事情来捕获和处理 C++ 异常，所以默认的 C++ 未捕获异常处理程序为我们调用std::terminate，发出出错的信号。

异步误差

生成异步错误有点棘手，因为实现会尽可能同步地检测和报告错误。同步错误更容易调试，因为它们发生在宿主程序中特定的起始点，所以只要有可能就应该优先考虑。不过，出于演示目的，生成异步错误的一种方法是在命令组提交中添加一个后备/辅助队列，并丢弃碰巧抛出的同步异常。图 5-3 显示了这样的代码，它调用我们的handle_async_error函数来允许我们进行实验。没有辅助/后备队列也可能发生和报告异步错误，因此请注意，辅助队列只是示例的一部分，绝不是异步错误的必要条件。

img/489625_1_En_5_Fig3_HTML.png

图 5-3

创建异步错误

应用程序错误处理策略

C++ 异常特性被设计成将程序中检测到错误的地方和可能处理错误的地方完全分开，这个概念非常适合 SYCL 中的同步和异步错误。通过throw和catch机制，可以定义处理程序的层次结构，这在生产应用程序中很重要。

构建一个能够以一致和可靠的方式处理错误的应用程序需要预先制定一个策略，并为错误管理构建一个软件架构。C++ 提供了灵活的工具来实现许多可供选择的策略，但是这种架构超出了本章的范围。有许多书籍和其他参考资料专门讨论这个主题，所以我们鼓励大家去查阅它们，以全面了解 C++ 错误管理策略。

也就是说，错误检测和报告并不总是需要生产规模的。如果目标只是在执行过程中检测错误并报告错误(但不一定是从错误中恢复)，那么可以通过最少的代码可靠地检测和报告程序中的错误。接下来的部分首先介绍了如果我们忽略错误处理并且什么都不做会发生什么(默认行为并没有那么糟糕！)，后面是推荐的错误报告，它在基本应用程序中很容易实现。

忽略错误处理

C++ 和 SYCL 旨在告诉我们，即使我们没有显式地处理错误，也会出现问题。未处理的同步或异步错误的默认结果是程序异常终止，操作系统应该告诉我们这一点。下面的两个例子分别模拟了如果我们不处理同步和异步错误时将会发生的行为。

图 5-4 显示了一个未处理的 C++ 异常的结果，例如，这可能是一个未处理的 SYCL 同步错误。我们可以使用这段代码来测试在这种情况下特定的操作系统会报告什么。

img/489625_1_En_5_Fig4_HTML.png

图 5-4

C++ 中未处理的异常

图 5-5 显示了被调用的std: :terminate的示例输出，这将是我们的应用程序中未处理的 SYCL 异步错误的结果。我们可以使用这段代码来测试在这种情况下特定的操作系统会报告什么。

img/489625_1_En_5_Fig5_HTML.png

图 5-5

std: :terminate在 SYCL 异步异常未处理时调用

虽然我们可能应该处理程序中的错误，但是由于未被捕获的错误将被捕获，程序将被终止，所以我们不需要担心程序会无声无息地失败！

同步错误处理

我们保持这一节非常短，因为 SYCL 同步错误只是 C++ 异常。SYCL 中添加的大多数额外错误机制都与异步错误有关，我们将在下一节中讨论，但是同步错误很重要，因为实现会尝试同步检测和报告尽可能多的错误，因为它们更容易推理和处理。

SYCL 定义的同步错误是从sycl::exception类型的std::exception衍生而来的一个类，它允许我们通过一个try-catch结构来捕捉 SYCL 错误，如图 5-6 所示。

img/489625_1_En_5_Fig6_HTML.png

图 5-6

具体要抓的模式sycl::exception

在 C++ 错误处理机制之上，SYCL 为运行时抛出的异常添加了一个sycl::exception类型。其他的都是标准的 C++ 异常处理，所以大多数开发人员都很熟悉。

图 5-7 提供了一个稍微完整的例子，其中处理了额外的异常类，以及通过从main()返回而结束的程序。

img/489625_1_En_5_Fig7_HTML.png

图 5-7

从代码块中捕捉异常的模式

异步错误处理

异步错误由 SYCL 运行时(或底层后端)检测，错误的发生与宿主程序中命令的执行无关。这些错误存储在 SYCL 运行时内部的列表中，只在程序员可以控制的特定点上进行处理。为了涵盖异步错误的处理，我们需要讨论两个主题:

当有未完成的异步错误要处理时调用的异步处理程序
当调用异步处理程序时

异步处理程序

异步处理程序是应用程序定义的函数，它向 SYCL 上下文和/或队列注册。在下一节定义的时间，如果有任何未处理的异步异常可供处理，那么 SYCL 运行时将调用异步处理程序，并向其传递这些异常的列表。

异步处理程序作为一个std::function传递给一个上下文或队列构造器，并且可以根据我们的偏好以常规函数、lambda 或仿函数等方式定义。处理程序必须接受一个sycl::exception_list参数，例如图 5-8 中所示的示例处理程序。

img/489625_1_En_5_Fig8_HTML.png

图 5-8

定义为 lambda 的异步处理程序实现示例

在图 5-8 中，std::rethrow_exception后接特定异常类型的 catch 提供了异常类型的过滤，在这种情况下只过滤到sycl::exception。我们还可以在 C++ 中使用其他过滤方法，或者选择处理所有异常，而不管其类型。

该处理程序在构建时与一个队列或上下文相关联(在第六章中详细介绍了底层细节)。例如，要用我们正在创建的队列注册图 5-8 中定义的处理程序，我们可以写


queue my_queue{ gpu_selector{}, handle_async_error };

同样，要用我们正在创建的上下文注册图 5-8 中定义的处理程序，我们可以写


context my_context{ handle_async_error };

大多数应用程序不需要显式创建或管理上下文(它们是在后台自动为我们创建的)，因此如果要使用异步处理程序，大多数开发人员应该将这种处理程序与为特定设备(而不是显式上下文)构建的队列相关联。

在定义异步处理程序时，大多数开发人员应该在队列中定义它们，除非出于其他原因已经显式地管理了上下文。

如果没有为队列或队列的父上下文定义异步处理程序，并且在该队列上(或上下文中)发生了必须处理的异步错误，则调用默认的异步处理程序。默认处理程序的运行方式如同图 5-9 所示的编码。

img/489625_1_En_5_Fig9_HTML.png

图 5-9

默认异步处理程序的行为示例

默认处理程序应该向用户显示一些异常列表中的错误信息，然后异常终止应用程序，这也会导致操作系统报告终止异常。

我们在异步处理程序中放什么由我们自己决定。它的范围可以从记录错误到应用程序终止，再到恢复错误条件，以便应用程序可以继续正常执行。常见的情况是通过调用sycl::exception::what()来报告错误的任何细节，然后终止应用程序。

尽管由我们来决定异步处理程序在内部做什么，但一个常见的错误是打印一条错误消息(在程序的其他消息中可能会被忽略)，然后完成处理程序函数。除非我们有适当的错误管理原则，允许我们恢复已知的程序状态，并确信继续执行是安全的，否则我们应该考虑在异步处理函数中终止应用程序。这减少了错误结果出现在程序中的机会，在该程序中检测到错误，但是应用程序被无意中允许继续执行。在许多程序中，一旦我们遇到异步异常，异常终止是首选结果。

如果没有全面的错误恢复和管理机制，在输出有关错误的信息后，考虑在异步处理程序中终止应用程序。

处理程序的调用

运行时在特定的时间调用异步处理程序。错误发生时不会立即报告，因为如果出现这种情况，错误管理和安全应用程序编程(尤其是多线程)将变得更加困难和昂贵。相反，异步处理程序在以下特定时间被调用:

当宿主程序调用特定队列上的queue::throw_asynchronous()时
当宿主程序调用特定队列上的queue::wait_and_throw()时
当宿主程序在特定事件上调用event::wait_and_throw()时
当一个queue被破坏时
当一个context被破坏时

方法 1–3 为宿主程序提供了一种机制来控制何时处理异步异常，以便可以管理特定于应用程序的线程安全和其他细节。它们有效地提供了异步异常进入宿主程序控制流的受控点，并且可以像处理同步错误一样进行处理。

如果用户没有显式调用方法 1-3 中的一个，那么当队列和上下文被销毁时，在程序拆卸期间通常会报告异步错误。这通常足以向用户发出信号，表明出了问题，程序结果不应该被信任。

然而，在程序拆卸期间依靠错误检测并不是在所有情况下都有效。例如，如果程序将仅在达到某些算法收敛标准时终止，并且如果这些标准仅可通过成功执行设备内核来实现，则异步异常可能发信号通知该算法将永远不会收敛并开始拆卸(将会注意到该错误)。在这些情况下，以及在有更完整的错误处理策略的生产应用中，在程序中的常规和受控点调用throw_asynchronous()或wait_and_throw()是有意义的(例如，在检查算法收敛是否发生之前)。

设备上的错误

本章中讨论的错误检测和处理机制是基于主机的。它们是一些机制，通过这些机制，主机程序可以检测和处理在主机程序中或者在设备上执行内核期间可能出现的错误。我们还没有介绍的是，如何从我们编写的设备代码中发出信号，表明有什么地方出错了。这种遗漏不是错误，而是相当故意的。

SYCL 明确禁止在设备代码中使用 C++ 异常处理机制(比如throw)，因为对于某些类型的设备来说，这是我们通常不想付出的性能代价。如果我们检测到设备代码中出现了错误，我们应该使用现有的非基于异常的技术发出错误信号。例如，我们可以写入一个记录错误的缓冲区，或者从我们定义的表示发生了错误的数值计算中返回一些无效的结果。在这些情况下，正确的策略是非常具体的应用。

摘要

在这一章中，我们介绍了同步和异步错误，讨论了如果我们对可能发生的错误无所作为时的默认行为，并讨论了在应用程序的受控点处理异步错误的机制。错误管理策略是软件工程中的一个主要话题，也是许多应用程序中编写的代码的重要组成部分。SYCL 集成了我们在错误处理方面已经掌握的 C++ 知识，并提供了灵活的机制来集成我们首选的错误管理策略。

Creative Commons

六、统一共享内存

img/489625_1_En_6_Figa_HTML.gif

接下来的两章深入探讨了如何管理数据。有两种互为补充的不同方法:统一共享内存(USM)和缓冲区。USM 公开了与缓冲区不同的内存抽象级别—USM 有指针，而缓冲区是更高级别的接口。本章重点介绍 USM。下一章将集中讨论缓冲区。

除非我们明确知道要使用缓冲区，否则 USM 是一个很好的起点。USM 是一个基于指针的模型，允许通过常规 C++ 指针读写内存。

我们为什么要使用 USM？

因为 USM 是基于 C++ 指针的，所以对于现有的基于指针的 C++ 代码来说，它是一个自然的起点。将指针作为参数的现有函数无需修改即可继续工作。在大多数情况下，唯一需要的改变是用特定于 USM 的分配例程替换现有的对malloc或new的调用，我们将在本章后面讨论这些例程。

分配类型

虽然 USM 基于 C++ 指针，但并非所有指针都是平等的。USM 定义了三种不同类型的分配，每种类型都有独特的语义。设备可能不支持所有类型(甚至任何类型)的 USM 分配。稍后我们将学习如何查询设备支持什么。图 6-1 总结了三种分配类型及其特点。

img/489625_1_En_6_Fig1_HTML.png

图 6-1

USM 分配类型

设备分配

这第一种类型的分配是我们需要的，以便有一个指向设备的附加存储器的指针，例如(G)DDR 或 HBM。设备分配可以由运行在设备上的内核读取或写入，但不能从主机上执行的代码直接访问。试图访问主机上的设备分配可能会导致数据不正确或程序因错误而崩溃。我们必须使用显式 USM memcpy机制在主机和设备之间复制数据，该机制指定了在两个位置之间必须复制多少数据，这将在本章的后面部分介绍。

主机分配

第二种类型的分配比设备分配更容易使用，因为我们不必在主机和设备之间手动拷贝数据。主机分配是主机内存中的分配，可在主机和设备上访问。这些分配虽然可以在设备上访问，但不能迁移到设备的附加内存。取而代之的是，读写这个内存的内核是远程完成的*，通常是通过较慢的总线，比如 PCI-Express。便利性和性能之间的权衡是我们必须考虑的。尽管主机分配会导致更高的访问成本，但仍然有充分的理由使用它们。示例包括很少访问的数据或无法容纳在设备附加内存中的大型数据集。*

*### 共享分配

最后一种分配结合了设备和主机分配的属性，将主机分配的程序员便利性与设备分配提供的更高性能结合在一起。与主机分配一样，共享分配在主机和设备上都是可访问的。它们之间的区别在于，共享分配可以在主机内存和设备连接内存之间自由迁移，自动进行，无需我们的干预。如果某个分配已经迁移到该设备，则在该设备上执行的任何内核访问该分配的性能都将优于从主机远程访问该分配。然而，共享分配并不能给我们所有的好处而没有任何缺点。

自动迁移可以通过多种方式实现。无论运行时选择哪种方式来实现共享分配，它们通常都要付出延迟增加的代价。通过设备分配，我们可以准确地知道需要复制多少内存，并可以安排尽快开始复制。自动迁移机制看不到未来，在某些情况下，直到内核试图访问数据时才开始移动数据。然后，内核必须等待或阻塞，直到数据移动完成，然后才能继续执行。在其他情况下，运行时可能不知道内核将访问多少数据，并且可能保守地移动比所需数量更多的数据，这也增加了内核的延迟。

我们还应该注意，虽然共享分配可以迁移，但这并不一定意味着 DPC++ 的所有实现都将迁移它们。我们预计大多数实现都将共享分配与迁移一起实现，但是一些设备可能更喜欢将它们实现为与主机分配相同。在这样的实现中，分配在主机和设备上仍然可见，但是我们可能看不到迁移实现可以提供的性能提升。

分配内存

USM 允许我们以各种不同的方式分配内存，以满足不同的需求和偏好。然而，在我们更详细地讨论所有方法之前，我们应该讨论 USM 分配与常规 C++ 分配有何不同。

我们需要知道什么？

常规的 C++ 程序可以通过多种方式分配内存:new、malloc或分配器。无论我们喜欢哪种语法，内存分配最终都是由主机操作系统中的系统分配器来执行的。当我们在 C++ 中分配内存时，唯一关心的是“我们需要多少内存？”以及“有多少内存可供分配？”但是，USM 需要额外的信息才能执行分配。

首先，USM 分配需要指定所需的分配类型:设备、主机或共享。请求正确的分配类型是很重要的，以便获得该分配所需的行为。接下来，每个 USM 分配必须指定一个context对象，分配将针对该对象进行。context对象还没有太多的讨论，所以这里值得说一点。上下文代表我们可以在其上执行内核的一个或一组设备。我们可以把上下文看作是一个方便的地方，让运行时保存一些关于它正在做什么的状态。在大多数 DPC++ 程序中，除了传递上下文之外，程序员不太可能直接与上下文交互。

USM 分配不能保证在不同的上下文中可用——所有 USM 分配、队列和内核共享同一个context对象是很重要的。通常，我们可以从用于向设备提交工作的队列中获得这个上下文。最后，device分配还要求我们指定哪个设备将为分配提供内存。这一点很重要，因为我们不想超额预订设备的内存(除非设备能够支持这一点，我们将在本章稍后讨论数据迁移时对此进行详细说明)。通过添加这些额外的参数，可以将 USM 分配例程与它们的 C++ 类似物区分开来。

多种风格

有时候，试图用一个单一的选项来取悦每个人被证明是一个不可能的任务，就像有些人喜欢咖啡胜过茶，或者喜欢emacs胜过vi.如果我们问程序员分配接口应该是什么样子，我们会得到几个不同的答案。USM 支持这种多样性的选择，并提供了几种不同风格的分配界面。这些不同的风格是 C 风格、C++ 风格和 C++ 分配器风格。我们现在将讨论每一个并指出它们的相似之处和不同之处。

c 级津贴

第一种类型的分配函数(在图 6-2 中列出，稍后在图 6-6 和 6-7 中显示的示例中使用)是在 C: malloc函数中的内存分配之后建模的，这些函数采用多个字节进行分配并返回一个void *指针。这种类型的函数是类型不可知的。我们必须指定要分配的总字节数，这意味着如果我们想要分配类型为X的N对象，我们必须请求N * sizeof(X)总字节数。返回的指针属于类型void *，这意味着我们必须将它转换为类型X的适当指针。这种样式非常简单，但是由于需要进行大小计算和类型转换，可能会很冗长。

我们可以进一步将这种分配方式分为两类:命名函数和单一函数。这两种风格的区别在于我们如何指定所需的 USM 分配类型。对于命名函数(malloc_device、malloc_host和malloc_shared)，USM 分配的类型编码在函数名中。单一功能malloc要求将 USM 分配类型指定为附加参数。没有一种味道比另一种更好，选择使用哪一种取决于我们的偏好。

我们不能在不简要提及对齐的情况下继续讨论。每个版本的malloc也有一个对应的aligned_alloc。malloc函数返回与设备默认行为一致的内存。它将返回一个具有有效对齐方式的合法指针，但是在某些情况下，我们可能更愿意手动指定对齐方式。在这些情况下，我们应该使用aligned_alloc变量中的一个，它也要求我们为分配指定期望的对齐。如果我们指定了一个非法的对齐，就不要指望程序能正常工作！合法对齐是 2 的幂。值得注意的是，在许多设备上，分配是最大限度地对齐的，以对应于硬件的功能，因此尽管我们可能要求分配是 4、8、16 或 32 字节对齐的，但实际上我们可能会看到更大的对齐，这给了我们所要求的，甚至更多。

img/489625_1_En_6_Fig2_HTML.png

图 6-2

c 风格的 USM 分配功能

C++ 分配

USM 分配函数的下一种风格(在图 6-3 中列出)与第一种非常相似，但更多的是 C++ 的外观和感觉。我们再次拥有了分配例程的命名和单个函数版本，以及我们的默认和用户指定的对齐版本。不同之处在于，现在我们的函数是 C++ 模板化的函数，它分配类型为T的Count对象，并返回类型为T *的指针。利用现代 C++ 简化了事情，因为我们不再需要以字节为单位手动计算分配的总大小，或者将返回的指针转换为适当的类型。这也有助于在代码中生成更紧凑、更不易出错的表达式。然而，我们应该注意到，与 C++ 中的“new”不同，malloc 风格的接口不为被分配的对象调用构造器——我们只是分配足够的字节来适应该类型。

这种类型的分配是用 USM 编写新代码的良好开端。对于已经大量使用 C 或 C++ malloc的现有 C++ 代码来说，前面的 C 风格是一个很好的起点，我们将在其中添加 USM 的使用。

img/489625_1_En_6_Fig3_HTML.png

图 6-3

C++ 风格的 USM 分配函数

C++ 分配器

USM 分配的最终版本(图 6-4 )比之前的版本更加拥抱现代 C++。这种风格基于 C++ 分配器接口，该接口定义了用于在容器(如std::vector)中直接或间接执行内存分配的对象。如果我们的代码大量使用可以对用户隐藏内存分配和释放细节的容器对象，这种分配器风格是最有用的，简化了代码并减少了出错的机会。

img/489625_1_En_6_Fig4_HTML.png

图 6-4

C++ 分配器风格的 USM 分配函数

释放内存

程序分配的任何东西最终都必须被释放。USM 定义了一个free方法来释放由malloc或aligned_malloc函数分配的内存。这个free方法还将分配内存的上下文作为一个额外的参数。队列也可以代替上下文。如果内存是用 C++ 分配器对象分配的，那么也应该使用该对象来释放内存。

img/489625_1_En_6_Fig5_HTML.png

图 6-5

三种分配方式

分配示例

在图 6-5 中，我们展示了如何使用刚刚描述的三种风格来执行相同的分配。在这个例子中，我们将N单精度浮点数分配为共享分配。第一次分配f1使用 C 风格的void *返回 malloc 例程。对于这种分配，我们显式地传递从队列中获得的设备和上下文。我们还必须将结果强制转换回一个float *。第二个分配f2做了同样的事情，但是使用了 C++ 风格的模板 malloc。因为我们将元素的类型float传递给分配例程，所以我们只需要指定我们想要分配多少个浮点数，而不需要对结果进行强制转换。我们还使用接受队列而不是设备和上下文的形式，产生了一个非常简单和紧凑的语句。第三个分配f3使用 USM C++ 分配器类。我们实例化适当类型的分配器对象，然后使用该对象执行分配。最后，我们展示如何正确地释放每个分配。

数据管理

现在我们已经了解了如何使用 USM 分配内存，我们将讨论如何管理数据。我们可以从两个方面来看这个问题:数据初始化和数据移动。

初始化

数据初始化涉及到在我们执行计算之前用值填充我们的内存。常见初始化模式的一个例子是在使用分配之前用零填充分配。如果我们要使用 USM 分配来做到这一点，我们可以通过多种方式来实现。首先，我们可以编写一个内核来做这件事。如果我们的数据集特别大，或者初始化需要复杂的计算，这是一种合理的方法，因为初始化可以并行执行(并且它使初始化的数据准备就绪，可以在设备上运行)。第二，我们可以在分配的所有元素上实现一个循环，将每个元素设置为零。然而，这种方法有一个潜在的问题。对于主机分配和共享分配，循环可以很好地工作，因为它们在主机上是可访问的。然而，因为设备分配在主机上是不可访问的，所以主机代码中的循环将不能写入它们。这让我们想到了第三个选择。

memset函数旨在有效地实现这种初始化模式。USM 提供了一个版本的memset，它是handler和queue类的成员函数。它有三个参数:表示我们要设置的内存基址的指针，表示要设置的字节模式的字节值，以及要设置为该模式的字节数。与主机上的循环不同，memset并行发生，也与device分配一起工作。

虽然memset是一个有用的操作，但它只允许我们指定一个字节模式来填充分配，这是相当有限的。USM 还提供了一个fill方法(作为handler和queue类的成员),让我们用任意模式填充内存。fill 方法是一个函数，它以我们想要写入分配的模式类型为模板。用一个int模板化它，我们可以用数字“42”填充一个分配。类似于memset , fill有三个参数:指向要填充的分配基址的指针、要填充的值以及我们希望将该值写入分配的次数。

数据传送

数据移动可能是 USM 需要理解的最重要的方面。如果正确的数据没有在正确的时间出现在正确的位置，我们的程序就会产生错误的结果。USM 定义了我们可以用来管理数据的两种策略:显式和隐式。选择我们想要使用的策略与我们的硬件支持的 USM 分配类型或我们想要使用的类型有关。

明确的

USM 提供的第一个策略是显式数据移动(图 6-6 )。这里，我们必须在主机和设备之间显式复制数据。我们可以通过调用handler和queue类中的memcpy方法来实现。memcpy方法有三个参数:一个指向目标内存的指针，一个指向源内存的指针，以及要在主机和设备之间复制的字节数。我们不需要指定复制应该在哪个方向发生—这在源和目标指针中是隐含的。

显式数据移动的最常见用法是在 USM 中向/从device分配拷贝数据，因为它们在主机上不可访问。必须插入数据的显式拷贝确实需要我们付出努力。此外，它也可能是错误的来源:副本可能被意外忽略，可能复制了不正确的数据量，或者源或目标指针可能不正确。

然而，显式数据移动不仅有缺点。这给了我们很大的优势:对数据移动的完全控制。在某些应用程序中，控制复制多少数据以及何时复制数据对于实现最佳性能非常重要。理想情况下，我们可以尽可能将计算与数据移动重叠，确保硬件以高利用率运行。

其他类型的 USM 分配host和shared都可以在主机和设备上访问，不需要显式复制到设备。这让我们想到了 USM 中的另一种数据移动策略。

img/489625_1_En_6_Fig6_HTML.png

图 6-6

USM 显式数据移动示例

隐形的

USM 提供的第二种策略是隐式数据移动(示例用法如图 6-7 所示)。在这种策略中，数据移动以隐含的方式发生，也就是说，不需要我们的输入。使用隐式数据移动，我们不需要插入对memcpy的调用，因为我们可以通过 USM 指针直接访问数据，无论我们想在哪里使用它。相反，确保数据在被使用时在正确的位置可用成为系统的工作。

对于主机分配，人们可能会争论它们是否真的会导致数据移动。根据定义，它们始终是指向主机内存的指针，因此由给定主机指针表示的内存不能存储在设备上。但是，在设备上访问主机分配时，确实会发生数据移动。我们读取或写入的值通过适当的接口传入或传出内核，而不是将内存迁移到设备。这对于数据不需要驻留在设备上的流式内核非常有用。

隐式数据移动主要涉及 USM 共享分配。这种类型的分配在主机和设备上都可以访问，更重要的是，可以在主机和设备之间迁移。关键在于，这种迁移是自动发生的，或者说是隐式发生的，只需访问不同位置的数据即可。接下来，我们将讨论在为共享分配进行数据迁移时需要考虑的几个问题。

img/489625_1_En_6_Fig7_HTML.png

图 6-7

USM 隐式数据移动示例

移动

通过显式数据移动，我们可以控制发生多少数据移动。使用隐式数据移动，系统会为我们处理这一点，但它可能不会这样高效。DPC++ 运行时不是一个 Oracle——它不能在应用程序访问数据之前预测应用程序将访问哪些数据。此外，指针分析对于编译器来说仍然是一个非常困难的问题，编译器可能无法准确地分析和识别内核中可能使用的每个分配。因此，隐式数据移动机制的实现可能会根据支持 USM 的设备的功能做出不同的决定，这既会影响共享分配的使用方式，也会影响其执行方式。

如果一个设备非常强大，它可能能够按需迁移内存。在这种情况下，在主机或设备尝试访问当前不在所需位置的分配后，会发生数据移动。按需数据极大地简化了编程，因为它提供了所需的语义，即 USM 共享指针可以在任何地方访问并正常工作。如果一个设备不支持按需迁移(第十二章解释了如何查询一个设备的能力)，它可能仍然能够保证相同的语义，但对如何使用共享指针有额外的限制。

USM 共享分配的受限形式决定了何时何地可以访问共享指针，以及共享分配可以有多大。如果设备不能按需迁移内存，这意味着运行时必须保守，并假设内核可以访问其设备附加内存中的任何分配。这带来了几个后果。

首先，这意味着主机和设备不应试图同时访问共享分配。应用程序应该分阶段交替访问。主机可以访问分配，然后内核可以使用该数据进行计算，最后主机可以读取结果。如果没有这种限制，主机可以自由地访问内核当前接触的分配的不同部分。这种并发访问通常发生在设备存储器页面的粒度上。主机可以访问一个页面，而设备可以访问另一个页面。原子地访问同一块数据将在第十九章中介绍。

这种受限形式的共享分配的下一个后果是，分配受到连接到设备的内存总量的限制。如果设备不能按需迁移内存，它就不能将数据迁移到主机来腾出空间引入不同的数据。如果设备支持按需迁移，则有可能超额订阅其连接的内存，允许内核计算超过设备内存正常容量的数据，尽管这种灵活性可能会因额外的数据移动而带来性能损失。

细粒度控制

当设备支持共享分配的按需迁移时，在当前不驻留内存的位置访问内存后，会发生数据移动。但是，在等待数据移动完成时，内核可能会停止。它执行的下一条语句甚至可能导致更多的数据移动，并给内核执行带来额外的延迟。

DPC++ 为我们提供了一种修改自动迁移机制性能的方法。它通过定义两个函数来做到这一点:prefetch和mem_advise。图 6-8 显示了每种方法的简单应用。这些函数让我们向运行时提示内核将如何访问数据，这样运行时就可以选择在内核试图访问数据之前开始移动数据。注意，这个例子使用了队列快捷方式方法，这些方法直接调用queue对象上的parallel_for，而不是在传递给submit方法(一个命令组)的 lambda 内部调用。

img/489625_1_En_6_Fig8_HTML.png

图 6-8

通过prefetch和mem_advise进行精细控制

对我们来说，最简单的方法就是调用prefetch。这个函数作为handler或queue类的成员函数被调用，并接受一个基指针和字节数。这让我们可以通知运行时，某个设备上将要使用某些数据，以便它可以急切地开始迁移这些数据。理想情况下，我们应该足够早地发出这些预取提示，以便在内核接触数据时，它已经驻留在设备上，从而消除我们之前描述的延迟。

DPC++ 提供的另一个函数是mem_advise。这个函数允许我们提供特定于设备的关于内核如何使用内存的提示。我们可以指定的这种可能的建议的一个例子是，数据将只在内核中读取，而不是写入。在这种情况下，系统可以意识到它可以复制设备上的数据，这样在内核完成后就不需要更新主机的版本。然而，传递给mem_advise的建议是特定于特定设备的，因此在使用该功能之前，请务必查看硬件文档。

问题

最后，并非所有设备都支持 USM 的所有功能。如果我们希望我们的程序可以在不同的设备上移植，我们不应该假设所有的 USM 特性都是可用的。USM 定义了我们可以查询的几项内容。这些查询可以分为两类:指针查询和设备能力查询。图 6-9 显示了每种方法的简单应用。

USM 中的指针查询回答了两个问题。第一个问题是“这个指针指向什么类型的 USM 分配？”get_pointer_type函数接受一个指针和 DPC++ 上下文并返回一个类型为usm::alloc的结果，它可以有四个可能的值:主机、设备、共享或未知。第二个问题是“这个 USM 指针是针对哪个设备分配的？”我们可以向函数get_pointer_device传递一个指针和一个上下文，并获取一个设备对象。这主要用于设备或共享 USM 分配，因为它对主机分配没有多大意义。

USM 提供的第二种类型的查询涉及设备的功能。USM 通过在设备对象上调用get_info来扩展可以查询的设备信息描述符列表。这些查询可用于测试设备支持哪些类型的 USM 分配。此外，我们可以通过本章前面介绍的方式查询设备上的共享分配是否受到限制。完整的查询列表如图 6-10 所示。在第十二章中，我们将更详细地了解查询机制。

img/489625_1_En_6_Fig10_HTML.png

图 6-10

USM 设备信息描述符

img/489625_1_En_6_Fig9_HTML.png

图 6-9

对 USM 指针和设备的查询

摘要

在这一章中，我们描述了统一共享内存，一种基于指针的数据管理策略。我们讨论了 USM 定义的三种分配类型。我们讨论了使用 USM 分配和取消分配内存的所有不同方式，以及如何由我们(程序员)对设备分配进行显式控制，或者由系统对共享分配进行隐式控制。最后，我们讨论了如何查询设备支持的不同 USM 功能，以及如何查询程序中关于 USM 指针的信息。

因为我们还没有在本书中详细讨论同步，所以在后面的章节中，当我们讨论调度、通信和同步时，会有更多关于 USM 的内容。具体来说，我们在第 8 、 9 和 19 章中涵盖了 USM 的这些额外考虑。

在下一章，我们将讨论数据管理的第二个策略:缓冲区。

Creative Commons

七、缓冲

img/489625_1_En_7_Figa_HTML.gif

在这一章中，我们将学习缓冲抽象。在前一章中，我们学习了统一共享内存(USM)，这是一种基于指针的数据管理策略。USM 迫使我们思考内存在哪里，什么应该在哪里可以访问。缓冲区抽象是一个更高级的模型，它对程序员隐藏了这一点。缓冲区只是表示数据，管理数据在内存中的存储和移动就成了运行时的工作。

本章介绍了管理数据的另一种方法。缓冲区和 USM 之间的选择通常取决于个人偏好和现有代码的风格，应用程序可以自由混合和匹配这两种风格来表示应用程序中的不同数据。

USM 只是公开了不同的内存抽象。USM 有指针，缓冲区是更高层次的抽象。缓冲区的抽象级别允许在应用程序中的任何设备上使用其中包含的数据，其中运行时管理使数据可用所需的任何内容。选择是好的，所以让我们进入缓冲区。

我们将更仔细地研究如何创建和使用缓冲区。如果不讨论存取器，对缓冲区的讨论将是不完整的。虽然缓冲区抽象了我们如何在程序中表示和存储数据，但我们并不使用缓冲区直接访问数据。相反，我们使用访问器对象来通知运行时我们打算如何使用我们正在访问的数据，并且访问器与任务图中强大的数据依赖机制紧密耦合。在我们讲述了我们可以用缓冲区做的所有事情之后，我们还将探索如何在我们的程序中创建和使用访问器。

缓冲

缓冲区是数据的高级抽象。缓冲区不一定局限于单个位置或虚拟内存地址。实际上，运行时可以自由地使用内存中许多不同的位置(甚至跨不同的设备)来表示一个缓冲区，但是运行时必须确保总是给我们一个一致的数据视图。可以在主机和任何设备上访问缓冲区。

img/489625_1_En_7_Fig1_HTML.png

图 7-1

缓冲类定义

buffer类是一个模板类，有三个模板参数，如图 7-1 所示。第一个模板参数是缓冲区将包含的对象的类型。按照 C++ 的定义，这个类型必须是可简单复制的，这基本上意味着不使用任何特殊的复制或移动构造器就可以安全地逐字节复制这个对象。下一个模板参数是一个描述缓冲区维数的整数。最后一个模板参数是可选的，默认值通常是所使用的值。该参数指定了一个 C++ 风格的分配器类，用于在主机上执行缓冲区所需的任何内存分配。首先，我们将研究创建缓冲区对象的许多方法。

创造

在下图中，我们展示了创建缓冲区对象的几种方法。如何在应用程序代码中创建缓冲区的选择是需要如何使用缓冲区和个人编码偏好的组合。让我们浏览一下这个例子，看看每个实例。

img/489625_1_En_7_Fig2_HTML.png

图 7-2

创建缓冲区，第一部分

我们在图 7-2 ，b1中创建的第一个缓冲区是一个包含十个整数的二维缓冲区。我们显式传递所有模板参数，甚至显式传递默认值buffer_allocator作为分配器类型。然而，使用现代 C++，我们可以更简洁地表达这一点。缓冲区b2也是一个使用默认分配器的十个整数的二维缓冲区。这里我们利用 C++17 的类模板参数演绎(CTAD)来自动推断我们必须表达的模板参数。CTAD 是一个要么全有要么全无的工具——它要么推断一个类的每个模板参数，要么一个都不推断。在这种情况下，我们使用一个带两个参数的范围来初始化b2来推断它是一个二维范围。分配器模板参数有一个默认值，所以我们在创建缓冲区时不需要显式地列出它。

通过 buffer b3，我们创建了一个 20 浮点的缓冲区，并使用默认构造的std: :allocator<float>来分配主机上任何必要的内存。当使用带有缓冲区的自定义分配器类型时，我们通常希望将实际的分配器对象传递给缓冲区来使用，而不是默认构造的分配器对象。Buffer b4展示了如何做到这一点，在对其构造器的调用中，在范围之后获取分配器对象。

对于我们示例中的前四个缓冲区，我们让缓冲区分配它需要的任何内存，并且不在创建它们时用任何值初始化数据。使用缓冲区有效地包装现有的 C++ 分配是一种常见的模式，这些分配可能已经用数据进行了初始化。我们可以通过将初始值的源传递给缓冲区构造器来实现这一点。这样做允许我们做几件事，我们将在下一个例子中看到。

img/489625_1_En_7_Fig3_HTML.png

图 7-3

创建缓冲区，第二部分

在图 7-3 中，buffer b5创建了一个四个双精度的一维缓冲区。除了指定缓冲区大小的范围之外，我们还将指向 C 数组myDoubles的主机指针传递给缓冲区构造器。这里我们可以充分利用 CTAD 来推断我们缓冲区的所有模板参数。我们传递的主机指针指向 doubles，这给了我们缓冲区的数据类型。维数是从一维范围自动推断出来的，一维范围本身是推断出来的，因为它是用一个数创建的。最后，使用默认的分配器，所以我们不必指定它。

传递一个主机指针有一些我们应该知道的分支。通过传递一个指向主机内存的指针，我们向运行时承诺，在缓冲区的生存期内，我们不会尝试访问主机内存。SYCL 实施不会(也不能)强制执行这一点——我们有责任确保不违反此合同。我们不应该在缓冲区活动时尝试访问该内存的一个原因是，缓冲区可能会选择使用主机上的不同内存来表示缓冲区内容，这通常是出于优化的原因。如果是这样，这些值将从主机指针复制到这个新的内存中。如果后续内核修改了缓冲区，则原始主机指针将不会反映更新后的值，直到某些指定的同步点。在这一章的后面，我们将更多地讨论数据何时被写回主机指针。

缓冲器b6与缓冲器b5非常相似，但有一个主要区别。这一次，我们用一个指向const double的指针初始化缓冲区。这意味着我们只能通过主机指针读取值，而不能写入值。然而，本例中我们的缓冲区类型仍然是double，而不是const double，因为扣除指南没有考虑const- ness。这意味着缓冲区可以被内核写入，但是在缓冲区过期后，我们必须使用不同的机制来更新主机(这将在本章后面讨论)。

也可以使用 C++ 共享指针对象初始化缓冲区。如果我们的应用程序已经使用了共享指针，这是很有用的，因为这种初始化方法将正确地计算引用，并确保内存不会被释放。Buffer b7从单个整数初始化一个缓冲区，并使用共享指针初始化。

img/489625_1_En_7_Fig4_HTML.png

图 7-4

创建缓冲区，第三部分

容器是现代 C++ 应用程序中常用的，例子包括std::array、std::vector、std::list或std::map。我们可以用两种不同的方式使用容器初始化一维缓冲区。第一种方式，如图 7-4 缓冲区b8所示，使用输入迭代器。我们将两个迭代器而不是主机指针传递给缓冲区构造器，一个表示数据的开始，另一个表示结束。缓冲区的大小是通过递增起始迭代器直到它等于结束迭代器返回的元素数来计算的。这对于任何实现 C++ InputIterator接口的数据类型都很有用。如果为缓冲区提供初始值的容器对象也是连续的，那么我们可以使用更简单的形式来创建缓冲区。Buffer b9通过简单地将向量传递给构造器来创建一个缓冲区。缓冲区的大小由用来初始化它的容器的大小决定，缓冲区数据的类型来自容器数据的类型。使用这种方法创建缓冲区是常见的，并推荐使用容器，如std::vector和std: :array。

缓冲区创建的最后一个例子说明了 buffer 类的另一个特性。可以从另一个缓冲区或子缓冲区创建一个缓冲区的视图。子缓冲区需要三样东西:对父缓冲区的引用、基索引和子缓冲区的范围。不能从子缓冲区创建子缓冲区。可以从同一个缓冲区创建多个子缓冲区，并且它们可以自由重叠。Buffer b10的创建与 buffer b2完全一样，是一个二维整数缓冲区，每行有五个整数。接下来，我们从缓冲区b10创建两个子缓冲区，子缓冲区b11和b12。子缓冲器b11从索引(0,0)开始，包含第一行中的每个元素。类似地，子缓冲区b12从索引(1,0)开始，包含第二行中的每个元素。这产生了两个不相交的子缓冲器。由于子缓冲区不重叠，不同的内核可以同时在不同的子缓冲区上运行，但是我们将在下一章更多地讨论调度执行图和依赖关系。

img/489625_1_En_7_Fig5_HTML.png

图 7-5

缓冲区属性

缓冲区也可以用改变其行为的特殊属性来创建。在图 7-5 中，我们将浏览三个不同可选缓冲属性的示例，并讨论如何使用它们。请注意，这些属性在大多数代码中相对不常见。

使用主机指针

在缓冲区创建期间，可以选择指定的第一个属性是use_host_ptr。当存在时，该属性要求缓冲区不在主机上分配任何内存，并且在缓冲区构造上传递或指定的任何分配器实际上都被忽略。相反，缓冲区必须使用传递给构造器的主机指针所指向的内存。请注意，这并不要求设备使用相同的内存来保存缓冲区的数据。设备可以自由地将缓冲区的内容缓存到它所连接的内存中。另请注意，该属性只能在将主机指针传递给构造器时使用。当程序希望完全控制所有主机内存分配时，此选项会很有用。

在图 7-5 的例子中，我们创建了一个缓冲区b，就像我们在前面的例子中看到的那样。接下来我们创建缓冲区b1，并用一个指向myInts的指针初始化它。我们还传递属性use_host_ptr，这意味着缓冲区b1将只使用myInts指向的内存，而不会分配任何额外的临时存储。

使用互斥锁

下一个属性use_mutex，涉及缓冲区和主机代码之间的细粒度内存共享。缓冲区b2是使用这个属性创建的。该属性引用了一个 mutex 对象，稍后可以从缓冲区中查询该对象，如我们在示例中所见。此属性还要求将一个主机指针传递给构造器，它让运行库确定何时通过提供的主机指针访问主机代码中的更新值是安全的。在运行时保证主机指针看到缓冲区的最新值之前，我们不能锁定互斥体。虽然这可以与use_host_ptr属性合并，但这不是必需的。use_mutex是一种机制，允许主机代码在缓冲区仍然存在时访问缓冲区内的数据，而不使用主机访问器机制(稍后描述)。一般来说，除非我们有特定的理由使用互斥体，否则应该首选主机访问器机制，特别是因为无法保证在成功锁定互斥体和数据准备好供主机代码使用之前需要多长时间。

上下文绑定

在我们的示例中，最后一个属性显示在缓冲区b3的创建中。这里，我们的 42 个整数的缓冲区是用context_bound属性创建的。属性采用对上下文对象的引用。通常，缓冲区可以在任何设备或上下文中自由使用。但是，如果使用此属性，它会将缓冲区锁定到指定的上下文。试图在另一个上下文中使用缓冲区将导致运行时错误。例如，通过识别内核可能被提交到错误队列的情况，这对于调试程序可能是有用的。实际上，我们并不期望在许多程序中使用这个属性，缓冲区在任何上下文中的任何设备上被访问的能力是缓冲区抽象的最强大的属性之一(这个属性撤销了这个属性)。

我们能用缓冲器做什么？

使用缓冲区对象可以做很多事情。我们可以查询缓冲区的特征，确定在缓冲区被破坏后是否有任何数据被写回主机内存以及在哪里，或者将缓冲区重新解释为具有不同特征的缓冲区。然而，有一件事是不能做的，那就是直接访问缓冲区所代表的数据。相反，我们必须创建访问器对象来访问数据，我们将在本章的后面了解这一点。

可以对缓冲区进行查询的示例包括它的范围、它所代表的数据元素的总数以及存储其元素所需的字节数。我们还可以查询缓冲区正在使用哪个分配器对象，以及该缓冲区是否是子缓冲区。

当缓冲区被破坏时更新主机内存是使用缓冲区时要考虑的一个重要方面。根据缓冲区的创建方式，在缓冲区销毁后，主机内存可能会更新，也可能不会更新计算结果。如果从指向非const数据的主机指针创建并初始化缓冲区，则当缓冲区被销毁时，用更新的数据更新该指针。然而，还有一种方法可以更新主机内存，而不管缓冲区是如何创建的。set_final_data方法是buffer的模板方法，可以接受原始指针、C++ OutputIterator或std::weak_ptr。当缓冲区被销毁时，缓冲区包含的数据将使用提供的位置写入主机。注意，如果缓冲区是从指向非const数据的主机指针创建和初始化的，就好像用那个指针调用了set_final_data。从技术上讲，原始指针是OutputIterator的特例。如果传递给set_final_data的参数是一个std::weak_ptr，如果指针已经过期或已经被删除，数据不会被写入主机。是否发生写回也可以由set_write_back方法控制。

附件

由缓冲区表示的数据不能通过缓冲区对象直接访问。相反，我们必须创建允许我们安全访问缓冲区数据的访问器对象。访问器通知运行时我们希望在哪里以及如何访问数据，允许运行时确保正确的数据在正确的时间位于正确的位置。这是一个非常强大的概念，尤其是当与部分基于数据依赖性来调度内核执行的任务图结合使用时。

访问器对象是从模板化的accessor类实例化的。这个类有五个模板参数。第一个参数是被访问数据的类型。这应该与相应缓冲区存储的数据类型相同。类似地，第二个参数描述了数据和缓冲区的维度，默认值为 1。

img/489625_1_En_7_Fig6_HTML.png

图 7-6

访问模式

接下来的三个模板参数是访问者独有的。第一个是访问模式。访问模式描述了我们打算如何在程序中使用访问器。图 7-6 中列出了可能的模式。我们将在第八章中学习如何使用这些模式来命令内核的执行和执行数据移动。如果没有指定或自动推断，则访问模式参数有默认值。如果我们没有另外指定，对于非const数据类型，访问器将默认为read_write访问模式，对于const数据类型，访问器将默认为read。这些默认值总是正确的，但是提供更准确的信息可能会提高运行时执行优化的能力。在开始应用程序开发时，简单地不指定访问模式是安全和简洁的，然后我们可以基于对应用程序的性能关键区域的分析来细化访问模式。

img/489625_1_En_7_Fig7_HTML.png

图 7-7

访问目标

下一个模板参数是访问目标。缓冲区是数据的抽象，并不描述数据存储在哪里以及如何存储。访问目标描述了我们正在访问什么类型的数据，以及哪个内存将包含这些数据。图 7-7 中列出了可能的访问目标。数据类型是两种类型之一:缓冲区或图像。本书中讨论了图像，但我们可以将它们视为专用缓冲区，为图像处理提供特定于域的操作。

访问目标的另一个方面是我们应该关注的。设备可能有不同类型的可用存储器。这些存储器由不同的地址空间表示。最常用的内存类型是设备的全局内存。内核中的大多数访问器将使用这个目标，所以 global 是默认目标(如果我们没有指定)。常量和本地缓冲区使用专用内存。顾名思义，常量内存用于存储在内核调用期间保持不变的值。本地内存是一个工作组可用的特殊内存，其他工作组无法访问。我们将在第九章中学习如何使用本地内存。另一个值得注意的目标是主机缓冲区，这是访问主机上的缓冲区时使用的目标。这个模板参数的默认值是global_buffer，所以在大多数情况下，我们不需要在代码中指定目标。

最后一个模板参数决定了一个访问器是否是一个占位符访问器。这不是一个程序员可能会直接设置的参数。占位符访问器是在命令组之外声明的，但是用于访问内核内部设备上的数据。一旦我们看了访问器创建的例子，我们将看到占位符访问器和非占位符访问器的区别。

虽然可以使用缓冲区对象的get_access方法从缓冲区对象中提取访问器，但是直接创建(构造)它们更简单。这是我们将在接下来的例子中使用的风格，因为它很容易理解，也很简洁。

访问者创建

图 7-8 显示了一个示例程序，其中包含了我们开始使用访问器所需的一切。在这个例子中，我们有三个缓冲器，A、B和C。我们提交给队列的第一个任务是为每个缓冲区创建访问器，并定义一个内核，使用这些访问器用一些值初始化缓冲区。每个访问器都是用它将访问的缓冲区的引用以及由我们提交给队列的命令组定义的处理程序对象来构造的。这有效地将访问器绑定到我们作为命令组的一部分提交的内核。常规访问器是设备访问器，因为默认情况下，它们的目标是存储在设备内存中的全局缓冲区。这是最常见的用例。

img/489625_1_En_7_Fig8_HTML.png

图 7-8

简单的访问器创建

我们提交的第二个任务也定义了三个缓冲区的访问器。然后我们在第二个内核中使用这些访问器将缓冲区A和B的元素添加到缓冲区C中。由于第二个任务与第一个任务操作相同的数据，运行时将在第一个任务完成后执行该任务。我们将在下一章详细了解这一点。

第三个任务展示了如何使用占位符访问器。在我们创建了缓冲区之后，访问器pC在图 7-8 中的例子的开头被声明。请注意，没有向构造器传递 handler 对象，因为我们没有要传递的对象。这让我们可以提前创建一个可重用的访问器对象。然而，为了在内核中使用这个访问器，我们需要在提交期间将它绑定到一个命令组。我们使用处理程序对象的require方法来完成这项工作。一旦我们将占位符访问器绑定到命令组，我们就可以像使用其他访问器一样在内核中使用它。

最后，我们创建一个host_accessor对象，以便在主机上读取我们的计算结果。请注意，这与我们在内核中使用的类型不同。主机访问器使用一个单独的host_accessor类来允许正确推断模板参数，提供一个简单的接口。请注意，本例中的主机访问器result也没有处理程序对象，因为我们也没有传递对象。主机访问器的特殊类型也让我们能够将它们与占位符区分开来。主机访问器的一个重要方面是，构造器仅在数据可供主机使用时才完成，这意味着主机访问器的构造可能需要很长时间。构造器必须等待任何产生要复制的数据的内核完成执行，以及等待复制本身完成。一旦主机访问器构造完成，就可以安全地在主机上直接使用它所访问的数据，并且我们可以保证在主机上获得最新版本的数据。

虽然这个例子是完全正确的，但是我们并没有说我们在创建访问器时打算如何使用它们。相反，我们对缓冲区中的非const int数据使用默认访问模式read-write。这可能会过度保守，并且可能会在操作之间创建不必要的依赖关系或多余的数据移动。如果一个运行时有更多关于我们计划如何使用我们创建的访问器的信息，它可能会做得更好。然而，在我们看一个这样做的例子之前，我们应该首先介绍另一个工具——访问标记。

访问标记是表达访问者所需的访问模式和目标组合的一种简洁方式。使用时，访问标记作为参数传递给访问器的构造器。可能的标签如图 7-9 所示。当用标记参数构造访问器时，C++ CTAD 可以正确地推导出所需的访问模式和目标，提供了一种简单的方法来覆盖那些模板参数的默认值。我们也可以手动指定所需的模板参数，但是标记提供了一种更简单、更紧凑的方式来获得相同的结果，而无需拼写出完全模板化的访问器。

img/489625_1_En_7_Fig9_HTML.png

图 7-9

访问标签

让我们以前面的例子为例，重写它以添加访问标记。这个新的改进示例如图 7-10 所示。

img/489625_1_En_7_Fig10_HTML.png

图 7-10

使用指定的用法创建访问者

我们首先声明我们的缓冲区，如图 7-8 所示。我们还创建了占位符访问器，我们将在后面使用。现在让我们看看提交给队列的第一个任务。以前，我们通过传递对命令组的缓冲区和处理程序对象的引用来创建我们的访问器。现在，我们向构造器调用添加两个额外的参数。第一个新参数是访问标记。因为这个内核正在为我们的缓冲区写初始值，所以我们使用了write_only访问标记。这让运行时知道这个内核正在产生新的数据，并且不会从缓冲区中读取。

第二个新参数是一个可选的访问器属性，类似于我们在本章前面看到的缓冲区的可选属性。我们传递的属性noinit让运行时知道缓冲区中以前的内容可以被丢弃。这很有用，因为它可以让运行时消除不必要的数据移动。在这个例子中，因为第一个任务是为我们的缓冲区写初始值，所以运行时没有必要在内核执行之前将未初始化的主机内存复制到设备上。noinit属性在这个例子中很有用，但是它不应该用于读-修改-写的情况，也不应该用于只能更新缓冲区中某些值的内核。

我们提交给队列的第二个任务与之前相同，但是现在我们向我们的访问器添加了访问标记。这里，我们给访问器aA和aB添加标签read_only，让运行时知道我们将只通过这些访问器读取缓冲区A和B的值。第三个访问器aC获得read_write访问标记，因为我们将A和B的元素之和累加到C中。我们在示例中显式地使用标签来保持一致，但是这是不必要的，因为默认的访问模式是read_write。

默认用法保留在第三个任务中，在这里我们使用占位符访问器。这与我们在图 7-8 中看到的简化示例保持不变。我们的最后一个访问器，主机访问器result，现在在我们创建它时会收到一个访问标记。因为我们只读取主机上的最终值，所以我们将read_only标记传递给构造器。如果我们以破坏主机访问器的方式重写程序，启动另一个在缓冲区C上运行的内核不需要将它写回设备，因为read_only标签让运行时知道它不会被主机修改。

我们可以用访问器做什么？

使用访问器对象可以完成许多事情。然而，我们能做的最重要的事情是在访问器的名字中拼写出来——访问数据。这通常是通过访问器的[]操作符来完成的。我们在图 7-8 和 7-10 的示例中使用了[]操作符。这个操作符要么接受一个可以正确索引多维数据的id对象，要么接受一个size_t。当访问者有多个维度时，使用第二种情况。它返回一个对象，然后用[]再次索引该对象，直到我们得到一个标量值，这在二维情况下将是a[i][j]的形式。请记住，访问器维度的排序遵循 C++ 的约定，其中最右边的维度是单位步长维度(迭代“最快”)。

访问器还可以返回指向基础数据的指针。这个指针可以按照正常的 C++ 规则直接访问。注意，关于这个指针的地址空间，可能涉及额外的复杂性。地址空间和它们的怪癖将在后面的章节中讨论。

许多东西也可以从访问器对象中查询。示例包括通过访问器可访问的元素数量、它所覆盖的缓冲区区域的字节大小或可访问的数据范围。

访问器为 C++ 容器提供了一个类似的接口，可以在许多容器被传递的情况下使用。访问器支持的容器接口包括data方法，相当于get_pointer，以及几种向前和向后迭代器。

摘要

在本章中，我们已经学习了缓冲区和存取器。缓冲区是对数据的抽象，它对程序员隐藏了内存管理的底层细节。他们这样做是为了提供一个更简单、更高层次的抽象。我们通过几个例子向我们展示了构造缓冲区的不同方法，以及可以被指定来改变它们的行为的不同可选属性。我们学习了如何用来自主机内存的数据初始化缓冲区，以及如何在使用完缓冲区时将数据写回主机内存。

因为我们不应该直接访问缓冲区，所以我们学习了如何使用访问器对象来访问缓冲区中的数据。我们了解了设备访问器和主机访问器之间的区别。我们讨论了不同的访问模式和目标，以及它们如何通知运行时程序将如何以及在哪里使用访问器。我们展示了使用默认访问模式和目标来使用访问器的最简单方法，并且我们学习了如何区分占位符访问器和非占位符访问器。然后，我们看到了如何通过向我们的访问器声明添加访问标记，为运行时提供更多关于我们的访问器用法的信息，从而进一步优化示例程序。最后，我们讨论了在程序中使用访问器的许多不同方式。

在下一章，我们将更详细地了解运行时如何使用我们通过访问器给它的信息来调度不同内核的执行。我们还将了解这些信息如何通知运行时缓冲区中的数据何时以及如何需要在主机和设备之间复制。我们将了解如何显式控制涉及缓冲区的数据移动——以及 USM 分配。

Creative Commons

C---数据并行教程-一-