CUDA編程模型概述

??CUDA是一種通用的并行計算平臺和編程模型，是在C語言上擴展的的。借助于CUDA，可以像編寫C語言程序一樣實現(xiàn)并行算法。
??CUDA編程模型提供了一個計算機架構(gòu)抽象作為應(yīng)用程序和其可用硬件之間的橋梁。下圖說明了程序和編程模型實現(xiàn)之間的抽象結(jié)構(gòu)。其中通信抽象是程序與編程模型實現(xiàn)之間的分界線，它通過專業(yè)的硬件原語和操作系統(tǒng)的編譯器或庫來實現(xiàn)。利用編程模型所編寫的程序指定了程序的各組成部分是如何共享信息及相互協(xié)作的。編程模型從邏輯上提供了一個特定的計算機架構(gòu)，通常它體現(xiàn)在編程語言或編程環(huán)境中。

??除了與其他并行編程模型共有的抽象外，CUDA編程模型還利用GPU架構(gòu)的計算能力提供了以下特有功能：
??>一種通過層次結(jié)構(gòu)在GPU中組織線程的方法；
??>一種通過層次結(jié)構(gòu)在GPU中訪問內(nèi)存的方法。

CUDA 編 程 結(jié) 構(gòu)
??CUDA編程模型使用由C語言擴展生成的注釋代碼在異構(gòu)計算系統(tǒng)中執(zhí)行應(yīng)用程序。在一個異構(gòu)環(huán)境中包含多個CPU和GPU，每個GPU和CPU的內(nèi)存都由一條PCI-Express總線分隔開，因此，需要注意區(qū)分主機和設(shè)備這兩個內(nèi)容：
??>主機：CPU及其內(nèi)存（主機內(nèi)存）；
??>設(shè)備：GPU及其內(nèi)存（設(shè)備內(nèi)存）。
??為了清楚地指明不同地內(nèi)存空間，主機內(nèi)存中的變量名與設(shè)備內(nèi)存中的變量名一般使用不同的前綴。
??內(nèi)核（kernel）是一個CUDA編程模型的一個重要組成部分，其代碼在GPU上運行。多數(shù)情況下，主機可以獨立地對設(shè)備進(jìn)行操作。內(nèi)核一旦被啟動，管理權(quán)立刻返回給主機，釋放CPU來執(zhí)行由設(shè)備上運行地并行代碼實現(xiàn)地額外的任務(wù)。CUDA編程模型主要是異步的，因此在GPU上進(jìn)行的運算可以與主機-設(shè)備通信重疊。一個典型的CUDA程序包括并行代碼互補的串行代碼。如下圖所示，串行代碼在CPU上執(zhí)行，而并行代碼在GPU上執(zhí)行。主機代碼按照ANSI C標(biāo)準(zhǔn)進(jìn)行編寫，而設(shè)備代碼使用CUDA C進(jìn)行編寫。NVIDIA的C編譯器（nvcc）為主機和設(shè)備生成可執(zhí)行代碼。

??一個典型的CUDA程序?qū)崿F(xiàn)流程遵循以下模式：
??>1.把數(shù)據(jù)從CPU內(nèi)存拷貝到GPU內(nèi)存；
??>2.調(diào)用核函數(shù)對存儲在GPU內(nèi)存中的數(shù)據(jù)進(jìn)行操作。
??>3.將數(shù)據(jù)從GPU內(nèi)存?zhèn)魉突氐紺PU內(nèi)存。

內(nèi) 存 管 理
??為了擁有充分的控制權(quán)并使系統(tǒng)達(dá)到最佳性能，CUDA運行時負(fù)責(zé)分配與釋放設(shè)備內(nèi)存，并且在主機內(nèi)存和設(shè)備內(nèi)存之間傳輸數(shù)據(jù)。C語言以及相應(yīng)的針對內(nèi)存操作的CUDA C函數(shù)如下表所示：

標(biāo)準(zhǔn)的C函數(shù)CUDA C函數(shù)標(biāo)準(zhǔn)的C函數(shù)CUDA C函數(shù)

malloc	cudaMalloc	memset	cudaMemset
memcpy	cudaMemcpy	free	cudaFree

