C---数据并行教程-三-C++ 数据并行教程（三）十六、CPU 编程内核编程最初流行是作为一种对 GPU 编程的方

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

原文：Data Parallel C++

协议：CC BY-NC-SA 4.0

十六、CPU 编程

img/489625_1_En_16_Figa_HTML.gif

内核编程最初流行是作为一种对 GPU 编程的方式。随着内核编程的推广，理解我们的编程风格如何影响代码到 CPU 的映射是很重要的。

CPU 已经发展了很多年。一个主要的转变发生在 2005 年左右，当时提高时钟速度带来的性能收益减少了。并行性成为最受欢迎的解决方案——CPU 生产商没有提高时钟速度，而是引入了多核芯片。计算机在同时执行多项任务时变得更加有效！

虽然多核成为提高硬件性能的主流途径，但在软件中释放这种优势需要付出巨大的努力。多核处理器要求开发人员提出不同的算法，这样硬件的改进才能引人注目，而这并不总是容易的。我们拥有的内核越多，就越难让它们高效地工作。DPC++ 是解决这些挑战的编程语言之一，它具有许多有助于利用 CPU(和其他架构)上各种形式的并行性的构造。

本章讨论了 CPU 架构的一些细节，CPU 硬件如何执行 DPC++ 应用程序，并提供了为 CPU 平台编写 DPC++ 代码的最佳实践。

性能警告

DPC++ 为并行化我们的应用程序或从头开始开发并行应用程序铺平了道路。当在 CPU 上运行时，程序的应用程序性能在很大程度上取决于以下因素:

内核代码的单次调用和执行的底层性能
在并行内核中运行的程序的百分比及其可伸缩性
CPU 利用率、有效的数据共享、数据局部性和负载平衡
工作项之间的同步和通信量
创建、恢复、管理、挂起、销毁和同步执行工作项的线程所带来的开销，串行到并行或并行到串行转换的数量会使这种开销变得更大
共享内存或错误共享内存导致的内存冲突
共享资源(如内存、写组合缓冲区和内存带宽)的性能限制

此外，与任何处理器类型一样，CPU 可能因供应商而异，甚至因产品的不同而不同。一个 CPU 的最佳实践可能不是另一个 CPU 和配置的最佳实践。

要在 CPU 上实现最佳性能，请尽可能多地了解 CPU 架构的特征！

通用 CPU 的基础知识

多核 CPU 的出现和快速发展推动了共享内存并行计算平台的广泛接受。CPU 提供了笔记本电脑、台式机和服务器级别的并行计算平台，使它们无处不在，并在几乎任何地方展示性能。最常见的 CPU 架构形式是缓存一致的非一致内存访问(cc-NUMA)，其特点是访问时间不完全一致。甚至很多小型双插槽通用 CPU 系统都有这种内存系统。这种架构已经成为主流，因为处理器中内核的数量以及插槽的数量都在不断增加。

在 cc-NUMA CPU 系统中，每个插槽都连接到系统中总内存的一个子集。缓存一致的互连将所有的套接字粘合在一起，并为程序员提供单一的系统视图。这种存储器系统是可扩展的，因为总的存储器带宽随着系统中插座的数量而扩展。互连的好处是应用程序可以透明地访问系统中的所有内存，而不管数据驻留在哪里。然而，这是有代价的:从存储器访问数据和指令的等待时间不再一致(例如，固定的访问等待时间)。相反，延迟取决于数据在系统中的存储位置。在一个好的例子中，数据来自直接连接到代码运行的套接字的内存。在糟糕的情况下，数据必须来自连接到系统中远处插槽的内存，并且由于 cc-NUMA CPU 系统上插槽之间互连的跳数，内存访问的成本可能会增加。

在图 16-1 中，显示了带有 cc-NUMA 存储器的通用 CPU 架构。这是一个简化的系统架构，包含当今通用多插槽系统中的内核和存储器组件。在本章的其余部分，该图将用于说明相应代码示例的映射。

为了实现最佳性能，我们需要确保理解特定系统的 cc-NUMA 配置的特征。例如，英特尔最近推出的服务器就采用了网状互连架构。在这种配置中，内核、高速缓存和内存控制器被组织成行和列。在努力实现系统的最高性能时，了解处理器与内存的连接是至关重要的。

img/489625_1_En_16_Fig1_HTML.png

图 16-1

通用多核 CPU 系统

图 16-1 中的系统有两个插槽，每个插槽有两个内核，每个内核有四个硬件线程。每个内核都有自己的一级(l 1)高速缓存。L1 缓存连接到共享的末级缓存，末级缓存连接到套接字上的内存系统。套接字内的内存访问延迟是一致的，这意味着它是一致的，并且可以准确预测。

这两个套接字通过缓存一致的互连进行连接。内存分布在整个系统中，但是可以从系统中的任何地方透明地访问所有内存。当访问不在进行访问的代码正在运行的套接字中的内存时，内存读写延迟是不一致的，这意味着当从远程套接字访问数据时，它可能会施加更长且不一致的延迟。然而，互连的一个关键方面是一致性。我们不需要担心数据在整个内存系统中变得不一致(这将是一个功能问题)，相反，我们只需要担心我们如何访问分布式内存系统对性能的影响。

CPU 中的硬件线程是执行工具。这些是执行指令流(CPU 术语中的线程)的单元。图 16-1 中的硬件线程从 0 到 15 连续编号，这是用于简化本章示例讨论的符号。除非另有说明，本章中所有提及的 CPU 系统都是指图 16-1 中所示的 cc-NUMA 系统。

SIMD 硬件基础

1996 年，第一个广泛部署的 SIMD(根据 Flynn 的分类，单指令，多数据)指令集是 x86 架构上的 MMX 扩展。此后，许多 SIMD 指令集扩展在英特尔架构和更广泛的行业中得到应用。CPU 核心通过执行指令来执行其工作，并且核心知道如何执行的特定指令由其实现的指令集(例如，x86、x86_64、AltiVec、NEON)和指令集扩展(例如，SSE、AVX、AVX-512)来定义。指令集扩展增加的许多操作都集中在 SIMD 指令上。

SIMD 指令通过使用比正在处理的基本数据单元更大的寄存器和硬件，允许在单个内核上同时执行多个计算。使用 512 位寄存器，我们可以用一条机器指令执行八次 64 位计算。

图 16-2 中所示的例子可以给我们带来八倍的速度提升。实际上，它可能会有所缩减，因为八倍加速的一部分用于消除一个瓶颈并暴露下一个瓶颈，如内存吞吐量。一般来说，使用 SIMD 的性能优势因具体场景而异，在少数情况下，它的性能甚至比更简单的非 SIMD 等效代码还要差。也就是说，当我们知道何时以及如何应用(或者让编译器应用)SIMD 时，在今天的处理器上可以获得相当大的收益。与所有性能优化一样，程序员应该在将一台典型的目标机器投入生产之前测量它的增益。在本章的以下部分中有关于预期性能增益的更多细节。

img/489625_1_En_16_Fig2_HTML.png

图 16-2

CPU 硬件线程中的 SIMD 执行

带有 SIMD 单元的 cc-NUMA CPU 架构构成了多核处理器的基础，该处理器可以从指令级并行开始，以五种不同的方式利用广泛的并行，如图 16-3 所示。

img/489625_1_En_16_Fig3_HTML.png

图 16-3

并行执行指令的五种方式

在图 16-3 中，指令级并行可以通过单个线程内标量指令的无序执行和 SIMD(单指令，多数据)数据并行来实现。线程级并行可以通过在同一内核或不同规模的多个内核上执行多个线程来实现。更具体地说，线程级并行性可以从以下方面得到体现:

现代 CPU 架构允许一个内核同时执行两个或多个线程的指令。
在每个处理器中包含两个或更多大脑的多核架构。操作系统将其每个执行核心视为一个独立的处理器，拥有所有相关的执行资源。
处理器(芯片)级的多重处理，可以通过执行完全独立的代码线程来实现。因此，处理器可以让一个线程从一个应用程序运行，另一个线程从一个操作系统运行，也可以让多个并行线程从一个应用程序中运行。
分布式处理，可以通过在计算机集群上执行由多个线程组成的进程来实现，这些进程通常通过消息传递框架进行通信。

为了充分利用多核处理器资源，编写软件时必须将其工作负载分散到多个内核中。这种方法被称为利用线程级并行或简单的线程。

随着多处理器计算机以及采用超线程(HT)技术和多核技术的处理器变得越来越普遍，将并行处理技术作为提高性能的标准实践变得非常重要。本章的后面几节将介绍 DPC++ 中的编码方法和性能调整技术，它们使我们能够在多核 CPU 上实现最高性能。

像其他并行处理硬件(例如，GPU)一样，为 CPU 提供足够大的数据元素集进行处理非常重要。为了展示利用多级并行处理处理大量数据的重要性，考虑一个简单的 C++ 流三元组程序，如图 16-4 所示。

img/489625_1_En_16_Fig4_HTML.png

图 16-4

流三元组 C++ 循环

A NOTE ABOUT STREAM TRIAD WORKLOAD

流三元组工作负载( www.cs.virginia.edu/stream )是一种重要且流行的基准测试工作负载，CPU 供应商使用它来展示高度调优的性能。我们使用 STREAM Triad 内核来演示并行内核的代码生成，以及通过本章描述的技术来显著提高性能的方法。STREAM Triad 是一个相对简单的工作负载，但足以以一种可理解的方式展示许多优化。

USE VENDOR-PROVIDED LIBRARIES!

当供应商提供一个函数的库实现时，使用它比将函数重新实现为并行内核更有益！

流三元组循环可以在使用单个 CPU 核心进行串行执行的 CPU 上被平凡地执行。一个好的 C++ 编译器会执行循环矢量化，为具有 SIMD 硬件的 CPU 生成 SIMD 代码，以利用指令级 SIMD 并行性。例如，对于支持 AVX-512 的英特尔至强处理器，英特尔 C++ 编译器生成如图 16-5 所示的 SIMD 代码。重要的是，编译器对代码的转换通过在运行时对每个循环迭代做更多的工作(SIMD 宽度和展开迭代),减少了执行时的循环迭代次数！

img/489625_1_En_16_Fig5_HTML.png

图 16-5

流三元组 C++ 循环的 AVX-512 代码

如图 16-5 所示，编译器能够以两种方式利用指令级并行性。首先是通过使用 SIMD 指令，利用指令级数据并行性，其中一条指令可以同时并行处理八个双精度数据元素(每条指令)。第二，基于硬件多路指令调度，编译器应用循环展开来获得这些指令之间没有依赖关系的乱序执行效果。

如果我们尝试在 CPU 上执行这个函数，它可能会运行得很好——虽然不是很好，因为它没有利用 CPU 的任何多核或线程功能，但对于小型阵列来说已经足够好了。但是，如果我们试图在一个 CPU 上用一个大的数组来执行这个函数，它的性能可能会很差，因为单线程将只利用一个 CPU 内核，当它达到该内核的内存带宽饱和时就会出现瓶颈。

利用线程级并行

为了提高 CPU 和 GPU 的 STREAM Triad 内核的性能，我们可以通过将循环转换为parallel_for内核来计算一系列可以并行处理的数据元素。

通过将流三元组内核提交到队列中进行并行执行，可以在 CPU 上轻松地执行流三元组内核。这个 STREAM Triad DPC++ 并行内核的主体看起来与在 CPU 上以串行 C++ 执行的 STREAM Triad 循环的主体完全一样，如图 16-6 所示。

img/489625_1_En_16_Fig6_HTML.png

图 16-6

DPC++ 流三元组parallel_for内核代码

尽管并行内核非常类似于编写为带循环的串行 C++ 的 STREAM Triad 函数，但它在 CPU 上运行得更快，因为parallel_for使阵列的不同元素能够在多个内核上并行处理。如图 16-7 所示，假设我们有一个系统，有一个插槽，四个内核，每个内核有两个超线程；有 1024 个双精度数据元素需要处理；并且在实现中，在每个包含 32 个数据元素的工作组中处理数据。这意味着我们有 8 个线程和 32 个工作组。工作组调度可以按循环顺序进行，即线程 id = 工作组 id mod 8。本质上，每个线程将执行四个工作组。每轮可以并行执行八个工作组。注意，在这种情况下，工作组是由 DPC++ 编译器和运行时隐式形成的一组工作项。

img/489625_1_En_16_Fig7_HTML.png

图 16-7

流三元组并行核的映射

请注意，在 DPC++ 程序中，不需要指定将数据元素分区并分配给不同处理器内核(或超线程)的确切方式。这为 DPC++ 实现提供了选择如何在特定 CPU 上最好地执行并行内核的灵活性。也就是说，一个实现可以为程序员提供某种程度的控制来实现性能调优。

虽然 CPU 可能会强加相对较高的线程上下文切换和同步开销，但是在处理器内核上驻留更多的软件线程是有益的，因为这为每个处理器内核提供了执行工作的选择。如果一个软件线程正在等待另一个线程产生数据，则处理器内核可以切换到准备运行的不同软件线程，而不会使处理器内核空闲。

CHOOSING HOW TO BIND AND SCHEDULE THREADS

选择一个有效的方案来在线程之间划分和调度工作，对于在 CPU 和其他设备类型上调优应用程序非常重要。后续部分将描述一些技术。

线程相似性洞察

线程关联性指定特定线程在其上执行的 CPU 核心。如果线程在内核之间移动，性能会受到影响，例如，如果线程不在同一个内核上执行，如果数据在不同内核之间来回切换，缓存局部性会变得低效。

DPC++ 运行时库支持通过环境变量 DPCPP_CPU_CU_AFFINITY、DPCPP_CPU_PLACES、DPCPP_CPU_NUM_CUS 和 DPCPP_CPU_SCHEDULE 将线程绑定到核心的几种方案，这些变量不是由 SYCL 定义的。

第一个是环境变量DPCPP_CPU_CU_AFFINITY。使用这些环境变量控件进行调优既简单又成本低廉，并且可以对许多应用程序产生巨大影响。该环境变量的描述如图 16-8 所示。

img/489625_1_En_16_Fig8_HTML.png

图 16-8

DPCPP_CPU_CU_AFFINITY环境变量

当环境变量DPCPP_CPU_CU_AFFINITY被指定时，软件线程通过以下公式绑定到超线程:

$sprea: boundHT=\left( tid\;\mathit{\operatorname{mod}}\; numHT\right)+\left( tid\;\mathit{\operatorname{mod}}\; numSocket\right)\times numHT$

$close: boundHT= tid\;\mathit{\operatorname{mod}}\;\left( numSocket\times numHT\right)$

在哪里

tid表示软件线程标识符。
boundHT表示线程tid绑定到的超级线程(逻辑核心)。
numHT表示每个套接字的超线程数量。
numSocket表示系统中插座的数量。

假设我们在双核双插槽超线程系统上运行一个具有八个线程的程序，换句话说，我们有四个内核，总共有八个超线程可供编程。图 16-9 显示了线程如何映射到不同DPCPP_CPU_CU_AFFINITY设置的超线程和内核的示例。

img/489625_1_En_16_Fig9_HTML.png

图 16-9

使用超线程将线程映射到内核

与环境变量DPCPP_CPU_CU_AFFINITY一起，还有其他支持 CPU 性能调整的环境变量:

DPC PP _ CPU _ NUM _ CUS=[n]，设置内核执行使用的线程数量。它的默认值是系统中硬件线程的数量。
DPC PP _ CPU _ PLACES=[sockets|numa_domains|cores|threads]，类似于 OpenMP 5.1 中的OMP_PLACES指定了将要设置亲缘关系的位置。默认设置为cores。
DPC PP _ CPU _ SCHEDULE=[dynamic|affinity|static]，指定了调度工作组的算法。其默认设置为dynamic。
- 动态:启用 TBB auto_partitioner，它通常执行足够的拆分来平衡工作线程之间的负载。
- 关联:启用 TBB affinity_partitioner，这提高了缓存关联，并在将子范围映射到工作线程时使用比例分割。
- 静态:启用 TBB static_partitioner，它尽可能均匀地在工作线程之间分配迭代。

