隨著多核CPU和眾核GPU的到來，并行編程已經(jīng)得到了業(yè)界越來越多的重視，CPU-GPU異構程序能夠極大提高現(xiàn)有計算機系統(tǒng)的運算性能，對于科學計算等運算密集型程序有著非常重要的意義。這一系列文章是根據(jù)《CUDA C語言編程指南》來整理的，該指南是NVIDIA公司提供的CUDA學習資料，介紹了CUDA編程最基本最核心的概念，是學習CUDA必不可少的閱讀材料。

初學CUDA，筆記錯誤之處在所難免，還請

隨著多核CPU和眾核GPU的到來，并行編程已經(jīng)得到了業(yè)界越來越多的重視，CPU-GPU異構程序能夠極大提高現(xiàn)有計算機系統(tǒng)的運算性能，對于科學計算等運算密集型程序有著非常重要的意義。這一系列文章是根據(jù)《CUDA C語言編程指南》來整理的，該指南是NVIDIA公司提供的CUDA學習資料，介紹了CUDA編程最基本最核心的概念，是學習CUDA必不可少的閱讀材料。

初學CUDA，筆記錯誤之處在所難免，還請發(fā)現(xiàn)問題的諸位讀者不吝賜教。

1. 什么是CUDA？

CUDA全稱是Compute Unified Device Architecture，中文名稱即統(tǒng)一計算設備架構，它是NVIDIA公司提出了一種通用的并行計算平臺和編程模型。使用CUDA，我們可以開發(fā)出同時在CPU和GPU上運行的通用計算程序，更加高效地利用現(xiàn)有硬件進行計算。為了簡化并行計算學習，CUDA為程序員提供了一個類C語言的開發(fā)環(huán)境以及一些其它的如FORTRAN、DirectCOmpute、OpenACC的高級語言/編程接口來開發(fā)CUDA程序。

2. CUDA編程模型如何擴展？

我們知道，不同的GPU擁有不同的核心數(shù)目，在核心較多的系統(tǒng)上CUDA程序運行的時間較短，而在核心較少的系統(tǒng)上CUDA程序的執(zhí)行時間較多。那么，CUDA是如何做到的呢？ 并行編程的中心思想是分而治之：將大問題劃分為一些小問題，再把這些小問題交給相應的處理單元并行地進行處理。在CUDA中，這一思想便體現(xiàn)在它的具有兩個層次的問題劃分模型。一個問題可以首先被粗粒度地劃分為若干較小的子問題，CUDA使用被稱為塊（Block）的單元來處理它們，每個塊都由一些CUDA線程組成，線程是CUDA中最小的處理單元，將這些較小的子問題進一步劃分為若干更小的細粒度的問題，我們便可以使用線程來解決這些問題了。對于一個普通的NVIDIA GPU，其CUDA線程數(shù)目通常能達到數(shù)千個甚至更多，因此，這樣的問題劃分模型便可以成倍地提升計算機的運算性能。 GPU是由多個流水多處理器構成的，流水處理器以塊（Block）為基本調(diào)度單元，因此，對于流水處理器較多的GPU，它一次可以處理的塊（Block）更多，從而運算速度更快，時間更短。而反之對于流水處理器較少的GPU，其運算速度便會較慢。這一原理可以通過下圖形象地看出來：

3. CUDA基本概念

本節(jié)將介紹CUDA的一些基本的編程概念，該節(jié)用到的例子來自于CUDA Sample中的VectorAdd項目。

3.1 內(nèi)核（Kernels）

CUDA C是C語言的一個擴展，它允許程序員定義一種被稱為內(nèi)核函數(shù)（Kernel Functions）的C函數(shù)，內(nèi)核函數(shù)運行在GPU上，一旦啟動，CUDA中的每一個線程都將會同時并行地執(zhí)行內(nèi)核函數(shù)中的代碼。

內(nèi)核函數(shù)使用關鍵字__global__來聲明，運行該函數(shù)的CUDA線程數(shù)則通過<<<...>>>執(zhí)行配置語法來設置。（參見章節(jié)"C語言擴展"），每一個執(zhí)行內(nèi)核函數(shù)的線程都由一個唯一的線程ID，這一ID可以通過在內(nèi)核函數(shù)中訪問threadIdx變量來得到。

下面通過一些示例代碼來展示剛剛提到的這些概念該如何應用在編程中：

[cpp] view plaincopy

//?Kernel?definition????

__global__?void?VecAdd(float*?A,?float*?B,?float*?C)?{????

????int?i?=?threadIdx.x;????

????C[i]?=?A[i]?+?B[i];????

}????

????

int?main()?{????

????...????

????//?Kernel?invocation?with?N?threads????

????VecAdd<<<1,?N>>>(A,?B,?C);????

???...????

}???

在上面的代碼中，N個線程將會并行地同時執(zhí)行加法運算。

3.2 線程層次（Thread Hierarchy）

CUDA的每一個線程都有其線程ID，線程的ID信息由變量threadIdx給出。threadIdx是CUDA C語言的內(nèi)建變量，通常它用一個三維數(shù)組來表示。使用三維數(shù)組的方便之處在于可以很方便地表示一維、二維和三維線程索引，進而方便地表示一維、二維和三維線程塊（thread block）。這樣，無論是數(shù)組、矩陣還是體積的計算，都可以很容易地使用CUDA進行運算。 線程的索引與線程ID之間存在著直接的換算關系，對于一個索引為(x, y, z)的線程來說： ? ? ? 1、如果線程塊（block）是一維的，則線程ID = x ? ? ? 2、如果線程塊是二維的，假設塊尺寸為（Dx，Dy），那么線程ID = x + y * Dx ? ? ? 3、如果線程塊是三維的，設其尺寸為（Dx，Dy，Dz），那么線程ID = x + y * Dx + z * Dx * Dy ? ? ? 下面的例子展示了兩個NxN矩陣相加的CUDA實現(xiàn)： [cpp] view plaincopy

//?Kernel?definition????

__global__?void?MatAdd(float?A[N][N],?float?B[N][N],?float?C[N][N])?{????

????int?i?=?threadIdx.x;????

????int?j?=?threadIdx.y;????

????C[i][j]?=?A[i][j]?+?B[i][j];????

}????

????

int?main()?{????

????...????

????//?Kernel?invocation?with?one?block?of?N?*?N?*?1?threads????

????int?numBlocks?=?1;????

????dim3?threadsPerBlock(N,?N);????

????MatAdd<<<numBlocks,?threadsPerBlock>>>(A,?B,?C);????

????...????

}????

每個線程塊（block）中的線程數(shù)量是有限制的，因為依據(jù)前面所說，同一線程塊（block）中的所有線程都會被分配到同一個處理器核上運行，共享有限的存儲資源，因此對于當前的GPU，線程塊所能包含的最大線程數(shù)目為1024。 上面的例子中numBlocks代表線程塊的數(shù)量，這里的值為1。在一般的CUDA程序中，這個值通常大于1，也就是說將會有多個線程塊被分配到多個處理器核中同時進行處理，這樣就大大提高了程序的并行性。
在CUDA中，線程塊包含在線程格（grid）當中，線程格可以是一維、二維或者三維的，線程格的尺寸一般根據(jù)待處理數(shù)據(jù)的規(guī)模或者處理器的數(shù)量來指定。線程格中所包含的線程塊數(shù)目通常遠遠大于GPU處理器核心的數(shù)目。下圖展示了線程格（grid）、線程塊（block）以及線程（thread）之間的關系：
內(nèi)核函數(shù)的調(diào)用可以簡化為kernel<<<A，B>>>(parameters)，在尖括號中，A代表線程格（grid）的尺寸，它可以是三維的，用類型dim3表示，也可以是一維的，用int類型表示。B代表線程塊（block）的尺寸，它與A類似，也可分別用dim3或int類型表示。 在內(nèi)核函數(shù)內(nèi)部，CUDA為我們內(nèi)建了一些變量用于訪問線程格、線程塊的尺寸和索引等信息，它們是： 1. gridDim：代表線程格（grid）的尺寸，gridDim.x為x軸尺寸，gridDim.y、gridDim.z類似。拿上圖來說，它的gridDim.x = 3，gridDim.y = 2，gridDim.z = 1。 2. blockIdx：代表線程塊（block）在線程格（grid）中的索引值，拿上圖來說，Block（1,1）的索引值為：blockIdx.x = 1，blockIdx.y = 1。 3. blockDim：代表線程塊（block）的尺寸，blockDIm.x為x軸尺寸，其它依此類推。拿上圖來說，注意到Block（1,1）包含了4 * 3個線程，因此blockDim.x = 4, blockDim.y = 3。 4. threadIdx：線程索引，前面章節(jié)已經(jīng)詳細探討過了，這里不再贅述。 明白了這些變量的含義，那么下面的矩陣加法程序便不難理解了： [cpp] view plaincopy

//?Kernel?definition????

__global__?void?MatAdd(float?A[N][N],?float?B[N][N],?float?C[N][N])?{????

????int?i?=?blockIdx.x?*?blockDim.x?+?threadIdx.x;????

????int?j?=?blockIdx.y?*?blockDim.y?+?threadIdx.y;????

????if?(i?<?N?&&?j?<?N)????

????????C[i][j]?=?A[i][j]?+?B[i][j];????

}????

????

int?main()?{????

????...????

????//?Kernel?invocation????

????dim3?threadsPerBlock(16,?16);????

????dim3?numBlocks(N?/?threadsPerBlock.x,?N?/?threadsPerBlock.y);????

????MatAdd<<<numBlocks,?threadsPerBlock>>>(A,?B,?C);????

????...????

}??

在上面的程序中，線程塊（block）的尺寸是16x16，這是CUDA編程中一個非常普遍的選擇。線程格（grid）包含了足夠多的線程塊（block）來進行計算。 線程塊（block）是獨立執(zhí)行的，在執(zhí)行的過程中線程塊之間互不干擾，因此它們的執(zhí)行順序是隨機的。 同一線程塊中的線程可以通過訪問共享內(nèi)存（shared memory）或者通過同步函數(shù)__syncthreads()來協(xié)調(diào)合作。這些概念將在以后的章節(jié)中詳細解釋。

3.3 內(nèi)存層次（Memory Hierarchy）

在GPU上CUDA線程可以訪問到的存儲資源有很多，每個CUDA線程擁有獨立的本地內(nèi)存（local Memory）；每一個線程塊（block）都有其獨立的共享內(nèi)存（shared memory），共享內(nèi)存對于線程塊中的每個線程都是可見的，它與線程塊具有相同的生存時間；同時，還有一片稱為全局內(nèi)存（global memory）的區(qū)域?qū)λ械腃UDA線程都是可訪問的。

除了上述三種存儲資源以外，CUDA還提供了兩種只讀內(nèi)存空間：常量內(nèi)存（constant memory）和紋理內(nèi)存（texture memory），同全局內(nèi)存類似，所有的CUDA線程都可以訪問它們。對于一些特殊格式的數(shù)據(jù)，紋理內(nèi)存提供多種尋址模式以及數(shù)據(jù)過濾方法來操作內(nèi)存。這兩類存儲資源主要用于一些特殊的內(nèi)存使用場合。

一個程序啟動內(nèi)核函數(shù)以后，全局內(nèi)存、常量內(nèi)存以及紋理內(nèi)存將會一直存在直到該程序結束。下面是CUDA的內(nèi)存層次圖：

3.4 異構編程（Heterogeneous Programming）

CUDA的異構編程模型假定CUDA線程都運行在一個可被看做CPU協(xié)處理器的芯片上，這就使得CUDA內(nèi)核函數(shù)可以和CPU端C程序的運行并行運行，從而加快程序的運行效率。為了達到這個效果，CUDA程序需要管理兩大塊由DRAM構成的內(nèi)存區(qū)域：CPU端可以訪問到的主機內(nèi)存（host memory）以及GPU端供CUDA內(nèi)核訪問到的設備內(nèi)存（device memory），設備內(nèi)存主要由全局內(nèi)存、常量內(nèi)存以及紋理內(nèi)存構成。現(xiàn)在，CUDA程序的運行機制便很明了了：CPU端代碼生成原始數(shù)據(jù)，通過CUDA運行時函數(shù)庫將這些原始數(shù)據(jù)傳輸?shù)紾PU上，在CPU端啟動CUDA內(nèi)核函數(shù)進行運算，然后將運算結果從設備端傳輸?shù)街鳈C端，計算任務便完成了。

4.?CUDA C語言編程接口

異構程序設計跟傳統(tǒng)的串行程序設計差別是很大的，學習起來也是非常不容易的。NVIDIA非常夠意思，為了簡化CUDA的學習曲線，它采用了絕大多數(shù)程序員都熟悉的C語言作為其根基，CUDA C是NVIDIA為程序員提供的一類編程接口，它實際上是一個C語言的擴展，在C的基礎上增加了一些新的語法和變量，并且提供了功能豐富的庫函數(shù)，方便程序員使用GPU進行異構計算。 除了前面章節(jié)提到的CUDA最基本、最核心的概念以外，CUDA C呈現(xiàn)給程序員的接口主要由兩大類API構成，它們分別是CUDA Runtime API和CUDA Driver API，Runtime API實際上是對于Driver API的封裝，其目的自然是方便程序員的代碼編寫工作。Driver API為用戶提供了更細一層的控制手段，通過它可以控制諸如CUDA Contexts（一種類似主機進程的概念）以及CUDA Modules（類似主機動態(tài)加載庫的概念）等更加底層的CUDA模塊。

4.1 NVCC編譯器