??cudaMalloc與標(biāo)準(zhǔn)C語言的malloc函數(shù)幾乎一樣，只是此函數(shù)在GPU的內(nèi)存里分配內(nèi)存。
??cudaMemcpy函數(shù)負(fù)責(zé)主機和設(shè)備之間的數(shù)據(jù)傳輸，其函數(shù)原型為：cudaError_t cudaMemcpy(void* dst,const void* src,size_t count, cudaMemcpyKind kind);此函數(shù)從src指向的源存儲區(qū)復(fù)制一定數(shù)量的字節(jié)到dst指向的目標(biāo)存儲區(qū)。復(fù)制方向由kind指定，其中的kind有以下幾種：
??>cudaMemcpyHostToHost
??>cudaMemcpyHostToDevice
??>cudaMemcpyDeviceToHost
??>cudaMemcpyDeviceToDevice
??這個函數(shù)以同步方式執(zhí)行，因為在cudaMemcpy函數(shù)返回以及傳輸操作完成之前主機應(yīng)用程序是阻塞的。除了內(nèi)核啟動之外的CUDA調(diào)用都會返回一個錯誤的枚舉類型cudaError_t。如果GPU內(nèi)存分配成功，函數(shù)返回cudaSuccess，否則返回cudaErrorMemoryAllocation。
??CUDA編程模型從GPU架構(gòu)中抽象出一個內(nèi)存層次結(jié)構(gòu)。如下圖所示，主要包含兩部分：全局內(nèi)存和共享內(nèi)存。

線 程 管 理
??當(dāng)核函數(shù)在主機端啟動時，它的執(zhí)行會移動到設(shè)備上，此時設(shè)備中會產(chǎn)生大量的線程，并且每個線程都執(zhí)行由核函數(shù)指定的語句。一個兩層的線程層次結(jié)構(gòu)如下圖所示，由線程塊和線程塊網(wǎng)格構(gòu)成。

??又一個內(nèi)核所產(chǎn)生的所有線程統(tǒng)稱為一個網(wǎng)格。同一網(wǎng)格中的所有線程共享相同的全局內(nèi)存空間。一個網(wǎng)格由多個線程塊構(gòu)成，一個線程塊包含一組線程，同一線程塊內(nèi)的線程協(xié)作可以通過同步、共享內(nèi)存等方式來實現(xiàn)，而不同塊內(nèi)的線程不能協(xié)作。
??線程依靠blockIdx(線程塊在線程格內(nèi)的索引)以及threadIdx(塊內(nèi)的線程索引)這兩個坐標(biāo)變量來區(qū)分彼此。這些變量是核函數(shù)中需要預(yù)初始化的內(nèi)置變量。當(dāng)執(zhí)行一個核函數(shù)時，CUDA運行時為每個線程分配坐標(biāo)變量blockIdx和threadIdx?；谶@些坐標(biāo)，我們可以將部分?jǐn)?shù)據(jù)分配給不同的線程。該坐標(biāo)變量是基于uint3定義的CUDA內(nèi)置的向量類型，是一個包含3個無符號整數(shù)的結(jié)構(gòu)，可以通過x、y、z三個字段來指定。
??CUDA可以組織三維的網(wǎng)格和塊。線程層次結(jié)構(gòu)如上面的圖所示，其結(jié)構(gòu)是一個包含二維塊的二維網(wǎng)格。網(wǎng)格和塊的維度由blockDim(線程塊的維度，用每個線程塊中的線程數(shù)來表示)、gridDim(線程格的維度，用每個線程格中的線程數(shù)來表示)兩個內(nèi)置變量指定，它們是dim3類型的變量，是基于uint3定義的整數(shù)型向量，用來表示維度。當(dāng)定義一個dim3類型的變量時，所有未指定的元素都被初始化為1.dim3類型變量中的每個組件可以通過它的x、y、z字段獲得。
??在CUDA程序中有兩組不同的網(wǎng)格和塊變量：手動定義的dim3數(shù)據(jù)類型和預(yù)定義的uint3數(shù)據(jù)類型。在主機端，作為內(nèi)核調(diào)用的一部分，可以使用dim3數(shù)據(jù)類型定義一個網(wǎng)格和塊的維度。當(dāng)執(zhí)行核函數(shù)時， CUDA運行時會生成相應(yīng)的內(nèi)置預(yù)初始化的網(wǎng)格、塊和線程變量，它們在核函數(shù)內(nèi)均可被訪問到且為uint3類型。手動定義的dim3類型的網(wǎng)格和塊變量僅在主機端可見，而uint3類型的內(nèi)置預(yù)初始化的網(wǎng)格和塊變量僅在設(shè)備端可見。
??由于一個內(nèi)核啟動的網(wǎng)格和塊的維數(shù)會影響性能，這一結(jié)構(gòu)為程序員優(yōu)化程序提供了一個額外的途徑。對于一個給定的數(shù)據(jù)大小，確定網(wǎng)格和塊尺寸的一般步驟為：確定塊的個數(shù)、在已知數(shù)據(jù)大小和塊大小的基礎(chǔ)上計算網(wǎng)格維度。網(wǎng)格和塊的維度存在幾個限制因素：對于塊大小的一個主要限制因素就是可利用的計算資源，如寄存器，共享內(nèi)存等。某些限制可以通過查詢GPU設(shè)備。

