充分利用 MultiGPU架构

越来越多的实例证明，借助单颗 GPU技术实现的速度提升可以比传统CPU高出几个数量级。如果在工作站或计算节点上插入两颗或四颗 GPU，那么其运算应用或游戏的性能可分别翻一番和翻两番。因此今天的人们越来越关注，在一个系统中，主处理器的多核功能与GPU的协同作用使整个系统性能的提升潜能。

虽然这一点非常诱人，但多 GPU 和混合型 CPU+GPU 的性能在很大程度上依赖于计算机主板上的 PCIe 总线装置。留心供应商会在 PCIe 总线上偷工减料! 配备了合适的 PCIe 芯片组，multiGPU 应用可以根据系统中 GPU 的数量相应地提高性能。如果系统内的芯片组不恰当，在 multiGPU 上的投资就不会有成效。为什么要浪费资金？ 确保 PCIe 芯片组能令你的 GPU 实现所有的性能！

在英伟达CUDA 4.0中就包含了一系列特性，能简化工作站或计算节点内多GPU的技术。爱尔兰高端计算中心 (ICHEC) phiGEMM库利用CUDA 4.0 特性实现了同时使用多 GPU 和主处理器的矩阵乘法计算。phiGEMM 的性能提升非常显著，单个GPU+CPU的性能与Linpack HPL矩阵乘法相等，后者用于评估世界前500强超级计算机。在四个GPU和一个主处理器之间运行单个矩阵乘法时，phiGEMM 可实现超过1万亿次浮点/秒 (10710 亿次浮点/秒) 的双精度矩阵乘法计算，其所在矩阵比任何一个GPU的内存都要大！

如图1展示了对于一个25000 *25000矩阵所实现的性能，两个2.93 GHz Intel X5670 处理器的12个核心的运算速度可稳定在1300亿次浮点/秒。在没有Intel核心带来性能提升的情况下，4颗英伟达GPU性能可达 9420 亿次64位浮点运算/秒。换言之，相比单个CPU，4核GPU的性能提高了3.4倍。phiGEMM库支持免费下载。

性能不良的PCIe总线芯片组造成的影响非常显著。例如，同样的phiGEMM矩阵乘法在共用PCIe总线上会慢17%，因为传送数据用的时间会更长。也就是说，每个GPU由于共用总线而只得到一半数据带宽。矩阵乘法是展现多GPU性能的良好平台，因为随着矩阵规模的增加，运行时间会受限于浮点性能而非数据传送。

很多供应商宣称他们的系统支持多高速 x16 PCIe插槽，这在某些性能假设下是符合事实的。为节省资金，某些供应商使用只在一个设备激活时才能实现全部性能的PCIe芯片组。在《CUDA Application Design and Development》中也曾提及，某些PCIe装置会区分对待多个GPU中的某一个。从英伟达Visual Profiler的图形输出上可以看到，数据传输在第二个设备 Device_1:Context_1 上用的时间显著增长。

英伟达 Visual Profiler输出显示，在用一条劣质PCIe总线连接的多GPU之间运行3D快速傅里叶变换 (FFT) 时，性能降低接近 60%。 性能之所以有如此显著的下降，是因为相比矩阵乘法，快速傅里叶变换每传送一个数据所进行的计算量要少。更少的计算量意味着相比计算吞吐量数据传输速度对应用的制约更明显。

性能不良的 PCIe 总线对这类应用运行时间的影响更加大。性能下降60%意义重大，因为这能抵消多GPU 带来的优势。Thrust C++ 数据并行应用程序接口在CUDA 4.0版中为标准配置。借助类属编程和仿函数 (类似于函数的C++对象)，可编写C++应用，在CUDA矢量和阵列中以高性能大规模并行方式运行。Thrust令CUDA编程变得简单，因为任何懂得C++的人已经知道如何为GPU编写程序。

从multiGPU和混合型CPU+GPU编程的角度看比较值得关注的是，Thrust可产生既能在多核处理器上又能在GPU上运行的代码！它供了两个分类符“__device__ __host__”。它会告诉编译器为仿函数产生既能在主机上又能在GPU设备上运行的代码。我使用Thrust功能编写的应用可利用一个工作站内所有可用的计算资源。

对于这些应用，我喜欢使用基于主机的仿函数和OpenMP (Open MultiProcessing) 指令明确指定主处理器的并行性。 作为一种应用程序接口，OpenMP可被大多数编译器支持，可用来创建支持多核处理器的并行应用。注意基于 Thrust 的应用可被透明编译，因而可全部在主多核处理器上运行。无需变更代码！ 相反，恰当地定义了一个特殊变量 “THRUST_DEVICE_BACKEND”以指明一个 OpenMP 后端。 Thrust 的编写方式是，编译器产生的代码能在主多核处理器上并行运行而无需 GPU。

使用MPI接口的程序员可使用该基于 Thrust 的功能创建分布式应用，在系统每个 GPU 和主处理器上运行一个单独的 MPI 过程。注意，分布式应用也可在很大程度上使用 PCIe 总线，也可像前面讨论过的 multiGPU 那样暴露供应商 PCIe 装置的缺点。

所以，最终问题的关键，还是在PCIe总线。寻找完全支持向多GPU传送数据的高性能 PCIe 总线，至关重要。即使你的应用暂时用不到多GPU，也要为潜在的应用做好准备。毕竟你可以通过添加额外的 GPU 把性能提升二到四倍。

而对于那些运行计算集群的人，性能欠佳的PCIe总线更会对性能产生负面影响。使用phiGEMM库、快速傅里叶变换或其它多GPU示例进行的基准测试都能立即暴露出问题。不要再迷信衡量PCIe到单个设备的带宽结果了，我们已经步入了一个multiGPU计算的高速并行时代。