任何一種程序設計語言都需要相應的編譯器將其編譯為二進制代碼，進而在目標機器上得到執(zhí)行。對于異構計算而言，這一過程與傳統(tǒng)程序設計語言是有一些區(qū)別的。為什么？因為CUDA它本質(zhì)上不是一種語言，而是一種異構計算的編程模型，使用CUDA C寫出的代碼需要在兩種體系結構完全不同的設備上執(zhí)行：1、CPU；2、GPU。因此，CUDA C的編譯器所做的工作就有點略多了。一方面，它需要將源代碼中運行在GPU端的代碼編譯得到能在CUDA設備上運行的二進制程序。另一方面，它也需要將源代碼中運行在CPU端的程序編譯得到能在主機CPU上運行的二進制程序。最后，它需要把這兩部分有機地結合起來，使得兩部分代碼能夠協(xié)調(diào)運行。 CUDA C為我們提供了這樣的編譯器，它便是NVCC。嚴格意義上來講，NVCC并不能稱作編譯器，NVIDIA稱其為編譯器驅(qū)動（Compiler Driver），本節(jié)我們暫且使用編譯器來描述NVCC。使用nvcc命令行工具我們可以簡化CUDA程序的編譯過程，NVCC編譯器的工作過程主要可以劃分為兩個階段：離線編譯（Offline Compilation）和即時編譯（Just-in-Time Compilation）。 離線編譯（Offline Compilation） ? ? 下面這幅圖簡單說明了離線編譯的過程：

? ? ? ? 在CUDA源代碼中，既包含在GPU設備上執(zhí)行的代碼，也包括在主機CPU上執(zhí)行的代碼。因此，NVCC的第一步工作便是將二者分離開來，這一過程結束之后： ? ? ? ? 1. 運行于設備端的代碼將被NVCC工具編譯為PTX代碼（GPU的匯編代碼）或者cubin對象（二進制GPU代碼）； ? ? ? ? 2. 運行于主機端的代碼將被NVCC工具改寫，將其中的內(nèi)核啟動語法（如<<<...>>>）改寫為一系列的CUDA Runtime函數(shù)，并利用外部編譯工具（gcc for linux，或者vc compiler for windows）來編譯這部分代碼，以得到運行于CPU上的可執(zhí)行程序。 ? ? ? ? 完事之后，NVCC將自動把輸出的兩個二進制文件鏈接起來，得到異構程序的二進制代碼。 即時編譯（Just-in-time Compile） ? ? ? ? 任何在運行時被CUDA程序加載的PTX代碼都會被顯卡的驅(qū)動程序進一步編譯成設備相關的二進制可執(zhí)行代碼。這一過程被稱作即時編譯（just-in-time compilation）。即時編譯增加了程序的裝載時間，但是也使得編譯好的程序可以從新的顯卡驅(qū)動中獲得性能提升。同時到目前為止，這一方法是保證編譯好的程序在還未問世的GPU上運行的唯一解決方案。 ? ? ? ? 在即時編譯的過程中，顯卡驅(qū)動將會自動緩存PTX代碼的編譯結果，以避免多次調(diào)用同一程序帶來的重復編譯開銷。NVIDIA把這部分緩存稱作計算緩存（compute cache），當顯卡驅(qū)動升級時，這部分緩存將會自動清空，以使得程序能夠自動獲得新驅(qū)動為即時編譯過程帶來的性能提升。 ? ? ? ? 有一些環(huán)境變量可以用來控制即時編譯過程： ? ? ? ? 1. 設置CUDA_CACHE_DISABLE為1將會關閉緩存功能 ? ? ? ? 2. CUDA_CACHE_MAXSIZE變量用于指定計算緩存的字節(jié)大小，默認情況下它的值是32MB，它最大可以被設置為4GB。任何大于緩存最大值得二進制代碼將不會被緩存。在需要的情況下，一些舊的二進制代碼可能被丟棄以騰出空間緩存新的二進制代碼。 ? ? ? ? 3. CUDA_CACHE_PATH變量用于指定計算緩存的存儲目錄地址，它的缺省值如下：
? ? ? ? 4. 設置CUDA_FORCE_PTX_JIT為1會強制顯卡驅(qū)動忽略應用程序中的二進制代碼并且即時編譯程序中的嵌入PTX代碼。如果一個內(nèi)核函數(shù)沒有嵌入的PTX代碼，那么它將會裝載失敗。該變量可以用來確認程序中存在嵌入的PTX代碼。同時，使用即時編譯（just-in-time Compilation）技術也可確保程序的向前兼容性。

4.2 兼容性

? ? ??1、二進制兼容性 ? ? ? 二進制代碼是設備相關的，使用NVCC編譯器編譯時，若指定-code選項，則會編譯產(chǎn)生目標設備的二進制cubin對象。例如，編譯時使用-code=sm_13會產(chǎn)生適用于計算能力1.3的二進制代碼。二進制代碼在CUDA計算設備上具有小版本的向前兼容性，但是在大版本上不具備兼容性。也就是說，對于計算能力X.y的硬件，使用-code=sm_Xy編譯后，程序能夠運行于計算能力X.z（其中z>=y）的硬件上，但不能運行在計算能力M.n（M!=X）的硬件上。 ? ? ??2、PTX代碼兼容性 ? ? ? 不同計算能力的設備所支持的PTX指令條數(shù)是不同的，一些PTX指令只在擁有較高計算能力的設備上被支持。例如，全局內(nèi)存（global Memory）的原子操作指令只能用于計算能力不小于1.1的設備；雙精度浮點運算指令只能用于計算能力不小于1.3的設備。在將C語言編譯為PTX代碼時，NVCC使用-arch編譯選項指定PTX代碼目標設備的計算能力。因此，要想使用雙精度運算，編譯時必須使用選項-arch=sm_13（或使用更高的計算能力），否則NVCC會自動將雙精度操作降級為單精度操作。 ? ? ? 為某一特定設備產(chǎn)生的PTX代碼，在運行時總是能夠被具有更高計算能力的設備JIT編譯為可執(zhí)行的二進制代碼。 ? ? ? 3、應用程序兼容性 ? ? ? 執(zhí)行CUDA程序有兩種方式，一種是直接加載編譯好的CUDA二進制代碼運行，另一種是首先加載程序中的PTX代碼，再執(zhí)行JIT編譯得到二進制的設備可執(zhí)行文件，然后運行。特別需要注意的是，為了讓程序運行具有更高計算能力的未來設備上，必須讓程序加載PTX代碼。 ? ? ? 事實上，在一個CUDA C程序中可以嵌入不止一個版本的PTX/二進制代碼。那么，具體執(zhí)行時哪一個版本的PTX或者二進制代碼會得到執(zhí)行呢？答案是：最兼容的那個版本。例如編譯一個名為x.cu的CUDA源代碼：
將會產(chǎn)生兼容計算能力1.1硬件的二進制代碼（第一排的-gencode選項）以及兼容計算能力1.1設備的PTX和二進制代碼，這些代碼都將會嵌入到編譯后的目標文件中。 ? ? ? 主機端將會產(chǎn)生一些額外的代碼，在程序運行時，這些代碼會自動決定裝載哪一個版本的代碼來執(zhí)行。對于上面的例子：

• 計算能力1.0的設備運行該程序?qū)b載1.0版本的二進制代碼
• 計算能力1.1、1.2或者1.3的設備運行該程序?qū)b載1.1版本的二進制代碼
• 計算能力2.0或者更高的設備運行該程序?qū)b載1.1版本的PTX代碼進而對其進行JIT編譯得到相應設備的二進制代碼

? ? ??同時，x.cu還可以在程序中使用一些特殊的宏來改變不同設備的代碼執(zhí)行路徑。例如，對于計算能力1.1的設備而言，宏__CUDA_ARCH__等于110，在程序中可以對該宏的值進行判斷，然后分支執(zhí)行程序。 ? ? ? NVCC用戶手冊列出了很多-arch，-code和-gencode等編譯選項的簡化書寫形式。例如，-arch=sm_13就是-arch=compute_13 -code=compute13, sm_13的簡化形式。更多詳盡的內(nèi)容請參閱該手冊。 ? ? ? 4、C/C++兼容性 ? ? ? NVCC編譯器前端使用C++語法啊規(guī)則來處理CUDA源文件。在主機端，CUDA支持完整的C++語法；而在設備端，只有部分C++語法是被支持的。這方面更為詳盡的討論請參見《CUDA C程序設計指南》的C/C++語言支持章節(jié)。 ? ? ? 5、64位兼容性 ? ? ? 64位版本的nvcc編譯器將設備代碼編譯為64位模式，即指針是64位的。運行64位設備代碼的先決條件是主機端代碼必須也使用64位模式進行編譯。同樣，32位版本的nvcc將設備代碼編譯為32位模式，這些代碼也必須與相應的32位主機端代碼相配合方能運行。 ? ? ? 32位nvcc編譯器可以使用-m64編譯選項將設備代碼編譯為64位模式。同時64位nvcc編譯器也可使用-m32編譯選項將設備代碼編譯為32位模式。

4.3 CUDA C Runtime

? ? ? ? CUDA C Runtime使用cudart動態(tài)鏈接庫實現(xiàn)（cudart.dll或者cudart.so），運行時中所有的入口函數(shù)都以cuda為前綴。

4.3.1 初始化

CUDA C Runtime函數(shù)庫沒有明確的初始化函數(shù)，在程序第一次調(diào)用Runtime庫函數(shù)時它會自動初始化。因此，在記錄Runtime函數(shù)調(diào)用時間和理解程序中第一個Runtime調(diào)用返回的錯誤代碼時，需要將初始化考慮在內(nèi)。 ? ? ? ? 在初始化期間，Runtime將會為系統(tǒng)中每一個設備創(chuàng)建一個CUDA上下文（類似CPU中進程的數(shù)據(jù)結構），這個上下文是設備的基本上下文，它被程序中所有的主機線程所共享。創(chuàng)建過程在后臺運行，并且，Runtime將隱藏基本上下文使之對Runtime API這一層次的程序員不可見。 ? ? ? ? 當一個主機線程調(diào)用cudaDeviceReset()函數(shù)時，它將會銷毀線程當前控制設備的基本上下文。也就是說，當線程下一次調(diào)用runtime函數(shù)時將會重啟初始化，一個新的CUDA基本上下文將被創(chuàng)建出來。

4.3.2 設備內(nèi)存

? ? ? ? 正如前面異構計算章節(jié)所講，CUDA編程模型假定系統(tǒng)是由主機和設備構成的，它們分別具有自己獨立的內(nèi)存空間。Runtime負責設備內(nèi)存的分配，回收，拷貝以及在主機和設備間傳輸數(shù)據(jù)的工作。

? ? ? ? 設備內(nèi)存可以有兩種分配方式：線性內(nèi)存或者CUDA數(shù)組

? ? ? ? CUDA數(shù)組是一塊不透明的內(nèi)存空間，它主要被優(yōu)化用于紋理存取。

? ? ? ? 線性內(nèi)存空間與平時我們訪問的內(nèi)存類似，對于計算能力1.x的設備來說，它存在于一個32位的地址空間。對于更高計算能力的設備而言，它存在于一個40位的地址空間中。因此，單獨分配的實體可以使用指針來相互應用。

? ? ? ? 我們通常使用cudaMalloc()函數(shù)分配線性內(nèi)存空間，使用cudaFree()函數(shù)釋放線性內(nèi)存空間，使用cudaMemcpy()函數(shù)在主機和設備之間傳輸數(shù)據(jù)。下面是CUDA Vector Add代碼示例的一些片段：

[cpp] view plaincopy

//?Device?code????

__global__?void?VecAdd(float?*A,?float?*B,?float?*C,?int?N)?{????

????int?i?=?blockDim.x?*?blockIdx.x?+?threadIdx.x;????

????if?(i?<?N)????

????????C[i]?=?A[i]?+?B[i];????

}????

????

//?Host?code????

int?main()?{????

????int?N?=?...;????

????size_t?size?=?N?*?sizeof(float);????

????

????//?Allocate?input?vectors?h_A?and?h_B?in?host?memory????

????float?*h_A?=?(float*)malloc(size);????

????float?*h_B?=?(float*)malloc(size);????

????

????//?Initialize?input?vectors????

????...????

?????

????//?Allocate?vectors?in?device?memory????

????float?*d_A,?*d_B,?*d_C;????

????cudaMalloc(&d_A,?size);????

????cudaMalloc(&d_B,?size);????

????cudaMalloc(&d_C,?size);????

????

????//?Copy?vectors?from?host?memory?to?device?memory????

????cudaMemcpy(d_A,?h_A,?size,?cudaMemcpyHostToDevice);????

????cudaMemcpy(d_B,?h_B,?size,?cudaMemcpyHostToDevice);????

????

????//?Invoke?kernel????

????int?threadsPerBlock?=?256;????

????int?blocksPerGrid?=?(N?+threadsPerBlock?-?1)?/?threadsPerBlock;????

????VecAdd<<<blocksPerGrid,?threadsPerBlock>>>(d_A,?d_B,?d_C,?N);????

????

????//?Copy?result?from?device?memory?to?host?Memory????

????cudaMemcpy(h_C,?d_C,?size,?cudaMemcpyDeviceToHost);????

????

????//?Free?device?memory????

????cudaFree(d_A);????

????cudaFree(d_B);????

????cudaFree(d_C);????

????

????//?Free?host?memory????

????...????

}???

片段展示了設備內(nèi)存的分配，傳輸以及回收過程。

? ? ? ? 除了上面展示的方法，我們還可以使用cudaMallocPitch()和cudaMalloc3D()函數(shù)來分配線性內(nèi)存。這些函數(shù)能夠確保分配的內(nèi)存滿足設備內(nèi)存訪問的對齊要求，對于行地址的訪問以及多維數(shù)組間的數(shù)據(jù)傳輸提供高性能保證，因此非常適合對于二維和三維數(shù)組內(nèi)存空間的分配。下面的代碼片段展示了分配和使用尺寸為width x height的二維數(shù)組的技術：

[cpp] view plaincopy

//?Host?code????

int?width?=?64,?height?=?64;????

float?*devPtr;????

size_t?pitch;????

cudaMallocPitch(&devPtr,?&pitch,?width?*?sizeof(float),?height);????

MyKernel<<<100,?512>>>(devPtr,?pitch,?width,?height);????

????

//?Device?code????

__global__?void?MyKernel(float*?devPtr,?size_t?pitch,?int?width,?int?height)?{????

????for?(int?r?=?0;?r?<?height;?++r)?{????

????????float*?row?=?(float*)((char*)devPtr?+?r?*?pitch);????

????????for?(int?c?=?0;?c?<?width;?++c)?{????

????????????float?element?=?row[c];????

????????}????

????}????

}????