啟 動 一 個 CUDA 核 函 數(shù)
??C語言函數(shù)調(diào)用語句：function_name(argument list);而CUDA內(nèi)核調(diào)用是對C語言函數(shù)調(diào)用語句的延申，kernal_name <<<grid,block>>>(argument list);其中<<<>>>運算符內(nèi)是核函數(shù)的執(zhí)行配置。執(zhí)行配置的第一個值是網(wǎng)絡(luò)維度，也就是啟動塊的數(shù)目，第二個值是塊維度，也就是每個塊中線程的數(shù)目。通過指定網(wǎng)格和塊的維度，可以進(jìn)行內(nèi)核中線程的數(shù)目以及內(nèi)核中使用的線程布局的配置。
??由于數(shù)據(jù)在全局內(nèi)存中是線性存儲的，因為可以用變量blockIdx.x和threadIdx.x來進(jìn)行以下操作：在網(wǎng)格中標(biāo)識一個唯一的線程、建立線程和數(shù)據(jù)元素之間的映射關(guān)系。不同于C語言的函數(shù)調(diào)用，所有的CUDA核函數(shù)的啟動都是異步的。CUDA內(nèi)核調(diào)用完成后，控制權(quán)立刻返回給CPU?？梢酝ㄟ^調(diào)用以下函數(shù)來強制主機端程序等待所有的核函數(shù)執(zhí)行結(jié)束：cudaError_t cudaMemcpy(void* dst,const void* src,size_t count,cudaMemcpykind kind);

編 寫 核 函 數(shù)
??核函數(shù)是在設(shè)備端執(zhí)行的代碼。在核函數(shù)中，需要為一個線程規(guī)定要進(jìn)行的計算以及要進(jìn)行的數(shù)據(jù)訪問。當(dāng)核函數(shù)被調(diào)用時，許多不同的CUDA線程并行執(zhí)行同一個計算任務(wù)。核函數(shù)必須有一個void返回類型。
??函數(shù)類型限定符指定了一個函數(shù)在主機上執(zhí)行還是在設(shè)備上執(zhí)行，以及可被主機調(diào)用還是被設(shè)備調(diào)用，下表總結(jié)了CUDA C程序中的函數(shù)類型限定符。

限定符執(zhí)行調(diào)用備注

global	在設(shè)備端執(zhí)行	可從主機端調(diào)用，也可以從計算能力為3的設(shè)備中調(diào)用	必須有一個void返回類型
device	在設(shè)備端執(zhí)行	僅能從設(shè)備端調(diào)用
host	在主機端執(zhí)行	僅能從主機端調(diào)用	可以省略

??__device__和__host__限定符可以一齊使用，這樣函數(shù)可以同時在主機和設(shè)備端進(jìn)行編譯。
??CUDA核函數(shù)的限制：只能訪問設(shè)備內(nèi)存、必須具有void返回類型、不支持可變數(shù)量的參數(shù)、不支持靜態(tài)變量、顯示異步行為。

驗 證 核 函 數(shù)
??除了許多可用的調(diào)試工具外，還有兩個非常簡單實用的方法可以驗證核函數(shù)：首先，你可以在Fermi及更高版本的設(shè)備端的核函數(shù)中使用printf函數(shù)；其次，可以將執(zhí)行參數(shù)設(shè)置為<<<1,1>>>,因此強制用一個塊和一個線程執(zhí)行核函數(shù)，這模擬了串行執(zhí)行程序。

處 理 錯 誤
??由于許多CUDA調(diào)用是異步的，所有有時很難確定某個錯誤是由哪一步程序引起的。定義一個錯誤處理宏封裝所有的CUDA API調(diào)用，這簡化了錯誤檢查過程：

#define CHECK(call) {const cudaError_t error = call;if(error != cudaSuccess){printf("Error: %s:%d",__FILE__,__LINE__);printf("code:%d, reson:%s\n",error,cudaGetErrorString(error));exit(1);} }

??完成錯誤檢查宏的定義之后，就可以在代碼中使用宏。比如：CHECK(cudaMemcpy(d_C,gpuRef,nBytes,cudaMemcpyHostToDevice));如果內(nèi)存拷貝或之前的異步操作產(chǎn)生了錯誤，這個宏會報告錯誤代碼，并且輸出一個可讀信息，然后停止程序。再比如：CHECK(cudaDeviceSynchronize())會阻塞主機端線程的運行直到設(shè)備端所有的請求都結(jié)束，并確保最后的核函數(shù)啟動部分不會出錯。

給核函數(shù)計時

用 CPU 計 時 器 計 時
??可以使用gettimeofday系統(tǒng)調(diào)用來創(chuàng)建一個CPU計時器，以獲取系統(tǒng)的時鐘時間，它需要包含sys/time.h頭文件。

