并行信号处理技术-异构并行实时处理平台

        计算体系包括组成和硬件两个方面，随着计算机系统的发展，逐渐由单个处理器的简单计算系统发展为由多个处理器，甚至多个不同架构处理器组成的复杂计算系统。这篇文章就计算体系结构展开讨论，从计算框架特别是异构系统下的计算模式和并行计算结构设计的角度，分析异构并行计算架构设计和软件编程技术。

并行计算体系结构

不同层次的并行化设计已成为现代计算体系设计的推动力量，而能耗和成本是并行化设计的主要约束条件。应用程序中出现的并行模式有数据级并行（DLP）和任务级并行（TLP）。而计算机硬件在实现这两种不同并行模式时常采用指令并行、向量化运算、多线程和并发请求四种方法。这四种方法可追溯至50年前由Michae Flynn于20世纪60年代提出的Flynn分类法，计算系统可以分为如下四类：

• 单指令单数据（SISD, Single Instruction Single Data）；
• 单指令多数据（SIMD, Single Instruction Multiple Data）；
• 多指令单数据（MISD, Multiple Instruction Single Data）；
• 多指令多数据（MISD, Multiple Instruction Multiple Data）。

SISD（单指令单数据）硬件不具备并行性，逐条指令串行执行，每次处理一个数据，因此被称为单指令单数据计算系统。早起的计算机均属于SISD计算系统，如采用冯诺.依曼架构的IBM PC机。

图1  SISD结构示意图

SIMD（单指令多数据）设备在单个指令控制下能够同时处理多项数据，因此适合于数字信号处理、图形图像处理等向量运算较多的领域。Intel CPU中的MMXTM、SSE、SSE2及SSE3扩展指令集就是完成这类向量化处理工作，过去雷达信号处理领域常用的TS201就属于SIMD处理设备。图形处理器GPU在广义上也属于SIMD设备。

图2  SIMD结构示意图

MISD（多指令单数据）设备指同时使用多个指令来对同一个数据进行不同的处理。这种计算结构在发展中只作为理论模型出现而没有实际应用。

图3  MISD结构示意图

MIMD（多指令多数据）设备可以同时使用多个指令同时对多个数据进行的处理，也就是众核处理设备，最新的多核计算平台就属于MIMD结构。目前常用的多核心处理器属于MIMD设备，例如目前雷达信号处理领域较常用的C6678处理器和多核CPU。

图4  MIMD结构示意图

并发多任务开发模型

现代计算系统通常包含了多个处理设备，可能是多个同架构或不同架构的微处理器，其间通过高速总线连接，构成一个统一的计算系统，广义上都属于MIMD计算体系。在这样的计算模型下，有一些行之有效软件设计方法和编程架构，主要包括消息传递模型、共享内存模型、流模型和异构混合编程模型。

消息传递模型（MPI）

        消息传递模型（Message Passing Interface, MPI）是一种基于消息传递的并行编程模型[33]。消息传递作为标准的编程接口是一套可移植API，不同组织依据这套API有不同的实现，包括开源的MPICH、LAM MPI和商用的Intel MPI。在进行并行程序开发时只需要显式的使用相应的应用程序编程接口完成数据和消息的传递工作，而不需要考虑程序在具体硬件运算设备上执行。
在MPI编程模型中计算任务被拆分为不同的执行单元，各执行单元以独立进程的形式存在，具有完全独立的堆栈空间和程序段。各子任务（进程）间的通信通过MPI API完成，因此MPI标准不仅适用于多线程程序开发，更适用于多进程程序的开发。消息传递的物理实现取决于MPI的具体实现，可以是共享内存方式，也可以是网络连接方式，因此MPI编程模型能够将计算任务分发到不同的运算处理器甚至不同的运算节点上，也就是通常所言的分布式计算。在第五章中讨论的软件化雷达系统从系统顶层结构来看就属于这类分布式系统，可以使用MPI模型完成各运算节点间的逻辑组织。

图5  MPI系统典型结构示意图

图5显示了一个典型的MPI分布式系统，以通信域、进程组、上下文标识和数据结构为特征。通信域指定了通信操作上下文以确定其执行范围；不同进程在使用相同的预定义数据结构在同一个通信域中进行消息的发送和接受；各通信域间互不干扰，通过上下文标识进行区分，而属于同一个通信域的所有进程被称为一个进程组。

