PS：最近不务正业对GPU产生了浓厚的兴趣，可能是AI需要软硬结合的原因？对这本书（2017年出版不算太旧）进行整理，以备不时之需。

# 前言

众所周知，由于架构和能源的限制，将多核心中央处理器（CPU）的核心数量增加到数千个是不可能的，而图形处理器（GPU）可以轻松地利用数千个简单有效的核心，每个核心执行单个指令线程。

对于图形应用程序对计算和带宽的巨大需求已经导致GPU成为主要的大规模并行架构。如今，具有浮点运算能力的通用GPU（GPGPU）在具有混合CPU-GPU架构的大多数顶级机器上广泛用于通用计算。这样的GPGPU新产品NVIDIA Pascal拥有3840个核心。

这样强大的计算芯片的主要问题是管理如此庞大的核心数量。为了满足未来的需求，需要新的架构技术来进一步增加核心数量，同时考虑影响其可靠性，性能和能耗特性的技术限制。而缺乏有效的编程模型和系统软件以利用这样大量的核心在实际应用中的全部性能，是最具挑战性的问题。

本书聚焦于基于GPU的系统的研究和实践，并试图解决它们的重要问题。涵盖的主题涵盖从硬件和架构问题到直接涉及应用程序或系统用户的高级问题，包括并行编程工具和中间件支持的等等这种计算系统。本书分为四个部分，每个部分都包含由知名研究人员撰写的章节，他们在该部分的主题下分享了他们的最新研究成果。

第1部分涉及不同的编程问题和工具。该部分由5个章节组成：1-5。第1章讨论了如何以可靠的方式编写GPU程序。尽管GPU程序在不同领域提供高计算吞吐量，但正确优化代码有些困难，这是由于GPU并发性的微妙性。因此，本章的重点是讨论并解决一些事情，以使GPU编程更容易。该章节还提供了一些关于GPU软件形式分析的最新进展。后者是一个重要问题，随着GPU编程的进步。因此，新的GPU程序员应该学习验证方法。

第2章介绍了异构集群的统一开放计算语言（OpenCL）框架（SnuCL）。它是一个免费的，开源的OpenCL框架，用于异构集群。基本上，OpenCL是一种针对异构并行计算系统的编程模型，它定义了传统处理器和加速器之间的抽象层。使用OpenCL的优点是程序员编写一个OpenCL应用程序，然后在任何支持OpenCL的系统上运行它。然而，它缺乏针对异构集群的目标能力。

SnuCL为程序员提供了一个单一、统一的OpenCL平台映像，用于集群。借助SnuCL，OpenCL应用程序能够利用计算节点中的计算设备，就像它们在主机节点中一样。此外，使用SnuCL，可以将来自不同供应商的多个OpenCL平台集成到单个平台中。因此，OpenCL对象在不同供应商的计算设备之间共享，能够实现异构系统的高性能和易编程性。

第3章探讨了在大规模并行GPU上的线程通信和同步。在GPU上，数千个线程同时运行，处理器的性能很大程度上取决于线程间通信和同步的效率。了解现代GPU支持哪些机制以及它们对算法设计的影响是编写有效GPU代码的关键问题。由于传统的GPGPU工作量是大规模并行的，线程之间几乎没有协作，早期的GPU只支持粗粒度的线程通信和同步。如今，当前的趋势是加速更多的多样化工作负载，这意味着粗粒度机制是利用并行性的主要限制因素。最新的工业标准编程框架OpenCL 2.0引入了细粒度线程通信和同步支持，以解决这个问题。本章讨论了现代GPU上可用的粗粒度和细粒度线程同步和通信机制。

