第7章 CUDA实践之道 摘录

7.2 串行编码与并行编码

7.2.1 CPU与GPU的设计目标

CPU使用的MIMD模型，而GPU采用SIMT指令模型。

CPU并行的方法是执行多个独立的指令流，并在这些指令流中尝试提取出指令级并行。因此，CPU的指令流水线非常长，CPU从指令流水线中找出可以送到独立执行单元的指令。这些执行单元通常包括至少一个浮点单元、至少一个整数单元、一个分支预测单元和至少一个加载/存储单元。

分支(if, branch prediction)是可以静态预测的。一些编译器通过对分支指令设置一个比特位来判断分支条件是否满足。一次，循环语句的跳转通常预测为采取“taken”(即跳回起始点)，而其中的条件判断通常预测为不采取"not taken"。这样做完全避免分化，同时也增加了优势，即后续指令已预提取到了缓存。

分支预测从简单但高效的静态模型演化成为记录当前分支历史的动态模型。现代处理器都采用了多层复杂分支的预测机制。

同分支预测技术一样，一种叫做预测执行的技术也被采纳。考虑到CPU可能正确地预测分支，那么此分支地址开始执行指令流就是正确的。

直到目前为止，CPU与GPU的一个主要区别就是缓存等级、数量的不同。随着芯片密度不断增长，指令周期缩短了。现在，内存访问成了现代处理器设计的主要瓶颈，目前一般通过多级缓存应对解决。

GPU设计者相信程序员会充分利用高速内存，将数据放在距离处理器近的地方，因此，他们在每个SM上都设计了一个共享内存。共享内存类似于传统处理器上的一级缓存，是一小块低延时、高带宽的内存。

费米架构的设计者将其片上内存扩展为64K，其中16K必须分配给一级缓存，还有一部分总是作为共享内存的形式存在。cudaFuncSetCachefig()函数设置一级与共享分配。

同CPU一样，费米架构的GPU将每次内存访问获取的128字节的数据放入一个缓存行(cache line)中，随后相邻线程的访问通常会命中缓存，即在缓存中找到需要的数据，避免了再次访问全局内存。

GPU设计与CPU设计的一个显著不同就是SIMT执行模型。在MIMD模型中，每个线程都有独立的硬件来操作整个独立的指令流。如果线程执行相同的指令流但不同的数据，那么这种方法非常浪费硬件资源。而GPU提高了一组硬件来运行这个N个线程，N个正好为当前线程束的大小32.

GPU中的SIMT模型实现类似于原先的矢量架构SIMD模型。SIMT模型解决了一个关键的问题，就是程序员不必再为每个执行相同路径的线程写代码。线程可以产生分支然后在之后的某一点汇集。但由于只有一组硬件设备执行不同的程序路径，因此程序的性能有所下降，不同程序路径在控制流汇集之前顺序或轮流执行。

CPU采用的顺序控制流执行方式。如果执行一个需要很多周期的指令，将会阻塞当前的线程。因此，Intel采用了超线程技术。当正在运行的线程受阻时，硬件会内部切换到另一个线程。单线程情形，当遇到这种由指令延迟与存储延迟引起的阻塞时，CPU处理器会闲置等待操作完成。而线程模型的设计就是用来隐藏这种延迟的。

GPU使用了惰性计算模式(lazy evaluation model)，这使它拥有了另一个优势，即它不用阻塞当前线程，除非需要访问一个有依赖性的寄存器。

操作的并行度取决于处理器上的可执行单元的数量。这些处理单元可以是多核处理器执行单元(以使用线程的处理器核的形式)，或者是大规模处理器设计中的执行单元。

7.2.2 CPU与GPU上最佳算法对比

虽然不是所有的算法都适合用来实施并解决并行问题，但是许多问题可以分解成在数据集上的操作。如果从数据或任务并行的角度看，多数情况下，这些操作都是内在可并行的。

目前，对具体架构的优化通常需要将应用程序与特定的硬件相绑定。

通常，MPI实现成CPU与GPU混合的模式，这里使用CPU处理网络通信以及硬盘的输入输出。通过一块共享内存，可以使用GPU Direct实现页面直接传输到InfiniBand网卡上。应该采用直接通过PCI-E总线点对点(P2P)地进行数据传输，无需使用CPU的主存。RDMA是开普勒架构中的一个特性，它使得GPU在网络中成了一个更加独立的节点。

GPU解决问题时需要考虑到使用的线程的数量是有限的。事实上，由于受寄存器使用的限制，一般合理且复杂的内核函数会将线程数限制为256或512。线程间的通信问题是分解问题的关键。线程间的通信通常通过使用高速共享内存实现，同一块内的线程可以快速通信，并且延迟很小。相比之下，线程块之间的通信只能通过重新调度内核和使用全局内存来实现。

GPU解决问题时需要考虑的另一个因素是设备可提供的内存。单个GPU能提供的内存不是很大，但可以通过使用多个GPU来解决，利用容纳多个PCI-E显卡的高端主板。

另外，GPU上递归也有其盲点。只有计算能力大于2.0的 GPU才支持递归运算，并且只有带__device__前缀的函数可以使用，带__global__前缀的函数无法使用。

然而，任何递归算法都可以描述成迭代算法。

GPU内核函数调用时需要的线程数量是固定的，尽管kepler架构之后有所改变。但之前仍然不能很好实现动态并行。动态并行的并行度随时间而改变。在快速排序问题中，它以2的指数级增长。而在寻找路径的问题中，发现一个新节点可以会引入30000多个其他路径。

