CUDA学习笔记（LESSON5）——GPU优化

CUDA系列笔记

CUDA学习笔记（LESSON1/2）——架构、通信模式与GPU硬件

CUDA学习笔记（LESSON3）——GPU基本算法（Part I）

CUDA学习笔记（LESSON4）——GPU基本算法（Part II）

CUDA学习笔记（LESSON5）——GPU优化

CUDA学习笔记（LESSON7）——常用优化策略&动态并行化

---

GPU优化

对于GPU的优化，我们有不同的方法，比如挑选一个好的并行算法、遵循高效的内存存取原则、优化存储体冲突（bank conflicts）以及位操作微观优化，后两者是两个ninja topic（ninja topic意思是很多时候在这方面投入精力进行优化并不能得到很大的效率提升，大部分时候不需要GPU程序员去关注），因此也不作为重点讨论。

下面讲一个GPU中常用的优化模式，称为APOD，一个好的GPU程序总是不断在分析（Analyze），并行化（Parallelize），优化（Optimize）与Deploy（应用）之间循环衍进的，我们往往会注重并行化跟优化的过程，但是需要注意的是应用与分析也是非常重要的一环，应用是观察最后能得到多少实际效果的检验，而分析则有利于我们做更进一步的优化。

我们还需要注意需要把优化热点（hotspot）放在不同的程序段上，例如我们对某一个程序段优化到一定程度以后再对它进行优化取得效率的提升将大大降低，因此这个时候继续选择这个优化热点并不是一个很好的选择，而跟聪明的方法是把热点放在其他程序段上

例子

下面让我们来看一个例子来阐述如何对一个程序进行优化，我们的任务是将一个矩阵内的元素进行转置操作。

下面我们来比较三种方法的运行时间，第一种方法是串行的方法，我们将转置的操作放在同一个线程中，用两层循环的形式来达到目的；第二种方法是按行并行化，意思是分配N个线程（假设矩阵为N×N），每个线程处理N个元素的转置；第三种方法是每个元素都开启一个线程来做转置操作。下面是最后的运行结果。

下面让我们想想这个0.67ms真的是达到速度的极限了吗？限制运行速度有两个因素，一个是对数据执行计算的时间，一个是内存存取的时间。我们现在所做的工作是将计算的时间压缩到了最小，那么内存存取的时间呢？我们可以通过deviceQuery这个内部接口来查看设备的吞吐量信息。我们可以分析得到理论的带宽与我们实际使用的带宽，从而能计算出带宽利用率（DRAM utilization），我们的带宽利用率并不是特别高。下面让我们来分析一下如何提高带宽利用率

提高global write efficiency

那我们怎么进行改进呢，很容易就想到coalescing。下面我们用nsight工具中的NVPP可以进行时间以及利用率的分析。我们可以看出方案三全局数据读取利用率（global read efficiency）达到了100%，但是全局数据写入利用率（global write efficiency）只有12.5%，这是为什么呢？原因就是我们实现了coalesced read, scattered write。在看代码之前我们要讲一下wrap的概念。

之前我们讲过了一个GPU中有很多个SMs（streaming multiprocessors），而一个SM中有很多个SPs（streaming processors），现在架构中一个SM中一般有8个SP，也就是说最多运行8个block，而SP运行线程的时候是以wrap为单位的，一个wrap中一般有32或16个线程，这些线程是完全并行运行的，只有当一个wrap运行完之后下一个wrap中的线程才会进入。而wrap中的线程是以x坐标为索引的，意思就是如果一个block中的线程是二维的（32×32），那么我们将取第一行32个线程作为一个wrap，然后取第二行32个线程作为一个wrap，以此类推。正是如此，当有一个wrap进行操作的时候，我们可以让其中的相邻线程采用coalescing的模式来提高存取效率。

下面让我们来看看第三种方案中核心的代码。

我们可以看出对于in数组而言，相邻的i对应的是相邻的线程，因此可以采用coaleced read模式，而写入out数组的时候相邻的i对应的线程是不相邻的，因此写入利用率会比较低。我们来看看如何解决这个问题。我们想到的方法就是将矩阵分为一个个块（tile），将它以行为索引，用coalesced read的方式读取到shared memory中，然后在shared memeory中做转置，之后再以行为索引用coalesced write的方式写入global memory中。

以下是核心代码，我们将K的大小由方案三的16改成了32以便满足一个wrap中 thread的个数。请注意在写入shared memory的过程中我们直接以转置的顺序写入了，以便节省一个同步操作。

但是我们这样设定以后方案三的运行时间增加了一倍（相比之前的0.67ms，变成了现在的1.17ms），而我们却没有对方案三进行修改，而只是修改了K的大小，而方案四也只比方案三快了一点点，那么这其中究竟发生了什么呢？

减小平均时延

在如上操作以后以后我们会发现我们的带宽利用率还是很低，那这到底是什么原因呢？我们首先来看一下限制GPU带宽的原则，叫做Litter's Law。它阐述了带宽、传输的字节与延时的关系。我们可以把传输的过程看做一个管道，因此当很多个线程同时在进行内存存取操作的时候才能够将管道填满以便提高传输的字节。我们最终的目标是想要提高最终的带宽利用率，请注意在下图中传输的字节中是有用的字节（useful bytes delivered），因此我们在方案四中做的改进是提高了字节的利用率，也就是提高了useful bytes delivered，借此方法我们想获得更高的带宽利用率，但是我们却发现提高并不是很明显，那么问题必然是出在了时延上。