第4章着重讨论GPU上的软件级任务调度。为了充分利用众核处理器的潜力，任务调度是一个关键问题。与CPU相反，GPU缺乏程序员或编译器控制调度的必要API。因此，使用现代GPU上的硬件调度器以灵活的方式是困难的。本章介绍了一个编译器和运行时框架，以自动转换和优化GPU程序，以实现对流式多处理器（SM）的可控任务调度。框架的中心是以SM为中心的转换，它解决了硬件调度器的复杂性并提供了调度能力。该框架为新的优化提供了许多机会，其中本章介绍了三个用于优化并行性、局部性和处理器分区的优化。广泛的实验表明，这些优化可以在多种场景下显著提高一组GPU程序的性能。

第5章研究了GPU上数据放置的复杂性，并介绍了PORPLE，一个软件框架，展示了如何自动解决一个给定的GPU应用程序的复杂性。数据放置的概念是一个重要的问题，因为现代GPU内存系统由许多具有不同特性的组件组成，问题是如何将数据放置在各种内存组件上。

第2部分介绍了一些针对GPU的有用算法和应用。它包括第6-14章。第6章关注生物序列分析。高通量的DNA测序技术导致生物数据库的指数增长。必须分析和解释这些生物序列。

为了确定生物数据的功能和结构，需要进行生物序列分析。然而，生物数据库的增长速度远远超过了单核处理器的性能。随着众核处理器的出现，可以利用生物序列分析工具。本章讨论了GPU在两个主要序列比较问题中的最新进展：序列对比和序列-配置文件比较。

第13章讨论了GPU上的图算法。首先介绍和比较了表示和分析图的主要数据结构和技术。接着讨论了GPU的高效图算法的理论和最新研究。本章所关注的算法主要是遍历算法（广度优先搜索）、单源最短路径（Dijkstra、Bellman-Ford、delta stepping、混合算法）和全源最短路径（Floyd-Warshall）等。随后，本章讨论负载平衡和存储器访问技术，概述其主要问题和管理技术。

第8章考虑了使用GPU对序列进行对齐，探讨了使用GPU对两个和三个序列进行最优对齐的方法。对两个序列进行对齐的问题通常称为序列对齐。该章节还介绍了在NVIDIA Tesla C2050上进行的实验结果。

第9章介绍了GPU上解决三对角系统的增强块Cimmino分布式（ABCD）算法。三对角系统的特殊结构在科学和工程问题中经常出现，例如交替方向隐式方法、流体模拟和泊松方程。本章介绍了在GPU上解决三对角系统的ABCD方法的并行化。在各种方面中，本章探讨了边界填充技术以消除GPU上的执行分支等。还采用了各种性能优化技术，如内存合并，以进一步提高性能。性能评估显示，GPU实现比传统CPU版本提高了超过24倍的速度。

第10章讨论了线性和混合整数规划方法。它表明运营研究（OR）社区中的复杂问题可以从GPU中获益。作者还通过突出不同作者如何克服实现困难，介绍了应用于线性和混合整数规划的GPU计算的主要贡献的调查结果。

第11章考虑了平面图最短路径计算的加速实现。针对两种最短路径问题，描述了三种算法及其相关的GPU实现。第一个算法解决了全源最短路径问题，而第二个算法则为更好的并行扩展属性和支持存储器访问而容易。第三个算法解决了单源最短路径查询问题。实现结果显示，同时利用256个GPU的计算能力。因此，改进间隔比现有的并行方法高出一个数量级。

第12章讨论了GPU上的排序算法。简要介绍了CUDA编程、记忆和计算模型以及通用GPU。

第三部分集中讨论架构和性能问题，其中包括15-18章。在第15章中，引入了一个框架来加速GPU应用程序中的瓶颈。困难的是，由于异构应用程序的要求不同，在GPU中资源的利用存在不平衡，因此在执行过程中出现了不同的瓶颈。在本章中，介绍了一个核心助攻瓶颈加速（CABA）框架，它利用空闲的片上资源来消除GPU执行中的不同瓶颈。 CABA提供灵活的机制来自动生成在GPU核心上执行特定任务的“辅助卷积”。换句话说，它使用空闲的计算单元和管道来缓解内存带宽瓶颈。因此，它提高了GPU的性能和效率。 CABA架构然后被深入地讨论和评估，以在GPU存储器子系统层次结构中有效和灵活地执行数据压缩。结果表明，使用CABA进行数据压缩，平均可提高41.7％（最大2.6倍）的性能，跨越宽范围的内存带宽敏感GPGPU应用程序。

