5 流水线模式

与指令装配线类似，装配线模式通过并行使用不同的硬件资源来获得高性能。如果多个指令既没有数据和控制依赖，也没有结构化依赖，则可以同时在处理器的多个装配线单元上执行，称为平行指令装配线。超标量和VLIW处理器通常有多个执行单元，可能是相同的或不同的。这些执行单元可以并行执行指令或装配线。通常，处理器的指令级并行技术要求可以动态地依赖于分析、顺序或随机发射、顺序或随机执行；编译器或开发人员也可以承担一些相关的工作，以提高程序的执行性能。

说明流水线优化粒度过细，需要深入了解处理器的流水线延迟、吞吐量和编译能力，不易使用。由于编译器在生成代码时可能会改动代码的执行顺序，即编译器生成的代码顺序和源代码的顺序并不完全一致，因此C语言并不适合用来进行指令流水线的优化，只有汇编语言才能满足这一要求。指令流水线优化通常用于常用函数，如矩阵乘法，对性能要求非常严格。

更好的指令流水线优化实践模式是修改源代码，使编译器生成所需的指令系列，然后根据编译器生成的汇编代码进行修改。

优化指令流水线的重点是消除指令之间的依赖关系。例如，代码列表1-1 所示的循环扩展代码依赖于变量sum。它可以优化为代码清单1-2 的版本。

代码清单1-1 循环展开代码

sum=0; for(int i = 0; i < len; i = 4){ sum = a[i]; sum = a[i 1]; sum = a[i 2]; sum = a[i 3]; }

代码清单1-2 优化后的代码

sum=sum1=sum2=sum3=0; for(int i = 0; i < len; i = 4){ sum = a[i]; sum1 = a[i 1]; sum2 = a[i 2]; sum3 = a[i 3]; } sum = sum1; sum2 = sum3; sum = sum2;

代码清单1-1 就循环中的四个句子而言，它们都依赖于sum变量更新，因此循环中的四个加法操作需要串行执行（读a操作可以平行执行）。优化后，四个句子不再依赖，因此加法操作和访问存储都可以平行操作。

当需要使用汇编语言并行优化指令级时，作者建议读者尝试使用内置函数生成所需的代码，然后修改内置函数版本生成的汇编代码，以生成所需的代码。

装配线模式是指令装配线平行的抽象和改进。流水线模式借用了指令流水线平行中多个不同单元同时执行的优点。同时，通过增加粒度去除指令流水线平行粒度过细的限制，使得流水线模式可以应用于日常编程任务，并广泛应用于许多服务器。

目前主流的异构系统CPU GPU为例，假设CPU计算能力为1Tflops，GPU计算能力为5Tflops；如果只使用CPU，所以较多可以获得1个计算性能Tflops；如果只使用GPU，较多可以获得5个计算性能Tflops，通过流水线模式同时使用，较大可获得6Tflops的性能。

流水线模式可以从多层次映射到计算系统，充分利用装配线模式可以提高多层次应用的性能。

在集群层面，许多应用程序需要在集群节点之间交换数据，即通信；与串行计算相比，通信是集群计算引入的额外消耗。如果不能消除通信消耗，应用程序的性能将随着处理器的增加而逐渐下降。计算时可采用流水线模式进行通信，然后通过计算隐藏，获得稳定的应用性能。MPI支持的异步通信功能可以在节点间计算数据传输，可以有效缓解这个问题。

在单个节点上，访问硬盘的延迟远远大于访问内存和计算。如果可以在访问硬盘时计算，则可以掩盖硬盘读写延迟。Linux 和 gcc的异步IO库都提供了相应的支持，但在大多数情况下，库都通过线程支持异步IO的概念。

在单个核心上，核心将通过多内存通道访问内存中的数据，合理安排访问存储顺序和依赖关系，更好地发挥多通道内存的带宽。如果内存读写和计算相互依赖，内存通道将处于等待状态，难以获得较佳性能。

以常见的隐藏硬盘读写延迟为例，其主要包含两个部分：从硬盘中加载数据，在数据上面进行计算。本章以此为例，展示在各种不同环境下如何实现流水线模式。

代码清单1-3 所示，总共迭代 numIter 每次迭代都会加载数据，然后计算。

代码清单1-3 串行实现流水线模式

float sum = 0.0f; for(int iter = 0; iter < numIter; iter ){ loadDataFromFile(file, iter, len, data); sum = computeSqureSum(len, data); }

