CUDA C编程（二）CUDA编程模型
CUDA編程模型概述
??CUDA是一種通用的并行計算平臺和編程模型，是在C語言上擴展的的。借助于CUDA，可以像編寫C語言程序一樣實現并行算法。
??CUDA編程模型提供了一個計算機架構抽象作為應用程序和其可用硬件之間的橋梁。下圖說明了程序和編程模型實現之間的抽象結構。其中通信抽象是程序與編程模型實現之間的分界線，它通過專業的硬件原語和操作系統的編譯器或庫來實現。利用編程模型所編寫的程序指定了程序的各組成部分是如何共享信息及相互協作的。編程模型從邏輯上提供了一個特定的計算機架構，通常它體現在編程語言或編程環境中。
??除了與其他并行編程模型共有的抽象外，CUDA編程模型還利用GPU架構的計算能力提供了以下特有功能：
??>一種通過層次結構在GPU中組織線程的方法；
??>一種通過層次結構在GPU中訪問內存的方法。
CUDA 編 程 結 構
??CUDA編程模型使用由C語言擴展生成的注釋代碼在異構計算系統中執行應用程序。在一個異構環境中包含多個CPU和GPU，每個GPU和CPU的內存都由一條PCI-Express總線分隔開，因此，需要注意區分主機和設備這兩個內容：
??>主機：CPU及其內存（主機內存）；
??>設備：GPU及其內存（設備內存）。
??為了清楚地指明不同地內存空間，主機內存中的變量名與設備內存中的變量名一般使用不同的前綴。
??內核（kernel）是一個CUDA編程模型的一個重要組成部分，其代碼在GPU上運行。多數情況下，主機可以獨立地對設備進行操作。內核一旦被啟動，管理權立刻返回給主機，釋放CPU來執行由設備上運行地并行代碼實現地額外的任務。CUDA編程模型主要是異步的，因此在GPU上進行的運算可以與主機-設備通信重疊。一個典型的CUDA程序包括并行代碼互補的串行代碼。如下圖所示，串行代碼在CPU上執行，而并行代碼在GPU上執行。主機代碼按照ANSI C標準進行編寫，而設備代碼使用CUDA C進行編寫。NVIDIA的C編譯器（nvcc）為主機和設備生成可執行代碼。
??一個典型的CUDA程序實現流程遵循以下模式：
??>1.把數據從CPU內存拷貝到GPU內存；
??>2.調用核函數對存儲在GPU內存中的數據進行操作。
??>3.將數據從GPU內存傳送回到CPU內存。
內 存 管 理
??為了擁有充分的控制權并使系統達到最佳性能，CUDA運行時負責分配與釋放設備內存，并且在主機內存和設備內存之間傳輸數據。C語言以及相應的針對內存操作的CUDA C函數如下表所示：
| malloc | cudaMalloc | memset | cudaMemset |
| memcpy | cudaMemcpy | free | cudaFree |
??cudaMalloc與標準C語言的malloc函數幾乎一樣，只是此函數在GPU的內存里分配內存。
??cudaMemcpy函數負責主機和設備之間的數據傳輸，其函數原型為：cudaError_t cudaMemcpy(void* dst,const void* src,size_t count, cudaMemcpyKind kind);此函數從src指向的源存儲區復制一定數量的字節到dst指向的目標存儲區。復制方向由kind指定，其中的kind有以下幾種：
??>cudaMemcpyHostToHost
??>cudaMemcpyHostToDevice
??>cudaMemcpyDeviceToHost
??>cudaMemcpyDeviceToDevice
??這個函數以同步方式執行，因為在cudaMemcpy函數返回以及傳輸操作完成之前主機應用程序是阻塞的。除了內核啟動之外的CUDA調用都會返回一個錯誤的枚舉類型cudaError_t。如果GPU內存分配成功，函數返回cudaSuccess，否則返回cudaErrorMemoryAllocation。
??CUDA編程模型從GPU架構中抽象出一個內存層次結構。如下圖所示，主要包含兩部分：全局內存和共享內存。
線 程 管 理
??當核函數在主機端啟動時，它的執行會移動到設備上，此時設備中會產生大量的線程，并且每個線程都執行由核函數指定的語句。一個兩層的線程層次結構如下圖所示，由線程塊和線程塊網格構成。
??又一個內核所產生的所有線程統稱為一個網格。同一網格中的所有線程共享相同的全局內存空間。一個網格由多個線程塊構成，一個線程塊包含一組線程，同一線程塊內的線程協作可以通過同步、共享內存等方式來實現，而不同塊內的線程不能協作。
??線程依靠blockIdx(線程塊在線程格內的索引)以及threadIdx(塊內的線程索引)這兩個坐標變量來區分彼此。這些變量是核函數中需要預初始化的內置變量。當執行一個核函數時，CUDA運行時為每個線程分配坐標變量blockIdx和threadIdx。基于這些坐標，我們可以將部分數據分配給不同的線程。該坐標變量是基于uint3定義的CUDA內置的向量類型，是一個包含3個無符號整數的結構，可以通過x、y、z三個字段來指定。
??CUDA可以組織三維的網格和塊。線程層次結構如上面的圖所示，其結構是一個包含二維塊的二維網格。網格和塊的維度由blockDim(線程塊的維度，用每個線程塊中的線程數來表示)、gridDim(線程格的維度，用每個線程格中的線程數來表示)兩個內置變量指定，它們是dim3類型的變量，是基于uint3定義的整數型向量，用來表示維度。當定義一個dim3類型的變量時，所有未指定的元素都被初始化為1.dim3類型變量中的每個組件可以通過它的x、y、z字段獲得。
??在CUDA程序中有兩組不同的網格和塊變量：手動定義的dim3數據類型和預定義的uint3數據類型。在主機端，作為內核調用的一部分，可以使用dim3數據類型定義一個網格和塊的維度。當執行核函數時， CUDA運行時會生成相應的內置預初始化的網格、塊和線程變量，它們在核函數內均可被訪問到且為uint3類型。手動定義的dim3類型的網格和塊變量僅在主機端可見，而uint3類型的內置預初始化的網格和塊變量僅在設備端可見。
??由于一個內核啟動的網格和塊的維數會影響性能，這一結構為程序員優化程序提供了一個額外的途徑。對于一個給定的數據大小，確定網格和塊尺寸的一般步驟為：確定塊的個數、在已知數據大小和塊大小的基礎上計算網格維度。網格和塊的維度存在幾個限制因素：對于塊大小的一個主要限制因素就是可利用的計算資源，如寄存器，共享內存等。某些限制可以通過查詢GPU設備。
啟 動 一 個 CUDA 核 函 數
??C語言函數調用語句：function_name(argument list);而CUDA內核調用是對C語言函數調用語句的延申，kernal_name <<<grid,block>>>(argument list);其中<<<>>>運算符內是核函數的執行配置。執行配置的第一個值是網絡維度，也就是啟動塊的數目，第二個值是塊維度，也就是每個塊中線程的數目。通過指定網格和塊的維度，可以進行內核中線程的數目以及內核中使用的線程布局的配置。
??由于數據在全局內存中是線性存儲的，因為可以用變量blockIdx.x和threadIdx.x來進行以下操作：在網格中標識一個唯一的線程、建立線程和數據元素之間的映射關系。不同于C語言的函數調用，所有的CUDA核函數的啟動都是異步的。CUDA內核調用完成后，控制權立刻返回給CPU。可以通過調用以下函數來強制主機端程序等待所有的核函數執行結束：cudaError_t cudaMemcpy(void* dst,const void* src,size_t count,cudaMemcpykind kind);
編 寫 核 函 數
??核函數是在設備端執行的代碼。在核函數中，需要為一個線程規定要進行的計算以及要進行的數據訪問。當核函數被調用時，許多不同的CUDA線程并行執行同一個計算任務。核函數必須有一個void返回類型。
??函數類型限定符指定了一個函數在主機上執行還是在設備上執行，以及可被主機調用還是被設備調用，下表總結了CUDA C程序中的函數類型限定符。
| global | 在設備端執行 | 可從主機端調用，也可以從計算能力為3的設備中調用 | 必須有一個void返回類型 |
| device | 在設備端執行 | 僅能從設備端調用 | |
| host | 在主機端執行 | 僅能從主機端調用 | 可以省略 |
??__device__和__host__限定符可以一齊使用，這樣函數可以同時在主機和設備端進行編譯。
??CUDA核函數的限制：只能訪問設備內存、必須具有void返回類型、不支持可變數量的參數、不支持靜態變量、顯示異步行為。
驗 證 核 函 數
??除了許多可用的調試工具外，還有兩個非常簡單實用的方法可以驗證核函數：首先，你可以在Fermi及更高版本的設備端的核函數中使用printf函數；其次，可以將執行參數設置為<<<1,1>>>,因此強制用一個塊和一個線程執行核函數，這模擬了串行執行程序。
處 理 錯 誤
??由于許多CUDA調用是異步的，所有有時很難確定某個錯誤是由哪一步程序引起的。定義一個錯誤處理宏封裝所有的CUDA API調用，這簡化了錯誤檢查過程：
??完成錯誤檢查宏的定義之后，就可以在代碼中使用宏。比如：CHECK(cudaMemcpy(d_C,gpuRef,nBytes,cudaMemcpyHostToDevice));如果內存拷貝或之前的異步操作產生了錯誤，這個宏會報告錯誤代碼，并且輸出一個可讀信息，然后停止程序。再比如：CHECK(cudaDeviceSynchronize())會阻塞主機端線程的運行直到設備端所有的請求都結束，并確保最后的核函數啟動部分不會出錯。
給核函數計時
用 CPU 計 時 器 計 時
??可以使用gettimeofday系統調用來創建一個CPU計時器，以獲取系統的時鐘時間，它需要包含sys/time.h頭文件。
??這樣可以用cpuSecond函數來測試核函數：
double iStart = cpuSecord(); kernel_name<<<grid,block>>>(argument list); cudaDeviceSynchronize(); double iElaps = cpuSecond() - iStart;用 nvprof 工 具 計 時
??自CUDA5.0以來，NVIDIA提供了一個名為nvprof的命令行分析工具，可以幫助從應用程序的CPU和GPU活動情況中獲取時間線信息，其包括內核執行、內存傳輸以及CUDA API的調用。
組織并行線程
??如果使用了合適的網絡和塊大小來正確地組織線程，那么可以對內核性能產生很大的影響。對于矩陣運算，傳統的方法是在內核中使用一個包含二維網格與二維塊的布局來組織線程，但是，這種傳統的方法無法獲得最佳性能。在矩形加法中使用由二維線程塊構成的二維網絡、由一維線程塊構成的一維網格、由一維線程塊構成的二維網格等布局可以取得不同性能的結果。
使 用 塊 和 線 程 建 立 矩 陣 索 引
??在一個矩陣加法核函數中，一個線程通常被分配一個數據元素來處理。首先要完成的任務是使用塊和線程索引從全局內存中訪問指定的數據。通常情況下，對一個二維示例來說，需要管理線程和塊索引、矩陣中給定點的坐標、全局線性內存中的偏移量三種索引。對于一個給定的線程，首先可以通過把線程和塊索引映射到矩陣坐標上來獲取線程塊和縣城索引的全局內存偏移量，然后將這些矩陣坐標映射到全局內存的存儲單元中。
??第一步，可以用以下公式把線程和塊索引映射到矩陣坐標上：
??第二步，可以用以下公式把矩陣坐標映射到全局內存中的索引/存儲單元上：
idx = iy * nx + ix;??其中，printThreadIndo函數被用于輸出關于每個線程的以下信息：線程索引、塊索引、矩陣坐標、線性全局內存偏移、相應元素的值。
設備管理
??查詢和管理GPU設備的兩種方法：
??>CUDA運行時API函數；
??>NVIDIA系統管理界面（nvidia-smi）命令行實用程序。
使 用 運 行 時 API 查 詢 GPU 信 息
??在CUDA運行時API中由很多函數可以幫助管理這些設備。可以使用以下函數來查詢關于GPU設備的所有信息：
??cudaDeviceProp結構體返回GPU設備的屬性。
確 定 最 優 GPU
??一些系統支持多GPU。在每個GPU都不同的情況下，選擇性能最好的GPU運行核函數是非常重要的。通過比較GPU包含的多處理器的數量選出計算能力最佳的GPU。可以使用以下代碼來選擇計算能力最優的設備：
使 用 nvidia-smi 查 詢 GPU 信 息
??nvidia-smi是一個命令行工具，用于管理和監控GPU設備，并允許查詢和修改設備狀態。可以從命令行調用nvidia-smi。
總結
以上是生活随笔為你收集整理的CUDA C编程（二）CUDA编程模型的全部內容，希望文章能夠幫你解決所遇到的問題。
- 上一篇： 常见的Markdownpad2运行破解以
- 下一篇： MyEclipse中SVN的常见的使用方