近期開始學習CUDA編程，需要閱讀很多資料，為了便于整理復習，特將閱讀筆記記錄，以備后用。

這一系列文章是根據NVIDIA公司官方文檔《CUDA C Best Practices》的內容來進行整理的，由于筆者剛開始進行CUDA的學習，而并行語言的學習不如串行語言如C、C++那樣容易入門，因此理解錯誤之處在所難免，歡迎讀到錯誤的各位批評指正。

1. 學習目的

? ? ? CUDA是一個C語言的擴充，學習它的目的是利用CUDA將程序中的可并行部分交由GPU來完成，以達到CPU與GPU協同工作的效果，極大提升程序性能。

2. 優化過程

? ? ? 學習CUDA的一個非常有用的作用在于程序員可以對現有的C/C++程序進行改寫，將其中適合在GPU上運行的并行部分代碼挖掘出來，改寫代碼使得它們可以在GPU上運行，從而極大地提升程序的運算性能。將串行代碼轉化為并行代碼的過程是迭代的，簡單來說，每一次迭代可以劃分為4個相對獨立的過程：

? ? ? 1、分析（Assess）

? ? ? 2、并行化（Parallelize）

? ? ? 3、優化（Optimize）

? ? ? 4、部署（Deploy）

? ? ? 這四個過程合起來稱作APOD，CUDA為每個過程提供了解決方案以改善程序性能。APOD是一個循環的過程：對程序進行初步分析并行優化后，可以繼續運用上述手段，分析程序中的可并行化段，利用CUDA對該段代碼進行改寫，優化改寫的程序，最后將改寫后的代碼部署到原有程序中。這一過程可以形象地用下圖來進行表示：

? ? ? 并行化過程是使用CUDA對原有C/C++程序進行初步改寫，CUDA提供了眾多的并行庫供程序員調用，例如cuBLAS、cuFFT以及Thrust等，除此之外，還可以運用一些預處理指令來優化編譯器的行為。
? ? ? 優化是對并行化后的程序重復進行APOD的過程，以使程序達到更好的性能。
? ? ? 在部署階段，程序員需要比較改寫后與改寫前程序運算的結果，以驗證改寫的正確性。驗證完畢后，將改寫程序加入到原有項目當中。

? ? ? 分析過程主要是找出程序性能提升的極限，找出它們的方法是運用所謂的Amdahl定理和Gustafson定理，這兩個定理隨后將會提到。

3. 異構計算

? ? ? ??異構計算是指運用多種不同架構的處理器來完成計算任務，使用CUDA，我們可以協調CPU和GPU，讓它們分工合作以達到計算的目的。

? ? ? 在CUDA中，CPU用主機（Host）來表示，GPU則用設備（Device）來表示。主機和設備之間是有一些區別的，這些區別的主要部分集中在線程模型和物理內存方面：

<1> 線程資源（Threading resources）

? ? ??CPU所支持的同時運行的線程數是極其有限的，一個擁有4個6核心處理器的服務器處理器在同一時刻只能同時運行24個線程（注意是同一時刻！），而現代NVIDIA GPU則可以支持同一時刻數千個活動線程同時運行。

<2> 線程（Threads）

? ? ??在CPU中線程切換由于涉及到上下文（Context）的改變，代價很大。與之相比，GPU中的線程則非常輕巧，切換幾乎沒有代價。簡言之，CPU的設計初衷是為了最小化線程切換的延遲，而GPU的設計理念是為了處理大量同時運行的輕量級線程以最大化吞吐量。

<3> 內存（RAM）

? ? ??主機和設備都擁有各自的物理內存，它們之間通過PCI-E總線來交換信息。為了使用CUDA，數據必須通過PCI-E總線從主機傳輸到設備上。傳輸的開銷是非?？捎^的，因此，為了獲得更好的性能表現，數據重用是非常重要的。簡單來說，數據應當盡可能久地保存在設備上以備運算所用。

4. 性能分析

? ? ??在很多項目中，完成了絕大多數工作任務的是相對較少的一部分核心代碼。使用性能分析器，開發人員可以識別這樣的熱點代碼，找到瓶頸，進而有針對性地對代碼進行優化。性能分析的工具非常多，典型的工具如gprof便是其中之一，它是一款Linux平臺上的開源性能分析器。下面則是其分析結果的部分截圖：