? ? ? ? 下面的代碼片段展示了一個尺寸為width x height x depth的三維數(shù)組的分配和使用方法：

[cpp] view plaincopy

//?Host?code????

int?width?=?64,?height?=?64,?depth?=?64;????

cudaExtent?extent?=?make_cudaExtent(width?*?sizeof(float),?height,?depth);????

cudaPitchedPtr?devPitchedPtr;????

cudaMalloc3D(&devPitchedPtr,?extent);????

MyKernel<<<100,?512>>>(devPitchedPtr,?width,?height,?depth);????

????

//?Device?code????

__global__?void?MyKernel(cudaPitchedPtr?devPitchedPtr,?int?width,?int?height,?int?depth)?{????

????char*?devPtr?=?devPitchedPtr.ptr;????

????size_t?pitch?=?devPitchedPtr.pitch;????

????size_t?slicePitch?=?pitch?*?height;????

????for?(int?z?=?0;?z?<?depth;?++z)?{????

????????char*?slice?=?devPtr?+?z?*?slicePitch;????

????????for?(int?y?=?0;?y?<?height;?++y)?{????

????????????float*?row?=?(float*)(slice?+?y?*?pitch);????

????????????for?(int?x?=?0;?x?<?width;?++x)????

????????????????float?element?=?row[x];????

????????}????

????}????

}????

更多詳細的內(nèi)容請查閱參考手冊。

? ? ? ? 下面的代碼示例展示了多種使用Runtime API訪問全局變量的技術：

[cpp] view plaincopy

__constant__?float?constData[256];????

float?data[256];????

cudaMemcpyToSymbol(constData,?data,?sizeof(data));????

cudaMemcpyFromSymbol(data,?constData,?sizeof(data));????

????

__device__?float?devData;????

float?value?=?3.14f;????

cudaMemcpyToSymbol(devData,?&value,?sizeof(float));????

????

__device__?float*?devPointer;????

float*?ptr;????

cudaMalloc(&ptr,?256?*?sizeof(float));????

cudaMemcpyToSymbol(devPointer,?&ptr,?sizeof(ptr));????

使用cudaGetSymbolAddress()函數(shù)可以獲得被聲明存儲在全局內(nèi)存中的變量地址。為了獲得分配內(nèi)存的大小，可以使用cudaGetSymbolSize()函數(shù)。

4.3 CUDA C Runtime

4.3.3 共享內(nèi)存（Shared Memory）

? ? ? ? 共享內(nèi)存是CUDA設備中非常重要的一個存儲區(qū)域，有效地使用共享內(nèi)存可以充分利用CUDA設備的潛能，極大提升程序性能。那么，共享內(nèi)存有哪些特點呢？ ? ? ? ? 1、共享內(nèi)存（shared Memory）是集成在GPU處理器芯片上的（on-chip），因此相比于存在于顯存顆粒中的全局內(nèi)存（global Memory）和本地內(nèi)存（local Memory），它具有更高的傳輸帶寬，一般情況下，共享內(nèi)存的帶寬大約是全局內(nèi)存帶寬的7-10倍。 ? ? ? ? 2、共享內(nèi)存的容量很小。根據(jù)NVIDIA官方文檔的說法，在計算能力1.x的設備中，每一個流多處理器（Streaming Multiprocessor）上的共享內(nèi)存容量為16KB。對于計算能力2.x、3.0及3.5的設備該參數(shù)為48KB。因此共享內(nèi)存是稀有資源。 ? ? ? ? 3、共享內(nèi)存在物理上被劃分為很多塊，每一塊被稱為一個存儲體（bank）。在同一時刻，CUDA設備可以同時訪問多個存儲體。因此，如果一次針對共享內(nèi)存的訪存操作需要讀取n個地址，而這n個地址恰好分布在n個不同的存儲體（bank）中，那么只需要一個存取周期就可以完成n個地址的訪存任務了。對于計算能力1.x的設備，共享內(nèi)存被平均劃分為16個存儲體。而對于計算能力2.x、3.0及3.5的設備此參數(shù)為32。在共享內(nèi)存中，相鄰兩塊32bit的數(shù)據(jù)分別屬于相鄰的兩個存儲體。存儲體每兩個時鐘周期可以傳輸32位數(shù)據(jù)。 ? ? ? ? 4、共享內(nèi)存既可以靜態(tài)分配，也可以動態(tài)分配。 ? ? ? ? 從共享內(nèi)存的這些特點中我們可以看出，它實際上相當于一個程序員可以操控的緩存（cache），下面，我們使用矩陣乘法的例子來說明如何有效使用共享內(nèi)存。 ? ? ? ? 首先，我們使用最直觀的方法來完成矩陣乘法C = A x B：讀取A的每一行和B的每一列，順次完成計算任務。矩陣乘法的示意圖如下所示：

下面是矩陣乘法的CUDA C主要實現(xiàn)代碼： [cpp] view plaincopy

//?Matrices?are?stored?in?row-major?order:????

//?M(row,?col)?=?*(M.elements?+?row?*?M.width?+?col)????

typedef?struct?{????

????int?width;????

????int?height;????

????float?*elements;????

}?Matrix;????

????

//?Thread?block?size????

#define?BLOCK_SIZE?16????

????

//?Forward?declaration?of?the?matrix?multiplication?kernel????

__global__?void?MatMulKernel(const?Matrix,?const?Matrix,?Matrix);????

????

//?Matrix?multiplication?-?Host?code????

//?Matrix?dimensions?are?assumed?to?be?multiples?of?BLOCK_SIZE????

void?MatMul(const?Matrix?A,?const?Matrix?B,?Matrix?C)?{????

????//?Load?A?and?B?to?device?memory????

????Matrix?d_A;????

????d_A.width?=?A.width;?d_A.height?=?A.height;????

????size_t?size?=?A.width?*?A.height?*?sizeof(float);????

????cudaMalloc(&d_A.elements,?size);????

????cudaMemcpy(d_A.elements,?A.elements,?size,?cudaMemcpyHostToDevice);????

????Matrix?d_B;????

????d_B.width?=?B.width;?d_B.height?=?B.height;????

????size?=?B.width?*?B.height?*?sizeof(float);????

????cudaMalloc(&d_B.elements,?size);????

????cudaMemcpy(d_B.elements,?B.elements,?size,?cudaMemcpyHostToDevice);????

????

????//?Allocate?C?in?device?memory????

????Matrix?d_C;????

????d_C.width?=?C.width;?d_C.height?=?C.height;????

????size?=?C.width?*?C.height?*?sizeof(float);????

????cudaMalloc(&d_C.elements,?size);????

????

????//?Invoke?kernel????

????dim3?dimBlock(BLOCK_SIZE,?BLOCK_SIZE);????

????dim3?dimGrid(B.width?/?dimBlock.x,?A.height?/?dimBlock.y);????

????MatMulKernel<<<dimGrid,?dimBlock>>>(d_A,?d_B,?d_C);????

????

????//?Read?C?from?device?memory????

????cudaMemcpy(C.elements,?d_c.elements,?size,?cudaMemcpyDeviceToHost);????

????

????//?Free?device?memory????

????cudaFree(d_A.elements);????

????cudaFree(d_B.elements);????

????cudaFree(d_C.elements);????

}????

????

//?Matrix?multiplication?kernel?called?by?MatMul()????

__global__?void?MatMulKernel(Matrix?A,?Matrix?B,?Matrix?C)?{????

????//?Each?thread?computes?one?element?of?C????

????//?by?accumulating?results?into?Cvalue????

????float?Cvalue?=?0;????

????int?row??=?blockIdx.y?*?blockDim.y?+?threadIdx.y;????

????int?col?=?blockIdx.x?*?blockDim.x?+?threadIdx.xl????

????for?(int?e?=?0;?e?<?A.width;?++e)????

????????Cvalue?+=?A.elements[row?*?A.width?+?e]?*?B.elements[e?*?B.width?+?col];????

????C.elements[row?*?C.width?+?col]?=?Cvalue;????

}??

可以看出，為了計算矩陣C的任何一個元素，程序都需要從全局內(nèi)存（global memory）中獲得矩陣A的一行和矩陣B的一列。因此，完成這一計算矩陣A被讀取了B.width次，矩陣B被讀取了A.height次。 現(xiàn)在我們來使用共享內(nèi)存（shared memory）實現(xiàn)矩陣乘法。假設矩陣C可以被劃分為若干個較小的子方陣Csub，我們使用一個線程塊（thread block）來負責某一子方陣的計算，線程塊中的每一個線程（thread）正好負責子方陣Csub中一個元素的計算。這樣劃分后，任何一個結果子方陣Csub'（尺寸為block_size * block_size）都是與該方陣具有相同行索引的尺寸為A.width * block_size的A的子矩陣Asub和與該方陣具有相同列索引的尺寸為block_size * B.height的B的子矩陣Bsub相乘所得到。 ? ? ? ? 為了匹配設備的計算資源，兩個子矩陣Asub和Bsub被劃分為盡可能多的分離的維度為block_size的子方陣，Csub的值便是這些子矩陣相乘后相加所得到的結果。子矩陣乘法的執(zhí)行順序都是首先將它們從全局內(nèi)存（global memory）拷貝到共享內(nèi)存（shared memory）（線程塊中的每一個線程正好負責方陣一個元素的拷貝），然后由線程自己完成相應元素的計算任務，利用寄存器存儲局部結果，最后將寄存器的內(nèi)容與新得到的計算結果依此累加起來得到最終運算結果并將其傳輸?shù)饺謨?nèi)存（global memory）中。 ? ? ? ? 通過使用這種分治的計算策略，共享內(nèi)存得到了很好的利用，采用這種方案計算完成時全局內(nèi)存中矩陣A被訪問的次數(shù)為B.width / block_size，矩陣B被訪問的次數(shù)為A.height / block_size，很明顯，這為我們節(jié)省了非常多的全局內(nèi)存帶寬。優(yōu)化后的矩陣計算示意圖如下所示：

















































? ? ? ? 為了提升計算效率，我們?yōu)轭愋蚆atrix增加了一個成員變量stride。__device__函數(shù)用來獲得和設置子矩陣的元素。下面是優(yōu)化后的代碼： [cpp] view plaincopy

//?Matrices?are?stored?in?row-major?order;????

//?M(row,?col)?=?*(M.elements?+?row?*?M.stride?+?col)????

typedef?struct?{????

????int?width;????

????int?height;????

????int?stride;????

????float*?elements;????

}?Matrix;????

????

//?Get?a?matrix?element????

__device__?float?GetElement(const?Matrix?A,?int?row,?int?col)?{????

????return?A.elements[row?*?A.stride?+?col];????

}????

????

//?Set?a?matrix?element????

__device__?void?SetElement(Matrix?A,?int?row,?int?col,?float?value)?{????

????A.elements[row?*?A.stride?+?col]?=?value;????

}????

????

//?Get?the?BLOCK_SIZExBLOCK_SIZE?sub-matrix?Asub?of?A?that?is????

//?located?col?sub-matrices?to?the?right?and?row?sub-matrices?down????

//?from?the?upper-left?corner?of?A????

__device__?Matrix?GetSubMatrix(Matrix?A,?int?row,?int?col)?{????

????Matrix?Asub;????

????Asub.width?=?BLOCK_SIZE;????

????Asub.height?=?BLOCK_SIZE;????

????Asub.stride?=?A.stride;????

????Asub.elements?=?&A.elements[A.stride?*?BLOCK_SIZE?*?row?+?BLOCK_SIZE?*?col];????

????return?Asub;????

}????

????

//?Thread?block?size????

#define?BLOCK_SIZE?16????

????

//?Forward?declaration?of?the?matrix?multiplication?kernel????

__global__?void?MatMulKernel(const?Matrix,?const?Matrix,?Matrix);????

????

//?Matrix?multiplication?-?Host?code????

//?Matrix?dimensions?are?assumed?to?be?multiples?of?BLOCK_SIZE????

void?MatMul(const?Matrix?A,?const?Matrix?B,?Matrix?C)?{????

????//?Load?A?and?B?to?device?memory????

????Matrix?d_A;????

????d_A.width?=?d_A.stride?=?A.width;????

????d_A.height?=?A.height;????

????size_t?size?=?A.width?*?A.height?*?sizeof(float);????

????cudaMalloc(&d_A.elements,?size);????

????cudaMemcpy(d_A.elements,?A.elements,?size,?cudaMemcpyHostToDevice);????

????Matrix?d_B;????

????d_B.width?=?d_B.stride?=?B.width;????

????d_B.height?=?B.height;????

????size?=?B.width?*?B.height?*?sizeof(float);????

????cudaMalloc(&d_B.elements,?size);????

????cudaMemcpy(d_B.elements,?B.elements,?size,?cudaMemcpyHostToDevice);????

????

????//?Allocate?C?in?device?memory????

????Matrix?d_C;????

????d_C.width?=?d_C.stride?=?C.width;????

????d_C.height?=?C.height;????

????size?=?C.width?*?C.height?*?sizeof(float);????

????cudaMalloc(&d_C.elements,?size);????

????

????//?Invoke?kernel????

????dim3?dimBlock(BLOCK_SIZE,?BLOCK_SIZE);????

????dim3?dimGrid(B.width?/?dimBlock.x,?A.height?/?dimBlock.y);????

????MatMulKernel<<<dimGrid,?dimBlock>>>(d_A,?d_B,?d_C);????

????

????//?Read?C?from?device?memory????

????cudaMemcpy(C.elements,?d_C.elements,?size,?cudaMemcpyDeviceToHost);????

????

????//?Free?device?memory????

????cudaFree(d_A.elements);????

????cudaFree(d_B.elements);????

????cudaFree(d_C.elements);????

}????

????

//?Matrix?multiplication?kernel?called?by?MatMul()????