问题就出在程序中的同步操作上，我们知道每个线程的运行时间是不一样的，这就意味着运行快的线程运行到同步操作的时候不得不停下来等待运行慢的线程，这样就导致了存取的时延只与最慢的线程有关，其他线程都得停下来等待最慢的线程。而解决方案有两个：减少一个block中的线程数（我们知道线程同步是在一个block内同步的）或者增加每个SM中的block数。

占用率（Occupancy）

占有率是指一个SM中实际运行的线程数与其可运行的最大线程数之比。对于每个SM，它的资源都是有限的，SM中最多运行的block、最多运行的thread、它自身的register的大小以及shared memory的大小都是有限的。可以在一个SM中同时运行的线程数目往往都取决于这几个参数。

根据一个SM中最大可运行的block数我们可以计算最大的thread运行的个数，从而能算出GPU的线程占用率是多少。

下面是一些常用的影响占用率的方法。

然而一味提高占用率也不是一件很好的事情。还是拿刚才转置矩阵做例子，如果减少了每个tile（block）中的线程数，那么一个tile内等待时间就会下降（总数量少了，从统计意义上讲长时间的线程就少了），带宽利用率也就提升了，与此同时一个tile内线程数少了，我们也就可以启用更多的线程来逼近占有率的极限（例如最大线程数是1536，如果一个tile中有1024个线程，那同时最多只能运行一个tile，占有率为66.7%；如果一个tile中有512个线程，那么同时可以运行3个tile，占有率也就提升至100%），我们可也以把提高占有率这件事情理解为同时运行的线程数多了，因此useful bytes delivered增加了，带宽利用率也就提升了。但是假如将一个tile划分的过小也就失去了划分tile的意义，也就是失去了coalescing存取模式的优势。因此这是一件需要权衡的事情，应该适当选择tile的大小。由此我们就可以解决之前提到的问题了，我们将K的大小改为16，就可以取得大约0.1ms速度的提升

下面让我们把程序在Udacity的IDE上运行来查看一下程序运行速度的变化以及带宽的利用率（之前得到的时间都是在笔记本上运行的结果）

由此我们也可以看出带宽利用率随着我们对代码不断改进而提高，而最后一步是解决了bank comflicts的问题，利用率有进一步提高，具体过程就不在这讲了。

优化计算性能

我们之前说了优化程序分两个方面：优化计算性能与优化存储性能，一般来说存储性能的低下是限制程序运行效率的瓶颈，但是也有的时候我们需要对程序的计算性能进行。主要方法就是以下两种，第一种我们已经谈过，这种方法不仅可以减少内存存取的时延，也可以减少计算过程中线程不必要的等待时间。

第二种方法是减少线程发散。我们在之前的博客也提过线程发散会导致不必要的等待时间，由于每个wrap中只有32个线程，因此由于线程发散导致的速度减慢最多是32倍。也就是说如果一个wrap中所有线程都采取同一分支，那么将不会产生发散。对于循环语句产生的发散也是类似的。

减少发散的方案有两种：尽量少写带分支的代码与注意线程之间工作量的不平衡。有的时候虽然我们的程序中有导致发散的switch与for语句，但是产生发散的线程实际上很少，这种情况程序的运行效率也不会下降太多，也就没有必要进行优化。所以，我们没有必要被线程发散吓到，而是应该具体情况具体分析，来进行适当的优化。

 

数学优化（math optimizations）

不同的数学操作也会花费不同的时间，例如，32位浮点数操作会比64位更快、用内置的函数例如__sin()可能会损失2-3b的精度但是却能让程序运行更快。这大概算half ninja topic，大部分时候我们无需太注意这件事。

 

主机-GPU交互

这部分主要讲的是系统层面的优化。我们要知道当要把一块CPU内存中的数据转移到GPU中时，我们首先要把这部分数据转移到一快staging area中，然后才能进行移动。为了避免这不必要的staging step，我们可以用cudaHostMalloc()，来分配一端pinned host memory，这段内存就能直接转移到GPU中了。或者我们可以用cudaHostRegister()，来把一段已经分配好的内存变成pinned host memory。

pinned host memory有两个好处，一个是上述说的运行更快，另一个就是允许cudaMemcpyAsync()操作，也就是允许在数据在CPU-GPU之间转移的时候Host中的程序能够继续运行。为了控制这个操作，我们可以引入流的概念。

流（Stream）是用来控制一系列按顺序执行的操作。这意味着同一个流内的操作是需要串行执行的，不同的流之间是可以并行执行的。对于下图右边程序的过程为：当程序运行到cudaMemcpyAsync时，会把它放进s1中，此时程序不会阻塞，而是继续运行，看到kernel A以后将其放进s2中，以此可以将下两个操作放入s3、s4中，此时这四个操作几乎是同时开始运行的。

流最主要的作用就是能够让数据的传输与计算同时进行，来提高GPU的效率。若不采用流，我们不得不先把数据搬到GPU，进行计算以后再搬回来，在这段搬运的过程中GPU实际上是闲置的，也就大大浪费了计算资源。