共享内存模型（OpenMP）

        共享存储器并行编程模型（Open Multi-Processing, OpenMP）是一种共享存储器的多线程程序设计方法。该模型要求所有存储器均被连接到同一个共享的地址段上，系统中所有处理器可以使用统一编址方式访问这些存储器。由于存储器共享，因此不同处理器间的数据可以被立即交换访问，大大提高了并行多任务处理程序的执行效率。OpenMP为开发人员提供了一种简单而安全的多线程开发方法，无需关心容易出错的线程操作。编程人员能够使用编译指导（Compiler Directing）或运行时库（Runtime Library）完成并行化。编译指导语句不改变程序源代码，在编译阶段完成并行域的识别拆分，具有较高的运行效率但对运行时的支持较少；运行时库需要对程序源代码进行并行优化，破坏了与原有串行程序的一致性，但提供了良好的运行时支持。

 

图6  OpenMP编程模型示意图

        OpenMP程序从主线程（Master Thread）开始运行，遇到并行域（Parallel Region）后分解成多个子线程开始并行执行。子线程在并行域执行完成后被销毁，程序回到主线程。本文在下述几章的程序设计过程中均使用这种并行设计的方法来提高程序运行效率。

流模型（Streaming）

流处理是对不间断数据连续处理的技术，其处理流程采用有向无环图（DAG）描述，DAG是对作业流程的图形表示，用来描述流处理任务的逻辑拓扑。
 

图7  流处理的DAG描述示意图

流处理作业系统有原生（Native）流处理和微批（Micro-batch）处理[39]两种实现方式。如图2.8，原生流处理对于一条记录（待处理数据）到达即立即处理，而微批流处理则会按照预定的规则（一般按照时间间隔）划分，积累成数据段进入流处理系统以采用突发方式进行处理。
    

图8  流处理系统处理流程

原生流处理（左） 微批流处理（右）

不同的流处理系统，在运行时、编程模型、函数式原语、状态管理、消息传输保障、容错、性能以及平台的成熟度和接受度方面有很大不同。目前主流的流处理系统主要有Apache Storm，Trident，Spark Streaming，Samza和Apache Flink。
Apache Storm由Nathan Marz及其团队于2010年在BackType开发，后被Twitter收购于2014年开源至Apache社区并逐渐成为工业标准。Storm属于原生流处理系统，提供低层次API，使用Thrift定义topology，支持包括Scala、Java在内的多种编程语言。
Trident是对Storm的高层次的抽象，简化了topology构建过程，增加窗口操作、聚合操作或者状态管理等高级操作，属于微批量流处理系统。对标于Storm的至少一次流传输机制，Trident提供了精确一次传输机制。Trident同样支持Java、Scala等编程语言。
Spark包含Spark SQL, MLlib和Spark Streaming，是目前最受欢迎的流处理框架。由于Spark的运行时（Runtime）是在批处理基础上构建的，因此Spark Streaming也依赖于批处理方式，属于微批流处理框架。Spark Streaming以高层次的声明式API供调用，支持包括Scala和Python在内的多种语言。
Samza是由LinkedIn公司开发的流处理框架，在基于消息队列的Kafka基础上构建的准实时流处理系统，提供组合式API，以作业作为基本处理流程串联处理。同样支持包括Scala在内的多种编程语言。
Apache Flink属于原生流处理系统，以高层次API提供给设计人员，也提供批处理方式，其批处理模式的下API与Spark接近，但其批处理原理与Spark不同。Flink将所有数据都抽象为一个流，更接近于现实逻辑。

表2  主流流处理系统特点比较

流模型	Storm	Trident	Spark	Samza	Flink
处理方式	原生	微批量	微批量	原生	原生
API	组合式	组合式	声明式	组合式	声明式
传输保障	至少一次	精确一次	精确一次	至少一次	精确一次
容错	反馈记录	反馈记录	RDD检查点	日志	检查点
状态管理	非内建	专用操作	专用流	状态操作	状态操作
延迟	极低	中	中	低	低
吞吐量	低	中	高	高	高
技术成熟度	高	高	高	中	低

