第9章 应用程序性能优化(1) 摘录

本章针对程序的性能优化，分解为如下7种策略：
        1： 理解问题以及分解问题为串行/并行的工作负载。
        2： 理解并优化内存带宽、延迟和缓存使用问题。
        3： 理解与主机端传输数据的玄机。考察锁页内存、零复制内存的表现和某些硬件的带宽限制。
        4： 理解线程结构和计算能力，并了解它们对性能的影响方式。
        5： 结合一些通用算法的优化实例，讨论如何实现算法。
        6： 关注性能分析，定位应用程序的瓶颈所在及其来源。
        7： 考察如何让应用程序根据各种硬件实现自我调优。

9.1 策略1： 并行/串行在GPU/CPU上的问题分解

9.1.1 分析问题

首先需要考虑的是，对问题尝试并行化是否是正确的解决方案。

9.1.2 时间

界定算法执行时间"可接受的"时间段是很重要的。目前可接受的并不一定意味着是最佳优化。

通常会发现任何优化活动都会有一定量所谓的“唾手可得的东西”。在优化中上述这些改变是很容易的且会带来一个合理的加速。当这些被处理掉，就逐渐变得更难找到优化之处，只有采用更复杂的技术才能进一步优化，这不仅花费更多时间而且引入更多潜在错误。

在设置适当的加速目标之后，在给定一组硬件的前提下，你必须知道什么是合理的。

在考虑可接受的时间段是多少，问问你自己要达到该时间段还需付出多少努力。如果没有超过2倍，往往就值得花时间来优化CPU的实现，而不是创建一个全新的、并行的方法来解决该问题。多线程引出了各种与依赖关系、死锁、同步、调试等有关问题。

在过去的30年，购买更快的硬件就可，例如IO瓶颈、内存带宽、计算吞吐量都可。

如果你决定GPU路线，那么通常应该将你的设计目标设置为当前程序执行时间的10倍。你所达到的实际量级取决于编程者的知识、可用的编程时间以及应用程序中的并行程度。最后一个因素具有巨大的决定性。

9.1.3 问题分解

这个问题可以被分解成并行运行的组块吗？如果答案是否定的，GPU则不是要考虑的方法。你需要看CPU优化技术，比如缓存优化、内存优化、SIMD优化等。

假设你能够把问题分为多个并发块，那接下来应该并发多少个并发块？CPU并行化地一个主要限制因素经常是没有足够大的粒度(或粗粒度)的并行工作要做。GPU运行成千上万的线程，所以问题需要被分解成千块，而不只是像CPU那样只执行少数并发任务。

问题分解应该总是先从数据开始，然后考虑执行的任务。你应该试图用输出数据集来表示问题。你能否构建一个公式，描述如何由对应的输入数据集中的数据转换为数据集的某个输出点吗？如果可以做到，那么问题转换到GPU空间就相对容易了。

这一方法的问题在于，为了取得最好的效益，你需要完全理解问题。你不能简单瞥一眼最高占用CPU的东西，然后把它们并行化。这一方法的真实益处在于把从输入数据点到输出数据点的链完全并行化。看到问题潜在的并行通常是第一个障碍。

优化通常用于执行数据的操作或函数。然而，随着计算能力相对于带宽的飞速增长，数据是首要考虑的因素。

如果你打算用多个GPU或多个GPU，这里最后一个考虑的如何在处理器元素上分解问题和数据集。以计算周期的角度来看，节点之间的通信将是非常昂贵，所以需要尽可能地将它最小化并跟计算重叠起来。

9.1.4 依赖性

依赖就是一些计算需要用到以前计算的结果，可以是针对问题域的计算也可能是数组下标的计算。

依赖关系主要由两种形式，要么一个元素是依赖于它附近若干元素，要么在多轮遍历数据集时下一轮依赖当前轮。

extern int a, c, d;
extern const int b, e;

void func_with_dependencies(void)
{
	a = 100 * b;
	c = 1000 * b;
	d = (a + c) * e;
}