TBB 划分器使用粒度来控制工作划分，默认粒度为 1，表示所有的工作组都可以独立执行。更多信息可在规格中找到。oneapi。com/versions/latest/elements/one TBB/source/algorithms。html #分区器。

缺少线程关联性调优并不一定意味着性能降低。性能通常更多地取决于并行执行的线程总数，而不是线程和数据的关联和绑定程度。使用基准测试应用程序是确定线程关联性是否会影响性能的一种方式。如图 16-1 所示的 DPC++ STREAM Triad 代码，在没有线程关联设置的情况下，开始时性能较低。通过控制亲缘性设置和使用通过环境变量的软件线程的静态调度(对于 Linux 在下面显示了导出)，性能得到了提高:


export DPCPP_CPU_PLACES=numa_domains
export DPCPP_CPU_CU_AFFINITY=close

通过使用numa_domains作为亲缘关系的位置设置，TBB 任务竞技场被绑定到 NUMA 节点或套接字，并且工作被均匀地分布在任务竞技场上。一般情况下，环境变量DPCPP_CPU_PLACES推荐与DPCPP_CPU_CU_AFFINITY一起使用。这些环境变量设置帮助我们在 Skylake 服务器系统上实现了约 30%的性能提升，该系统具有 2 个插槽，每个插槽有 28 个双向超线程内核，运行频率为 2.5 GHz。但是，我们仍然可以做得更好，以进一步提高这台 CPU 的性能。

注意记忆的第一次接触

内存存储在第一次接触(使用)的地方。由于我们示例中的初始化循环没有并行化，它由主机线程串行执行，导致所有内存都与主机线程运行所在的套接字相关联。其他套接字的后续访问将从附加到初始套接字(用于初始化)的内存中访问数据，这显然对性能不利。我们可以通过并行化初始化循环来控制套接字之间的首次接触效应，从而在 STREAM Triad 内核上实现更高的性能，如图 16-10 所示。

img/489625_1_En_16_Fig10_HTML.png

图 16-10

控制首次触摸效果的流三元组并行初始化内核

利用初始化代码中的并行性可以提高内核在 CPU 上运行时的性能。在这种情况下，我们在英特尔至强处理器系统上实现了大约 2 倍的性能提升。

本章最近的章节已经表明，通过开发线程级并行，我们可以有效地利用 CPU 内核和超线程。然而，我们还需要在 CPU 核心硬件中利用 SIMD 矢量级并行，以实现最高性能。

DPC++ 并行内核受益于内核和超线程的线程级并行性！

CPU 上的 SIMD 向量化

虽然一个没有交叉工作项依赖的编写良好的 DPC++ 内核可以在 CPU 上有效地并行运行，但我们也可以将矢量化应用到 DPC++ 内核，以利用 SIMD 硬件，类似于第十五章中描述的 GPU 支持。本质上，CPU 处理器可以通过利用大多数数据元素通常在连续的存储器中并且通过数据并行内核采取相同的控制流路径的事实，使用 SIMD 指令来优化存储器加载、存储和操作。例如，在具有语句a[i] = a[i] + b[i]的内核中，通过在多个数据元素之间共享硬件逻辑并将它们作为一组来执行，每个数据元素以相同的指令流加载、加载、添加和存储来执行，这可以自然地映射到硬件的 SIMD 指令集。具体来说，一条指令可以同时处理多个数据元素。

由一条指令同时处理的数据元素的数量有时被称为指令或执行它的处理器的向量长度(或 SIMD 宽度)。在图 16-11 中，我们的指令流以四路 SIMD 执行方式运行。

img/489625_1_En_16_Fig11_HTML.png

图 16-11

SIMD 执行的指令流

CPU 处理器不是唯一实现 SIMD 指令集的处理器。GPU 等其他处理器在处理大型数据集时会执行 SIMD 指令来提高效率。与其他类型的处理器相比，英特尔至强 CPU 处理器的一个关键区别在于它拥有三个固定大小的 SIMD 寄存器宽度 128 位 XMM、256 位 YMM 和 512 位 ZMM，而不是可变长度的 SIMD 宽度。当我们使用子组或向量类型编写具有 SIMD 并行的 DPC++ 代码时，我们需要注意硬件中的 SIMD 宽度和 SIMD 向量寄存器的数量。

确保 SIMD 执行的合法性

语义上，DPC++ 执行模型确保 SIMD 执行可以应用于任何内核，并且每个工作组(即子组)中的一组工作项目可以使用 SIMD 指令并发执行。一些实现可能改为选择使用 SIMD 指令在内核中执行循环，但是当且仅当所有原始数据依赖性被保留，或者数据依赖性被编译器基于私有化和归约语义解决时，这才是可能的。

使用工作组内的 SIMD 指令，单个 DPC++ 内核执行可以从单个工作项的处理转换为一组工作项。在 ND-range 模型下，增长最快的(单位步长)维由编译器矢量器选择，在其上生成 SIMD 代码。本质上，要在给定 ND-range 的情况下启用矢量化，同一子组中的任何两个工作项之间都不应该有跨工作项依赖，或者编译器需要保留同一子组中的跨工作项向前依赖。

当工作项的内核执行被映射到 CPU 上的线程时，细粒度的同步代价很高，线程上下文切换开销也很高。因此，在为 CPU 编写 DPC++ 内核时，消除工作组内工作项之间的依赖性是一项重要的性能优化。另一个有效的方法是将这种依赖性限制在一个子组中的工作项上，如图 16-12 中的写前读依赖性所示。如果在 SIMD 执行模型下执行子组，则编译器可以将内核中的子组屏障视为空操作，并且在运行时不会产生真正的同步成本。

img/489625_1_En_16_Fig12_HTML.png

图 16-12

使用子群向量化具有前向相关性的循环

