CUDA學習筆記之程序優化

標簽： cuda優化conflict存儲算法數學計算 2010-01-05 17:18 5035人閱讀 評論(4) 收藏 舉報 分類： CUDA（6）

版權聲明：本文為博主原創文章，未經博主允許不得轉載。

CUDA程序優化

CUDA程序優化應該考慮的點：

精度：只在關鍵步驟使用雙精度，其他部分仍然使用單精度浮點以獲得指令吞吐量和精度的平衡；

?????????? 延遲：需要首先緩沖一部分數據，緩沖的大小應該可以保證每個內核程序處理的一批數據能夠讓GPU慢負荷工作；

?????????? 計算量：計算量太小的程序使用CUDA很不合算；當需要計算的問題的計算密集度很低的時候，執行計算的時間遠遠比IO花費的時間短，整個程序的瓶頸出現在PCI-E帶寬上。

優秀的CUDA程序特征：

在給定的數據規模下，選用算法的計算復雜度不明顯高于最優的算法；

Active warp的數量能夠讓SM滿載，并且active block的數量大于2，能夠有效地隱藏訪存延遲；

????????????? ? 當瓶頸出現在運算指令時，指令流的效率已經過了充分優化；

???????? 當瓶頸出現在訪問IO時，程序已經選用了恰當的存儲器來儲存數據，并使用了適當的存儲器訪問方式，以獲得最大帶寬；

CUDA的編寫與優化需要解決的問題：

???????? 確定任務中的串行和并行的部分，選擇合適的算法；

????????????? ?按照算法確定數據和任務的劃分方式，將每個需要實現的步驟映射為一個滿足CUDA兩層并行模型的內核函數，讓每個SM上至少有6個活動warp和至少2個活動block；

??????? 在精度不組或者發生一處時必須使用雙精度浮點或者更長的整數類型；

??????? 優化顯存訪問：合并采用相同block和grid的kernel；盡力避免將線程私有變量分配到local memory；

??????? 優化指令流：在誤差可接受的情況下，使用CUDA算術指令集中的快速指令；避免多余的同步；在只需要少量線程進行操作的情況下，使用類似“if threaded<N”的方式，避免多個線程同時運行占用更長時間或者產生錯誤結果；

??????? 資源均衡：調整每個線程處理的數據量，shared memory和register和使用量；通過調整block大小，修改算法和指令以及動態分配shared memory，都可以提高shred的使用效率；register的多少是由內核程序中使用寄存器最多的時刻的用量決定的，因此減小register的使用相對困難；

??????? 節約register方法：使用shared memory存儲變量；使用括號明確地表示每個變量的生存周期；使用占用寄存器較小的等效指令代替原有指令；

???????? 與主機通信優化：盡量減少CPU與GPU間的傳輸：使用cudaMallocHost分配主機端存儲器，可以獲得更大帶寬；一次緩存較多的數據后再一次傳輸，可以獲得較高的貸款；需要將結果顯示到屏幕的時候，直接使用與圖形學API互操作的功能；使用流和異步處理隱藏與主機的通信時間；使用zero-memory技術和Write-Combined memory提高可用帶寬；

測量程序運行時間：

??????? CUDA內核程序的運行時間：可以在設備端測量，也可以在主機端測量；

??????? CUDA API的運行時間：只能在主機端測量；使用CUDA runtime API時，會在第一次調用runtime API函數時啟動CUDA環境，計時的時候應該避免將這一部分計入，因此在正式測試之前應當首先及你選哪個一側包含數據輸入輸出地就愛上你，使得GPU從平時的節能模式進入工作狀態，使得測試結果更加可靠；

設備端測量時間：

?? ???調用clock（）函數：返回的是GPU的時鐘周期，需要除以GPU的運行頻率才能得到以秒為單位的時間；

使用CUDA API事件管理功能；