__global__?void?MatMulKernel(Matrix?A,?Matrix?B,?Matrix?C)?{????

????//?Block?row?and?column????

????int?blockRow?=?blockIdx.y;????

????int?blockCol?=?blockIdx.x;????

????

????//?Each?thread?block?computes?one?sub-matrix?Csub?of?C????

????Matrix?Csub?=?GetSubMatrix(C,?blockRow,?blockCol);????

????

????//?Each?thread?computes?one?element?of?Csub????

????//?by?accumulating?results?into?Cvalue????

????float?Cvalue?=?0;????

????

????//?Thread?row?and?column?within?Csub????

????int?row?=?threadIdx.y;????

????int?col?=?threadIdx.x;????

????

????//?Look?over?all?the?sub-matrices?of?A?and?B?that?are?required?to?compute?Csub????

????//?Multiply?each?pair?of?sub-matrices?together?and?accumulate?the?results????

????for?(int?m?=?0;?m?<?(A.width?/?BLOCK_SIZE);?++m)?{????

????????//?Get?sub-matrix?Asub?of?A????

????????Matrix?Asub?=?GetSubMatrix(A,?blockRow,?m);????

????????????

????????//?Get?sub-matrix?Bsub?of?B????

????????Matrix?Bsub?=?GetSubMatrix(B,?m,?blockCol);????

????

????????//?Shared?memory?used?to?store?Asub?and?Bsub?respectively????

????????__shared__?float?As[BLOCK_SIZE][BLOCK_SIZE];????

????????__shared__?float?Bs[BLOCK_SIZE][BLOCK_SIZE];????

????

????????//?Load?Asub?and?Bsub?from?device?memory?to?shared?memory????

????????//?Each?thread?loads?one?element?of?each?sub-matrix????

????????As[row][col]?=?GetElement(Asub,?row,?col);????

????????Bs[row][col]?=?GetElement(Bsub,?row,?col);????

????

????????//?Synchronize?to?make?sure?the?sub-matrices?are?loaded????

????????//?before?starting?the?computation????

????????__syncthreads();????

????

????????//?Multiply?Asub?and?Bsub?together????

????????for?(int?e?=?0;?e?<?BLOCK_SIZE;?++e)????

????????????Cvalue?+=?As[row][e]?*?Bs[e][col];????

????

????????//?Synchronize?to?make?sure?that?the?preceding?computation?is?done?before????

????????//?loading?two?new?sub-matrices?of?A?and?B?in?the?next?iteration????

????????__syncthreads();????

????}????

????

????//?Write?Csub?to?device?memory????

????//?Each?thread?writes?one?element????

????SetElement(Csub,?row,?col,?Cvalue);????

}???

異步并行執(zhí)行

主機和設備間并行執(zhí)行

? ? ? ? 為了支持主機和設備的并行執(zhí)行，CUDA提供了一些異步函數(shù)。異步是指設備在完成功能執(zhí)行之前就將控制權交還給主機線程，以便主機線程繼續(xù)執(zhí)行。這些函數(shù)有： ? ? ? ? 1、內(nèi)涵啟動（Kernel Launches）； ? ? ? ? 2、同一設備內(nèi)存中兩個地址塊之間的數(shù)據(jù)傳輸； ? ? ? ? 3、從主機到設備的小于或等于64KB的一次數(shù)據(jù)塊傳輸； ? ? ? ? 4、使用Async前綴函數(shù)進行的數(shù)據(jù)傳輸； ? ? ? ? 5、內(nèi)存置位函數(shù)調(diào)用（Memory set functions calls）。 程序員可以通過設置環(huán)境變量CUDA_LAUNCH_BLOCKING來開啟或者關閉內(nèi)核啟動（Kernel Launch）的異步功能。但是這一方法僅限于調(diào)試，在任何產(chǎn)品代碼中不應當關閉異步內(nèi)核啟動。 ? ? ? ? 內(nèi)核啟動在下面這些情況下則是同步的： ? ? ? ? 1、應用程序通過調(diào)試器或者內(nèi)存檢查器運行在計算能力為1.x的設備上。 ? ? ? ? 2、硬件計數(shù)器信息正被性能分析器收集。

將內(nèi)核啟動與數(shù)據(jù)傳輸重疊起來

? ? ? ? 對于一些計算能力等于或高于1.1的設備，它們可以將內(nèi)核啟動任務和鎖頁內(nèi)存到設備內(nèi)存的數(shù)據(jù)傳輸任務并行執(zhí)行。應用程序可以檢查設備屬性中的asyncEngineCount項來確定設備是否支持這一功能。當該項值大于0時代表設備支持這一層次的并行。對于計算能力1.x的設備，該功能不支持通過cudaMallocPitch()函數(shù)分配的CUDA數(shù)組或2D數(shù)組。

并行內(nèi)核執(zhí)行

? ? ? ? 一些計算能力2.x或更高的設備可以同時并行執(zhí)行多個內(nèi)核函數(shù)。應用程序可以檢查設備屬性中的concurrentKernels項來確定設備是否支持這一功能，值為1代表支持。運算能力3.5的設備在同一時刻能夠并行執(zhí)行的最大內(nèi)核函數(shù)數(shù)量為32，運算能力小于3.5的硬件則最多支持同時啟動16個內(nèi)核函數(shù)的執(zhí)行。同時需要注意的是，在一個CUDA上下文中的內(nèi)核函數(shù)不能與另一個CUDA上下文中的內(nèi)核函數(shù)同時執(zhí)行。使用很多紋理內(nèi)存或者大量本地內(nèi)存的內(nèi)核函數(shù)也很可能無法與其它內(nèi)核函數(shù)并行執(zhí)行。

并行數(shù)據(jù)傳輸

? ? ? ? 一些計算能力為2.x或更高的設備可以將鎖頁內(nèi)存到設備內(nèi)存的數(shù)據(jù)傳輸和設備內(nèi)存到鎖頁內(nèi)存的數(shù)據(jù)傳輸并行執(zhí)行。應用程序可檢查設備屬性中的asyncEngineCount項來確定這一功能的支持程度，等于2時表示支持。

流（Streams）

? ? ? ? 應用程序通過流來管理并行。一個流是一個順次執(zhí)行的命令序列。不同的流之間并行執(zhí)行，沒有固定的執(zhí)行順序。

1、流的創(chuàng)建與銷毀

? ? ? ? 定義一個流的過程通常包括：創(chuàng)建一個流對象，然后指定它為內(nèi)核啟動或者主機設備間數(shù)據(jù)傳輸?shù)牧鲄?shù)。下面的一段代碼創(chuàng)建了兩個流并且在鎖頁內(nèi)存中分配了一塊float類型的數(shù)組hostPtr：

[cpp] view plaincopy

cudaStream_t?stream[2];??

for?(int?i?=?0;?i?<?2;?++i)??

????cudaStreamCreate(&stream[i]);??

float?*hostPtr;??

cudaMallocHost(&hostPtr,?2?*?size);??

下面的代碼定義了每一個流的行為：從主機端拷貝數(shù)據(jù)到設備端，內(nèi)核啟動，從設備端拷貝數(shù)據(jù)到主機端：
[cpp] view plaincopy

for?(int?i?=?0;?i?<?2;?++i)?{??

????cudaMemcpyAsync(inputDevPtr?+?i?*?size,?hostPtr?+?i?*?size,?size,?cudaMemcpyHostToDevice,?stream[i]);??

????MyKernel<<<100,?512,?0,?stream[i]>>>(outputDevPtr?+?i?*?size,?inputDevPtr?+?i?*?size,?size);??

????cudaMemcpyAsync(hostPtr?+?i?*?size,?outputDevPtr?+?i?*?size,?size,?cudaMemcpyDeviceToHost,?stream[i]);??

}??

這部分代碼中有一點需要注意：為了并行化數(shù)據(jù)拷貝和內(nèi)核執(zhí)行，主機端內(nèi)存必須分配為鎖頁（page-locked）內(nèi)存。

要銷毀一個流需要調(diào)用函數(shù)cudaStreamDestroy()

[cpp] view plaincopy

for?(int?i?=?0;?i?<?2;?++i)??

????cudaStreamDestroy(stream[i]);??

cudaStreamDestroy() 函數(shù)等待之前流中的指令序列運行完成，然后銷毀指定流，將控制權返還給主機端。

2、默認流（Default stream）

? ? ? ? 在內(nèi)核啟動或者數(shù)據(jù)拷貝過程中如果不指定流，或者設置流參數(shù)為0，則相應的指令將會運行在默認流上，它們也因此而順次執(zhí)行。

3、明同步（Explicit Synchronization）

? ? ? ? 在CUDA中有很多種方式可以用來同步流的執(zhí)行：

? ? ? ??cudaDeviceSynchronize()函數(shù)使得主機端線程阻塞直到所有流中的指令執(zhí)行完成。

? ? ? ?cudaStreamSynchronize()函數(shù)將一個流對象作為輸入?yún)?shù)，用以等待指定流中的所有指令執(zhí)行完成。

? ? ? ?cudaStreamWaitEvent()函數(shù)將一個流對象和一個事件作為輸入?yún)?shù)，它將延遲該函數(shù)調(diào)用后在指定流中所有新加入的命令的執(zhí)行直到指定的事件完成為止。流參數(shù)可以為0，在該情形下所有流中的任何新加入的指令都必須等待指定事件的發(fā)生，然后才可以執(zhí)行。

? ? ? ?cudaStreamQuery()函數(shù)為應用程序提供了一個檢測指定流中之前指令是否執(zhí)行完成的方法。

? ? ? ? 為了避免同步帶來的性能下降，所有上述同步函數(shù)最好用于計時目的或者分離錯誤的內(nèi)核執(zhí)行或數(shù)據(jù)拷貝。

4、暗同步（Implicit Synchronization）

? ? ? ? 如果任何一個流中正在執(zhí)行以下操作，那么其它流是不能與其并行運行的：

? ? ? ? a. 分配鎖頁內(nèi)存空間

? ? ? ? b. 設備內(nèi)存分配

? ? ? ? c. 設備內(nèi)存置位

? ? ? ? d. 同一設備兩個不同地址間正在進行數(shù)據(jù)拷貝

? ? ? ? e. 默認流中有指令正在執(zhí)行

? ? ? ? f. L1/shared內(nèi)存配置的轉(zhuǎn)換

? ? ? ? 對于支持并行內(nèi)核執(zhí)行并且計算能力3.0或以下的設備來說，任何一個需要檢查依賴性以確定流內(nèi)核啟動是否完成的操作：

? ? ? ? a. 只有當前CUDA上下文中所有流中所有之前的內(nèi)核啟動之后才能夠啟動執(zhí)行。

? ? ? ? b. 將會阻塞所有當前CUDA上下文中的任意流中新加入的內(nèi)核調(diào)用直到內(nèi)核檢查完成。

? ? ? ? 需要進行依賴性檢查的操作包括執(zhí)行檢查的內(nèi)核啟動所在流中的其它指令以及任何在該流上對cudaStreamQuery()函數(shù)的調(diào)用。因此，應用程序可以遵照以下指導原則來提升潛在并行性：

? ? ? ? （1）所有非依賴操作應當比依賴性操作提前進行

? ? ? ? （2）任何類型的同步越遲越好

5、重疊行為（Overlapping Behavior）

? ? ? ? 兩個流間重疊行為的數(shù)量取決于以下幾個因素：

? ? ? ? （1）每個流中命令發(fā)出的次序

? ? ? ? （2）設備是否支持內(nèi)核啟動與數(shù)據(jù)傳輸并行

? ? ? ? （3）設備是否支持多內(nèi)核并行啟動

? ? ? ? （4）設備是否支持多數(shù)據(jù)傳輸并行

? ? ? ? 例如，在不支持并行數(shù)據(jù)傳輸?shù)脑O備上，“流的創(chuàng)建與銷毀”章節(jié)中代碼樣例中的操作就不能并行，因為在stream[0]中發(fā)出設備端到主機端的數(shù)據(jù)拷貝后，stream[1]又發(fā)出主機端到設備端的數(shù)據(jù)拷貝命令，這兩個命令式不能重疊執(zhí)行的。假設設備支持數(shù)據(jù)傳輸與內(nèi)核啟動并行，那么如下代碼：

[cpp] view plaincopy

for?(int?i?=?0;?i?<?2;?++i)??

????cudaMemcpyAsync(inputDevPtr?+?i?*?size,?hostPtr?+?i?*?size,?size,?cudaMemcpyHostToDevice,?stream[i]);??

for?(int?i?=?0;?i?<?2;?++i)??

????MyKernel<<<100,?512,?0,?stream[i]>>>(outputDevPtr?+?i?*?size,?inputDevPtr?+?i?*?size,?size);??

for?(int?i?=?0;?i?<?2;?++i)??

????cudaMemcpyAsync(hostPtr?+?i?*?size,?outputDevPtr?+?i?*?size,?size,?cudaMemcpyDeviceToHost,?stream[i]);??

可將stream[0]的內(nèi)核啟動和stream[1]從主機端到設備端的數(shù)據(jù)拷貝重疊起來并行執(zhí)行。

6、回調(diào)函數(shù)

? ? ? ? CUDA運行時提供了cudaStreamAddCallback()函數(shù)以在流中的任意位置插入一個回調(diào)函數(shù)點。回調(diào)函數(shù)運行于主機端，如果在默認流中插入回調(diào)函數(shù)，那么它將等待所有其它流中的命令執(zhí)行完成之后才會開始執(zhí)行。

? ? ? ? 下面的代碼展示了回調(diào)函數(shù)技術的應用：

[cpp] view plaincopy

void?CUDART_CB?MyCallback(cudaStream_t?stream,?cudaError_t?status,?void?**data)?{??

????printf("Inside?callback?%d\n",?(int)data);??

}??

...??

for?(int?i?=?0;?i?<?2;?++i)?{??

????cudaMemcpyAsync(devPtrIn[i],?hostPtr[i],?size,?cudaMemcpyHostToDevice,?stream[i]);??

????MyKernel<<<100,?512,?0,?stream[i]>>>(devPtrOut[i],?devPtrIn[i],?size);??

????cudaMemcpyAsync(hostPtr[i],?devPtrOut[i],?size,?cudaMemcpyDeviceToHost,?stream[i]);??

????cudaStreamAddCallback(stream[i],?MyCallback,?(void**)i,?0);??

}??