第16章考虑通过神经算法变换来加速GPU。基本上，通过利用可编程网格，可扩展的设计，具有灵活的内存模型和具有高度优化的编译器，实现GPU的高性能增益。因此，算法变换技术可以有效地调整程序以更好地利用GPU资源，从而提高性能和效率。本章重点介绍了Algorithm Transformation Toolkit（ATT），它是一个用于加速GPU应用程序的工具包，该工具包实现了一组算法变换，可以帮助程序员更有效地优化GPU应用程序。 ATT的评估表明，它可以提供高达3.00倍的性能改进，平均改进1.51倍。

实现大量数据级并行和采用单指令多线程（SIMT）执行模型。许多应用领域都受益于GPU，包括识别、游戏、数据分析、天气预测和多媒体等。大多数这些应用都是适合进行近似计算的优秀候选者。因此，采用近似计算技术可以提高GPU的性能和能效。在各种近似技术中，神经加速器能够带来显著的性能和效率提升。本章介绍了一种神经加速的GPU（NGPU）架构，将神经加速嵌入到GPU加速器中，同时不影响它们的SIMT执行。

作者在第17章介绍了一种带有动态带宽分配的异构光子网络芯片，用于GPU。未来的多核芯片预计将拥有数百个异构组件，包括处理节点、分布式内存、定制逻辑、GPU单元和可编程织物。由于未来芯片预计会同时运行多个不同的并行工作负载，因此不同的通信核心将需要不同的带宽。因此，异构网络芯片（NoC）架构的存在是必要的。最近的研究表明，光子互连能够实现高带宽和能效性能较好的芯片内数据传输。本章讨论了一种动态异构光子NoC（d-HetPNOC）架构，采用动态带宽分配，以实现与同质光子NoC架构相比更好的性能和能效。

第一章 可靠GPU编程的形式化分析技术：现有解决方案和后续的改进方向

一、GPUs IN SUPPORT OF PARALLEL COMPUTING

图形处理单元（GPU）是当前应用对强制功耗和有限电线延迟硬件的计算需求的自然产物 [2]。GPU通过使用较简单的核心来实现比CPU更高的计算效率，并通过切换停滞的线程来隐藏内存延迟。

总体而言，GPU面向吞吐量的执行模型非常适合许多数据并行应用程序。

GPU的演进速度是引人注目的：从1997年的3M晶体管Nvidia NV3开始，到2012年的7B晶体管Nvidia GK110（Kepler）[3]。Nvidia的CUDA编程模型于2007年推出，提供了比编写像素和顶点着色器的巴洛克符号更高的一步，而最近的CUDA 7.5 [4]提供了多功能的并发原语。OpenCL是一种业界标准的编程模型 [5]，得到了包括Nvidia、AMD、ARM和Imagination Technologies在内的所有主要设备厂商的支持，对数万个核心的计算意图提供了简单、可移植的映射。

GPU现在远远超出了图形应用程序，成为我们在游戏、网络搜索、基因测序和高性能超级计算等各个领域追求并行性的重要组成部分。

1、并行计算和GPU代码中的错误和不足

编写正确的程序一直是计算机科学的一个基本挑战，即使在图灵 [6] 的时代也是如此。使用CPU线程或消息传递（如MPI）编写的并行程序比顺序程序更容易出错，因为程序员必须编码线程之间的同步和通信逻辑，并安排共享资源（主要是内存）。这种情况导致错误很难定位和修复。