异构混合编程模型（MOC）

本文硬件平台是典型的CPU+GPU异构系统，由于两种处理器的设计思路和实现结构的不同，适合于不同的处理。如下图，程序中的串行部分运行于CPU，而适合数据并行的部分运行于GPU。
 

图9  GPU加速的异构运算系统

借鉴图2.10显示的“天河一号”混合编程模式[40]，考虑算法运算组件的隔离性和可移植性，通常使用独立进程的方式运行，而算法内部使用CPU多线程并或GPU加速的方式，因此异构系统常采用MPI/OpenMP/CUDA的混合模式，充分利用了MPI在进程间通信优势和便利性以及OpenMP轻量线程组的开发方式，在运算密集部分充分利用GPU的数据并行优势使用CUDA加速[41]，保证了MIMD系统的高效性和开发的便利性，不同并行模型优势互补，本文称之为MOC（MPI OpenMP CUDA）模型。这样的异构并行化方式拥有三级并行粒度和存储结构，针对不同的数据和处理类型使用不同粒度的并行方式进行加速以达到最高的性能[32]。
 

图10  天河-1混合编程模型

MOC模型在编程上采用如下结构，不同计算任务间采用MPI解决粗粒度的数据传递和结果汇总，在子任务级别上进行了并行和流式设计；计算节点内部采用共享内存的OpenMP屏蔽复杂的线程操作，以轻量级线程组方式解决子模块级别的并行加速问题；而最高效的数据并行则采用CUDA加速。这种并行模型由于拆分合理、结构灵活而不受制于设备数量，具有良好的扩展性和通用性，非常适合第五章的软件化框架设计[33]。
 

图11  MOC混合编程模式下的三级并行粒度

GPGPU并行计算及CUDA架构

GPGPU编程模型

由于GPGPU的广泛应用，各厂商和组织为其开发了众多编程模型和设计框架，主流的有NVidia的CUDA，AMD的steam，微软的C++ AMP以及开源的OpenCL和OpenACC等。

(1) OpenCL

开放运算语言（Open Computing Language, OpenCL）是第一个兼容通用的开源开发平台标准，由Apple联合AMD、IBM、Intel、NVidia开发，其目的是统一CPU、GPU、DSP、FPGA等所有硬件加速设备组成的异构计算平台[44]，提供统一的并行程序开发标准。
 

图14  OpenCL框架

OpenCL框架由平台接口、运行时接口和内核编程语言组成。平台（Platform）包含宿主机（Host），运算设备（Device）和框架（Framework），多个平台可共存于同一计算系统内；平台API的作用是为运行时创建运行上下文（Context）；运行时（Runtime）为使用上下文的各种功能提供了调用接口，其首要任务是为硬件设备建立运行于主机或设备内的命令队列，提供内核对象、程序对象和存储器对象的调用依据。内核编程语言是C99标准的扩展集，用于编写直接控制硬件运算的内核代码，即运行于对应设备上的函数，该函数遵循SIMT编程模型，为运算任务指定了工作项（Work-item）、工作组（Work-group）和ND区域（ND-range）。
 

图15  OpenCL编程模型

(2) CUDA

统一计算设备架构（Compute Unified Device Architecture, CUDA）是由NVidia推出的GPU计算框架，该框架兼容OpenCL编译器和NVidia自定义C99扩展集，可以被编译为PTX代码交由GPU进行计算。在逻辑结构上CUDA将GPU硬件抽象为并行线程[47]，而不需要将数学运算映射到图形API，极大的方便了并行程序设计。CUDA拥有极其丰富的编程语言支持和运算库，大大方便了高层次的算法设计开发。
 

图16  CUDA编程框架的生态环境

CUDA由开发库、运行时环境和驱动组成。开发库包含了CUDA提供的各种应用支持，以API的形式提供给开发人员；运行时环境提供丰富的API以对接底层设备和运行时，包括存储器管理、设备管理、运行调度等；驱动通过对GPU硬件抽象对上提供了软件访问接口，图2.17显示了这三者和应用程序的关系。
 

图17  CUDA体系结构的组成

(3) OpenACC

