C++ 数据并行教程（二）
八、调度内核和数据移动
我们需要讨论一下我们作为平行项目的音乐会指挥的角色。并行程序的适当编排是一件美妙的事情——代码全速运行而不等待数据，因为我们已经安排所有数据在适当的时间到达和离开。代码分解良好，以保持硬件最大限度地忙碌。这是构成梦想的东西！
快车道上的生活——不仅仅是一条车道！—要求我们认真对待指挥工作。为了做到这一点，我们可以根据任务图来考虑我们的工作。
因此，在这一章中，我们将讨论任务图，这种机制用于正确有效地运行复杂的内核序列。在应用程序中有两件事情需要排序:内核和数据移动。任务图是我们用来实现正确排序的机制。
首先，我们将快速回顾如何使用依赖关系来排序第三章中的任务。接下来，我们将介绍 DPC++ 运行时如何构建图形。我们将讨论 DPC++ 图形的基本构件，命令组。然后，我们将举例说明构建常见模式图的不同方法。我们还将讨论数据移动，无论是显式的还是隐式的，是如何在图中表示的。最后，我们将讨论使我们的图表与主机同步的各种方法。
什么是图表调度？
在第三章中，我们讨论了数据管理和数据使用的排序。那一章描述了 DPC++ 中图形背后的关键抽象:依赖性。内核之间的依赖关系基本上是基于内核访问的数据。内核在计算输出之前需要确定它读取了正确的数据。
我们描述了对确保正确执行很重要的三种类型的数据依赖。第一种是写后读(RAW ),发生在一个任务需要读取另一个任务产生的数据时。这种类型的依赖描述了两个内核之间的数据流。第二种依赖发生在一个任务需要在另一个任务读取数据后更新数据的时候。我们称这种类型的依赖为读后写(WAR)依赖。最后一种类型的数据依赖发生在两个任务试图写入相同的数据时。这就是所谓的写后写(WAW)依赖性。
数据相关性是我们用来构建图表的基础。这组依赖关系是我们表达简单的线性核链和具有数百个具有复杂依赖关系的核的大型复杂图所需要的全部。无论计算需要哪种类型的图，DPC++ 图都可以确保程序根据所表达的依赖关系正确执行。然而，确保一个图正确地表达程序中的所有依赖关系是程序员的责任。
图形如何在 DPC++ 中工作
一个命令组可以包含三种不同的东西:一个动作、它的依赖项和各种各样的主机代码。在这三件事情中，最需要的是行动，因为没有它，指挥组真的什么也做不了。大多数命令组也会表达依赖性，但也有不表达的情况。一个这样的例子是在程序中提交的第一个动作。它不依赖于任何东西来开始执行；因此，我们不会指定任何依赖关系。命令组中可能出现的另一个东西是在主机上执行的任意 C++ 代码。这是完全合法的，并且有助于指定动作或其依赖项，并且在创建命令组时执行该代码(而不是在基于已满足的依赖项执行动作时)。
命令组通常表示为传递给 submit 方法的 C++ lambda 表达式。命令组也可以通过队列对象上的快捷方式来表达，队列对象采用一个内核和一组基于事件的依赖关系。
命令组操作
命令组可以执行两种类型的操作:内核和显式内存操作。一个命令组只能执行一个动作。正如我们在前面章节中看到的，内核是通过调用parallel_for或single_task方法来定义的，并表达我们想要在设备上执行的计算。显式数据移动操作是第二种类型的操作。USM 的例子包括memcpy、memset和fill操作。缓冲器的例子包括copy、fill和update_host。
命令组如何声明依赖关系
命令组的另一个主要组成部分是在组定义的动作可以执行之前必须满足的依赖集。DPC++ 允许以多种方式指定这些依赖关系。
如果程序使用有序 DPC++ 队列，队列的有序语义指定连续排队的命令组之间的隐式依赖关系。在之前提交的任务完成之前，一个任务无法执行。
基于事件的依赖性是指定在命令组可以执行之前必须完成什么的另一种方式。这些基于事件的依赖性可以用两种方式来指定。当命令组被指定为传递给队列的submit方法的 lambda 时，使用第一种方法。在这种情况下，程序员调用命令组处理程序对象的depends_on方法，将事件或事件向量作为参数传递。当从队列对象上定义的快捷方法创建命令组时，使用另一种方法。当程序员直接调用队列上的parallel_for或single_task时，事件或事件向量可能会作为额外的参数传递。
指定依赖关系的最后一种方法是通过创建访问器对象。访问器指定如何使用它们在缓冲区对象中读取或写入数据，让运行时使用这些信息来确定不同内核之间存在的数据依赖关系。正如我们在本章开始时所回顾的，数据依赖的例子包括一个内核读取另一个内核产生的数据，两个内核写入相同的数据，或者一个内核在另一个内核读取数据后修改数据。
例子
现在，我们将用几个例子来说明我们刚刚学到的一切。我们将展示如何用几种方式表达两种不同的依赖模式。我们将说明的两种模式是线性依赖链，其中一个任务在另一个任务之后执行，以及“Y”模式，其中两个独立的任务必须在连续的任务之前执行。
这些依赖模式的图表可以在图 8-1 和 8-2 中看到。图 8-1 描绘了一个线性依赖链。第一个节点表示数据的初始化，而第二个节点表示将数据累积到单个结果中的归约操作。图 8-2 描绘了一个“Y”模式，我们独立地初始化两个不同的数据。数据初始化后，加法核将把两个向量加在一起。最后，图中的最后一个节点将结果累积成一个值。
图 8-2
“Y”型依赖图
图 8-1
线性相关链图
对于每种模式，我们将展示三种不同的实现。第一个实现将使用有序队列。第二种将使用基于事件的依赖关系。最后一个实现将使用缓冲区和存取器来表达命令组之间的数据依赖性。
图 8-3 显示了如何使用有序队列表达线性依赖链。这个例子非常简单，因为有序队列的语义已经保证了命令组之间的执行顺序。我们提交的第一个内核将数组的元素初始化为 1。然后，下一个内核获取这些元素，并将它们汇总到第一个元素中。因为我们的队列是有序的，所以我们不需要做任何其他事情来表示第二个内核应该在第一个内核完成之前不执行。最后，我们等待队列执行完所有任务，并检查我们是否获得了预期的结果。
图 8-3
具有有序队列的线性相关链
图 8-4 显示了使用无序队列和基于事件的依赖关系的相同例子。这里，我们捕获第一次调用parallel_for返回的事件。然后，第二个内核能够通过将它作为参数传递给depends_on来指定对该事件及其所代表的内核执行的依赖。我们将在图 8-6 中看到如何使用定义内核的快捷方法之一来缩短第二个内核的表达式。
图 8-4
事件线性相关链
图 8-5 使用缓冲区和存取器代替 USM 指针重写了我们的线性依赖链示例。这里我们再次使用无序队列，但是使用通过访问器指定的数据依赖关系，而不是基于事件的依赖关系来排序命令组的执行。第二个内核读取第一个内核产生的数据，运行时可以看到这一点，因为我们基于相同的底层缓冲区对象声明了访问器。与前面的例子不同，我们不等待队列执行完所有任务。相反，我们声明一个主机访问器，它定义了第二个内核的输出和我们的断言(我们在主机上计算了正确的答案)之间的数据依赖关系。请注意，虽然主机访问器为我们提供了主机上数据的最新视图，但它并不保证原始主机内存已经更新(如果在创建缓冲区时指定了任何内存)。我们不能安全地访问原始主机内存，除非缓冲区首先被破坏，或者除非我们使用更高级的机制，如第七章中描述的互斥机制。
图 8-5
具有缓冲器和附件的线性相关链
图 8-6 显示了如何使用有序队列表达一个“Y”模式。在这个例子中，我们声明了两个数组，data1和data2。然后我们定义两个内核，每个内核初始化一个数组。这些内核并不相互依赖，但是因为队列是有序的，所以内核必须一个接一个地执行。注意，在这个例子中交换这两个内核的顺序是完全合法的。在第二个内核执行之后，第三个内核将第二个数组的元素添加到第一个数组的元素中。最终的内核将第一个数组的元素相加，计算出与我们在线性依赖链的例子中相同的结果。这个求和核依赖于前面的核，但是这个线性链也被有序队列捕获。最后，我们等待所有内核完成，并验证我们成功地计算了我们的幻数。
图 8-6
具有有序队列的“Y”型模式
图 8-7 显示了我们的“Y”模式示例，使用无序队列代替有序队列。由于队列的顺序，依赖性不再是隐式的，我们必须使用事件显式地指定命令组之间的依赖性。如图 8-6 所示，我们从定义两个没有初始依赖关系的独立内核开始。我们用两个事件来表示这些内核，e1和e2。当我们定义第三个内核时，我们必须指定它依赖于前两个内核。我们这样做是因为它依赖于事件e1和e2在执行之前完成。然而，在这个例子中，我们使用一种快捷方式来指定这些依赖关系，而不是处理程序的depends_on方法。这里，我们将事件作为额外参数传递给parallel_for。因为我们想一次传递多个事件，所以我们使用接受一个std::vector事件的表单，但是幸运的是，现代 C++ 通过自动将表达式{e1, e2}转换成适当的向量，为我们简化了这个过程。
图 8-7
事件的“Y”型模式
在我们最后的例子中，如图 8-8 所示，我们再次用缓冲区和访问器替换 USM 指针和事件。这个例子将两个数组data1和data2表示为缓冲对象。我们的内核不再使用快捷方式来定义内核，因为我们必须将访问器与命令组处理程序相关联。同样，第三个内核必须捕获对前两个内核的依赖。在这里，这是通过为我们的缓冲区声明访问器来实现的。因为我们之前已经为这些缓冲区声明了访问器，所以运行时能够正确地排序这些内核的执行。此外，当我们声明访问器b时，我们还在这里向运行时提供额外的信息。我们添加了访问标签read_only来让运行时知道我们只是要读取这些数据，而不是产生新的值。正如我们在线性依赖链的缓冲区和存取器示例中看到的，我们的最终内核通过更新第三个内核中产生的值来进行自我排序。我们通过声明一个主机访问器来检索我们计算的最终值，该主机访问器将等待最终内核完成执行，然后将数据移回主机，在那里我们可以读取数据并断言我们计算了正确的结果。
图 8-8
带存取器的“Y”型模式
CG 的各个部分是什么时候执行的？
因为任务图是异步的，所以想知道命令组何时被执行是有意义的。到目前为止，应该很清楚，一旦满足了内核的依赖性，就可以执行内核，但是命令组的主机部分会发生什么情况呢？
当一个命令组被提交到一个队列时，它会立即在主机上执行(在submit调用返回之前)。命令组的主机部分只执行一次。命令组中定义的任何内核或显式数据操作都将在设备上排队等待执行。
数据传送
数据移动是 DPC++ 中图形的另一个非常重要的方面，对于理解应用程序性能至关重要。但是，如果数据移动是在程序中隐式发生的，无论是使用缓冲区和访问器还是使用 USM 共享分配，这一点经常会被意外忽略。接下来，我们将研究在 DPC++ 中数据移动影响图形执行的不同方式。
明确的
显式数据移动的优点是它在图中显式地出现，让程序员清楚地看到图的执行过程。我们将把显式数据操作分为 USM 操作和缓冲区操作。
正如我们在第六章中了解到的，当我们需要在设备分配和主机之间拷贝数据时，USM 中会发生显式数据移动。这是通过在队列和处理程序类中都可以找到的memcpy方法来完成的。提交操作或命令组会返回一个事件，该事件可用于与其他命令组一起订购副本。
通过调用命令组处理程序对象的copy或update_host方法，使用缓冲区进行显式数据移动。copy方法可用于在主机内存和设备上的访问器对象之间手动交换数据。这样做有多种原因。一个简单的例子是对长时间运行的计算序列进行检查点操作。使用拷贝方法，数据可以以单向方式从设备写入任意主机内存。如果这是使用缓冲区完成的，大多数情况下(即缓冲区不是用use_host_ptr创建的)需要先将数据复制到主机，然后从缓冲区的存储器复制到所需的主机存储器。
update_host方法是copy的一种非常特殊的形式。如果在主机指针周围创建了缓冲区，此方法会将访问器表示的数据复制回原始主机内存。如果一个程序用一个用特殊的use_mutex属性创建的缓冲区手动同步主机数据，这可能是有用的。然而，这种用例不太可能在大多数程序中出现。
隐形的
隐式数据移动可能会对 DPC++ 中的命令组和任务图产生隐藏的后果。通过隐式数据移动，数据通过 DPC++ 运行时或硬件和软件的某种组合在主机和设备之间复制。在任一情况下，复制都是在没有用户明确输入的情况下进行的。让我们再次分别看一下 USM 和 buffer 案例。
使用 USM，隐式数据移动随着host和shared分配而发生。正如我们在第六章中了解到的，host分配并不真正移动数据，而是远程访问数据，shared分配可能会在主机和设备之间迁移。因为这种迁移是自动发生的，所以 USM 隐式数据移动和命令组真的没有什么可考虑的。然而，关于shared的分配有一些细微差别值得记住。
prefetch操作的工作方式与memcpy相似，目的是让运行时在内核尝试使用共享分配之前开始迁移它们。然而，与为了确保正确结果而必须复制数据的memcpy不同，预取通常被视为对运行时的提示以提高性能，并且预取不会使内存中的指针值无效(就像复制到新的地址范围时的复制一样)。如果在内核开始执行之前预取没有完成，程序仍将正确执行，并且许多代码可能选择使图形中的命令组不依赖于预取操作，因为它们不是功能需求。
缓冲区也有一些细微差别。使用缓冲区时，命令组必须为缓冲区构造指定如何使用数据的访问器。这些数据依赖关系表达了不同命令组之间的顺序，并允许我们构建任务图。然而，带有缓冲区的命令组有时还有另一个用途:它们指定数据移动的要求。
访问器指定内核将读取或写入缓冲区。由此得出的推论是，数据也必须在设备上可用，如果不可用，运行时必须在内核开始执行之前将数据转移到设备上。因此，DPC++ 运行时必须跟踪缓冲区的当前版本，以便可以调度数据移动操作。访问器创建有效地在图中创建了一个额外的隐藏节点。如果数据移动是必要的，运行时必须首先执行它。只有这样，提交的内核才能执行。
让我们再看看图 8-8 。在这个例子中，我们的前两个内核需要将缓冲区data1和data2复制到设备中；运行时隐式创建额外的图形节点来执行数据移动。当提交第三个内核的命令组时，这些缓冲区很可能仍然在设备上，因此运行时不需要执行任何额外的数据移动。第四个内核的数据也可能不需要任何额外的数据移动，但是主机访问器的创建需要运行时在访问器可用之前安排将缓冲区data1移回主机。
与主机同步
我们将讨论的最后一个主题是如何与主机同步图形执行。我们已经在这一章中谈到了这一点，但是我们现在将检查一个程序可以做到这一点的所有不同方式。
主机同步的第一种方法是我们在前面的许多例子中使用过的:等待一个队列。队列对象有两个方法，wait和wait_and_throw，它们阻塞执行，直到提交给队列的每个命令组都完成为止。这是一个非常简单的方法，可以处理许多常见的情况。但是，值得指出的是，这种方法是非常粗粒度的。如果需要更细粒度的同步，我们将讨论的另一种方法可能更适合应用程序的需求。
主机同步的下一种方法是对事件进行同步。这比同步队列更加灵活，因为它允许应用程序只同步特定的操作或命令组。这是通过调用事件上的wait方法或者调用事件类上的静态方法wait来完成的，后者可以接受事件的向量。
我们已经看到了图 8-5 和 8-8 中使用的下一个方法:主机访问器。主机访问者执行两个功能。首先，顾名思义，它们使主机上的数据可供访问。第二，它们通过在当前访问的图和主机之间定义新的依赖关系来与主机同步。这确保了复制回主机的数据是图形正在执行的计算的正确值。但是，我们再次注意到，如果缓冲区是从现有的主机内存中构造的，则不能保证这个原始内存包含更新后的值。
请注意，主机访问者正在阻塞。在数据可用之前，主机上的执行可能不会超过主机访问器的创建。同样，当主机访问器存在并保持其数据可用时，不能在设备上使用缓冲区。一种常见的模式是在附加的 C++ 范围内创建主机访问器，以便在不再需要主机访问器时释放数据。这是下一种主机同步方法的示例。
DPC++ 中的某些对象在被销毁时有特殊的行为，它们的析构函数被调用。我们刚刚了解了主机访问者如何使数据保留在主机上，直到它们被销毁。缓冲区和图像在被销毁或离开作用域时也有特殊的行为。当一个缓冲区被销毁时，它会等待所有使用该缓冲区的命令组完成执行。一旦缓冲区不再被任何内核或内存操作使用，运行时可能必须将数据复制回主机。如果缓冲区是用主机指针初始化的，或者如果主机指针被传递给方法set_final_data，就会发生这种复制。然后，运行库将复制回该缓冲区的数据，并在对象被销毁之前更新主机指针。
与主机同步的最后一个选项涉及一个在第七章中首次描述的不常见功能。回想一下，缓冲区对象的构造器可以选择接受一个属性列表。创建缓冲区时可以传递的有效属性之一是use_mutex。当以这种方式创建缓冲区时，它增加了缓冲区所拥有的内存可以与宿主应用程序共享的要求。对这个内存的访问是由用来初始化缓冲区的互斥体控制的。当访问与缓冲区共享的内存是安全的时，主机能够获得互斥锁。如果无法获得锁，用户可能需要将内存移动操作排入队列，以便与主机同步数据。这种用法非常特殊，不太可能在大多数 DPC++ 应用程序中找到。
摘要
在这一章中，我们已经学习了图形以及在 DPC++ 中如何构建、调度和执行图形。我们详细介绍了什么是命令组以及它们的功能。我们讨论了命令组中可能包含的三样东西:依赖性、动作和各种主机代码。我们回顾了如何使用事件以及通过访问器描述的数据依赖性来指定任务之间的依赖性。我们了解到命令组中的单个操作可以是内核操作，也可以是显式内存操作，然后我们看了几个例子，这些例子展示了我们可以构建通用执行图模式的不同方式。接下来，我们回顾了数据移动是 DPC++ 图形的一个重要部分，并且我们了解了它是如何在图形中显式或隐式出现的。最后，我们研究了所有将图形的执行与主机同步的方法。
理解程序流可以使我们理解如果我们有运行时故障要调试时可以打印的那种调试信息。第十三章在“调试运行时故障”一节中有一个表格，考虑到我们在书中学到的知识，这个表格会更有意义一些。然而，本书并不试图详细讨论这些高级编译器转储。
希望这让您感觉自己像一个图形专家，能够构建复杂的图形，从线性链到具有数百个节点和复杂数据和任务依赖关系的巨大图形！在下一章中，我们将开始深入到对提高特定设备上的应用程序的性能有用的底层细节。
开放存取本章根据知识共享署名 4.0 国际许可证(http://Creative Commons . org/licenses/by/4.0/)的条款获得许可，该许可证允许以任何媒体或格式使用、共享、改编、分发和复制，只要您适当注明原作者和来源，提供知识共享许可证的链接并指明是否进行了更改。
本章中的图像或其他第三方材料包含在本章的知识共享许可中，除非在材料的信用额度中另有说明。如果材料不包括在本章的知识共享许可中，并且您的预期使用不被法律法规允许或超出了允许的使用范围，您将需要直接从版权所有者处获得许可。
九、通信和同步
在第四章中，我们讨论了表达并行性的方法，要么使用基本的数据并行内核，显式 ND-range 内核，要么使用分层并行内核。我们讨论了基本的数据并行内核如何独立地对每一块数据应用相同的操作。我们还讨论了显式 ND-range 内核和分层并行内核如何将执行范围划分为工作项目的工作组。
在这一章中，我们将在继续寻求平行思考的过程中，重新审视如何将问题分解成小块的问题。本章提供了关于显式 ND-range 内核和分层并行内核的更多细节，并描述了如何使用工作项分组来提高某些类型算法的性能。我们将描述工作项组如何为并行工作的执行提供额外的保证，并且我们将介绍支持工作项组的语言特性。在第 15 、 16 和 17 章中优化特定设备的程序时，以及在第十四章中描述常见的并行模式时，这些想法和概念中的许多都很重要。
工作组和工作项
回想一下第四章中的内容，显式 ND-range 和分层并行内核将工作项组织成工作组，并且工作组中的工作项保证同时执行。这个属性很重要，因为当工作项被保证并发执行时，一个工作组中的工作项可以合作解决一个问题。
图 9-1 显示了一个分为多个工作组的 ND 范围，每个工作组用不同的颜色表示。每个工作组中的工作项保证并发执行，因此一个工作项可以与共享相同颜色的其他工作项进行通信。
图 9-1
二维 ND-大小范围(8，8)分为四个大小工作组(4，4)
因为不同工作组中的工作项目不能保证同时执行，所以具有一种颜色的工作项目不能与具有不同颜色的工作项目可靠地通信，并且如果一个工作项目试图与当前没有执行的另一个工作项目通信，则内核可能会死锁。因为我们希望我们的内核完成执行，我们必须确保当一个工作项目与另一个工作项目通信时，它们在同一个工作组中。
高效沟通的构建模块
本节描述支持组中工作项之间高效通信的构建块。一些是基本的构建模块，支持定制算法的构建，而另一些是更高级的，描述许多内核使用的通用操作。
通过屏障实现同步
沟通最基本的构件是屏障功能。屏障功能有两个主要目的:
首先，barrier 函数同步组中工作项的执行。通过同步执行，一个工作项可以确保另一个工作项在使用该操作的结果之前已经完成了该操作。或者，在另一个工作项使用操作结果之前，给一个工作项时间来完成其操作。
第二，barrier 函数同步每个工作项如何看待内存的状态。这种类型的同步操作被称为强制内存一致性或防护内存(更多细节在第十九章)。存储器一致性至少与同步执行一样重要，因为它确保了在屏障之前执行的存储器操作的结果对于屏障之后的其他工作项目是可见的。没有内存一致性，一个工作项中的操作就像森林中倒下的一棵树，声音可能被其他工作项听到，也可能听不到！
图 9-2 显示了一个组中的四个工作项，它们在一个障碍函数中同步。尽管每个工作项的执行时间可能不同，但是直到所有工作项都执行了屏障，才可以执行越过屏障的工作项。在执行屏障函数之后，所有的工作项都有一个一致的内存视图。
图 9-2
一个组中的四个工作项在屏障函数处同步
WHY ISN’T MEMORY CONSISTENT BY DEFAULT?
对于许多程序员来说，内存一致性的想法——以及不同的工作项可以有不同的内存视图——可能感觉非常奇怪。如果默认情况下所有工作项的内存都是一致的，不是更容易吗？简而言之，答案是肯定的，但实施起来也会非常昂贵。通过允许工作项目具有不一致的存储器视图，并且在程序执行期间仅要求在定义的点处的存储器一致性，加速器硬件可能更便宜，可能执行得更好，或者两者兼而有之。
因为屏障函数同步执行，所以要么组中的所有工作项目都执行屏障，要么组中没有工作项目执行屏障，这一点至关重要。如果组中的一些工作项绕过任何障碍函数，组中的其他工作项可能会永远等待障碍——或者至少直到用户放弃并终止程序！
COLLECTIVE FUNCTIONS
当一个功能需要由一个组中的所有工作项目执行时，它可以被称为一个集合功能，因为该操作是由该组执行的，而不是由该组中的单个工作项目执行的。屏障函数不是 SYCL 中唯一可用的集合函数。其他集合函数将在本章后面介绍。
工作组本地存储器
工作组屏障功能足以协调工作组中工作项目之间的通信，但是通信本身必须通过记忆发生。通信可以通过 USM 或缓冲区进行，但这可能不方便且效率低下:它需要专用于通信的分配，并且需要在工作组之间划分分配。
为了简化内核开发并加速工作组中工作项之间的通信，SYCL 定义了一个特殊的本地内存空间，专门用于工作组中工作项之间的通信。
在图 9-3 中，显示了两个工作组。两个工作组都可以访问全局内存空间中的 USM 和缓冲区。每个工作组可以访问自己的本地内存空间中的变量，但不能访问另一个工作组的本地内存中的变量。
图 9-3
每个工作组可以访问所有全局内存，但只能访问自己的本地内存
当一个工作组开始时，它的本地内存的内容是未初始化的，并且在一个工作组完成执行后，本地内存不再存在。由于这些特性，当一个工作组正在执行时，本地存储器只能用于临时存储。
对于一些设备，例如对于许多 CPU 设备，本地存储器是软件抽象，并且使用与全局存储器相同的存储器子系统来实现。在这些设备上，使用本地内存主要是一种方便的通信机制。一些编译器可以使用内存空间信息进行编译器优化，但是在其他方面，使用本地内存进行通信并不会比通过这些设备上的全局内存进行通信的性能更好。
但是对于其他设备，如许多 GPU 设备，本地内存有专用资源，在这些设备上，通过本地内存进行通信将比通过全局内存进行通信性能更好。
当使用本地内存时，一个工作组中的工作项之间的通信会更加方便和快捷！
我们可以使用设备查询info::device: :local_mem_type来确定加速器是否有专用于本地存储器的资源，或者本地存储器是否被实现为全局存储器的软件抽象。有关查询设备属性的更多信息，请参考第十二章；有关如何为 CPU、GPU 和 FPGAs 实现本地存储器的更多信息，请参考第 15 、 16 和 17 章。
使用工作组障碍和本地记忆
既然我们已经确定了工作项之间有效通信的基本构件，我们可以描述如何在内核中表达工作组障碍和本地内存。请记住，工作项之间的通信需要工作项分组的概念，因此这些概念只能针对 ND 范围内核和分层内核来表达，而不包括在基本数据并行内核的执行模型中。
本章将在第四章介绍的简单矩阵乘法核心示例的基础上，介绍执行矩阵乘法的工作组中工作项目之间的通信。在许多设备上——但不一定是全部！—通过本地内存进行通信将提高矩阵乘法内核的性能。
A NOTE ABOUT MATRIX MULTIPLICATION
在本书中，矩阵乘法内核用于演示内核的变化如何影响性能。虽然使用本章介绍的技术可以提高某些设备的矩阵乘法性能，但矩阵乘法是一种非常重要和常见的运算，许多供应商已经实现了高度优化的矩阵乘法版本。厂商投入大量的时间和精力来实现和验证特定设备的功能，并且在某些情况下可能使用在标准并行内核中难以或不可能使用的功能或技术。
USE VENDOR-PROVIDED LIBRARIES!
当供应商提供一个函数的库实现时，使用它比将函数重新实现为并行内核更有益！对于矩阵乘法，人们可以将 oneMKL 作为英特尔 oneAPI 工具包的一部分，来寻找适合 DPC++ 程序员的解决方案。
图 9-4 显示了我们将要开始的朴素的矩阵乘法内核，摘自第四章。
图 9-4
第四章中的简单矩阵乘法内核
在第四章中，我们观察到矩阵乘法算法具有高度的重用性，并且对工作项进行分组可以提高访问的局部性，从而提高缓存命中率。在这一章中，我们没有依靠隐式缓存行为来提高性能，而是使用本地内存作为显式缓存，以保证访问的局部性。
对于许多算法来说，将本地内存视为显式缓存是有帮助的。
图 9-5 是第四章的修改图，显示了一个由单行组成的工作组，这使得使用本地存储器的算法更容易理解。注意，对于结果矩阵的一行中的元素，每个结果元素都是使用来自输入矩阵之一的唯一数据列计算的，以蓝色和橙色显示。因为这个输入矩阵没有数据共享，所以它不是本地内存的理想选择。但是，请注意，该行中的每个结果元素都访问另一个输入矩阵中完全相同的数据，以绿色显示。因为这些数据是重用的，所以它是受益于工作组本地内存的绝佳候选对象。
图 9-5
矩阵乘法到工作组和工作项的映射
因为我们想要乘可能非常大的矩阵，并且因为工作组本地存储器可能是有限的资源，所以我们修改的内核将处理每个矩阵的子部分，我们将这些子部分称为矩阵块。对于每个图块，我们修改后的内核会将图块的数据加载到本地内存中，同步组中的工作项，然后从本地内存而不是全局内存中加载数据。第一个图块的访问数据如图 9-6 所示。
图 9-6
处理第一个图块:绿色输入数据(X 的左侧)被重用并从本地内存中读取，蓝色和橙色输入数据(X 的右侧)从全局内存中读取
在我们的内核中，我们选择了与工作组大小相等的瓦片大小。这不是必需的，但是因为它简化了进出本地存储器的传输，所以选择工作组大小的倍数的切片大小是常见且方便的。
ND-Range 核中的工作组障碍和局部记忆
本节描述了工作组障碍和局部记忆是如何在 ND-range 核中表示的。对于 ND-range 内核，表示是显式的:内核声明并操作表示本地地址空间中的分配的本地存取器，并调用屏障函数来同步工作组中的工作项目。
本地访问者
要声明在 ND-range 内核中使用的本地内存，使用一个本地访问器。像其他访问器对象一样，本地访问器是在命令组处理程序中构造的，但是与第 3 和 7 章中讨论的访问器对象不同，本地访问器不是从缓冲区对象创建的。相反，通过指定类型和描述该类型元素数量的范围来创建局部访问器。像其他访问器一样，局部访问器可以是一维、二维或三维的。图 9-7 展示了如何声明本地访问器并在内核中使用它们。
图 9-7
声明和使用本地访问器
请记住，当每个工作组开始时，本地内存是未初始化的，并且在每个工作组完成后不会持续存在。这意味着本地访问器必须总是read_write，因为否则内核将无法分配本地内存的内容或查看分配的结果。但是，本地访问器也可以是原子的，在这种情况下，通过访问器对本地存储器的访问是原子地执行的。原子访问将在第十九章中详细讨论。
同步功能
为了同步 ND-range 内核工作组中的工作项，调用nd_item类中的barrier函数。因为屏障函数是nd_item类的成员，所以它只对 ND-range 内核可用，对基本数据并行内核或分层内核不可用。
barrier 函数目前接受一个参数来描述要同步的内存空间或 fence ，但是随着内存模型在 SYCL 和 DPC++ 中的发展，barrier 函数的参数将来可能会改变。然而，在所有情况下，屏障函数的参数提供了关于同步的内存空间或内存同步的范围的额外控制。
当没有参数传递给屏障函数时，屏障函数将使用功能正确且保守的默认值。本章中的代码示例使用这种语法以获得最大的可移植性和可读性。对于高度优化的内核，建议精确描述哪些内存空间或哪些工作项必须同步，这样可以提高性能。
一个完整的 ND 范围内核示例
现在我们知道了如何声明一个本地内存访问器，并使用屏障函数同步对它的访问，我们可以实现一个 ND-range 内核版本的矩阵乘法，它协调工作组中工作项之间的通信，以减少全局内存的流量。完整的示例如图 9-8 所示。
图 9-8
用 ND-range parallel_for和工作组本地存储器表示平铺矩阵乘法内核
这个内核中的主循环可以被认为是两个不同的阶段:在第一阶段，工作组中的工作项协作将共享数据从 A 矩阵加载到工作组本地内存中；在第二种情况下，工作项使用共享数据执行自己的计算。为了确保所有的工作项在进入第二阶段之前已经完成了第一阶段，这两个阶段通过调用barrier来同步所有的工作项并提供一个内存栅栏来分开。这种模式很常见，在内核中使用工作组本地内存几乎总是需要使用工作组屏障。
注意，还必须调用barrier来同步当前图块的计算阶段和下一个矩阵图块的加载阶段之间的执行。如果没有这种同步操作，当前矩阵片的一部分可能会在另一个工作项完成计算之前被工作组中的一个工作项覆盖。一般来说，每当一个工作项在本地内存中读取或写入由另一个工作项读取或写入的数据时，就需要同步。在图 9-8 中，同步是在循环结束时进行的，但是在每次循环迭代开始时进行同步也同样正确。
等级核中的工作组障碍和局部记忆
本节描述了如何在分层内核中表达工作组障碍和本地记忆。与 ND-range 内核不同，分层内核中的本地内存和屏障是隐式的，不需要特殊的语法或函数调用。一些程序员会发现分层内核表示更加直观和易于使用，而其他程序员会喜欢 ND-range 内核提供的直接控制。在大多数情况下，可以使用两种表示来描述相同的算法，因此我们可以选择我们认为最容易开发和维护的表示。
本地内存和屏障的范围
回想一下第四章中的，分层内核通过使用parallel_for_work_group和parallel_for_work_item函数表达了两个级别的并行执行。并行执行的这两个级别或范围用于表示变量是否在工作组本地存储器中并且在工作组中的所有工作项之间共享，或者变量是否在每个工作项的私有存储器中，该私有存储器不在工作项之间共享。这两个作用域还用于同步一个工作组中的工作项，并加强内存一致性。
图 9-9 显示了一个示例层次内核，它在本地内存的工作组范围内声明一个变量，加载到其中，然后在工作项范围内使用该变量。在工作组范围内写入本地内存和在工作项范围内从本地内存读取之间存在一个隐含的障碍。
图 9-9
具有本地存储器变量的分层内核
分层内核表示的主要优点是它看起来非常类似于标准的 C++ 代码，其中一些变量可能在一个作用域中赋值，而在一个嵌套的作用域中使用。当然，这也可能被认为是一个缺点，因为它并不直接清楚哪些变量在本地存储器中，以及何时必须由分层内核编译器插入屏障。对于屏障昂贵的设备来说尤其如此！
一个完整的分层内核示例
现在我们知道了如何在分层内核中表达本地内存和屏障，我们可以编写一个分层内核，实现与图 9-7 中 ND-range 内核相同的算法。该内核如图 9-10 所示。
图 9-10
作为分层内核实现的平铺矩阵乘法内核
虽然分层内核与 ND-range 内核非常相似，但有一个关键的区别:在 ND-range 内核中，矩阵乘法的结果在写入内存中的输出矩阵之前被累积到每个工作项变量sum中，而分层内核则累积到内存中。我们也可以在分层内核中累加到每个工作项的变量中，但是这需要一个特殊的private_memory语法来在工作组范围内声明每个工作项的数据，我们选择使用分层内核语法的原因之一是为了避免特殊语法！
分层内核不需要特殊的语法来声明工作组本地内存中的变量，但是它们需要特殊的语法来声明工作项私有内存中的一些变量！
为了避免特殊的每工作项数据语法，分层内核中工作项循环的常见模式是将中间结果写入工作组本地内存或全局内存。
图 9-10 中内核最后一个有趣的属性与循环迭代变量kk有关:由于循环是在工作组范围内，循环迭代变量kk可以在工作组本地内存之外分配，就像tileA数组一样。不过在这种情况下，由于kk的值对于工作组中的所有工作项都是相同的，所以智能编译器可能会选择从每个工作项的内存中分配kk,特别是对于工作组本地内存是稀缺资源的设备。
子群体
到目前为止，根据内核的编写方式，通过工作组本地内存交换数据，以及通过隐式或显式屏障函数进行同步，工作项已经与工作组中的其他工作项进行了通信。
在第四章中，我们讨论了另一组工作项目。子组是工作组中工作项目的实现定义的子集，它们在相同的硬件资源上一起执行或者具有额外的调度保证。因为实现决定了如何将工作项分组为子组，所以子组中的工作项可能能够比任意工作组中的工作项更有效地进行通信或同步。
本节描述了子组中工作项之间通信的构建块。注意，子组目前仅针对 ND-range 内核实现，并且子组不能通过分层内核来表达。
通过子群障碍的同步
就像 ND-range 内核中的工作组中的工作项目可以如何使用工作组屏障函数来同步一样，子组中的工作项目可以使用子组屏障函数来同步。工作组中的工作项通过调用nd_item类中的group_barrier函数或barrier函数进行同步，子组中的工作项通过调用特殊sub_group类中的group_barrier函数或barrier函数进行同步，该类可从nd_item类中查询，如图 9-11 所示。
图 9-11
查询和使用sub_group类
与工作组屏障一样，子组屏障可以接受可选参数，以更精确地控制屏障操作。不管子组屏障功能是同步全局存储器还是本地存储器，仅同步子组中的工作项可能比同步工作组中的所有工作项更便宜。
在子组内交换数据
与工作组不同，子组没有用于交换数据的专用内存空间。相反，子组中的工作项可以通过工作组本地内存、全局内存或者更常见的通过使用子组集合函数来交换数据。
如前所述，集合函数是描述由一组工作项目而不是单个工作项目执行的操作的函数，并且因为屏障同步函数是由一组工作项目执行的操作，所以它是集合函数的一个例子。
其他集合函数表示常见的通信模式。我们将在本章后面详细描述许多集合函数的语义，但是现在，我们将简要描述我们将使用子组实现矩阵乘法的broadcast集合函数。
broadcast集合函数从组中的一个工作项中获取一个值，并将其传递给组中的所有其他工作项。示例如图 9-12 所示。注意，broadcast 函数的语义要求标识组中哪个值要通信的local_id对于组中的所有工作项必须是相同的，确保 broadcast 函数的结果对于组中的所有工作项也是相同的。
图 9-12
由broadcast功能处理
如果我们查看本地内存矩阵乘法内核的最内层循环，如图 9-13 所示，我们可以看到对矩阵块的访问是一种广播，因为组中的每个工作项从矩阵块中读取相同的值。
图 9-13
矩阵乘法内核包括一个广播操作
我们将使用子组广播函数来实现一个矩阵乘法内核，它不需要工作组本地内存或屏障。在许多设备上，子组广播比带有工作组本地内存和障碍的广播更快。
一个完整的子群 ND-Range 核示例
图 9-14 是一个使用子群实现矩阵乘法的完整例子。请注意，这个内核不需要工作组本地内存或显式同步，而是使用子组广播集合函数在工作项之间传递矩阵平铺的内容。
图 9-14
用 ND-range parallel_for和子群集合函数表示的平铺矩阵乘法核
集体职能
在本章的“子组”一节中，我们描述了集体函数以及集体函数如何表达常见的通信模式。我们特别讨论了 broadcast collective 函数，它用于将一个组中的一个工作项的值传递给组中的其他工作项。本节描述附加的集合函数。
虽然本节中描述的集合功能可以使用诸如原子、工作组本地存储器和屏障之类的特性直接在我们的程序中实现，但是许多设备都包括专用硬件来加速集合功能。即使设备不包含专用硬件，供应商提供的集合函数的实现也可能针对运行它们的设备进行了调整，因此调用内置的集合函数通常会比我们编写的通用实现执行得更好。
使用通用通信模式的集合函数来简化代码和提高性能！
工作组和子组都支持许多集合功能。其他集合功能仅支持子组。
广播
broadcast函数允许一个组中的一个工作项与该组中的所有其他工作项共享一个变量的值。图 9-12 中显示了广播功能的工作原理。工作组和子组都支持broadcast功能。
投票
any_of和all_of函数(以下统称为“投票”函数)使工作项能够比较其组中布尔条件的结果:any_of如果组中至少一个工作项的条件为真，则返回真，只有当组中所有工作项的条件为真时，all_of才返回真。图 9-15 显示了这两个功能的比较。
图 9-15
any_of和all_of功能的比较
工作组和子组都支持any_of和all_of投票功能。
洗牌
子组最有用的特性之一是能够在单个工作项之间直接通信，而不需要显式的内存操作。在许多情况下，例如子组矩阵乘法内核，这些混洗操作使我们能够从内核中移除工作组本地内存使用和/或避免对全局内存的不必要的重复访问。有几种风格的随机播放功能可用。
最通用的混洗功能称为shuffle，如图 9-16 所示，它允许子组中任何一对工作项之间的任意通信。然而，这种通用性可能是以性能为代价的，我们强烈鼓励尽可能使用更专业的随机播放函数。
图 9-16
基于预先计算的置换索引，使用通用的shuffle对x值进行排序
在图 9-16 中，使用预先计算的排列索引，使用通用混洗来对子组的x值进行排序。对于子组中的一个工作项目显示了箭头，其中混洗的结果是工作项目的 x 值，其中local_id等于 7。
注意，子组broadcast函数可以被认为是通用shuffle函数的特殊版本，其中混洗索引对于子组中的所有工作项都是相同的。当混洗索引对于子组中的所有工作项都是相同的时，使用broadcast而不是shuffle为编译器提供了额外的信息，并且可以提高某些实现的性能。
shuffle_up和shuffle_down功能有效将子组的内容向给定方向移动固定数量的元素，如图 9-17 所示。注意，返回到子组中最后五个工作项的值是未定义的，在图 9-17 中显示为空白。移位对于并行化具有循环相关性的循环或实现通用算法(如互斥或包含扫描)非常有用。
图 9-17
使用shuffle_down将子组的x值移动五项
shuffle_xor函数交换两个工作项的值，这由应用于工作项的子组本地 id 和固定常量的 XOR 运算的结果指定。如图 9-18 和 9-19 所示，几种常见的通信模式可以用异或来表示:例如，交换相邻值对
图 9-19
使用shuffle_xor反转x的值
图 9-18
使用shuffle_xor交换相邻的x对
或者反转子组值。
SUB-GROUP OPTIMIZATIONS USING BROADCAST, VOTE, AND COLLECTIVES
应用于子组的 broadcast、vote 和其他集合函数的行为与应用于工作组时是相同的，但它们值得额外关注，因为它们可能会在某些编译器中实现激进的优化。例如，编译器可能能够减少向子组中的所有工作项广播的变量的寄存器使用，或者可能能够基于any_of和all_of函数的使用来推断控制流分歧。
装载和存储
子组加载和存储功能有两个目的:第一，通知编译器子组中的所有工作项正在加载从内存中相同(统一)位置开始的连续数据，第二，使我们能够请求大量连续数据的优化加载/存储。
对于 ND-range parallel_for，编译器可能不清楚不同工作项计算的地址如何相互关联。例如，如图 9-20 所示，从索引[0，32]访问一个连续的内存块，从每个工作项的角度来看，似乎有一个跨步的访问模式。
图 9-20
访问四个连续块的子组的存储器访问模式
一些体系结构包括专用硬件来检测子组中的工作项目何时访问连续数据并组合它们的存储器请求，而其他体系结构要求提前知道这一点并将其编码在加载/存储指令中。子组加载和存储在任何平台上都不是正确性所必需的，但在某些平台上可能会提高性能，因此应被视为一种优化提示。
摘要
本章讨论了一个组中的工作项如何通信和协作来提高某些类型内核的性能。
我们首先讨论了 ND-range 内核和分层内核是如何支持将工作项分组到工作组中的。我们讨论了将工作项分组到工作组中是如何改变并行执行模型的，从而保证工作组中的工作项并发执行，并支持通信和同步。
接下来，我们讨论了一个工作组中的工作项如何使用屏障进行同步，以及屏障如何针对 ND-range 内核进行显式表达，或者针对分层内核在工作组和工作项范围之间进行隐式表达。我们还讨论了如何通过工作组本地内存执行工作组中工作项之间的通信，以简化内核并提高性能，我们还讨论了如何使用 ND-range 内核的本地访问器或分层内核的工作组范围内的分配来表示工作组本地内存。
我们讨论了 ND-range 内核中的工作组如何进一步划分为工作项目的子组，其中工作项目的子组可以支持额外的通信模式或调度保证。
对于工作组和子组，我们讨论了如何通过使用集体功能来表达和加速常见的交流模式。
本章中的概念是理解第十四章中描述的常见并行模式以及理解如何针对第 15 、 16 和 17 章中的特定器件进行优化的重要基础。
开放存取本章根据知识共享署名 4.0 国际许可证(http://Creative Commons . org/licenses/by/4.0/)的条款获得许可，该许可证允许以任何媒体或格式使用、共享、改编、分发和复制，只要您适当注明原作者和来源，提供知识共享许可证的链接并指明是否进行了更改。
本章中的图像或其他第三方材料包含在本章的知识共享许可中，除非在材料的信用额度中另有说明。如果材料不包括在本章的知识共享许可中，并且您的预期使用不被法律法规允许或超出了允许的使用范围，您将需要直接从版权所有者处获得许可。
十、定义内核
到目前为止，在本书中，我们的代码示例已经使用 C++ lambda 表达式表示了内核。Lambda 表达式是一种简洁而方便的方法，可以在使用它的地方表示内核，但它不是 SYCL 中表示内核的唯一方法。在这一章中，我们将详细探索定义内核的各种方法，帮助我们选择最适合我们 C++ 编码需求的内核形式。
本章解释并比较了表示内核的三种方式:
-
λ表达式
-
命名函数对象(仿函数)
-
与通过其他语言或 API 创建的内核的互操作性
本章最后讨论了如何在一个程序对象中显式地操作内核来控制内核何时以及如何被编译。
为什么用三种方式来表示一个内核？
在深入细节之前，让我们先总结一下为什么有三种定义内核的方法，以及每种方法的优缺点。图 10-1 给出了一个有用的总结。
请记住，内核是用来表示一个计算单元的，许多内核实例通常会在一个加速器上并行执行。SYCL 支持多种方式来表达内核，以自然、无缝地集成到各种代码库中，同时在各种加速器类型上高效执行。
图 10-1
表示内核的三种方式
作为 Lambda 表达式的内核
C++ lambda 表达式，也称为匿名函数对象、未命名函数对象、闭包，或者简称为 lambdas ，是一种在使用内核时表达内核权利的便捷方式。本节描述如何将内核表示为 C++ lambda 表达式。这扩展了第一章中关于 C++ lambda 函数的介绍性复习，其中包括一些带有输出的编码示例。
C++ lambda 表达式非常强大，并且具有表达性语法，但是在表达内核时，只需要(并且支持)完整 C++ lambda 表达式语法的特定子集。
图 10-2
使用 lambda 表达式定义的内核
内核 Lambda 表达式的元素
图 10-2 显示了一个以典型的 lambda 表达式编写的内核——本书中的代码示例已经使用了这种语法。
图 10-3 中的插图显示了更多可用于内核的 lambda 表达式的元素，但这些元素中有许多并不典型。在大多数情况下，lambda 缺省值就足够了，所以一个典型的内核 lambda 表达式看起来更像图 10-2 中的 lambda 表达式，而不是图 10-3 中更复杂的 lambda 表达式。
图 10-3
内核 lambda 表达式的更多元素，包括可选元素
-
lambda 表达式的第一部分描述 lambda 捕获。从周围的作用域中捕获一个变量使它能够在 lambda 表达式中使用，而不需要显式地将它作为参数传递给 lambda 表达式。
C++ lambda 表达式支持通过复制变量或创建对变量的引用来捕获变量，但对于内核 lambda 表达式，变量只能通过复制来捕获。一般的做法是简单地使用默认的捕获模式
[=]，它通过值隐式地捕获所有变量，尽管也可以显式地命名每个捕获的变量。内核中使用的任何变量如果没有被值捕获，都会导致编译时错误。 -
lambda 表达式的第二部分描述传递给 lambda 表达式的参数，就像传递给命名函数的参数一样。
对于内核 lambda 表达式，参数取决于内核是如何被调用的，并且通常标识并行执行空间中工作项的索引。有关各种并行执行空间以及如何标识每个执行空间中工作项的索引的更多详细信息，请参考第四章。
-
lambda 表达式的最后一部分定义了 lambda 函数体。对于内核 lambda 表达式，函数体描述了应该在并行执行空间中的每个索引处执行的操作。
内核支持 lambda 表达式的其他部分，但这些部分要么是可选的，要么很少使用:
-
一些说明符(如
mutable)可能会受到支持，但不建议使用它们，并且在 SYCL(在临时 SYCL 2020 或 DPC++ 的未来版本中可能会删除支持。示例代码中没有显示任何内容。 -
支持异常规范，但是如果提供的话必须是
noexcept，因为内核不支持异常。 -
λ属性被支持，并且可以用来控制内核如何被编译。例如，
reqd_work_group_size属性可用于要求内核的特定工作组大小。 -
可以指定返回类型，但是如果提供的话必须是
void，因为内核不支持非void返回类型。
LAMBDA CAPTURES: IMPLICIT OR EXPLICIT?
一些 C++ 风格指南建议不要对 lambda 表达式进行隐式(或默认)捕获，因为可能会出现悬空指针问题，尤其是当 lambda 表达式跨越范围边界时。当使用 lambda 表示内核时，可能会出现相同的问题，因为内核 lambda 在设备上异步执行，与主机代码分离。
因为隐式捕获有用且简洁，所以它是 SYCL 内核的常见实践，也是我们在本书中使用的约定，但最终是我们决定是喜欢隐式捕获的简洁还是显式捕获的清晰。
命名内核 Lambda 表达式
当内核被写成 lambda 表达式时，在某些情况下还必须提供一个元素:因为 lambda 表达式是匿名的，有时 SYCL 需要一个显式的内核名称模板参数来唯一地标识被写成 lambda 表达式的内核。
图 10-4
命名内核 lambda 表达式
当内核由单独的设备代码编译器编译时，命名内核 lambda 表达式是主机代码编译器识别调用哪个内核的一种方式。命名一个内核 lambda 还支持编译后内核的运行时自省，或者通过名字构建一个内核，如图 10-9 所示。
为了在不需要内核名称模板参数时支持更简洁的代码，DPC++ 编译器支持通过-fsycl-unnamed-lambda编译器选项省略内核 lambda 的内核名称模板参数。使用该选项时，不需要显式的内核名称模板参数，如图 10-5 所示。
图 10-5
使用未命名的内核 lambda 表达式
因为 lambda 表达式的内核名称模板参数在大多数情况下是不需要的，所以我们通常可以从一个未命名的 lambda 开始，只有在需要内核名称模板参数的特定情况下才添加内核名称。
当不需要内核名称模板参数时，最好使用未命名的内核 lambdas 来减少冗余。
作为命名函数对象的内核
命名函数对象，也称为函子，是 C++ 中的一种既定模式，它允许对任意数据集合进行操作，同时保持定义良好的接口。当用于表示内核时，命名函数对象的成员变量定义内核可以操作的状态，并且为并行执行空间中的每个工作项目调用重载函数调用operator()。
命名函数对象需要比 lambda 表达式更多的代码来表达内核，但是额外的冗长提供了更多的控制和额外的能力。例如，分析和优化表示为命名函数对象的内核可能更容易，因为内核使用的任何缓冲区和数据值都必须显式传递给内核，而不是自动捕获。
最后，因为命名函数对象就像任何其他 C++ 类一样，表达为命名函数对象的内核可以是模板化的，这与表达为 lambda 表达式的内核不同。表示为命名函数对象的内核也更容易重用，并且可以作为单独头文件或库的一部分提供。
内核命名函数对象的元素
图 10-6 中的代码描述了一个被命名为函数对象的内核元素。
图 10-6
作为命名函数对象的内核
当一个内核被表示为一个命名的函数对象时，命名的函数对象类型必须遵循 C++11 规则，以便能够简单地复制。非正式地，这意味着命名的函数对象可以被安全地逐字节复制，使得命名的函数对象的成员变量能够被传递给在设备上执行的内核代码并由其访问。
重载函数调用operator()的参数取决于内核如何启动，就像用 lambda 表达式表示的内核一样。
因为函数对象是命名的，所以宿主代码编译器可以使用函数对象类型与设备代码编译器生成的内核代码关联，即使函数对象是模板化的。因此，不需要额外的内核名称模板参数来命名内核函数对象。
与其他 API 的互操作性
当 SYCL 实现建立在另一个 API 之上时，该实现可能能够与使用底层 API 机制定义的内核进行互操作。这使得应用程序可以轻松地、渐进地将 SYCL 集成到现有的代码库中。
因为 SYCL 实现可能位于许多其他 API 之上，所以本节描述的功能是可选的，并且可能不是所有实现都支持。根据具体的设备类型或设备供应商，底层 API 甚至可能有所不同！
概括地说，一个实现可能支持两种互操作性机制:来自 API 定义的源或中间表示(IR)或来自 API 特定的句柄。在这两种机制中，从 API 定义的源或中间表示创建内核的能力更容易移植，因为一些源或 IR 格式受多个 API 支持。例如，OpenCL C 内核可以被许多 API 直接使用，或者可以被编译成 API 可以理解的某种形式，但是来自一个 API 的特定于 API 的内核句柄不太可能被不同的 API 理解。
请记住，所有形式的互操作性都是可选的！
不同的 SYCL 实现可能支持从不同的 API 特定句柄创建内核——或者根本不支持。
请务必查看文档以了解详细信息！
与 API 定义的源语言的互操作性
通过这种形式的互操作性，内核的内容被描述为源代码，或者使用 SYCL 没有定义的中间表示，但是内核对象仍然是使用 SYCL API 调用创建的。这种形式的互操作性允许重用用其他源语言编写的内核库，或者使用特定领域语言(DSL)以中间表示形式生成代码。
实现必须理解内核源代码或中间表示，才能利用这种形式的互操作性。例如，如果内核是使用 OpenCL C 以源代码形式编写的，那么实现必须支持从 OpenCL C 内核源代码构建 SYCL 程序。
图 10-7 显示了如何将 SYCL 内核写成 OpenCL C 内核源代码。
图 10-7
从 OpenCL C 内核源代码创建的内核
在这个例子中，内核源字符串在 SYCL 主机 API 调用的同一个文件中被表示为 C++ 原始字符串文字，但并不要求必须如此，一些应用程序可能会从文件中读取内核源字符串，甚至实时生成它。
因为 SYCL 编译器无法看到用 API 定义的源语言编写的 SYCL 内核，所以任何内核参数都必须使用set_arg()或set_args()接口显式传递。SYCL 运行时和 API 定义的源语言必须就将对象作为内核参数传递的约定达成一致。在这个例子中，访问器dataAcc作为全局指针内核参数data被传递。
build_with_source()接口支持传递可选的 API 定义的构建选项，以精确控制内核的编译方式。在本例中，程序构建选项-cl-fast-relaxed-math用于指示内核编译器可以使用精度宽松的更快的数学库。程序构建选项是可选的，如果不需要构建选项，可以省略。
与 API 定义的内核对象的互操作性
有了这种形式的互操作性，内核对象本身在另一个 API 中创建，然后导入 SYCL。这种形式的互操作性使应用程序的一部分能够使用底层 API 直接创建和使用内核对象，而应用程序的另一部分能够使用 SYCL APIs 重用相同的内核。图 10-8 中的代码显示了如何从 OpenCL 内核对象创建 SYCL 内核。
图 10-8
从 OpenCL 内核对象创建的内核
与其他形式的互操作性一样，SYCL 编译器无法看到 API 定义的内核对象。因此，必须使用set_arg()或set_args()接口显式传递内核参数，并且 SYCL 运行时和底层 API 必须就传递内核参数的约定达成一致。
程序对象中的内核
在前面的章节中，当内核从 API 定义的表示或者从 API 特定的句柄创建时，内核分两步创建:首先通过创建一个程序对象，然后通过从程序对象创建内核。程序对象是作为一个单元编译的内核和它们调用的函数的集合。
对于表示为 lambda 表达式或命名函数对象的内核，包含内核的程序对象通常是隐式的，对应用程序不可见。对于需要更多控制的应用程序，应用程序可以显式地管理内核和封装它们的程序对象。为了描述为什么这可能是有益的，简单看一下有多少 SYCL 实现管理实时(JIT)内核编译是有帮助的。
虽然规范没有要求，但许多实现都“懒惰地”编译内核这通常是一个好策略，因为它确保了应用程序的快速启动，并且不会不必要地编译从不执行的内核。这种策略的缺点是内核的第一次使用通常比随后的使用需要更长的时间，因为它包括编译内核所需的时间，加上提交和执行内核所需的时间。对于一些复杂的内核，编译内核所需的时间可能会很长，因此需要在应用程序执行期间将编译转移到不同的点，例如当应用程序正在加载时，或者在单独的后台线程中。
一些内核也可能受益于实现定义的“构建选项”,以精确控制内核的编译方式。例如，对于某些实现，可以指示内核编译器使用精度更低、性能更好的数学库。
为了更好地控制内核编译的时间和方式，应用程序可以使用特定的编译选项，明确请求在使用内核之前编译内核。然后，像往常一样，可以将预编译的内核提交到队列中执行。图 10-9 显示了这是如何工作的。
图 10-9
使用构建选项编译内核 lambdas
在这个例子中，一个程序对象是从 SYCL 上下文中创建的，由指定的模板参数定义的内核是使用build_with_kernel_type函数构建的。对于这个例子，程序构建选项-cl-fast-relaxed-math表明内核编译器可以使用具有宽松精度的更快的数学库，但是程序构建选项是可选的，如果不需要特殊的程序构建选项，可以省略。在这种情况下，命名内核 lambda 的模板参数是必需的，以标识要编译的内核。
程序对象也可以从上下文和设备的特定列表中创建，而不是从上下文中的所有设备中创建，从而允许一组设备的程序对象用与另一组设备的另一程序对象不同的构建选项来编译。
除了通常的内核 lambda 表达式之外，还使用get_kernel函数将之前编译的内核传递给parallel_for。这确保了使用宽松数学库构建的先前编译的内核得到使用。如果先前编译的内核没有被传递给parallel_for，那么内核将被再次编译，没有任何特殊的编译选项。这可能在功能上是正确的，但肯定不是预期的行为！
在许多情况下，例如在前面显示的简单示例中，这些额外的步骤不太可能对应用程序的行为产生明显的改变，为了清楚起见，可以省略这些步骤，但是在针对性能调整应用程序时，应该考虑这些步骤。
IMPROVING INTEROPERABILITY AND PROGRAM OBJECT MANAGEMENT
尽管本章中描述的 SYCL 互操作性和程序对象管理接口非常有用，但它们可能会在 SYCL 和 DPC++ 的未来版本中得到改进和增强。请参考最新的 SYCL 和 DPC++ 文档，查找本书中没有提供或不够稳定的更新！
摘要
在这一章中，我们探索了定义内核的不同方法。我们描述了如何通过将内核表示为 C++ lambda 表达式或命名函数对象来无缝集成到现有的 C++ 代码库中。对于新的代码库，我们还讨论了不同内核表示的优缺点，以帮助根据应用程序或库的需求选择定义内核的最佳方式。
我们还描述了如何与其他 API 进行互操作，或者通过从 API 定义的源语言或中间表示创建内核，或者通过从内核的 API 表示的句柄创建内核对象。互操作性使应用程序能够随着时间的推移从较低级别的 API 迁移到 SYCL，或者与为其他 API 编写的库接口。
最后，我们描述了如何在 SYCL 应用程序中编译内核，以及如何直接操作程序对象中的内核来控制编译过程。尽管大多数应用程序不需要这种级别的控制，但在调优应用程序时，这是一项需要注意的有用技术。
开放存取本章根据知识共享署名 4.0 国际许可证(http://Creative Commons . org/licenses/by/4.0/)的条款获得许可，该许可证允许以任何媒体或格式使用、共享、改编、分发和复制，只要您适当注明原作者和来源，提供知识共享许可证的链接并指明是否进行了更改。
本章中的图像或其他第三方材料包含在该章的知识共享许可中，除非该材料的信用额度中另有说明。如果材料未包含在本章的知识共享许可中，并且您的预期用途不被法定法规允许或超出了允许的用途，您将需要直接从版权所有者处获得许可。
十一、向量
向量是数据的集合。这可能很有用，因为我们计算机中的并行性来自计算硬件的集合，并且数据通常在相关分组中处理(例如，RGB 像素中的颜色通道)。听起来像是天作之合吗？这非常重要，我们将用一章来讨论向量类型的优点以及如何利用它们。在这一章中，我们不会深入探讨矢量化，因为矢量化会因设备类型和实现而异。矢量化将在第 15 和 16 章节中介绍。
本章试图解决以下问题:
-
什么是向量类型？
-
我真的需要了解多少关于 vector 接口的知识？
-
是否应该用向量类型来表示并行性？
-
什么时候应该使用向量类型？
我们使用工作代码示例讨论可用向量类型的优点和缺点，并强调利用向量类型的最重要的方面。
如何思考向量
当我们与并行编程专家交谈时，向量是一个令人惊讶的有争议的话题，根据作者的经验，这是因为不同的人以不同的方式定义和思考这个术语。
有两种广义的方式来考虑向量数据类型(数据集合):
-
作为一种方便的类型，它将您可能想要引用和操作的数据分组为一组，例如，将一个像素的颜色通道(如 RGB、YUV)分组为一个变量(如
float3)，该变量可以是一个向量。我们可以定义一个 pixel 类或结构，并在其上定义像+这样的数学运算符，但是向量类型可以方便地为我们开箱即用。便利类型可以在许多用于编程 GPU 的着色器语言中找到，因此这种思维方式在许多 GPU 开发人员中已经很常见。 -
作为描述代码如何映射到硬件中的 SIMD 指令集的机制。例如，在一些语言和实现中，
float8上的操作理论上可以映射到硬件中的八通道 SIMD 指令。向量类型在多种语言中被用作针对特定指令集的 CPU 特定 SIMD 内部函数的一种方便的高级替代方法，因此这种思维方式在许多 CPU 开发人员中已经很普遍了。
虽然这两种解释非常不同，但当 SYCL 和其他语言变得适用于 CPU 和 GPU 时，它们无意中被结合在一起并混淆在一起。SYCL 1.2.1 规范中的 vector 与这两种解释都是兼容的(我们将在后面重新讨论这一点)，所以在进一步讨论之前，我们需要澄清一下我们在 DPC++ 中推荐的思路。
在本书中，我们讨论了如何将工作项组合在一起，以公开强大的通信和同步原语，例如子组障碍和洗牌。为了使这些操作在向量硬件上有效，假设子组中的不同工作项组合并映射到 SIMD 指令。换句话说，多个工作项被编译器组合在一起，此时它们可以映射到硬件中的 SIMD 指令。请记住第四章中的内容，这是在矢量硬件上运行的 SPMD 编程模型的基本前提，其中单个工作项构成了硬件中可能是 SIMD 指令的通道*，而不是定义硬件中 SIMD 指令的整个操作的工作项。当在硬件中映射到 SIMD 指令时，当使用 DPC++ 编译器以 SPMD 风格编程时，您可以认为编译器总是跨工作项进行矢量化。*
对于本书中描述的功能和硬件，向量主要用于本节的第一种解释——向量是方便的类型，不应被视为映射到硬件中的 SIMD 指令。在适用的平台(CPU、GPU)上，工作项被组合在一起形成硬件中的 SIMD 指令。向量应该被认为是提供了方便的操作符，如 swizzles 和数学函数，使我们的代码中对数据组的常见操作变得简洁(例如，添加两个 RGB 像素)。
对于来自没有向量的语言或来自 GPU 着色语言的开发人员，我们可以将 SYCL 向量视为工作项的本地向量，因为如果添加两个四元素向量，该添加可能需要硬件中的四条指令(从工作项的角度来看，它将被标量化)。向量的每个元素将通过硬件中不同的指令/时钟周期相加。根据这种解释，向量是一种便利，因为我们可以在源代码中的一次操作中添加两个向量，而不是在源代码中执行四次标量操作。
对于来自 CPU 背景的开发人员，我们应该知道，在编译器中默认情况下，隐式向量化到 SIMD 硬件以几种独立于向量类型的方式发生。编译器跨工作项执行这种隐式矢量化，从格式良好的循环中提取矢量操作，或者在映射到矢量指令时支持矢量类型——有关更多信息，请参见第十六章。
OTHER IMPLEMENTATIONS POSSIBLE!
SYCL 和 DPC++ 的不同编译器和实现在理论上可以对代码中的向量数据类型如何映射到向量硬件指令做出不同的决定。我们应该阅读供应商的文档和优化指南，以了解如何编写能够映射到高效 SIMD 指令的代码。这本书主要是针对 DPC++ 编译器编写的，因此记录了围绕它构建的思维和编程模式。
CHANGES ARE ON THE HORIZON
我们刚刚说过，在考虑映射到设备上的硬件时，将向量类型视为便利类型，并期望跨工作项的矢量化。这有望成为 DPC++ 编译器和工具链未来的默认解释。然而，有两个额外的前瞻性变化需要注意。
首先，我们可以期待一些未来的 DPC++ 特性，将允许我们编写直接映射到硬件中 SIMD 指令的显式矢量代码，特别是对于那些希望针对特定架构调整代码细节并从编译器矢量器中获得控制权的专家。这是一个很少开发人员会使用的利基特性，但是我们可以预期编程机制最终会在可能的地方存在。这些编程机制将非常清楚哪些代码是以显式矢量风格编写的，因此我们今天编写的代码和新的更显式(且可移植性更差)的风格之间不会混淆。
第二，这本书的这一部分(讨论向量的解释)强调了对向量的含义存在混淆，这将在未来的 SYCL 中得到解决。在 SYCL 2020 临时规范中对此有所暗示，其中描述了一种数学数组类型(marray)，这显然是本节的第一种解释——一种与矢量硬件指令无关的方便类型。我们应该期待另一种类型也最终出现来覆盖第二种解释，很可能与 C++ std::simd模板一致。由于这两种类型与 vector 数据类型的特定解释明确相关，我们作为程序员的意图将从我们编写的代码中变得清晰。这将更不容易出错，更不容易混淆，甚至可能减少专家开发者之间的激烈讨论，当问题出现时“什么是是向量？”
向量类型
SYCL 中的 Vector 类型是跨平台的类模板，可以在设备和主机 C++ 代码中高效工作，并允许在主机及其设备之间共享 vector。Vector 类型包括允许从一组重组的组件元素构建新 vector 的方法，这意味着新 vector 的元素可以以任意顺序从旧 vector 的元素中选取。vec是一种 vector 类型，可以在目标设备后端编译成内置的 vector 类型，并在主机上提供兼容的支持。
vec类根据其元素数量和元素类型进行模板化。元素数参数numElements可以是 1、2、3、4、8 或 16 中的一个。任何其他值都将导致编译失败。元素类型参数dataT必须是设备代码支持的基本标量类型之一。
SYCL vec类模板提供了与由vector_t定义的底层向量类型的互操作性，后者仅在为设备编译时可用。vec类可以从vector_t的实例构建，并且可以隐式转换为vector_t的实例，以便支持与来自内核函数的本地 SYCL 后端(例如 OpenCL 后端)的互操作性。当元素数量为 1 时，为了使单元素向量和标量易于互换，vec类模板的实例也可以隐式转换为数据类型的实例。
为了编程方便，SYCL 提供了许多形式为using <type><elems> = vec<<storage-type>、<elems>>的类型别名，其中<elems>是 2 、 3 、 4 、 8 和 16 ，整数类型的<type>和<storage-type>的配对是char【⇔int8_t、uchar ⇔ uint8_t、uint8_t uint ⇔ uint32_t，long ⇔ int64_t，and ulong ⇔ uint64_t对于浮点型half，float，double。 例如，uint4是vec < uint32_t，4 > ，float16是vec < float，16 > 的别名。
矢量接口
向量类型的功能通过类vec公开。vec类表示一组组合在一起的数据元素。vec类模板的构造器、成员函数和非成员函数的接口如图 11-1 、 11-4 和 11-5 所示。
图 11-2 中列出的 XYZW 成员仅在numElements <= 4时可用。RGBA 会员仅在numElements == 4时可用。
图 11-3 中的lo、hi、odd、even成员仅在numElements > 1时可用。
图 11-5
vec非成员函数
图 11-4
vec成员函数
图 11-3
vec运算符界面
图 11-2
swizzled_vec成员函数
图 11-1
vec类声明和成员函数
加载和存储成员函数
向量加载和存储操作是vec类的成员，用于加载和存储向量的元素。这些操作可以对与向量的通道类型相同的元素数组进行。示例如图 11-6 所示。
图 11-6
使用加载和存储成员函数。
在vec类中，dataT和numElements是反映vec的组件类型和维度的模板参数。
load()成员函数模板将从multi_ptr地址的内存中读取dataT类型的值，在dataT的元素中偏移numElements*offset，并将这些值写入 vec 的通道。
store()成员函数模板将读取 vec 的通道，并将这些值写入 multi_ptr 地址的内存，在dataT的元素中偏移numElements*offset。
该参数是一个multi_ptr而不是一个访问器，这样本地创建的指针和从主机传递的指针都可以使用。
multi_ptr的数据类型是dataT,``vec类专门化的组件的数据类型。这要求传递给load()或store()的指针必须匹配vec实例本身的类型。
调酒业务
在图形应用中，重组意味着重新排列向量的数据元素。例如，如果a = {1, 2, 3, 4,}，并且知道一个四元素向量的分量可以称为{x, y, z, w}，我们可以写b = a.wxyz().，变量b中的结果将是{4, 1, 2, 3}。这种形式的代码在 GPU 应用中很常见，在这些应用中有高效的硬件来执行这种操作。调酒有两种方式:
-
通过调用一个
vec的 swizzle 成员函数，该函数接受在0和numElements-1之间的可变数量的整数模板参数，指定 swizzle 索引 -
通过调用一个简单的 swizzle 成员函数，比如
XYZW_SWIZZLE和RGBA_SWIZZLE
请注意，简单的 swizzle 函数仅适用于最多四个元素的向量，并且仅在包含sycl.hpp之前定义宏SYCL_SIMPLE_SWIZZLES时可用。在这两种情况下，返回类型总是一个__swizzled_vec__的实例，一个实现定义的临时类，表示原始vec实例的重组。swizzle 成员函数模板和简单的 swizzle 成员函数都允许重复使用 swizzle 索引。图 11-7 显示了__swizzled_vec__的简单用法。
图 11-7
使用__swizzled_vec__类的例子
并行内核中的向量执行
如章节 4 和 9 所述，一个工作项是并行层次结构的叶节点，代表一个内核函数的单个实例。工作项目可以以任何顺序执行，并且不能彼此通信或同步，除非通过对本地和全局存储器的原子存储器操作或通过组集合函数(例如，shuffle、barrier)。
正如本章开始时所描述的，DPC++ 中的 vector 应该被解释为方便我们编写代码。每个向量对于单个工作项来说是局部的(而不是与硬件中的矢量化相关)，因此可以被认为相当于我们工作项中的私有数组numElements。例如，“float4 y4”申报的存储相当于float y4[4]。考虑图 11-8 所示的例子。
图 11-8
向量执行示例
对于标量变量 x，在具有 SIMD 指令的硬件(例如，CPU、GPU)上具有多个工作项目的内核执行的结果可能使用向量寄存器和 SIMD 指令，但是矢量化是跨工作项目的，并且与我们代码中的任何向量类型无关。每个工作项可以在隐式vec_x中的不同位置上操作，如图 11-9 所示。工作项中的标量数据可以被认为是跨同时执行的工作项隐式矢量化(组合到 SIMD 硬件指令中)，在一些实现中和在一些硬件上，但是我们编写的工作项代码不以任何方式对此进行编码——这是 SPMD 编程风格的核心。
图 11-9
从标量变量x到vec_x[8]的向量扩展
如图 11-9 所示，通过编译器从标量变量x到vec_x[8]的隐式向量扩展，编译器从出现在多个工作项中的标量操作在硬件中创建 SIMD 操作。
对于向量变量y4，多个工作项的内核执行结果，例如八个工作项，不通过使用硬件中的向量运算来处理 vec4。相反，每个工作项独立地看到自己的向量，向量上元素的操作跨多个时钟周期/指令发生(向量被编译器标量化)，如图 11-10 所示。
图 11-10
垂直扩展到相当于八个工作项的y4的vec_y[8][4]
每个工作项都可以看到 y4 的原始数据布局，这为推理和调整提供了一个直观的模型。性能下降是编译器必须为 CPU 和 GPU 生成聚集/分散内存指令，如图 11-11 所示(向量在内存中是连续的，相邻的工作项并行操作不同的向量)，因此当编译器跨工作项(例如，跨子组)进行矢量化时，标量通常是一种比显式向量更有效的方法。详见第十五章和第十六章。
图 11-11
带有地址转义的矢量代码示例
当编译器能够证明y4的地址没有从当前内核工作项中转义或者所有被调用函数都将被内联时，编译器可以执行优化，就像使用一组向量寄存器从y4到vec_y[4][8]进行水平单位步长扩展一样，如图 11-12 所示。在这种情况下，编译器无需为 CPU 和 GPU 生成聚集/分散 SIMD 指令，就能获得最佳性能。编译器优化报告为程序员提供了关于这种类型的转换的信息，无论它是否发生，并且可以提供关于如何调整我们的代码以提高性能的提示。
图 11-12
水平单位步幅扩展到y4的vec_y[4][8]
向量并行性
尽管 DPC++ 源代码中的向量应该被解释为只局限于单个工作项的便利工具，但是如果没有提到硬件中的 SIMD 指令是如何操作的，这一章关于向量的内容是不完整的。这一讨论与我们源代码中的向量无关，但提供了正交背景，这将有助于我们进入本书后面描述特定设备类型(GPU、CPU、FPGA)的章节。
现代的 CPU 和 GPU 包含 SIMD 指令硬件，其对包含在一个向量寄存器或寄存器文件中的多个数据值进行操作。例如，借助英特尔 x86 AVX-512 和其他现代 CPU SIMD 硬件，SIMD 指令可用于开发数据并行性。在提供 SIMD 硬件的 CPU 和 GPU 上，我们可以考虑一个向量加法运算，比如对一个八元素向量，如图 11-13 所示。
图 11-13
八路数据并行的 SIMD 加法
这个例子中的向量加法可以在向量硬件上的单个指令中执行，将向量寄存器vec_x和vec_y与 SIMD 指令并行相加。
以独立于硬件的方式展示潜在的并行性，确保我们的应用可以扩展(或缩小)以适应不同平台的功能，包括那些具有矢量硬件指令的平台。在应用程序开发过程中，在工作项目和其他形式的并行性之间取得正确的平衡是我们都必须面对的挑战，这将在第 15 、 16 和 17 章中详细介绍。
摘要
在编程语言中，术语向量有多种解释，当我们想要编写高性能和可伸缩的代码时，理解特定语言或编译器的解释是非常重要的。DPC++ 和 DPC++ 编译器是围绕这样的思想构建的，即源代码中的向量是工作项本地的便利函数，编译器跨工作项的隐式向量化可以映射到硬件中的 SIMD 指令。当我们想要编写直接映射到 vector 硬件的代码时，我们应该查看供应商文档以及 SYCL 和 DPC++ 的未来扩展。使用多个工作项(例如 ND-range)编写我们的内核并依靠编译器跨工作项进行矢量化应该是大多数应用程序的编写方式，因为这样做利用了 SPMD 的强大抽象，它提供了一个易于理解的编程模型，并提供了跨设备和架构的可扩展性能。
本章描述了vec接口，当我们想要对相似类型的数据进行分组操作时，它提供了开箱即用的便利(例如，一个像素有多个颜色通道)。它还简要介绍了硬件中的 SIMD 指令，为我们在第 15 和 16 章节中更详细的讨论做准备。
开放存取本章根据知识共享署名 4.0 国际许可证(http://Creative Commons . org/licenses/by/4.0/)的条款获得许可，该许可证允许以任何媒体或格式使用、共享、改编、分发和复制，只要您适当注明原作者和来源，提供知识共享许可证的链接并指明是否进行了更改。
本章中的图像或其他第三方材料包含在该章的知识共享许可中，除非该材料的信用额度中另有说明。如果材料未包含在本章的知识共享许可中，并且您的预期用途不被法定法规允许或超出了允许的用途，您将需要直接从版权所有者处获得许可。
十二、设备信息
第二章向我们介绍了将工作导向特定设备的机制——控制代码执行的*。*在本章中，我们将探讨如何适应运行时出现的设备。
我们希望我们的程序是可移植的。为了便于携带，我们需要我们的程序适应设备的功能。我们可以将程序参数化，只使用现有的功能，并根据设备的具体情况调整代码。如果我们的程序不是为适应而设计的，那么就会发生不好的事情，包括执行缓慢或程序失败。
幸运的是，SYCL 规范的创建者考虑到了这一点，并给了我们接口让我们解决这个问题。SYCL 规范定义了一个device类，它封装了一个可以执行内核的设备。查询设备类的能力，使我们的程序能够适应设备的特性和能力，是本章所教授的核心。
我们中的许多人会从有逻辑来弄清楚“是否有 GPU 存在？”通知我们的程序在执行时将做出的选择。这是本章内容的开始。正如我们将看到的，有更多的信息可以帮助我们使我们的程序健壮和高性能。
将程序参数化有助于正确性、可移植性、性能可移植性和未来的检验。
本章深入探讨了最重要的查询以及如何在我们的程序中有效地使用它们。
特定于设备的属性可以使用get_info, but来查询。DPC++ 与 SYCL 1.2.1 的不同之处在于，它完全重载了get_info,以减少使用get_work_group_info来获取工作组信息的需要，而工作组信息实际上是特定于设备的信息。DPC++ 不支持使用get_work_group_info。这一变化意味着特定于设备的内核和工作组属性可以作为特定于设备的属性的查询被正确地找到(get_info)。这纠正了继承自 OpenCL 的 SYCL 1.2.1 中仍然存在的令人困惑的历史异常。
精炼内核代码，使其更具规范性
考虑到我们的编码，一个内核一个内核地，大致可以分为三类:
-
通用内核代码:在任何地方运行，不针对特定的设备类别。
-
特定于设备类型的内核代码:在一种类型的设备(例如，GPU、CPU、FPGA)上运行，不针对设备类型的特定型号进行调整。这非常有用，因为许多设备类型具有共同的特性，所以可以放心地做出一些假设，这些假设不适用于为所有设备编写的完全通用的代码。
-
特定于设备的调优内核代码:在一种设备上运行，针对设备的特定参数进行调优——这涵盖了从少量调优到非常详细的优化工作的广泛可能性。
作为程序员，我们的工作是确定不同的设备类型何时需要不同的模式。我们用第十四章、 15 章、 16 章和 17 章来阐明这一重要思想。
最常见的是从实现通用内核代码开始，让它工作起来。第二章专门讨论了在开始内核实现时什么方法最容易调试。一旦我们有了一个工作的内核，我们就可以对它进行改进，以针对特定设备类型或设备型号的功能。
第十四章提供了一个思考框架，在我们深入考虑设备之前，先考虑并行性。我们选择的模式(也就是算法)决定了我们的代码，作为程序员，我们的工作就是决定不同的设备何时需要不同的模式。第 15 (GPU)、 16 (CPU)和 17 (FPGA)章节更深入地探究了区分这些设备类型的品质，并激发了对使用模式的选择。当不同设备类型上的方法(模式选择)不同时，正是这些品质促使我们考虑为不同设备编写不同版本的内核。
当我们为特定类型的设备(例如，特定的 CPU、GPU、FPGA 等)编写内核时。)，使其适应特定厂商甚至此类设备的型号是合乎逻辑的。良好的编码风格是基于特性(例如，从设备查询中找到的项目大小支持)对代码进行参数化。
我们应该编写代码来查询描述设备实际功能的参数，而不是其营销信息；查询设备的型号并对其做出反应是非常糟糕的编程实践——这样的代码可移植性较差。
通常为我们想要支持的每种设备类型编写不同的内核(内核的 GPU 版本和内核的 FPGA 版本，可能还有内核的通用版本)。当我们变得更具体时，为了支持特定的设备供应商或甚至设备模型，当我们可以参数化内核而不是复制它时，我们可能会受益。如果我们认为合适，我们可以自由选择。有太多参数调整的代码可能难以阅读，或者在运行时负担过重。然而，参数可以整齐地放入内核的一个版本是很常见的。
当算法大体相同，但针对特定设备的功能进行了调整时，参数化最有意义。当使用完全不同的方法、模式或算法时，编写不同的内核要干净得多。
如何枚举设备和功能
第二章列举并解释了选择执行设备的五种方法。本质上，方法#1 是最不规范的在某个地方运行它，我们进化到最规范的方法#5，它考虑在一系列设备中的一个相当精确的设备模型上执行。介于两者之间的列举方法混合了灵活性和规定性。图 12-1 、 12-2 和 12-3 有助于说明我们如何选择器件。
图 12-1 显示，即使我们允许实现为我们选择一个默认设备(第二章中的方法#1)，我们仍然可以查询关于所选设备的信息。
图 12-2 展示了我们如何尝试使用一个特定的设备(在这个例子中，是一个 GPU)来建立一个队列，但是如果没有 GPU 可用的话，就明确地回到主机上。这给了我们一些设备选择的控制权。如果我们简单地使用默认队列，我们可能会以意外的设备类型(例如，DSP、FPGA)结束。如果我们明确地想要在没有 GPU 设备的情况下使用主机设备，这段代码会为我们做到这一点。回想一下，主机设备总是保证存在的，所以我们不需要担心使用host_selector。
不建议我们使用图 12-2 所示的解决方案。除了看起来有点吓人和容易出错之外，图 12-2 并没有让我们控制选择哪一个 GPU，因为如果有多个可用的 GPU，我们得到哪一个取决于实现。尽管既有启发性又有实用性，但还有更好的方法。建议我们编写自定义的设备选择器，如下一个代码示例所示(图 12-3 )。
自定义设备选择器
图 12-3 使用自定义设备选择器。自定义设备选择器在第二章中首次讨论，作为选择代码运行位置的方法#5(图 2-15 )。定制设备选择器使其operator()，如图 12-3 所示，为应用可用的每个设备调用。所选设备是得分最高的设备。 1 在这个例子中，我们将为我们的选择器开一点玩笑:
-
拒绝供应商名称包含“Martian”(return–1)字样的 GPU。
-
偏爱供应商名称包含单词“ACME”的 GPU(返回 824)。
-
其他任何 GPU 都是好的(return 799)。
-
如果没有 GPU，我们选择主机设备(返回 99)。
-
忽略所有其他设备(return–1)。
下一部分，“好奇:get_info<>”深入到get_devices(), get_platforms()和get_info<>提供的丰富信息。这些接口打开了我们可能想要用来挑选设备的任何类型的逻辑，包括图 2-15 和 12-3 所示的简单的供应商名称检查。
图 12-1
默认情况下分配给我们的设备
关于设备的查询依赖于安装的软件(特殊的用户级驱动程序)来响应关于设备的查询。SYCL 和 DPC++ 依赖于此，就像操作系统需要驱动程序来访问硬件一样——仅仅将硬件安装在机器上是不够的。
图 12-3
定制设备选择器—我们的首选解决方案
图 12-2
如果可能，使用 try-catch 选择 GPU 设备，否则选择主机设备
好奇:get_info<>
为了让我们的程序“知道”哪些设备在运行时可用，我们可以让我们的程序从 device 类中查询可用的设备，然后我们可以使用get_info<>查询特定的设备来了解更多的细节。我们提供了一个简单的程序，名为好奇(见图 12-4 )，它使用这些接口将信息转储出来让我们直接查看。这对于在开发或调试使用这些接口的程序时进行健全性检查非常有用。这个程序不能按预期工作通常可以告诉我们，我们需要的软件驱动程序没有正确安装。图 12-5 显示了该程序的示例输出，其中包含关于当前设备的高级信息。
图 12-5
好奇. cpp 的示例输出
图 12-4
设备查询机制的简单使用
更好奇:详细的枚举代码
我们提供了一个程序，我们将其命名为 verycurious.cpp(图 12-6 ，来说明使用get_info<>可以获得的一些详细信息。同样，我们发现自己编写这样的代码有助于开发或调试程序。图 12-5 显示了该程序的输出样本，以及关于当前设备的底层信息。
现在我们已经展示了如何访问信息，我们将讨论在应用程序中查询和操作最重要的信息字段。
图 12-6
设备查询机制的更详细的使用:query 好奇. cpp
好奇:get_info<>
has_extension()接口允许程序直接测试一个特性，而不是像前面的代码示例那样遍历来自get_info <info::platform::extensions>的扩展列表。SYCL 2020 临时规范定义了新的机制来查询设备的扩展和详细方面，但我们不会在本书中涵盖这些功能(这些功能刚刚完成)。更多信息请参考在线 oneAPI DPC++ 语言参考。
设备信息描述符
本章前面使用的“好奇”程序示例利用了最常用的 SYCL 设备类成员函数(即is_host, is_cpu, is_gpu, is_accelerator, get_info, has_extension)。这些成员函数记录在 SYCL 规范的“SYCL 设备类的成员函数”表中(在 SYCL 1.2.1 中，是表 4.18)。
“好奇”程序示例也使用get_info成员函数查询信息。包括主机设备在内的所有 SYCL 设备都必须支持一组查询。SYCL 规范中题为“器件信息描述符”的表格描述了此类项目的完整列表(在 SYCL 1.2.1 中为表 4.20)。
设备特定的内核信息描述符
像平台和设备一样，我们可以使用get_info函数查询关于我们内核的信息。这些信息(例如，支持的工作组大小、首选的工作组大小、每个工作项所需的私有内存量)是特定于设备的，因此kernel类的get_info成员函数接受一个device作为参数。
DEVICE-SPECIFIC KERNEL INFORMATION IN SYCL 1.2.1
出于 OpenCL 命名的历史原因，SYCL 继承了名为kernel::get_info和kernel::get_work_group_info的查询组合，分别返回关于内核对象的信息和关于内核在特定设备上执行的信息。
在 DPC++ 和 SYCL(从 2020 年临时版本开始)中使用重载允许通过单一的get_info API 支持这两种类型的信息。
细节:那些“正确”的细节
我们将把细节分为关于必要条件(正确性)的信息和对调优有用但对正确性不必要的信息。
在这第一个正确性类别中，我们将列举内核正常启动应该满足的条件。不遵守这些设备限制将导致程序失败。图 12-7 显示了我们如何获取这些参数中的一部分，使得这些值可以在主机代码和内核代码中使用(通过 lambda 捕获)。我们可以修改代码来利用这些信息；例如，它可以指导我们关于缓冲区大小或工作组大小的代码。
图 12-7
获取可用于塑造内核的参数
提交不满足这些条件的内核将会产生错误。
设备查询
device_type: cpu, gpu, accelerator, custom, 2 automatic, host, all。这些最常由is_host(), is_cpu, is_gpu(),等测试(见图 12-6 ):
-
max_work_item_sizes``:``nd_range工作组每个维度允许的最大工作项数。非定制设备的最小值为(1, 1, 1)。 -
在单个计算单元上执行内核的工作组中允许的最大工作项目数。最小值为 1。
-
global_mem_size:全局内存的大小，以字节为单位。 -
local_mem_size:本地内存的大小，以字节为单位。除定制设备外，最小尺寸为 32 K。 -
在 SYCL 规范中没有详细说明的设备特定信息，通常是供应商特定的，如我们的
verycurious程序所示(图 12-6 )。 -
max_compute_units:表示设备上可用的并行数量——由实施定义，请小心解读！ -
sub_group_sizes:返回设备支持的子组大小集合。 -
如果该设备支持显式 USM 中描述的设备分配，则
usm_device_allocations:返回true。 -
如果该设备可以访问主机分配，则
usm_host_allocations:返回true。 -
如果该设备支持共享分配，则
usm_shared_allocations:返回true。 -
如果该设备支持由设备上的“受限 USM”的限制所管理的共享分配，则
usm_restricted_shared_allocations:返回true。该属性要求属性usm_shared_allocations为该设备返回true。 -
如果系统分配器可以代替 USM 分配机制用于该设备上的共享分配，则
usm_system_allocator:返回true。
我们建议在程序逻辑中避免 max_compute_units。
我们发现应该避免查询计算单元的最大数量，部分原因是这个定义不够清晰，无法用于代码调优。大多数程序应该表达它们的并行性，并让运行时将其映射到可用的并行性上，而不是使用max_compute_units。依赖于max_compute_units的正确性只有在增加了特定于实现和设备的信息时才有意义。专家可能会这样做，但大多数开发人员没有也不需要这样做！在这种情况下，让运行时完成它的工作！
内核查询
执行这些内核查询需要第十章“程序对象中的内核”中讨论的机制:
-
work_group_size:返回可用于在特定设备上执行内核的最大工作组大小 -
compile_work_group_size:返回由内核指定的工作组大小(如果适用)；否则返回(0，0，0) -
如果适用，返回由内核指定的子组大小；否则返回 0
-
如果适用，返回由内核指定的子组的数量；否则返回 0
-
max_sub_group_size:返回以指定工作组大小启动的内核的最大子组大小 -
max_num_sub_groups:返回内核子组的最大数量
细节:那些“调整/优化”
有几个额外的参数可以考虑作为我们内核的微调参数。这些可以被忽略，而不会危及程序的正确性。这些允许我们的内核真正利用硬件的细节来提高性能。
关注这些查询的结果有助于优化缓存(如果存在的话)。
设备查询
-
global_mem_cache_line_size:全局内存缓存行的大小，以字节为单位。 -
global_mem_cache_size:全局内存缓存的大小，以字节为单位。 -
local_mem_type:支持的本地存储器类型。这可以是暗示专用本地存储器存储的info::local_mem_type::local，例如 SRAM 或info::local_mem_type::global。后一种类型意味着本地内存只是作为全局内存之上的一种抽象来实现，没有任何性能提升。对于自定义设备(仅限)，本地内存类型也可以是info::local_mem_type::none，表示不支持本地内存。
内核查询
-
preferred_work_group_size:在特定设备上执行内核的首选工作组规模。 -
在特定设备上执行内核的首选工作组规模
运行时与编译时属性
本章中描述的查询是通过运行时 API(get_info)执行的，这意味着直到运行时才知道结果。这涵盖了许多用例，但 SYCL 规范也正在努力提供编译时的属性查询，当工具链知道它们时，允许更高级的编程技术，如基于设备属性的内核模板化。对于现有的运行时查询，基于查询的代码编译时适应是不可能的，这种能力对于高级优化或编写使用一些扩展的内核非常重要。在编写本书时，这些接口还没有定义得足够好来描述这些接口，但是我们可以期待 SYCL 和 DPC++ 中即将出现的更强大的查询和代码适应机制！查看在线 oneAPI DPC++ 语言参考和 SYCL 规范以获取更新。
摘要
最具移植性的程序会查询系统中可用的设备，并根据运行时信息调整它们的行为。这一章打开了通向丰富信息的大门，这些信息允许对我们的代码进行这样的裁剪，以适应运行时出现的硬件。
通过将我们的应用程序参数化以适应硬件的特性，我们的程序可以变得更加可移植，性能更加可移植，并且更加经得起未来的考验。我们还可以测试当前的硬件是否在我们在程序设计中所做的任何假设的范围内，并且当发现硬件超出我们的假设范围时，发出警告或中止。
开放存取本章根据知识共享署名 4.0 国际许可证(http://Creative Commons . org/licenses/by/4.0/)的条款获得许可，该许可证允许以任何媒体或格式使用、共享、改编、分发和复制，只要您适当注明原作者和来源，提供知识共享许可证的链接并指明是否进行了更改。
本章中的图像或其他第三方材料包含在本章的知识共享许可中，除非在材料的信用额度中另有说明。如果材料不包括在本章的知识共享许可中，并且您的预期使用不被法律法规允许或超出了允许的使用范围，您将需要直接从版权所有者处获得许可。
Footnotes 1如果我们的设备选择器只返回负值，那么在图 12-2 中的非 GPU 系统上，my_selector()将抛出一个runtime_error异常。因为我们为主机返回一个正值，所以这不会发生在图 12-3 中。
2
本书不讨论定制设备。如果我们发现自己正在编写一个使用自定义类型来标识自己的设备，我们将需要研究该设备的文档来了解更多信息。
十三、实用技巧
这一章包含了许多有用的信息、实用的技巧、建议和技术，它们在编程 SYCL 和使用 DPC++ 时被证明是有用的。这些主题没有一个是详尽的，所以目的是提高认识和鼓励更多的学习。
获取 DPC++ 编译器和代码示例
第一章讲述了如何获得 DPC++ 编译器(oneapi . com/implementations或 github. com/ intel/ llvm )以及从哪里获得代码示例(www . a press . com/9781484255735—寻找本书的服务:源代码)。再次提到这一点是为了强调尝试这些示例是多么有用(包括进行修改！)来获得实践经验。加入那些知道图 1-1 中的代码实际打印出什么的人的俱乐部吧！
在线论坛和文档
英特尔开发人员专区举办了一个论坛，用于讨论 DPC++ 编译器、DPC++ 库(第章第十八部分)、DPC++ 兼容性工具(用于 CUDA 迁移——将在本章稍后讨论)以及 oneAPI 工具包中包含的 gdb(本章也涉及调试)。这是一个张贴关于编写代码的问题(包括可疑的编译器错误)的绝佳位置。你会在这个论坛上找到一些作者的帖子，尤其是在写这本书的时候。论坛可在线访问 https://software.intel.com/en-us/forums/oneapi-data-parallel-c-compiler 。
在线 oneAPI DPC++ 语言参考是一个很好的资源，可以找到类和成员定义的完整列表、编译器选项的详细信息等等。
平台模型
SYCL 或 DPC++ 编译器被设计成和我们曾经使用过的任何其他 C++ 编译器一样的行为和感觉。一个显著的区别是，常规 C++ 编译器只为 CPU 生成代码。在高层次上理解内部工作是值得的，它使编译器能够为主机 CPU 和设备产生代码。
SYCL 和 DPC++ 使用的平台模型(图 13-1 )指定了一个主机来协调和控制在设备上执行的计算工作。第二章描述了如何给设备分配工作，第四章深入探讨了如何给设备编程。第十二章描述了在不同的特性级别使用平台模型。
正如我们在第二章中讨论的，总有一个设备对应着主机，称为主机设备。为设备代码提供这个保证可用的目标，允许在假设至少有一个设备可用的情况下编写设备代码，即使它是主机本身！选择在哪些设备上运行设备代码是在程序控制之下的——作为程序员，如果我们想在特定的设备上执行代码，以及如何执行代码，这完全是我们的选择。
图 13-1
平台模型:可以抽象使用，也可以具体使用
多架构二进制文件
因为我们的目标是用一个单一的源代码来支持一个异构的机器，所以自然希望得到一个单一的可执行文件。
多架构二进制文件(又名胖二进制文件)是一个单一的二进制文件，它已经被扩展为包含我们的异构机器所需的所有编译和中间代码。多架构二进制文件的概念并不新鲜。例如，一些操作系统支持多架构 32 位和 64 位库和可执行文件。多架构二进制代码的行为就像我们习惯的任何其他a.out或A.exe一样——但是它包含了异构机器所需的一切。这有助于为特定设备选择正确的运行代码。正如我们接下来讨论的，fat 二进制中设备代码的一种可能形式是一种中间格式，它将设备指令的最终创建推迟到运行时。
编译模型
SYCL 和 DPC++ 的单源特性允许编译的行为和感觉像普通的 C++ 编译。我们不需要为设备调用额外的通道或者处理绑定设备和主机代码。这些都是由编译器自动为我们处理的。当然，理解正在发生的事情的细节是很重要的，原因有几个。如果我们想要更有效地针对特定的架构，这是很有用的知识，并且了解我们是否需要调试编译过程中发生的故障也很重要。
我们将回顾编译模型，以便我们在需要这些知识的时候得到教育。由于编译模型支持同时在一个主机和潜在的几个设备上执行的代码，编译器、链接器和其他支持工具发出的命令比我们习惯的 C++ 编译更复杂(只针对一种架构)。欢迎来到异类世界！
DPC++ 编译器故意对我们隐藏了这种异构的复杂性，并且“正好可以工作”
DPC++ 编译器可以生成类似于传统 C++ 编译器的特定于目标的可执行代码(提前 (AOT)编译，有时也称为离线内核编译)，或者它可以生成一个中间表示，可以在运行时即时 (JIT)编译到特定的目标。
如果设备目标提前已知(在我们编译程序的时候)，编译器只能提前编译。推迟即时编译提供了更多的灵活性，但是需要编译器和运行时在我们的应用程序运行时执行额外的工作。
DPC++ 编译可以是“提前”的，也可以是“及时”的。
默认情况下，当我们为大多数设备编译代码时，设备代码的输出以中间形式存储。在运行时，系统上的设备处理程序将即时将中间形式编译成在设备上运行的代码，以匹配系统上可用的内容。
我们可以要求编译器提前为特定的设备或设备类别进行编译。这有节省运行时间的优点，但是也有增加编译时间和二进制文件的缺点！提前编译的代码不如实时编译的代码可移植，因为它不能在运行时进行调整。我们可以将两者都包含在我们的二进制文件中，以获得两者的好处。
提前针对特定设备进行编译还有助于我们在构建时检查我们的程序是否应该在该设备上运行。使用即时编译，程序可能会在运行时编译失败(使用第五章中的机制可以发现这一点)。在本章接下来的“调试”部分有一些调试技巧，第五章详细介绍了如何在运行时捕捉这些错误，以避免要求我们的应用程序中止。
图 13-2 说明了从源代码到 fat 二进制(可执行)的 DPC++ 编译过程。我们选择的任何组合都组合成一个胖二进制。当应用程序执行时，运行时使用 fat 二进制文件(这是我们在主机上执行的二进制文件！).有时，我们可能希望在单独的编译中为特定设备编译设备代码。我们希望这样一个单独编译的结果最终被合并到我们的胖二进制文件中。当完全编译(进行完全综合布局布线)时间可能非常长时，这对于 FPGA 开发可能非常有用，并且事实上这是 FPGA 开发的要求，以避免要求在运行时系统上安装综合工具。图 13-3 显示了支持此类需求的捆绑/拆分活动的流程。我们总是可以选择一次编译所有内容，但是在开发过程中，选择分解编译会非常有用。
每个 SYCL 和 DPC++ 编译器都有一个目标相同的编译模型，但是具体的实现细节会有所不同。这里显示的图表是针对 DPC++ 编译器工具链的。
一个特定于 DPC++ 的组件如图 13-2 所示，作为本书中不再提及的集成头生成器。我们甚至不需要知道它是什么或做什么就可以编程。然而，为了满足好奇心，这里有一些信息:集成头文件生成器生成一个头文件，提供关于翻译单元中 SYCL 内核的信息。这包括 SYCL 内核类型的名称如何映射到符号名称，以及关于内核参数及其在相应的 lambda 或 functor 对象中的位置的信息，这些对象是由编译器创建来捕获它们的。integration header 是一种机制，用于通过 C++ lambda/functor 对象实现从主机代码调用内核的便捷方式，这将我们从设置单个参数、按名称解析内核等耗时的任务中解放出来。
图 13-3
编译过程:卸载捆绑器/解捆绑器
图 13-2
编译过程:提前和及时选项
向现有 C++ 程序添加 SYCL
向现有的 C++ 程序添加适当的并行性是使用 SYCL 的第一步。如果一个 C++ 应用程序已经在利用并行执行，这可能是一个意外收获，也可能是一个令人头疼的问题。这是因为我们将应用程序的工作划分为并行执行的方式极大地影响了我们可以用它做什么。当程序员谈论重构一个程序时，他们指的是重新安排程序内的执行和数据流，以使其准备好利用并行性。这是一个复杂的话题，我们只简单地谈一下。关于如何为并行化准备应用，没有一个通用的答案，但是有一些提示值得注意。
当向 C++ 应用程序添加并行性时，一个简单的方法是在程序中找到一个孤立点，在那里并行性的机会最大。我们可以从这里开始修改，然后根据需要继续在其他领域添加并行性。一个复杂的因素是重构(例如，重新安排程序流和重新设计数据结构)可以提高并行性的机会。
一旦我们在程序中找到一个最有可能实现并行的孤立点，我们就需要考虑如何在程序中的这个点上使用 SYCL。这就是本书其余部分所教导的。
概括地说，引入并行性的关键步骤包括
-
并发安全(在传统 CPU 编程中通常称为线程安全):调整所有共享的可变数据(可以改变并被并发共享的数据)以便并发使用
-
引入并发性和/或并行性
-
针对并行性进行调整(最佳扩展，针对吞吐量或延迟进行优化)
首先考虑步骤 1 是很重要的。许多应用程序已经针对并发性进行了重构，但许多还没有。由于 SYCL 是并行性的唯一来源，我们重点关注内核中使用的数据以及可能与主机共享的数据的安全性。如果我们的程序中有其他技术(OpenMP、MPI、TBB 等)。)引入了并行性，这是我们 SYCL 编程的另一个关注点。需要注意的是，在一个程序中使用多种技术是可以的——SYCL 不需要成为一个程序中唯一的并行来源。这本书没有涵盖与其他并行技术混合的高级主题。
排除故障
本节给出了一些适度的调试建议，以缓解调试并行程序所特有的挑战，尤其是针对异构机器的调试。
我们永远不要忘记，当我们的应用程序在主机设备上运行时，我们可以选择调试它们。该调试提示在第二章中被描述为方法#2。因为设备的架构通常包含较少的调试挂钩，所以工具通常可以更精确地探测主机上的代码。在主机上运行 everything 的另一个好处是，许多与同步相关的错误将会消失，包括在主机和设备之间来回移动内存。虽然我们最终需要调试所有这样的错误，但这允许增量调试，因此我们可以在其他错误之前解决一些错误。
运行在主机设备上的调试提示是一个强大的调试工具。
当在主机上运行所有代码时，并行编程错误，特别是数据竞争和死锁，通常更容易被工具检测和消除。令我们懊恼的是，当在主机和设备的组合上运行时，我们将最经常地看到由于这种并行编程错误而导致的程序失败。当这样的问题出现时，记住回退到 host-only 是一个强大的调试工具是非常有用的。幸运的是，SYCL 和 DPC++ 经过精心设计，让我们可以使用这个选项，并且易于访问。
调试提示如果一个程序死锁，检查主机访问器是否被正确销毁。
当我们开始调试时，下面的 DPC++ 编译器选项是一个好主意:
-
-g:输出调试信息。 -
-ferror-limit=1:将 C++ 与 SYCL/DPC++ 等模板库一起使用时保持理智。 -
让编译器强制执行良好的编码，以帮助避免在运行时产生错误的代码来调试。
我们真的不需要为了使用 DPC++ 而陷入修复迂腐警告的困境，所以选择不使用-Wpedantic是可以理解的。
当我们让代码在运行时被及时编译时，我们就可以检查代码了。这高度依赖于我们的编译器所使用的层，因此查看编译器文档以获得建议是一个好主意。
调试内核代码
调试内核代码时，首先在主机设备上运行(如第二章所述)。第二章中设备选择器的代码可以很容易地修改，以接受运行时选项，或编译时选项，在我们调试时将工作重定向到主机设备。
在调试内核代码时，SYCL 定义了一个可以在内核内部使用的 C++ 风格的stream(图 13-4 )。DPC++ 还提供了一个 C 风格printf的实验性实现，它有一些有用的功能，但有一些限制。更多详情请见在线 oneAPI DPC++ 语言参考。
图 13-4
sycl::stream
调试内核代码时，经验鼓励我们将断点放在parallel_for之前或parallel_for,内部，但实际上不要放在parallel_for上。放置在parallel_for上的断点可以多次触发断点，即使在执行下一个操作之后。这个 C++ 调试建议适用于许多模板扩展，如 SYCL 中的模板扩展，其中模板调用上的断点在被编译器扩展时会转化为一组复杂的断点。可能有一些实现可以缓解这种情况，但这里的关键点是，我们可以通过不在parallel_for本身上精确设置断点来避免所有实现上的一些混淆。
调试运行时故障
当在编译时发生运行时错误时，我们要么是在处理编译器/运行时错误，要么是我们意外地编写了无意义的程序，直到它在运行时出错并产生难以理解的运行时错误消息时才被发现。深入这些 bug 可能有点吓人，但即使粗略地看一下，也可能让我们更好地了解导致特定问题的原因。它可能会产生一些额外的知识来指导我们避免这个问题，或者它可能只是帮助我们向编译器团队提交一个简短的错误报告。无论哪种方式，知道一些工具的存在是很重要的。
表明运行时失败的程序输出可能如下所示:
origin>: error: Invalid record (Producer: 'LLVM9.0.0' Reader: 'LLVM 9.0.0')
terminate called after throwing an instance of 'cl::sycl::compile_program_error'
看到这里提到的这个抛出让我们知道我们的宿主程序可以被构造来捕捉这个错误。虽然这可能不能解决我们的问题，但它确实意味着运行时编译器故障不需要中止我们的应用程序。第五章深入探讨这个话题。
当我们看到一个运行时故障并且很难快速调试它时，简单地尝试使用提前编译进行重建是值得的。如果我们的目标设备有提前编译选项，这可能是一件容易尝试的事情，可能会产生更容易理解的诊断。如果我们的错误可以在编译时而不是在 JIT 或运行时被看到，通常会在来自编译器的错误消息中发现更多有用的信息，而不是我们通常在 JIT 或运行时看到的少量错误信息。具体选项，查看在线 oneAPI DPC++ 文档进行提前编译。
当我们的 SYCL 程序运行在 OpenCL 运行时之上并使用 OpenCL 后端时，我们可以使用 OpenCL 拦截层运行我们的程序:github . com/Intel/OpenCL-Intercept-Layer。这是一个可以检查、记录和修改应用程序(或高级运行时)生成的 OpenCL 命令的工具。它支持很多控件，但是最初设置的好控件是ErrorLogging、BuildLogging，可能还有CallLogging(尽管它会生成很多输出)。使用DumpProgramSPIRV可以进行有用的转储。OpenCL Intercept 层是一个独立的实用程序，不属于任何特定的 OpenCL 实现，因此它可以与许多 SYCL 编译器一起工作。
对于采用英特尔 GPU 的 Linux 系统上的可疑编译器问题，我们可以转储英特尔图形编译器的中间编译器输出。我们通过将环境变量IGC_ShaderDumpEnable设置为 1(对于某些输出)或者将环境变量IGC_ShaderDumpEnableAll设置为 1(对于大量输出)来实现这一点。倾销的产品进入/tmp/IntelIGC。这种技术可能不适用于所有的图形驱动程序，但值得一试，看看它是否适用于我们的系统。
图 13-5 列出了编译器或运行时支持的这些和一些额外的环境变量(在 Windows 和 Linux 上支持),以帮助高级调试。这些是依赖于 DPC++ 实现的高级调试选项，用于检查和控制编译模型。本书没有讨论或利用它们。在线 oneAPI DPC++ 语言参考是了解更多信息的好地方。
图 13-5
DPC++ 高级调试选项
这些选项在本书中没有详细描述，但是在这里提到它们是为了根据需要打开高级调试的通道。这些选项可能让我们深入了解如何解决问题或错误。有可能我们的源代码无意中触发了一个问题，这个问题可以通过更正源代码来解决。否则，使用这些选项是为了对编译器本身进行非常高级的调试。因此，他们更多地与编译器开发人员联系在一起，而不是编译器的用户。一些高级用户发现这些选项很有用；因此，它们在这里被提及，在本书中不再提及。为了更深入地挖掘，GitHub for DPC++ 在llvm/sycl/doc/environment variables 下有一个针对所有环境变量的文档。md 。
调试技巧当其他选项用尽，我们需要调试一个运行时问题时，我们会寻找一些转储工具，这些工具可能会给我们一些提示。
初始化数据和访问内核输出
在这一节中，我们将深入探讨一个让 SYCL 新用户感到困惑的话题，这也是我们作为 SYCL 开发新手遇到的最常见的错误。
简而言之，当我们从主机内存分配(例如，数组或向量)创建缓冲区时，我们不能直接访问主机分配，直到缓冲区被销毁。在缓冲区的整个生命周期内，缓冲区拥有在构造时传递给它的任何主机分配。很少使用做的机制让我们在缓冲区仍然存在时访问主机分配(例如，缓冲区互斥)，但这些高级功能对这里描述的早期错误没有帮助。
如果我们从一个主机内存分配中构造一个缓冲区，在缓冲区被销毁之前，我们不能直接访问主机内存分配！当缓冲区处于活动状态时，它拥有分配。
当缓冲区仍然拥有主机分配时，主机程序访问该分配时会出现一个常见的错误。一旦发生这种情况，一切都完了，因为我们不知道缓冲区使用分配的目的是什么。如果数据不正确，不要感到惊讶——我们试图从中读取输出的内核可能还没有开始运行！如第 3 和 8 章所述，SYCL 是围绕异步任务图机制构建的。在我们尝试使用来自任务图操作的输出数据之前，我们需要确保我们已经到达了代码中的同步点，在那里图形已经执行并使数据对主机可用。缓冲区销毁和主机访问器的创建都是导致这种同步的操作。
图 13-6 显示了我们经常编写的一种常见代码模式，其中我们通过关闭定义缓冲区的块范围来销毁缓冲区。通过使缓冲区超出范围并被销毁，我们可以通过传递给缓冲区构造器的原始主机分配安全地读取内核结果。
图 13-6
通用模式—从主机分配创建缓冲区
如图 13-6 所示，将缓冲区与现有主机内存关联有两个常见原因:
-
简化缓冲区中数据的初始化。我们可以从我们(或应用程序的另一部分)已经初始化的主机内存中构造缓冲区。
-
减少键入的字符，因为用'
}'结束作用域比创建缓冲区的host_accessor更简洁(尽管更容易出错)。
如果我们使用主机分配来转储或验证内核的输出值，我们需要将缓冲区分配放入块范围(或其他范围)，以便我们可以控制它何时被销毁。然后，在我们访问主机分配以获得内核输出之前，我们必须确保缓冲区被销毁。图 13-6 显示这是正确完成的，而图 13-7 显示了一个常见的错误，即当缓冲区仍然存在时，输出被访问。
图 13-7
常见错误:在缓冲区生存期内直接从主机分配中读取数据
高级用户可能更喜欢使用缓冲区销毁将结果数据从内核返回到主机内存分配中。但是对于大多数用户，尤其是新开发人员，建议使用限定了作用域的主机访问器。
更喜欢使用主机访问器而不是缓冲区的作用域，尤其是在入门时。
为了避免这些错误，我们建议在开始使用 SYCL 和 DPC++ 时使用主机访问器而不是缓冲区范围。主机访问器提供从主机到缓冲区的访问，一旦它们的构造器已经完成运行，我们保证任何先前对缓冲区的写入(例如，来自在host_accessor被创建之前提交的内核)已经执行并且是可见的。本书混合使用了这两种风格(即，主机访问器和传递给缓冲区构造器的主机分配),以使读者熟悉这两种风格。在开始使用时，使用主机访问器往往不容易出错。图 13-8 展示了如何使用主机访问器从内核中读取输出，而不需要首先破坏缓冲区。
图 13-8
建议:使用主机访问器读取内核结果
只要缓冲区是活动的，就可以使用主机访问器，比如在典型缓冲区生命周期的两端——用于缓冲区内容的初始化和从内核读取结果。图 13-9 显示了这种模式的一个例子。
图 13-9
建议:使用主机访问器进行缓冲区初始化和结果读取
要提到最后一个细节是，主机访问器有时会在应用程序中引起相反的错误，因为它们也有生存期。当一个缓冲区的host_accessor处于活动状态时，运行时将不允许任何设备使用该缓冲区！运行时不分析我们的宿主程序来确定它们何时可能访问宿主访问器，所以它知道宿主程序已经完成访问缓冲区的唯一方法是运行host_accessor析构函数。如图 13-10 所示，如果我们的主机程序正在等待一些内核运行(例如queue::wait()或获取另一个主机访问器),并且如果 DPC++ 运行时正在等待我们的早期主机访问器被销毁，然后才能运行使用缓冲区的内核，这可能会导致应用程序看起来挂起。
图 13-10
Bug(挂！)来自host_accessors的不当使用
使用主机访问器时，请确保在内核或其他主机访问器不再需要解锁缓冲区时销毁它们。
多个翻译单元
当我们想要调用内核中定义在不同翻译单元中的函数时，这些函数需要用SYCL_EXTERNAL标记。如果没有这个属性，编译器将只编译一个在设备代码之外使用的函数(从设备代码内部调用这个外部函数是非法的)。
如果我们在同一个翻译单元中定义函数，那么对于SYCL_EXTERNAL函数有一些限制是不适用的:
-
SYCL_EXTERNAL只能用于函数。 -
SYCL_EXTERNAL函数不能使用原始指针作为参数或返回类型。必须改用显式指针类。 -
SYCL_EXTERNAL函数不能调用parallel_for_work_item方法。 -
SYCL_EXTERNAL不能从parallel_for_work_group范围内调用函数。
如果我们试图编译一个内核，它调用的函数不在同一个翻译单元内，也没有用SYCL_EXTERNAL声明，那么我们可能会遇到类似如下的编译错误
error: SYCL kernel cannot call an undefined function without SYCL_EXTERNAL attribute
如果函数本身在没有SYCL_EXTERNAL属性的情况下被编译，我们可能会看到链接或运行时失败，比如
terminate called after throwing an instance of 'cl::sycl::compile_program_error'
...error: undefined reference to ...
DPC++ 支持SYCL_EXTERNAL. SYCL 不要求编译器支持SYCL_EXTERNAL；一般来说，这是一个可选的功能。
多个翻译单元对性能的影响
编译模型的一个含义(见本章前面)是，如果我们将设备代码分散到多个翻译单元中，这可能会比我们的设备代码位于同一位置时触发更多的即时编译调用。这高度依赖于实现，并且随着实现的成熟会随着时间的推移而变化。
这种对性能的影响很小，在我们的大多数开发工作中可以忽略，但是当我们进行微调以最大限度地提高代码性能时，我们可以考虑两件事情来减轻这些影响:(1)将设备代码组合在同一个翻译单元中，以及(2)使用提前编译来完全避免即时编译的影响。由于这两者都需要我们付出一些努力，所以我们只有在完成开发并试图充分利用应用程序的性能时才会这样做。当我们求助于这种详细的调优时，有必要测试这些变化，以观察它们对我们正在使用的 SYCL 实现的影响。
当匿名的兰姆达需要名字的时候
SYCL 提供了指定定义为 lambdas 的名称，以备工具需要和用于调试目的(例如，根据用户定义的名称启用显示)。在本书的大部分内容中，匿名 lambda 被用于内核，因为使用 DPC++ 时不需要名字(除了编译选项的传递，如第十章中关于 lambda 命名的讨论所述)。从 SYCL 2020 暂定开始，它们也是可选的。
当我们有在一个代码库上混合来自多个供应商的 SYCL 工具的高级需求时，工具可能要求我们命名为 lambdas。这是通过在使用 lambda 的 SYCL 动作构造中添加一个<class uniquename>(例如parallel_for)来实现的。这种命名允许来自多个厂商的工具在一次编译中以一种定义的方式进行交互，并且还可以通过显示我们在调试工具和层中定义的内核名称来提供帮助。
从 CUDA 迁移到 SYCL
将 CUDA 代码迁移到 SYCL 或 DPC++ 在本书中没有详细介绍。有一些工具和资源可以探索如何做到这一点。移植 CUDA 代码相对简单，因为它是一种基于内核的并行方法。一旦用 SYCL 或 DPC++ 编写，这个新程序就能针对比 CUDA 单独支持的更多的设备。新增强的程序仍然可以针对 NVIDIA GPU，使用支持 NVIDIA GPU 的 SYCL 编译器。
迁移到 SYCL 打开了 SYCL 支持的设备多样性的大门，这远远超出了 GPU。
当使用 DPC++ 兼容性工具时，--report-type= value选项提供了关于移植代码的非常有用的统计信息。这本书的一位评论家称之为“英特尔dpct提供的一面美丽的旗帜。”根据项目的源代码组织，在移植 CUDA 代码时,--in-root选项可以证明非常有用。
要了解关于 CUDA 迁移的更多信息，有两个资源是很好的起点:
-
英特尔的 DPC++ 兼容工具将 CUDA 应用转化为 DPC++ 代码(tinyurl . com/cudatodcppp)。
-
Codeplay 教程“从 CUDA 迁移到 SYCL”(tinyurl . com/codeplayCUDAtoSY CL)。
摘要
今天的流行文化经常把小费称为生活窍门。不幸的是，编程文化经常赋予黑客一个负面的含义，所以作者避免将这一章命名为“SYCL 黑客”毫无疑问，本章只是触及了使用 SYCL 和 DPC++ 的一些实用技巧。更多的技巧可以在在线论坛上分享，我们一起学习如何用 DPC++ 充分利用 SYCL。
开放存取本章根据知识共享署名 4.0 国际许可证(http://Creative Commons . org/licenses/by/4.0/)的条款获得许可，该许可证允许以任何媒体或格式使用、共享、改编、分发和复制，只要您适当注明原作者和来源，提供知识共享许可证的链接并指明是否进行了更改。
本章中的图像或其他第三方材料包含在本章的知识共享许可中，除非在材料的信用额度中另有说明。如果材料不包括在本章的知识共享许可中，并且您的预期使用不被法律法规允许或超出了允许的使用范围，您将需要直接从版权所有者处获得许可。
十四、常见的并行模式
当我们处于程序员的最佳状态时，我们识别工作中的模式，并应用被时间证明是最佳解决方案的技术。并行编程也不例外，如果不研究在这个领域中已经被证明有用的模式，那将是一个严重的错误。考虑大数据应用采用的 MapReduce 框架；他们的成功很大程度上源于基于两个简单而有效的并行模式——map和 reduce 。
并行编程中有许多常见的模式，它们一次又一次地出现，与我们使用的编程语言无关。这些模式是通用的，可以在任何并行级别(例如子组、工作组、完整设备)和任何设备(例如 CPU、GPU、FPGAs)上使用。然而，模式的某些属性(比如它们的可伸缩性)可能会影响它们对不同设备的适用性。在某些情况下，使应用程序适应新设备可能只需要选择适当的参数或微调模式的实现；在其他情况下，我们可以通过选择完全不同的模式来提高性能。
理解如何、何时以及在哪里使用这些常见的并行模式是提高我们对 DPC++(以及一般的并行编程)的熟练程度的关键部分。对于那些已经有并行编程经验的人来说，了解这些模式在 DPC++ 中是如何表达的，是一种快速提高和熟悉该语言功能的方法。
本章旨在回答以下问题:
-
理解哪些模式是最重要的？
-
这些模式与不同设备的功能有什么关系？
-
哪些模式已经作为 DPC++ 函数和库提供了？
-
如何使用直接编程来实现这些模式？
理解模式
这里讨论的模式是麦克库尔等人在《结构化并行编程》一书中描述的并行模式的一个子集。我们不讨论与并行的类型相关的模式(例如，fork-join，branch-and-bound ),而是集中讨论对编写数据并行内核最有用的算法模式。
我们完全相信，理解并行模式的子集对于成为一名有效的 DPC++ 程序员至关重要。图 14-1 中的表格提供了不同模式的高级概述，包括它们的主要用例、关键属性以及它们的属性如何影响它们与不同硬件设备的关联性。
图 14-1
并行模式及其对不同设备类型的相似性
地图
map 模式是所有模式中最简单的并行模式，具有函数式编程语言经验的读者会很快熟悉它。如图 14-2 所示，通过应用一些函数，一个范围的每个输入元素被独立地映射到一个输出。许多数据并行操作可以表示为映射模式的实例(例如，向量加法)。
图 14-2
地图图案
由于函数的每个应用程序都是完全独立的，map 的表达式通常非常简单，依靠编译器和/或运行时来完成大部分困难的工作。我们应该期望写入 map 模式的内核适用于任何设备，并且这些内核的性能能够很好地与可用的硬件并行性相适应。
然而，在决定将整个应用程序重写为一系列地图内核之前，我们应该仔细考虑！这种开发方法效率很高，并保证应用程序可以移植到各种各样的设备类型，但鼓励我们忽略可能显著提高性能的优化(例如，提高数据重用、融合内核)。
蜡纸
模板图案与贴图图案密切相关。如图 14-3 所示，一个函数应用于一个输入和一组由模板描述的相邻输入，以产生一个输出。模板图案经常出现在许多领域，包括科学/工程应用(例如有限差分代码)和计算机视觉/机器学习应用(例如图像卷积)。
图 14-3
模板图案
当模板模式不在适当位置执行时(即，将输出写入单独的存储位置)，该功能可以独立地应用于每个输入。现实世界中调度模板通常比这更复杂:计算相邻的输出需要相同的数据，多次从内存中加载该数据会降低性能；我们可能希望就地应用模板(即，覆盖原始输入值)以减少应用程序的内存占用。
因此，模板内核对不同设备的适用性高度依赖于模板的属性和输入问题。根据经验法则:
-
小模板可以从 GPU 的暂存存储中受益。
-
大型模板可以受益于(相对)较大的 CPU 缓存。
-
通过在 FPGAs 上实现脉动阵列，对小输入进行操作的小模板可以实现显著的性能增益。
由于模板很容易描述，但实现起来很复杂，因此模板是领域特定语言(DSL)开发中最活跃的领域之一。已经有几个嵌入式 DSL 利用 C++ 的模板元编程功能在编译时生成高性能的模板内核，我们希望这些框架移植到 DPC++ 只是时间问题。
减少
归约是一种常见的并行模式，使用通常为关联和交换的操作符(例如，加法)组合来自内核调用的每个实例的部分结果。缩减的最普遍的例子是计算总和(例如，当计算点积时)或计算最小/最大值(例如，使用最大速度来设置时间步长)。
图 14-4 显示了通过树归约实现的归约模式，这是一种流行的实现方式，需要对一系列 N 输入元素进行 log 2 ( N )组合运算。虽然树归约很常见，但其他实现也是可能的——一般来说，我们不应该假设归约以特定的顺序组合值。
图 14-4
还原模式
在现实生活中，内核很少是令人尴尬的并行，即使它们是并行的，它们也经常与 Reduce 成对出现(如在 MapReduce 框架中)来总结它们的结果。这使得 reductions 成为需要理解的最重要的并行模式之一，也是我们必须能够在任何设备上高效执行的模式。
针对不同设备调整缩减是计算部分结果所花费的时间和组合它们所花费的时间之间的微妙平衡；使用太少的并行会增加计算时间，而使用太多的并行会增加组合时间。
通过使用不同的设备执行计算和组合步骤来提高整体系统利用率可能很有吸引力，但是这种调整工作必须仔细考虑在设备之间移动数据的成本。在实践中，我们发现在数据产生时直接在同一设备上执行缩减通常是最佳方法。因此，使用多个设备来提高归约模式的性能不依赖于任务并行性，而是依赖于另一个级别的数据并行性(即，每个设备对部分输入数据执行归约)。
扫描
扫描模式使用二元关联运算符计算广义前缀和，输出的每个元素代表一个部分结果。如果元素 i 的部分和是范围【0， i 中所有元素的和(即包括 i 的和*),则称一次扫描为包含。如果元素 i 的部分和是范围[0， i ]中所有元素的和(即不包括 i* 的和*，则称扫描为互斥。*
乍一看，扫描似乎是一个固有的串行操作，因为每个输出的值取决于前一个输出的值！虽然 scan 确实比其他模式具有更少的并行机会(因此可扩展性可能更差)，但图 14-5 显示了在相同数据上使用多次扫描来实现并行扫描是可能的。
图 14-5
扫描模式
因为扫描操作中的并行机会有限，所以执行扫描的最佳设备在很大程度上取决于问题大小:较小的问题更适合 CPU，因为只有较大的问题才会包含足以使 GPU 饱和的数据并行度。对于 FPGAs 和其他空间架构来说，问题的大小并不重要，因为扫描自然有助于流水线并行。与缩减的情况一样，一个很好的经验法则是在生成数据的同一设备上执行扫描操作，在优化期间考虑扫描操作适合应用程序的位置和方式通常会比单独优化扫描操作产生更好的结果。
打包和拆包
打包和解包模式与扫描密切相关，通常在扫描功能的基础上实现。我们在这里将它们分开讨论，因为它们实现了可能与前缀和没有明显联系的常见操作的高性能实现(例如，添加到列表中)。
包装
如图 14-6 所示，填充模式基于布尔条件丢弃输入范围的元素，将未被丢弃的元素填充到输出范围的连续位置。该布尔条件可以是预先计算的掩码，或者可以通过对每个输入元素应用某个函数来在线计算。
图 14-6
包装模式
与扫描一样，打包操作具有内在的串行性质。给定要打包/复制的输入元素，计算其在输出范围中的位置需要关于有多少先前元素也被打包/复制到输出中的信息。该信息相当于对驱动包的布尔条件的排他扫描。
解除…的负担
如图 14-7 所示(顾名思义)，解包模式与打包模式相反。输入范围的连续元素被解包到输出范围的非连续元素中，其他元素保持不变。这种模式最明显的用例是解包先前打包的数据，但是它也可以用来填充先前计算产生的数据中的“空隙”。
图 14-7
解包模式
使用内置函数和库
这些模式中有许多可以直接使用 DPC++ 的内置功能或供应商提供的用 DPC++ 编写的库来表达。在真正的大型软件工程项目中，利用这些函数和库是平衡性能、可移植性和生产率的最佳方式。
DPC++ 简化库
DPC++ 提供了一种方便的抽象来描述具有归约语义的变量，而不是要求我们每个人都维护自己的可移植的高性能归约内核库。这种抽象简化了归约核的表达式，并使归约被执行的事实显式化，从而允许实现为设备、数据类型和归约操作的不同组合在不同的归约算法之间进行选择。
图 14-8 中的内核展示了一个使用归约库的例子。注意，内核体不包含任何对归约的引用——我们必须指定的是，内核包含一个归约，它使用plus仿函数组合了sum变量的实例。这为自动生成优化的缩减序列的实现提供了足够的信息。
图 14-8
使用归约库表示为 ND 范围数据并行核的归约
在撰写本文时，归约库只支持具有单个归约变量的内核。DPC++ 的未来版本有望支持同时执行多个归约的内核，方法是在传递给parallel_for的nd_range和仿函数参数之间指定多个归约，并将多个归约器作为内核仿函数的参数。
在内核完成之前，不保证归约的结果会被写回原始变量。除了这个限制之外，访问归约结果的行为与访问 SYCL 中的任何其他变量的行为相同:访问存储在缓冲区中的归约结果需要创建适当的设备或主机访问器，而访问存储在 USM 分配中的归约结果可能需要显式同步和/或内存移动。
DPC++ 归约库不同于其他语言中的归约抽象的一个重要方面是，它限制了我们在内核执行期间对归约变量的访问——我们不能检查归约变量的中间值，并且我们被禁止使用除指定组合函数之外的任何东西来更新归约变量。这些限制防止我们犯难以调试的错误(例如，在试图计算最大值时添加缩减变量)，并确保缩减可以在各种不同的设备上有效地实现。
reduction类
reduction类是我们用来描述内核中的缩减的接口。构造归约对象的唯一方法是使用图 14-9 所示的函数之一。
图 14-9
reduction函数的函数原型
该函数的第一个版本允许我们指定归约变量和用于合并每个工作项贡献的操作符。第二个版本允许我们提供一个与归约操作符相关联的可选标识值——这是对用户定义的归约的一个优化，我们稍后将再次讨论。
注意，reduction函数的返回类型是未指定的，而reduction类本身完全是实现定义的。尽管这对于 C++ 类来说可能有点不寻常，但它允许一个实现使用不同的类(或者一个具有任意数量模板参数的类)来表示不同的归约算法。DPC++ 的未来版本可能会决定重新考虑这种设计，以便使我们能够在特定的执行上下文中显式地请求特定的归约算法。
reducer类
reducer类的一个实例封装了一个归约变量，公开了一个有限的接口，确保我们不能以任何实现认为不安全的方式更新归约变量。图 14-10 中显示了reducer等级的简化定义。像reduction类一样，reducer类的精确定义是实现定义的——缩减器的类型将取决于缩减是如何执行的，为了最大化性能，在编译时知道这一点很重要。然而，允许我们更新归约变量的函数和操作符是定义良好的，并且保证受任何 DPC++ 实现的支持。
图 14-10
reducer类的简化定义
具体来说，每个 reducer 都提供了一个combine()函数，它将部分结果(来自单个工作项)与 reduction 变量的值结合起来。这个组合函数的行为是由实现定义的，但这不是我们在编写内核时需要担心的。根据归约运算符，还需要一个归约运算符来使其他运算符可用；例如，+=运算符是为plus归约而定义的。提供这些额外的运算符只是为了方便程序员并提高可读性；在它们可用的地方，这些操作符具有与直接调用combine()相同的行为。
用户定义的缩减
几个常见的归约算法(例如，树归约)并不看到每个工作项直接更新单个共享变量，而是在私有变量中累积一些部分结果，这些部分结果将在将来的某个时刻被组合。这样的私有变量引入了一个问题:实现应该如何初始化它们？将变量初始化为每个工作项的第一个贡献具有潜在的性能影响，因为需要额外的逻辑来检测和处理未初始化的变量。相反，将变量初始化为归约运算符的标识可以避免性能损失，但只有在标识已知的情况下才有可能。
当归约操作在简单算术类型上并且归约运算符是标准函子(例如，plus)时，DPC++ 实现只能自动确定要使用的正确标识值。对于用户定义的约简(即那些对用户定义的类型进行操作和/或使用用户定义的函子的约简)，我们可以通过直接指定标识值来提高性能。
对用户定义的归约的支持仅限于普通的可复制类型和没有副作用的组合函数，但这足以支持许多现实生活中的用例。例如，图 14-11 中的代码演示了使用用户定义的归约来计算向量中的最小元素及其位置。
图 14-11
使用用户定义的约简，通过 ND-range 核找到最小值的位置
oneAPI DPC++ 库
C++ 标准模板库(STL)包含了几个与本章讨论的并行模式相对应的算法。STL 中的算法通常适用于由成对迭代器指定的序列，并且从 C++17 开始，支持一个执行策略参数，表示它们应该顺序执行还是并行执行。
oneAPI DPC++ 库(oneDPL)利用这一执行策略参数来提供一种高效的并行编程方法，这种方法利用了在幕后用 DPC++ 编写的内核。如果一个应用程序可以单独使用 STL 算法的功能来表达，那么 oneDPL 就可以在不编写任何 DPC++ 内核代码的情况下利用我们系统中的加速器！
图 14-12 中的表格显示了 STL 中可用的算法如何与本章中描述的并行模式相关联，以及如何与传统串行算法(在 C++17 之前可用)相关联。关于如何在 DPC++ 应用中使用这些算法的更详细的解释可以在第十八章中找到。
图 14-12
将并行模式与 C++17 算法库相关联
群组功能
DPC++ 设备代码中对并行模式的支持由单独的组函数库提供。这些组函数利用特定工作项目组(即工作组或子组)的并行性来在有限的范围内实现通用并行算法，并且可以用作构建块来构造其他更复杂的算法。
与 oneDPL 一样，DPC++ 中组函数的语法基于 C++ 中算法库的语法。每个函数的第一个参数接受一个group或sub_group对象来代替执行策略，C++ 算法的任何限制都适用。组功能由指定组中的所有工作项目协作执行，因此必须类似于组屏障来处理——组中的所有工作项目必须在聚合控制流中遇到相同的算法(即，组中的所有工作项目必须类似地遇到或不遇到算法调用)，并且所有工作项目必须提供相同的功能参数，以便确保它们在正在执行的操作上达成一致。
在撰写本文时，reduce、exclusive_scan和inclusive_scan函数仅限于支持原始数据类型和最常见的归约运算符(例如，plus、minimum和maximum)。这对于许多用例来说已经足够了，但是 DPC++ 的未来版本有望将集体支持扩展到用户定义的类型和操作符。
直接编程
尽管我们建议尽可能地利用库，但是我们可以通过查看如何使用“本地”DPC++ 内核实现每个模式来学习很多东西。
本章剩余部分中的内核不应期望达到与高度调优的库相同的性能水平，但有助于更好地理解 DPC++ 的功能，甚至可以作为构建新库功能原型的起点。
USE VENDOR-PROVIDED LIBRARIES!
当供应商提供一个函数的库实现时，使用它比将函数重新实现为内核几乎总是有益的！
地图
由于其简单性，map 模式可以直接实现为一个基本的并行内核。图 14-13 所示的代码显示了这样一个实现，使用 map 模式计算一个范围内每个输入元素的平方根。
图 14-13
在数据并行内核中实现 map 模式
蜡纸
如图 14-14 所示，将模板直接实现为具有多维缓冲区的多维基本数据并行内核，简单易懂。
图 14-14
在数据并行内核中实现模板模式
然而，模板模式的这种表达非常幼稚，不应该期望表现得很好。正如本章前面提到的，众所周知，需要利用局部性(通过空间或时间分块)来避免从内存中重复读取相同的数据。图 14-15 显示了一个使用工作组本地内存的简单空间分块示例。
图 14-15
使用工作组本地内存在 ND-range 内核中实现模板模式
为给定模板选择最佳优化需要编译时对块大小、邻域和模板函数本身进行自省，这需要比这里讨论的更复杂的方法。
减少
通过利用在工作项之间提供同步和通信能力的语言特性(例如，原子操作、工作组和子组功能、子组洗牌)，可以在 DPC++ 中实现归约内核。图 14-16 和 14-17 中的内核显示了两种可能的归约实现:使用基本parallel_for的简单归约和每个工作项的原子操作；还有一个稍微聪明一点的缩减，分别使用 ND-range parallel_for和 work-group reduce函数来利用局部性。我们将在第十九章更详细地回顾这些原子操作。
图 14-17
实现一个表示为 ND-range 核的简单约简
图 14-16
实现表示为数据并行内核的简单约简
有许多其他方式来编写归约内核，并且不同的设备可能会偏好不同的实现，这是由于对原子操作的硬件支持、工作组本地存储器大小、全局存储器大小、快速设备范围屏障的可用性，或者甚至专用归约指令的可用性的差异。在某些架构上，它甚至可能更快(或者是必要的！)使用 log 2 ( N )个单独的内核调用来执行树缩减。
我们强烈建议，只有在 DPC++ reduction 库不支持的情况下，或者当针对特定设备的功能对内核进行微调时，才考虑手动实现 reduction——即使这样，也只有在 100%确定 reduction 库性能不佳之后！
扫描
正如我们在本章前面所看到的，实现并行扫描需要对数据进行多次扫描，并且在每次扫描之间进行同步。由于 DPC++ 不提供同步 ND 范围内所有工作项的机制，因此必须使用多个内核来直接实现设备范围的扫描，这些内核通过全局内存来传递部分结果。
如图 14-18 、 14-19 和 14-20 所示的代码展示了使用几个内核实现的包容性扫描。第一个内核跨工作组分发输入值，在工作组本地内存中计算工作组本地扫描(注意，我们可以使用工作组inclusive_scan函数代替)。第二个内核使用单个工作组计算局部扫描，这次是基于每个块的最终值。第三个内核组合这些中间结果来最终确定前缀和。这三个内核对应着图 14-5 中图的三层。
图 14-20
在 ND 范围内核中实现全局包含扫描的第 3 阶段(最终阶段)
图 14-19
在 ND-range 内核中实现全局包含扫描的第 2 阶段:扫描每个工作组的结果
图 14-18
在 ND-range 内核中实现全局包含扫描的第 1 阶段:跨每个工作组进行计算
图 14-18 和 14-19 非常相似；唯一的区别是范围的大小以及输入和输出值的处理方式。这种模式的实际实现可以使用一个带有不同参数的函数来实现这两个阶段，出于教学原因，这里只将它们作为不同的代码。
打包和拆包
打包和解包也称为收集和分散操作。这些操作处理数据在内存中的排列方式以及我们希望将其呈现给计算资源的方式之间的差异。
包装
由于 pack 依赖于独占扫描，所以实现适用于 ND-range 的所有元素的 pack 也必须通过全局内存并在几个内核入队的过程中进行。但是，有一种常见的包装用例，它不要求将操作应用于 ND 范围的所有元素，即只在特定工作组或子组中的项目上应用包装。
图 14-21 中的片段显示了如何在独占扫描的基础上实现组包操作。
图 14-21
在独占扫描的基础上实现组打包操作
图 14-22 中的代码演示了如何在内核中使用这种打包操作来构建需要一些额外后处理的元素列表(在未来的内核中)。所示的例子基于来自分子动力学模拟的真实内核:分配给粒子 i 的子组中的工作项合作识别在 i 的固定距离内的所有其他粒子，并且只有该“邻居列表”中的粒子将用于计算作用在每个粒子上的力。
图 14-22
使用子组打包操作来构建需要附加后处理的元素列表
请注意，pack 模式不会对元素进行重新排序——打包到输出数组中的元素会按照它们在输入中的顺序出现。pack 的这个属性很重要，它使我们能够使用 pack 功能来实现其他更抽象的并行算法(比如std::copy_if和std::stable_partition)。然而，有其他的并行算法可以在不需要维护顺序的包功能之上实现(例如std::partition)。
解除…的负担
与 pack 一样，我们可以使用 scan 实现 unpack。图 14-23 显示了如何在独占扫描之上实现子组解包操作。
图 14-23
在独占扫描之上实现子组解包操作
图 14-24 中的代码演示了如何使用这样的子组解包操作来改善具有分散控制流的内核中的负载平衡(在这种情况下，计算 Mandelbrot 集)。每个工作项被分配一个单独的像素进行计算，并进行迭代，直到达到收敛或最大迭代次数。然后使用解包操作用新像素替换完成的像素。
图 14-24
使用子组解包操作来改善具有不同控制流的内核的负载平衡
这种方法提高效率(和减少执行时间)的程度高度依赖于应用程序和输入，因为检查完成和执行解包操作都会引入一些开销！因此，在实际的应用程序中成功地使用这种模式将需要基于存在的差异量和正在执行的计算进行一些微调(例如，只有当活动工作项目的数量低于某个阈值时，才引入启发式方法来执行解包操作)。
摘要
本章展示了如何使用 DPC++ 和 SYCL 特性实现一些最常见的并行模式，包括内置函数和库。
SYCL 和 DPC++ 生态系统仍在开发中，随着开发人员从生产级应用程序和库的开发中获得更多的语言经验，我们期望发现这些模式的新的最佳实践。
更多信息
-
《结构化并行编程:高效计算的模式》,作者:迈克尔·麦克库尔、阿奇·罗宾逊和詹姆斯·赖因德斯，2012 年，由摩根·考夫曼出版，ISBN 978-0-124-15993-8
-
英特尔 oneAPI DPC++ 库指南，
https://software.intel.com/en-us/oneapi-dpcpp-library-guide -
算法库，C++ 参考，
https://en.cppreference.com/w/cpp/algorithm
开放存取本章根据知识共享署名 4.0 国际许可证(http://Creative Commons . org/licenses/by/4.0/)的条款获得许可，该许可证允许以任何媒体或格式使用、共享、改编、分发和复制，只要您适当注明原作者和来源，提供知识共享许可证的链接并指明是否进行了更改。
本章中的图像或其他第三方材料包含在本章的知识共享许可中，除非在材料的信用额度中另有说明。如果材料不包括在本章的知识共享许可中，并且您的预期使用不被法律法规允许或超出了允许的使用范围，您将需要直接从版权所有者处获得许可。
十五、GPU 编程
在过去的几十年里，图形处理单元(GPU)已经从能够在屏幕上绘制图像的专用硬件设备发展成为能够执行复杂并行内核的通用设备。如今，几乎每台计算机都在传统 CPU 旁边包含一个 GPU，许多程序可以通过将并行算法的一部分从 CPU 卸载到 GPU 来加速。
在本章中，我们将描述典型的 GPU 如何工作，GPU 软件和硬件如何执行 SYCL 应用程序，以及在为 GPU 编写和优化并行内核时需要记住的技巧和技术。
性能警告
与任何处理器类型一样，GPU 因供应商而异，甚至因产品而异；因此，一种设备的最佳实践可能不是另一种设备的最佳实践。本章中的建议可能会让许多 GPU 受益，无论是现在还是将来，但是…
为了实现特定 GPU 的最佳性能，请始终查阅 GPU 供应商的文档！
本章末尾提供了许多 GPU 供应商的文档链接。
GPU 如何工作
本节描述典型的 GPU 如何工作，以及 GPU 与其他类型的加速器有何不同。
GPU 构建模块
图 15-1 显示了一个非常简化的 GPU，由三个高级构建模块组成:
-
执行资源:GPU 的执行资源是执行计算工作的处理器。不同的 GPU 供应商对其执行资源使用不同的名称，但所有现代 GPU 都由多个可编程处理器组成。处理器可能是异构的并专门用于特定的任务，或者它们可能是同构的并可互换。大多数现代 GPU 的处理器都是同类且可互换的。
-
固定功能:GPU 固定功能是比执行资源更不可编程的硬件单元，专门用于单一任务。当 GPU 用于图形时，图形管道的许多部分(如光栅化或光线跟踪)都是使用固定功能来执行的，以提高能效和性能。当 GPU 用于数据并行计算时，固定函数可以用于诸如工作负荷调度、纹理采样和依赖性跟踪之类的任务。
-
Caches and memory: Like other processor types , GPUs frequently have caches to store data accessed by the execution resources. GPU caches may be implicit , in which case they require no action from the programmer, or may be explicit scratchpad memories, in which case a programmer must purposefully move data into a cache before using it. Many GPUs also have a large pool of memory to provide fast access to data used by the execution resources.
图 15-1
典型的 GPU 构建模块—不符合比例！
更简单的处理器(但更多)
传统上，在执行图形操作时，GPU 会处理大量数据。例如，典型的游戏帧或渲染工作负载涉及数千个顶点，每帧产生数百万个像素。为了保持交互式帧速率，必须尽可能快地处理这些大批量数据。
一个典型的 GPU 设计权衡是从构成执行资源的处理器中删除一些功能，以加速单线程性能，并使用这些节省来构建更多的处理器，如图 15-2 所示。例如，GPU 处理器可能不包括由其他类型的处理器使用的复杂无序执行能力或分支预测逻辑。由于这些权衡，单个数据元素在 GPU 上的处理速度可能比在另一个处理器上慢，但更多的处理器使 GPU 能够快速高效地处理许多数据元素。
图 15-2
GPU 处理器更简单，但数量更多
为了在执行内核时利用这种权衡，给 GPU 足够大的数据元素处理范围是很重要的。为了证明卸载大量数据的重要性，考虑一下我们在本书中一直在开发和修改的矩阵乘法内核。
A REMINDER ABOUT MATRIX MULTIPLICATION
在本书中，矩阵乘法内核用于演示内核的变化或其调度方式如何影响性能。尽管使用本章中描述的技术可以显著提高矩阵乘法的性能，但矩阵乘法是一种非常重要和常见的运算，许多硬件(GPU、CPU、FPGA、DSP 等)都无法使用它。)供应商已经实现了包括矩阵乘法在内的许多例程的高度优化版本。这些厂商投入了大量的时间和精力来实现和验证特定设备的功能，并且在某些情况下可能使用在标准内核中难以或不可能使用的功能或技术。
USE VENDOR-PROVIDED LIBRARIES!
当供应商提供一个函数的库实现时，使用它比将函数重新实现为内核几乎总是有益的！对于矩阵乘法，人们可以将 oneMKL 作为英特尔 oneAPI 工具包的一部分，来寻找适合 DPC++ 程序员的解决方案。
通过将矩阵乘法内核作为单个任务提交到队列中，可以在 GPU 上轻松执行矩阵乘法内核。这个矩阵乘法内核的主体看起来就像一个在主机 CPU 上执行的函数，如图 15-3 所示。
图 15-3
单任务矩阵乘法看起来很像 CPU 主机代码
如果我们尝试在 CPU 上执行这个内核，它可能会执行得很好——不是很好，因为它不会利用 CPU 的任何并行能力，但对于小矩阵大小来说可能足够好了。如图 15-4 所示，如果我们试图在一个 GPU 上执行这个内核，它的性能可能会非常差，因为单个任务将只使用一个 GPU 处理器。
图 15-4
GPU 上的单个任务内核会导致许多执行资源闲置
表达平行
为了提高这个内核对于 CPU 和 GPU 的性能，我们可以通过将一个循环转换为一个parallel_for来提交一系列数据元素进行并行处理。对于矩阵乘法内核，我们可以选择提交代表两个最外层循环之一的数据元素范围。在图 15-5 中，我们选择并行处理结果矩阵的行。
图 15-5
有点并行的矩阵乘法
CHOOSING HOW TO PARALLELIZE
选择哪个维度进行并行化是针对 GPU 和其他设备类型调整应用的一种非常重要的方式。本章的后续部分将描述为什么在一个维度上进行并行化可能比在不同维度上进行并行化性能更好的一些原因。
尽管有些并行的内核与单任务内核非常相似，但它应该在 CPU 上运行得更好，在 GPU 上运行得更好。如图 15-6 所示，parallel_for使代表结果矩阵行的工作项能够在多个处理器资源上并行处理，因此所有执行资源都保持忙碌。
图 15-6
有些并行的内核会占用更多的处理器资源
请注意，没有指定行被分区和分配到不同处理器资源的确切方式，这为实现提供了选择如何在设备上最好地执行内核的灵活性。例如，实现可以选择在同一处理器上执行连续的行，而不是在一个处理器上执行单独的行，以获得局部性好处。
表达更多并行性
通过选择并行处理两个外部循环，我们可以进一步并行化矩阵乘法内核。因为parallel_for可以表达多达三维的平行循环，这是简单明了的，如图 15-7 所示。在图 15-7 中，注意传递给parallel_for的范围和表示并行执行空间中索引的项目现在都是二维的。
图 15-7
甚至更多的并行矩阵乘法
当在 GPU 上运行时，展示额外的并行性可能会提高矩阵乘法内核的性能。即使当矩阵行数超过 GPU 处理器数时，这种情况也可能发生。接下来的几节描述了出现这种情况的可能原因。
简化的控制逻辑(SIMD 指令)
许多 GPU 处理器通过利用大多数数据元素倾向于采用相同的控制流路径通过内核来优化控制逻辑。例如，在矩阵乘法内核中，由于循环边界不变，每个数据元素执行最内层循环的次数相同。
当数据元素采用相同的控制流路径通过内核时，处理器可以通过在多个数据元素之间共享控制逻辑并将它们作为一组来执行来降低管理指令流的成本。做到这一点的一种方法是实现一个单指令、多数据或 SIMD 指令集，其中多个数据元素由一个单指令同时处理。
THREADS VS. INSTRUCTION STREAMS
在许多并行编程环境和 GPU 文献中，术语“线程”用来表示“指令流”在这些环境中,“线程”不同于传统的操作系统线程，并且通常更加轻量级。然而，情况并不总是这样，在某些情况下,“线程”被用来描述完全不同的东西。
由于术语“线程”被过度使用并且容易被误解，本章使用术语“指令流”来代替。
图 15-8
四宽 SIMD 处理器:四个 alu 共享提取/解码逻辑
一条指令同时处理的数据元素的数量有时被称为该指令或执行该指令的处理器的 SIMD 宽度。在图 15-8 中，四个 alu 共享相同的控制逻辑，因此这可以被描述为一个四宽 SIMD 处理器。
GPU 处理器并不是唯一实现 SIMD 指令集的处理器。其他处理器类型也实现 SIMD 指令集，以提高处理大型数据集时的效率。GPU 处理器与其他处理器类型的主要区别在于，GPU 处理器依靠并行执行多个数据元素来实现良好的性能，并且 GPU 处理器可能比其他处理器类型支持更宽的 SIMD 宽度。例如，GPU 处理器支持 16、32 或更多数据元素的 SIMD 宽度并不少见。
PROGRAMMING MODELS: SPMD AND SIMD
虽然 GPU 处理器实现了不同宽度的 SIMD 指令集，但这通常是一个实现细节，对于在 GPU 处理器上执行数据并行内核的应用程序是透明的。这是因为许多 GPU 编译器和运行时 API 实现了单程序、多数据或 SPMD 编程模型，其中 GPU 编译器和运行时 API 确定最有效的一组数据元素，以便用 SIMD 指令流进行处理，而不是显式表达 SIMD 指令。第九章的“子组”部分探讨了数据元素分组对应用程序可见的情况。
在图 15-9 中，我们扩大了每个执行资源以支持四宽 SIMD，允许我们并行处理四倍多的矩阵行。
图 15-9
在 SIMD 处理器上执行某种程度上并行的内核
使用并行处理多个数据元素的 SIMD 指令是图 15-5 和 15-7 中的并行矩阵乘法内核的性能能够超越处理器数量的方式之一。通过在同一处理器上执行连续的数据元素，SIMD 指令的使用在许多情况下还提供了自然的局部性优势，包括矩阵乘法。
内核受益于处理器间的并行和处理器内的并行！
预测和掩蔽
只要所有数据元素通过内核中的条件代码采用相同的路径，在多个数据元素之间共享指令流就能很好地工作。当数据元素通过条件代码采取不同的路径时，控制流被称为分叉。当控制流在 SIMD 指令流中分叉时，通常两条控制流路径都被执行，一些通道被屏蔽或者被断言。这确保了正确的行为，但是正确性是以性能为代价的，因为被屏蔽的通道不执行有用的工作。
为了展示预测和屏蔽是如何工作的，考虑图 15-10 中的内核，它将每个具有“奇数”索引的数据元素乘以 2，并将每个具有“偶数”索引的数据元素递增 1。
图 15-10
具有发散控制流的内核
假设我们在图 15-8 所示的四宽 SIMD 处理器上执行这个内核，我们在一个 SIMD 指令流中执行前四个数据元素，在另一个 SIMD 指令流中执行接下来的四个数据元素，依此类推。图 15-11 显示了通道可能被屏蔽和执行可能被预测的方法之一，以正确执行这个具有不同控制流的内核。
图 15-11
发散核的可能通道掩码
SIMD 效率
SIMD 效率衡量 SIMD 指令流与同等标量指令流相比表现如何。在图 15-11 中，由于控制流将通道划分为两个相等的组，因此在分叉控制流中的每个指令都以一半的效率执行。在最坏的情况下，对于高度分散的内核，效率可能会因处理器的 SIMD 宽度而降低。
实现 SIMD 指令集的所有处理器都将遭受影响 SIMD 效率的发散惩罚，但是因为 GPU 处理器通常比其他处理器类型支持更宽的 SIMD 宽度，所以当优化 GPU 的内核时，重构算法以最小化发散控制流并最大化收敛执行可能特别有益。这并不总是可能的，但是作为一个例子，选择沿着一个执行更集中的维度进行并行化可能比沿着一个执行高度分散的不同维度进行并行化性能更好。
SIMD 效率和项目组
到目前为止，本章中的所有内核都是基本的数据并行内核，没有指定执行范围内的任何项目分组，这为设备选择最佳分组提供了实现自由。例如，具有较宽 SIMD 宽度的设备可能偏好较大的分组，但是具有较窄 SIMD 宽度的设备可能适合较小的分组。
当一个内核是具有显式工作项分组的 ND 范围内核时，应该注意选择最大化 SIMD 效率的 ND 范围工作组大小。当一个工作组的大小不能被处理器的 SIMD 宽度整除时，工作组的一部分可能会在整个内核运行期间禁用通道。内核preferred_work_group_size_multiple查询可以用来选择有效的工作组规模。有关如何查询设备属性的更多信息，请参阅第十二章。
选择由单个工作项组成的工作组规模可能会执行得很差，因为许多 GPU 会通过屏蔽除一个通道之外的所有 SIMD 通道来实现单个工作项工作组。例如，图 15-12 中的内核可能会比图 15-5 中非常相似的内核性能差得多，尽管两者之间唯一显著的区别是从基本的数据并行内核转变为低效的单工作项 ND-range 内核(nd_range<1>{M, 1})。
图 15-12
低效的单项、有点并行的矩阵乘法
切换工作以隐藏延迟
许多 GPU 实现了另一种技术来简化控制逻辑，最大化执行资源，并提高性能:许多 GPU 允许多个指令流同时驻留在处理器上，而不是在处理器上执行单个指令流。
让多个指令流驻留在一个处理器上是有益的，因为它给每个处理器一个执行工作的选择。如果一个指令流正在执行长等待时间的操作，例如从内存中读取，处理器可以切换到另一个准备运行的指令流，而不是等待操作完成。有了足够的指令流，到处理器切换回原始指令流时，长等待时间操作可能已经完成，而根本不需要处理器等待。
图 15-13 显示了处理器如何使用多个同步指令流来隐藏延迟并提高性能。尽管第一个指令流与多个指令流一起执行的时间稍长，但通过切换到其他指令流，处理器能够找到准备好执行的工作，并且永远不需要空闲地等待长时间的操作完成。
图 15-13
切换指令流以隐藏延迟
GPU 剖析工具可以使用诸如占用率之类的术语来描述 GPU 处理器当前正在执行的指令流的数量与指令流的理论总数。
低占用率并不一定意味着低性能，因为少量的指令流可能会使处理器忙碌。同样，高占用率并不一定意味着高性能，因为如果所有指令流都执行低效、长等待时间的操作，GPU 处理器可能仍然需要等待。在其他条件相同的情况下，增加占用率可以最大限度地提高 GPU 处理器隐藏延迟的能力，通常会提高性能。增加占用率是图 15-7 中使用更多并行内核可以提高性能的另一个原因。
这种在多个指令流之间切换以隐藏延迟的技术特别适合 GPU 和数据并行处理。回想一下图 15-2 中，GPU 处理器通常比其他类型的处理器简单，因此缺乏复杂的延迟隐藏特性。这使得 GPU 处理器更容易受到延迟问题的影响，但由于数据并行编程涉及处理大量数据，GPU 处理器通常有大量的指令流要执行！
将内核卸载到 GPU
本节描述应用程序、SYCL 运行时库和 GPU 软件驱动程序如何协同工作，在 GPU 硬件上卸载内核。图 15-14 中的图表显示了具有这些抽象层的典型软件栈。在许多情况下，这些层的存在对应用程序是透明的，但在调试或分析我们的应用程序时，理解并考虑它们是很重要的。
图 15-14
将并行内核卸载到 GPU(简化)
SYCL 运行时库
SYCL 运行时库是 SYCL 应用程序与之交互的主要软件库。运行时库负责实现queues、buffers、accessors等类以及这些类的成员函数。运行时库的一部分可能在头文件中，因此直接编译成应用程序可执行文件。运行时库的其他部分是作为库函数实现的，它们作为应用程序构建过程的一部分与应用程序可执行文件相链接。运行时库通常不是特定于设备的，同一个运行时库可以协调卸载到 CPU、GPU、FPGAs 或其他设备。
GPU 软件驱动程序
虽然从理论上讲，SYCL 运行时库可以直接卸载到 GPU，但实际上，大多数 SYCL 运行时库都与 GPU 软件驱动程序接口，以向 GPU 提交工作。
GPU 软件驱动程序通常是 API 的实现，如 OpenCL、Level Zero 或 CUDA。大多数 GPU 软件驱动程序都是在 SYCL 运行时调用的用户模式驱动程序库中实现的，用户模式驱动程序可能会调用操作系统或内核模式驱动程序来执行系统级任务，如分配内存或向设备提交工作。用户模式驱动程序也可以调用其他用户模式库；例如，GPU 驱动程序可以调用 GPU 编译器将内核从中间表示即时编译成 GPU ISA(指令集架构)。这些软件模块以及它们之间的交互如图 15-15 所示。
图 15-15
典型的 GPU 软件驱动模块
GPU 硬件
当运行时库或 GPU 软件用户模式驱动程序被明确请求提交工作时，或者当 GPU 软件试探性地确定应该开始工作时，它通常会通过操作系统或内核模式驱动程序调用，以开始在 GPU 上执行工作。在某些情况下，GPU 软件用户模式驱动程序可能会直接向 GPU 提交工作，但这种情况不太常见，可能不是所有设备或操作系统都支持。
当在 GPU 上执行的工作的结果被主机处理器或另一个加速器消耗时，GPU 必须发出信号来指示工作完成。工作完成中涉及的步骤与工作提交的步骤非常相似，只是执行顺序相反:GPU 可能会向操作系统或内核模式驱动程序发出信号，表明它已完成执行，然后用户模式驱动程序将得到通知，最后运行时库将通过 GPU 软件 API 调用观察到工作已完成。
这些步骤中的每一步都会引入延迟，在许多情况下，运行时库和 GPU 软件会在更低的延迟和更高的吞吐量之间进行权衡。例如，更频繁地向 GPU 提交工作可以减少延迟，但是频繁地提交也会由于每次提交的开销而减少吞吐量。收集大量工作会增加延迟，但会将提交开销分摊到更多工作上，并为并行执行带来更多机会。运行时和驱动程序被调整以做出正确的权衡，并且通常做得很好，但是如果我们怀疑驱动程序试探法低效地提交工作，我们应该查阅文档，看看是否有方法使用特定于 API 甚至特定于实现的机制来覆盖默认的驱动程序行为。
当心卸货的成本！
尽管 SYCL 实现和 GPU 供应商在不断创新和优化，以降低将工作卸载到 GPU 的成本，但在 GPU 上开始工作和在主机或其他设备上观察结果时，总会涉及开销。当选择在何处执行算法时，既要考虑在设备上执行算法的好处，也要考虑将算法及其所需的任何数据移动到设备的成本。在某些情况下，使用主机处理器执行并行操作可能是最有效的，或者在 GPU 上低效地执行算法的串行部分，以避免将算法从一个处理器移动到另一个处理器的开销。
从整体上考虑我们算法的性能——在一个设备上低效地执行算法的一部分可能比将执行转移到另一个设备上更有效！
与设备内存之间的传输
在具有专用内存的 GPU 上，要特别注意专用 GPU 内存和主机或其他设备上的内存之间的传输成本。图 15-16 显示了系统中不同内存类型之间的典型内存带宽差异。
图 15-16
设备内存、远程内存和主机内存之间的典型差异
回想一下第三章，GPU 更喜欢在专用设备内存上运行，这可以快一个数量级或更多，而不是在主机内存或另一个设备的内存上运行。尽管访问专用设备内存比访问远程内存或系统内存快得多，但如果数据不在专用设备内存中，则必须对其进行复制或迁移。
只要数据将被频繁访问，将它移动到专用设备内存是有益的，特别是当 GPU 执行资源忙于处理另一个任务时，传输可以异步执行。当数据很少或不可预测地被访问时，即使每次访问的成本更高，也可以节省传输成本并远程或在系统内存中操作数据。第六章描述了控制内存分配的方法，以及将数据复制和预取到专用设备内存的不同技术。这些技术在为 GPU 优化程序执行时非常重要。
GPU 内核最佳实践
前面几节描述了传递给parallel_for的分派参数如何影响内核如何分配给 GPU 处理器资源，以及在 GPU 上执行内核所涉及的软件层和开销。本节描述了内核在 GPU 上执行时的最佳实践。
从广义上讲，内核要么是受内存限制的，这意味着它们的性能受到进出 GPU 上执行资源的数据读写操作的限制，要么是受计算限制的，这意味着它们的性能受到 GPU 上执行资源的限制。为 GPU 和许多其他处理器优化内核的良好开端！—确定我们的内核是内存受限还是计算受限，因为改善内存受限内核的技术通常不会使计算受限内核受益，反之亦然。GPU 供应商通常提供分析工具来帮助做出这一决定。
根据我们的内核是内存受限还是计算受限，需要不同的优化技术！
因为 GPU 倾向于拥有许多处理器和较宽的 SIMD 宽度，内核倾向于更多地受到内存的限制，而不是计算的限制。如果我们不确定从哪里开始，检查我们的内核如何访问内存是一个很好的第一步。
访问全局内存
高效地访问全局内存对于优化应用程序性能至关重要，因为工作项或工作组操作的几乎所有数据都源自全局内存。如果内核对全局内存的操作效率很低，那么它的性能几乎总是很差。尽管 GPU 通常包括专用硬件收集和分散单元，用于读取和写入内存中的任意位置，但对全局内存的访问性能通常由数据访问的位置决定。如果工作组中的一个工作项目正在访问存储器中的一个元素，该元素与工作组中的另一个工作项目所访问的元素相邻，则全局存储器访问性能可能是好的。如果一个工作组中的工作项改为访问步进或随机的内存，则全局内存访问性能可能会更差。一些 GPU 文档将对邻近内存访问的操作描述为合并内存访问。
回想一下，对于我们在图 15-15 ,中有些并行的矩阵乘法内核，我们可以选择是并行处理结果矩阵的一行还是一列，我们选择并行处理结果矩阵的行。这被证明是一个糟糕的选择:如果一个id等于m的工作项与一个 id 等于m-1或m+1的相邻工作项被分组，那么用于访问matrixB的索引对于每个工作项都是相同的，但是用于访问matrixA的索引相差K，这意味着访问是高度跨越的。matrixA的访问模式如图 15-17 所示。
图 15-17
对matrixA的访问速度很快，效率很低
相反，如果我们选择并行处理结果矩阵的列，则访问模式具有更好的局部性。图 15-18 中的内核在结构上与图 15-5 中的内核非常相似，唯一的区别是图 15-18 中的每个工作项操作结果矩阵的一列，而不是结果矩阵的一行。
图 15-18
并行计算结果矩阵的列，而不是行
尽管这两个内核在结构上非常相似，但在许多 GPU 上操作数据列的内核将明显优于操作数据行的内核，这纯粹是因为更高效的内存访问:如果一个 id 等于n的工作项与一个 id 等于n-1或n+1的相邻工作项分组，则每个工作项用于访问matrixA的索引现在是相同的，并且用于访问matrixB的索引是连续的。matrixB的访问模式如图 15-19 所示。
图 15-19
对matrixB的访问是连续且高效的
对连续数据的访问通常非常高效。一个很好的经验法则是，一组工作项对全局内存的访问性能是被访问的 GPU 缓存线数量的函数。如果所有访问都在单个高速缓存行内，则访问将以最高性能执行。如果访问需要两个高速缓存行，比如说通过访问每隔一个元素或者从高速缓存未对齐的地址开始，则访问可能以一半的性能运行。当组中的每个工作项目访问一个唯一的高速缓存行时，比方说对于非常快速或随机的访问，该访问可能以最低的性能运行。
PROFILING KERNEL VARIANTS
对于矩阵乘法，选择沿一维并行显然会导致更高效的内存访问，但对于其他内核，选择可能不那么明显。对于实现最佳性能至关重要的内核，如果不清楚要并行化哪个维度，有时有必要开发和分析沿每个维度并行化的不同内核变体，以了解哪种内核更适合设备和数据集。
访问工作组本地内存
在上一节中，我们描述了对全局内存的访问如何受益于位置，从而最大化缓存性能。正如我们所看到的，在某些情况下，我们可以设计我们的算法来有效地访问内存，例如通过选择在一个维度而不是另一个维度进行并行化。然而，这种技术并不是在所有情况下都可行。本节描述了我们如何使用工作组本地内存来有效地支持更多的内存访问模式。
回想一下第九章，工作组中的工作项可以通过工作组本地内存通信和使用工作组屏障同步来合作解决问题。这种技术对 GPU 尤其有益，因为典型的 GPU 都有专门的硬件来实现屏障和工作组本地内存。不同的 GPU 供应商和不同的产品可能会以不同的方式实现工作组本地内存，但与全局内存相比，工作组本地内存通常有两个好处:本地内存可能支持比访问全局内存更高的带宽和更低的延迟，即使当全局内存访问命中缓存时也是如此，并且本地内存通常被分成不同的内存区域，称为库。只要一个组中的每个工作项目访问不同的存储体，本地存储器访问就以全性能执行。分库访问允许本地内存支持比全局内存多得多的具有最高性能的访问模式。
许多 GPU 厂商会将连续的本地内存地址分配给不同的存储体。这确保了连续的存储器访问总是以全性能运行，而不管起始地址如何。然而，当存储器访问被跨越时，一个组中的一些工作项目可能访问分配给同一存储体的存储器地址。当这种情况发生时，它被认为是一个存储体冲突，并导致串行访问和较低的性能。
为了获得最高的全局内存性能，请尽量减少访问的缓存线数量。
为了获得最大的本地内存性能，请尽量减少存储体冲突的数量！
图 15-20 总结了全局存储器和本地存储器的访问模式和预期性能。假设当ptr指向全局内存时，指针与 GPU 缓存行的大小对齐。从缓存对齐的地址开始连续访问内存，可以获得访问全局内存的最佳性能。访问未对齐的地址可能会降低全局内存性能，因为访问可能需要访问额外的高速缓存行。因为访问未对齐的本地地址不会导致额外的存储体冲突，所以本地存储器性能不会改变。
跨越的情况值得更详细地描述。访问全局内存中的所有其他元素需要访问更多的缓存行，这可能会降低性能。访问本地内存中的所有其他元素可能会导致内存块冲突和性能下降，但前提是内存块的数量能被 2 整除。如果银行的数量是奇数，这种情况下也将满负荷运行。
当访问之间的跨度很大时，每个工作项访问一个唯一的缓存行，从而导致最差的性能。然而对于本地存储器，性能取决于步幅和存储体的数量。当跨距N等于存储体数量时，每次访问都会导致存储体冲突，所有访问都是串行的，导致性能最差。然而，如果步幅M和存储体的数量没有共同因素，则访问将以全性能运行。出于这个原因，许多优化的 GPU 内核将在本地内存中填充数据结构，以选择减少或消除存储体冲突的步长。
图 15-20
不同访问模式、全局和本地内存的可能性能
用子组完全避免本地存储
正如在第九章中所讨论的，子组集合函数是一种在组中的工作项之间交换数据的替代方法。对于许多 GPU 来说，子组代表由单个指令流处理的工作项的集合。在这些情况下，子组中的工作项可以在不使用工作组本地内存的情况下廉价地交换数据和同步。许多性能最好的 GPU 内核使用子组，因此对于昂贵的内核，我们的算法是否可以重新制定以使用子组集合函数是非常值得研究的。
使用小数据类型优化计算
本节描述了在消除或减少内存访问瓶颈后优化内核的技术。要记住的一个非常重要的观点是，GPU 传统上被设计为在屏幕上绘制图片。尽管随着时间的推移，GPU 的纯计算能力已经得到了发展和提高，但在某些领域，它们的图形继承仍然显而易见。
例如，考虑对内核数据类型的支持。许多 GPU 针对 32 位浮点运算进行了高度优化，因为这些运算在图形和游戏中很常见。对于可以处理较低精度的算法，许多 GPU 也支持较低精度的 16 位浮点类型，以精度换取更快的处理速度。相反，尽管许多 GPU 支持 64 位双精度浮点运算，但额外的精度是有代价的，32 位运算的性能通常比 64 位运算好得多。
整数数据类型也是如此，32 位整数数据类型的性能通常比 64 位整数数据类型好，16 位整数的性能甚至可能更好。如果我们可以使用更小的整数来构建我们的计算，我们的内核可能会执行得更快。需要特别注意的一个方面是寻址操作，它通常对 64 位size_t数据类型进行操作，但有时可以重新安排使用 32 位数据类型来执行大多数计算。在某些本地内存情况下，16 位索引就足够了，因为大多数本地内存分配都很小。
优化数学函数
另一个内核可能为了性能而牺牲准确性的领域涉及 SYCL 内置函数。SYCL 包括一组丰富的数学函数，在一系列输入中具有明确的精度。大多数 GPU 本身不支持这些功能，而是使用一长串其他指令来实现它们。虽然数学函数的实现通常针对 GPU 进行了很好的优化，但是如果我们的应用程序可以容忍较低的精度，我们应该考虑一种精度较低、性能较高的不同实现。有关 SYCL 内置函数的更多信息，请参见第十八章。
对于常用的数学函数，SYCL 库包括fast或native函数变量，具有降低的或实现定义的精度要求。对于一些 GPU 来说，这些函数可以比它们精确的对等函数快一个数量级，所以如果它们对算法来说有足够的精度，那么它们是非常值得考虑的。例如，许多图像后处理算法具有定义明确的输入，可以容忍较低的精度，因此非常适合使用fast或native数学函数。
如果一个算法可以容忍较低的精度，我们可以使用较小的数据类型或较低精度的数学函数来提高性能！
专用功能和扩展
为 GPU 优化内核的最后一个考虑是许多 GPU 中常见的专用指令。举个例子，几乎所有的 GPU 都支持在单个时钟内执行两个操作的mad或fma乘加指令。GPU 编译器通常非常擅长识别和优化单个乘法和加法，以使用单个指令来代替，但 SYCL 也包括可以显式调用的mad和fma函数。当然，如果我们希望我们的 GPU 编译器为我们优化乘法和加法，我们应该确保我们不会通过禁用浮点收缩来阻止优化！
其他专用 GPU 指令可能只能通过编译器优化或 SYCL 语言扩展来获得。例如，一些 GPU 支持专门的点积累加指令，编译器会尝试识别并优化这些指令，或者直接调用这些指令。有关如何查询 GPU 实现所支持的扩展的更多信息，请参考第十二章。
摘要
在这一章中，我们首先描述了典型的 GPU 是如何工作的，以及 GPU 与传统 CPU 有何不同。我们描述了 GPU 如何针对大量数据进行优化，方法是用处理器特性来加速额外处理器的单个指令流。
我们描述了 GPU 如何使用宽 SIMD 指令并行处理多个数据元素，以及 GPU 如何使用 SIMD 指令使用预测和屏蔽来执行具有复杂流控制的内核。我们讨论了预测和屏蔽如何降低高度发散的内核的 SIMD 效率和性能，以及如何选择沿一个维度与另一个维度并行化可以减少 SIMD 发散。
由于 GPU 有如此多的处理资源，我们讨论了给予 GPU 足够的工作以保持高占用率是多么重要。我们还描述了 GPU 如何使用指令流来隐藏延迟，这使得让 GPU 执行大量工作变得更加重要。
接下来，我们讨论了将内核卸载到 GPU 所涉及的软件和硬件层，以及卸载的成本。我们讨论了在单个设备上执行算法如何比将执行从一个设备转移到另一个设备更有效。
最后，我们描述了内核在 GPU 上执行时的最佳实践。我们描述了有多少内核从内存限制开始，以及如何有效地访问全局内存和本地内存，或者如何通过使用子组操作来完全避免本地内存。相反，当内核受计算限制时，我们描述了如何通过用较低的精度换取较高的性能或使用定制的 GPU 扩展来访问专门的指令来优化计算。
更多信息
关于 GPU 编程还有很多要学的，这一章只是触及了皮毛！
GPU 规格和白皮书是了解特定 GPU 和 GPU 架构更多信息的绝佳途径。许多 GPU 供应商提供了关于他们的 GPU 以及如何编程的非常详细的信息。
在撰写本文时，可以在软件上找到关于主要 GPU 的相关阅读资料。英特尔。com ， devblogs。英伟达。com ，以及 amd。com 。
有些 GPU 厂商有开源驱动或驱动组件。如果可能的话，检查或单步执行驱动程序代码可能会有所帮助，从而了解应用程序中哪些操作是昂贵的，或者哪里可能存在开销。
本章完全专注于通过缓冲存取器或统一共享内存对全局内存的传统访问，但大多数 GPU 也包括一个固定功能的纹理采样器，可以加速图像操作。有关图像和采样器的更多信息，请参考 SYCL 规范。
开放存取本章根据知识共享署名 4.0 国际许可证(http://Creative Commons . org/licenses/by/4.0/)的条款获得许可，该许可证允许以任何媒体或格式使用、共享、改编、分发和复制，只要您适当注明原作者和来源，提供知识共享许可证的链接并指明是否进行了更改。
本章中的图像或其他第三方材料包含在本章的知识共享许可中，除非在材料的信用额度中另有说明。如果材料不包括在本章的知识共享许可中，并且您的预期使用不被法律法规允许或超出了允许的使用范围，您将需要直接从版权所有者处获得许可。