? ? ? 可以看到函數genTimeStep的運行時間幾乎占總時間的三分之一，是程序的瓶頸所在。 ? ? ? 當然，看到瓶頸并進行針對性優化的同時，我們也要注意定量分析程序可達的最佳性能，這方面有兩個問題值得注意：

4.1 強標度與Amdahl定律（Strong Scaling and Amdahl's Law）

? ? ??強標度是在問題總規模不變的情況下，衡量時間如何隨處理器數量的增減而變化的方法。根據Amdahl定律，對于一個規模固定的的問題，通過增加處理器數量所能達到的加速比S可用下面的式子來表示：
其中P是可并行化部分代碼執行時間與總執行時間的比率，N代表執行可并行化代碼的處理器總數?？梢钥吹?#xff0c;N越大，則S越大，當N趨向于無窮大時，S也達到了其最大值S = 1 / (1 - P)。

4.2 弱標度與Gustafson定律（Weak Scaling and Gustafson's Law）

? ? ??弱標度是在每一處理器可處理問題規模恒定的情況下，衡量時間如何隨處理器數量的增減而變化的方法。弱標度常使用Gustafson定律來計算，根據弱標度的定義，系統所處理的問題規模會隨著處理器的增加而增加，因此，加速比S可以用下式來表示：
這里P代表串行執行時問題中可并行化部分運行時間與總運行時間的比率，N代表處理器的數量。 對于Gustafson定律，我們可以這樣理解：在弱標度的情形下，問題規模并不是恒定的，總運行時間才是恒定的。設串行運行的總時間為T，每個處理器處理耗時為t，那么單處理器串行執行時需要T/t個處理器時間，將處理器數量增加到N，這時部分代碼由于并行化的緣故是可以重疊執行的，在t時間中，相當于完成了串行執行時PNt + (1 - P)t時間的任務，此時執行總時間為T / (PNt + (1 - P)t)，將串行執行時間與N處理器并行執行時間相比，得到加速比為：S = PN + 1 - P，也就是Gustafson定律括號展開后得到的結果。 ? ? ? 綜上，Amdahl定律和Gustafson定律可以幫助開發人員找到問題優化的極限，理解它們是十分重要的。

5. 得到正解

? ? ??得到正確的計算結果是我們的最初目的，但是使用CUDA這種并行編程模型是很容易出錯的，這時我們就需要一些方法和工具來幫助我們驗證計算結果的正確性，同時，在CUDA編程中也有一些值得我們注意的問題。

5.1 正確性驗證

? ? ? 正確性驗證主要有兩種方法：

? ? ??1、引用比較（Reference Comparison）

? ? ? 引用比較的核心思想是使用未并行化代碼產生的一些具有代表性的結果與并行化后的程序運行結果，當它們的絕對差在可接受范圍內時，就認為并行化的結果是正確的。注意，改寫前和改寫后代碼運行結果的不一致是由浮點數表示的不確切性造成的。

? ? ? 第一步完成后，我們使用APOD（Assess、Parallelize、Optimize、Deployment）過程對并行化代碼實施進一步優化，我們只要保證每一步優化的引用比較結果正確，那么對于最終的并行化程序，其結果的正確性是可以得到保證的。

? ? ??2、單元測試（Unit Testing）

? ? ? 單元測試與引用比較的方法是相輔相成的。單元測試是指開發人員在編寫代碼時就將項目代碼組織成單元級別，然后運用一定的技術手段對各個單元分別測試其正確性。在CUDA中，我們可以把內核（kernels）寫成一系列小的__device__函數的組合而不是將代碼封裝到一個龐大的__global__函數中。這樣，我們就可以在連接各部分代碼之前對各個設備（device）代碼進行單獨測試。

5.2 調試（Debugging）

? ? ? 調試CUDA需要一些特別的工具：

? ? ? 1. CUDA-GDB?

? ? ? CUDA-GDB是Linux和Mac環境中GNU調試器的一個端口（Port），具體信息參見：GNU-GDB

? ? ? 2. NVIDIA Parallel Nsight

? ? ? NVIDIA Parallel Nsight調試和性能分析器是可以在Windows環境下作為Microsoft Visual Studio的插件使用，具體參見：NVIDIA Parallel Nsight

? ? ? 3. 一些第三方調試工具

