本文系轉載，介紹了常識性的cuda優化策略，雖然有些已經過時，但思想不會過時。
1. memory coalescing，保證內存融合。因為global memory在CC為1.x上是按照half wrap進行訪問讀寫的，而在2.x上是按照wrap進行訪問讀寫的。在顯存中，有多個存儲器控制器，負責對顯存的讀寫，因此，一定要注意存儲器控制器的負載均衡問題。每一個存儲器控制器所控制的那片顯存中的地址空間稱為一個分區。連續的256Byte數據位於同一個分區，相鄰的另一組256Byte數據位於另一個分區。訪問global memory就是要讓所有的分區同時工作。合並訪問就是要求同一half-wrap中的thread按照一定byte長度訪問對齊的段。在1.0和1.1上，half-wrap中的第k個thread必須訪問段里的第k個字，並且half-wrap訪問的首地址必須是字長的16倍，這是因為1.0和1.1按照half-wrap進行訪問global memory，如果訪問的是32bit字，比如說一個float，那么half-wrap總共訪問就需要16個float長，因此，每個half-wrap的訪問首地址必須是字長的16倍。1.0和1.x只支持對32bit、64bit和128bit的合並訪問，如果不能合並訪問，就會串行16次。1.2和1.3改進了1.0和1.1的訪問要求，引進了斷長的概念，與1.0和1.1上的端對齊長度概念不同，支持8bit-段長32Byte、16bit-段長64Byte、32bit-64bit-128bit-段長128Byte的合並訪問。對1.2和1.3而言，只要half-wrap訪問的數據在同一段中，就是合並訪問，不再像1.0和1.1那樣，非要按照順序一次訪問才算合並訪問。如果訪問的數據首地址沒有按照段長對齊，那么half-wrap的數據訪問會分兩次進行訪問，多訪問的數據會被丟棄掉。所以，下面的情況就很容易理解：對1.0和1.1，如果thread的ID與訪問的數據地址不是順序對應的，而是存在交叉訪問，即：沒有與段對齊，那么，就會16次串行訪問，而對1.2和1.3來講，會判斷這half-wrap所訪問的數據是不是在同一個128Byte的段上，如果是，則一次訪問即可，否則，如果half-wrap訪問地址連續，但橫跨兩個128Byte，則會產生兩次 傳輸，一個64Byte，一個32Byte。當然，有時還要考慮wrap的ID的奇偶性。1.2和1.3放寬了對合並訪問的條件，最快的情況下的帶寬是最好的情況下的帶寬的1/2，然而，如果half-wrap中的連續thread訪問的顯存地址相互間有一定的間隔時，性能就會灰常差。比如，half-wrap按列訪問矩陣元素，如果thread的id訪問2*id的地址空間數據，那么，半個wrap訪問的數據剛好是128Byte，一次訪問可以搞定，但是，有一半數據會丟失，所以，也表示浪費了帶寬，這一點一定要注意。如果不是2倍，而是3倍、4倍，那么，有效帶寬繼續下降。在程序優化時，可以使用share memory來避免間隔訪問顯存。
2. bank conflict，bank沖突。先說一下，share memory在沒有bank conflict情況下，訪問速度是global和local的100倍呢，你懂的。類似global memory的分區，share memory進行了bank划分。如果half-wrap內的很多thread同時要求訪問同一個bank，那么就是bank conflict，這時，硬件就會將這些訪問請求划分為獨立的請求，然后再執行訪問。但是，如果half-wrap內所有thread都訪問同一個bank，那么會產生一次broadcast廣播，只需要一次就可以相應所有訪問的請求。每個bank寬度長度為32bit。對於1.x來講，一個SM中的share memory被划分為16個bank，而2.x是32個bank。1.x的bank conflict和2.x的bank conflict是不一樣的。對1.x來講，多個thread訪問同一個bank，就會出現bank conflict，half-wrap內所有thread訪問同一個bank除外。但是，對2.x來講，多個thread訪問同一個bank已經不再是bank conflict了。比如：
char data = Sdata[BaseIndex+tid];
在1.x上屬於bank conflict，因為，0~3thread訪問同一個bank，4~7訪問同一個bank，類推，這種情況屬於4-way bank conflict。但是，對於2.x來講，這種情況已經不是bank conflict了，以為2.x采用了broadcast機制，牛吧，哈哈。 這里要多看看矩陣乘積和矩陣轉置例子中的share memory的使用，如何保證memory coalescing和避免bank conflict的。
3. texture memory是有cache的，但是，如果同一個wrap內的thread的訪問地址很近的話，那么性能更高。
4.以下是要注意的：
(1)在2.x的CC上，L1 cache比texture cache具有更高的數據帶寬。所以，看着使用哈。
(2)對global memory的訪問，1.0和1.1的設備，容易造成memory uncoalescing，而1.2和1.3的設備，容易造成bandwidth waste。 而對2.x的設備而言，相比1.2和1.3，除了多了L1 cache，沒有其他的特別之處。
(3）采用-maxrregcount=N阻止complier分配過多的register。
(4)occupancy是每個multiprocessor中active wrap的數目與可能active wrap的最大數目的比值。higher occupancy並不意味着higher performance，因為畢竟有一個點，超過這個點，再高的occupancy也不再提高性能了。
5.影響occupancy的一個因素，就是register的使用量。比如，對於1.0和1.1的device來講，每個multiprocessor最多有8192個register，而最多的simultaneous thread個數為768個，那么對於一個multiprocessor，如果occupancy達到100%的話，每個thread最多可以分配10個register。另外，如果在1.0和1.1上，一個kernel里面的一個block有128個thread，每個thread使用register個數為12，那么，occupancy為83%，這是因為一個block有128個thread，則，由於multiprocessor里面最大的simultaneous thread為768，根據這個數目計算，最多同時有6個active block，但是6個active block，就會導致總共thread個數為128*6*12個，嚴重超過了8192，所以不能為6，得為5，因為128*5<768, and 128*5*12<8192, 5是滿足要求的最大的整數。如果一個kernel里面的一個block有256個thread，同樣一個thread用12個register，那么occupancy為66%，因為active block為2。可以在編譯選項里面加入--ptxas-options=-v查看kernel中每個thread使用register的數量。同時，NV提供了CUDA_Ocuppancy_calculator.xls作為occupancy計算的輔助工具。順便說一下，對於1.2和1.3的device來講，每個multiprocessor最多的simultaneous thread個數為1024個。
6. 為了隱藏由於register dependent寄存器依賴造成的訪問延遲latency，最小要保證25%的occupancy，也就是說，對於1.x的device來講，一個multiprocessor最少得發起192個thread。對於1.0和1.1來講， occupancy為192/768=25%，達到要求，但是對於1.2和1.3而言，192/1024=18.75%，不過，也只能這樣。對於2.x系列的device來講，由於是dual-issue，一個multiprocessor最多發起simultaneous thread個數為1536個，所以，一個multiprocessor最少同時發起384個thread時，occupancy為384/1536=25%，又達到了25%。
7. 對於block和thread的分配問題，有這么一個技巧，每個block里面的thread個數最好是32的倍數，因為，這樣可以讓計算效率更高，促進memory coalescing。其實，每個grid里面block的dimension維度和size數量，以及每個block里面的thread的dimension維度和size數量，都是很重要的。維度呢，采用合適的維度，可以更方便的將並行問題映射到CUDA架構上，但是，對性能不會有太大改進。所以，size才是最重要的，記住叻! 其實，訪問延遲latency和occupancy占有率，都依賴於每個multiprocessor中的active wrap的數量，而active wrap的數量，又依賴於register和share memory的使用情況。首先，grid中block的數目要大於multiprocessor的數目，以保證每個multiprocessor里面最少有一個block在執行，而且，最好有幾個active block，使得blocks不要等着__syncthreads()，而是占用了hardware。其次，block里面的thread的數目也很重要。對於1.0和1.1的設備來講，如果一個kernel里面block的大小為512個thread，那么，occupancy為512/768=66%，並且一個multiprocessor中只有一個active block，然而，如果block里面的thread為256個thread，那么，768/256=3，是整數，因此，occupancy為100%，一個multiprocessor里面有3個active block。但是，記住了，higher occupancy don't mean better performance更高的占有率並不意味着更好的性能。還是剛才那個例子，100%的occupancy並不比66%的occupancy的性能高很多，因為，更低的occupancy使得thread可以有更多的register可以使用，而不至於不夠用的register分配到local memory中，降低了變量存取訪問速度。一般來講啊，只要occupancy達到了50%，再通過提高occupancy來提高性能的可能性不是很大，不如去考慮如何register和share memory的使用。保證memory coalescing和防止bank conflict。記住如下幾點：
(1)block里面thread個數最好為wrap大小的倍數，即：32的倍數。使得計算效率更高，保證memory coalescing。
(2)如果multiprocessor中有多個active block時，每個block里面的thread個數最好為64的倍數。
(3)當選擇不同的block大小時，可以先確定block里面thread個數為128到256之間，然后再調整grid中block大小。
(4)如果是讓問延遲latency造成程序性能下降時，考慮在一個block里面采用小block划分，不要在一個multiprocessor中分配一個很大的block，盡量分配好幾個比較小的block，特別是程序中使用了__syncthreads()，這個函數是保證block里面所有wrap到這里集合，所以，block里面的thread越少越好，最好是一個wrap或者兩個wrap，這樣就可以減少__syncthreads()造成的訪問延遲。
(5)如果如果一個block里面分配的register超過了multiprocessor的最大極限時，kernel的launch就會fail。
8. share memory的使用量也是影響occupancy的一個重要因子。thread與share memory的元素之間，沒有必要是一對一的。一個線程可以一起負責處理share memory數組中的第一個、第二個以及第三個元素，都ok的。第一個thread處理share memory中的第一個元素，第二個thread負責處理第二個元素，類推如下，這種情況不是必須的，有時也沒必要這么做。在代碼里面，采用一個thread負責處理share memory數組中的多個元素的方法，是非常好的策略。這是因為如果share memory里面各個元素要進行相同的操作的話，比如乘以2，那么，這些操作可以被負責處理多個元素的一個thread一次搞定，分攤了thread處理share memory元素數量的成本費用。
9. 當上面那些high level級別的優化策略都檢查使用過以后，就可以考慮low level級別的優化：instruction optimization指令集優化。這個也可以很好的提高性能的。指令集的優化，可以稍微總結如下：
(1)盡量使用shift operation位移運算來取代expensive昂貴的division除法和modulo取余運算，這里說的都是integer運算，float不行的。如果n是2冪數，(i/n)=(i>>log2(n)), (i%n)=(i&(n-1)). 其實，這只是一個量的問題，對於1.x的設備而言，如果一個kernel里面使用了十多個tens of這樣的指令，就要考慮用位移運算來取代了;對於2.x的設備而言，如果一個kernel里面使用了20個這樣的指令，也要考慮使用位移運算來取代除法和取余運算。其實，compiler有時會自動做這些轉換的，如果n是2的冪數。
(2)reciprocal square root，對於平方根倒數1.0f/sqrtf(x)，編譯器會采用rsqrtf(x)來取代，因為硬件做了很多優化。當然，對於double型的平方根倒數，就采用rsqrt(x)啦。呵呵，記住了。
(3)編譯器有時會做一些指令轉化。在要做計算的單精度浮點型常數后面，一定要加入f，否則，會被當做雙精度浮點運算來計算，對於2.x以上的設備來講，這一點很重要，記好了。
(4)如果追求速度speed，而不是精度precision，那么盡量使用fast math library。比如，__sinf(x)、__expf(x)比sinf(x)和expf(x)有更快的速度，但是，精度卻差一些。如果是__sinf(x-2)則比sinf(x-2)的速度要快一個數量級，因為x-2運算用到了local memory，造成太高的訪問延遲。當然，在compiler option中使用-use_fast_math可以讓compiler強制將sinf(x)和expf(x)轉化為__sinf(x)和__expf(x)進行計算。對於transcendental function超越函數，作用對象是單精度浮點型數據時，經常這么用，其他類型的數據，性能提升不大。
(5)對於2和10為底做指數運算，一定要采用exp2()或者expf2()以及exp10()或者expf10()，不要采用pow()和powf()，因為后者會消耗更多的register和instruction指令。 另外，exp2()、expf2()、exp10()、expf10()的性能和exp()以及expf()性能差不太多，當然比pow()和powf()要快10多倍呢。加好了哈。
(6)減少global memory的使用，盡量將global memory數據加載到share memory，再做訪問。因為訪問uncached的顯存數據，需要400~600個clock cycle的內存延遲。
10. 下一個就是control flow了。一定要避免在同一個wrap里面發生different execution path。盡量減少if、swith、do、for、while等造成同一個wrap里面的thread產生diverge。因為，一旦有divergence，不同的execution path將會順序的串行的執行一遍，嚴重影響了並行性。但是：
{
case 0：
break;
case 1:
break;
...
case 31:
break;
}
上面這個例子，則不會發生divergence，因為控制條件剛好和wrap里面的thread相對應。
其實，有時，compiler會采用branch predication分支預測來打開loop循環或者優化if和switch語句， 這時，wrap就不會出現divergence了。在寫code時，我們也可以自己采用#pragma uroll來打開loop循環。在使用branch predication時，所有指令都將會執行，其實，只有預測正確的真正的執行了，而預測錯誤的，其實就是thread，不會去讀取該instruction的地址和數據，也根本不會寫結果。其實，編譯器做分制預測，是有條件的，只有分支條件下的指令instruction的個數小於等於某個閾值的時候，才會做分支預測branch predication。如果編譯器覺得可能會產生多個divergent wrap，那么閾值為7，否則為4。(這里很不理解7和4是怎么來的)。
11. 在loop循環的counter，盡量用signed integer，不要用unsigned integer。比如：for(i = 0; i < n; i++) {out[i] = in[offset+stride*i];} 這里呢，stride*i可以會超過32位integer的范圍，如果i被聲明為unsigned，那么stride*i這個溢出語句就會阻止編譯器做一些優化，比如strength reduction。相反，如果聲明為signed，也沒有溢出語句時，編譯器會對很多地方做優化。所以，loop counter盡量設置為int，而不是unsigned int。
12. 在1.3及其以上的device上，才支持double-precision floating-point values，即：64位雙精度浮點運算。當使用double時，在編譯器選項里面添加：-arch=sm_13
13. 還有一點需要注意，如果A、B、C都是float，那么A+(B+C)並不一定等於(A+B)+C。
14. 先看下面兩個語句：float a; a = a * 1.02;
對於1.2及其以下的device來講，或者1.3及其以上device，但是沒有打開支持double運算的選項，那么，由於不支持double，所以，1.02*a這個乘積是一個float;
對於1.3及其以上的device來講，如果打開了支持double運算的選項，那么，a*1.02是一個double，而將乘積賦值給a，這個結果是float，所以，是先做了從float到double的promotion擴展，然后做了從double到float的truncation截取。
15. 多GPU編程。如果有p個GPU同時並行，那么，程序中就需要p個CPU threads。這些thread可以用OpenMP(小規模)或者MPI(大規模)進行管理。GPU之間的數據拷貝，必須通過CPU實現。對於OpenMP，是這樣的：一個CPU thread將數據從對應的GPU中拷貝到host端的share memory region中，然后另一個CPU thread將數據從host端的share memory region拷貝到對應的GPU中。也就是說：OpenMP是通過share memory進行數據拷貝的。而對於MPI而言，數據是通過message passing進行傳遞的。一個CPU thread使用cudaMemcpy將數據從device拷貝到host，然后通過MPI_Sendrecv()，另一個CPU thread就使用cudaMemcpy將數據從host端拷貝到呃device端。編譯選項，記着采用nvcc -Xcompiler /openmp或者nvcc -Xcompiler mpicc。