CUDA编程
近期開始學習CUDA編程，需要閱讀很多資料，為了便于整理復習，特將閱讀筆記記錄，以備后用。
這一系列文章是根據(jù)NVIDIA公司官方文檔《CUDA C Best Practices》的內(nèi)容來進行整理的，由于筆者剛開始進行CUDA的學習，而并行語言的學習不如串行語言如C、C++那樣容易入門，因此理解錯誤之處在所難免，歡迎讀到錯誤的各位批評指正。
1. 學習目的
? ? ? CUDA是一個C語言的擴充，學習它的目的是利用CUDA將程序中的可并行部分交由GPU來完成，以達到CPU與GPU協(xié)同工作的效果，極大提升程序性能。
2. 優(yōu)化過程
? ? ? 學習CUDA的一個非常有用的作用在于程序員可以對現(xiàn)有的C/C++程序進行改寫，將其中適合在GPU上運行的并行部分代碼挖掘出來，改寫代碼使得它們可以在GPU上運行，從而極大地提升程序的運算性能。將串行代碼轉(zhuǎn)化為并行代碼的過程是迭代的，簡單來說，每一次迭代可以劃分為4個相對獨立的過程：
? ? ? 1、分析（Assess）
? ? ? 2、并行化（Parallelize）
? ? ? 3、優(yōu)化（Optimize）
? ? ? 4、部署（Deploy）
? ? ? 這四個過程合起來稱作APOD，CUDA為每個過程提供了解決方案以改善程序性能。APOD是一個循環(huán)的過程：對程序進行初步分析并行優(yōu)化后，可以繼續(xù)運用上述手段，分析程序中的可并行化段，利用CUDA對該段代碼進行改寫，優(yōu)化改寫的程序，最后將改寫后的代碼部署到原有程序中。這一過程可以形象地用下圖來進行表示：
? ? ? 并行化過程是使用CUDA對原有C/C++程序進行初步改寫，CUDA提供了眾多的并行庫供程序員調(diào)用，例如cuBLAS、cuFFT以及Thrust等，除此之外，還可以運用一些預處理指令來優(yōu)化編譯器的行為。
? ? ? 優(yōu)化是對并行化后的程序重復進行APOD的過程，以使程序達到更好的性能。
? ? ? 在部署階段，程序員需要比較改寫后與改寫前程序運算的結(jié)果，以驗證改寫的正確性。驗證完畢后，將改寫程序加入到原有項目當中。
? ? ? 分析過程主要是找出程序性能提升的極限，找出它們的方法是運用所謂的Amdahl定理和Gustafson定理，這兩個定理隨后將會提到。
3. 異構(gòu)計算
? ? ? ??異構(gòu)計算是指運用多種不同架構(gòu)的處理器來完成計算任務，使用CUDA，我們可以協(xié)調(diào)CPU和GPU，讓它們分工合作以達到計算的目的。
? ? ? 在CUDA中，CPU用主機（Host）來表示，GPU則用設備（Device）來表示。主機和設備之間是有一些區(qū)別的，這些區(qū)別的主要部分集中在線程模型和物理內(nèi)存方面：
<1> 線程資源（Threading resources）
? ? ??CPU所支持的同時運行的線程數(shù)是極其有限的，一個擁有4個6核心處理器的服務器處理器在同一時刻只能同時運行24個線程（注意是同一時刻！），而現(xiàn)代NVIDIA GPU則可以支持同一時刻數(shù)千個活動線程同時運行。
<2> 線程（Threads）
? ? ??在CPU中線程切換由于涉及到上下文（Context）的改變，代價很大。與之相比，GPU中的線程則非常輕巧，切換幾乎沒有代價。簡言之，CPU的設計初衷是為了最小化線程切換的延遲，而GPU的設計理念是為了處理大量同時運行的輕量級線程以最大化吞吐量。
<3> 內(nèi)存（RAM）
? ? ??主機和設備都擁有各自的物理內(nèi)存，它們之間通過PCI-E總線來交換信息。為了使用CUDA，數(shù)據(jù)必須通過PCI-E總線從主機傳輸?shù)皆O備上。傳輸?shù)拈_銷是非常可觀的，因此，為了獲得更好的性能表現(xiàn)，數(shù)據(jù)重用是非常重要的。簡單來說，數(shù)據(jù)應當盡可能久地保存在設備上以備運算所用。
4. 性能分析
? ? ??在很多項目中，完成了絕大多數(shù)工作任務的是相對較少的一部分核心代碼。使用性能分析器，開發(fā)人員可以識別這樣的熱點代碼，找到瓶頸，進而有針對性地對代碼進行優(yōu)化。性能分析的工具非常多，典型的工具如gprof便是其中之一，它是一款Linux平臺上的開源性能分析器。下面則是其分析結(jié)果的部分截圖：
? ? ? 可以看到函數(shù)genTimeStep的運行時間幾乎占總時間的三分之一，是程序的瓶頸所在。 ? ? ? 當然，看到瓶頸并進行針對性優(yōu)化的同時，我們也要注意定量分析程序可達的最佳性能，這方面有兩個問題值得注意：
4.1 強標度與Amdahl定律（Strong Scaling and Amdahl's Law）
? ? ??強標度是在問題總規(guī)模不變的情況下，衡量時間如何隨處理器數(shù)量的增減而變化的方法。根據(jù)Amdahl定律，對于一個規(guī)模固定的的問題，通過增加處理器數(shù)量所能達到的加速比S可用下面的式子來表示：其中P是可并行化部分代碼執(zhí)行時間與總執(zhí)行時間的比率，N代表執(zhí)行可并行化代碼的處理器總數(shù)。可以看到，N越大，則S越大，當N趨向于無窮大時，S也達到了其最大值S = 1 / (1 - P)。
4.2 弱標度與Gustafson定律（Weak Scaling and Gustafson's Law）
? ? ??弱標度是在每一處理器可處理問題規(guī)模恒定的情況下，衡量時間如何隨處理器數(shù)量的增減而變化的方法。弱標度常使用Gustafson定律來計算，根據(jù)弱標度的定義，系統(tǒng)所處理的問題規(guī)模會隨著處理器的增加而增加，因此，加速比S可以用下式來表示：這里P代表串行執(zhí)行時問題中可并行化部分運行時間與總運行時間的比率，N代表處理器的數(shù)量。 對于Gustafson定律，我們可以這樣理解：在弱標度的情形下，問題規(guī)模并不是恒定的，總運行時間才是恒定的。設串行運行的總時間為T，每個處理器處理耗時為t，那么單處理器串行執(zhí)行時需要T/t個處理器時間，將處理器數(shù)量增加到N，這時部分代碼由于并行化的緣故是可以重疊執(zhí)行的，在t時間中，相當于完成了串行執(zhí)行時PNt + (1 - P)t時間的任務，此時執(zhí)行總時間為T / (PNt + (1 - P)t)，將串行執(zhí)行時間與N處理器并行執(zhí)行時間相比，得到加速比為：S = PN + 1 - P，也就是Gustafson定律括號展開后得到的結(jié)果。 ? ? ? 綜上，Amdahl定律和Gustafson定律可以幫助開發(fā)人員找到問題優(yōu)化的極限，理解它們是十分重要的。
5. 得到正解
? ? ??得到正確的計算結(jié)果是我們的最初目的，但是使用CUDA這種并行編程模型是很容易出錯的，這時我們就需要一些方法和工具來幫助我們驗證計算結(jié)果的正確性，同時，在CUDA編程中也有一些值得我們注意的問題。
5.1 正確性驗證
? ? ? 正確性驗證主要有兩種方法：
? ? ??1、引用比較（Reference Comparison）
? ? ? 引用比較的核心思想是使用未并行化代碼產(chǎn)生的一些具有代表性的結(jié)果與并行化后的程序運行結(jié)果，當它們的絕對差在可接受范圍內(nèi)時，就認為并行化的結(jié)果是正確的。注意，改寫前和改寫后代碼運行結(jié)果的不一致是由浮點數(shù)表示的不確切性造成的。
? ? ? 第一步完成后，我們使用APOD（Assess、Parallelize、Optimize、Deployment）過程對并行化代碼實施進一步優(yōu)化，我們只要保證每一步優(yōu)化的引用比較結(jié)果正確，那么對于最終的并行化程序，其結(jié)果的正確性是可以得到保證的。
? ? ??2、單元測試（Unit Testing）
? ? ? 單元測試與引用比較的方法是相輔相成的。單元測試是指開發(fā)人員在編寫代碼時就將項目代碼組織成單元級別，然后運用一定的技術(shù)手段對各個單元分別測試其正確性。在CUDA中，我們可以把內(nèi)核（kernels）寫成一系列小的__device__函數(shù)的組合而不是將代碼封裝到一個龐大的__global__函數(shù)中。這樣，我們就可以在連接各部分代碼之前對各個設備（device）代碼進行單獨測試。
5.2 調(diào)試（Debugging）
? ? ? 調(diào)試CUDA需要一些特別的工具：
? ? ? 1. CUDA-GDB?
? ? ? CUDA-GDB是Linux和Mac環(huán)境中GNU調(diào)試器的一個端口（Port），具體信息參見：GNU-GDB
? ? ? 2. NVIDIA Parallel Nsight
? ? ? NVIDIA Parallel Nsight調(diào)試和性能分析器是可以在Windows環(huán)境下作為Microsoft Visual Studio的插件使用，具體參見：NVIDIA Parallel Nsight
? ? ? 3. 一些第三方調(diào)試工具
? ? ? 一些第三方工具也支持CUDA的調(diào)試，具體參見：Debugging solutions
5.3 數(shù)值精度問題
? ? ? 由于CUDA使用的是浮點數(shù)進行運算，涉及到一些精確度問題，需要我們在編程時注意一下。
1、單精度與雙精度問題
? ? ? 單精度和雙精度浮點運算的結(jié)果差別是很大的，在CUDA中，只有運算能力大于或等于1.3的硬件才能本地支持雙精度運算。執(zhí)行運算時，程序員務必搞清進行的是哪一種運算以獲得正確的結(jié)果。nvcc編譯命令行中使用-arch=sm_13可開啟雙精度運算。 ? ? ?
2、浮點運算不遵從結(jié)合律
? ? ? 對于三個浮點數(shù)A、B以及C，需要注意的是，(A+B)+C并不等于A+(B+C)。
3、雙精度擴展和單精度截斷
? ? ? 由于單精度浮點數(shù)和雙精度浮點數(shù)運算結(jié)果的不一致性，CUDA程序員應注意避免一些精度細節(jié)問題，如下面這段代碼：
[cpp] view plaincopy在C語言中，1.02會被隱含解釋為double類型，那么第二個式子右邊的a將會被擴展為雙精度浮點數(shù)而執(zhí)行double乘法，得到的結(jié)果為雙精度浮點數(shù)，最后再將這一結(jié)果截斷為單精度浮點數(shù)賦值給式子左邊的a。這會帶來隱患，解決這一問題的方法是用1.02f來表示單精度浮點常數(shù)。
6. 性能度量（Performance Metrics）
? ? ? 為了優(yōu)化CUDA程序的性能，我們需要一種定量的方法來對程序性能進行衡量，同時，我們也需要明晰帶寬（bandwidth）在性能衡量中所扮演的角色。下面我們將依依介紹這些概念。
6.1 定時
CUDA調(diào)用和內(nèi)核執(zhí)行可以使用CPU和GPU定時器來計時，使用這些計時器需要注意相關(guān)的問題：
1、使用CPU定時器
使用CPU定時器需要注意很多CUDA API函數(shù)是異步的，即這些函數(shù)在任務完成之前就會把控制權(quán)返回給CPU。它們一般都會在函數(shù)名之前帶有Async，遇到這類函數(shù)時需要特別小心。為了達到精準定時的效果，我們需要同步CPU和GPU，這可以通過在開始和停止定時器前調(diào)用函數(shù)cudaDeviceSynchronize()來實現(xiàn)。cudaDeviceSynchronize()阻塞當前CPU線程，直到指定流（Stream）中所有之前的CUDA調(diào)用運行完畢之后CPU線程才開始繼續(xù)執(zhí)行。類似地，CPU線程也可以與GPU流或者GPU事件進行同步，使用cudaStreamSynchronize()函數(shù)或者cudaEventSynchronize()函數(shù)可以達到這樣的效果。然而需要注意的是，上述兩個函數(shù)不適合在默認流以外的其他流中使用，因為驅(qū)動程序可能安排不同流中的代碼交錯運行，這樣就會造成計時錯誤。而默認流（Stream 0）在設備上表現(xiàn)為串行運行，因此可以使用上面兩個函數(shù)進行計時。
2、使用CUDA GPU定時器
CUDA event API提供了用于創(chuàng)建和消除以及記錄事件（使用timestamp）的調(diào)用函數(shù)，通過timestamp之間的差值換算，我們可以得到毫秒級的GPU運行時間信息。下面的代碼片段展示了這一技術(shù)：
How to time code using CUDA events
[cpp] view plaincopy
6.2 帶寬
目前對此理解尚淺，留個占位符以待今后補充。總結(jié)
- 上一篇： 监督学习 | ID3 决策树原理及Pyt
- 下一篇： 经纬度坐标与距离的相互转换及其实现