? ? ? 一些第三方工具也支持CUDA的調試，具體參見：Debugging solutions

5.3 數值精度問題

? ? ? 由于CUDA使用的是浮點數進行運算，涉及到一些精確度問題，需要我們在編程時注意一下。

1、單精度與雙精度問題

? ? ? 單精度和雙精度浮點運算的結果差別是很大的，在CUDA中，只有運算能力大于或等于1.3的硬件才能本地支持雙精度運算。執行運算時，程序員務必搞清進行的是哪一種運算以獲得正確的結果。nvcc編譯命令行中使用-arch=sm_13可開啟雙精度運算。 ? ? ?

2、浮點運算不遵從結合律

? ? ? 對于三個浮點數A、B以及C，需要注意的是，(A+B)+C并不等于A+(B+C)。

3、雙精度擴展和單精度截斷

? ? ? 由于單精度浮點數和雙精度浮點數運算結果的不一致性，CUDA程序員應注意避免一些精度細節問題，如下面這段代碼：

[cpp] view plaincopy

float?a;????

a?=?a?*?1.02;???

在C語言中，1.02會被隱含解釋為double類型，那么第二個式子右邊的a將會被擴展為雙精度浮點數而執行double乘法，得到的結果為雙精度浮點數，最后再將這一結果截斷為單精度浮點數賦值給式子左邊的a。這會帶來隱患，解決這一問題的方法是用1.02f來表示單精度浮點常數。

6. 性能度量（Performance Metrics）

? ? ? 為了優化CUDA程序的性能，我們需要一種定量的方法來對程序性能進行衡量，同時，我們也需要明晰帶寬（bandwidth）在性能衡量中所扮演的角色。下面我們將依依介紹這些概念。

6.1 定時

CUDA調用和內核執行可以使用CPU和GPU定時器來計時，使用這些計時器需要注意相關的問題：

1、使用CPU定時器

使用CPU定時器需要注意很多CUDA API函數是異步的，即這些函數在任務完成之前就會把控制權返回給CPU。它們一般都會在函數名之前帶有Async，遇到這類函數時需要特別小心。為了達到精準定時的效果，我們需要同步CPU和GPU，這可以通過在開始和停止定時器前調用函數cudaDeviceSynchronize()來實現。cudaDeviceSynchronize()阻塞當前CPU線程，直到指定流（Stream）中所有之前的CUDA調用運行完畢之后CPU線程才開始繼續執行。類似地，CPU線程也可以與GPU流或者GPU事件進行同步，使用cudaStreamSynchronize()函數或者cudaEventSynchronize()函數可以達到這樣的效果。然而需要注意的是，上述兩個函數不適合在默認流以外的其他流中使用，因為驅動程序可能安排不同流中的代碼交錯運行，這樣就會造成計時錯誤。而默認流（Stream 0）在設備上表現為串行運行，因此可以使用上面兩個函數進行計時。

2、使用CUDA GPU定時器

CUDA event API提供了用于創建和消除以及記錄事件（使用timestamp）的調用函數，通過timestamp之間的差值換算，我們可以得到毫秒級的GPU運行時間信息。下面的代碼片段展示了這一技術：

How to time code using CUDA events

[cpp] view plaincopy

cudaEvent_t?start,?stop;????

float?time;????

cudaEventCreate(&start);????

cudaEventCreate(&stop);????

????

cudaEventRecord(start,?0);????

kernel<<<grid,?threads>>>(d_odata,?d_idata,?size_x,?size_y,?NUM_REPS);????

cudaEventRecord(stop,?0);????

cudaEventSynchronize(stop);????

????

cudaEventElapsedTime(&time,?start,?stop);????

cudaEventDestroy(start);????

cudaEventDestroy(stop);??

這里cudaEventRecord()用于將start和stop事件放置到默認流（stream 0），當設備運行到事件處時，將會為該事件記錄一個時間戳（timestamp）。cudaEventElapsedTime()函數返回記錄中start和stop事件之間的毫秒級時間差。由于記錄是發生在GPU上的，其時鐘是獨立的，因此時鐘分辨率也是操作系統無關的。GPU計時詳細請參見：《CUDA Toolkit Reference Manual》.

6.2 帶寬

目前對此理解尚淺，留個占位符以待今后補充。

總結

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

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