不失一般性，我们只假设loadDataFromFile函数从文件中加载数据，computeSqureSum计算函数。使用loadDataFromFile存储器系统主要用于函数加载数据computeSqureSum计算主要使用计算组件，在计算系统中，两者都使用不同的计算机组件，因此潜在的是并行操作。

很明显的是loadDataFromFile和computeSqureSum不同的单元使用计算系统，但它们依赖于共同的变量data因此，在代码清单上 这两个函数只能串行执行。

从代码清单1-3 串行版可见：

1. 第n次迭代，loadDataFromFile和computeSqureSum流水线模式不能用于依赖两个函数。

2. n-1次迭代时computeSqureSum和n次迭代loadDataFromFile依赖关系(依赖关系)data变量)，但两者操作的数据不同，是伪依赖，可以通过变量重命名去除。重命名后没有相关性，可以并行做。

具体实现多线程、事件机制等。本节采用较直接、较容易编码的异步IO机制，如代码清单1-4 所示，关于异步IO机制的细节可参考glibc文档。确保循环体内的文档。computeSqureSum和loadDataFromFile能执行流水线，把loadDataFromFile前提一个循环。

代码清单1-4 异步IO实现流水线模式

S0 loadDataFromFileAsync(file, 0, len, data0); for(int iter = 0; iter < numIter-1; iter ){ S1 dataBuff = iter % 2 ? data0 : data1; S2 loadDataFromFileAsync(file, iter 1, len, dataBuff); S3 syncPreviousLoad(); S4 data = iter % 2 ? data1 : data0; sum = computeSqureSum(len, data); } S5 sum = computeSqureSum(len, (numIter-1)%2 ? data1 : data0);

S2处的loadDataFromFileAsync使用异步IO函数从文件中加载数据，不会阻塞，发出异步IO操作后，控制器将立即返回，CPU会接着执行S3处的函数。S3处的函数syncPreviousLoad()目的是保证前一个循环的异步IO操作（loadDataFromFileAsync）完成。使用异步IO，算法采用两个缓冲区(有人称这种方法为双缓冲区)，S1和S4是为当前迭代选择缓冲区的逻辑。与指令流水线并行一样，流水线模式也存在于流水线的建立和完成阶段，对应于代码S0和S5。

对于C 对于11线程，从线程本身的操作是异步的，即主线程不需要等待从线程完成。代码列表1-3 我们可以使用单独的线程来执行computeSqureSum此时函数需要单独的函数data保存计算结果，然后在计算完成后叠加每个线程的结果，如代码清单1-5 所示。

代码清单1-5 C 实现11线程流水线模式

int NT = 8;//we use 8 threads float sum = 0.0f; std::future<float> f[NT]; for(int iter = 0; iter < numIter; iter = NT){ for(int i = 0; i < NT; i ) { loadDataFromFile(file, iter, len, data i); f[i] = std::async(computeSqureSum, len, data i); } for(int i = 0; i < NT; i ) { sum = f[i].get(); } } sum = computeSqureSum(len, (numIter-1)%2 ? data1 : data0);

为了方便和不堵塞后面的调用，我们使用了它std::async函数，std::async默认函数会自动创建线程，并返回一个std::future类对象，std::future类的get线程执行函数的返回结果可以通过函数获得。

该代码的具体延迟隐藏行为和函数loadDataFromFile和函数compueSqureSum执行时间密切相关，感兴趣的读者可以自己分析。具体执行行为与代码列表1-4 差异明显。

对于此代码，实现流水线的另一种方式是操作同一段数据loadDataFromFile和computeSqureSum给一个线程整体执行，看起来更简单，但会带来其他问题，读者可以尝试分析。

使用GPU运算版和异步版没有本质区别，只会CPU计算转移到上面GPU上面，如代码清单1-6 所示。在CUDA 中，GPU 内核是异步的，不需要异步就可以设计。IO如代码清单1-7 所示。

代码清单1-6 CUDA 异步IO实现流水线模式

loadDataFromFileAsync(file, 0, len, data0); for(int iter = 0; iter < numIter-1; iter ){ dataBuff = iter % 2 ? data0 : data1; loadDataFromFileAsync(file, iter 1, len, dataBuff); syncPreviousLoad(); data = iter % 2 ? data1 : data0; computeSqureSumGPU(len, data); } S5 sum = computeSqureSumGPU(len, (numIter-1)%2 ? data1 : data0);

由于CUDA内核执行是异步的，也就是说它会在那里GPU在计算执行完成前返回，主线程可以向下执行。