为统一加速器并行软件开发环境，解决如CUDA、OpenCL等低层次语言在硬件隔离上的不彻底带来的编程难度，PGI、Cray、CAPS和NVidia四家公司联合制定了OpenACC应用编程接口标准，该标准借鉴了PGI Accelerator编程模型，同时融入了CUDA相关术语，在标准行文上与OpenMP相似[46]。
目前大多数的算法描述都针对于串行处理结构，因此在人工并行化时不仅需要对算法的运算原理和实现细节有非常深入的了解，而且需要对并行硬件平台及其软件架构的映射关系也有充分理解，才能有针对性的对算法进行并行化重设计以适应特定的并行硬件平台和软件框架，然而这种跨领域人才的缺失严重阻碍了并行计算在各领域的拓展，而OpenACC使用并行导语解决了这个矛盾。
 

图18  OpenACC的导语并行模式

OpenACC使用导语设计模式，允许编译器根据开发者在原有的串行代码上添加的编译导语产生低层代码，而开发者无需了解相应加速器的硬件结构即可完成程序的并行化设计，在不破坏原代码的前提下获得了良好的加速效果，相比于CUDA和OpenCL等低级语言显著提高了开发效率。由于OpenACC在设计初期就考虑了平台兼容性，因此同一份代码只需要经过重新编译即可无痛移植到不同的加速器平台。
 

图19  OpenACC在多平台上的一致通用性

统一计算设备架构CUDA

CUDA作为高层并行计算平台和编程模型，提供对NVidia GPU的高层次抽象，以利用GPU众多的处理单元进行高效数值计算。本节从编程模型、线程结构、存储结构和程序结构四个角度探讨CUDA的特性。

(1) 弹性编程模型

GPU在开发应用软件上的挑战是如何透明地扩展其运算规模以利用越来越多的处理器核心。CUDA并行编程模型采用三个关键的抽象[48]：线程组、共享存储和屏障同步的分层结构，解决了这个问题。
这种分层次抽象引导程序员将问题划分为逐个可解决的子问题，并将子问题进一步划分为一组线程可以协同解决的精细处理片段，并解决每个子问题时保持语言描述能力，并且同时实现自动可伸缩性[49]。可以并发或顺序地在GPU内的任何可用多处理器上调度每个线程块，使编译的CUDA程序可以在如图2.20所示的任意数量的多处理器上执行。
 

图20  CUDA的弹性编程模型

(2) 线程结构

并行程序以线程为基础单元，而线程束（warp）是GPU上线程组的最小单位，同一个线程束中的线程同步执行。一条指令只需要被读取一次就可以广播给同一个线程束中的所有线程，因此CUDA的指令运行模式被称为单指令多线程模式[50]。
线程根据最合适的数据共享方式被分为一维、二维或三维的线程块（Block），同一线程块中的线程必须位于同一流多处理器中，其线程间可通过Share Memory交换数据。线程块被组合为线程网格（Grid），网格中的线程块的数量由程序的并行设计需求确定而不受制于物理处理器的数量。线程由三维变量threadIdx标识，可以使用1D、2D或3D索引来识别线程，线程的索引及其ID以线性方式彼此对应。
 

图21  GPU的线程模型

这种分层的线程结构决定其应用的并行方式的两个层次：块间粗粒度并行和块内线程间细粒度并行。
细粒度并行：细粒度并行指块内线程并行，块内线程可通过Share Memory通信，并可使用同步命令__syncthreads()进行节奏同步，同时CUDA允许开发人员定义块内线程的结构。
粗粒度并行：粗粒度并行指块间并行，块间可以通过Global Memory通信，而线程块组成线程网格，同样CUDA允许开发人员定义线程网格的结构。

(3) 存储结构

CUDA设备拥有多个不同层次的存储器：全局内存（Global Memory）、本地内存（Local Memory）、共享内存（Share Memory）、常量内存（Constant Memory）、纹理内存（Texture Memory）和寄存器（Register），这些存储器的开放级别和访问规则均不相同[51]。每个线程有私有本地缓存和寄存器；每个线程块有共享内存，共享内存对块内线程可见；全局存储器对所有线程可见；常量内存和纹理内存是两个对所有线程可见的只读存储器。纹理内存可提供特殊的格式转换和寻址模式，可进行一些数据过滤操作。全局内存、常量内存和纹理内存与CUDA上下文具有相同的生命周期，可被多个核函数使用。
 