与通用并发程序相比，GPU程序是“尴尬的并行”的，线程受到受控规则的控制和同步。然而，GPU程序提出了某些独特的调试难题，这些挑战尚未得到足够的关注，正如第2和第3节中所详细讨论GPU错误的情况所示。

如果不加检查，GPU错误可能会成为程序无法运行的阻碍，使由价值数百万美元的科学项目产生的模拟结果毫无用处。突然的不可解释的崩溃、不可重复的执行以及不可重复的科学结果，通常被忽视在exascale计算的嘈杂声中。

然而这些错误确实会发生，并且严重地使专家们感到担忧，他们常常艰难地启动新投入的机器，或者需要停止他们正在进行的有用科学研究而进行修复。

本章的目的是描述这些挑战的一些解决方案，提供对正在开发的解决方案的理解，并描述仍需完成的工作内容。

在高层次介绍GPU后（第2节），我们以能够使程序分析和验证社区受益为目的，以一种回顾某些关键GPU正确性问题的方式进行调查（第3节）。问题是：什么是正确性挑战，我们如何从共同努力中受益以解决可扩展性和可靠性问题？然后回答了一个问题：我们如何建立可以处理今天和即将到来的异构并发形式的严格的正确性检查工具？为此，我们讨论了已有的有助于建立正确性的工具，提供了它们如何运作的高层次描述，并总结了它们的限制（第4节）。我们通过呼吁行动结束本章，展现了我们进一步开展这项工作的观点，这需要通过（a）研究驱动的GPU加速软件正确性检查方法的进展来解决开放性问题，(b) 通过传播和技术转移活动来增加该工业分析工具的使用率（第5节）。

2、快速介绍GPU

GPU通常被用作并行协处理器，在一个异构系统中由主机CPU控制。在这种设置下，具有丰富并行性的任务可以作为核函数(offloaded to the GPU)放入GPU内：一个指定任意线程行为的模板。图1展示了一个CUDA核函数，用于执行两个向量的并行点乘，取自CUDA 5.0 SDK [7]，我们将其用作一个示例说明。我们对GPU编程模型进行简要概述。我们介绍每个概念和组件时，给出该概念或组件的CUDA术语，并在括号中用OpenCL术语（如果有不同），之后使用CUDA术语。

2.1 线程的组织

内核在GPU上通过许多轻量级线程(工作项)以一种分层的方式作为一个网格(NDRange)的线程块(work-groups)组织起来执行，如图所示。

这段代码实现了一个向量点积运算。其中ACCUM_N定义了每个线程块中共享内存的大小。主要分为两个部分：

1.计算部分和：在每个线程块内，每个线程都会计算一段连续的部分和，并保存在共享内存的accumResult数组中。具体地，每个线程从自己的tid开始，每次加上一个块内线程数bdim，直到遍历所有ACCUM_N个元素。对于每个元素，内部执行一个循环，将d_A和d_B对应下标的元素相乘，并累加到sum中，最终将sum存储在accumResult对应的位置。

2.归约汇总：在所有线程都计算完部分和后，需要使用归约操作得到整个向量的点积。具体地，每个线程对相邻的两个部分和执行累加操作，最终得到整个向量的点积。归约过程中需要保证线程同步，使用__syncthreads()函数实现。最后，主线程将整个向量的点积存储在d_c中，并返回。

1. 定义一个名为 ACCUM_N 的宏，其值为 1024。

2. 定义一个名为 dotProduct 的全局函数，使用 GPU 加速计算矩阵 A 和矩阵 B 的向量点积，结果存储在 d_c 指针指向的位置。

3. 使用 __shared__ 关键字定义一个名为 accumResult 的共享数组，长度为 ACCUM_N，用来存储每个线程计算出的部分和。

4. 获取当前线程在块内的线程编号。

5. 获取块内的线程数。

6. 计算出当前线程需要计算的元素下标的增量 ACCUM_N。

7. 遍历每个线程需要计算的部分和。

8. 对每个元素进行计算，将结果存储在名为 sum 的变量中。