double cpuSecord(){struct timeval tp;gettimeofday(&tp,NULL);return ((double)tp.tv_sec + (double)tp.tv_usec*1.e-6); }

??這樣可以用cpuSecond函數(shù)來測試核函數(shù)：

double iStart = cpuSecord(); kernel_name<<<grid,block>>>(argument list); cudaDeviceSynchronize(); double iElaps = cpuSecond() - iStart;

用 nvprof 工 具 計 時
??自CUDA5.0以來，NVIDIA提供了一個名為nvprof的命令行分析工具，可以幫助從應(yīng)用程序的CPU和GPU活動情況中獲取時間線信息，其包括內(nèi)核執(zhí)行、內(nèi)存?zhèn)鬏斠约癈UDA API的調(diào)用。

組織并行線程

??如果使用了合適的網(wǎng)絡(luò)和塊大小來正確地組織線程，那么可以對內(nèi)核性能產(chǎn)生很大的影響。對于矩陣運算，傳統(tǒng)的方法是在內(nèi)核中使用一個包含二維網(wǎng)格與二維塊的布局來組織線程，但是，這種傳統(tǒng)的方法無法獲得最佳性能。在矩形加法中使用由二維線程塊構(gòu)成的二維網(wǎng)絡(luò)、由一維線程塊構(gòu)成的一維網(wǎng)格、由一維線程塊構(gòu)成的二維網(wǎng)格等布局可以取得不同性能的結(jié)果。
使 用 塊 和 線 程 建 立 矩 陣 索 引
??在一個矩陣加法核函數(shù)中，一個線程通常被分配一個數(shù)據(jù)元素來處理。首先要完成的任務(wù)是使用塊和線程索引從全局內(nèi)存中訪問指定的數(shù)據(jù)。通常情況下，對一個二維示例來說，需要管理線程和塊索引、矩陣中給定點的坐標(biāo)、全局線性內(nèi)存中的偏移量三種索引。對于一個給定的線程，首先可以通過把線程和塊索引映射到矩陣坐標(biāo)上來獲取線程塊和縣城索引的全局內(nèi)存偏移量，然后將這些矩陣坐標(biāo)映射到全局內(nèi)存的存儲單元中。
??第一步，可以用以下公式把線程和塊索引映射到矩陣坐標(biāo)上：

ix = threadIdx.x + blockIdx.x * blockDim.x;iy = threadIdx.y + blockIdx.y * blockDim.y;

??第二步，可以用以下公式把矩陣坐標(biāo)映射到全局內(nèi)存中的索引/存儲單元上：

idx = iy * nx + ix;

??其中，printThreadIndo函數(shù)被用于輸出關(guān)于每個線程的以下信息：線程索引、塊索引、矩陣坐標(biāo)、線性全局內(nèi)存偏移、相應(yīng)元素的值。

設(shè)備管理

??查詢和管理GPU設(shè)備的兩種方法：
??>CUDA運行時API函數(shù)；
??>NVIDIA系統(tǒng)管理界面（nvidia-smi）命令行實用程序。
使 用 運 行 時 API 查 詢 GPU 信 息
??在CUDA運行時API中由很多函數(shù)可以幫助管理這些設(shè)備。可以使用以下函數(shù)來查詢關(guān)于GPU設(shè)備的所有信息：

cudaError_t cudaGetDeviceProperties(cudaDeviceProp* prop,int device);

??cudaDeviceProp結(jié)構(gòu)體返回GPU設(shè)備的屬性。
確 定 最 優(yōu) GPU
??一些系統(tǒng)支持多GPU。在每個GPU都不同的情況下，選擇性能最好的GPU運行核函數(shù)是非常重要的。通過比較GPU包含的多處理器的數(shù)量選出計算能力最佳的GPU。可以使用以下代碼來選擇計算能力最優(yōu)的設(shè)備：

int numDevices = 0; cudaGetDeviceCount(&numDevices); if(numDevices > 1) {int maxMultiprocessors = 0, maxDevice = 0;for(int device=0;device <numDevices; device++){cudaDeviceProp props;cudaGetDeviceProperties(&props,device);if(maxMultiprocessors <props.multiProcessorCount){maxMultiprocessors = props.multiProcessorCount;maxDevice = device;}} }

使 用 nvidia-smi 查 詢 GPU 信 息
??nvidia-smi是一個命令行工具，用于管理和監(jiān)控GPU設(shè)備，并允許查詢和修改設(shè)備狀態(tài)。可以從命令行調(diào)用nvidia-smi。

總結(jié)

以上是生活随笔為你收集整理的CUDA C编程（二）CUDA编程模型的全部內(nèi)容，希望文章能夠幫你解決所遇到的問題。

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