CUDA C 程式設計權威指南 學習筆記：第二章 CUDA程式設計模型

詳細細節參考書籍或者這個部落格《CUDA C程式設計權威指南》——第2章 CUDA程式設計模型 2.1 CUDA程式設計模型概述

2.1.1 CUDA程式設計結構

CUDA程式設計模型還利用GPU架構的計算能力提供了以下幾個特有功能。
1. 一種通過層次結構在GPU中組織執行緒的方法
2. 一種通過層次結構在GPU中訪問記憶體的方法

CUDA程式設計模型主要是非同步的，因此在GPU上進行的運算可以與主機-裝置通訊重疊。一個典型的CUDA程式包括由並行程式碼互補的序列程式碼。如圖2-2所示，序列程式碼（及任務並行程式碼）在主機CPU上執行，而並行程式碼在GPU上執行。主機程式碼按照ANSI C標準進行編寫，而裝置程式碼使用CUDA C進行編寫。一個典型的CUDA程式實現流程遵循以下模式。

1. 把資料從CPU記憶體拷貝到GPU記憶體。
2. 呼叫核函式對儲存在GPU記憶體中的資料進行操作。
3. 將資料從GPU記憶體傳送回到CPU記憶體。
   

2.1.2 記憶體管理

udaMalloc與標準C語言中的malloc函式幾乎一樣，只是此函式在GPU的記憶體裡分配記憶體。通過充分保持與標準C語言執行庫中的介面一致性，可以實現CUDA應用程式的輕鬆接入。

cudaMemcpy函式負責主機和裝置之間的資料傳輸，這個函式以同步方式執行，因為在cudaMemcpy函式返回以及傳輸操作完成之前主機應用程式是阻塞的。

2.1.3 執行緒管理

CUDA明確了執行緒層次抽象的概念以便於你組織執行緒。這是一個兩層的執行緒層次結構，由執行緒塊和執行緒塊網格構成，如圖2-5所示。

• 同步
• 共享記憶體

不同塊內的執行緒不能協作。 執行緒依靠以下兩個座標變數來區分彼此。

• blockIdx（執行緒塊線上程格內的索引）
• threadIdx（塊內的執行緒索引）

網格和塊的維度由下列兩個內建變數指定。

• blockDim（執行緒塊的維度，用每個執行緒塊中的執行緒數來表示）
• gridDim（執行緒格的維度，用每個執行緒格中的執行緒數來表示）

2.1.4 啟動核函式

同一個塊中的執行緒之間可以相互協作，不同塊內的執行緒不能協作。核函式呼叫結束後，控制權立刻返回給主機端。你可以呼叫以下函式來強制主機端程式等待所有的核函式執行結束：cudaDeviceSynchronize(). 一些CUDA執行時API在主機和裝置之間是隱式同步的。當使用cudaMemcpy函式在主機和裝置之間拷貝資料時，主機端隱式同步，即主機端程式必須等待資料拷貝完成後才能繼續執行程式。

2.2 給核函式計時

2.2.1 使用CPU計時

由於GPU的執行是非同步的，所以使用CPU計時就必須要顯示做同步，cudaDeviceSynchronize()

2.2.2 使用nvprof計時

CUDA還提供了一個nvprof命令列分析工具

2.3 組織並行執行緒

從矩陣加法的例子中可以看出：

• 改變執行配置對核心效能有影響
• 傳統的核函式實現一般不能獲得最佳效能
• 對於一個給定的核函式，嘗試使用不同的網格和執行緒塊大小可以獲得更好的效能