依赖：鉴于一个指令的返回结果必须注入到下一个指令，这种类型的代码排放方式只允许较小的并行性并会导致流水线的失速。处于失速时，CPU与GPU将处于闲置状态。CPU与GPU都使用多线程来隐藏这个问题。

在CPU方面，来自其他虚拟CPU核的指令流填补指令流水线的空隙(如超线程技术)。然而，这要求CPU知道流水线中的指令属于那个线程，但这会使硬件变得复杂。在GPU上，也使用多线程，但采用时间切换方式，这样算术运算的延迟时间被以极小甚至可以没有代价的隐藏掉。 事实上，GPU上你需要20个时钟来隐藏这样的延迟。然而，这种延迟不一定来自另一个线程。

extern int a, c, d， f, g, h;
extern const int b, e;

void func_with_dependencies(void)
{
	a = 100 * b;
	c = 1000 * b;

    f = b * 101;
    g = b * 1001;

	d = (a + c) * e;
    h = (f + g) * e;

}

对上述代码，注意，如果你在计算变量a，c与使用它们计算d之间的位置插入一些独立的指令，将需要更久的时间才能获得d的计算结果。计算f，g和h的值与计算d是重叠的。实际上，你是通过重叠非依赖指令达到隐藏算术运算的延迟。

循环融合(loop fusion')是一种处理依赖关系并引入额外非依赖指令的技术。如下所示：

void func_with_dependencies(void)
{
	unsigned int i, j;
    a = 0;
    for ()
    { a = ???}

    for ()
    { d = ???}

    for ()
    { a =???}
}

两个计算的迭代空间是互相重叠的。因此可以将一个计算的一部分移动到另一个计算的循环体内部。这样就可以引入额外的，无依赖性的指令，另外能够降低总体的循环次数。循环迭代不是免费的，因为他们需要一个循环迭代值和一个分支。因此，降低的迭代次数会为我们在减少执行的指令数方面带来显著益处。

现在的GPU，很可能将这些循环展开，放到线程内，并由单个内核程序计算。然而，谨慎使用这种方法。通过执行这些操作，你同时减少了可用于线程/线程块调度的整体并行度。如果这个数目很小的话，会浪费执行时间。另外要注意，使用融合的内核时，通常会消耗更多的临时寄存器。由于寄存器的使用增加了，会限制一个SM上可调度的线程块数目，从而可能会限制实际可融合的数量。

最后，你应该好好考虑需要多轮遍历的算法。他们通常被实现为一些内核调用的序列，每一次调用在数据上循环一遍。由于每轮要读/写全局数据，效率通常较低下。许多这样的算法可以写成只涉及单个或少量目标数据点的内核程序。这为把数据放入共享内存或寄存器提供了可能，并且相较给定内核需要多次全局内存访问的方式，可以大大提高完成的工作量，这将明显改善多数内核的执行时间。

9.1.5 数据集大小

数据集的大小使选取问题的解决方案差别巨大：

一个典型的CPU实现可以分为以下几块：
数据集小于一级缓存	16KB~32KB
数据集小于二级缓存	256KB~1MB
数据集小于三级缓存	512KB~16MB
数据集小于单台主机内存大小	1GB~128GB
数据集小于主机端持久性存储大小	500GB~20TB
数据集分布在多台机器上	>20TB
一个典型的GPU实现可以分为以下几块：
数据集小于一级缓存	16KB~48KB
数据集小于二级缓存	512KB~1536MB
数据集小于GPU内存大小	1GB~128GB
数据集小于主机端持久性存储大小	500GB~20TB
数据集分布在多台机器上	>20TB

对于非常小的问题集，可以增加更多CPU核，可能会带来超线性加速比。如果将问题从内存移到三级缓存，或从三级缓存移到二级缓存。会看到一个很明显的加速。是因为使用的缓存有高得多的存储带宽。

GPU的主要问题不是缓存，而是你能在一张卡上保存多少数据。将数据从主机系统传入或传出会耗费大量计算时间。为了隐藏这个时间，你应该把计算与数据传输重叠起来执行。更好是利用主机的锁页内存同时做到传入与传出数据。由于锁页内存不会被虚拟内存管理系统换出，所以它必须是存放于主机的真正的DRAM内存。

在商业硬件上，你可用空间比总量要少一些。

在主机端，你的内存至少需要与输入和输出缓冲区分配的锁页内存等量。由于你通常使用最多2GB的锁页内存，因此剩余的内存量可用轻松地支持多个GPU，多数系统支持至少2个GPU卡。

当数据集可能因为计算、内存、存储或者能源方面的因素，无法置于单台机器时，你必须考虑使用多个节点。这就需要节点间的通信了。节点间的通信是非常耗时的，相比于任何内部的数据通信至少慢一个数量级。此外，你还必须掌握一套API。尽可能避免节点间通信这一步骤。

9.1.6 分辨率

提高问题的分辨率是否比提高速度更有吸引力？一个更精准的结果在你的问题里能得到什么？

9.1.7 识别瓶颈

1. Amdahl定律

Amdahl定律告诉我们，当数据流中任然存在串行执行元素时，将限制速度的提升。

无限扩展程序的唯一办法是消除程序执行中的所有串行瓶颈。

某程序突然遇到一个串行点或同步点，一切都堵塞了。针对这类问题，把瓶颈部分并行化就能解决。

在考虑直方图计算，你会看到如果把所有线程都加入同一个桶，就形成了同样的瓶颈。通常会采用原子操作，这样一组并行线程就要串行执行。相反，如果分配给每个线程属于它自己的一组桶，然后再将这些桶合并起来，就能消除串行瓶颈问题。

2. 分析

分析是确定你当前在哪儿以及应该在什么地方多花点时间的最有用的任务之一。

优化应该根据确切的数字和事实，而不是猜测那可能是最应该优化的地方。Nvidia提供了CUDA Profiler和Parallel Nsight，以提供分析信息。

分析器用过读取硬件计算器，来发现代码花费的时间和在GPU上的占有率。它会提供非常有用的数据，如总共合并读和写次数、缓存命中/失败率、分支频率、线程束串行化程度。

使用分析器做完一个初步的检查之后，你应该先查看花费总时间最多的代码段。典型的未优化程序，80%的时间花费在20%的代码上的。优化20%的代码是有效减少使用时间的关键，分析器是确定这20%代码所在的一把钥匙。

当然，一旦上述问题已被优化为最佳，如果不进行重新设计，后面为提供加速化的进一步优化将会变得越来越耗时。

9.1.8 CPU与GPU的任务分组

事实上，最好的应用程序往往可以充分利用CPU与GPU两者的优势，并相应地划分数据。任何基于GPU的优化也应该考虑CPU，因为这对于总的应用程序时间很重要。可以使用的CPU核数越多，通过分流一些工作给CPU的潜在收益越大。

如果说CPU可以处理GPU的工作1/10，那么当仅需要3个CPU核时，你的GPU就获得额外30%的吞吐量。

对于IO限制而言，因为引入更多的CPU的线程或进程，经常可以显著提高整体的IO吞吐量。这似乎很奇怪。因为IO设备的输入输出上限决定了它的吞吐量。在现代大内存的计算机中，大多数的IO操作都是进行缓存的。因此，IO操作大都在内存上移动而不是设备上移动。

独立的CPU进程或线程可以创建一个独立的GPU上下文，且启动它们自己的内核到该GPU中。这些额外的核经常以队列的方式在GPU中去执行。当所需的资源变为可用时，内核开始执行。

GPU的空闲时间比CPU空闲时间更昂贵，因为它的吞吐量通常在CPU时间的10倍以上。

通过在一个GPU中放入多个内核，这些内核就可以伺机占用空闲硬件槽位。这将在一定程度上增加第一组内核的延迟，但会大大提高应用程序的整体吞吐量。

通过引进一对进程，巧妙的重叠了IO，CPU，GPU和传输时间、整体吞吐量获得了显著的改善。

进程允许设置为处理器关联(processor affinity)，可以把进程绑定到一个给定的CPU核。这样做往往会提供性能，因为它可以更好地重用该核的缓存。

选择线程还是进程，在很大程度上取决于CPU之间需要同步任务的个数。

在权衡CPU/GPU的使用过程中，也需要知道如何最优地划分任务。当数据是稀疏分布的或者是小数据集的时候，CPU很擅长处理这类串行任务。

有时会看到CPU用在规约操作的最后阶段。通常几轮迭代之后，归约操作涉及的元素数会下降为原来的一半。可供调度的线程数量就小于一个GPU可供调度的最大线程数了。如果再继续迭代几轮，一些SM就开始闲置。

因此，一种优化策略是，当迭代到一定的阈值，剩余部分的计算就转交给CPU来完成。不过，从费米架构之后，英伟达解决了上述问题。能够让那些空闲的SM在一个排队的内核中使用。但是，要让SM变得空闲，必须要保证其上的所有线程块已经完成它们的任务。一些非最优的内核可能残留一个或者数个活跃线程(即使在规约操作的最后一层)，从而导致内核牵制住该SM，直至整个规约操作完成。

对于类似归约操作的一些算法，请确保每次迭代都在减少活跃线程束的数量，而不单单是活跃线程的数目。

9.1.9 本节小结

        1.理解问题并基于你的编码时间和熟练程度定义你的加速目标。
        2.识别问题中的并行性，并思考如何以最佳方式在CPU和一个或多个GPU之间分配。
        3.考虑一下，是较少的执行时间还是处理数据以获得更高分辨率更重要。
        4.理解任何串行代码的实现，并思考如何处理它们最合适。
        5.分析你的应用程序，以确保你的理解确实反映了实际情况。如果可以帮助你加强理解，请重复你之前的分析。

9.2 策略2：内存因素

9.2.1 内存带宽

内存带宽和延迟是所有应用程序都要考虑的关键因素，尤其是GPU应用程序。带宽是指与某个给定目标之间传输的数据量。在GPU的情况下，我们主要关心的是全局内存带宽。延迟则是指操作完成所用的时间。

GPU上的内存延迟设计为由运行在其他线程束中的线程所隐藏。当线程束访问的内存位置不可用时，硬件向内存提交一次读或写的请求。如果同一个线程束上其他线程访问的是相邻内存位置并且内存区域的开始位置是对齐的，那么该请求会自动与这些线程的请求组合或者合并。

我们需要考虑的一个关键领域是运行过程中的内存事务的数量。每一个内存事务被送入一个队列中然后由内存子系统单独执行。这当然会有一些开销。一个线程一次提交对4个浮点数或整型数的一个读操作比提交4个单独的读操作花费的代价更小。

为接近峰值带宽：
        1.使用线程束完全加载SM，实现接近100%的占用率。
        2.通过float4/int4向量类型使用128位读操作，此时占用率小了很多，但仍然能达到100%的峰值内存带宽。

9.1.2 限制的来源

内核通常被两个关键因素限制：内存延迟/带宽和指令延迟/带宽。正确理解这两类关键限制因素中哪一种正在限制系统的性能，对于指导你合理的分配精力是很关键的。

最简单的能够看到代码平衡位置的方法，是简单地注释掉所有算术运算，然后直接赋值成结果代替。算术指令包括所有的计算，分支，循环等操作。如果存在某种归约操作，你只需要将它替换成普通的求和操作。一定要确保包括所有从内存中读取到最终输出的参数，否则编译器将删除明显冗长的内存的读写操作。对内核重新定时，你会看到花在算术和算法部分的近似的百分比。如果这个百分比很高，那么你就受到了算术限制。反之，你受到内存限制。

此外，如果内存模式没能很好地合并，GPU将不得不串行执行指令流以支持分散的内存读写。如果是这种情况，那么有可能需要重新安排内存模式，以允许GPU将线程的内存访问模式合并。

是否可以扩大一个单一线程处理的输出数据集的元素数目呢？这通常同时有助于内存受限型和算术受限型的内核。如果你这样做，请不要在线程中引入循环，而是要通过复制代码实现。如果代码是很重要，这也可以作为设备函数或宏来实现。确保将读取操作提前到内核开始处，这样在需要数据时就已经完成了对它们的读取。这将增加寄存器的使用，所以一定要监控正在被调度的线程束个数以确保它们不会突然地退出。

至于算术限制的内核，查看源代码并思考如何将其翻译成PTX汇编代码。不要害怕实际产生的PTX代码。数组索引通常被替换为基于指针的代码，将速度较慢的乘法替换成更快的加法。使用2的幂次的除法和乘法指令分别可以被替换成速度更快的右移或左移位的运算。循环体中的所有常量(不变量)应该被移到循环体外。如果线程包含一个循环，那么展开循环通常会实现加速。

9.2.3 内存组织

在许多GPU应用程序中，使用正确的内存模式往往是关键的考虑因素。CPU程序通常在内存中以行的方式安排数据。我们必须尝试安排内存模式以使连续线程对内存的访问以列的方式进行。此原则同时适用于全局内存和共享内存。

cudaMalloc函数以128字节对齐的块为单位分配内存。如果使用结构会越过这个边界，有两个办法。首先，你可以在结构中添加填充的字节。或者，你可以使用cudaMallocPitch函数。

对齐是一个很重要的的标准，它将决定内存事务或缓存行需要获取一次还是两次。

通常，使用共享内存作为临时缓冲是明智的。然后，可以将其用于对全局内存进行合并的读写操作。

9.2.4 内存访问以计算比率

内存操作与计算操作的比率是值得思考的问题。你所期望的理想比例至少是10:1。也就是说，对于每一个内核，从全局内存执行的读取操作需要执行10条或更多的指令。这些指令可能是数组索引计算、循环计算、分支或条件判断。每个指令都应该对有效地输出起到一定的贡献。特别是循环没展开时，它经常会增加指令开销但并不会助于任何有用的工作。

因此，在每个周期内线程束调度器提交2条指令或4条指令。由于这些指令来自于不同的线程束，它们之间来自于不同的线程束，它们之间是相互独立的，因此将它们放入执行单元(CUDA核，SFU和LSU)流水线中。

基于切换其他线程束的能力，使用最少个数的常驻线程束无法隐藏内存指令或指令延迟。指令流的失速实际将会使CUDA核失速，这是我们非常不愿意看见的。实际上，多个线程块会被分配到一个SM上，以试图确保这个问题永远不会发生并且更重要的是生成各种形式的混合指令。

第二个要点是共享的资源(SPU，LSU)限制了持续执行相同操作的能力。由于CUDA核和LSU都被纳入了流水线中，但是它们只有16个单元宽度。因此，将线程束调度到这两个单元之一会花费两个周期。

当数据流中有全局内存写操作时，你需要将读操作提前至内核开始处。试用一下代码：

int tid = blockIdx.x * blockDim.x + threadIdx.x;
int a0 = a[tid];
int b0 = b[tid];

data[tid] = a0 * b0;

我们有两个办法，标量方法或向量方法。GPU只在硬件上支持向量的加载和保存而不支持向量操作。因此，乘法操作实际上要像C++中的重载操作符一样完成并且只是将两个相互独立的整数相乘。然而，向量方法分别执行两个64位加载和一个单独的64位存储而不是非向量版本的4个独立的32位加载和一个32位存储。因此，40%的内存事务内节省了。内存带宽使用是相同的，但是更少的内存事务意味着更小的内存延迟。因此，等待内存的总体失速时间减少了。

为了使用向量类型，只需声明数组为向量类型int32。被支持的类型为int2、int3、float2...。当然可以创建自己的类型，并且定义自己的操作符。每一个向量类型实际上都是一个对齐的结构体，包含N个声明为基类型的成员元素。

因此，希望你能真正看到不同类型的指令之间是需要平衡的。

重要的是要认识到，指令流需要足够的计算密度以充分利用SM上的CUDA核。内核只是简单地执行加载/存储操作和少量其他工作，因此无法达到设备的峰值性能。通过每个元线程处理2个、4个或8个元素，从而扩展这些内核使其包括独立的指令流。因此尽可能使用向量操作。

9.2.5 循环融合和内核融合

另一个可以显著地节省内存带宽的技术是基于9.2.4节提到的循环融合。循环融合是只两个明显独立的循环在一段范围内交错地执行。

内核融合是循环融合的演变。如果你有一系列按顺序执行的内核(一个接着一个执行)，这些内核的元素能否融合？ 对于那些还未完全理解的内核，这样做的千万要小心。调用两个连续的内核会在它们之间生成隐式地同步。

开发内核时，将操作分解成几个阶段或几轮是常见的。例如，第一轮你看你针对整个数据集计算结果。第二轮，你可以使用特定的标准对数据进行过滤，然后在特定的点进行深入的处理。如果第二轮能够本地化一个线程块，那么第一轮和第二轮能够组合成一个单独的内核。这就消除了将第一个内核写入主存，随后读取第二个内核的操作及调用内核的额外开销。如果第一轮能够将结果写入共享内存，那么只在第二轮需要这些结构，这样就完全消除了对全局内存的读取、写入。归约操作经常被划分到这一类并且能从这样的优化中显著的受益，因为第二阶段的输出通常比第一阶段的输出小很多，因此它显著节约了内存带宽。

内核融合技术如此有效的原因是它所带来的的数据重用。一旦数据存储到共享内存或寄存器集中，那么尽可能重用它。

9.2.6 共享内存和告诉缓存的使用

相比于全局内存，使用共享内存可由提供10:1的速度提升。但是共享内存的大小是受限的。

在数据集上迭代的内核如果没有重用数据，那么需要意识到它们可能正在以低效地方式使用缓存或者共享内存。

与在一个大型数据集执行多轮不同，内核融合这样的技术可用于在数据间移动而非多次传入它。思考一下输出数据的问题而不是输入数据。构建该问题是将线程分配给输出数据项而不是输入数据项。在数据流方面，建立流入而非流出。优先选择聚集(收集数据)(gather primitive)原语而不是分散原语(scatter primitive)。GPU会同时从全局内存和二级缓存直接将数据广播到每个SM，这一点支持高速度聚集型的操作。

如果数据项很小，则内存事务可以逐步将规模减少读取，直至每次访问32字节。因此，从十分分散的内存区域访问一个数据元素的内核，在任何基于缓存的架构，包括CPU和GPU，表现会十分糟糕。原因在于单个元素的读取会载入128字节的数据。对于大多数程序而言，存入缓存的数据会在下一次循环迭代中命中，这是由于程序常常访问与之前访问的数据临近的数据。因此对于大多数程序，这是一个显著的优点。但是，对于那些只需要单个数据元素的程序来说，剩余的124字节是多余的。对于这种内核，你需要为内存子系统去取所需的内存事务而不是缓存行的大小，只能在变异的时候通过-Xptxas-dlcm = cg标志来完成此工作。这将所有的访问减少到每次事务32字节并且令一级缓存失效。对于只读数据，考虑使用纹理内存或者常量内存。

9.2.7 本节小结

        1. 仔细考虑你的内核处理的数据并且如何将其以最佳的方式安排在内存中。
        2.针对128字节的合并访问，优化访存模式，对齐到128字节的内存读取大小和一级缓存行大小。
        3.注意权衡单精度和双精度对其内存使用的影响。
        4.在适当的时候将多个内核合并成单内核。
        5.以最适当的方式使用共享内存和缓存，以确保你能充分利用更高计算能力设备上扩展容量。