图22 CUDA的存储器模型

(4) 程序结构

CUDA采用异构编程方式，主机（CPU）和设备（GPU）分别有各自的存储器，主程序由CPU执行，当遇到描述数据并行的核函数（Kernel），将被编译成设备代码传送到GPU由多个CUDA线程并行执行。主程序调用核函数时须进行执行配置，以确定该核函数执行的线程结构和存储器分配。
 

图23 CUDA框架的主从程序结构

CUDA框架将主机代码和设备代码混合在同一工程中，由编译器和处理器进行区分，并分别编译为不同的执行文件[56]。nvcc将把两部分代码做预编译处理，使用不用的编译工具链进行编译链接，形成不同的二进制执行文件在不同平台上执行。
 

图24 CUDA混合应用程序的编译链接流程

CUDA提供Runtime API和Driver API两套封装。Driver API支持对底层硬件更加细致的控制，其调用硬件速度比Runtime API快且向后兼容，但两种API对于常见功能都有实现。在Runtime API接口之上，NVidia提供众多软件加速库，包括c++ STL的并行实现thrust API、blas并行实现cublas、FFTW并行实现cuFFT等。
 

图25 CUDA的软件层次

基于GPGPU的异构并行计算平台

本文异构并行运算平台采用CPU、GPU和FPGA共同协作完成计算密集型处理，三者通过PCI-e总线连接，CPU、GPU和FPGA拥有各自的存储系统，其间通过DMA操作实现数据的传递。这里CPU采用高性能的Intel至强多核系列，高速DDR4作为片外高速缓存；GPU采用NVidia高性能计算专用卡K80，以GDDR5作为片外高速缓存；FPGA采用Xilinx的K7系列，挂载DDR3作为高速片外缓存。采用SATA SSD组成的RAID阵列作为持久性数据存储介质，通过高性能RAID控制器接入PCI-e总线系统。对外接口采用千兆以太网和多通道自定义光纤，实现控制命令、处理结果合原始数据传输通道，平台采用如下结构：

图26  使用GPU加速的通用高性能计算平台

计算平台由商用1U机箱承载，具有CPU计算核心24个，CUDA计算核心4992个，DDR4内存64GB，GDDR5内存24GB，DDR3内存4GB，磁盘阵列采用8块SATA SSD组成的RAID5阵列，具有最高2.4GB/s的存取带宽，4.8TB的存储容量，兼顾了高速大容量存储的存储需求和数据安全，具有4组PCIe 3.0 x16总线接口，具有32GB/s的高速传输带宽。以24组最高5Gbps的可变速率光纤通道作为外部传输总线，4路千兆以太网用作外部控制和结果报送，兼顾了高速原始数据的出入和控制通路的灵活性。高性能通用计算平台具有全软件定义、网络中心控制、可扩展、可重配置、多功能、集成化、一体化等众多特点，是应用软件化信号处理的良好平台。
由于雷达信号数据率较大，并考虑到现有雷达前端多采用自定义光纤通道，本系统采用自定义光纤采集卡实现高带宽数字信号数据传输。自定义光纤通道具有可定制性强、带宽高、稳定性高和可靠性高等优势。当雷达信号进入采集卡后，由FPGA进行合并和原始数据解析，通过板载缓存进行数据缓冲，避免由于系统调度或吞吐量波动导致的数据丢失和数据溢出。由主机控制逻辑将数据通过PCIe接口传输至系统内存或采用GPU Direct技术直接进入GPU计算卡上高速存储器中。
光纤采集卡采用PCIe 2.0 x8接口设计，具有理论4GB/s，实测2.89GB/s的传输带宽。采集卡数据经系统内存可以同时分发给RAID阵列和GPU计算卡，多个模块间工作相互独立、并行工作，可以实现原始数据的同步存储和实时信号处理。RAID阵列卡可以将大吞吐量的雷达信号数据进行缓冲并带有冗余地写入多个磁盘中，既保证了写入的速度，也保障了数据安全性。目前SATA SSD能够达到450MB/s的速度，且相比过去的机械硬盘极大地减少了寻道时间。同步的雷达信号存储将有利于研究数据积累，存储的原始数据能够给算法开发提供大量可靠的实际数据来源。便于未来研究过程中的数据回放、数据仿真、系统模拟运行等功能。
 

