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實際上是閑置的，也就大大浪費了計算資源。