主機端測量時間：使用c標準庫中的clock_t()函數測試，由于其精度很低，因此應該運行多次然后求平均運行時間；注意異步函數（比如內核函數和帶有asyn后綴的存儲器拷貝函數），在GPU上執行完成之前，CPU線程已經得到了它的返回值；從主機測量一系列CUDA調用需要的時間的時候，要首先調用cudaThreadSynchronize()函數等，使得GPU線程執行完畢后，進入CPU線程，從而得到正確的執行效果；在一串流中的第一個流（ID為0的流）的行為總是同步的，因此使用這些函數對0號流進行測時，得到的記過是可靠的。

任務劃分原則：

????? ??在兩次主機—設備通信之間進行盡量多的計算；考慮使用流運算隱藏主機—設備通信時間，通過Pinned memory、zero—copy、write—combined memory等手段提高實際傳輸帶寬；

?????? ?盡量使得每個block中線程數量是32的整數倍，最好保持在64~256之間，并根據任務的具體情況確定每個維度上的大小，以減少計算訪存地址時的整數除法和求模運算；

????????????? 對一個block的任務進行劃分后，再按照block的維度和尺寸要求對grid進行劃分：每個block 的訪存均勻分布在顯存的各個分區中；block 間的負載可以存在一定程度的不均衡；

Grid和Block的維度設計：

??? ????首先考慮block的尺寸，grid的尺寸一般越大越好；

??????? 每個SM中至少要有6個active warp用于隱藏流水線延遲，并且擁有至少2個active block；

計算每個SM上active warp和active block的數量：

確定每個SM使用的資源數量：使用nvcc的—keep編譯選項，或者在.cu編譯規則（cuda build rule）中選擇保留中間文件，得到.cubin文件，用寫字板打開后可以看到imem 和reg分別代表內核函數中每個線程使用的local memory和register數量；

根據硬件確定SM上的可用資源：可以用SDK中的deviceQuery 獲得每個SM中的資源；根據內核不同，SM上的warp總數上限，block總數上限，寄存器數量，shared memory數量都不同；

每個block中的線程數量不能超過512；

計算每個block使用的資源，并確定active block和active warp數量：

???? e.g. 每個block中有64個線程，每個block使用256 Byte shared memory，8個寄存器文件，

那么：每個人block使用的shared memory： 256 Byte；

????? 每個block使用的寄存器文件數量： 8*64 = 512；

????? 每個block中使用的warp數量：64/32 = 2;

???????????????? 如果在G80/92 GPU中運行這個內核程序：

??????????????????????? 由shared memory數量限制的active block數量： 16384、256? = 64；

????????????????????????????????????????? ?由寄存器數量限制active block數量：8192/512 = 16;

????? ??????????????????由warp數量限制的active block數量 24/2 = 12;

??????????????????????? 每個SM中的最大active block數量：8；

這些計算可以由NVIDIA在CUDA SDK中提供的 CUDA occupancy calculator完成；

?? Block 的維度和每個維度上的尺寸的主要作用是避免做整數除法和求模運算，對執行單元效率沒有什么顯著影響；

計算grid中各個維度上block的數量：grid在x軸上的block數量 = （問題在x軸上的尺寸+每個block在x軸上的尺寸-1）/每個block在x軸上的尺寸；

?存儲器訪問優化：????

?主機—設備通信優化：

????? 目前一條PCI—E 2.0*16通道的理論帶寬是每向8GB/s, 遠小于顯存和GPU片內存儲器帶寬；

???? ?Pinned memory：強制讓操作系統在物理內存中完成內存申請和釋放工作，不用參會頁交換，因此速度比pageable memory快；

???????????????????????? 聲明這些內存會占用操作系統的可用內存，可能會影響到操作系統運行需要的物理內存；

???????????????????????? 需要合理規劃CPU和GPU各自使用的內存，使整個系統達到最優；

異步執行：

????? 內核啟動和顯存內的數據拷貝（Device to Device）總是異步的；

????? 內存和顯存間的數據拷貝函數有異步和同步兩個版本：

????????????? 同步(順序執行)： cudaMemcpy(a_d,a_h,size,cudaMemcpyHostToDevice);

???????????????????? ??????????cpuFunction();

????????????? 異步(同時執行)： cudaMemcpyAsync（…………）；

????? ???????????????????????????????? ????????cpuFunction()；

????? 屬于同一個流中的內核啟動總是同步的；