然而，GPU如何实现这些算法呢？有许多方式，最简单的就是当知道并行度扩展的方式后，例如快速排序，可以简单地调用一个内核，这个内核执行算法迭代中的第一个层级或第N个层级，将这些内核一个接着送入一个执行流。当某一个层级结束时，将其状态信息写回全局内存，然后选择下一层级的内核执行。

若每次迭代并行度增加的数量不确定时，依然可以将状态存储在全局内存。但此时需要与主机端进行通信以确定下次迭代需要开启的线程数目。我们可以使用原子写操作将信息写到块内的共享内存中，然后在线程块执行完毕之前做原子加法操作将信息写回全局内存。接着使用memcpy将数据复制回主机端以调整启动下个内核需要的参数。

当使用支持动态并行的GPU卡，可以最开始几层用递归调用来计算，直到到达一定深度的层数，以保证有足够的内核函数可供调度。另外一种办法就是先在CPU上进行一些初始计算，当需要大规模并行时再将数据复制到GPU进行计算。不要以为所有的计算都必须在GPU上执行。尤其是在并行度不高的任务时，CPU是一个很好的伙伴。

这类问题的另一种解决方法是使用一种叫做分段扫描的特殊扫描操作，利用分段扫描，可以对数据集进行有规律的扫描操作(最大值，最小值，总和等等)，并附上一个额外的数组，将原来的数组分成不同大小的块。每块分配一个多个线程进行计算。由于附加的数组可以在运行时更新，因此如果分段扫描能保持在一个单独线程块内执行，就可以减少调用多内核的需要。否则，则需要采取一种更简单的解决办法。

如今GPU的递归问题通常都使用迭代问题的框架。

7.3 数据集处理

一种更好的方法是将其分成若干个表，然后在完成相应操作之后将它们合并。大多数并行处理算法都以不同的方式采用这个方法以避免更新共享数据结构而带来的序列化瓶颈。

过滤操作实际上是一种常见的并行模式，一种分割操作。分割操作基于一些关键因素将已知的数据集划分成N个部分。分割操作会产生两个表，一个满足标准，另一个不满足。

当并行执行这种标准时候，有许多问题需要我们考虑。第一个问题就是，满足标准的数据的数量以及不满足标准的数量都是未知的。第二个问题就是需要处理的元素的数量是未知的。第二个问题就是需要处理的元素的数量是未知的。而它的决定着输出集的建立。最后，原始数据集的排序必须保留。

__ballot函数主要判断由给定线程传入的一个判定值。此处，该判定值可以简单地认为是一个真、假布尔值。如果该值为零，则返回一个只有第N个比特位为1的值，其中N表示线程索引(threadIdx.x)。该函数除非与atomicOr联合使用，否则ballot的用处不会立刻显现。

有许多方法可以用于求和。如果线程块数目很少，可以将每块的结构传回CPU，用CPU来计算总和。

另一种方法就是把所有线程块计算得到的部分和写回GPU的全局内存中。然而，要完成对所有单独线程块的求和，需要等待所有SM中所有的线程块都完成了判定。唯一的解决办法就是结束当前的内核函数然后调用另一个。然后之前写回全局内存的值通过诸如并行规约的方式读入，并计算得到最终的总和。

在Fermi架构上，由于每个线程块对总和只更新一次，所以可以使用atomicAdd函数对于全局内存的总和值进行简单的累加操作。从统计学的角度来说，多个线程块同时到达执行原子加法指令的可能性非常小。考虑到读取源数据的内存操作会按照顺序到达，实际上SM中的执行流就实现了很好的序列化，并且能够保证进行原子加法操作时不会与其他SM产生竞争。

7.4 性能分析

7.5 一个使用AES的示例

7.5.10 AES总结

总结如下：
        1. 针对CPU与GPU的理想内存模式是不同的。优化CPU的访存模式受益甚大。对于越早开发的GPU，内存模式的优化更为关键。
        2. 对于最早的GPU设备，申请只读内存需显式声明为常量内存，而不是由编译器自动指定。
        3. 为了使编译器更容易自动优化性能，内核程序的重新排序或者变化是必不可少的。
        4. 有效的寄存器使用方式以及合适的使用数量，对高性能CUDA程序设计至关重要。
        5. 可以通过二级缓存在线程块间共享只读数据，也可以使用共享内存保存这些数据。但后者需要保存N份数据，这里N是可用的线程块数。
        6. 复杂的、含线程分支的算法，如gmul函数的解码，使用没有线程分支的内存查找代替，查找表驻留于缓存或共享内存。缓存就是专门为此类数据驱动方式的分散内存访问模式而增添的。
        7. 确认变量的分配没有溢出寄存器，尽早消除栈或本地内存的使用。
        8. 务必尽早在解决方案之处验证其正确性。最好相互独立地开发代码的不同版本。
        9. 务必查看程序的实际执行时间。你的心智模型(mental model)关于事情运转方式的认识未必正确，常常会忽略一些东西。经常查看数据，观测每次变化后的效果。

7.6 本章小结

应该通过具体操作过程，学会面向多种计算能力的硬件时的权衡之术和设计要点。你也可以看到，项目开发过程中早期的设计决策是如何影响到结果的。在你开始写代码之处，你应该全盘考虑和理解寄存器、共享内存、缓存的用法以及全局内存的访问模式等设计的主要方面。

要实现优异的性能表现，而不只是平均表现，必须要理解且彻底理解，你正在的开发环境。诸如不同内存层次结构的概念。线程和本地线程存储的概念。