本章針對程式的性能優化，分解為如下7個步驟：
政策1： 了解問題，并正确分解問題為串行/并行的工作負載。
政策2： 了解并優化記憶體帶寬、延遲和緩存使用問題。
政策3： 了解與主機端傳輸資料的玄機。考察鎖頁記憶體、零複制記憶體的表現和某些硬體的帶寬限制。
政策4： 了解線程結構和計算能力，并了解它們對性能的影響方式。
政策5： 結合一些通用算法的優化執行個體，讨論如何實作算法。
政策6： 關注性能分析，定位應用程式的瓶頸所在及其來源。
政策7： 考察如何讓應用程式根據各種硬體實作自我調優。
9.1 政策1： 并行/串行在GPU/CPU上的問題分解
9.1.1 分析問題
首先需要考慮的是，對問題嘗試并行化是否是正确的解決方案。
9.1.2 時間
界定算法執行時間"可接受的"時間段是很重要的。目前可接受的并不一定意味着是最佳優化。
通常會發現任何優化活動都會有一定量所謂的“唾手可得的東西”。在優化中上述這些改變是很容易的且會帶來一個合理的加速。當這些被處理掉，就逐漸變得更難找到優化之處，隻有采用更複雜的充足才能進一步優化，這不僅花費更多時間而且引入更多潛在錯誤。
在設定适當的加速目标之後，在給定一組硬體的前提下，你必須知道什麼是合理的。
在考慮可接受的時間段是多少，問問你自己要達到該時間段還需付出多少努力。如果沒有超過2倍，往往就值得花時間來優化CPU的實作，而不是建立一個全新的、并行的方法來解決該問題。多線程引出了各種與依賴關系、死鎖、同步、調試等有關問題。
在過去的30年，購買更快的硬體就可。如果IO瓶頸、記憶體帶寬、計算吞吐量都可。
如果你決定指定GPU路線，那麼通常應該講你的設計目标設定為目前程式執行時間的10倍。你所達到的實際量級取決于程式設計者的知識、可用的程式設計時間以及應用程式中的并行程度。最後一個因素具有巨大的決定性。
9.1.3 問題分解
這個問題可以被分解成并行運作地組塊嗎？如果答案是否定的，GPU則不是要考慮的方法。你需要看CPU優化技術，比如緩存優化、記憶體優化、SIMD優化等。
假設你能夠把問題分為多個并發塊，那接下來應該并發多少個并發塊？CPU并行化地一個主要限制因素經常是沒有足夠大的粒度(或粗粒度)地并行工作要做。GPU運作成千上萬的線程，是以問題需要被分解成千塊，而不隻是像CPU那樣隻執行少數并發任務。
問題分解應該總是先從資料開始，然後考慮執行的任務。你應該試圖用輸出資料集來表示問題。你能否建構一個公式，描述如何由對應的輸入資料集中的資料轉換為資料集的某個輸出點嗎？如果可以做到，那麼問題轉換到GPU空間就相對容易了。
這一方法的問題在于，為了取得最好的效益，你需要完全了解問題。你不能簡單瞥一眼最高占用CPU的東西，然後把它們并行化。這一方法的真實益處在于把從輸入資料點到輸出資料點的鍊完全并行化。看到問題潛在的并行通常是第一個障礙。
優化通常用于執行資料的操作或函數。然而，随着計算能力相對于帶寬的飛速增長，資料是現在首要考慮的因素。
如果你打算用多個GPU或多個GPU，這裡最後一個考慮的如何在處理器元素上分解問題和資料集。以計算周期的角度來看，節點之間的通信将是非常昂貴，是以需要盡可能地将它最小化并跟計算重疊起來。
9.1.4 依賴性
依賴就是一些計算需要用到以前計算的結果，可以是針對問題域的計算也可能是數組下标的計算。
依賴關系主要由兩種形式，要麼一個元素是依賴于它附近若幹元素，要麼在多輪周遊資料集時下一輪依賴目前輪。
extern int a, c, d;
extern const int b, e;
void func_with_dependencies(void)
{
a = 100 * b;
c = 1000 * b;
d = (a + c) * e;
}
依賴：鑒于一個指令的傳回結果必須注入到下一個指令，這種類型的代碼排放方式隻允許較小的并行性并會導緻流水線的失速。處于失速時，CPU與GPU将處于閑置狀态。CPU與GPU都使用多線程來隐藏這個問題。
在CPU方面，來自其他虛拟CPU核的指令流填補指令流水線的空隙(如超線程技術)。然而，這要求CPU知道流水線中的指令屬于那個線程，但這會使硬體變得複雜。在GPU上，也使用多線程，但采用時間切換方式，這樣算術運算的延遲時間被以極小甚至可以沒有代價的隐藏掉。 事實上，GPU上你需要20個時鐘來隐藏這樣的延遲。然而，這種延遲不一定來自另一個線程。
extern int a, c, d， f, g, h;
extern const int b, e;
void func_with_dependencies(void)
{
a = 100 * b;
c = 1000 * b;
f = b * 101;
g = b * 1001;
d = (a + c) * e;
h = (f + g) * e;
}
對上上述代碼，注意，如果你在計算變量a，c與使用它們計算d之間的位置插入一些獨立的指令，将需要更久的時間才能獲得d的計算結果。計算f，g和h的值與計算d是重疊的。實際上，你是通過重疊非依賴指令達到隐藏算術運算的延遲。
循環融合(loop fusion')是一種處理依賴關系并引入額外非依賴指令的技術。如下所示：
void func_with_dependencies(void)
{
unsigned int i, j;
a = 0;
for ()
{ a = ???}
for ()
{ d = ???}
for ()
{ a =???}
}
兩個計算的疊代空間是互相重疊的。是以可以将一個計算的一部分移動到另一個計算的循環體内部。這樣就可以引入額外的，無依賴性的指令，另外能夠降低總體的循環次數。循環疊代不是免費的，因為他們需要一個循環疊代值和一個分支。是以，降低的疊代次數會為我們在減少執行的指令數方面帶來顯著益處。
現在的GPU，很可能将這些循環展開，放到線程内，并由單個核心程式計算。然而，謹慎使用這種方法。通過執行這些操作，你同時減少了可用于線程/線程塊排程的整體并行度。如果這個數目很小的話，會浪費執行時間。另外要注意，使用融合的核心時，通常會消耗更多的臨時寄存器。由于寄存器的使用增加了，會限制一個SM上可排程的線程塊數目，進而可能會限制實際可融合的數量。
最後，你應該好好考慮需要多輪周遊的算法。他們通常被實作為一些核心調用的序列，每一次調用在資料上循環一遍。由于每輪要讀/寫全局資料，效率通常較低下。許多這樣的算法可以寫成隻涉及單個或少量目标資料點的核心程式。這為把資料放入共享記憶體或寄存器提供了可能，并且相較給定核心需要多次全局記憶體通路的方式，可以大大提高完成的工作量，這将明顯改善多數核心的執行時間。
9.1.5 資料集大小
| 資料集的大小使選取問題的解決方案差别巨大 |
| 一個典型的CPU實作可以分為以下幾塊： | |
| 資料集小于一級緩存 | 16KB~32KB |
| 資料集小于二級緩存 | 256KB~1MB |
| 資料集小于三級緩存 | 512KB~16MB |
| 資料集小于單台主機記憶體大小 | 1GB~128GB |
| 資料集小于主機端持久性存儲大小 | 500GB~20TB |
| 資料集分布在多台機器上 | >20TB |
| 一個典型的GPU實作可以分為以下幾塊： | |
| 資料集小于一級緩存 | 16KB~48KB |
| 資料集小于二級緩存 | 512KB~1536MB |
| 資料集小于GPU記憶體大小 | 1GB~128GB |
| 資料集小于主機端持久性存儲大小 | 500GB~20TB |
| 資料集分布在多台機器上 | >20TB |
對于非常小的問題集，可以增加更多CPU核，可能會帶來超線性加速比。如果将問題從記憶體移到三級緩存，或從三級緩存移到二級緩存。會看到一個很明顯的加速。是因為使用的緩存有高得多的記憶體帶寬。
GPU的主要問題不是緩存，而是你能在一張卡上儲存多少資料。将資料從主機系統傳入或傳出會耗費大量計算時間。為了隐藏這個時間，你應該把計算與資料傳輸重疊起來執行。更好是利用主機的鎖頁記憶體同時做到傳入與傳出資料。由于鎖頁記憶體不會被虛拟記憶體管理系統換出，是以它必須是存放于主機的真正的DRAM記憶體。
在商業硬體上，你可用空間比總量要少一些。
在主機端，你的記憶體至少需要與輸入和輸出緩沖區配置設定的鎖頁記憶體等量。由于你通常使用最多2GB的鎖頁記憶體，是以剩餘的記憶體量可用輕松地支援多個GPU，多數系統支援至少2個GPU卡。
當資料集可能因為計算、記憶體、存儲或者能源方面的因素，無法置于單台機器時，你必須考慮使用多個節點。這就需要節點間的通信了。節點間的通信是非常耗時的，相比于任何内部的資料通信至少慢一個數量級。此外，你還必須掌握一套API。盡可能避免節點間通信這一步驟。
9.1.6 分辨率
提高問題的分辨率是否比提高速度更有吸引力？一個更精準的結果在你的問題裡能得到什麼？
9.1.7 識别瓶頸
1. Amdahl定律
Amdahl定律告訴我們，當資料流中任然存在串行執行元素時，将限制速度的提升。
無限擴充程式的唯一辦法是消除程式執行中的所有串行瓶頸。
某程式通過運作大量工作，做的很好，當突然遇到一個串行點或同步點，一切都堵塞了。針對這類問題，把瓶頸部分并行化就能解決。
在考慮直方圖計算，你會看到如果把所有線程都加入同一個桶，就形成了同樣的瓶頸。通常會采用原子操作，這樣一組并行線程就要串行執行。相反，如果配置設定給每個線程屬于它自己的一組桶，然後再将這些桶合并起來，就能消除串行瓶頸問題。
2. 分析
分析是确定你目前在哪兒以及應該在什麼地方多花點時間的最有用的任務之一。
優化應該根據确切的數字和事實，而不是猜測那可能是最應該優化的地方。Nvidia提供了CUDA Profiler和Parallel Nsight，以提供分析資訊。
分析器用過讀取硬體電腦，來發現代碼花費的時間和在GPU上的占有率。它會提供非常有用的資料，如總共合并讀和寫次數、緩存命中/失敗率、分支頻率、線程束串行化程度。
使用分析器做完一個初步的檢查之後，你應該先檢視花費總時間最多的代碼段。使用分析器做完一個初步的檢查之後，你應該先檢視花費總時間最多的代碼段。典型的未優化的程式，80%的時間花費在20%的代碼上的。優化20%的代碼是有效減少使用時間的關鍵，分析器是确定這20%代碼所在的一把鑰匙。
當然，一旦上述問題已被優化為最佳，如果不進行重新設計，後面的為提供加速化的進一步優化将會變得越來越耗時。
9.1.8 CPU與GPU的任務分組
事實上，最好的應用程式往往可以充分利用CPU與GPU兩者的優勢，并相應的劃分資料。任何基于GPU的優化也應該考慮CPU，因為這對于總的應用程式時間很重要。可以使用的CPU核數越多，通過分流一些工作給CPU的潛在收益越大。
如果說CPU可以處理GPU的工作1/10，那麼當僅需要3個CPU核時，你的GPU就獲得額外30%的吞吐量。
對于IO限制而言，因為引入更多的CPU的線程或程序，經常可以顯著提高整體的IO吞吐量。這似乎很奇怪。因為IO裝置的輸入輸出上限決定了它的吞吐量。在現代大記憶體的計算機中，大多數的IO操作都是進行緩存的。是以，IO操作大都在記憶體上移動而不是裝置上移動。
獨立的CPU程序或線程可以建立一個獨立的GPU上下文，且啟動它們自己的核心到該GPU中。這些額外的核經常以隊列的方式在GPU中去執行。當所需的資源變為可用時，核心開始執行。
GPU的空閑時間比CPU空閑時間更昂貴，因為它的吞吐量通常在CPU時間的10倍以上。
通過在一個GPU中放入多個核心，這些核心就可以伺機占用空閑硬體槽位。這将在一定程度上增加第一組核心的延遲，但會大大提高應用程式的整體吞吐量。
通過引進一對程序，巧妙的重疊了IO，CPU，GPU和傳輸時間、整體吞吐量獲得了顯著的改善。
程序允許設定為處理器關聯(processor affinity)，可以把程序綁定到一個給定的CPU核。這樣做往往會提供性能，因為它可以更好地重用該核的緩存。
選擇線程還是程序，在很大程度上取決于CPU之間需要同步任務的個數。
在權衡CPU/GPU的使用過程中，也需要知道如何最優地劃分任務。當資料時稀疏分布的或者是小資料集的時候，CPU很擅長處理這類串行任務。
有時會看到CPU用在規約操作的最後階段。通常幾輪疊代之後，歸約操作涉及的元素數會下降為原來的一半。可供排程的線程數量就小于一個GPU可供排程的最大線程數了。如果再繼續疊代幾輪，一些SM就開始閑置。
是以，一種優化政策是，當疊代到一定的門檻值，剩餘部分的計算就轉交給CPU來完成。不過，從費米架構之後，英偉達解決了上述問題。能夠讓那些空閑的SM在一個排隊的核心中使用。但是，要讓SM變得空閑，必須要保證其上的所有線程塊已經完成它們的任務。一些非最優的核心可能殘留一個或者數個活躍線程(即使在規約操作的最後一層)，進而導緻核心牽制住該SM，直至整個規約操作完成。
對于類似歸約操作的一些算法，請確定每次疊代都在減少活躍線程束的數量，而不單單是活躍線程的數目。
9.1.9 本節小結
了解問題并基于你的編碼時間和熟練程度定義你的加速目标。
識别問題中的并行性，并思考如何以最佳方式在CPU和一個或多個GPU之間配置設定。
考慮一下，是較少的執行時間還是處理資料以獲得更高分辨率更重要。
了解任何串行代碼的實作，并思考如何處理它們最合适。
分析你的應用程式，以確定你的了解确實反映了實際情況。如果可以幫助你加強了解，請重複你之前的分析。
9.2 政策2：記憶體因素
9.2.1 記憶體帶寬
記憶體帶寬和延遲是所有應用程式都要考慮的關鍵因素，尤其是GPU應用程式。帶寬是指與某個給定目标之間傳輸的資料量。在GPU的情況下，我們主要關心的是全局記憶體帶寬。延遲則是指操作完成所用的時間。
GPU上的記憶體延遲設計為由運作在其他線程束中的線程所隐藏。當線程束通路的記憶體位置不可用時，硬體向記憶體送出一次讀或寫的請求。如果同一個線程束上其他線程通路的是相鄰記憶體位置并且記憶體區域的開始位置是對齊的，那麼該請求會自動與這些線程的請求組合或者合并。
我們需要考慮的一個關鍵領域是運作過程中的記憶體事務的數量。每一個記憶體事務被送入一個隊列中然後由記憶體子系統單獨執行。這當然會有一些開銷。一個線程一次送出對4個浮點數或整型數的一個讀操作比送出4個單獨的讀操作花費的代價更小。
為接近峰值帶寬，你可以采取兩個方法。方法1，使用線程束完全加載處理器，實作接近100%的占用率。方法2，通過float4/int4向量類型使用128為讀操作，此時占用率小了很多，但任然能達到100%的峰值記憶體帶寬。
9.1.2 限制的來源
核心通常被兩個關鍵因素限制：記憶體延遲/帶寬和指令延遲/帶寬。正确了解這兩類關鍵限制因素中哪一種正在限制系統的性能，對于指導你合理的配置設定精力是很關鍵的。
最簡單的能夠看到代碼平衡位置的方法，是簡單地注釋掉所有算術運算，然後直接指派成結果代替。算術指令包括所有的計算，分支，循環等操作。如果存在某種歸約操作，你隻需要将它替換成普通的求和操作。一定要確定包括所有從記憶體中讀取到最終輸出的參數，否則編譯器将删除明顯冗長的記憶體的讀寫操作。對核心重新定時，你會看到花在算術和算法部分的近似的百分比。如果這個百分比很高，那麼你就受到了算術限制。反之，你受到記憶體限制。
此外，如果記憶體模式沒能很好地合并，GPU将不得不串行執行指令流以支援分散的記憶體讀寫。如果是這種情況，那麼有可能需要重新安排記憶體模式，以允許GPU将線程的記憶體通路模式合并。
是否可以擴大一個單一線程處理的輸出資料集的元素數目呢？這通常同時有助于記憶體受限型和算術受限型的核心。如果你這樣做，請不要線上程中引入循環，而是要通過複制代碼實作。如果代碼是很重要，這也可以作為裝置函數或宏來實作。確定将讀取操作提前到核心開始處，這樣在需要資料時就已經完成了對它們的讀取。這将增加寄存器的使用，是以一定要監控正在被排程的線程束個數以確定它們不會突然地退出。
至于算術限制的核心，檢視源代碼并思考如何将其翻譯成PTX彙編代碼。不要害怕實際産生的PTX代碼。數組索引通常被替換為基于指針的代碼，将速度較慢的乘法替換成更快的加法。使用2的幂次的除法和乘法指令分别可以被替換成速度更快的右移或左移位的運算。循環體中的所有常量(不變量)應該被移到循環體外。如果線程包含一個循環，那麼展開循環通常會實作加速。
9.2.3 記憶體組織
在許多GPU應用程式中，使用正确的記憶體模式往往是關鍵的考慮因素。CPU程式通常在記憶體中以行的方式安排資料。我們必須嘗試安排記憶體模式以使連續線程對記憶體的通路以列的方式進行。此原則同時适用于全局記憶體和共享記憶體。
cudaMalloc函數以128位元組對齊的塊為機關配置設定記憶體。如果使用結構會越過這個邊界，有兩個辦法。首先，你可以在結構中添加填充的位元組。或者，你可以使用cudaMallocPitch函數。
對齊是一個很重要的的标準，它将決定記憶體事務或緩存行需要擷取一次還是兩次。
通常，使用共享記憶體作為臨時緩沖是明智的。然後，可以将其用于對全局記憶體進行合并的讀寫操作。
9.2.4 記憶體通路以計算比率
記憶體操作與計算操作的比率是值得思考的問題。你所期望的理想比例至少是10:1。也就是說，對于每一個核心，從全局記憶體執行的讀取操作需要執行10條或更多的指令。這些指令可能是數組索引計算、循環計算、分支或條件判斷。每個指令都應該對有效地輸出起到一定的貢獻。特别是循環沒展開時，它經常會增加指令開銷但并不會助于任何有用的工作。
是以，在每個周期内線程束排程器送出2條指令或4條指令。由于這些指令來自于不同的線程束，它們之間來自于不同的線程束，它們之間是互相獨立的，是以将它們放入執行單元(CUDA核，SFU和LSU)流水線中。
基于切換其他線程束的能力，使用最少個數的常駐線程束無法隐藏記憶體指令或指令延遲。指令流的失速實際将會使CUDA核失速，這是我們非常不願意看見的。實際上，多個線程塊會被配置設定到一個SM上，以試圖確定這個問題永遠不會發生并且更重要的是生成各種形式的混合指令。
第二個要點是共享的資源(SPU，LSU)限制了持續執行相同操作的能力。由于CUDA核和LSU都被納入了流水線中，但是它們隻有16個單元寬度。是以，将線程束排程到這兩個單元之一會花費兩個周期。
當資料流中有全局記憶體寫操作時，你需要将讀操作提前至核心開始處。試用一下代碼：
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int a0 = a[tid];
int b0 = b[tid];
data[tid] = a0 * b0;
我們有兩個辦法，标量方法或向量方法。GPU隻在硬體上支援向量的加載和儲存而不支援向量操作。是以，乘法操作實際上要像C++中的重載操作符一樣完成并且隻是将兩個互相獨立的整數相乘。然而，向量方法分别執行兩個64位加載和一個單獨的64位存儲而不是非向量版本的4個獨立的32位加載和一個32位存儲。是以，40%的記憶體事務内節省了。記憶體帶寬使用是相同的，但是更少的記憶體事務意味着更小的記憶體延遲。是以，等待記憶體的總體失速時間減少了。
為了使用向量類型，隻需聲明數組為向量類型int32。被支援的類型為int2、int3、float2...。當然可以建立自己的類型，并且定義自己的操作符。每一個向量類型實際上都是一個對齊的結構體，包含N個聲明為基類型的成員元素。
是以，希望你能真正看到不同類型的指令之間是需要平衡的。
重要的是要認識到，指令流需要足夠的計算密度以充分利用SM上的CUDA核。核心隻是簡單地執行加載/存儲操作和少量其他工作，是以無法達到裝置的峰值性能。通過每個元線程處理2個、4個或8個元素，進而擴充這些核心使其包括獨立的指令流。是以盡可能使用向量操作。
9.2.5 循環融合和核心融合
另一個可以顯著地節省記憶體帶寬的技術是基于9.2.4節提到的循環融合。循環融合是隻兩個明顯獨立的訓話在一段範圍内交錯地執行。
核心融合是循環融合的演變。如果你有一系列按順序執行的核心(一個接着一個執行)，這些核心的元素能否融合？ 對于那些還未完全了解的核心，這樣做的千萬要小心。調用兩個連續的核心會在它們之間生成隐式地同步。
開發核心時，将操作分解成幾個階段或幾輪是常見的。例如，第一輪你看你針對整個資料集計算結果。第二輪，你可以使用特定的标準對資料進行過濾，然後在特定的點進行深入的處理。如果第二輪能夠本地化一個線程塊，那麼第一輪和第二輪能夠組合成一個單獨的核心。這就消除了将第一個核心寫入主存，随後讀取第二個核心的操作及調用核心的額外開銷。如果第一輪能夠将結果寫入共享記憶體，那麼隻在第二輪需要這些結構，這樣就完全消除了對全局記憶體的讀取、寫入。歸約操作經常被劃分到這一類并且能從這樣的優化中顯著的受益，因為第二階段的輸出通常比第一階段的輸出小很多，是以它顯著節約了記憶體帶寬。
核心融合技術如此有效的原因是它所帶來的的資料重用。一旦資料存儲到共享記憶體或寄存器集中，那麼盡可能重用它。
9.2.6 共享記憶體和告訴緩存的使用
相比于全局記憶體，使用共享記憶體可由提供10:1的速度提升。但是共享記憶體的大小是受限的。
在資料集上疊代的核心如果沒有重用資料，那麼需要意識到它們可能正在以低效地方式使用緩存或者共享記憶體。
與在一個大型資料集執行多輪不同，核心融合這樣的技術可用于在資料間移動而非多次傳入它。思考一下輸出資料的問題而不是輸入資料。建構該問題是将線程配置設定給輸出資料項而不是輸入資料項。在資料流方面，建立流入而非流出。優先選擇聚集(收集資料)(gather primitive)原語而不是分散原語(scatter primitive)。GPU會同時從全局記憶體和二級緩存直接将資料廣播到每個SM，這一點支援高速度聚集型的操作。
如果資料項很小，則記憶體事務可以逐漸将規模減少讀取，直至每次通路32位元組。是以，從十分分散的記憶體區域通路一個資料元素的核心，在任何基于緩存的架構，包括CPU和GPU，表現會十分糟糕。原因在于單個元素的讀取會載入128位元組的資料。對于大多數程式而言，存入緩存的資料會在下一次循環疊代中命中，這是由于程式常常通路與之前通路的資料臨近的資料。是以對于大多數程式，這是一個顯著的優點。但是，對于那些隻需要單個資料元素的程式來說，剩餘的124位元組是多餘的。對于這種核心，你需要為記憶體子系統制度去所需的記憶體事務而不是緩存行的大小，隻能在變異的時候通過-Xptxas-dlcm = cg标志來完成此工作。這将所有的通路減少到每次事務32位元組并且令一級緩存失效。對于隻讀資料，考慮使用紋理記憶體或者常量記憶體。
9.2.7 本節小結
仔細考慮你的核心處理的資料并且如何将其以最佳的方式安排在記憶體中。
針對128位元組的合并通路，優化訪存模式，對齊到128位元組的記憶體讀取大小和一級緩存行大小。
注意權衡單精度和雙精度對其記憶體使用的影響。
在适當的時候将多核心合并成單核心。
以最适當的方式使用共享記憶體和緩存，以確定你能充分利用更高計算能力裝置上擴充容量。
![POJ 1284 Primitive Roots (歐拉函數&原根定理)[圖]](data:image/gif;base64,R0lGODlhAQABAIAAAP///wAAACwAAAAAAQABAAACAkQBADs=)