????? 如果幾次內核啟動屬于不同的流，那么他們的執行可能是亂序的；

利用異步提高計算效率：

????? 使用流和異步是CPU和GPU同事進行運算；

????? 利用不同流之間的異步執行，使流之間的傳輸和運算能夠同時執行，更好地利用GPU資源；

全局存儲器訪問優化：

???? ??需要考慮half-warp訪問的對齊問題，不同的硬件要求不同；（存疑？？？？？？？？）

?????? 采用合并訪問；

?????? 盡量避免間隔訪問：比如按列訪問矩陣，可以借助shared memory來實現這一點；

Shared memory訪問優化：

?????? 共享存儲器被組織為16個可以被同時訪問的存貯器模塊，稱為bank；

Bank組織方式：寬度32bit，相鄰的32bit字被組織在相鄰的bank中，每個bank在每個時鐘周期可以提供32bit的帶寬；

一個warp被分為兩個half-warp進行訪問；

避免bank conflict：在SDK中，使用寬度為17或則會threadDim.x+1的行來避免bank conflict；（存疑？？？？？？？？）

Shared memory采用了廣播機制：在相應一個對同一個地址的讀請求時，一個32bit字可以被讀取并同時廣播給不同的線程；

當一個half-warp中有多個線程讀取同一個32bit字地址中的數據時，可以減少bank conflict的數量；

如果half-warp中的線程全都讀取同一地址中的數據時，此時完全不會發生bank conflict；

如果half-warp內有多個線程要對同一地址進行讀寫操作，此時會產生不確定結果，這種情況應該使用shared memory的原子操作；

共享存儲器保存著加載kernel時傳遞過來的參數，以及kernel執行配置參數，如果參數列表很長，應該將其中一部分參數放入constant memory；

使用紋理存儲器：

?? ?????主要用于存放圖像和查找表：不用嚴格遵守合并訪問條件，就能達到較高帶寬；

????????????????????????????????? 對于少量數據的隨機訪問，效率不會太差；

????????????????????????????????? 可以使用線性濾波和自動類型轉換等功能調用硬件的不可編程計算資源，不必占用可編程計算單元；

使用常數存儲器：

??????? 主要用于存放指令中的常數；速度低于shared memory；

指令流優化：

增大吞吐量手段：

避免使用地吞吐量指令；

?????? 優化每種類型的存儲器，有效利用帶寬；

?????? 允許線程調度單元精良用多的數學計算來覆蓋訪存延遲，需要有教導的算術密度；

吞吐量：每個多處理器在一個時鐘周期下執行的操作數目；

算術指令：盡量使用單精度浮點單元進行運算，在計算能力小于等于1.2的設備中，每個雙精度的變量將會轉換成單精度格式，雙精度運算也會轉為單精度算術運算；

????????? 單精度浮點基本算術運算：加，乘，乘加運算的吞吐量是每個時鐘周期8個操作；

????????? 求導數運算：每個時鐘周期2個操作；

????????? 單精度除法：每個時鐘周期0.88個操作；

????????? 單精度浮點倒數平方根：2；

???????? ?平方根：1；

????????? 對數：2；

????????? 正弦余弦：參數較大的時候，采用歸約操作將x的絕對值減小；有快路徑和慢路徑（大參數）；

????????? 整數算術運算：整數加法（8），乘（2）；除法和取模開銷特別大，盡量地避免或者用位運算代替；

????????? 比較，min，max：（8）；

????????? 位運算（8）；

????????? 類型轉換（8）；

控制流指令： If, switch, do, for, while 可能引起一個warp線程跳轉到不同的分支，嚴重影響指令吞吐量；

訪存指令：包括任何讀寫memory的指令；

對于local memory只有在register不夠用或者編譯器無法解析的時候才會發生；

將較大的數據（float，double）拆分成每個線程32bit，或者將多個[u]char,[u]short合并成每個線程32bit的形式訪問；

在訪問local/global memory時候，會有額外的400~600個時鐘周期的訪問延遲；

同步指令： _syncthreads() 的吞吐量是每時鐘周期 8 個操作；

總結

以上是生活随笔為你收集整理的CUDA学习笔记之程序优化的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。