图27  计算平台上的典型数据流程和控制流程

上图显示了在本软件化通用信号处理系统中典型数据流框图，绿色代表存储数据流，红色代表运算数据流，蓝色代表控制和结果数据流。GPU计算卡执行高度优化后的并行算法，随后通知雷达系统控制逻辑，数据通过PCIe通道被读取至系统内存，交由CPU进行下一步处理。GPU计算卡多用于多线程并行数据计算，CPU上算法的开发相对较为灵活，并且擅长于分支跳转等处理，对于数据量小、并行性不高、分支判断复杂的逻辑处理具有GPU无法比拟的优势。后期雷达信号处理，如航迹显示、目标跟踪、雷达数据记录等均可由CPU进行处理，少量大粒度的并行计算可以通过多线程技术在多路多核芯主板上提升系统计算性能[4]。同时也可以由CPU控制将其通过千兆以太网汇总至雷达显控终端。
在现有市场条件下如上所述的软件化信号处理平台在纯商用设备上实现（表2.2），采用AMAX XG-23201GK机架式服务器一台，主要运算部件为两颗Intel E5-2643v3和NVidia Kepler K80加速卡，配备内存64GB，采用2U标准服务器机箱承载。

表2  COTS通用计算平台产品清单

品名	类别	说明	单机部件
AMAX XG-23201GK	处理器	Intel Haswell E5-2643v3	2
	缓存模块	8GB DDR4 2133MHz ECC Reg	8
	存储模块	2.5" Intel SSD S3510 800G/SATA	8
	计算模块	NVidia Kepler K80	1

异构平台并行编程及优化技术

基于上节描述的硬件平台，本文采用独立存储器的异构编程方式，基于总线的通信方法，将三种差异性较大的平台，多核CPU、GPU和FPGA，尽可能的解耦合，减少不同平台间的数据通信，在有限的通信带宽下节约通信时间。
混合编程部分主要是多CPU和多GPU，对于CPU部分采用OpenMP的共享内存并行编程方式，OpenMP由编译指导语句、运行时库函数组成。对于GPU加速程序，采用NVidia的CUDA框架进行设计。官方提供了不同层次的API供开发者使用，CUDA的内核编程语言作为标准C99的扩展子集，可以让开发人员快速上手开发。下图显示了CUDA针对于不同语言的编程接口和不同级别的应用程序接口。
 

图28 采用CUDA加速的异构平台编程框架

基于CUDA的GPU并行应用程序，其执行流程可归纳为下载数据、执行核函数、上传结果三大步骤。具体来说有如下图显示的四个处理流程能：分配主机内存和设备内存。将运算数据从主机内存拷贝至设备内存。加载、启动核函数并执行运算。将设备内存中的运算结果拷贝回主机内存。根据运算流程，CUDA应用程序的性能瓶颈主要在访存带宽、主机接口带宽和指令吞吐量三个方面。下面就基于Tesla架构下CUDA程序的优化在这三个方面做出说明。
 

图29 CUDA架构编程模型

CUDA架构并行程序优化的目标是在最短时间、允许误差下完成运算任务[59]。这里“最短时间”特指数据吞吐量，而不是数据延迟。下面就影响程序加速性质的三个方面分别给出说明。

并行运算量影响加速性能

运算量的大小决定了指令吞吐量和I/O时间的隐藏。计算量过小导致GPU长时间处于空闲状态以等待数据I/O，因此小计算量的程序在性能上主要受限于指令吞吐量和IO操作的瓶颈[60]。具体来说，应用程序进行GPU加速后需要花相对大量的时间把数据从主机内存拷贝到设备内存，却只花了极少的时间用于数据的处理。这样的程序运行速度显然受限于PCIe的主机连接带宽。从另一方面来讲，算法的运算量中能够得到有效加速的部分是算法可并行部分，可并行部分占整个程序运行时间的比例决定了加速效果的上限，根据Amdahl定理，在固定负载（即运算总量不变）的前提下，公式(2-1)定义了程序的加速比。

  (2-1)