代码清单1-7 CUDA实现流水线模式

loadDataFromFile(file, 0, len, data0); for(int iter = 0; iter < numIter-1; iter ){ data = iter % 2 ? data1 : data0; computeSqureSumGPU(len, data); syncPreviousKernel(); dataBuff = iter % 2 ? data0 : data1; loadDataFromFile(file, 0, len, dataBuf); } S5 sum = computeSqureSumGPU(len, (numIter-1)%2 ? data1 : data0);

首先加载一次迭代的数据到data0中，iter=0时，computeSqureSumGPU函数使用data0运算，由于computeSqureSumGPU函数是异步的，因此CPU不会阻塞，会立刻返回接着执行加载数据到data1中，故此时 GPU 计算和数据加载同时进行。

由于loadDataFromFile使用的dataBuff是上一次computeSqureSumGPU使用的，因此需要同步，syncPreviousKernel函数即是为了解决这一问题。

在单个节点的多个进程间，如果进程需要从文件系统中加载数据进行计算，通常有两种模式可以使用：

1. 使用MPI的并行文件IO。这需要文件系统和MPI实现的优化，软件开发人员无法干涉。

2. 某个进程负责加载数据，然后将数据分发给其它计算进程。

考虑到展示的目的，本节使用2的方式，如代码清单1-8 所示。

代码清单1-8 流水线模式MPI实现

int NT; float psum = 0.0f; float sum; MPI_Comm_rank(&rank, MPI_COMM_WORLD); MPI_Comm_size(&NT, MPI_COMM_WORLD); MPI_Request request; if(0 == rank) loadDataFromFile(file, iter, lenNT, data0); MPI_Scatter(data0, len, MPI_FLOAT, data0, len, MPI_FLOAT, 0, MPI_COMM_WORLD); for(int iter = 0; iter < numIter-NT; iter += NT){ dataBuff = iter/NT % 2 ? data0 : data1; if(0 == rank) loadDataFromFile(file, iter, lenNT, dataBuffer); MPI_Iscatter(dataBuffer, len, MPI_FLOAT, dataBuffer, len, MPI_FLOAT, 0, MPI_COMM_WORLD, &request); data = iter/NT-1 % 2 ? data1 : data0; psum += computeSqureSum(len, data); MPI_Wait(&request, NULL); } data = iter/NT-1 % 2 ? data1 : data0; psum += computeSqureSum(len, data); MPI_Reduce(&psum, &sum, 1, MPI_FLOAT, MPI_SUM, 0, MPI_COMM_WORLD);

代码使用了MPI3新增加的异步 通信方式，保证异步的通信和计算能够同时进行，进而通过计算隐藏通信的延迟。不过实际上，这个设计存在严重的负载不平衡问题，进程0从硬盘中加载数据可能是较慢的步骤，而这一步骤的延迟却没有被隐藏。

对于流水线模式来说，通常存在如下问题：

n 无法流水，当计算或访存存在某些依赖时，则其计算必须顺序处理。在现实应用中，无法流水的情况基本上不存在。在笔者的印象中，几乎所有一开始认为不可以使用流水线模式的应用，较终都可以通过流水线模式改写。

n 访存和计算时间差距太大，流水没有效果。如果计算和访存时间差距太大，那么即使流水线完美建立，能够隐藏的时间也是两者中更小的，而对较终时间的影响并不大。换一种说法是：如果计算和访存时间差距太大，那么没有必要使用流水线模式。假设MPI通信时间为1s，计算时间为100s，那么隐藏延迟后总时间是100s，并不会比串行计算快多少。

n 库函数可能没办法异步，或者使用线程时，多个线程使用的资源之间存在竞争。异步IO接口和MPI的异步通信接口之间就不能组合出来一个新的异步接口。如果使用多线程达到异步效果时，多个线程可能要同时/不同时使用同一个资源，资源竞争可能会导致性能结果不如预期，而这类性能问题实际上很难调试。

对于流水线模式来说，一个不是问题的问题是：复杂的流水线模式很难构建；以基于CPU+GPU的异构计算为例，需要同时考虑CPU计算、GPU计算，PCIe数据双向传输，可能需要设计一个四阶段的流水线，这是一个巨大的挑战。

本节以NVIDIA CUDA为例，展示如何在GPU CPU和联接两者的PCIE总线间通过流水线模式发挥三者的计算能力。