上面的代碼定義了兩個流的操作，每個流都完成一次主機端到設備端的數(shù)據(jù)拷貝，一次內(nèi)核啟動，一次設備端到主機端的數(shù)據(jù)拷貝，最后增加了一個加入回調(diào)函數(shù)的操作。當設備端代碼運行到回調(diào)函數(shù)點的時候，設備將控制權交還給主機端，主機端運行完成以后再將控制權返還給設備端，然后設備端繼續(xù)運行。

? ? ? ? 值得注意的是，在一個回調(diào)函數(shù)中，一定不能進行任何CUDA API的調(diào)用，直接的或者間接的都是不可以的。

發(fā)現(xiàn)問題的諸位讀者不吝賜教。

1. 什么是CUDA？

CUDA全稱是Compute Unified Device Architecture，中文名稱即統(tǒng)一計算設備架構，它是NVIDIA公司提出了一種通用的并行計算平臺和編程模型。使用CUDA，我們可以開發(fā)出同時在CPU和GPU上運行的通用計算程序，更加高效地利用現(xiàn)有硬件進行計算。為了簡化并行計算學習，CUDA為程序員提供了一個類C語言的開發(fā)環(huán)境以及一些其它的如FORTRAN、DirectCOmpute、OpenACC的高級語言/編程接口來開發(fā)CUDA程序。

2. CUDA編程模型如何擴展？

我們知道，不同的GPU擁有不同的核心數(shù)目，在核心較多的系統(tǒng)上CUDA程序運行的時間較短，而在核心較少的系統(tǒng)上CUDA程序的執(zhí)行時間較多。那么，CUDA是如何做到的呢？ 并行編程的中心思想是分而治之：將大問題劃分為一些小問題，再把這些小問題交給相應的處理單元并行地進行處理。在CUDA中，這一思想便體現(xiàn)在它的具有兩個層次的問題劃分模型。一個問題可以首先被粗粒度地劃分為若干較小的子問題，CUDA使用被稱為塊（Block）的單元來處理它們，每個塊都由一些CUDA線程組成，線程是CUDA中最小的處理單元，將這些較小的子問題進一步劃分為若干更小的細粒度的問題，我們便可以使用線程來解決這些問題了。對于一個普通的NVIDIA GPU，其CUDA線程數(shù)目通常能達到數(shù)千個甚至更多，因此，這樣的問題劃分模型便可以成倍地提升計算機的運算性能。 GPU是由多個流水多處理器構成的，流水處理器以塊（Block）為基本調(diào)度單元，因此，對于流水處理器較多的GPU，它一次可以處理的塊（Block）更多，從而運算速度更快，時間更短。而反之對于流水處理器較少的GPU，其運算速度便會較慢。這一原理可以通過下圖形象地看出來：

3. CUDA基本概念

本節(jié)將介紹CUDA的一些基本的編程概念，該節(jié)用到的例子來自于CUDA Sample中的VectorAdd項目。

3.1 內(nèi)核（Kernels）

CUDA C是C語言的一個擴展，它允許程序員定義一種被稱為內(nèi)核函數(shù)（Kernel Functions）的C函數(shù)，內(nèi)核函數(shù)運行在GPU上，一旦啟動，CUDA中的每一個線程都將會同時并行地執(zhí)行內(nèi)核函數(shù)中的代碼。

內(nèi)核函數(shù)使用關鍵字__global__來聲明，運行該函數(shù)的CUDA線程數(shù)則通過<<<...>>>執(zhí)行配置語法來設置。（參見章節(jié)"C語言擴展"），每一個執(zhí)行內(nèi)核函數(shù)的線程都由一個唯一的線程ID，這一ID可以通過在內(nèi)核函數(shù)中訪問threadIdx變量來得到。

下面通過一些示例代碼來展示剛剛提到的這些概念該如何應用在編程中：

[cpp] view plaincopy

//?Kernel?definition????

__global__?void?VecAdd(float*?A,?float*?B,?float*?C)?{????

????int?i?=?threadIdx.x;????

????C[i]?=?A[i]?+?B[i];????

}????

????

int?main()?{????

????...????

????//?Kernel?invocation?with?N?threads????

????VecAdd<<<1,?N>>>(A,?B,?C);????

???...????

}???

在上面的代碼中，N個線程將會并行地同時執(zhí)行加法運算。

3.2 線程層次（Thread Hierarchy）

CUDA的每一個線程都有其線程ID，線程的ID信息由變量threadIdx給出。threadIdx是CUDA C語言的內(nèi)建變量，通常它用一個三維數(shù)組來表示。使用三維數(shù)組的方便之處在于可以很方便地表示一維、二維和三維線程索引，進而方便地表示一維、二維和三維線程塊（thread block）。這樣，無論是數(shù)組、矩陣還是體積的計算，都可以很容易地使用CUDA進行運算。 線程的索引與線程ID之間存在著直接的換算關系，對于一個索引為(x, y, z)的線程來說： ? ? ? 1、如果線程塊（block）是一維的，則線程ID = x ? ? ? 2、如果線程塊是二維的，假設塊尺寸為（Dx，Dy），那么線程ID = x + y * Dx ? ? ? 3、如果線程塊是三維的，設其尺寸為（Dx，Dy，Dz），那么線程ID = x + y * Dx + z * Dx * Dy ? ? ? 下面的例子展示了兩個NxN矩陣相加的CUDA實現(xiàn)： [cpp] view plaincopy

//?Kernel?definition????

__global__?void?MatAdd(float?A[N][N],?float?B[N][N],?float?C[N][N])?{????

????int?i?=?threadIdx.x;????

????int?j?=?threadIdx.y;????

????C[i][j]?=?A[i][j]?+?B[i][j];????

}????

????

int?main()?{????

????...????

????//?Kernel?invocation?with?one?block?of?N?*?N?*?1?threads????

????int?numBlocks?=?1;????

????dim3?threadsPerBlock(N,?N);????

????MatAdd<<<numBlocks,?threadsPerBlock>>>(A,?B,?C);????

????...????

}????

每個線程塊（block）中的線程數(shù)量是有限制的，因為依據(jù)前面所說，同一線程塊（block）中的所有線程都會被分配到同一個處理器核上運行，共享有限的存儲資源，因此對于當前的GPU，線程塊所能包含的最大線程數(shù)目為1024。 上面的例子中numBlocks代表線程塊的數(shù)量，這里的值為1。在一般的CUDA程序中，這個值通常大于1，也就是說將會有多個線程塊被分配到多個處理器核中同時進行處理，這樣就大大提高了程序的并行性。
在CUDA中，線程塊包含在線程格（grid）當中，線程格可以是一維、二維或者三維的，線程格的尺寸一般根據(jù)待處理數(shù)據(jù)的規(guī)模或者處理器的數(shù)量來指定。線程格中所包含的線程塊數(shù)目通常遠遠大于GPU處理器核心的數(shù)目。下圖展示了線程格（grid）、線程塊（block）以及線程（thread）之間的關系：
內(nèi)核函數(shù)的調(diào)用可以簡化為kernel<<<A，B>>>(parameters)，在尖括號中，A代表線程格（grid）的尺寸，它可以是三維的，用類型dim3表示，也可以是一維的，用int類型表示。B代表線程塊（block）的尺寸，它與A類似，也可分別用dim3或int類型表示。 在內(nèi)核函數(shù)內(nèi)部，CUDA為我們內(nèi)建了一些變量用于訪問線程格、線程塊的尺寸和索引等信息，它們是： 1. gridDim：代表線程格（grid）的尺寸，gridDim.x為x軸尺寸，gridDim.y、gridDim.z類似。拿上圖來說，它的gridDim.x = 3，gridDim.y = 2，gridDim.z = 1。 2. blockIdx：代表線程塊（block）在線程格（grid）中的索引值，拿上圖來說，Block（1,1）的索引值為：blockIdx.x = 1，blockIdx.y = 1。 3. blockDim：代表線程塊（block）的尺寸，blockDIm.x為x軸尺寸，其它依此類推。拿上圖來說，注意到Block（1,1）包含了4 * 3個線程，因此blockDim.x = 4, blockDim.y = 3。 4. threadIdx：線程索引，前面章節(jié)已經(jīng)詳細探討過了，這里不再贅述。 明白了這些變量的含義，那么下面的矩陣加法程序便不難理解了： [cpp] view plaincopy

//?Kernel?definition????

__global__?void?MatAdd(float?A[N][N],?float?B[N][N],?float?C[N][N])?{????

????int?i?=?blockIdx.x?*?blockDim.x?+?threadIdx.x;????

????int?j?=?blockIdx.y?*?blockDim.y?+?threadIdx.y;????

????if?(i?<?N?&&?j?<?N)????

????????C[i][j]?=?A[i][j]?+?B[i][j];????

}????

????

int?main()?{????

????...????

????//?Kernel?invocation????

????dim3?threadsPerBlock(16,?16);????

????dim3?numBlocks(N?/?threadsPerBlock.x,?N?/?threadsPerBlock.y);????

????MatAdd<<<numBlocks,?threadsPerBlock>>>(A,?B,?C);????

????...????

}??

在上面的程序中，線程塊（block）的尺寸是16x16，這是CUDA編程中一個非常普遍的選擇。線程格（grid）包含了足夠多的線程塊（block）來進行計算。 線程塊（block）是獨立執(zhí)行的，在執(zhí)行的過程中線程塊之間互不干擾，因此它們的執(zhí)行順序是隨機的。 同一線程塊中的線程可以通過訪問共享內(nèi)存（shared memory）或者通過同步函數(shù)__syncthreads()來協(xié)調(diào)合作。這些概念將在以后的章節(jié)中詳細解釋。

3.3 內(nèi)存層次（Memory Hierarchy）

在GPU上CUDA線程可以訪問到的存儲資源有很多，每個CUDA線程擁有獨立的本地內(nèi)存（local Memory）；每一個線程塊（block）都有其獨立的共享內(nèi)存（shared memory），共享內(nèi)存對于線程塊中的每個線程都是可見的，它與線程塊具有相同的生存時間；同時，還有一片稱為全局內(nèi)存（global memory）的區(qū)域?qū)λ械腃UDA線程都是可訪問的。

除了上述三種存儲資源以外，CUDA還提供了兩種只讀內(nèi)存空間：常量內(nèi)存（constant memory）和紋理內(nèi)存（texture memory），同全局內(nèi)存類似，所有的CUDA線程都可以訪問它們。對于一些特殊格式的數(shù)據(jù)，紋理內(nèi)存提供多種尋址模式以及數(shù)據(jù)過濾方法來操作內(nèi)存。這兩類存儲資源主要用于一些特殊的內(nèi)存使用場合。

一個程序啟動內(nèi)核函數(shù)以后，全局內(nèi)存、常量內(nèi)存以及紋理內(nèi)存將會一直存在直到該程序結束。下面是CUDA的內(nèi)存層次圖：

3.4 異構編程（Heterogeneous Programming）

CUDA的異構編程模型假定CUDA線程都運行在一個可被看做CPU協(xié)處理器的芯片上，這就使得CUDA內(nèi)核函數(shù)可以和CPU端C程序的運行并行運行，從而加快程序的運行效率。為了達到這個效果，CUDA程序需要管理兩大塊由DRAM構成的內(nèi)存區(qū)域：CPU端可以訪問到的主機內(nèi)存（host memory）以及GPU端供CUDA內(nèi)核訪問到的設備內(nèi)存（device memory），設備內(nèi)存主要由全局內(nèi)存、常量內(nèi)存以及紋理內(nèi)存構成。現(xiàn)在，CUDA程序的運行機制便很明了了：CPU端代碼生成原始數(shù)據(jù)，通過CUDA運行時函數(shù)庫將這些原始數(shù)據(jù)傳輸?shù)紾PU上，在CPU端啟動CUDA內(nèi)核函數(shù)進行運算，然后將運算結果從設備端傳輸?shù)街鳈C端，計算任務便完成了。

4.?CUDA C語言編程接口

異構程序設計跟傳統(tǒng)的串行程序設計差別是很大的，學習起來也是非常不容易的。NVIDIA非常夠意思，為了簡化CUDA的學習曲線，它采用了絕大多數(shù)程序員都熟悉的C語言作為其根基，CUDA C是NVIDIA為程序員提供的一類編程接口，它實際上是一個C語言的擴展，在C的基礎上增加了一些新的語法和變量，并且提供了功能豐富的庫函數(shù)，方便程序員使用GPU進行異構計算。 除了前面章節(jié)提到的CUDA最基本、最核心的概念以外，CUDA C呈現(xiàn)給程序員的接口主要由兩大類API構成，它們分別是CUDA Runtime API和CUDA Driver API，Runtime API實際上是對于Driver API的封裝，其目的自然是方便程序員的代碼編寫工作。Driver API為用戶提供了更細一層的控制手段，通過它可以控制諸如CUDA Contexts（一種類似主機進程的概念）以及CUDA Modules（類似主機動態(tài)加載庫的概念）等更加底層的CUDA模塊。

4.1 NVCC編譯器

任何一種程序設計語言都需要相應的編譯器將其編譯為二進制代碼，進而在目標機器上得到執(zhí)行。對于異構計算而言，這一過程與傳統(tǒng)程序設計語言是有一些區(qū)別的。為什么？因為CUDA它本質(zhì)上不是一種語言，而是一種異構計算的編程模型，使用CUDA C寫出的代碼需要在兩種體系結構完全不同的設備上執(zhí)行：1、CPU；2、GPU。因此，CUDA C的編譯器所做的工作就有點略多了。一方面，它需要將源代碼中運行在GPU端的代碼編譯得到能在CUDA設備上運行的二進制程序。另一方面，它也需要將源代碼中運行在CPU端的程序編譯得到能在主機CPU上運行的二進制程序。最后，它需要把這兩部分有機地結合起來，使得兩部分代碼能夠協(xié)調(diào)運行。 CUDA C為我們提供了這樣的編譯器，它便是NVCC。嚴格意義上來講，NVCC并不能稱作編譯器，NVIDIA稱其為編譯器驅(qū)動（Compiler Driver），本節(jié)我們暫且使用編譯器來描述NVCC。使用nvcc命令行工具我們可以簡化CUDA程序的編譯過程，NVCC編譯器的工作過程主要可以劃分為兩個階段：離線編譯（Offline Compilation）和即時編譯（Just-in-Time Compilation）。 離線編譯（Offline Compilation） ? ? 下面這幅圖簡單說明了離線編譯的過程：

