背景
隨着時鍾頻率的發展陷入停滯，集成更多的計算邏輯和計算核心成為了獲取更高算力的主要途徑。多核處理器可以視作多路平台的自然演化，而 GPGPU 的出現利用大規模並行架構為多核 CPU 難以解決的問題提供了顛覆性的解決方案。
現代計算機模糊了 Flynn 分類的界限，將 MIMD 與 SIMD 結合以獲得更高性能。現代計算機的發展趨勢主要有兩種：一種是增加片上內核數目與緩存容量並結合專用 SIMD 指令集，另一種采用異構架構並分別處理不同任務。異構計算主要是指使用不同類型指令集和體系架構的計算單元組成系統的計算方式。常見的計算單元類別包含 CPU、GPU 等協處理器、DSP、ASIC、FPGA 等。而 CPU+GPU 在這些異構架構中最為常見。
相較於 CPU，GPU 有更少的片上緩存和大量能夠並行執行的簡單 ALU。CPU 在設計上針對由短序列計算操作和難以預測的控制流程構成的動態工作負載進行優化，而 GPU 則被設計以處理計算任務主導且帶有簡單控制流的工作負載。因此CPU 計算適合處理控制密集型任務，而 GPU 計算適合處理包含數據並行的計算密集型任務。二者的功能互補性促成了 CPU+GPU 的異構並行計算架構的發展。
當然，一切多核與眾核芯片帶來的性能提升都不是免費的，需要我們對大量傳統算法進行重新設計。只有充分理解 CPU 與 GPU 架構，用並行思維思考，才能設計和編寫高效的異構並行計算程序。
CUDA 是一種異構計算平台。在 CUDA 平台上，可以使用標准程序語言的擴展、API、編譯器指令、CUDA 加速庫等編寫程序，以利用 NVIDIA GPU 高效地處理復雜的並行計算問題。
CUDA 編程模型
CUDA 采用全局串行局部並行的編程模型，GPU 作為協處理器對程序的部分進行加速。我們希望將程序中的可並發部分分解為成百上千個線程，交由 GPU 並發執行以充分發揮其性能。
事實上，CPU 與 GPU 中的線程略有區別：CPU 線程是重量級實體，操作系統交替執行線程，線程上下文切換花銷很大，而 GPU 線程是輕量級的，GPU 應用一般包含成千上萬的線程，多數在排隊狀態，線程之間切換基本沒有開銷。CPU 的核被設計用來盡可能減少一個或兩個線程運行時間的延遲，而 GPU 核則充分為大量線程做足准備，最大幅度提高吞吐量。這些線程會被調度到數以千計的流式處理單元上執行，並通過合理的分層結構來控制。在執行模型中，我們會進一步說明。
如何從編程意義上組織成百上千的並發線程？在 CUDA 中，每次核函數調用產生一組線程，它們使用公用的函數參數來執行相同的功能，通過內部結構變量定義自己在線程結構中的位置。具體地，CUDA 將線程組織成 6 維的超立體結構，它由兩層 3 維組織結構嵌套而成。線程組織成 3 維的線程塊結構，線程塊再組織成 3 維的線程網格結構。每個線程通過兩個 3 維的內部結構變量 threadIdx, blockIdx 結合 threadDim, blockDim 來定義自己在線程結構中的位置，並實現位置信息與所分配數據子集的映射。事實上，這種結構的設計綜合了軟件與硬件方面的考慮。
線程在編程時以函數體現，在 GPU 上執行的函數稱為核函數。帶有 __global__ 修飾的函數只允許在設備（即 GPU）上執行，它們只能由主機調用。主機端調用的核函數沒有返回值，而核函數的輸入和輸出都存儲在設備內存中，需要顯式地與主機內存之間進行拷貝操作，后續章節中會詳細討論。
我們常用 host 指代 CPU 及其內存，而用 device 指代 GPU 及其內存。CUDA 程序中包含 host 程序和 device 程序分別在 CPU 和 GPU 上運行。它們之間通過以內存拷貝為主的方式進行通信。
典型的 CUDA 程序執行流程如下：
- 分配 host 內存，並進行數據初始化；
- 分配 device 內存，並從 host 將數據拷貝到 device 上；
- 調用 CUDA 的核函數在 device 上完成指定的運算；
- 將 device 上的運算結果拷貝到 host 上；
- 釋放 device 和 host 上分配的內存。
在 CUDA 編程中通過函數類型限定詞區別 host 和 device 函數，主要的三個函數類型限定詞如下：
__global__：在 device 上執行，一般從 host 中調用；__device__：在 device 上執行，單僅可以從 device 中調用；__host__：在 host 上執行，僅可以從 host 上調用，一般省略不寫。
部分限定詞可以同時使用，讀者可查閱有關資料進一步了解。
Host 程序在調用核函數時通過特定語法 <<<g,b>>> 來指定線程的組織層次。使用 dim3 數據類型來表示包含三個元素的整型向量，如果提供數據不足 3 個，則缺省補 1。只需要使用一維時，可以簡寫為標量。
簡舉一例，用核函數進行向量加法。限於篇幅，以下程序段僅展示核函數的定義與調用，完整實例程序請參閱 NVIDIA 官方示例或查閱其它資料。
__global__ void vectorAdd(const float *A, const float *B, float *C, int numElements)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < numElements) {
C[i] = A[i] + B[i];
}
}
int main()
{
// ...
cudaMemcpy(d_A, A, length, cudaMemcpyHostToDevice);
// ...
vectorAdd <<<blocksPerGrid, threadsPerBlock >>>(d_A, d_B, d_C, numElements);
// ...
cudaMemcpy(C, d_C, length, cudaMemcpyDeviceToHost);
// ...
}
最后，需要指出的是，在編程模型中，盡管我們努力地不去談論任何硬件的組織結構，有一個問題無法避免：計算能力（Compute Capability）。計算能力反映了架構的性能特點，但名不副實的是，它的數值並非一個絕對的性能指標，而更像是架構的版本號，有時我們也將它稱為 SM version。計算能力對線程與線程塊的組織方式，內存的分配限制給出了詳細的界限。因此，在編程時，我們務必考慮目標平台的計算能力，或者提供較大的可擴放性。
CUDA 執行模型
CUDA 采用單指令多線程（SIMT）執行模型。SIMT 非常類似於 SIMD，而不同主要在於處理單元處理的“向量”大小是軟件定義的線程塊的大小。核函數在設備上運行時，相同的指令序列會被大量流處理單元（SP）部分同步地執行。
流式多處理器（SM）是在同一控制單元下執行的一組流處理單元。GPU 中通常包括數十個 SM ，每個 SM 支持數百線程並發執行。
線程以線程塊為單位，調度到不同的 SM 上執行。SM 和線程塊是一對多的關系：線程塊被調度到一個 SM 並保持在其上直到執行結束，而單個 SM 可以同時容納多個線程塊。
線程的調度單位是線程塊，而執行單位則是進一步細分的線程束。換言之，同一線程塊內的所有線程未必在物理意義上並行執行。線程塊被調度到 SM 后，進一步細分為若干固定大小的線程束，這一大小與硬件相關，通常為 32。
同束線程被共同的控制單元執行，它們必須同時執行相同的指令。作為 SIMT 與 SIMD 的不同之處，線程束中的線程擁有獨立的寄存器（物理上是寄存器堆中的條目）、指令地址計數器等狀態部件，並可以擁有獨立的執行路徑。
實現上，SM 獲取指令，將指令廣播道所有 SP 中，所有 SP 一起執行。 SM 通常擁有兩套以上的線程束調度器和指令分派器。線程束調度器選擇一個線程束，將指令發送到一個流處理單元組中。
CUDA 存儲模型
為了取得大容量、高性能與生產成本間的平衡，必須依靠內存模型來獲得最佳的延遲和帶寬。
回憶在 CPU 中，我們為了使得 CPU 能全速運行，多需要為任意隨機地址的快速訪問提供保障，因此提供了較大的多級片上緩存。而非通用 GPU 的工作以過濾和傳輸圖形信息為典型，在處理一次讀取的大量數據后，這些數據通常不必保持在片上。因此，原先的 GPU 需要高速的數據總線，而只需要少量的片上緩存。
通用計算時代，GPGPU 正在顛覆這一點。如今的 GPU 帶有越來越大的片上內存與緩存。與 CPU 不同的式，GPU 的內存層次結構對用戶並不透明。用戶需要部分對內存類型選擇、數據移動進行選擇和介入。
在 CUDA 中，我們關注的存儲層次結構主要包括寄存器、本地內存、共享內存、緩存、全局內存、常量內存、紋理內存。在邏輯結構中的示意圖如下：
它們的位置、訪問范圍與生命周期的簡要對比如下表所示：
| 類型 | 位置 | 邏輯訪問范圍 | 生命周期 |
|---|---|---|---|
| 寄存器 | 片上 | 線程 | 線程 |
| 本地內存 | 片外 | 線程 | 線程 |
| 共享內存 | 片上 | 塊 | 塊 |
| 全局內存 | 片外 | 網格 | 主機控制 |
| 常量/紋理內存 | 片外 | 網格 | 主機控制 |
寄存器用於保存局部變量，從而減少對全局內存或共享內存的訪問，加快操作的處理速度。CUDA 中的寄存器以寄存器文件（register file，寄存器堆）的形式存在，它是由多個寄存器組成的陣列，通常由 SRAM 實現。在現代 CPU 中由於寄存器重命名技術的使用，架構寄存器對應的寄存器堆中的物理存儲條目也是動態的。
若局部變量需要的空間超過計算能力的限制，則會溢出到位於本地內存的運行時棧中。注意本地內存之所以稱為本地，是因為它僅能被特定線程訪問，然而，它在物理上和全局內存一樣位於片外。
共享內存是可以被一個 SM 上所有 SP 共享的 RAM，它可以用於存儲那些本可以放在全局內存中但需要頻繁使用的數據或者全局內存部分數據的副本，以及用於在 SM 的各 SP 間共享數據。某種意義上可以將共享內存視作特殊的 L1$，事實上在 Fermi 和 Kepler 架構中共享內存與 L1$ 是同一塊片上 RAM，由用戶編程指定如何切分。
例如，若線程頻繁對某個數據進行讀寫操作，可以設置操作的數據為 __shared__，使其位於共享內存中（也稱其常駐緩存），且同一個線程塊內的所有線程共享該內存區域。當出現多個線程對同一個內存區域進行操作時，需要對線程進行同步操作。__syncthreads() 為我們提供了簡潔高效的屏障同步操作。
全局內存是設備端片外內存的主要部分，容量高但速度低。主機端與設備端的主要通過全局內存進行，因為它是主機通過 CUDA 庫函數可以訪問的唯一部分。
常量內存也是設備端片外內存的一部分，它可以被緩存，同時支持將單個數據廣播給線程束中的所有線程。
紋理內存通過對數據的大小和格式限制來實現圖像數據的快速訪問與變換。正因如此，它在 GPGPU 領域的應用較為有限。
后記
CUDA 平台為異構並行計算提供了豐富的工具，但高效並行程序的設計仍然面臨着諸多的問題與艱巨的挑戰。對 CUDA 異構並行程序設計的學習應當從硬件平台的體系結構、並行算法設計的方法與技術、並行編程的平台與工具三個方面充分着力，並主動尋求與運用領域的深入融合，充分發掘異構並行計算的潛能與優勢。