GPU运算使用GPU硬件内部的计算单元，而CPU计算使用CPU核心和对应的存储器层次，通过PCIE总线的数据传输会使用独立的DMA（Direct Memory Access, 直接内存存取）存取引擎，这三种运算使用三个不同的硬件单元，因此是完全能够同时运算的。

NVIDIA CUDA通过流的概念很好的利用了流水线计算模式，CUDA 流可以看成一个队列，这个队列负责把其中的操作交给硬件执行。本节笔者就以NVIDIA CUDA为例，展示如何使用流水线模式加速计算。

在CUDA编程语言中，内核的执行是异步的，即调用它的线程会在它执行完成之前返回，接着向下执行，这引入了许多有趣的特性。本节将会展示如何利用此特性来让CPU和GPU同时计算，如代码清单1-9 所示。

代码清单1-9 CPU GPU同时计算

20 gpu_computing<<<grid, block, 0, stream>>>(len, d_in, d_out); 21 cpu_computing(len, in, out); 22 cudaStreamSynchronize(stream);

　　在CUDA 流 stream中执行的计算gpu_computing还在GPU上执行的时候，调用线程已经在接着执行cpu_computing计算了。有可能调用线程已经执行完cpu_computing的计算，但是GPU还没有执行完，此时需要调用cudaStreamSynchronize函数以等待GPU执行完成。20和21两行代码不能换位置，一旦调换位置gpu_computing就必须要在cpu_computing计算完成后才能开始进行。

对于CUDA来说，不同的流执行的操作可以并发执行。可以将计算交给一个流，通过PCIe总线的内存 操作交给另一个流，这两个流的操作便可同时执行，如代码清单1-10 所示。

代码清单1-10 GPU PCIe同时计算

20 gpu_computing<<<grid, block, 0, stream>>>(len, d_in, d_out); 21 cudaMemcpyAsync(d_out2, d_in2, size, cudaMemcpyDefault, stream2); 22 cudaStreamSynchronize(stream); 23 cudaStreamSynchronize(stream2);

调用线程启动gpu_computing计算后会立刻返回，接着去执行cudaMemcpyAsync，为了保证cudaMemcpyAsync不会等待gpu_computing的完成，这两个运算需要不同的流。

调用线程调用cudaMemcpyAsync之后，会接着执行，并不会等待这两个操作的完成，如果需要等待其完成，需要显式的调用cudaStreamSynchronize函数来等待两个操作的完成。

要保证CPU和PCIE同时计算，理论上可以有两种办法：

一、把CPU计算交给一个独立的子线程去做；子线程和主线程执行的PCIE操作可以同时执行。

二、使用CUDA流使PCIE操作异步进行。

为了简便起见，本节使用CUDA流完成，如代码清单1-9 所示。

代码清单1-11 CPU PCIE同时计算

20 cudaMemcpyAsync(d_out, d_in, size, cudaMemcpyDefault, stream); 21 cpu_computing(len, in, out); 22 cudaStreamSynchronize(stream);

调用线程发起的在流stream中执行的cudaMemcpyAsync操作会立刻返回，调用线程会接着执行cpu_computing，为了保证cudaMemcpyAsync的结果对调用线程可见，使用cudaStreamSynchronize保证了这点。20和21这两句的顺序不能调换，因为调用线程只有执行完成cpu_computing才会接着执行。

要想让GPU、PCIE和CPU能够同时执行，则必须要让三者中两两互不等待。对于GPU运算和PCIE操作都可以异步，但是它们之间不能等待，也不能等待CPU操作，这意味着：

一、同步的操作要在异步的操作后面，以保证同步操作不会阻碍异步操作完成。

二、两个异步操作之间要可以同时进行，即两者要使用不同的CUDA流。

满足这些条件的一个示例如代码清单1-12 所示。

代码清单1-12 CPU GPU PCIE同时运算

20 gpu_computing<<<grid, block, 0, stream>>>(len, d_in, d_out); 21 cudaMemcpyAsync(d_out2, d_in2, size, cudaMemcpyDefault, stream2); 22 cpu_computing(len, in, out); 23 cudaStreamSynchronize(stream); 24 cudaStreamSynchronize(stream2);

gpu_computing和cudaMemcpyAsync分别在stream和stream2中执行，两者互不依赖，能够同时执行，且调用线程不会等待将接着执行cpu_computing操作。

20和 21两者可以交换位置，因为一方面两者没有相互依赖，另一方面两者也不会阻塞调用线程。而22不能放到20或21的位置，因为其会阻塞调用线程直到计算结束。