? ? ? ? 在CUDA源代碼中，既包含在GPU設備上執(zhí)行的代碼，也包括在主機CPU上執(zhí)行的代碼。因此，NVCC的第一步工作便是將二者分離開來，這一過程結束之后： ? ? ? ? 1. 運行于設備端的代碼將被NVCC工具編譯為PTX代碼（GPU的匯編代碼）或者cubin對象（二進制GPU代碼）； ? ? ? ? 2. 運行于主機端的代碼將被NVCC工具改寫，將其中的內(nèi)核啟動語法（如<<<...>>>）改寫為一系列的CUDA Runtime函數(shù)，并利用外部編譯工具（gcc for linux，或者vc compiler for windows）來編譯這部分代碼，以得到運行于CPU上的可執(zhí)行程序。 ? ? ? ? 完事之后，NVCC將自動把輸出的兩個二進制文件鏈接起來，得到異構程序的二進制代碼。 即時編譯（Just-in-time Compile） ? ? ? ? 任何在運行時被CUDA程序加載的PTX代碼都會被顯卡的驅(qū)動程序進一步編譯成設備相關的二進制可執(zhí)行代碼。這一過程被稱作即時編譯（just-in-time compilation）。即時編譯增加了程序的裝載時間，但是也使得編譯好的程序可以從新的顯卡驅(qū)動中獲得性能提升。同時到目前為止，這一方法是保證編譯好的程序在還未問世的GPU上運行的唯一解決方案。 ? ? ? ? 在即時編譯的過程中，顯卡驅(qū)動將會自動緩存PTX代碼的編譯結果，以避免多次調(diào)用同一程序帶來的重復編譯開銷。NVIDIA把這部分緩存稱作計算緩存（compute cache），當顯卡驅(qū)動升級時，這部分緩存將會自動清空，以使得程序能夠自動獲得新驅(qū)動為即時編譯過程帶來的性能提升。 ? ? ? ? 有一些環(huán)境變量可以用來控制即時編譯過程： ? ? ? ? 1. 設置CUDA_CACHE_DISABLE為1將會關閉緩存功能 ? ? ? ? 2. CUDA_CACHE_MAXSIZE變量用于指定計算緩存的字節(jié)大小，默認情況下它的值是32MB，它最大可以被設置為4GB。任何大于緩存最大值得二進制代碼將不會被緩存。在需要的情況下，一些舊的二進制代碼可能被丟棄以騰出空間緩存新的二進制代碼。 ? ? ? ? 3. CUDA_CACHE_PATH變量用于指定計算緩存的存儲目錄地址，它的缺省值如下：
? ? ? ? 4. 設置CUDA_FORCE_PTX_JIT為1會強制顯卡驅(qū)動忽略應用程序中的二進制代碼并且即時編譯程序中的嵌入PTX代碼。如果一個內(nèi)核函數(shù)沒有嵌入的PTX代碼，那么它將會裝載失敗。該變量可以用來確認程序中存在嵌入的PTX代碼。同時，使用即時編譯（just-in-time Compilation）技術也可確保程序的向前兼容性。

4.2 兼容性

? ? ??1、二進制兼容性 ? ? ? 二進制代碼是設備相關的，使用NVCC編譯器編譯時，若指定-code選項，則會編譯產(chǎn)生目標設備的二進制cubin對象。例如，編譯時使用-code=sm_13會產(chǎn)生適用于計算能力1.3的二進制代碼。二進制代碼在CUDA計算設備上具有小版本的向前兼容性，但是在大版本上不具備兼容性。也就是說，對于計算能力X.y的硬件，使用-code=sm_Xy編譯后，程序能夠運行于計算能力X.z（其中z>=y）的硬件上，但不能運行在計算能力M.n（M!=X）的硬件上。 ? ? ??2、PTX代碼兼容性 ? ? ? 不同計算能力的設備所支持的PTX指令條數(shù)是不同的，一些PTX指令只在擁有較高計算能力的設備上被支持。例如，全局內(nèi)存（global Memory）的原子操作指令只能用于計算能力不小于1.1的設備；雙精度浮點運算指令只能用于計算能力不小于1.3的設備。在將C語言編譯為PTX代碼時，NVCC使用-arch編譯選項指定PTX代碼目標設備的計算能力。因此，要想使用雙精度運算，編譯時必須使用選項-arch=sm_13（或使用更高的計算能力），否則NVCC會自動將雙精度操作降級為單精度操作。 ? ? ? 為某一特定設備產(chǎn)生的PTX代碼，在運行時總是能夠被具有更高計算能力的設備JIT編譯為可執(zhí)行的二進制代碼。 ? ? ? 3、應用程序兼容性 ? ? ? 執(zhí)行CUDA程序有兩種方式，一種是直接加載編譯好的CUDA二進制代碼運行，另一種是首先加載程序中的PTX代碼，再執(zhí)行JIT編譯得到二進制的設備可執(zhí)行文件，然后運行。特別需要注意的是，為了讓程序運行具有更高計算能力的未來設備上，必須讓程序加載PTX代碼。 ? ? ? 事實上，在一個CUDA C程序中可以嵌入不止一個版本的PTX/二進制代碼。那么，具體執(zhí)行時哪一個版本的PTX或者二進制代碼會得到執(zhí)行呢？答案是：最兼容的那個版本。例如編譯一個名為x.cu的CUDA源代碼：
將會產(chǎn)生兼容計算能力1.1硬件的二進制代碼（第一排的-gencode選項）以及兼容計算能力1.1設備的PTX和二進制代碼，這些代碼都將會嵌入到編譯后的目標文件中。 ? ? ? 主機端將會產(chǎn)生一些額外的代碼，在程序運行時，這些代碼會自動決定裝載哪一個版本的代碼來執(zhí)行。對于上面的例子：

• 計算能力1.0的設備運行該程序?qū)b載1.0版本的二進制代碼
• 計算能力1.1、1.2或者1.3的設備運行該程序?qū)b載1.1版本的二進制代碼
• 計算能力2.0或者更高的設備運行該程序?qū)b載1.1版本的PTX代碼進而對其進行JIT編譯得到相應設備的二進制代碼

? ? ??同時，x.cu還可以在程序中使用一些特殊的宏來改變不同設備的代碼執(zhí)行路徑。例如，對于計算能力1.1的設備而言，宏__CUDA_ARCH__等于110，在程序中可以對該宏的值進行判斷，然后分支執(zhí)行程序。 ? ? ? NVCC用戶手冊列出了很多-arch，-code和-gencode等編譯選項的簡化書寫形式。例如，-arch=sm_13就是-arch=compute_13 -code=compute13, sm_13的簡化形式。更多詳盡的內(nèi)容請參閱該手冊。 ? ? ? 4、C/C++兼容性 ? ? ? NVCC編譯器前端使用C++語法啊規(guī)則來處理CUDA源文件。在主機端，CUDA支持完整的C++語法；而在設備端，只有部分C++語法是被支持的。這方面更為詳盡的討論請參見《CUDA C程序設計指南》的C/C++語言支持章節(jié)。 ? ? ? 5、64位兼容性 ? ? ? 64位版本的nvcc編譯器將設備代碼編譯為64位模式，即指針是64位的。運行64位設備代碼的先決條件是主機端代碼必須也使用64位模式進行編譯。同樣，32位版本的nvcc將設備代碼編譯為32位模式，這些代碼也必須與相應的32位主機端代碼相配合方能運行。 ? ? ? 32位nvcc編譯器可以使用-m64編譯選項將設備代碼編譯為64位模式。同時64位nvcc編譯器也可使用-m32編譯選項將設備代碼編譯為32位模式。

4.3 CUDA C Runtime

? ? ? ? CUDA C Runtime使用cudart動態(tài)鏈接庫實現(xiàn)（cudart.dll或者cudart.so），運行時中所有的入口函數(shù)都以cuda為前綴。

4.3.1 初始化

CUDA C Runtime函數(shù)庫沒有明確的初始化函數(shù)，在程序第一次調(diào)用Runtime庫函數(shù)時它會自動初始化。因此，在記錄Runtime函數(shù)調(diào)用時間和理解程序中第一個Runtime調(diào)用返回的錯誤代碼時，需要將初始化考慮在內(nèi)。 ? ? ? ? 在初始化期間，Runtime將會為系統(tǒng)中每一個設備創(chuàng)建一個CUDA上下文（類似CPU中進程的數(shù)據(jù)結構），這個上下文是設備的基本上下文，它被程序中所有的主機線程所共享。創(chuàng)建過程在后臺運行，并且，Runtime將隱藏基本上下文使之對Runtime API這一層次的程序員不可見。 ? ? ? ? 當一個主機線程調(diào)用cudaDeviceReset()函數(shù)時，它將會銷毀線程當前控制設備的基本上下文。也就是說，當線程下一次調(diào)用runtime函數(shù)時將會重啟初始化，一個新的CUDA基本上下文將被創(chuàng)建出來。

4.3.2 設備內(nèi)存

? ? ? ? 正如前面異構計算章節(jié)所講，CUDA編程模型假定系統(tǒng)是由主機和設備構成的，它們分別具有自己獨立的內(nèi)存空間。Runtime負責設備內(nèi)存的分配，回收，拷貝以及在主機和設備間傳輸數(shù)據(jù)的工作。

? ? ? ? 設備內(nèi)存可以有兩種分配方式：線性內(nèi)存或者CUDA數(shù)組

? ? ? ? CUDA數(shù)組是一塊不透明的內(nèi)存空間，它主要被優(yōu)化用于紋理存取。

? ? ? ? 線性內(nèi)存空間與平時我們訪問的內(nèi)存類似，對于計算能力1.x的設備來說，它存在于一個32位的地址空間。對于更高計算能力的設備而言，它存在于一個40位的地址空間中。因此，單獨分配的實體可以使用指針來相互應用。

? ? ? ? 我們通常使用cudaMalloc()函數(shù)分配線性內(nèi)存空間，使用cudaFree()函數(shù)釋放線性內(nèi)存空間，使用cudaMemcpy()函數(shù)在主機和設備之間傳輸數(shù)據(jù)。下面是CUDA Vector Add代碼示例的一些片段：

[cpp] view plaincopy

//?Device?code????

__global__?void?VecAdd(float?*A,?float?*B,?float?*C,?int?N)?{????

????int?i?=?blockDim.x?*?blockIdx.x?+?threadIdx.x;????

????if?(i?<?N)????

????????C[i]?=?A[i]?+?B[i];????

}????

????

//?Host?code????

int?main()?{????

????int?N?=?...;????

????size_t?size?=?N?*?sizeof(float);????

????

????//?Allocate?input?vectors?h_A?and?h_B?in?host?memory????

????float?*h_A?=?(float*)malloc(size);????

????float?*h_B?=?(float*)malloc(size);????

????

????//?Initialize?input?vectors????

????...????

?????

????//?Allocate?vectors?in?device?memory????

????float?*d_A,?*d_B,?*d_C;????

????cudaMalloc(&d_A,?size);????

????cudaMalloc(&d_B,?size);????

????cudaMalloc(&d_C,?size);????

????

????//?Copy?vectors?from?host?memory?to?device?memory????

????cudaMemcpy(d_A,?h_A,?size,?cudaMemcpyHostToDevice);????

????cudaMemcpy(d_B,?h_B,?size,?cudaMemcpyHostToDevice);????

????

????//?Invoke?kernel????

????int?threadsPerBlock?=?256;????

????int?blocksPerGrid?=?(N?+threadsPerBlock?-?1)?/?threadsPerBlock;????

????VecAdd<<<blocksPerGrid,?threadsPerBlock>>>(d_A,?d_B,?d_C,?N);????

????

????//?Copy?result?from?device?memory?to?host?Memory????

????cudaMemcpy(h_C,?d_C,?size,?cudaMemcpyDeviceToHost);????

????

????//?Free?device?memory????

????cudaFree(d_A);????

????cudaFree(d_B);????

????cudaFree(d_C);????

????

????//?Free?host?memory????

????...????

}???

片段展示了設備內(nèi)存的分配，傳輸以及回收過程。

? ? ? ? 除了上面展示的方法，我們還可以使用cudaMallocPitch()和cudaMalloc3D()函數(shù)來分配線性內(nèi)存。這些函數(shù)能夠確保分配的內(nèi)存滿足設備內(nèi)存訪問的對齊要求，對于行地址的訪問以及多維數(shù)組間的數(shù)據(jù)傳輸提供高性能保證，因此非常適合對于二維和三維數(shù)組內(nèi)存空間的分配。下面的代碼片段展示了分配和使用尺寸為width x height的二維數(shù)組的技術：

[cpp] view plaincopy

//?Host?code????

int?width?=?64,?height?=?64;????

float?*devPtr;????

size_t?pitch;????

cudaMallocPitch(&devPtr,?&pitch,?width?*?sizeof(float),?height);????

MyKernel<<<100,?512>>>(devPtr,?pitch,?width,?height);????

????

//?Device?code????

__global__?void?MyKernel(float*?devPtr,?size_t?pitch,?int?width,?int?height)?{????

????for?(int?r?=?0;?r?<?height;?++r)?{????

????????float*?row?=?(float*)((char*)devPtr?+?r?*?pitch);????

????????for?(int?c?=?0;?c?<?width;?++c)?{????

????????????float?element?=?row[c];????

????????}????

????}????

}????

? ? ? ? 下面的代碼片段展示了一個尺寸為width x height x depth的三維數(shù)組的分配和使用方法：

[cpp] view plaincopy

//?Host?code????

int?width?=?64,?height?=?64,?depth?=?64;????

cudaExtent?extent?=?make_cudaExtent(width?*?sizeof(float),?height,?depth);????

cudaPitchedPtr?devPitchedPtr;????

cudaMalloc3D(&devPitchedPtr,?extent);????

MyKernel<<<100,?512>>>(devPitchedPtr,?width,?height,?depth);????

????

//?Device?code????