内核被矢量化(矢量长度为 8 ，其 SIMD 执行如图 16-13 所示。以组大小(1，8)形成一个工作组，内核内部的循环迭代分布在这些子组工作项上，并以八向 SIMD 并行方式执行。

img/489625_1_En_16_Fig13_HTML.png

图 16-13

具有前向相关性的循环的 SIMD 向量化

在本例中，如果内核中的循环控制了性能，那么允许子组中的 SIMD 向量化将会显著提高性能。

使用并行处理数据元素的 SIMD 指令是让内核性能超越 CPU 内核和超线程数量的一种方式。

SIMD 掩蔽和成本

在实际应用中，我们可能会遇到像if语句这样的条件语句，像a = b > a? a: b这样的条件表达式，迭代次数可变的循环，switch语句，等等。任何有条件的都可能导致标量控制流不执行相同的代码路径，就像在 GPU 上一样(第十五章)，可能导致性能下降。SIMD 掩码是一组值为1或0的位，由内核中的条件语句生成。考虑具有A={1, 2, 3, 4}, B={3, 7, 8, 1}和比较表达式a < b.的示例，比较返回具有四个值{1, 1, 1, 0}的掩码，其可以存储在硬件掩码寄存器中，以指示后面的 SIMD 指令的哪些通道应该执行由比较保护(启用)的代码。

如果内核包含条件代码，它将通过基于与每个数据元素相关联的屏蔽位(SIMD 指令中的 lane)执行的屏蔽指令进行矢量化。每个数据元素的屏蔽位是屏蔽寄存器中的相应位。

使用掩码可能会导致比相应的非掩码代码更低的性能。这可能是由以下原因造成的

每次加载时的额外遮罩混合操作
对目的地的依赖

屏蔽是有代价的，所以只在必要的时候使用它。当内核是 ND 范围内核，并且在执行范围内具有工作项目的显式分组时，在选择 ND 范围工作组大小时应该小心，以便通过最小化屏蔽成本来最大化 SIMD 效率。当一个工作组的大小不能被处理器的 SIMD 宽度整除时，工作组的一部分可以在内核屏蔽的情况下执行。

图 16-14 显示了使用合并屏蔽如何产生对目标寄存器的依赖:

img/489625_1_En_16_Fig14_HTML.png

图 16-14

内核屏蔽的三代屏蔽代码

在没有屏蔽的情况下，处理器每个周期执行两次乘法(vmulps)。
使用合并屏蔽，处理器每四个周期执行两次乘法，因为乘法指令(vmulps)将结果保存在目的寄存器中，如图 16-17 所示。
零屏蔽不依赖于目标寄存器，因此每个周期可以执行两次乘法(vmulps)。

访问缓存对齐的数据比访问非对齐的数据提供更好的性能。在许多情况下，地址在编译时是未知的，或者是已知的但没有对齐。在这些情况下，可以实现对存储器访问的剥离，以通过并行内核中的多版本化技术，使用屏蔽的访问来处理最初的几个元素，直到第一个对齐的地址，然后处理未屏蔽的访问，随后是屏蔽的剩余部分。这种方法增加了代码量，但总体上改善了数据处理。

避免结构数组以提高 SIMD 效率

AOS(结构阵列)结构导致聚集和分散，这既会影响 SIMD 效率，又会为内存访问带来额外的带宽和延迟。硬件聚集-分散机制的存在并没有消除这种转换的需要——聚集-分散访问通常需要比连续加载高得多的带宽和延迟。给定一个struct {float x; float y; float z; float w;} a[4],的 AOS 数据布局，考虑一个在其上运行的内核，如图 16-15 所示。

img/489625_1_En_16_Fig15_HTML.png

图 16-15

SIMD 聚在一个内核里

当编译器沿着一组工作项目*、*对内核进行矢量化处理时，由于需要非单位步长的内存访问，它会导致 SIMD 收集指令的生成。例如，a[0].x、a[1].x、a[2].x和a[3].x的步距是 4，而不是更有效的单位步距 1。

img/489625_1_En_16_Figb_HTML.gif

在内核中，我们通常可以通过消除内存聚集-分散操作来实现更高的 SIMD 效率。一些代码受益于数据布局的变化，这种变化将以结构数组(AOS)表示形式编写的数据结构转换为数组结构(SOA)表示形式，也就是说，每个结构字段都有单独的数组，以在执行 SIMD 矢量化时保持内存访问的连续性。例如，考虑如下所示的struct {float x[4]; float y[4]; float z[4]; float w[4];} a;的 SOA 数据布局:

img/489625_1_En_16_Figc_HTML.gif

如图 16-16 所示，内核可以使用单位步长(连续)向量加载和存储对数据进行操作，即使是在矢量化的情况下！

img/489625_1_En_16_Fig16_HTML.png

图 16-16

内核中 SIMD 单位步幅向量加载

SOA 数据布局有助于在访问跨数组元素的结构的一个字段时防止聚集，并有助于编译器对与工作项相关联的连续数组元素上的内核进行矢量化。请注意，考虑到使用这些数据结构的所有地方，这种 AOS 到 SOA 或 AOSOA 的数据布局转换预计将在程序级别(由我们)完成。仅仅在循环层次上这样做将会涉及到循环前后格式之间代价高昂的转换。然而，我们也可以依靠编译器来执行向量加载和洗牌优化 AOS 数据布局，这需要一些成本。当 SOA(或 AOS)数据布局的成员具有向量类型时，编译器矢量化将根据底层硬件执行水平扩展或垂直扩展，如第十一章所述，以生成最佳代码。

数据类型对 SIMD 效率的影响

每当 C++ 程序员知道数据适合 32 位有符号类型时，他们通常使用整数数据类型，这通常会产生如下代码


int id = get_global_id(0); a[id] = b[id] + c[id];

然而，鉴于get_global_id(0)的返回类型是size_t (无符号整数，通常是 64 位)，在某些情况下，转换会减少编译器可以合法执行的优化。例如，当编译器在内核中对代码进行矢量化处理时，这可能会导致 SIMD 收集/分散指令

读取a[get_global_id(0)]导致 SIMD 单位步幅向量加载。
对a[(int)get_global_id(0)]的读取导致非单位步幅采集指令。

这种微妙的情况是由从size_t到int(或uint)的数据类型转换的绕回行为(C++ 标准中未指定的行为和/或定义良好的绕回行为)引入的，这主要是基于 C 语言发展的历史产物。具体来说，跨某些转换的溢出是未定义的行为，这实际上允许编译器假设这种情况永远不会发生，并进行更积极的优化。图 16-17 为那些想了解细节的人展示了一些例子。

img/489625_1_En_16_Fig17_HTML.png

图 16-17

整数type value回绕的例子

SIMD 收集/分散指令比 SIMD 单位步长向量加载/存储操作慢。为了实现最佳的 SIMD 效率，无论使用哪种编程语言，避免聚集/分散对于应用来说都是至关重要的。

大多数 SYCL get_*_id()系列函数都有相同的细节，尽管许多情况都符合MAX_INT的范围，因为可能的返回值是有限的(例如，一个工作组内的最大 id)。因此，只要合法，DPC++ 编译器将假设相邻工作项块上的单位步长内存地址，以避免聚集/分散。如果由于全局 id 的值和/或全局 id 的导数值可能溢出，编译器不能安全地生成线性单位步长向量内存加载/存储，编译器将生成聚集/分散。

在为用户提供最佳性能的理念下，DPC++ 编译器假定没有溢出，并且在实践中几乎总是捕捉真实情况，因此编译器可以生成最佳 SIMD 代码以实现良好的性能。但是，DPC++ 编译器提供了一个覆盖编译器宏—D _ _ SYCL _ DISABLE _ ID _ TO _ INT _ conv _ _—来告诉编译器将会有溢出，并且从 ID 查询中导出的矢量化访问可能不安全。这可能会对性能产生很大影响，只要不安全，就应该使用这种方法来假设没有溢出。

SIMD 执行使用`single_task`

在单任务执行模型下，与向量类型和函数相关的优化取决于编译器。编译器和运行时可以自由地在single_task内核中启用显式 SIMD 执行或选择标量执行，结果将取决于编译器的实现。例如，DPC++ CPU 编译器支持向量类型，并为 CPU SIMD 执行生成 SIMD 指令。vec load、store 和 swizzle 函数将直接对向量变量执行操作，通知编译器数据元素正在访问从内存中相同(统一)位置开始的连续数据，并使我们能够请求连续数据的优化加载/存储。

img/489625_1_En_16_Fig18_HTML.png

图 16-18

在single_task内核中使用矢量类型和 swizzle 操作

在图 16-18 所示的例子中，在单任务执行下，声明了一个带有三个数据元素的向量。使用old_v.abgr().执行 swizzle 操作如果 CPU 为一些 swizzle 操作提供 SIMD 硬件指令，我们可以在应用程序中使用 swizzle 操作来获得一些性能优势。

SIMD VECTORIZATION GUIDELINES

CPU 处理器实现具有不同 SIMD 宽度的 SIMD 指令集。在许多情况下，这是一个实现细节，对于在 CPU 上执行内核的应用程序是透明的，因为编译器可以确定一组有效的数据元素，以特定的 SIMD 大小进行处理，而不是要求我们显式地使用 SIMD 指令。子组可以用于更直接地表达数据元素的分组应该在内核中经受 SIMD 执行的情况。

考虑到计算的复杂性，选择最适合矢量化的代码和数据布局可能最终会带来更高的性能。选择数据结构时，请尝试选择数据布局、对齐方式和数据宽度，以便最频繁执行的计算能够以 SIMD 友好的方式以最大的并行度访问内存，如本章所述。

摘要

为了充分利用 CPU 上的线程级并行和 SIMD 矢量级并行，我们需要牢记以下目标:

熟悉所有类型的 DPC++ 并行性和我们希望瞄准的底层 CPU 架构。
在与硬件资源最匹配的线程级别上，利用适量的并行性，不多也不少。使用供应商工具，如分析器和剖析器，来帮助指导我们的调优工作，以实现这一目标。
请注意线程关联和内存首次接触对程序性能的影响。
设计具有数据布局、对齐方式和数据宽度的数据结构，以便最频繁执行的计算能够以 SIMD 友好的方式访问内存，并具有最大的 SIMD 并行性。
注意平衡屏蔽和代码分支的成本。
使用清晰的编程风格，最大限度地减少潜在的内存混淆和副作用。
请注意使用 vector 类型和接口的可伸缩性限制。如果编译器实现将它们映射到硬件 SIMD 指令，固定的向量大小可能无法在多代 CPU 和来自不同供应商的 CPU 之间很好地匹配 SIMD 寄存器的 SIMD 宽度。

Creative Commons

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

十七、FPGAs 编程

img/489625_1_En_17_Figa_HTML.gif

基于内核的编程最初流行起来是作为访问 GPU 的一种方式。因为它现在已经被推广到许多类型的加速器，所以了解我们的编程风格如何影响代码到 FPGA 的映射也很重要。

现场可编程门阵列(FPGA)对于大多数软件开发人员来说是陌生的，部分原因是大多数台式计算机在典型的 CPU 和 GPU 旁边不包括 FPGA。但是 FPGA值得了解，因为它们在许多应用中提供了优势。我们需要问与其他加速器相同的问题，比如“我什么时候应该使用 FPGA？”，“我的应用的哪些部分应该卸载到 FPGA？”，以及“如何编写在 FPGA 上运行良好的代码？”

本章为我们提供了开始回答这些问题的知识，至少在这一点上，我们可以决定 FPGA 是否适合我们的应用，并了解哪些结构通常用于实现性能。这一章是一个起点，我们可以从这里开始阅读供应商文档，以填充特定产品和工具链的细节。我们首先概述程序如何映射到 FPGA 等空间架构，然后讨论 FPGA 作为加速器的一些特性，最后介绍用于提高性能的编程结构。

本章中的“如何考虑 FPGA”一节适用于考虑任何 FPGA。SYCL 允许供应商指定 CPU 和 GPU 之外的设备，但没有具体说明如何支持 FPGA。FPGA 的特定供应商支持目前是 DPC++ 独有的，即 FPGA 选择器和管道。FPGA 选择器和管道是本章中使用的唯一的 DPC++ 扩展。人们希望供应商们能够在支持 FPGAs 的相似或兼容的方法上达成一致，这也是 DPC++ 作为一个开源项目所鼓励的。

性能警告

与任何处理器或加速器一样，FPGA 设备因供应商而异，甚至因产品的不同而不同；因此，一种设备的最佳实践可能不是另一种设备的最佳实践。无论是现在还是将来，本章中的建议都可能使许多 FPGA 器件受益，但是…

…要实现特定 FPGA 的最佳性能，请务必查阅供应商的文档！

如何看待 FPGAs

FPGAs 通常被归类为空间架构。与使用指令集架构(ISA)的设备(包括大多数人更熟悉的 CPU 和 GPU)相比，它们受益于非常不同的编码风格和并行形式。为了开始了解 FPGAs，我们将简要介绍基于 ISA 的加速器的一些想法，以便强调关键差异。

就我们的目的而言，基于 ISA 的加速器是指设备可以执行许多不同的指令，一次执行一条或几条。这些指令通常相对简单，比如“从内存地址 A 加载”或“添加下列数字”一连串的操作串在一起形成一个程序，处理器在概念上一个接一个地执行指令。

在基于 ISA 的加速器中，芯片的单个区域(或整个芯片)在每个时钟周期执行来自程序的不同指令。指令在固定的硬件架构上执行，该架构可以在不同的时间运行不同的指令，如图 17-1 所示。例如，提供加法运算的内存加载单元可能与提供减法运算的内存加载单元相同。类似地，同一个算术单元可能既用于执行加法指令，也用于执行减法指令。随着程序的执行，芯片上的硬件被不同的指令重用。

img/489625_1_En_17_Fig1_HTML.png

图 17-1

简单的基于 ISA 的(临时)处理:随着时间的推移重用硬件(区域)

空间建筑是不同的。它们不是基于在共享硬件上执行各种指令的机器，而是从相反的角度出发。一个程序的空间实现在概念上把整个程序作为一个整体，并把它一次放在设备上。设备的不同区域在程序中执行不同的指令。这在许多方面与随时间推移在指令之间共享硬件(例如 ISA)的观点相反——在空间架构中，每个指令接收其自己的专用硬件，该硬件可以与实现其他指令的硬件同时执行(相同的时钟周期)。图 17-2 显示了这个想法，它是整个程序(在这个例子中是一个非常简单的程序)的空间实现。

img/489625_1_En_17_Fig2_HTML.png

图 17-2

空间处理:每次操作使用设备的不同区域

对程序的空间实现的这种描述过于简单，但它抓住了这样一个思想，即在空间架构中，程序的不同部分在设备的不同部分上执行，而不是随着时间的推移被发布到一组共享的更通用的硬件上。

由于 FPGA 的不同区域被编程为执行不同的操作，一些通常与基于 ISA 的加速器相关联的硬件是不必要的。例如，图 17-2 显示我们不再需要取指令或解码单元、程序计数器或寄存器文件。空间架构将一条指令的输出连接到下一条指令的输入，而不是将数据存储在寄存器文件中供未来指令使用，这就是空间架构通常被称为数据流架构的原因。

我们介绍的到 FPGA 的映射出现了几个明显的问题。首先，由于程序中的每条指令都要占用设备空间的一定比例，如果程序需要的空间超过 100%，会发生什么呢？一些解决方案提供了资源共享机制，使更大的程序能够以性能成本来适应，但 FPGAs 确实有程序适应的概念。这既是优点也是缺点:

好处:如果一个程序使用 FPGA 上的大部分区域，并且每个时钟周期都有足够的工作让所有硬件忙碌，那么在设备上执行程序会非常高效，因为它具有极高的并行性。更一般的架构每个时钟周期可能有大量未使用的硬件，而对于 FPGA，面积的使用可以针对特定应用进行完美定制，不会造成浪费。这种定制可以让应用通过大规模并行运行得更快，通常具有令人信服的能效。
不利的一面是:大型程序可能需要调整和重构，以适应设备。编译器的资源共享功能有助于解决这一问题，但通常会降低性能，从而降低使用 FPGA 的优势。基于 ISA 的加速器是非常高效的资源共享实现方式，事实证明，当应用的架构能够充分利用大部分可用区域时，FPGAs 对计算最有价值。

在极端情况下，FPGA 上的资源共享解决方案导致一种看起来像基于 ISA 的加速器的架构，但它内置于可重新配置的逻辑，而不是在固定硅中进行优化。相对于固定芯片设计，可重新配置的逻辑会带来额外开销，因此，通常不会选择 FPGA 来实现 ISA，当应用能够利用资源来实现高效的数据流算法时，FPGA 具有最大的优势，我们将在接下来的章节中讨论这一点。

流水线并行性

图 17-2 中经常出现的另一个问题是程序的空间实现与时钟频率的关系，以及程序从头到尾执行的速度。在所示的例子中，很容易相信可以从内存中加载数据，执行乘法和加法运算，并将结果存储回内存中，速度非常快。随着程序变得越来越大，可能在整个 FPGA 器件上有成千上万个操作，很明显，对于一个接一个地操作所有指令(操作通常取决于先前操作产生的结果)，考虑到每个操作引入的处理延迟，可能需要很长时间。

操作之间的中间结果在图 17-3 所示的空间架构中随时间更新(传播)。例如，加载执行，然后将其结果传递给乘法器，然后将其结果传递给加法器，依此类推。一段时间后，中间数据一直传播到操作链的末端，最终结果可用或存储到内存中。

img/489625_1_En_17_Fig3_HTML.png

图 17-3

一个简单空间计算实现的传播时间

如图 17-3 所示的空间实现是非常低效的，因为大部分硬件只在一小部分时间里做有用的工作。大多数情况下，乘法之类的操作要么等待来自加载的新数据，要么保持其输出，以便链中稍后的操作可以使用其结果。大多数空间编译器和实现通过流水线来解决这种低效率，这意味着单个程序的执行分散在许多时钟周期中。这是通过在一些操作之间插入寄存器(硬件中的数据存储原语)来实现的，其中每个寄存器在一个时钟周期内保存一个二进制值。通过保存操作输出的结果，以便链中的下一个操作可以看到并操作该保存的值，前一个操作可以自由地进行不同的计算，而不会影响后续操作的输入。

算法流水线的目标是让每个操作(硬件单元)在每个时钟周期都保持忙碌。图 17-4 显示了前面简单例子的流水线实现。请记住，编译器会为我们完成所有的流水线操作和平衡工作！我们讨论这个主题是为了让我们能够理解如何在接下来的部分中用工作填充管道，而不是因为我们需要担心在代码中手工管道化任何东西。

img/489625_1_En_17_Fig4_HTML.png

图 17-4

计算的流水线化:阶段并行执行

当空间实现被流水线化时，它变得非常高效，就像工厂流水线一样。每个管道阶段只执行整体工作的一小部分，但是它执行得很快，然后立即开始处理下一个工作单元。从开始到结束，流水线处理一个单个计算需要许多时钟周期，但是流水线可以同时对不同数据计算许多不同的计算实例。

当足够多的工作开始在流水线中执行时，经过足够多的连续时钟周期，然后每个单独的流水线阶段以及程序中的操作可以在每个时钟周期期间执行有用的工作，这意味着整个空间设备同时执行工作。这就是空间架构的强大之处之一——整个设备可以一直并行工作。我们称之为流水线并行。

流水线并行是在 FPGAs 上实现性能的主要并行形式。

PIPELINING IS AUTOMATIC

在用于 FPGA 的 DPC++ 的英特尔实现中，以及在用于 FPGA 的其他高级编程解决方案中，算法的流水线由编译器自动执行。大致了解空间架构上的实现是有用的，如本节所述，因为这样可以更容易地构建应用程序来利用管道并行性。应该明确的是，流水线寄存器插入和平衡是由编译器执行的，而不是由开发人员手动执行的。

真实的程序和算法通常具有控制流(例如，if/else 结构),该控制流使程序的某些部分在时钟周期的某个百分比不活动。FPGA 编译器通常在可能的情况下组合分支两端的硬件，以最小化浪费的空间面积，并在控制流分流期间最大化计算效率。这使得控制流分歧比其他架构，尤其是向量化的架构，成本更低，开发问题更少。

内核消耗芯片“面积”

在现有的实现中，DPC++ 应用中的每个内核都会生成一个空间流水线，该流水线会消耗 FPGA 的一些资源(我们可以将此视为设备上的空间或区域)，这在概念上如图 17-5 所示。

img/489625_1_En_17_Fig5_HTML.png

图 17-5

同一个 FPGA 二进制文件中的多个内核:内核可以并发运行

由于内核在设备上使用自己的区域，不同的内核可以并发执行。如果一个内核正在等待诸如存储器访问之类的事情，FPGA 上的其他内核可以继续执行，因为它们是芯片上其他地方的独立流水线。这种想法，正式描述为内核之间的独立向前进展，是 FPGA 空间计算的一个关键属性。

何时使用 FPGA

与任何加速器架构一样，预测 FPGA 何时是加速器的正确选择，何时是替代方案，通常取决于对架构、应用特性和系统瓶颈的了解。本节描述了要考虑的应用程序的一些特征。

很多很多的工作

像大多数现代计算加速器一样，实现良好的性能需要执行大量的工作。如果从单个数据元素计算单个结果，那么利用加速器可能根本没有用(任何种类的)。这和 FPGAs 没什么区别。知道 FPGA 编译器利用流水线并行性后，这一点变得更加明显。一个算法的流水线实现有许多阶段，通常有数千个或更多，每个阶段在任何时钟周期内都应该有不同的工作。如果没有足够的工作在大部分时间占据大部分流水线阶段，那么效率就会很低。我们将一段时间内流水线阶段的平均利用率称为流水线的占用率。这和优化 GPU 等其他架构时使用的占有率定义是不一样的！

有多种方法可以在 FPGA 上生成工作来填充流水线阶段，我们将在接下来的章节中讨论。

自定义操作或操作宽度

FPGAs 最初设计用于执行有效的整数和按位运算，并作为粘合逻辑，可以调整其他芯片的接口以相互配合工作。虽然 FPGAs 已经发展成为计算能力强大的解决方案，而不仅仅是粘合逻辑解决方案，但它们在按位运算、自定义数据宽度或类型的整数数学运算以及数据包报头中任意位字段的运算方面仍然非常高效。

本章末尾描述的 FPGA 的细粒度架构意味着可以高效地实现新颖和任意的数据类型。例如，如果我们需要 33 位整数乘法器或 129 位加法器，FPGAs 可以非常高效地提供这些自定义操作。由于这种灵活性，FPGAs 通常用于快速发展的领域，例如最近的机器学习，其中数据宽度和操作的变化速度超过了 ASICs 的内置速度。

标量数据流

从图 17-4 可以明显看出，FPGA 空间流水线的一个重要方面是，操作之间的中间数据不仅留在片内(不存储到外部存储器)，而且每个流水线级之间的中间数据都有专用的存储寄存器。FPGA 并行性来自于计算的流水线操作，使得许多操作同时执行，每个操作在流水线的不同阶段执行。这不同于向量架构，在向量架构中，多个计算作为共享向量指令的通道来执行。

空间管道中并行性的标量性质对于许多应用程序来说是重要的，因为即使在工作单元之间存在紧密的数据依赖性时，它仍然适用。这些数据依赖可以在不损失性能的情况下处理，我们将在本章后面讨论循环携带的依赖时讨论。其结果是，空间管道，因此 FPGAs，对于跨工作单元(如工作项)的数据依赖性不能被破坏并且必须进行细粒度通信的算法来说是有吸引力的。针对其他加速器的许多优化技术关注于通过各种技术打破这些依赖性，或者通过诸如子组之类的特性在受控规模上管理通信。相反，FPGAs 可以很好地执行紧密依赖的通信，并且应该被考虑用于存在这种模式的算法类别。

LOOPS ARE FINE!

对数据流架构的一个常见误解是，具有固定或动态迭代计数的循环会导致数据流性能不佳，因为它们不是简单的前馈管道。至少对于英特尔 DPC++ 和 FPGA 工具链来说，情况并非如此。相反，循环迭代是在管道内产生高占用率的好方法，编译器是围绕允许多个循环迭代以重叠方式执行的概念构建的。循环提供了一种简单的机制来让管道忙于工作！

低延迟和丰富的连接

利用器件上丰富的输入和输出收发器的 FPGAs 的更多传统用途同样适用于使用 DPC++ 的开发人员。例如，如图 17-6 所示，一些 FPGA 加速卡具有网络接口，可以将数据直接传输到设备中，进行处理，然后将结果直接传输回网络。当需要最小化处理等待时间时，以及当通过操作系统网络栈的处理太慢或需要卸载时，通常寻求这样的系统。

img/489625_1_En_17_Fig6_HTML.png

图 17-6

低延迟 I/O 流:FPGA 紧密连接网络数据和计算

当考虑通过 FPGA 收发器进行直接输入/输出时，机会几乎是无限的，但选择确实取决于构成加速器的电路板上的可用器件。由于依赖于特定的加速卡和各种这样的用途，除了在下一节中描述管道语言构造之外，本章不深入这些应用程序。相反，我们应该阅读与特定加速卡相关的供应商文档，或者搜索符合我们特定接口需求的加速卡。

定制的内存系统

FPGA 上的存储器系统，例如函数私有存储器或工作组本地存储器，是由小块片上存储器构建而成的。这很重要，因为每个内存系统都是为使用它的算法或内核的特定部分定制的。FPGAs 具有显著的片内存储器带宽，结合定制存储器系统的形成，它们可以在具有非典型存储器访问模式和结构的应用中表现出色。图 17-7 显示了在 FPGA 上实现内存系统时，编译器可以执行的一些优化。

img/489625_1_En_17_Fig7_HTML.png

图 17-7

FPGA 内存系统是由编译器为我们的特定代码定制的

其他架构，如 GPU，有固定的内存结构，有经验的开发人员很容易推理，但在许多情况下也很难优化。例如，其他加速器上的许多优化都集中在内存模式修改上，以避免存储体冲突。如果我们的算法可以受益于定制的存储器结构，例如每个存储体的不同数量的访问端口或不同寻常的存储体数量，那么 FPGAs 可以提供立竿见影的优势。从概念上讲，区别在于编写代码以有效地使用固定的内存系统(大多数其他加速器)和让编译器定制设计的内存系统对我们的特定代码有效(FPGA)。

在 FPGA 上运行

在 FPGA 上运行内核有两个步骤(与任何提前编译加速器一样):

将源代码编译成可以在我们感兴趣的硬件上运行的二进制文件
在运行时选择我们感兴趣的正确加速器

要编译内核以便它们可以在 FPGA 硬件上运行，我们可以使用命令行:


dpcpp -fintelfpga my_source_code.cpp -Xshardware

该命令告诉编译器将my_source_code.cpp中的所有内核转换成可以在英特尔 FPGA 加速器上运行的二进制文件，然后将它们打包到生成的主机二进制文件中。当我们执行主机二进制程序时(例如，通过在 Linux 上运行./a.out，在执行提交的内核之前，运行时将根据需要自动编程任何附加的 FPGA，如图 17-8 所示。

img/489625_1_En_17_Fig8_HTML.png

图 17-8

FPGA 在运行时自动编程

FPGA 编程二进制文件嵌入在我们在主机上运行的已编译的 DPC++ 可执行文件中。FPGA 是在幕后自动为我们配置的。

当我们运行一个主机程序并提交第一个内核以便在 FPGA 上执行时，在内核开始执行之前可能会有一点延迟，因为 FPGA 正在被编程。为额外的执行重新提交内核不会看到同样的延迟，因为内核已经被编程到设备中并准备好运行。

第二章介绍了运行时 FPGA 器件的选择。我们需要告诉主机程序我们希望内核在哪里运行，因为除了 FPGA 之外，通常还有多个加速器选项可用，例如 CPU 和 GPU。为了快速回顾在程序执行期间选择 FPGA 的一种方法，我们可以使用如图 17-9 所示的代码。

img/489625_1_En_17_Fig9_HTML.png

图 17-9

使用fpga_selector在运行时选择 FPGA 器件

编译时间

传言称，编译 FPGA 设计可能需要很长时间，比编译基于 ISA 的加速器要长得多。谣言是真的！本章结尾概述了 FPGA 的细粒度架构元素，这些元素既带来了 FPGA 的优势，也带来了计算密集型编译(布局布线优化)，在某些情况下可能需要数小时。

从源代码到 FPGA 硬件执行的编译时间足够长，我们不想专门在硬件中开发和迭代我们的代码。FPGA 开发流程提供了几个阶段，最大限度地减少了硬件编译的数量，使我们在硬件编译时间有限的情况下也能高效工作。图 17-10 显示了典型的阶段，在这些阶段，我们的大部分时间都花在提供快速周转和快速迭代的早期步骤上。

img/489625_1_En_17_Fig10_HTML.png

图 17-10

大多数验证和优化发生在冗长的硬件编译之前

来自编译器的仿真和静态报告是 DPC++ 中 FPGA 代码开发的基石。仿真器就像 FPGA 一样工作，包括支持相关扩展和仿真执行模型，但运行在主机处理器上。因此，编译时间与我们预期的编译到 CPU 设备的时间相同，尽管我们不会看到在实际 FPGA 硬件上执行所带来的性能提升。模拟器对于在应用程序中建立和测试功能正确性非常有用。

像仿真一样，静态报告由工具链快速生成。它们报告编译器创建的 FPGA 结构和编译器识别的瓶颈。这两者都可以用来预测我们的设计在 FPGA 硬件上运行时是否会获得良好的性能，并用来优化我们的代码。请阅读供应商的文档以获得关于报告的信息，这些报告通常会随着工具链的发布而不断改进(请参阅文档以了解最新和最棒的特性！).供应商提供了关于如何根据报告进行解释和优化的大量文档。这些信息将是另一本书的主题，所以我们不能在这一章中深入讨论细节。

FPGA 仿真器

仿真主要用于从功能上调试我们的应用程序，以确保它的行为符合预期并产生正确的结果。没有理由在编译时间更长的实际 FPGA 硬件上进行这种级别的开发。通过从dpcpp编译命令中移除-Xshardware标志，同时在我们的主机代码中使用INTEL::fpga_emulator_selector而不是INTEL::fpga_selector来激活仿真流。我们将使用如下命令进行编译


dpcpp -fintelfpga my_source_code.cpp

同时，我们将在运行时使用如图 17-11 所示的代码选择 FPGA 仿真器。通过使用fpga_emulator_selector, which使用主机处理器来仿真 FPGA，我们在必须为实际 FPGA 硬件进行更长时间的编译之前，保持了快速的开发和调试过程。

img/489625_1_En_17_Fig11_HTML.png

图 17-11

使用fpga_emulator_selector进行快速开发和调试

如果我们经常在硬件和仿真器之间切换，在程序中使用宏从命令行在设备选择器之间切换是有意义的。如果需要，请查看供应商的文档和在线 FPGA DPC++ 代码示例。

FPGA 硬件编译“提前”发生

图 17-10 中的完全编译和硬件剖析阶段在 SYCL 术语中是一个提前编译的*。这意味着内核到设备二进制文件的编译发生在我们最初编译程序的时候，而不是程序提交到设备上运行的时候。在 FPGA 上，这一点尤为重要，因为*

编译需要很长时间，这是我们在运行应用程序时通常不希望发生的。
DPC++ 程序可以在没有主处理器的系统上执行。FPGA 二进制文件的编译过程得益于具有大量附加存储器的快速处理器。提前编译让我们可以很容易地选择编译发生的位置，而不是在程序部署的系统上运行。

A LOT HAPPENS BEHIND THE SCENES WITH DPC++ ON AN FPGA!

传统的 FPGA 设计(不使用高级语言)可能非常复杂。除了编写我们的内核之外，还有许多步骤，例如构建和配置与片外存储器通信的接口，以及通过插入所需的寄存器来关闭时序，以使编译后的设计运行得足够快，从而与某些外设通信。DPC++ 为我们解决了这一切，让我们不需要了解任何常规 FPGA 设计的细节就可以实现工作应用！该工具将我们的内核视为优化和提高设备效率的代码，然后自动处理与片外外设对话、关闭时序和为我们设置驱动程序的所有细节。

与任何其他加速器一样，在 FPGA 上实现最高性能仍然需要详细的架构知识，但与传统 FPGA 流程相比，使用 DPC++ 从代码到工作设计的步骤要简单得多，效率也更高。

为 FPGAs 编写内核

一旦我们决定在我们的应用中使用 FPGA，或者只是决定试用一种，了解如何编写代码以获得良好的性能是非常重要的。本节描述了突出重要概念的主题，并涵盖了一些经常引起混淆的主题，以使入门更快。

暴露并行性

我们已经了解了如何利用流水线并行在 FPGA 上高效执行工作。另一个简单的管道示例如图 17-12 所示。

img/489625_1_En_17_Fig12_HTML.png

图 17-12

具有五个阶段的简单流水线:六个时钟周期处理一个数据元素

在这条管道中，有五个阶段。每个时钟周期，数据从一个阶段移动到下一个阶段一次，因此在这个非常简单的例子中，从数据进入阶段 1 到从阶段 5 退出需要 6 个时钟周期。

流水线的一个主要目标是使多个数据元素能够在流水线的不同阶段同时被处理。为了确保这一点是清楚的，图 17-13 显示了一个没有足够工作的流水线(在这种情况下只有一个数据元素)，这导致每个流水线级在大多数时钟周期内都没有被使用。这是对 FPGA 资源的低效使用，因为大部分硬件在大部分时间都是空闲的。

img/489625_1_En_17_Fig13_HTML.png

图 17-13

如果只处理单个工作元素，流水线阶段通常是不使用的

为了更好地占用管道阶段，想象一下在第一阶段之前等待的未启动工作队列是有用的，第一阶段向管道提供信息。每个时钟周期，流水线可以从队列中消耗并启动多一个工作元素，如图 17-14 所示。在一些初始启动周期之后，流水线的每个阶段都被占用，并在每个时钟周期做有用的工作，从而有效利用 FPGA 资源。

img/489625_1_En_17_Fig14_HTML.png

图 17-14

当每个流水线级保持忙碌时，就实现了高效利用

接下来的两个部分将介绍一些方法，这些方法可以让队列中充满准备好开始的工作。我们会看看

nd-range 核
环

在这些选项之间进行选择会影响运行在 FPGA 上的内核的基本架构。在某些情况下，算法很适合这种或那种风格，而在其他情况下，程序员的偏好和经验决定了应该选择哪种方法。

使用 ND-range 保持流水线繁忙

第章第四部分描述了 ND-range 分级执行模型。图 17-15 说明了关键的概念:一个 ND-range 执行模型，其中有一个工作项目的层次分组，工作项目是内核定义的基本工作单元。该模型最初被开发来实现 GPU 的高效编程，其中工作项目可以在执行模型层级的不同级别上并发执行。为了匹配 GPU 硬件高效的工作类型，ND 范围工作项在大多数应用程序中不会频繁地相互通信。

img/489625_1_En_17_Fig15_HTML.png

图 17-15

ND-range 执行模型:工作项目的层次分组

使用 ND-range 可以非常有效地填充 FPGA 空间流水线。FPGA 完全支持这种编程风格，我们可以将其视为图 17-16 所示，在每个时钟周期，不同的工作项进入流水线的第一阶段。

img/489625_1_En_17_Fig16_HTML.png

图 17-16

ND-range 供给空间管道

什么时候我们应该使用工作项在 FPGA 上创建一个 ND-range 内核来保持流水线被占用？很简单。每当我们可以将我们的算法或应用程序构建为不需要经常通信(或者理想情况下根本不需要通信)的独立工作项时，我们都应该使用 ND-range！如果工作项确实需要经常通信，或者如果我们不自然地考虑 ND 范围，那么循环(在下一节中描述)也提供了一种表达我们算法的有效方式。

如果我们可以构建我们的算法，使得工作项不需要太多的通信(或者根本不需要),那么 ND-range 是一个生成工作以保持空间管道满的好方法！

一个很好的例子是一个随机数发生器，在这个例子中，序列中的数字的创建独立于先前生成的数字。

图 17-17 显示了一个 ND-range 内核，它将为 16 × 16 × 16 范围内的每个工作项调用一次随机数生成函数。注意随机数生成函数是如何将工作项 id 作为输入的。

img/489625_1_En_17_Fig17_HTML.png

图 17-17

多个工作项(16 × 16 × 16)调用一个随机数生成器

这个例子展示了一个使用了一个range的parallel_for调用，只指定了一个全局大小。我们也可以使用带有nd_range的parallel_for调用风格，其中指定了全局工作大小和本地工作组大小。FPGAs 可以非常有效地从片内资源实现工作组本地存储器，因此只要有意义就可以随意使用工作组，因为我们需要工作组本地存储器，或者因为有了工作组 id 可以简化我们的代码。

PARALLEL RANDOM NUMBER GENERATORS

图 17-17 中的例子假设generate_random_number_from_ID(I)是一个随机数发生器，当以并行方式调用时，它是安全和正确的。例如，如果在parallel_for范围内的不同工作项执行该函数，我们期望每个工作项创建不同的序列，每个序列都遵循生成器期望的分布。并行随机数生成器本身就是一个复杂的主题，所以使用库或者通过诸如块跳转算法之类的技术来了解这个主题是一个好主意。

管道不介意数据依赖！

当对一些工作项目作为向量指令通道一起执行的向量架构(例如，GPU)进行编程时，挑战之一是在工作项目之间没有大量通信的情况下构造高效的算法。有些算法和应用程序非常适合矢量硬件，有些则不适合。映射不佳的一个常见原因是算法需要大量共享数据，这是由于数据依赖于在某种意义上相邻的其他计算。如第十四章所述，子组通过在同一个子组中的工作项之间提供有效的通信来解决向量架构上的一些挑战。

对于不能分解成独立工作的算法，FPGAs 起着重要的作用。FPGA 空间流水线不是跨工作项矢量化，而是跨流水线阶段执行连续的工作项。这种并行性的实现意味着工作项(甚至是不同工作组中的工作项)之间的细粒度通信可以在空间管道中轻松有效地实现！

一个例子是随机数发生器，其中输出 N+1 取决于知道输出 N 是什么。这在两个输出之间产生了数据依赖性，并且如果每个输出都是由 nd 范围中的工作项生成的，那么在一些架构上，工作项之间存在数据依赖性，这可能需要复杂并且通常昂贵的同步。当串行编码这样的算法时，通常会编写一个循环，其中迭代 N+1 使用来自迭代 N 的计算，如图 17-18 所示。每次迭代依赖于前一次迭代计算的状态。这是一种非常常见的模式。

img/489625_1_En_17_Fig18_HTML.png

图 17-18

循环携带的数据相关性(state)

空间实现可以非常有效地将结果在管道中向后传递给在稍后的周期中开始的工作(即，在管道中的较早阶段工作)，并且空间编译器围绕该模式实现了许多优化。图 17-19 显示了从阶段 5 到阶段 4 的数据反向通信的想法。空间管道不会跨工作项进行矢量化。这通过在管道中向后传递结果实现了高效的数据依赖通信！

img/489625_1_En_17_Fig19_HTML.png

图 17-19

反向通信支持高效的数据相关通信

向后传递数据(传递到管道中的早期阶段)的能力是空间架构的关键，但如何编写利用这一点的代码并不明显。有两种方法可以轻松表达这种模式:

环
具有 ND 范围内核的内核内管道

第二种选择是基于管道的，我们将在本章后面介绍，但它远不如循环那样常见，所以为了完整起见我们提到了它，但在这里就不详述了。供应商文档提供了关于管道方法的更多细节，但是更容易坚持下面描述的循环，除非有理由这样做。

循环的空间流水线实现

当对具有数据依赖性的算法进行编程时，循环是一种自然的选择。循环经常表示迭代之间的依赖关系，即使在最基本的循环例子中，决定循环何时退出的计数器也是跨迭代执行的(图 17-20 中的变量i)。

img/489625_1_En_17_Fig20_HTML.png

图 17-20

具有两个循环携带依赖项(即i和a)的循环

在图 17-20 的简单循环中，a= a + i右侧的a的值反映了前一次循环迭代存储的值，如果是循环的第一次迭代，则为初始值。当空间编译器实现一个循环时，循环的迭代可以用来填充流水线的各个阶段，如图 17-21 所示。注意，现在准备开始的工作队列包含循环迭代，而不是工作项！

img/489625_1_En_17_Fig21_HTML.png

图 17-21

由循环的连续迭代供给的流水线阶段

图 17-22 显示了一个修改后的随机数发生器示例。在这种情况下，不是根据工作项的 id 生成一个数字，如图 17-17 所示，生成器将之前计算的值作为一个参数。

img/489625_1_En_17_Fig22_HTML.png

图 17-22

取决于先前生成的值的随机数生成器

这个例子使用了single_task而不是parallel_for，因为重复的工作是由单个任务中的一个循环表示的，所以没有理由在这个代码中也包含多个工作项(通过parallel_for)。single_task中的循环使得将之前计算的temp值传递给随机数生成函数的每次调用变得更加容易(便于编程)。

在如图 17-22 的情况下，FPGA 可以有效地实现循环。在许多情况下，它可以保持一个完全占用的管道，或者至少可以通过报告告诉我们要改变什么来增加占用率。考虑到这一点，很明显，如果用工作项替换循环迭代，这个算法将更加难以描述，其中一个工作项生成的值需要传递给另一个工作项，以便在增量计算中使用。代码的复杂性会迅速增加，特别是如果工作不能被批处理，那么每个工作项实际上是在计算它自己独立的随机数序列。

循环启动间隔

从概念上讲，我们可能认为 C++ 中的循环迭代是一个接一个地执行，如图 17-23 所示。这就是编程模型，也是思考循环的正确方式。然而，在实现中，只要程序的大多数行为(即，定义的和无竞争的行为)没有明显改变，编译器就可以自由地执行许多优化。不考虑编译器优化，重要的是循环看起来执行*，就好像*图 17-23 是如何发生的。

img/489625_1_En_17_Fig23_HTML.png

图 17-23

从概念上讲，循环迭代一个接一个地执行

从空间编译器的角度来看，图 17-24 显示了循环流水线优化，其中循环迭代的执行在时间上是重叠的。不同的迭代将执行彼此不同的空间流水线阶段，并且跨流水线阶段的数据依赖性可以由编译器管理，以确保程序看起来好像迭代是连续的一样执行(除了循环将更快地完成执行！).

img/489625_1_En_17_Fig24_HTML.png

图 17-24

循环流水线操作允许循环的迭代在流水线阶段之间重叠

认识到循环迭代中的许多结果可能在循环迭代完成其所有工作之前完成计算，并且在空间流水线中，当编译器决定这样做时，结果可以被传递到更早的流水线阶段，循环流水线是容易理解的。图 17-25 显示了这个想法，其中阶段 1 的结果在流水线中被反馈，允许未来的循环迭代在前一次迭代完成之前尽早使用该结果。

img/489625_1_En_17_Fig25_HTML.png

图 17-25

增量随机数发生器的流水线实现

使用循环流水线，一个循环的多次迭代的执行有可能重叠。这种重叠意味着，即使存在循环携带的数据依赖性，循环迭代仍然可以用来填充工作管道，从而实现高效利用。图 17-26 显示了在如图 17-25 所示的同一条简单流水线中，循环迭代如何重叠它们的执行，即使有循环携带的数据依赖。

img/489625_1_En_17_Fig26_HTML.png

图 17-26

循环流水线同时处理多个循环迭代的部分

在实际算法中，通常不可能在每个时钟周期启动新的循环迭代，因为数据相关性可能需要多个时钟周期来计算。如果存储器查找，特别是从片外存储器的查找，在相关性计算的关键路径上，这经常发生。结果是流水线只能每N个时钟周期启动一次新的循环迭代，我们称之为N周期的启动间隔 ( II)。示例如图 17-27 所示。两个循环启动间隔(II)意味着新的循环迭代可以每隔一个周期开始，这导致流水线级的次优占用。

img/489625_1_En_17_Fig27_HTML.png

图 17-27

流水线级的次优占用

大于 1 的II会导致流水线效率低下，因为每一级的平均占用率会降低。从图 17-27 中可以明显看出，其中II=2和管道级未使用的比例很大(50%！)的时间。有多种方法可以改善这种情况。

编译器会尽可能地执行大量优化来减少 II，因此它的报告还会告诉我们每个循环的初始间隔是多少，如果发生这种情况，还会告诉我们为什么它大于 1。基于报告在循环中重构计算通常可以减少II，特别是因为作为开发人员，我们可以进行编译器不允许的循环结构更改(因为它们会被观察到)。阅读编译器报告，了解如何在特定情况下减少II。

另一种降低大于 1 的II的低效率的方法是通过嵌套循环，这可以通过将外部循环迭代与具有II>1的内部循环迭代交错来填充所有流水线阶段。有关使用这种技术的详细信息，请查阅供应商文档和编译器报告。

管道

空间和其他架构中的一个重要概念是先进先出(FIFO)缓冲器。FIFOs 之所以重要，有很多原因，但在考虑编程时，有两个属性特别有用:

伴随数据的还有隐含的控制信息。这些信号告诉我们 FIFO 是空的还是满的，这在将问题分解成独立的部分时很有用。
FIFOs 有存储容量。这使得在存在动态行为(如访问内存时高度可变的延迟)的情况下实现性能变得更加容易。

图 17-28 显示了一个 FIFO 操作的简单例子。

img/489625_1_En_17_Fig28_HTML.png

图 17-28

一段时间内 FIFO 的操作示例

在 DPC++ 中，FIFOs 是通过一个名为管道的特性公开的。在编写 FPGA 程序时，我们应该关心管道的主要原因是，它们允许我们将问题分解为更小的部分，以便以更模块化的方式专注于开发和优化。它们还允许利用 FPGA 丰富的通信功能。图 17-29 用图形显示了这两种情况。

img/489625_1_En_17_Fig29_HTML.png

图 17-29

管道简化了模块化设计和对硬件外设的访问

请记住，FPGA 内核可以同时存在于器件上(位于芯片的不同区域),在高效设计中，内核的所有部分在每个时钟周期都是活动的。这意味着优化 FPGA 应用需要考虑内核或部分内核之间的交互方式，而管道提供了一种抽象来简化这一过程。

管道是使用 FPGA 上的片内存储器实现的 FIFOs，因此它们允许我们在运行的内核之间和内部进行通信，而无需将数据移动到片外存储器。这提供了廉价的通信，并且与管道(空/满信号)耦合的控制信息提供了轻量级的同步机制。

DO WE NEED PIPES?

不。不使用管道也可以编写高效的内核。我们可以使用所有的 FPGA 资源，并使用没有管道的传统编程风格实现最高性能。但是对于大多数开发人员来说，编程和优化更模块化的空间设计更容易，管道是实现这一点的好方法。

如图 17-30 所示，有四种通用类型的管道可供选择。在本节的剩余部分，我们将讨论第一种类型(内核间管道)，因为它们足以说明什么是管道以及如何使用管道。管道还可以在单个内核中通信，并与主机或输入/输出外围设备通信。关于管道的形式和用途的更多信息，请查阅供应商文档，我们在这里没有空间深入讨论。

img/489625_1_En_17_Fig30_HTML.png

图 17-30

DPC++ 中管道连接的类型

一个简单的例子如图 17-31 所示。在这种情况下，有两个内核通过管道进行通信，每个读或写操作在一个int单元上进行。

img/489625_1_En_17_Fig31_HTML.png

图 17-31

两个内核之间的管道:(1) ND-range 和(2)具有循环的单个任务

从图 17-31 中可以观察到几点。首先，两个内核使用管道相互通信。如果内核之间没有访问器或事件依赖，DPC++ 运行时将同时执行两个*，允许它们通过管道而不是完整的 SYCL 内存缓冲区或 USM 进行通信。*

*使用基于类型的方法识别管道，其中每个管道使用管道类型的参数化进行识别，如图 17-32 所示。管道类型的参数化标识了特定的管道。对同一管道类型的读取或写入是对同一 FIFO 的。有三个模板参数共同定义了管道的类型和标识。

img/489625_1_En_17_Fig32_HTML.png

图 17-32

管道类型的参数化

建议使用类型别名来定义我们的管道类型，如图 17-31 的第一行代码所示，以减少编程错误，提高代码可读性。

使用类型别名来标识管道。这简化了代码，并防止意外创建意外管道。

管道有一个min_capacity参数。它默认为 0，这是自动选择，但是如果指定，它保证至少该数量的字可以被写入管道而不被读出。此参数在以下情况下很有用

与管道通信的两个内核不同时运行，我们需要管道中有足够的容量供第一个内核在第二个内核开始运行并从管道中读取数据之前写入其所有输出。
如果内核以突发方式生成或消耗数据，那么增加管道的容量可以在内核之间提供隔离，使它们的性能彼此分离。例如，产生数据的内核可以继续写入(直到管道容量变满)，即使消耗该数据的内核很忙，还没有准备好消耗任何东西。这提供了内核相对于彼此执行的灵活性，代价仅仅是 FPGA 上的一些存储器资源。

阻塞和非阻塞管道入口

和大多数 FIFO 接口一样，管道有两种接口风格:阻塞和非阻塞。阻塞访问等待(阻塞/暂停执行！)以使操作成功，而非阻塞访问会立即返回，并设置一个布尔值来指示操作是否成功。

成功的定义很简单:如果我们从管道中读取，并且有数据可供读取(管道不为空)，那么读取成功。如果我们正在写并且管道还没有满，那么写成功。图 17-33 显示了 pipe 类访问成员函数的两种形式。我们看到管道的成员函数允许它被写入或读取。回想一下，对管道的访问可以是阻塞的，也可以是非阻塞的。

img/489625_1_En_17_Fig33_HTML.png

图 17-33

允许对其进行读写的管道成员函数

阻塞访问和非阻塞访问都有其用途，这取决于我们的应用程序试图实现的目标。如果内核在从管道中读取数据之前不能做更多的工作，那么使用阻塞读取可能是有意义的。相反，如果内核希望从一组管道中的任何一个读取数据，并且不确定哪个管道可能有可用的数据，那么使用非阻塞调用从管道读取数据更有意义。在这种情况下，内核可以从管道中读取数据并处理数据(如果有数据的话),但是如果管道是空的，它可以继续尝试从下一个可能有数据可用的管道中读取数据。

有关管道的更多信息

在这一章中，我们只能触及管道的表面，但是我们现在应该对它们有一个概念，以及如何使用它们的基本知识。FPGA 供应商文档提供了更多信息，以及它们在不同类型应用中的使用示例，因此，如果我们认为管道与我们的特定需求相关，我们应该查看一下。

定制存储系统

在为大多数加速器编程时，大部分优化工作都倾向于使内存访问更加高效。FPGA 设计也是如此，尤其是当输入和输出数据通过片外存储器时。

FPGA 上的存储器访问值得优化有两个主要原因:

减少所需的带宽，特别是在一些带宽使用效率低下的情况下
修改导致空间流水线中不必要的停顿的存储器上的访问模式

有必要简单谈谈空间管道中的失速。编译器内置了关于从特定类型的内存中读取或向其写入所需时间的假设，并相应地优化和平衡了流水线，从而隐藏了进程中的内存延迟。但是，如果我们以低效的方式访问内存，我们可能会引入更长的延迟，并且作为一种副产品在流水线中停止，其中早期阶段无法取得执行进展，因为它们被等待某些东西(例如，内存访问)的流水线阶段阻塞。图 17-34 显示的就是这样一种情况，负载上方的管线停滞不前，无法前进。

img/489625_1_En_17_Fig34_HTML.png

图 17-34

内存停顿如何导致早期管道阶段也停顿

内存系统优化可以在几个方面进行。像往常一样，编译器报告是我们了解编译器为我们实现了什么，以及什么可能值得调整或改进的主要指南。我们在这里列出了几个优化主题，以突出我们可用的一些自由度。优化通常可以通过显式控件和修改代码来实现，以允许编译器推断出我们想要的结构。编译器静态报告和供应商文档是内存系统优化的关键部分，有时会在硬件执行期间与评测工具结合使用，以捕获实际的内存行为，用于验证或最终的调整阶段。

静态合并(Static coalescing):编译器会将内存访问合并成数量更少、范围更广的访问。这降低了存储器系统在流水线中加载或存储单元的数量、存储器系统上的端口、仲裁网络的大小和复杂性以及其它存储器系统细节方面的复杂性。一般来说，我们希望尽可能启用静态合并，这可以通过编译器报告来确认。简化内核中的寻址逻辑有时足以让编译器执行更积极的静态合并，因此请始终检查报告，确保编译器已经推断出我们所期望的内容！
内存访问 风格:编译器为内存访问创建加载或存储单元，这些单元针对被访问的内存技术(例如片上与 DDR 或 HBM)以及从源代码推断的访问模式(例如流、动态合并/加宽，或者可能受益于特定大小的高速缓存)而定制。编译器报告告诉我们编译器推断出了什么，并允许我们在相关的地方修改或添加控制到我们的代码，以提高性能。
内存系统 结构:内存系统(片内和片外)可以具有由编译器实现的分组结构和众多优化。有许多控件和模式修改可用于控制这些结构和调整空间实现的特定方面。

一些结束主题

在与初学 FPGAs 的开发人员交谈时，我们发现从较高的层面理解组成器件的元件以及提到时钟频率(这似乎是一个困惑点)通常会有所帮助。我们以这些话题结束这一章。

FPGA 构建模块

为了帮助理解工具流(特别是编译时)，有必要提一下构成 FPGA 的构建模块。这些构建块是通过 DPC++ 和 SYCL 抽象出来的，它们的知识在典型的应用程序开发中不起作用(至少在使代码功能化的意义上)。然而，它们的存在确实会影响空间架构优化和工具流的直觉发展，例如，在为我们的应用程序选择理想的数据类型时，偶尔会影响高级优化。

一个非常简化的现代 FPGA 器件由五个基本元件组成。

查找表:有几根二进制输入线并产生二进制输出的基本块。相对于输入的输出通过编程到查找表中的条目来定义。这些是非常原始的模块，但是在用于计算的典型现代 FPGA 上有许多(数百万)这样的模块。这些是我们大部分设计实现的基础！
数学引擎:对于常见的数学运算，如单精度浮点数的加法或乘法，FPGAs 有专门的硬件来使这些运算非常高效。一个现代的 FPGA 有数千个这样的模块——有些设备有超过 8000 个——这样至少这些浮点原语操作可以在每个时钟周期并行执行*！大多数 FPGAs 将这些数学引擎命名为数字信号处理器(DSP)。*
片内存储器:这是 FPGAs 区别于其他加速器的一个方面，存储器有两种类型(更确切地说，但我们不会在这里讨论):(1)寄存器，用于操作和其他一些目的之间的管道传输，以及(2)块存储器，提供遍布整个设备的小型随机存取存储器。现代 FPGA 可以拥有数百万个寄存器位和超过 10，000 个 20 Kbit RAM 存储模块。由于每一个时钟周期都可以激活，因此如果有效利用，片内存储器容量和带宽会非常可观。
片外硬件接口:FPGA 的发展在一定程度上是因为其非常灵活的收发器和输入/输出连接，允许与从片外存储器到网络接口等几乎任何东西进行通信。
路由结构 在 之间所有其他元素:在一个典型的 FPGA 上，前面提到的每个元素都有很多，它们之间的连接是不固定的。复杂的可编程路由结构允许信号在构成 FPGA 的细粒度元素之间传递。

给定每种特定类型的 FPGA 上的块的数量(一些块以百万计)和那些块的精细粒度，例如查找表，在生成 FPGA 配置比特流时看到的编译时间可能更有意义。不仅需要将功能分配给每个细粒度的资源，还需要在它们之间配置路由。在优化开始之前，大部分编译时间来自于找到我们的设计到 FPGA 结构的第一个合法映射！

时钟频率

FPGA 非常灵活且可配置，且与加固到 CPU 或任何其他固定计算机架构中等效设计相比，这种可配置性对 FPGA 运行的频率带来了一些成本。但这不是问题！FPGA 的空间架构不仅弥补了时钟频率，因为有如此多的独立操作同时发生，分布在 FPGA 的整个区域。简而言之，由于可配置的设计，FPGA 的频率低于其他架构，但每个时钟周期会发生更多的变化，从而平衡频率。在基准测试和比较加速器时，我们应该比较计算吞吐量(例如，每秒操作数)而不是原始频率。

也就是说，当 FPGA 上的资源利用率接近 100%时，工作频率可能会开始下降。这主要是由于设备上的信号路由资源被过度使用。有一些方法可以解决这个问题，通常是以增加编译时间为代价。但是，对于大多数应用，最好避免使用 FPGA 上超过 80–90%的资源，除非我们愿意深入细节以抵消频率下降。

经验法则是尽量不要超过 FPGA 上任何资源的 90%,当然也不要超过多个资源的 90%。超出可能会导致路由资源耗尽，从而导致工作频率降低，除非我们愿意深入研究较低级别的 FPGA 细节来抵消这一点。

摘要

在本章中，我们介绍了流水线如何将算法映射到 FPGA 的空间架构。我们还讨论了一些概念，这些概念可以帮助我们判断 FPGA 对我们的应用是否有用，并且可以帮助我们更快地启动和运行开发代码。从这一点出发，我们应该能够浏览供应商的编程和优化手册，并开始编写 FPGA 代码！FPGAs 提供的性能和支持的应用在其他加速器上没有意义，所以我们应该把它们放在我们大脑工具箱的前端！

Creative Commons

十八、库

img/489625_1_En_18_Figa_HTML.gif

我们花了整本书来宣传编写我们自己的代码的艺术。现在我们终于承认，一些伟大的程序员已经写出了我们可以直接使用的代码。库是我们完成工作的最好方式。这不是懒惰的问题，而是有比重新发明他人的工作更好的事情要做的问题。这是一块值得拥有的拼图。

开源的 DPC++ 项目包括一些库。这些库可以帮助我们继续使用 libstdc++、libc++ 和 MSVC 库函数，甚至是在我们的内核代码中。这些库包含在英特尔的 DPC++ 和 oneAPI 产品中。这些库不依赖于 DPC++ 编译器，因此它们可以与任何 SYCL 编译器一起使用。

DPC++ 库为创建异构应用程序和解决方案的程序员提供了另一种选择。它的 API 基于熟悉的标准——c++ STL、并行 STL (PSTL)和 SYCL——为程序员提供高生产率的 API。这可以最大限度地减少 CPU、GPU 和 FPGAs 之间的编程工作量，同时实现可移植的高性能并行应用。

SYCL 标准定义了一组丰富的内置函数，为主机和设备代码提供功能，也值得考虑。DPC++ 和许多 SYCL 实现用数学库实现了关键的数学内置。

本章讨论的库和内置是编译器不可知的。换句话说，它们同样适用于 DPC++ 编译器或 SYCL 编译器。fpga_device_policy类是用于 FPGA 支持的 DPC++ 特性。

由于在命名和功能上有重叠，本章将从 SYCL 内置函数的简单介绍开始。

内置函数

DPC++ 针对各种数据类型提供了一组丰富的 SYCL 内置函数。这些内置函数在主机和设备上的sycl命名空间中可用，根据编译器选项对目标设备提供低、中、高精度支持，例如 DPC++ 编译器提供的-mfma、-ffast-math和-ffp-contract=fast。主机和设备上的这些内置功能可分为以下几类:

浮点数学函数:asin、acos、log, sqrt、floor等。在图 18-2 中列出。
整数函数:abs、max, min等。如图 18-3 所示。
常用功能:clamp、smoothstep等。如图 18-4 所示。
几何函数:cross、dot、distance等。在图 18-5 中列出。
关系函数:isequal、isless、isfinite等。在图 18-6 中列出。

如果 C++ std 库提供了一个函数，如图 18-8 所列，以及一个 SYCL 内置函数，那么 DPC++ 程序员可以使用其中任何一个。图 18-1 展示了主机和设备的 C++ std::log函数和 SYCL 内置sycl::log函数，两个函数产生相同的数值结果。在示例中，内置的关系函数sycl::isequal用于比较std:log和sycl:的结果。

img/489625_1_En_18_Fig1_HTML.png

图 18-1

使用std::log和sycl:: log

除了 SYCL 中支持的数据类型，DPC++ 设备库还支持将std:complex作为一种数据类型，以及 C++ std 库中定义的相应数学函数。

对内置函数使用前缀`sycl::`

调用 SYCL 内置函数时，应该在名字前添加一个明确的sycl::。根据当前的 SYCL 规范，即使使用了“using namespace sycl;”，也不能保证只调用sqrt()就能调用所有实现中内置的 SYCL。

调用 SYCL 内置函数时，应该在内置名称前加上一个明确的sycl::。不遵循这个建议可能会导致奇怪和不可移植的结果。

如果内置函数名与我们应用程序中的非模板化函数冲突，在许多实现中(包括 DPC++ ),我们的函数将优先，这要归功于 C++ 重载决策规则，它更喜欢非模板化函数而不是模板化函数。然而，如果我们的代码有一个与内置名相同的函数名，最方便的做法是避免using namespace sycl;或者确保没有实际冲突发生。否则，一些 SYCL 编译器会因为实现中无法解决的冲突而拒绝编译代码。这样的冲突不会沉寂。因此，如果我们的代码今天编译，我们可以安全地忽略未来问题的可能性。

img/489625_1_En_18_Fig6_HTML.png

图 18-6

内置关系函数

img/489625_1_En_18_Fig5_HTML.png

图 18-5

内置几何函数

img/489625_1_En_18_Fig4_HTML.png

图 18-4

内置常用功能

img/489625_1_En_18_Fig3a_HTML.png

图 18-3

内置整数函数

img/489625_1_En_18_Fig2_HTML.png

图 18-2

内置数学函数

DPC++ 库

DPC++ 库由以下组件组成:

一组经过测试的 C++ 标准 API——我们只需要包含相应的 C++ 标准头文件，并使用std名称空间。
包含相应头文件的并行 STL。我们简单地使用#include <dpstd/...>来包含它们。DPC++ 库将名称空间dpstd用于扩展的 API 类和函数。

DPC++ 中的标准 c++ API

DPC++ 库包含一组经过测试的标准 c++ API。已经开发了许多 C++ 标准 API 的基本功能，使得这些 API 可以在设备内核中使用，类似于它们在典型 C++ 主机应用程序的代码中的使用方式。图 18-7 显示了如何在设备代码中使用std::swap的例子。

img/489625_1_En_18_Fig7_HTML.png

图 18-7

在设备代码中使用std::swap

我们可以使用下面的命令来构建和运行程序(假设它驻留在stdswap.cpp文件中):


dpcpp –std=c++17 stdswap.cpp –o stdswap.exe
./stdswap.exe

打印结果是:


8, 9
9, 8

图 18-8 列出了带有“ Y 的 C++ 标准 API，以表明在撰写本文时，这些 API 已经过测试，可用于 CPU、和 FPGA 设备的 DPC++ 内核。空白表示本书出版时覆盖范围不完整(并非所有三种设备类型)。作为在线 DPC++ 语言参考指南的一部分，还包括一个表，该表将随着时间的推移而更新——dpc++ 中的库支持将继续扩展其支持。

在 DPC++ 库中，一些 C++ std函数是基于它们在设备上对应的内置函数实现的，以达到与这些函数的 SYCL 版本相同的性能水平。

img/489625_1_En_18_Fig8a_HTML.png img/489625_1_En_18_Fig8b_HTML.png

图 18-8

包含 CPU/GPU/FPGA 的库支持(在图书出版时)

测试的标准 C++ API 在libstdc++ (GNU)的gcc 7.4.0 和libc++ (LLVM)的clang 10.0 中得到支持，MSVC 标准 c++ 库的 Microsoft Visual Studio 2017 也支持主机 CPU。

在 Linux 上，GNU libstdc++是 DPC++ 编译器的默认 C++ 标准库，因此不需要编译或链接选项。如果我们想使用libc++，使用编译选项-stdlib=libc++ -nostdinc++来利用libc++，并且不包括系统中的 C++ std 头文件。已经在 Linux 上的 DPC++ 内核中使用libc++验证了 DPC++ 编译器，但是需要使用libc++而不是libstdc++来重建 DPC++ 运行时。详情在 https://intel.github.io/llvm-docs/GetStartedGuide.html#build-dpc-toolchain-with-libc-library 。由于这些额外的步骤，libc++不是我们通常使用的推荐的 C++ 标准库。

在 FreeBSD 上，libc++是默认的标准库，不需要-stdlib=libc++选项。更多详情请见 https://libcxx.llvm.org/docs/UsingLibcxx.html 。在 Windows 上，只能使用 MSVC C++ 库。

为了实现跨架构的可移植性，如果一个 std 函数在图 18-8 中没有标注“Y”，我们在编写设备函数的时候就需要牢记可移植性！

DPC++ 并行 STL

并行 STL 是支持执行策略的 C++ 标准库算法的实现，如 ISO/IEC 14882:2017 标准(通常称为 C++17)中所规定的。现有实现还支持在并行 ts 版本 2 中指定的未排序执行策略，并在 C++ 工作组论文 P1001R1 中为下一版本的 C++ 标准提出了该策略。

当使用算法和执行策略时，如果没有 C++17 标准库的特定于供应商的实现，则指定名称空间std::execution，否则指定pstl::execution。

对于任何已实现的算法，我们可以传递值seq、unseq、par或par_unseq中的一个作为算法调用的第一个参数，以指定所需的执行策略。这些策略具有以下含义:

执行策略

意为

| | --- | --- | | seq | 顺序执行。 | | unseq | 未排序的 SIMD 处决。该政策要求所提供的所有功能在 SIMD 均可安全执行。 | | par | 多线程并行执行。 | | par_unseq | unseq和par.的综合效果 |

DPC++ 的并行 STL 通过使用特殊的执行策略扩展了对 DPC++ 设备的支持。DPC++ 执行策略指定了并行 STL 算法在哪里以及如何运行。它继承了标准的 C++ 执行策略，封装了 SYCL 设备或队列，并允许我们设置可选的内核名称。DPC++ 执行策略可以与所有支持符合 C++17 标准的执行策略的标准 C++ 算法一起使用。

DPC++ 执行策略

目前，DPC++ 库只支持并行未排序策略(par_unseq)。为了使用 DPC++ 执行策略，有三个步骤:

将#include <dpstd/execution>添加到我们的代码中。
通过提供标准策略类型、作为模板参数的唯一内核名称的类类型(可选)以及下列构造器参数之一来创建策略对象:
- 赛 CL 队列
- SYCL 设备
- SYCL 设备选择器
- 具有不同内核名称的现有策略对象
将创建的策略对象传递给并行 STL 算法。

一个dpstd::execution::default_policy对象是一个预定义的device_policy,使用默认内核名和默认队列创建。这可用于创建自定义策略对象，或者在调用算法时直接传递(如果默认选项足够的话)。

图 18-9 显示了引用policy类和函数时假设使用using namespace dpstd::execution;指令的例子。

img/489625_1_En_18_Fig9_HTML.png

图 18-9

创建执行策略

FPGA 执行策略

fpga_device_policy类是一种 DPC++ 策略，旨在 FPGA 硬件设备上实现更好的并行算法性能。在 FPGA 硬件或 FPGA 仿真设备上运行应用程序时使用策略:

定义在 FPGA 设备上运行的_PSTL_FPGA_DEVICE宏，以及在 FPGA 仿真设备上运行的_PSTL_FPGA_EMU。
将#include 添加到我们的代码中。
通过提供唯一内核名称的类类型和展开因子(见第十七章)作为模板参数(两者都是可选的)和下列构造器参数之一，创建一个策略对象:
- 为 FPGA 选择器的构建的 SYCL 队列(该行为在任何其他设备类型中都没有定义)
- 具有不同内核名称和/或展开因子的现有 FPGA 策略对象
将创建的策略对象传递给并行 STL 算法。

fpga_device_policy的默认构造器创建一个对象，该对象带有为fpga_selector构造的 SYCL 队列，或者如果定义了_PSTL_FPGA_EMU的话，为fpga_emulator_selector构造的 SYCL 队列。

dpstd::execution::fpga_policy是用默认内核名和默认展开因子创建的fpga_device_policy类的预定义对象。使用它来创建定制的策略对象，或者在调用算法时直接传递它。

图 18-10 中的代码假设using namespace dpstd::execution;用于策略，而using namespace sycl;用于队列和设备选择器。

为策略指定展开因子可以在算法实现中实现循环展开。默认值为 1。要了解如何选择更好的值，请参见第十七章。

img/489625_1_En_18_Fig10_HTML.png

图 18-10

使用 FPGA 策略

使用 DPC++ 并行 STL

为了使用 DPC++ 并行 STL，我们需要通过添加以下行的子集来包含并行 STL 头文件。这些行取决于我们打算使用的算法:

#include <dpstd/algorithm>
#include <dpstd/numeric>
#include <dpstd/memory>

dpstd::begin和dpstd::end是特殊的帮助函数，允许我们将 SYCL 缓冲区传递给并行 STL 算法。这些函数接受 SYCL 缓冲区，并返回满足以下要求的未指定类型的对象:

是CopyConstructible、CopyAssignable，与操作者==、!=可比。
以下表达式有效:a + n、a – n和a – b，其中a和b是类型的对象，n是整数值。
有一个没有参数的get_buffer方法。该方法返回传递给dpstd::begin和dpstd::end函数的 SYCL 缓冲区。

要使用这些助手函数，请将#include <dpstd/iterators>添加到我们的代码中。参见图 18-11 和 18-12 中使用std::fill函数的代码，作为使用开始/结束帮助器的示例。

img/489625_1_En_18_Fig11_HTML.png

图 18-11

使用std::fill

REDUCE DATA COPYING BETWEEN THE HOST AND DEVICE

并行 STL 算法可以用普通(主机端)迭代器调用，如图 18-11 中的代码示例所示。

在这种情况下，会创建一个临时 SYCL 缓冲区，并将数据复制到该缓冲区。在对设备上的临时缓冲区的处理完成之后，数据被复制回主机。建议尽可能直接使用现有的 SYCL 缓冲区，以减少主机和设备之间的数据移动，以及任何不必要的缓冲区创建和销毁开销。

img/489625_1_En_18_Fig12_HTML.png

图 18-12

使用默认策略的std::fill

图 18-13 显示了一个为所提供的搜索序列中的每个值执行输入序列二分搜索法的例子。作为对搜索序列的第i ^个元素的搜索结果，指示在输入序列中是否找到搜索值的布尔值被分配给结果序列的第i ^个个元素。该算法返回一个迭代器，该迭代器指向分配了结果的结果序列的最后一个元素之后的元素。该算法假设输入序列已经由所提供的比较器排序。如果没有提供比较器，那么将使用一个使用operator<来比较元素的函数对象。

前面描述的复杂性强调了我们应该尽可能地利用库函数，而不是编写我们自己的类似算法的实现，这可能需要大量的调试和调优时间。我们可以利用的库的作者通常是他们正在编码的设备架构内部的专家，并且可能访问我们不知道的信息，所以当优化的库可用时，我们应该总是利用它们。

图 18-13 所示的代码示例展示了使用 DPC++ 并行 STL 算法的三个典型步骤:

创建 DPC++ 迭代器。
从现有策略创建命名策略。
调用并行算法。

img/489625_1_En_18_Fig13_HTML.png

图 18-13

使用binary_search

图 18-13 中的示例使用dpstd::binary_search算法根据我们的器件选择在 CPU、GPU 或 FPGA 上执行二分搜索法。

使用 USM 并行 STL

以下示例描述了将并行 STL 算法与 USM 结合使用的两种方式:

通过 USM 指针
通过 USM 分配器

如果我们有一个 USM 分配，我们可以将指向分配起点和终点的指针传递给一个并行算法。务必确保执行策略和分配本身是为同一队列或上下文创建的，以避免运行时出现未定义的行为。

如果相同的分配要由几个算法处理，要么使用有序队列，要么在下一个算法中使用相同的分配之前明确等待每个算法完成(这是使用 USM 时的典型操作排序)。同样等待完成后再访问主机上的数据，如图 18-14 所示。

或者，我们可以使用带有 USM 分配器的std::vector，如图 18-15 所示。

img/489625_1_En_18_Fig15_HTML.png

图 18-15

通过 USM 分配器使用并行 STL

img/489625_1_En_18_Fig14_HTML.png

图 18-14

使用带有 USM 指针的并行 STL

DPC++ 执行策略的错误处理

如第五章所述，DPC++ 错误处理模型支持两种类型的错误。对于同步错误，运行时抛出异常，而异步错误只在程序执行期间的指定时间在用户提供的错误处理程序中处理。

对于用 DPC++ 策略执行的并行 STL 算法，处理所有的错误，不管是同步的还是异步的，都是调用者的责任。明确地

算法不会显式抛出异常。
由运行时在主机 CPU 上引发的异常(包括 DPC++ 同步异常)被传递给调用方。
并行 STL 不处理 DPC++ 异步错误，所以必须由调用应用程序处理(如果需要处理的话)。

若要处理 DPC++ 异步错误，必须使用错误处理程序对象创建与 DPC++ 策略关联的队列。预定义的策略对象(default_policy和其他)没有错误处理程序，所以如果我们需要处理异步错误，我们应该创建自己的策略。

摘要

DPC++ 库是 DPC++ 编译器的配套产品。它使用预构建和调优的通用函数和并行模式库，帮助我们解决部分异构应用程序的问题。DPC++ 库允许在内核中显式使用 C++ STL API，它通过并行 STL 算法扩展简化了跨架构编程，并通过自定义迭代器增加了并行算法的成功应用。除了支持熟悉的库(libstdc++、libc++、MSVS)，DPC++ 还提供了对 SYCL 内置函数的全面支持。本章概述了利用他人成果而不是自己编写所有内容的方法，我们应该尽可能使用这种方法来简化应用程序开发，并经常实现卓越的性能。

Creative Commons

十九、内存模型和原子

img/489625_1_En_19_Figa_HTML.gif

如果我们想成为优秀的并行程序员，内存一致性并不是一个深奥的概念。它是我们难题的关键部分，帮助我们确保数据在我们需要的时候出现在我们需要的地方，并且它的值是我们所期望的。这一章揭示了我们需要掌握的关键东西，以确保我们的程序正确运行。这个主题不是 SYCL 或 DPC++ 所独有的。

对编程语言的内存(一致性)模型有一个基本的理解对于任何想要允许对内存进行并发更新的程序员来说都是必要的(无论这些更新来自同一个内核中的多个工作项、多个设备，还是两者都有)。不管内存是如何分配的，这都是正确的，无论我们选择使用缓冲区还是 USM 分配，本章的内容对我们来说都是同样重要的。

在前面的章节中，我们已经关注了简单内核的开发，其中程序实例要么对完全独立的数据进行操作，要么使用结构化的通信模式共享数据，这些模式可以使用语言和/或库功能直接表达。随着我们向编写更复杂和更现实的内核迈进，我们很可能会遇到程序实例需要以更少结构化的方式进行通信的情况——理解内存模型如何与 DPC++ 语言特性和我们所针对的硬件功能相关联是设计正确、可移植和高效程序的必要前提。

标准 C++ 的内存一致性模型足以编写完全在主机设备上执行的应用程序，但 DPC++ 对其进行了修改，以解决在对异构系统进行编程时以及在讨论不能完全映射到 C++ 线程概念的程序实例时可能出现的复杂性。具体来说，我们需要能够

系统中的哪些设备可以访问哪些类型的内存分配:使用缓冲区和 USM。
在内核执行期间防止不安全的并发内存访问(数据竞争):使用屏障和原子。
启用执行相同内核的程序实例之间的安全通信和不同设备之间的安全通信:使用屏障、栅栏、原子、内存顺序和内存范围。
防止可能会以不符合我们预期的方式改变并行应用程序行为的优化:使用屏障、栅栏、原子、内存顺序和内存范围。
启用依赖于程序员意图知识的优化:使用内存顺序和内存范围。

内存模型是一个复杂的话题，但有一个很好的理由——处理器架构师关心让处理器和加速器尽可能高效地执行我们的代码！在这一章中，我们努力打破这种复杂性，突出最关键的概念和语言特征。这一章让我们不仅对内存模型了如指掌，还享受了很多人不知道的并行编程的一个重要方面。如果在阅读了这里的描述和示例代码后仍有疑问，我们强烈建议访问本章末尾列出的网站或参考 C++、SYCL 和 DPC++ 语言规范。

内存模型中有什么？

本节详细阐述了编程语言包含内存模型的动机，并介绍了并行程序员应该熟悉的几个核心概念:

数据竞争和同步
栅栏和围栏
原子操作
内存排序

要理解这些概念在 C++、SYCL 和 DPC++ 中的表达和用法，从高层次理解它们是必要的。在并行编程，尤其是使用 C++ 方面有丰富经验的读者可能希望跳过这一步。

数据竞争和同步

我们在程序中编写的操作通常不会直接映射到单个硬件指令或微操作。一个简单的加法操作，如data[i] += x，可以分解成一系列的指令或微操作:

将data[i]从内存加载到一个临时寄存器中。
计算x和data[i].相加的结果
将结果存储回data[i].

这不是我们在开发顺序应用程序时需要担心的事情——加法的三个阶段将按照我们期望的顺序执行，如图 19-1 所示。

img/489625_1_En_19_Fig1_HTML.png

图 19-1

data[i] += x的顺序执行分为三个独立的操作

切换到并行应用程序开发带来了额外的复杂性:如果我们有多个操作并发地应用于相同的数据，我们如何确定他们对该数据的观点是一致的？考虑图 19-2 所示的情况，其中data[i] += x的两次执行已经交错。如果两次执行使用不同的i值，应用程序将正确执行。如果它们使用相同的i值，两者都从内存中加载相同的值，并且一个结果被另一个覆盖！这只是调度它们的操作的许多可能方式之一，我们的应用程序的行为取决于哪个程序实例先获得哪个数据——我们的应用程序包含一个数据竞争。

img/489625_1_En_19_Fig2_HTML.png

图 19-2

并发执行的data[i] += x的一个可能交错

图 19-3 中的代码和图 19-4 中的输出显示了这在实践中是多么容易发生。如果M大于等于N，则j在每个程序实例中的值是唯一的；否则，j的值将会冲突，更新可能会丢失。我们说可能会丢失,因为包含数据竞争的程序仍然可以在某些时候或所有时间产生正确的答案(取决于实现和硬件如何调度工作)。编译器和硬件都不可能知道这个程序打算做什么，或者N和M的值在运行时可能是什么——作为程序员，我们有责任了解我们的程序是否可能包含数据竞争，以及它们是否对执行顺序敏感。

img/489625_1_En_19_Fig4_HTML.png

图 19-4

图 19-3 中N和M的小值代码输出示例

img/489625_1_En_19_Fig3_HTML.png

图 19-3

包含数据竞争的内核

一般来说，当开发大规模并行应用程序时，我们不应该关心单个工作项执行的确切顺序——希望有数百个(或数千个！)并发执行的工作项，试图对它们强加特定的顺序将会对可伸缩性和性能产生负面影响。相反，我们的重点应该是开发正确执行的可移植应用程序，这可以通过向编译器(和硬件)提供有关程序实例何时共享数据、共享发生时需要什么保证以及哪些执行顺序是合法的信息来实现。

大规模并行应用程序不应该关心单个工作项执行的确切顺序！

栅栏和围栏

防止同一组中的工作项之间的数据竞争的一种方法是使用工作组屏障和适当的内存栅栏在不同的程序实例之间引入同步。我们可以使用一个工作组屏障来排序我们对data[i]的更新，如图 19-5 所示，我们示例内核的更新版本如图 19-6 所示。请注意，因为工作组屏障不同步不同组中的工作项，所以只有将我们自己限制在一个工作组中，我们的简单示例才能保证正确执行！

img/489625_1_En_19_Fig6_HTML.png

图 19-6

使用屏障避免数据竞争

img/489625_1_En_19_Fig5_HTML.png

图 19-5

被屏障隔开的两个data[i] += x实例

虽然使用屏障来实现这种模式是可能的，但是通常并不鼓励这样做——它会强制一个组中的工作项按照特定的顺序依次执行，这可能会导致在出现负载不平衡的情况下长时间处于不活动状态。它也可能引入比严格必要的更多的同步——如果不同的程序实例碰巧使用不同的i值，它们仍将被迫在关卡处同步。

屏障同步是一个有用的工具，可以确保工作组或子组中的所有工作项在进入下一个阶段之前完成内核的某个阶段，但是对于细粒度(并且可能依赖于数据)同步来说太笨重了。对于更一般的同步模式，我们必须关注原子操作。

原子操作

原子操作允许对存储器位置的并发访问，而不会引入数据竞争。当多个原子操作访问同一个内存时，保证它们不会重叠。请注意，如果只有一些访问使用原子性，那么这种保证就不适用，作为程序员，我们有责任确保我们不会使用具有不同原子性保证的操作同时访问相同的数据。

同时在相同的内存位置上混合原子和非原子操作会导致未定义的行为！

如果我们的简单加法是用原子操作来表达的，结果可能如图 19-8 所示——每次更新现在都是一个不可分割的工作块，我们的应用程序将总是产生正确的结果。相应的代码如图 19-7 所示——我们将在本章后面重新讨论atomic_ref类及其模板参数的含义。

img/489625_1_En_19_Fig8_HTML.png

图 19-8

与原子操作同时执行的data[i] += x的交错

img/489625_1_En_19_Fig7_HTML.png

图 19-7

使用原子操作避免数据竞争

但是，需要注意的是，这仍然只是一种可能的执行顺序。使用原子操作保证了两个更新不会重叠(如果两个实例使用相同的值i)，但是仍然不能保证两个实例中的哪一个将首先执行。更重要的是，对于不同程序实例中的任何非原子操作，无法保证这些原子操作的顺序。

内存排序

即使在顺序应用中，如果优化编译器和硬件不改变应用的可观察行为，它们也可以自由地重新排序操作。换句话说，应用程序必须像程序员编写的那样运行。

不幸的是，这种假设保证不足以帮助我们对并行程序的执行进行推理。我们现在有两个重新排序的来源要担心:编译器和硬件可能会重新排序每个顺序程序实例中语句的执行，程序实例本身可能会以任何顺序(可能是交错的)执行。为了设计和实现程序实例之间的安全通信协议，我们需要能够约束这种重新排序。通过向编译器提供关于我们期望的内存顺序的信息，我们可以防止与我们的应用程序的预期行为不兼容的重新排序优化。

三种常用的内存排序是

一个放松的记忆排序
一个获取-释放或释放-获取内存排序
一个顺序一致的内存排序

在宽松的存储器排序下，存储器操作可以被重新排序而没有任何限制。宽松内存模型最常见的用法是增加共享变量(例如，单个计数器、直方图计算期间的值数组)。

在获取-释放内存排序下，一个程序实例释放一个原子变量，而另一个程序实例获取相同的原子变量，这充当这两个程序实例之间的同步点，并保证由释放实例发出的任何先前对内存的写入对获取实例可见。非正式地，我们可以认为原子操作将其他内存操作的副作用释放给其他程序实例，或者获取内存操作对其他程序实例的副作用。如果我们希望通过内存在程序实例对之间传递值，就需要这样的内存模型，这可能比我们想象的更常见。当一个程序获得一个锁时，它通常会继续执行一些额外的计算，并在最终释放锁之前修改一些内存——只有锁变量会被自动更新，但我们希望由锁保护的内存更新能够避免数据竞争。这种行为依赖于获取-释放内存顺序来保证正确性，试图使用宽松的内存顺序来实现锁是行不通的。

在顺序一致的存储器排序下，获取-释放排序的保证仍然成立，但是另外存在所有原子操作的单个全局顺序。这种内存排序的行为是三种行为中最直观的，也是最接近我们在开发顺序应用程序时习惯依赖的原始假设保证的行为。有了顺序一致性，推理程序实例组(而不是对)之间的通信就变得容易多了，因为所有程序实例必须在所有原子操作的全局排序上达成一致。

了解编程模型和设备的组合支持哪些内存顺序是设计可移植并行应用程序的必要部分。明确描述我们的应用程序所需的内存顺序，可以确保当我们所需的行为不受支持时，它们会以可预测的方式失败(例如，在编译时)，并防止我们做出不安全的假设。

内存模型

到目前为止，本章已经介绍了理解内存模型所需的概念。本章的剩余部分详细解释了内存模型，包括

如何表达我们内核的内存排序需求
如何查询特定设备支持的内存顺序
关于不相交的地址空间和多个设备，存储器模型如何表现
内存模型如何与障碍、栅栏和原子相互作用
缓冲区和 USM 之间原子操作的使用有何不同

内存模型基于标准 C++ 的内存模型，但在一些重要方面有所不同。这些差异反映了我们的长期愿景，即 DPC++ 和 SYCL 应该有助于为未来的 C++ 标准提供信息:类的默认行为和命名与 C++ 标准库紧密结合，旨在扩展标准 C++ 功能，而不是限制它。

图 19-9 中的表格总结了不同的内存模型概念如何在标准 C++ (C++11、C++14、C++17、C++20)与 SYCL 和 DPC++ 中作为语言特性公开。C++14、C++17 和 C++20 标准还包括一些影响 C++ 实现的说明。这些澄清不应该影响我们编写的应用程序代码，所以我们在这里不涉及它们。

img/489625_1_En_19_Fig9_HTML.png

图 19-9

比较标准 C++ 和 SYCL/DPC++ 内存模型

`memory_order`枚举类

内存模型通过memory_order枚举类的六个值公开了不同的内存顺序，这些值可以作为参数提供给栅栏和原子操作。为一个操作提供一个内存顺序参数，告诉编译器相对于该操作的所有其他内存操作*(对任何地址)需要什么样的内存顺序保证，如下所述:*

memory_order::relaxed

读写操作可以在操作之前或之后重新排序，没有任何限制。没有订购保证。
memory_order::acquire

程序中出现在操作之后的读和写操作必须发生在该操作之后(即，它们不能在操作之前重新排序)。
memory_order::release

出现在程序中的操作之前的读和写操作必须发生在它之前(即，它们不能在操作之后被重新排序)，并且之前的写操作保证对于已经被相应的获取操作(即，使用相同变量和memory_order::acquire或屏障函数的原子操作)同步的其他程序实例是可见的。
memory_order::acq_rel

该操作同时充当获取和释放。读取和写入操作不能围绕操作重新排序，必须使之前的写入可见，如之前针对memory_order::release所述。
memory_order::seq_cst

根据是读、写还是读-修改-写操作，该操作分别充当获取、释放或两者。具有这种记忆顺序的所有操作都是以连续一致的顺序观察的。

每个操作支持的内存顺序有几个限制。图 19-10 中的表格总结了哪些组合是有效的。

img/489625_1_En_19_Fig10_HTML.png

图 19-10

用memory_order支持原子操作

加载操作不将值写入内存，因此与释放语义不兼容。类似地，存储操作不从内存中读取值，因此与获取语义不兼容。其余的读-修改-写原子操作和栅栏与所有存储器排序兼容。

MEMORY ORDER IN C++

C++ 内存模型还包括memory_order::consume，其行为类似于memory_order::acquire。然而，C++17 标准不鼓励使用它，指出它的定义正在被修改。因此，它在 DPC++ 中的包含被推迟到未来的版本中。

`memory_scope`枚举类

标准 C++ 内存模型假设应用程序在具有单一地址空间的单一设备上执行。这些假设对于 DPC++ 应用程序都不成立:应用程序的不同部分在不同的设备上执行(即，主机设备和一个或多个加速器设备)；每个设备具有多个地址空间(即私有、本地和全局)；并且每个设备的全局地址空间可能不相交，也可能不相交(取决于 USM 支持)。

为了解决这个问题，DPC++ 扩展了 C++ 的内存顺序概念，以包括原子操作的范围，表示给定内存顺序约束适用的最小工作项集。范围集是通过一个memory_scope枚举类来定义的:

memory_scope::work_item

内存排序约束仅适用于调用工作项。这个作用域只对图像操作有用，因为一个工作项中的所有其他操作已经保证按程序顺序执行。
memory_scope::sub_group, memory_scope::work_group

内存排序约束仅适用于与调用工作项在同一子组或工作组中的工作项。
memory_scope::device

内存排序约束仅适用于在与调用工作项相同的设备上执行的工作项。
memory_scope::system

内存排序约束适用于系统中的所有工作项目。

除了设备能力所施加的限制，所有内存范围都是所有原子和隔离操作的有效参数。但是，在以下三种情况下，范围参数可能会自动降级到更窄的范围:

如果一个原子操作更新了工作组本地内存中的一个值，那么任何比memory_scope::work_group更宽的范围都会变窄(因为本地内存只对同一工作组中的工作项可见)。
如果一个设备不支持 USM，指定memory_scope::system总是等同于memory_scope::device(因为多个设备不能并发访问缓冲区)。
如果一个原子操作使用memory_order::relaxed，没有排序保证，内存范围参数实际上被忽略了。

查询设备功能

为了确保与以前版本的 SYCL 支持的设备兼容并最大限度地提高可移植性，DPC++ 支持 OpenCL 1.2 设备和其他可能不支持完整 C++ 内存模型的硬件(例如，某些类别的嵌入式设备)。DPC++ 提供设备查询来帮助我们推断系统中可用设备支持的内存顺序和内存范围:

atomic_memory_order_capabilities

atomic_fence_order_capabilities

返回特定设备上原子和隔离操作支持的所有内存排序的列表。要求所有设备至少支持memory_order::relaxed，要求主机设备支持所有内存排序。
atomic_memory_scope_capabilities

atomic_fence_scope_capabilities

返回特定设备上原子和隔离操作支持的所有内存范围的列表。要求所有设备至少支持memory_order::work_group，要求主机设备支持所有内存范围。

起初可能很难记住功能和设备能力的哪些组合支持哪些存储器顺序和范围。在实践中，我们可以通过遵循下面概述的两种开发方法之一来避免这种复杂性:

开发具有顺序一致性和系统防护的应用程序。

仅考虑在性能调优期间采用不太严格的内存顺序。
用宽松的一致性和工作组界限开发应用程序。

只有在正确性需要时，才考虑采用更严格的内存顺序和更宽的内存范围。

第一种方法确保所有原子操作和栅栏的语义与标准 C++ 的默认行为相匹配。这是最简单、最不容易出错的选项，但是具有最差的性能和可移植性。

第二种方法更符合以前版本的 SYCL 和 OpenCL 等语言的默认行为。虽然更复杂——因为它要求我们更加熟悉不同的内存顺序和范围——但它确保了我们编写的大部分 DPC++ 代码可以在任何设备上工作，而不会影响性能。

栅栏和围栏

到目前为止，本书中所有以前使用的障碍和栅栏都忽略了记忆顺序和范围的问题，依赖于默认行为。

DPC++ 中的每个组屏障都充当调用工作项可访问的所有地址空间的获取-释放栅栏，并使之前的写入至少对同一组中的所有其他工作项可见。这确保了一个障碍后一组工作项中的内存一致性，符合我们对同步含义的直觉(以及 C++ 中的 synchronizes-with 关系的定义)。

atomic_fence函数给了我们比这更细粒度的控制，允许工作项以指定的内存顺序和范围执行栅栏。在 DPC++ 的未来版本中，组栅栏可能同样接受可选参数来调整与栅栏相关联的获取-释放栅栏的内存范围。

DPC++ 中的原子操作

DPC++ 支持对各种数据类型的多种原子操作。所有器件都保证支持常见操作的原子版本(例如加载、存储、算术运算符)，以及实现无锁算法所需的原子比较和交换操作。该语言为所有基本整数、浮点和指针类型定义了这些操作，所有设备都必须支持 32 位类型的这些操作，但 64 位类型的支持是可选的。

`atomic`类

C++11 的std::atomic类提供了一个创建和操作原子变量的接口。原子类的实例拥有自己的数据，不能移动或复制，只能使用原子操作进行更新。这些限制大大减少了错误使用该类和引入未定义行为的机会。不幸的是，它们也阻止了该类在 DPC++ 内核中的使用——不可能在主机上创建原子对象并将它们传输到设备上！我们可以继续在我们的主机代码中使用std::atomic,但是试图在设备内核中使用它会导致编译错误。

ATOMIC CLASS DEPRECATED IN SYCL 2020 AND DPC++

SYCL 1.2.1 规范包含一个cl::sycl::atomic类，它松散地基于 C++11 的std::atomic类。我们笼统地说为，因为这两个类的接口之间存在一些差异，最明显的是 SYCL 1.2.1 版本不拥有自己的数据，默认采用宽松的内存排序。

DPC++ 完全支持cl::sycl::atomic类，但是为了避免混淆，不鼓励使用它。我们建议使用atomic_ref类(将在下一节中介绍)来代替它。

`atomic_ref`类

C++20 的std::atomic_ref类为原子操作提供了另一个接口，它比std::atomic提供了更大的灵活性。这两个类最大的区别是std::atomic_ref的实例不拥有它们的数据，而是从一个现有的非原子变量中构造的。创建原子引用实际上是一种承诺，即在引用的生命周期内，被引用的变量只能被原子地访问。这些正是 DPC++ 所需要的语义，因为它们允许我们在主机上创建非原子数据，将这些数据传输到设备上，并且只有在传输之后才将其视为原子数据。因此，DPC++ 内核中使用的atomic_ref类是基于std::atomic_ref的。

我们说是基于，因为该类的 DPC++ 版本包括三个额外的模板参数，如图 19-11 所示。

img/489625_1_En_19_Fig11_HTML.png

图 19-11

atomic_ref类的构造器和静态成员

如前所述，不同的 DPC++ 设备的功能各不相同。为 DPC++ 的原子类选择默认行为是一个困难的命题:默认标准 C++ 行为(即memory_order::seq_cst, memory_scope::system)将代码限制为只能在最有能力的设备上执行；另一方面，打破 C++ 惯例，默认使用最小公分母(即memory_order::relaxed, memory_scope::work_group)可能会导致在迁移现有 C++ 代码时出现意外行为。DPC++ 采用的设计提供了一个折衷方案，允许我们将我们想要的默认行为定义为对象类型的一部分(使用DefaultOrder和DefaultScope模板参数)。其他排序和作用域可以作为运行时参数提供给我们认为合适的特定原子操作——DefaultOrder和DefaultScope只影响我们没有或不能覆盖默认行为的操作(例如，当使用像+=这样的简写操作符时)。最后一个模板参数表示被引用对象所分配的地址空间。

原子引用根据其引用的对象类型为不同的操作提供支持。图 19-12 显示了所有类型支持的基本操作，提供了将数据自动移入和移出内存的能力。

img/489625_1_En_19_Fig12_HTML.png

图 19-12

用atomic_ref对所有类型进行基本操作

对整型和浮点型对象的原子引用扩展了可用的原子操作集，以包括算术运算，如图 19-13 和 19-14 所示。要求设备支持原子浮点类型，而不管它们在硬件中是否具有对浮点原子的本机支持，并且许多设备被期望使用原子比较交换来模拟原子浮点加法。这种模拟是在 DPC++ 中提供性能和可移植性的一个重要部分，只要算法需要，我们就可以在任何地方自由使用浮点原子——生成的代码将在任何地方都能正确工作，并将受益于浮点原子硬件的未来改进，而无需任何修改！

img/489625_1_En_19_Fig14_HTML.png

图 19-14

使用atomic_ref的附加操作仅适用于浮点类型

img/489625_1_En_19_Fig13_HTML.png

图 19-13

使用atomic_ref的附加操作仅适用于整型

使用带缓冲区的原子

正如上一节所讨论的，在 DPC++ 中没有办法分配原子数据并在主机和设备之间移动它。要将原子操作与缓冲区结合使用，我们必须创建一个非原子数据的缓冲区，然后通过原子引用访问该数据。

img/489625_1_En_19_Fig15_HTML.png

图 19-15

通过显式创建的atomic_ref访问缓冲区

图 19-15 中的代码是一个使用显式创建的原子引用对象在 DPC++ 中表达原子性的例子。缓冲区存储普通整数，我们需要一个具有读写权限的访问器。然后，我们可以为每个数据访问创建一个atomic_ref实例，使用+=操作符作为fetch_add成员函数的简写替代。

如果我们希望在同一个内核中混合对缓冲区的原子和非原子访问，以避免在不需要原子操作时支付原子操作的性能开销，这种模式非常有用。如果我们知道缓冲区中只有一个内存位置的子集将被多个工作项同时访问，那么我们只需要在访问那个子集时使用原子引用。或者，如果我们知道同一工作组中的工作项仅在内核的一个阶段(即，在两个工作组屏障之间)并发访问本地存储器，那么我们只需要在该阶段使用原子引用。

有时我们很乐意为每次访问支付原子性的开销，要么是因为每次访问都必须是原子性的，以保证正确性，要么是因为我们对生产率比对性能更感兴趣。对于这种情况，DPC++ 提供了一种声明访问器必须总是使用原子操作的简写方式，如图 19-16 所示。

img/489625_1_En_19_Fig16_HTML.png

图 19-16

通过原子访问器隐式创建的atomic_ref访问缓冲区

缓冲区像以前一样存储普通的整数，但是我们用一个特殊的atomic_accessor类型替换了常规的访问器。使用这种原子访问器会自动使用原子引用包装缓冲区的每个成员，从而简化内核代码。

直接使用原子引用类还是通过访问器使用原子引用类是最好的，这取决于我们的用例。为了在原型开发和初始开发过程中简单起见，我们的建议是从访问器开始，只有在性能调优过程中有必要时(例如，如果分析显示原子操作是性能瓶颈)或者只有在定义明确的内核阶段(例如，在本章后面的直方图代码中)才需要原子性时，才使用更显式的语法。

在统一共享内存中使用原子

如图 19-17 (转载自图 19-7 )所示，我们可以从 USM 中存储的数据中构造原子引用，就像我们对缓冲区所做的一样。事实上，这段代码与图 19-15 所示代码的唯一区别在于 USM 代码不需要缓冲区或存取器。

img/489625_1_En_19_Fig17_HTML.png

图 19-17

通过显式创建的atomic_ref访问 USM 分配

没有办法只使用标准的 DPC++ 特性来模仿原子访问器为 USM 指针提供的速记语法。然而，我们期望 DPC++ 的未来版本将提供一个建立在为 C++23 提出的mdspan类之上的简写。

在现实生活中使用原子

原子的潜在用法如此广泛和多样，以至于我们不可能在本书中提供每种用法的例子。我们提供了两个具有代表性的例子，它们在各个领域都有广泛的适用性:

计算直方图
实现设备范围的同步

计算直方图

图 19-18 中的代码演示了如何使用宽松原子结合工作组障碍来计算直方图。屏障将内核分为三个阶段，每个阶段都有自己的原子性需求。请记住，屏障同时充当同步点和获取-释放栅栏——这确保了一个阶段中的任何读取和写入对于后面阶段中工作组中的所有工作项都是可见的。

第一阶段将一些工作组本地内存的内容设置为零。每个工作组中的工作项通过设计来更新工作组本地内存中的独立位置——不会出现竞争情况，并且不需要原子性。

第二阶段在本地存储器中累积部分直方图结果。同一个工作组中的工作项可能会更新工作组本地内存中的相同位置，但是同步可以推迟到阶段结束时——我们可以使用memory_order::relaxed和memory_scope::work_group来满足原子性需求。

第三阶段将部分直方图结果贡献给存储在全局存储器中的总数。相同工作组中的工作项保证从工作组本地内存中的独立位置读取，但是可以更新全局内存中的相同位置——我们不再需要工作组本地内存的原子性，并且可以像以前一样使用memory_order::relaxed和memory_scope::system来满足全局内存的原子性要求。

img/489625_1_En_19_Fig18_HTML.png

图 19-18

使用不同存储空间中的原子引用计算直方图

实现设备范围的同步

回到第四章，我们警告过不要编写试图跨工作组同步工作项目的内核。然而，我们完全期望本章的几个读者将渴望在原子操作之上实现他们自己的设备范围的同步例程，并且我们的警告将被忽略。

设备范围的同步目前是不可移植的，最好留给专业程序员。该语言的未来版本将解决这个问题。

本节中讨论的代码是危险的，不应该期望在所有设备上都能工作，因为在调度和并发保证方面存在潜在的差异。由原子提供的存储器排序保证与前向进度保证正交；而且，在撰写本文时，SYCL 和 DPC++ 中的工作组调度完全是由实现定义的。形式化讨论执行模型和调度保证所需的概念和术语是当前活跃的学术研究领域，DPC++ 的未来版本有望在此基础上提供额外的调度查询和控制。目前，这些话题应该被认为是专家专有的。

图 19-19 显示了一个简单的设备范围闩锁(一个一次性栅栏)的实现，图 19-20 显示了其使用的一个简单例子。每个工作组选择一个工作项来通知该组到达闩锁处，并使用一个简单的旋转循环等待其他组的到达，而其他工作项使用工作组屏障等待所选择的工作项。正是这种自旋循环使得设备范围的同步不安全；如果任何工作组还没有开始执行，或者当前正在执行的工作组没有得到公平的调度，代码可能会死锁。

在没有独立的前向进度保证的情况下，仅仅依靠内存顺序来实现同步原语可能会导致死锁！

为了使代码正确工作，必须满足以下三个条件:

原子操作必须使用至少与所示一样严格的内存顺序，以保证生成正确的栅栏。
ND 范围中的每个工作组必须能够向前进展，以避免循环中旋转的单个工作组使尚未递增计数器的工作组饥饿。
The device must be capable of executing all work-groups in the ND-range concurrently, in order to ensure that all work-groups in the ND-range eventually reach the latch.

图 19-20

使用图 19-19 中的全设备锁

图 19-19

在原子引用的基础上构建一个简单的设备级锁存器

虽然不能保证这段代码是可移植的，但我们在这里包含它是为了强调两个要点:1) DPC++ 的表达能力足以支持特定于设备的调优，有时会牺牲可移植性；以及 2) DPC++ 已经包含了实现高级同步例程所必需的构件，这些构件可能包含在该语言的未来版本中。

摘要

本章提供了对内存模型和原子类的高级介绍。了解如何使用(以及如何不使用！)这些类是开发正确的、可移植的和高效的并行程序的关键。

内存模型是一个极其复杂的主题，我们在这里的重点是为编写真正的应用程序建立一个基础。如果需要更多的信息，有几个网站、书籍和讲座专门介绍下面提到的内存模型。

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