9. 将计算得到的部分和存储在共享数组 accumResult 中。

10. 完成对每个线程的部分和的计算。

11. 进行规约操作，将所有线程计算出的部分和相加，得到最终的结果。

12. 使用 __syncthreads() 函数同步所有线程。

13. 对于每个 stride，每个线程计算出当前位置的元素和距离当前位置 stride 的位置的元素的和。

14. 将结果存储在 accumResult 中。

15. 完成规约操作。

16. 如果当前线程是编号为 0 的线程，则将最终计算结果存储在 d_c 指针指向的位置中。

例如，使用每个256个线程的4个线程块的网格将产生1024个线程，每个线程都运行核函数的副本。图1中的__global__注释表示dotProduct是一个核函数。

在核函数内部，一个线程可以使用内置函数（如threadIdx（get_local_id）和blockDim（get_local_size））查询其在网格层次结构中的位置（以及网格和线程块的尺寸）。网格和块可以是多维的，例如，blockDim.x（get_local_size（0））和threadIdx.x（get_local_id（0））分别提供线程块的大小和第一维度内的线程的ID。这使线程可以对不同的数据进行操作，并通过核函数跟踪不同的执行路径。

2.2 内存空间

GPU上的线程可以访问多个内存空间中的数据，这些空间按照反映线程组织的层次结构排列，如图2所示。按下降顺序排列，它们是大小、范围-可见性和延迟，它们是：

。在CUDA中，一个网格中的所有线程都能看到的大型全局内存。在图1中，CUDA指针内核参数指向全局内存数组，因此d_a、d_b和d_c是全局数组。内核计算d_A和d_B数组的点积并将结果存储在d_c中（一个单元数组）。

。每个线程块之间可见的线程块共享（本地）内存。图1中的__shared__注释指定accumResult数组驻留在共享存储器中。每个线程块都有此数组的不同副本，用于累加部分点积值。

。一个小的每线程私有内存。图1中的循环变量i、j和stride驻留在私有内存中。每个线程都有这些变量的单独副本。程序员有责任在全局和共享内存空间之间协调数据移动。内存合并是出于性能原因的重要属性。当相邻的线程访问连续的内存位置时，硬件可以将这些访问合成更少的内存事务，从而增加带宽。

虽然线程共享某些内存空间，但内存写入不会立即对所有线程可见（因为这会严重影响性能）。在计算机体系结构中，内存一致性模型的概念用于明确解释线程何时可以观察到其他线程的写入。GPU编程模型指定了一个弱内存一致性模型[4,5]。也就是说，给定线程的更新和访问顺序不能保证由其他线程观察到。

2.3 屏障同步（跟flink的streaming水位机制类似）

屏障可以安全地在同一线程块的线程之间进行通信。

屏障操作会导致线程在所有线程块的所有线程都到达同一屏障之前停顿。实际上，所有线程必须在相同的控制流下达到屏障，才能避免屏障分歧问题，这会导致未定义的行为。点积核使用屏障同步（indicated by __syncthreads()）来确保更新共享数组accumResult的正确顺序。

屏障不能用于不同线程块之间的线程同步。相反，这要求程序员将工作负载分成多个在序列中调用的内核。不同线程块中的线程可以使用原子操作通过全局内存进行通信。

2.4 线程束和锁步执行

在Nvidia GPU上，硬件通过动态分区将线程块分成一组线程束；AMD GPU有类似波前的概念。目前，Nvidia指定一个线程束是线程块的32个相邻线程的集合。同一线程束中的线程以锁步模式执行，因此隐式地进行同步；我们在第3节中对此现象的机会和风险进行评论。

2.5 点积的例子

现在我们已经有了适当的装备来讨论图1中的内核。考虑使用128个线程块（即blockDim.x = 128）调用内核时，输入数组的长度n等于4096。