__global__?void?MyKernel(cudaPitchedPtr?devPitchedPtr,?int?width,?int?height,?int?depth)?{????

????char*?devPtr?=?devPitchedPtr.ptr;????

????size_t?pitch?=?devPitchedPtr.pitch;????

????size_t?slicePitch?=?pitch?*?height;????

????for?(int?z?=?0;?z?<?depth;?++z)?{????

????????char*?slice?=?devPtr?+?z?*?slicePitch;????

????????for?(int?y?=?0;?y?<?height;?++y)?{????

????????????float*?row?=?(float*)(slice?+?y?*?pitch);????

????????????for?(int?x?=?0;?x?<?width;?++x)????

????????????????float?element?=?row[x];????

????????}????

????}????

}????

更多詳細的內(nèi)容請查閱參考手冊。

? ? ? ? 下面的代碼示例展示了多種使用Runtime API訪問全局變量的技術：

[cpp] view plaincopy

__constant__?float?constData[256];????

float?data[256];????

cudaMemcpyToSymbol(constData,?data,?sizeof(data));????

cudaMemcpyFromSymbol(data,?constData,?sizeof(data));????

????

__device__?float?devData;????

float?value?=?3.14f;????

cudaMemcpyToSymbol(devData,?&value,?sizeof(float));????

????

__device__?float*?devPointer;????

float*?ptr;????

cudaMalloc(&ptr,?256?*?sizeof(float));????

cudaMemcpyToSymbol(devPointer,?&ptr,?sizeof(ptr));????

使用cudaGetSymbolAddress()函數(shù)可以獲得被聲明存儲在全局內(nèi)存中的變量地址。為了獲得分配內(nèi)存的大小，可以使用cudaGetSymbolSize()函數(shù)。

4.3 CUDA C Runtime

4.3.3 共享內(nèi)存（Shared Memory）

? ? ? ? 共享內(nèi)存是CUDA設備中非常重要的一個存儲區(qū)域，有效地使用共享內(nèi)存可以充分利用CUDA設備的潛能，極大提升程序性能。那么，共享內(nèi)存有哪些特點呢？ ? ? ? ? 1、共享內(nèi)存（shared Memory）是集成在GPU處理器芯片上的（on-chip），因此相比于存在于顯存顆粒中的全局內(nèi)存（global Memory）和本地內(nèi)存（local Memory），它具有更高的傳輸帶寬，一般情況下，共享內(nèi)存的帶寬大約是全局內(nèi)存帶寬的7-10倍。 ? ? ? ? 2、共享內(nèi)存的容量很小。根據(jù)NVIDIA官方文檔的說法，在計算能力1.x的設備中，每一個流多處理器（Streaming Multiprocessor）上的共享內(nèi)存容量為16KB。對于計算能力2.x、3.0及3.5的設備該參數(shù)為48KB。因此共享內(nèi)存是稀有資源。 ? ? ? ? 3、共享內(nèi)存在物理上被劃分為很多塊，每一塊被稱為一個存儲體（bank）。在同一時刻，CUDA設備可以同時訪問多個存儲體。因此，如果一次針對共享內(nèi)存的訪存操作需要讀取n個地址，而這n個地址恰好分布在n個不同的存儲體（bank）中，那么只需要一個存取周期就可以完成n個地址的訪存任務了。對于計算能力1.x的設備，共享內(nèi)存被平均劃分為16個存儲體。而對于計算能力2.x、3.0及3.5的設備此參數(shù)為32。在共享內(nèi)存中，相鄰兩塊32bit的數(shù)據(jù)分別屬于相鄰的兩個存儲體。存儲體每兩個時鐘周期可以傳輸32位數(shù)據(jù)。 ? ? ? ? 4、共享內(nèi)存既可以靜態(tài)分配，也可以動態(tài)分配。 ? ? ? ? 從共享內(nèi)存的這些特點中我們可以看出，它實際上相當于一個程序員可以操控的緩存（cache），下面，我們使用矩陣乘法的例子來說明如何有效使用共享內(nèi)存。 ? ? ? ? 首先，我們使用最直觀的方法來完成矩陣乘法C = A x B：讀取A的每一行和B的每一列，順次完成計算任務。矩陣乘法的示意圖如下所示：

下面是矩陣乘法的CUDA C主要實現(xiàn)代碼： [cpp] view plaincopy

//?Matrices?are?stored?in?row-major?order:????

//?M(row,?col)?=?*(M.elements?+?row?*?M.width?+?col)????

typedef?struct?{????

????int?width;????

????int?height;????

????float?*elements;????

}?Matrix;????

????

//?Thread?block?size????

#define?BLOCK_SIZE?16????

????

//?Forward?declaration?of?the?matrix?multiplication?kernel????

__global__?void?MatMulKernel(const?Matrix,?const?Matrix,?Matrix);????

????

//?Matrix?multiplication?-?Host?code????

//?Matrix?dimensions?are?assumed?to?be?multiples?of?BLOCK_SIZE????

void?MatMul(const?Matrix?A,?const?Matrix?B,?Matrix?C)?{????

????//?Load?A?and?B?to?device?memory????

????Matrix?d_A;????

????d_A.width?=?A.width;?d_A.height?=?A.height;????

????size_t?size?=?A.width?*?A.height?*?sizeof(float);????

????cudaMalloc(&d_A.elements,?size);????

????cudaMemcpy(d_A.elements,?A.elements,?size,?cudaMemcpyHostToDevice);????

????Matrix?d_B;????

????d_B.width?=?B.width;?d_B.height?=?B.height;????

????size?=?B.width?*?B.height?*?sizeof(float);????

????cudaMalloc(&d_B.elements,?size);????

????cudaMemcpy(d_B.elements,?B.elements,?size,?cudaMemcpyHostToDevice);????

????

????//?Allocate?C?in?device?memory????

????Matrix?d_C;????

????d_C.width?=?C.width;?d_C.height?=?C.height;????

????size?=?C.width?*?C.height?*?sizeof(float);????

????cudaMalloc(&d_C.elements,?size);????

????

????//?Invoke?kernel????

????dim3?dimBlock(BLOCK_SIZE,?BLOCK_SIZE);????

????dim3?dimGrid(B.width?/?dimBlock.x,?A.height?/?dimBlock.y);????

????MatMulKernel<<<dimGrid,?dimBlock>>>(d_A,?d_B,?d_C);????

????

????//?Read?C?from?device?memory????

????cudaMemcpy(C.elements,?d_c.elements,?size,?cudaMemcpyDeviceToHost);????

????

????//?Free?device?memory????

????cudaFree(d_A.elements);????

????cudaFree(d_B.elements);????

????cudaFree(d_C.elements);????

}????

????

//?Matrix?multiplication?kernel?called?by?MatMul()????

__global__?void?MatMulKernel(Matrix?A,?Matrix?B,?Matrix?C)?{????

????//?Each?thread?computes?one?element?of?C????

????//?by?accumulating?results?into?Cvalue????

????float?Cvalue?=?0;????

????int?row??=?blockIdx.y?*?blockDim.y?+?threadIdx.y;????

????int?col?=?blockIdx.x?*?blockDim.x?+?threadIdx.xl????

????for?(int?e?=?0;?e?<?A.width;?++e)????

????????Cvalue?+=?A.elements[row?*?A.width?+?e]?*?B.elements[e?*?B.width?+?col];????

????C.elements[row?*?C.width?+?col]?=?Cvalue;????

}??

可以看出，為了計算矩陣C的任何一個元素，程序都需要從全局內(nèi)存（global memory）中獲得矩陣A的一行和矩陣B的一列。因此，完成這一計算矩陣A被讀取了B.width次，矩陣B被讀取了A.height次。 現(xiàn)在我們來使用共享內(nèi)存（shared memory）實現(xiàn)矩陣乘法。假設矩陣C可以被劃分為若干個較小的子方陣Csub，我們使用一個線程塊（thread block）來負責某一子方陣的計算，線程塊中的每一個線程（thread）正好負責子方陣Csub中一個元素的計算。這樣劃分后，任何一個結果子方陣Csub'（尺寸為block_size * block_size）都是與該方陣具有相同行索引的尺寸為A.width * block_size的A的子矩陣Asub和與該方陣具有相同列索引的尺寸為block_size * B.height的B的子矩陣Bsub相乘所得到。 ? ? ? ? 為了匹配設備的計算資源，兩個子矩陣Asub和Bsub被劃分為盡可能多的分離的維度為block_size的子方陣，Csub的值便是這些子矩陣相乘后相加所得到的結果。子矩陣乘法的執(zhí)行順序都是首先將它們從全局內(nèi)存（global memory）拷貝到共享內(nèi)存（shared memory）（線程塊中的每一個線程正好負責方陣一個元素的拷貝），然后由線程自己完成相應元素的計算任務，利用寄存器存儲局部結果，最后將寄存器的內(nèi)容與新得到的計算結果依此累加起來得到最終運算結果并將其傳輸?shù)饺謨?nèi)存（global memory）中。 ? ? ? ? 通過使用這種分治的計算策略，共享內(nèi)存得到了很好的利用，采用這種方案計算完成時全局內(nèi)存中矩陣A被訪問的次數(shù)為B.width / block_size，矩陣B被訪問的次數(shù)為A.height / block_size，很明顯，這為我們節(jié)省了非常多的全局內(nèi)存帶寬。優(yōu)化后的矩陣計算示意圖如下所示：

















































? ? ? ? 為了提升計算效率，我們?yōu)轭愋蚆atrix增加了一個成員變量stride。__device__函數(shù)用來獲得和設置子矩陣的元素。下面是優(yōu)化后的代碼： [cpp] view plaincopy

//?Matrices?are?stored?in?row-major?order;????

//?M(row,?col)?=?*(M.elements?+?row?*?M.stride?+?col)????

typedef?struct?{????

????int?width;????

????int?height;????

????int?stride;????

????float*?elements;????

}?Matrix;????

????

//?Get?a?matrix?element????

__device__?float?GetElement(const?Matrix?A,?int?row,?int?col)?{????

????return?A.elements[row?*?A.stride?+?col];????

}????

????

//?Set?a?matrix?element????

__device__?void?SetElement(Matrix?A,?int?row,?int?col,?float?value)?{????

????A.elements[row?*?A.stride?+?col]?=?value;????

}????

????

//?Get?the?BLOCK_SIZExBLOCK_SIZE?sub-matrix?Asub?of?A?that?is????

//?located?col?sub-matrices?to?the?right?and?row?sub-matrices?down????

//?from?the?upper-left?corner?of?A????

__device__?Matrix?GetSubMatrix(Matrix?A,?int?row,?int?col)?{????

????Matrix?Asub;????

????Asub.width?=?BLOCK_SIZE;????

????Asub.height?=?BLOCK_SIZE;????

????Asub.stride?=?A.stride;????

????Asub.elements?=?&A.elements[A.stride?*?BLOCK_SIZE?*?row?+?BLOCK_SIZE?*?col];????

????return?Asub;????

}????

????

//?Thread?block?size????

#define?BLOCK_SIZE?16????

????

//?Forward?declaration?of?the?matrix?multiplication?kernel????

__global__?void?MatMulKernel(const?Matrix,?const?Matrix,?Matrix);????

????

//?Matrix?multiplication?-?Host?code????

//?Matrix?dimensions?are?assumed?to?be?multiples?of?BLOCK_SIZE????

void?MatMul(const?Matrix?A,?const?Matrix?B,?Matrix?C)?{????

????//?Load?A?and?B?to?device?memory????

????Matrix?d_A;????

????d_A.width?=?d_A.stride?=?A.width;????

????d_A.height?=?A.height;????

????size_t?size?=?A.width?*?A.height?*?sizeof(float);????

????cudaMalloc(&d_A.elements,?size);????

????cudaMemcpy(d_A.elements,?A.elements,?size,?cudaMemcpyHostToDevice);????

????Matrix?d_B;????

????d_B.width?=?d_B.stride?=?B.width;????

????d_B.height?=?B.height;????

????size?=?B.width?*?B.height?*?sizeof(float);????

????cudaMalloc(&d_B.elements,?size);????

????cudaMemcpy(d_B.elements,?B.elements,?size,?cudaMemcpyHostToDevice);????

????

????//?Allocate?C?in?device?memory????

????Matrix?d_C;????

????d_C.width?=?d_C.stride?=?C.width;????

????d_C.height?=?C.height;????

????size?=?C.width?*?C.height?*?sizeof(float);????

????cudaMalloc(&d_C.elements,?size);????

????

????//?Invoke?kernel????

????dim3?dimBlock(BLOCK_SIZE,?BLOCK_SIZE);????

????dim3?dimGrid(B.width?/?dimBlock.x,?A.height?/?dimBlock.y);????

????MatMulKernel<<<dimGrid,?dimBlock>>>(d_A,?d_B,?d_C);????

????

????//?Read?C?from?device?memory????

????cudaMemcpy(C.elements,?d_C.elements,?size,?cudaMemcpyDeviceToHost);????

????

????//?Free?device?memory????

????cudaFree(d_A.elements);????

????cudaFree(d_B.elements);????

????cudaFree(d_C.elements);????

}????

????

//?Matrix?multiplication?kernel?called?by?MatMul()????