其中， 为串行部分时间； 为可并行部分串行运算的时间； 为处理器数量，这里可以理解为CUDA线程数量。
定义 为串行部分占比：

  (2-2)

其中 为程序运行总时间。
那么在理想情况下，并行化后程序的理想上界为：

  (2-3)

实际情况下，并行化后需引入数据转存、同步、归约、通信等额外开销，假设额外开销时间为 ，那么实际中并行化后程序的加速比为：

  (2-4)

显然可并行部分占比越大，加速比上界越高；并行中可隐藏的额外开销时间越多，加速比上界越高[62]。这需要待优化的程序具有良好的可并行性，且可并行部分的运算量足够大以使并行化后的运算时间足以隐藏额外开销时间。

高精度处理影响加速效果

目前的GPU架构更擅长于进行单精度甚至半精度处理，而高精度的运算则较为受限，同时除法、求模、求余数等运算的指令吞吐量也较低。然而在科学计算中往往需要更高的计算精度以保证数值安全和迭代的收敛性。这时需要权衡运算精度和加速效果。根据Gustafson定律，在使用相同数量的并行处理设备的前提下，要获得更高的运算精度，需要花费更多的时间。这种模型下运算加速比可定义为如下形式。

  (2-5)

这意味着加速比与并行处理器（CUDA线程数）的增加是线性的，同样在考虑额外开销的前提下可得实际加速比为：

  (2-6)

理论上这个加速比是不存在上限的，但实际系统中，受限于系统可提供的并行处理器，在CUDA中也就是线程数量的限制。

微批处理造成延迟

在现行的CUDA版本中线程最小的执行单元是一个wrap，即32个线程同步开始工作，对于确定运算量的算法CUDA更希望以2.2.3节中描述的微批量的流模式进行运算。而根据系统实时性要求选择合理的微批长度称为影响算法实现性能、实时性和吞吐量的重要指标。根据Sun-Ni定理[64]，在存储器可增加的情况下，增加微批大小以获得更高的单次运算密度，从而提高整体吞吐量，能够获得理论加速比如公式(2-7)所示，但是随之带来的缺点是更大的运算延迟，加速比可定义为如下形式。

  (2-7)

其中G(p)表示存储器增加p倍后增加的运算量。
因此微批大小的选择需要更具实际系统的实时性指标，在运算延迟能够满足实时性要求的前提下尽可能选择较大的微批规模，这样有助于提高运算、I/O比，减少数据传输交互次数，提升整体性能。

2优化总结

在需要较大运算吞吐量的系统应使用GPU进行加速处理，而对于实时性和吞吐量均由严格要求的系统，则需要ASIC、FPGA等逻辑实现才能满足要求，但是会带来更高的开发成本核研制周期。
综上所述，对于充分优化的CUDA程序具有如下几个特点：给定数据规模下，并行算法计算复杂度不显著高于传统最优算法。有足够的并行线程以获得高指令吞吐量。恰当了选择了合理的分级存储以获得最高的IO带宽。
因此，CUDA程序的开发和优化可遵循如下步骤进行：
1． 分解计算任务重的串行部分和可并行部分，对运算任务进行重新划分以确定使用的并行算法。
2． 按照任务划分确定数据的切割方式，设计实现高效并行算法确保获得较高的指令吞吐量。
3． 按照以上初步思路编写结果正确的并行程序作为优化的起点，确保正确使用存储器栅栏、线程同步等技术保证数据产品的准确性。
4． 利用寄存器、共享内存、纹理内存等技术优化访存，避免访存带宽成为性能瓶颈[65]。在访问瓶颈问题解决之前，更细致的优化方式很难取得理想的效果。
5． 通过线程标记、循环展开、原子操作、合理的同步等技术优化指令流，以获得最高的指令吞吐量[66]。
6． 调整存储资源分配，保证每个线程的数据使用量以确保在SM中有足够的线程在运行而不受制于资源不足而需要上下文切换导致性能损失[67]。
7． 由于与主机之间的PCIe带宽远比于主频和缓存速度小，因此需要优化主机通信以减少CPU和GPU之间的数据传输或者通过流式处理方式[68]隐藏数据传输时间以获得更高的运算效率。
8． 在极端情况下可以使用PTX汇编编写算法的关键部分，以对访存方式和指令吞吐量有细致的控制。