内核有两个并行阶段。在第一阶段，部分积累到共享数组accumResult中。外部循环为accumResult数组的每个元素分配一个不同的线程。由于元素(ACCUM_N)比线程多，外循环为每个线程分配了ACCUM_N/blockDim.x = 8个元素。例如，线程0分配元素0、128、256、…、896。输出每个元素i的结果在线程私有变量sum中累加，并由内循环累加。内循环在ACCUM_N间隔内执行n /ACCUM_N = 4个部分积。例如，当i = 0时，线程0的内循环将计算{j≣0、1024、2048、3072}中的∑aj bj 的积。存取的步长stride确保内存在全局存储器中访问是合并的。

在第二个阶段，部分积被缩减为最终期望结果。内核使用并行树缩减，而不是逐个序列求和ACCUM_N元素。缩减是使用逻辑树执行的。给出了8个元素的简化缩减树，如图3所示。循环的每次迭代——使用降幂二步长值——对应于树的不同级别。屏障确保给定级别的更新在下一级的任何访问之前有序，并因此成为线程间通信的一种形式。

3、GPU编程中的正确性问题

现在我们概述四个关键的正确性问题，这些问题可能使从CPU到GPU的转换困难。

3.1 数据竞争

不足或错误的障碍物barrier可能导致数据竞争。如果两个线程访问相同的内存位置（在全局或共享内存中），其中至少一个访问是非原子访问，至少一个访问修改该位置，并且没有涉及两个线程的干预屏障同步，那么就会发生数据竞争（在OpenCL 2.0中，某些原子操作之间的同步也可以用于避免竞速;我们在这里不深入讨论细节）。例如，考虑图1中点积内核的原地树规约。第20行的屏障确保给定迭代中（当stride = k时，例如）的所有访问在下一次迭代（当stride = k / 2时）的任何访问之前排序。如果省略该屏障，则可能发生数据竞争。例如，带有stride 2的线程0和带有stride 4的线程2将在累积结果[8]上进行竞争。

在大多数并发程序中，竞争是某些代码严重问题的明显指标，包括由于不充分的原子性而可能导致不确定结果或语义上不连贯的更新。 GPU编程模型（例如，OpenCL）要求程序员编写没有数据竞争的代码。包含数据竞争的程序具有未定义的语义。因此，许多编译器优化假定无竞争性;在存在竞争的情况下，它们可能会产生意外的结果。

可以通过保守的屏障放置来防止数据竞争，但屏障的过度使用可能有两个问题。首先，屏障的执行成本很高，因此不必要的同步引入不必要的性能开销。第二，将屏障放置在条件代码中可能是危险的，因为它涉及到第2节中讨论的屏障分歧问题。

在Nvidia GPU上，许多CUDA程序员选择在仅需要同步同一warp中的线程之间时省略屏障。例如，点积内核中第二阶段的树规约循环的最后六次迭代（请参见图1）可能被替换为以下语句序列，以避免当stride小于或等于32时的显式同步：

if(tid < 32) accumResult[tid] += accumResult[tid + 32];

if(tid < 16) accumResult[tid] += accumResult[tid + 16];

if(tid < 8) accumResult[tid] += accumResult[tid + 8];

if(tid < 4) accumResult[tid] += accumResult[tid + 4];

if(tid < 2) accumResult[tid] += accumResult[tid + 2];

if(tid < 1) accumResult[tid] += accumResult[tid + 1];

此做法依赖于编译器在优化过程中保留隐式的warp内同步。目前尚不清楚是否存在这种情况。在CUDA编程指南5.0版中，有关warp内同步的建议已被删除，并且从业者对当前平台在这方面提供的保证存在分歧（例如，有关该主题的NVIDIA论坛讨论[9]）。

尽管如此，一些随CUDA 5.0 SDK一起提供的示例依赖于隐式warp内同步。更令人惊讶的是，开源基准套件Parboil [10]和SHOC [11]中的OpenCL内核删除了前面提到的屏障。这显然是错误的，因为锁步warp的概念并不是OpenCL规范的一部分。