__global__?void?MatMulKernel(Matrix?A,?Matrix?B,?Matrix?C)?{????

????//?Block?row?and?column????

????int?blockRow?=?blockIdx.y;????

????int?blockCol?=?blockIdx.x;????

????

????//?Each?thread?block?computes?one?sub-matrix?Csub?of?C????

????Matrix?Csub?=?GetSubMatrix(C,?blockRow,?blockCol);????

????

????//?Each?thread?computes?one?element?of?Csub????

????//?by?accumulating?results?into?Cvalue????

????float?Cvalue?=?0;????

????

????//?Thread?row?and?column?within?Csub????

????int?row?=?threadIdx.y;????

????int?col?=?threadIdx.x;????

????

????//?Look?over?all?the?sub-matrices?of?A?and?B?that?are?required?to?compute?Csub????

????//?Multiply?each?pair?of?sub-matrices?together?and?accumulate?the?results????

????for?(int?m?=?0;?m?<?(A.width?/?BLOCK_SIZE);?++m)?{????

????????//?Get?sub-matrix?Asub?of?A????

????????Matrix?Asub?=?GetSubMatrix(A,?blockRow,?m);????

????????????

????????//?Get?sub-matrix?Bsub?of?B????

????????Matrix?Bsub?=?GetSubMatrix(B,?m,?blockCol);????

????

????????//?Shared?memory?used?to?store?Asub?and?Bsub?respectively????

????????__shared__?float?As[BLOCK_SIZE][BLOCK_SIZE];????

????????__shared__?float?Bs[BLOCK_SIZE][BLOCK_SIZE];????

????

????????//?Load?Asub?and?Bsub?from?device?memory?to?shared?memory????

????????//?Each?thread?loads?one?element?of?each?sub-matrix????

????????As[row][col]?=?GetElement(Asub,?row,?col);????

????????Bs[row][col]?=?GetElement(Bsub,?row,?col);????

????

????????//?Synchronize?to?make?sure?the?sub-matrices?are?loaded????

????????//?before?starting?the?computation????

????????__syncthreads();????

????

????????//?Multiply?Asub?and?Bsub?together????

????????for?(int?e?=?0;?e?<?BLOCK_SIZE;?++e)????

????????????Cvalue?+=?As[row][e]?*?Bs[e][col];????

????

????????//?Synchronize?to?make?sure?that?the?preceding?computation?is?done?before????

????????//?loading?two?new?sub-matrices?of?A?and?B?in?the?next?iteration????

????????__syncthreads();????

????}????

????

????//?Write?Csub?to?device?memory????

????//?Each?thread?writes?one?element????

????SetElement(Csub,?row,?col,?Cvalue);????

}???

異步并行執(zhí)行

主機和設備間并行執(zhí)行

? ? ? ? 為了支持主機和設備的并行執(zhí)行，CUDA提供了一些異步函數(shù)。異步是指設備在完成功能執(zhí)行之前就將控制權交還給主機線程，以便主機線程繼續(xù)執(zhí)行。這些函數(shù)有： ? ? ? ? 1、內(nèi)涵啟動（Kernel Launches）； ? ? ? ? 2、同一設備內(nèi)存中兩個地址塊之間的數(shù)據(jù)傳輸； ? ? ? ? 3、從主機到設備的小于或等于64KB的一次數(shù)據(jù)塊傳輸； ? ? ? ? 4、使用Async前綴函數(shù)進行的數(shù)據(jù)傳輸； ? ? ? ? 5、內(nèi)存置位函數(shù)調(diào)用（Memory set functions calls）。 程序員可以通過設置環(huán)境變量CUDA_LAUNCH_BLOCKING來開啟或者關閉內(nèi)核啟動（Kernel Launch）的異步功能。但是這一方法僅限于調(diào)試，在任何產(chǎn)品代碼中不應當關閉異步內(nèi)核啟動。 ? ? ? ? 內(nèi)核啟動在下面這些情況下則是同步的： ? ? ? ? 1、應用程序通過調(diào)試器或者內(nèi)存檢查器運行在計算能力為1.x的設備上。 ? ? ? ? 2、硬件計數(shù)器信息正被性能分析器收集。

將內(nèi)核啟動與數(shù)據(jù)傳輸重疊起來

? ? ? ? 對于一些計算能力等于或高于1.1的設備，它們可以將內(nèi)核啟動任務和鎖頁內(nèi)存到設備內(nèi)存的數(shù)據(jù)傳輸任務并行執(zhí)行。應用程序可以檢查設備屬性中的asyncEngineCount項來確定設備是否支持這一功能。當該項值大于0時代表設備支持這一層次的并行。對于計算能力1.x的設備，該功能不支持通過cudaMallocPitch()函數(shù)分配的CUDA數(shù)組或2D數(shù)組。

并行內(nèi)核執(zhí)行

? ? ? ? 一些計算能力2.x或更高的設備可以同時并行執(zhí)行多個內(nèi)核函數(shù)。應用程序可以檢查設備屬性中的concurrentKernels項來確定設備是否支持這一功能，值為1代表支持。運算能力3.5的設備在同一時刻能夠并行執(zhí)行的最大內(nèi)核函數(shù)數(shù)量為32，運算能力小于3.5的硬件則最多支持同時啟動16個內(nèi)核函數(shù)的執(zhí)行。同時需要注意的是，在一個CUDA上下文中的內(nèi)核函數(shù)不能與另一個CUDA上下文中的內(nèi)核函數(shù)同時執(zhí)行。使用很多紋理內(nèi)存或者大量本地內(nèi)存的內(nèi)核函數(shù)也很可能無法與其它內(nèi)核函數(shù)并行執(zhí)行。

并行數(shù)據(jù)傳輸

? ? ? ? 一些計算能力為2.x或更高的設備可以將鎖頁內(nèi)存到設備內(nèi)存的數(shù)據(jù)傳輸和設備內(nèi)存到鎖頁內(nèi)存的數(shù)據(jù)傳輸并行執(zhí)行。應用程序可檢查設備屬性中的asyncEngineCount項來確定這一功能的支持程度，等于2時表示支持。

流（Streams）

? ? ? ? 應用程序通過流來管理并行。一個流是一個順次執(zhí)行的命令序列。不同的流之間并行執(zhí)行，沒有固定的執(zhí)行順序。

1、流的創(chuàng)建與銷毀

? ? ? ? 定義一個流的過程通常包括：創(chuàng)建一個流對象，然后指定它為內(nèi)核啟動或者主機設備間數(shù)據(jù)傳輸?shù)牧鲄?shù)。下面的一段代碼創(chuàng)建了兩個流并且在鎖頁內(nèi)存中分配了一塊float類型的數(shù)組hostPtr：

[cpp] view plaincopy

cudaStream_t?stream[2];??

for?(int?i?=?0;?i?<?2;?++i)??

????cudaStreamCreate(&stream[i]);??

float?*hostPtr;??

cudaMallocHost(&hostPtr,?2?*?size);??

下面的代碼定義了每一個流的行為：從主機端拷貝數(shù)據(jù)到設備端，內(nèi)核啟動，從設備端拷貝數(shù)據(jù)到主機端：
[cpp] view plaincopy

for?(int?i?=?0;?i?<?2;?++i)?{??

????cudaMemcpyAsync(inputDevPtr?+?i?*?size,?hostPtr?+?i?*?size,?size,?cudaMemcpyHostToDevice,?stream[i]);??

????MyKernel<<<100,?512,?0,?stream[i]>>>(outputDevPtr?+?i?*?size,?inputDevPtr?+?i?*?size,?size);??

????cudaMemcpyAsync(hostPtr?+?i?*?size,?outputDevPtr?+?i?*?size,?size,?cudaMemcpyDeviceToHost,?stream[i]);??

}??

這部分代碼中有一點需要注意：為了并行化數(shù)據(jù)拷貝和內(nèi)核執(zhí)行，主機端內(nèi)存必須分配為鎖頁（page-locked）內(nèi)存。

要銷毀一個流需要調(diào)用函數(shù)cudaStreamDestroy()

[cpp] view plaincopy

for?(int?i?=?0;?i?<?2;?++i)??

????cudaStreamDestroy(stream[i]);??

cudaStreamDestroy() 函數(shù)等待之前流中的指令序列運行完成，然后銷毀指定流，將控制權返還給主機端。

2、默認流（Default stream）

? ? ? ? 在內(nèi)核啟動或者數(shù)據(jù)拷貝過程中如果不指定流，或者設置流參數(shù)為0，則相應的指令將會運行在默認流上，它們也因此而順次執(zhí)行。

3、明同步（Explicit Synchronization）

? ? ? ? 在CUDA中有很多種方式可以用來同步流的執(zhí)行：

? ? ? ??cudaDeviceSynchronize()函數(shù)使得主機端線程阻塞直到所有流中的指令執(zhí)行完成。

? ? ? ?cudaStreamSynchronize()函數(shù)將一個流對象作為輸入?yún)?shù)，用以等待指定流中的所有指令執(zhí)行完成。

? ? ? ?cudaStreamWaitEvent()函數(shù)將一個流對象和一個事件作為輸入?yún)?shù)，它將延遲該函數(shù)調(diào)用后在指定流中所有新加入的命令的執(zhí)行直到指定的事件完成為止。流參數(shù)可以為0，在該情形下所有流中的任何新加入的指令都必須等待指定事件的發(fā)生，然后才可以執(zhí)行。

? ? ? ?cudaStreamQuery()函數(shù)為應用程序提供了一個檢測指定流中之前指令是否執(zhí)行完成的方法。

? ? ? ? 為了避免同步帶來的性能下降，所有上述同步函數(shù)最好用于計時目的或者分離錯誤的內(nèi)核執(zhí)行或數(shù)據(jù)拷貝。

4、暗同步（Implicit Synchronization）

? ? ? ? 如果任何一個流中正在執(zhí)行以下操作，那么其它流是不能與其并行運行的：

? ? ? ? a. 分配鎖頁內(nèi)存空間

? ? ? ? b. 設備內(nèi)存分配

? ? ? ? c. 設備內(nèi)存置位

? ? ? ? d. 同一設備兩個不同地址間正在進行數(shù)據(jù)拷貝

? ? ? ? e. 默認流中有指令正在執(zhí)行

? ? ? ? f. L1/shared內(nèi)存配置的轉(zhuǎn)換

? ? ? ? 對于支持并行內(nèi)核執(zhí)行并且計算能力3.0或以下的設備來說，任何一個需要檢查依賴性以確定流內(nèi)核啟動是否完成的操作：

? ? ? ? a. 只有當前CUDA上下文中所有流中所有之前的內(nèi)核啟動之后才能夠啟動執(zhí)行。

? ? ? ? b. 將會阻塞所有當前CUDA上下文中的任意流中新加入的內(nèi)核調(diào)用直到內(nèi)核檢查完成。

? ? ? ? 需要進行依賴性檢查的操作包括執(zhí)行檢查的內(nèi)核啟動所在流中的其它指令以及任何在該流上對cudaStreamQuery()函數(shù)的調(diào)用。因此，應用程序可以遵照以下指導原則來提升潛在并行性：

? ? ? ? （1）所有非依賴操作應當比依賴性操作提前進行

? ? ? ? （2）任何類型的同步越遲越好

5、重疊行為（Overlapping Behavior）

? ? ? ? 兩個流間重疊行為的數(shù)量取決于以下幾個因素：

? ? ? ? （1）每個流中命令發(fā)出的次序

? ? ? ? （2）設備是否支持內(nèi)核啟動與數(shù)據(jù)傳輸并行

? ? ? ? （3）設備是否支持多內(nèi)核并行啟動

? ? ? ? （4）設備是否支持多數(shù)據(jù)傳輸并行

? ? ? ? 例如，在不支持并行數(shù)據(jù)傳輸?shù)脑O備上，“流的創(chuàng)建與銷毀”章節(jié)中代碼樣例中的操作就不能并行，因為在stream[0]中發(fā)出設備端到主機端的數(shù)據(jù)拷貝后，stream[1]又發(fā)出主機端到設備端的數(shù)據(jù)拷貝命令，這兩個命令式不能重疊執(zhí)行的。假設設備支持數(shù)據(jù)傳輸與內(nèi)核啟動并行，那么如下代碼：

[cpp] view plaincopy

for?(int?i?=?0;?i?<?2;?++i)??

????cudaMemcpyAsync(inputDevPtr?+?i?*?size,?hostPtr?+?i?*?size,?size,?cudaMemcpyHostToDevice,?stream[i]);??

for?(int?i?=?0;?i?<?2;?++i)??

????MyKernel<<<100,?512,?0,?stream[i]>>>(outputDevPtr?+?i?*?size,?inputDevPtr?+?i?*?size,?size);??

for?(int?i?=?0;?i?<?2;?++i)??

????cudaMemcpyAsync(hostPtr?+?i?*?size,?outputDevPtr?+?i?*?size,?size,?cudaMemcpyDeviceToHost,?stream[i]);??

可將stream[0]的內(nèi)核啟動和stream[1]從主機端到設備端的數(shù)據(jù)拷貝重疊起來并行執(zhí)行。

6、回調(diào)函數(shù)

? ? ? ? CUDA運行時提供了cudaStreamAddCallback()函數(shù)以在流中的任意位置插入一個回調(diào)函數(shù)點。回調(diào)函數(shù)運行于主機端，如果在默認流中插入回調(diào)函數(shù)，那么它將等待所有其它流中的命令執(zhí)行完成之后才會開始執(zhí)行。

? ? ? ? 下面的代碼展示了回調(diào)函數(shù)技術的應用：

[cpp] view plaincopy

void?CUDART_CB?MyCallback(cudaStream_t?stream,?cudaError_t?status,?void?**data)?{??

????printf("Inside?callback?%d\n",?(int)data);??

}??

...??

for?(int?i?=?0;?i?<?2;?++i)?{??

????cudaMemcpyAsync(devPtrIn[i],?hostPtr[i],?size,?cudaMemcpyHostToDevice,?stream[i]);??

????MyKernel<<<100,?512,?0,?stream[i]>>>(devPtrOut[i],?devPtrIn[i],?size);??

????cudaMemcpyAsync(hostPtr[i],?devPtrOut[i],?size,?cudaMemcpyDeviceToHost,?stream[i]);??

????cudaStreamAddCallback(stream[i],?MyCallback,?(void**)i,?0);??

}??

上面的代碼定義了兩個流的操作，每個流都完成一次主機端到設備端的數(shù)據(jù)拷貝，一次內(nèi)核啟動，一次設備端到主機端的數(shù)據(jù)拷貝，最后增加了一個加入回調(diào)函數(shù)的操作。當設備端代碼運行到回調(diào)函數(shù)點的時候，設備將控制權交還給主機端，主機端運行完成以后再將控制權返還給設備端，然后設備端繼續(xù)運行。

? ? ? ? 值得注意的是，在一個回調(diào)函數(shù)中，一定不能進行任何CUDA API的調(diào)用，直接的或者間接的都是不可以的。

總結

以上是生活随笔為你收集整理的CUDA编程指南阅读笔记的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網(wǎng)站內(nèi)容還不錯，歡迎將生活随笔推薦給好友。