【IT168 文档】总体性能优化策略

　　性能优化始终围绕三个基本策略：

　　1、最大化并行执行以获得最大利用率;

　　2、优化存储器使用以获得最大存储器吞量;

　　3、优化指令使用获得最大指令吞吐量。

　　对于应用的某个特定部分，什么策略会产生最好的性能收益依赖于这些部分的性能限制;如，优化存储器访问限制的内核的指令使用不会产生明显的效果。优化努力应当不变地被测试到和监视到的性能限制所定向，如使用CUDA profiler。另外比较某个特定内核浮点操作吞吐量或存储器吞吐量(更有意义)与对应的理论峰值吞吐量的差异可指出性能提升空间。

　　最大化利用率

　　尽量展现并行性并且有效的将并行性映射到系统的各个组件上以保证它们大部分时间都在工作，为了最大利用率应用应当以这为准则构建。

　　应用层次

　　在高层次上，应用应当通过使用异步函数调用和3.2.6节描述的流来最大化主机，设备和连接主机和设备的总线的并行执行。应当把每个处理器最擅长的任务分配给它：串行工作分配给主机;并行工作分配给设备。

　　对于并行工作，在算法中，由于某些线程为了与其它线程共享数据而同步导致并行性中断的点，有两种情况：或那些线程属于同一个块，这种情况下，它们只要使用__syncthreads()和在同一个内核调用中使用共享存储器共享数据，或属于不同的块，这种情况下，它们必须使用两个不同的内核调用以通过全局存储器共享数据，一个内核将数据写入全局存储器，另一个从全局存储器中读。第二种情况优化得比较差，因为它增加了额外的内核调用消耗和全局存储器通信量。为了减弱这种影响，应当以一种将线程间通信局限在一个块的方式将算法映射到CUDA编程模型。

　　设备层次

　　在低层次，应用应当最大化设备内多处理器间的并行执行。

　　对于计算能力1.x的设备，一次只允许一个在内核在设备上执行，所以内核至少要发射的块数要和多处理器数一样多。

　　对于计算能力2.0的设备，多个内核可在设备上并发执行，因此可以使用流来发射多个内核并发执行(参见3.2.6节)来获得最大利用率。

　　多处理器层次

　　在一个更低的层次，应用应当最大化多处理器内部的各种功能单元的并行执行。

　　如4.2节所描述的，GPU多处理器依赖线程级并行来最大化其功能单元的利用。利用率和常驻束数量直接相关。在每次指令发射时，束调度器选择一个已准备好执行的束并将下个指令发射给束内的活动线程。束准备执行下一条指令花费的时钟周期数称为延迟，如果在延迟期间，每个时针周期内，束调度器有一些指令可为某些束发射就可获得完全的利用，或换句话说，当每个束的延迟可被其它束完全隐藏。要多少条指令才能隐藏延迟依赖指令吞吐量。例如，使用基本的单精度浮点算术运算指令隐藏L个时钟周期的延迟(已调度到CUDA核心上)。

　　1、对于计算能力1.x要L/4(最近取整)条指令，因为多处理器每四个时针周期为每束发射一条该指令，如G.3.1节。

　　2、对于计算能力2.0要L/2(最近取整)条指令，因为多处理器每两个时钟周期为两个束发射两条指令，如G.4.1节。

　　束没有准备好执行它下条指令的最常见原因是指令的输入操作为没有准备好。

　　如果所有输入操作数是寄存器，延迟是因为寄存器依赖，如，一些输入操作数是由一些还没有完成的在前面的指令写的。在紧接着的寄存器依赖的情况下(一些输入操作数是由前面的指令写的)，延迟等于前面指令的执行时间且在此期间束调度器必须为不同的束调度指令。指令不同，执行时间也会变化，但是典型地大约22个时钟周期，等于计算能力1.x设备上的6个束，等于计算能力2.0设备上的11个束。

　　如果某些输入操作数在片下存储器中，延迟更高：400到800时针周期。在如此高的延迟期保持束调度器繁忙要求的束数依赖于内核代码;一般，如果没有片下存储器的操作数的指令(如，大多数时候算术指令)数与片下存储器的操作的指令比率小的话(这个比率经常称为程序的运算密度)，要求更多束。比如，如果比率是10，为了隐藏大约600个时钟周期的延迟，对于计算能力1.x的设备大约要15个束，对于计算能力2.0的设备大约30个束。

　　另一个束没有准备好执行下一条指令的原因是在等待一些存储器栅栏(见b.5节)或同步节(见B.6节)。一个同步点能强制多处理器空闲，因为许多束要等待同一块内其它束完成同步点之前的指令。一个多处理器上有多个常驻块能够减少这种情况下的闲置，因为不同块内的束不用在同步点等待。

　　对于给定的内核调用，一个多处理器上的块和束数目依赖调用时的执行配置(参见B.13节)，多处理器的存储器资源，和如4.2节描述的内核资源要求。为了帮助程序员基于寄存器和共享存储器要求选择合适的线程块尺寸，CUDA软件开发工具包提供了一个称为CUDA占有率计算器的电子表格，在这里，占有率定义为常驻块数目和最大常驻块数目之比(附录G为各种计算能力给出了数据)。

　　内核使用的寄存器数目对常驻束数目有显著影响。例如，计算能力1.2的设备，如果内核使用了16个寄存器且每个块有512个线程且要求非常少的共享存储器，这样多处理器可常驻两个块，因为它们要求2*512*16个寄存器，这匹配多处理器的可用寄存器数目。但是如果只要内核多使用一个寄存器，就只有一个块常驻多处理器上了，因为两个块要求2*512*7个寄存器，这大于多处理器的可用寄存器数目。因此编译器试图在保证寄存器溢出(原文应当是这样，是否错误?)的前提下最小化寄存器使用和指令数目。可以使用-maxrregcount编译选项或如B.14节描述的发射绑定控制寄存器使用。

　　一个块内需要的总共享存储器数目等于静态分配的和动态分配的共享存储器之和，另外在计算能力1.x的设备上，还包括传输内核参数需要的共享存储器用量(参见B.1.4节)。

　　寄存器、本地存储器、共享存储器和常量存储器的使用可在编译时使用—ptxas-options=-v选项报告。注意每个双精度变量(在支持本地双精度的设备上，如计算能力1.2或更高的设备)和每个long long变量使用两个寄存器。但是计算能力1.2或以上的寄存器总量是低计算能力的至少两倍。

　　如果可能每个块的线程数应当是束尺寸的整数倍以避免因为束占用不足而浪费计算资源。

　　如上面提到地，对于一个给定的内核，执行配置对性能的影响依赖于内核代码。推荐实验后确定。应用也基于寄存器文件和共享存储器尺寸参数化执行配置，这依赖于设备的计算能力，也依赖于多处理器的数目和设备存储器带宽，所有的这些都可以可以使用运行时API和驱动API查询。