最近的研究表明，Nvidia和AMD GPU表现出弱内存行为（如第2节所讨论的），可以观察到非顺序一致的执行方式，这是引起微妙软件缺陷的来源[12]。因为弱内存效应而产生的错误将越来越相关，因为GPU应用程序将向利用细粒度并发代替屏障同步移动。

3.2 缺乏前向保证

GPU架构中线程调度不公平是另一个缺陷的来源。图4展示了在CUDA中实现inter-block barrier的尝试。这个众所周知的策略背后的思想如下。每个线程在块内同步（第3行），之后每个线程块的领袖（threadIdx.x = 0的线程）原子性地减少一个计数器（第5行）。假设计数器初始值为总线程块数（在第1行进行注释），那么旋转直到计数器达到零（第6行）可能会足以确保每个领袖已通过减量（第5行）。 （请注意，计数器值是通过向计数器原子添加零来检索的； atomicAdd返回操作的位置的旧值。使用原子操作而不是普通的load操作可以避免访问计数器时的数据竞争。）如果旋转直到计数器达到零确实可以确保每个领袖已通过减量，则已实现全局同步，领袖可以重新与其块中的其余部分同步（第8行），从而允许执行继续。栅栏（第2和第9行）的目的是确保全局同步点之前的内存访问操作在全局同步点之后的内存访问操作之前生效。关于此策略及其问题的社区讨论，请参见例如Ref. [13]。

这个屏障的问题在于它假定了线程块之间的前进。然而，如果请求足够数量的CUDA块，则它们将被分批调度：一定数量的块将被调度。

并且必须在计算单元释放用于安排更多块之前完全运行。这种情况会导致死锁：在第一轮调度的块的领导者会在第六行处自旋，等待计数器由附加块的领导者递减；这些块反过来将不能被调度，直到初始块完成执行。尝试在OpenCL中实现全局同步的类似方法也无法工作，原因是相同的。

这个例子说明了线程块之间的不公平调度。由于线程块被分成了warp，使得相同warp中的线程以锁步方式执行，执行相同的指令序列，因此在CUDA线程块内的线程之间也会出现缺乏进展的问题。如果程序员试图强制一个线程在同一warp中的另一个线程上忙等待，则会导致死锁；进展需要线程在运行时真正执行不同的指令。这破坏了基于忙等待的块内关键部分的天真实现。有关此问题的示例和社区讨论，请参见参考文献[14]。

3.3 浮点精度

对于在浮点数据上计算的内核而言，实现与参考实现的等价性可能会更具挑战性，这在高性能计算中很常见。

在设计并行算法时，假设浮点算符具有实数的代数性质可能很方便。例如，求和归约操作，其中将数组的所有元素相加，可以通过计算树并行化，如图5所示。如果并行处理元素的数量超过数组的大小，基于树的方法允许在对数步骤中进行归约。该图显示了一个8个元素数组的两个可能的计算树。如果加法满足结合律，即对于所有浮点值x、y和z，满足法律（x + y）+ z = x +（y + z），其中+表示浮点加法，则计算树将产生相同的结果。

然而，正向顺序实现得到的左关联和的结果与并行算法的结果相同。然而，众所周知，浮点加法不可结合，因此结果预计会因并行算法的结构而有所不同。

虽然这些问题也适用于中央处理器，但对于图形处理器可能更为重要的原因在于它们以吞吐量为导向，并且缓存较小。因此，在GPU上使用双精度算术代替几乎所有的单精度算术会增加惩罚，特别是由于GPU上双精度单元的数量通常较少。与GPU相关的浮点精度语言特定问题也存在。例如，在OpenCL中，有关是否准确表示非正规化数的问题是实现定义的，并且OpenCL提供了一种半精度数据类型，仅为相关运算符指定了最小（而不是精确）精度要求。实现差异使编写在多个GPU平台上行为精度可接受的高性能浮点代码变得困难。

（待补充中）