《并行程序設(shè)計(jì)》實(shí)驗(yàn)指導(dǎo)書(shū)之五

《并行程序設(shè)計(jì)》實(shí)驗(yàn)指導(dǎo)書(shū)之五實(shí)驗(yàn)5.1運(yùn)行CUDA向量加法的樣例代碼實(shí)驗(yàn)?zāi)康?．掌握cuda開(kāi)發(fā)環(huán)境的使用與配置；2．掌握利用cuda調(diào)試和運(yùn)行；3．掌握cuda并行計(jì)算的原理。實(shí)驗(yàn)要求1．熟練掌握C++語(yǔ)言；2．掌握VisualStudio*.NET*集成開(kāi)發(fā)環(huán)境的使用；3．掌握cuda開(kāi)發(fā)環(huán)境。實(shí)驗(yàn)原理CUDA(Compute

Unified

Device

Architecture)是由NVIDIA公司創(chuàng)立的基于他們公司生產(chǎn)的圖形處理器GPUs（Graphics

Processing

Units，可以通俗的理解為顯卡）的一個(gè)并行計(jì)算平臺(tái)和編程模型。通過(guò)CUDA，GPUs可以很方便地被用來(lái)進(jìn)行通用計(jì)算（有點(diǎn)像在CPU中進(jìn)行的數(shù)值計(jì)算等等）。在沒(méi)有CUDA之前，GPUs一般只用來(lái)進(jìn)行圖形渲染（如通過(guò)OpenGL，DirectX）。有了CUDA之后，開(kāi)發(fā)人員可以通過(guò)調(diào)用CUDA的API，來(lái)進(jìn)行并行編程，達(dá)到高性能計(jì)算的目的。CUDA中常見(jiàn)的概念及名稱：主機(jī)：CPU及系統(tǒng)的內(nèi)存（內(nèi)存條）；設(shè)備：

GPU及GPU本身的顯示內(nèi)存；線程(Thread)：資源調(diào)度單元，一般交給GPU的一個(gè)核去處理（有一維，二維，三維）。線程塊(Block)：多個(gè)線程組合在一起就是線程塊；各block并行執(zhí)行，互相之間不能通信，執(zhí)行時(shí)無(wú)法指定順序；線程塊有數(shù)量限制，最多可以有65535個(gè)線程塊。線程格(Grid)：由多個(gè)線程塊組成。線程、線程塊、線程格的邏輯結(jié)構(gòu)如下所示：線程束：線程束是一個(gè)集合，其中包含32個(gè)線程。這個(gè)集合里的線程被“組合在一起”并且“步調(diào)一致”地進(jìn)行執(zhí)行。對(duì)程序的每一行，線程束里的線程都將在不同數(shù)據(jù)上分別執(zhí)行。GPU內(nèi)存分類1.全局內(nèi)存：就是指設(shè)備內(nèi)存。2.共享內(nèi)存：存儲(chǔ)在全局內(nèi)存中，使用時(shí)要添加關(guān)鍵字__shared__到變量聲明中。CUDA對(duì)GPU上啟動(dòng)的每個(gè)線程塊，都會(huì)保存一個(gè)共享變量的副本。在同一個(gè)線程塊內(nèi)的所有線程都要共享這塊內(nèi)存，但線程卻不能看到因此也不能修改其他線程塊的副本。這就實(shí)現(xiàn)了一個(gè)線程塊中的多個(gè)線程能在計(jì)算上進(jìn)行通信以及協(xié)作。3.常量?jī)?nèi)存：存儲(chǔ)在全局內(nèi)存中，使用時(shí)要添加關(guān)鍵字__constant__到變量聲明中。常量?jī)?nèi)存用來(lái)保存在核函數(shù)執(zhí)行期間不會(huì)發(fā)生變化的數(shù)據(jù)，變量是只讀的，通過(guò)特殊的處理方式，有時(shí)候使用常量?jī)?nèi)存替代全局內(nèi)存可以減少內(nèi)存帶寬，對(duì)性能提升有幫助；當(dāng)需要拷貝數(shù)據(jù)到常量?jī)?nèi)存中必須使用cudaMemcpyToSymbol，如果使用cudaMemcpy會(huì)將數(shù)據(jù)復(fù)制到全局內(nèi)存。4.紋理內(nèi)存：存儲(chǔ)在全局內(nèi)存中。面向訪問(wèn)內(nèi)存具有空間聚簇性的程序(例如圖像處理方面的計(jì)算程序)設(shè)計(jì)，互相臨近的線程所讀取的數(shù)據(jù)在物理存儲(chǔ)上也是臨近的，可以減少訪存次數(shù)，節(jié)約帶寬，以此來(lái)提升效率。紋理內(nèi)存有一維與二維兩種：一維紋理內(nèi)存的聲明方式是texture<類型>，使用cudaBindTexture()函數(shù)綁定紋理內(nèi)存，cudaUnbindTexture()函數(shù)解除綁定，讀取內(nèi)存數(shù)據(jù)時(shí)要使用tex1D()函數(shù)；二維紋理內(nèi)存的聲明方式是texture<類型,數(shù)字>，使用cudaBindTexture2D()函數(shù)綁定紋理內(nèi)存，cudaUnbindTexture()函數(shù)解除綁定，讀取內(nèi)存數(shù)據(jù)時(shí)要使用tex2D()函數(shù)。5.固定內(nèi)存：存儲(chǔ)在主機(jī)內(nèi)存中，又稱為不可分頁(yè)或頁(yè)鎖定內(nèi)存。對(duì)于固定內(nèi)存，操作系統(tǒng)不會(huì)對(duì)其分頁(yè)，也不會(huì)交換到磁盤(pán)上，可以確保它始終駐留在物理內(nèi)存上。在編寫(xiě)程序時(shí)可以直接訪問(wèn)這塊物理地址，因?yàn)樗粫?huì)被破壞或遷移。固定內(nèi)存是為了提高訪問(wèn)速度而被設(shè)計(jì)出來(lái)的。GPU如果知道主機(jī)中的物理地址，就可通過(guò)DMA方式來(lái)復(fù)制主機(jī)與GPU之間的數(shù)據(jù)。當(dāng)然，用戶編寫(xiě)程序時(shí)要注意不可一味使用固定內(nèi)存，這樣將導(dǎo)致物理內(nèi)存迅速消耗完。在使用固定內(nèi)存時(shí)，一般將調(diào)用cudaMemcpy()函數(shù)時(shí)使用的源內(nèi)存或目的內(nèi)存設(shè)置為固定內(nèi)存，在調(diào)用完后不再需要時(shí)立即釋放掉。分配固定內(nèi)存需要使用cudaHostAlloc()函數(shù)；釋放固定內(nèi)存需要使用cudaFreeHost()函數(shù)。注意復(fù)制固定內(nèi)存是異步的方式。核函數(shù)（Kernel）核函數(shù)在GPU上面運(yùn)行，在GPU上運(yùn)行的函數(shù)都可視作是核函數(shù)；核函數(shù)可以使用標(biāo)識(shí)符來(lái)修飾，通常使用的是__global__標(biāo)識(shí)符。調(diào)用的方法與C語(yǔ)言有所區(qū)別，通常是通過(guò)<<<參數(shù)n1,參數(shù)n2>>>，調(diào)用時(shí)必須聲明內(nèi)核函數(shù)的執(zhí)行參數(shù)；核函數(shù)需要通過(guò)線程格（Grid）來(lái)組織，線程格下包含若干線程塊（block），而線程塊下又包含若干個(gè)線程（thread）；核函數(shù)的執(zhí)行單位是線程塊（block）；編寫(xiě)程序時(shí)要注意的一點(diǎn)是，對(duì)kernel函數(shù)中所需要使用的數(shù)組或變量，一定要在調(diào)用前提前分配好空間，否則在GPU進(jìn)行計(jì)算的時(shí)候會(huì)發(fā)生錯(cuò)誤，例如越界錯(cuò)誤，也有可能導(dǎo)致藍(lán)屏。在編寫(xiě)程序需要知道如下這些CUDA對(duì)C語(yǔ)言的擴(kuò)展：1.指定函數(shù)執(zhí)行器件的擴(kuò)展：定義了函數(shù)類型限定符，可用來(lái)確定函數(shù)是在CPU上執(zhí)行或是在GPU上執(zhí)行，以及是從CPU上調(diào)用還是從GPU上調(diào)用；有如下幾種函數(shù)類型限定符：__device__，由它所修飾的函數(shù)從GPU上調(diào)用，且在GPU上執(zhí)行，這樣的函數(shù)能夠由__device__或__global__修飾的函數(shù)調(diào)用。由它所修飾的函數(shù)的使用有限制，比如不允許使用函數(shù)指針；__global__，由它所修飾的函數(shù)從CPU上調(diào)用，且在GPU上執(zhí)行，也就是前面提到過(guò)的內(nèi)核(kernel)函數(shù)；它只能由主機(jī)調(diào)用，不是一個(gè)完整的程序，只是表示數(shù)據(jù)并行的步驟，其指令流由多個(gè)線程執(zhí)行；__host__，由它所修飾的函數(shù)從CPU上調(diào)用，且在CPU上執(zhí)行，這與傳統(tǒng)的C函數(shù)沒(méi)什么區(qū)別；2.變量存儲(chǔ)位置的擴(kuò)展定義了變量類型限定符，在傳統(tǒng)的CPU程序中，變量的存儲(chǔ)位置是由編譯器來(lái)負(fù)責(zé)完成的，編寫(xiě)程序的人并不需要指定存儲(chǔ)位置；但在CUDA架構(gòu)中，需要用變量類型限定符來(lái)規(guī)定變量的存儲(chǔ)位置，有以下幾種可供選擇：GPU里的緩存、共享存儲(chǔ)器、寄存器；設(shè)備的顯存；主機(jī)的內(nèi)存等等?；谶@些，CUDA中抽象出8種不同的存儲(chǔ)器，下面列出的是常用的變量類型限定符以及其含義。__device__：這個(gè)修飾符聲明的數(shù)據(jù)的存放位置是顯存，主機(jī)通過(guò)運(yùn)行時(shí)庫(kù)可以對(duì)其進(jìn)行訪問(wèn)，所有的線程都可以訪問(wèn)這樣聲明的數(shù)據(jù)；__shared__：這個(gè)修飾符聲明的數(shù)據(jù)的存放位置是共享存儲(chǔ)器，不是所有線程都能訪問(wèn)這樣聲明的數(shù)據(jù)，只有它所在的塊里的線程可以對(duì)其進(jìn)行訪問(wèn)；__constant__：這個(gè)修飾符聲明的數(shù)據(jù)的存放位置是常量存儲(chǔ)器，主機(jī)通過(guò)運(yùn)行時(shí)庫(kù)可以對(duì)其進(jìn)行訪問(wèn)，所有的線程都可以訪問(wèn)這樣聲明的數(shù)據(jù)；3.執(zhí)行配置擴(kuò)展定義了執(zhí)行配置運(yùn)算符<<<>>>，調(diào)用內(nèi)核函數(shù)時(shí)要將執(zhí)行配置傳遞過(guò)去，<<<>>>就是用來(lái)傳遞執(zhí)行配置的。執(zhí)行配置的組成是4個(gè)參數(shù)：1.網(wǎng)格大小2.塊大小，3.共享存儲(chǔ)器大?。梢院雎?，默認(rèn)為0），4.執(zhí)行的流（可以忽略，默認(rèn)為0）。4.內(nèi)建變量擴(kuò)展內(nèi)建變量用來(lái)在運(yùn)行時(shí)獲取線程索引以及塊和網(wǎng)格的尺寸等信息。CUDA里一共有5個(gè)內(nèi)建變量：1.gridDim，包含grid的維度。這是一個(gè)結(jié)構(gòu)體，它包含三個(gè)元素x，y，z，用來(lái)表示網(wǎng)格在這幾個(gè)方向上的尺寸，雖然CUDA設(shè)計(jì)的是三維，但目前只能使用二維；2.blockDim，包含block的維度。和gridDim一樣也是一個(gè)由x，y，z三個(gè)元素組成的結(jié)構(gòu)體，它表示的是塊在這幾個(gè)方向上的尺寸；3.blockIdx，包含網(wǎng)格的緯度。和gridDim一樣也是一個(gè)由x，y，z三個(gè)元素組成的結(jié)構(gòu)體，這幾個(gè)元素分別表示當(dāng)前線程所在的塊在網(wǎng)格中x，y，z三個(gè)方向上的索引；4.threadIdx，和gridDim一樣也是一個(gè)由x，y，z三個(gè)元素組成的結(jié)構(gòu)體，分別表示當(dāng)前線程在其所在的塊中x，y，z三個(gè)方向上的索引；5.warpSize，它表明warp的尺寸。常用的GPU內(nèi)存函數(shù)cudaMalloc()(1)原型：cudaError_tcudaMalloc(void**devPtr,size_tsize)(2)參數(shù)：devPtr——指向分配的設(shè)備內(nèi)存；size——需要分配的內(nèi)存大小；(3)作用：與malloc()函數(shù)類似，只是此函數(shù)在GPU的內(nèi)存中進(jìn)行內(nèi)存分配。從設(shè)備上分配size比特的連續(xù)內(nèi)存并返回一個(gè)指向此內(nèi)存空間的指針*devPtr。分配的內(nèi)存可用于存儲(chǔ)任何類型的變量；(4)返回值：分配成功返回cudaSuccess，萬(wàn)一分配失敗的話返回cudaErrorMemoryAllocationcudaMemcpy()原型：cudaError_tcudaMemcpy(void*dst,constvoid*src,size_tcount, enumcudaMemcpyKindkind)參數(shù)：dst

——目標(biāo)內(nèi)存地址，src——源內(nèi)存地址，count——拷貝的字節(jié)數(shù)目，kind——數(shù)據(jù)拷貝的方向(3)作用：與memcpy()函數(shù)類似，從src指針指向的內(nèi)存空間中拷貝count字節(jié)數(shù)據(jù)到dst指針指向的空間中。kind有以下幾種取值可能：cudaMemcpyHostToHost，cudaMemcpyHostToDevice，cudaMemcpyDeviceToHost，或cudaMemcpyDeviceToDevice，指明了數(shù)據(jù)拷貝的方向。如果dst和src指向的空間不符合kind的方向，由此產(chǎn)生的行為是沒(méi)有定義的。以同步方式執(zhí)行，即當(dāng)函數(shù)返回時(shí)，復(fù)制操作就已經(jīng)完成了，并且在輸出緩沖區(qū)中包含了復(fù)制進(jìn)去的內(nèi)容。(4)返回值：cudaSuccess，cudaErrorInvalidValue，cudaErrorInvalidDevicePointer，cudaErrorInvalidMemcpyDirection。cudaFree()(1)原型：cudaError_tcudaFree(void*devPtr)(2)參數(shù)：devPtr——指向要釋放的內(nèi)存空間；(3)作用：與free()函數(shù)類似，只是此函數(shù)釋放的是cudaMalloc()分配的內(nèi)存；(4)返回值：cudaSuccess，cudaErrorInvalidDevicePointer，cudaErrorInitializationError實(shí)驗(yàn)內(nèi)容實(shí)驗(yàn)步驟：打開(kāi)VisualStudio，新建一個(gè)項(xiàng)目，模版要選擇NVIDIA下的CUDA；在生成項(xiàng)目中例子代碼就是向量加法，直接編譯運(yùn)行，記錄實(shí)驗(yàn)結(jié)果；實(shí)驗(yàn)4.2基于CUDA優(yōu)化矩陣乘法實(shí)驗(yàn)?zāi)康?．掌握基于cuda優(yōu)化科學(xué)計(jì)算的方法;2．比較各種cuda優(yōu)化方式的性能提升效果；3．掌握利用加速比、運(yùn)行時(shí)間、效率等測(cè)度分析并行程序性能；實(shí)驗(yàn)要求1．熟練掌握C++語(yǔ)言；2、掌握VisualStudio*.NET*集成開(kāi)發(fā)環(huán)境的使用；3．掌握cuda開(kāi)發(fā)環(huán)境；實(shí)驗(yàn)原理設(shè)A為

的矩陣，B為

的矩陣，那么稱

的矩陣C為矩陣A與B的乘積，記作

，其中矩陣C中的第

i行第

j列元素可以表示為：本次實(shí)驗(yàn)中將A矩陣和B矩陣假設(shè)成大小相同的方形矩陣，這不影響性能分析。同時(shí)不做算法上層面的優(yōu)化，直接使用三重for循環(huán)。程序邏輯圖matgen函數(shù)的輸入是矩陣首地址以及行數(shù)、列數(shù)，產(chǎn)生介于0與1之間的浮點(diǎn)隨機(jī)數(shù)填充矩陣中的每個(gè)元素；matmult函數(shù)的輸入是需要相乘的兩個(gè)矩陣，對(duì)它們運(yùn)行矩陣乘法程序。它是傳統(tǒng)的CPU程序。既可以用來(lái)與其它方式進(jìn)行矩陣乘法所得到的結(jié)果進(jìn)行對(duì)比，判斷正確與否，也可以用來(lái)作為性能對(duì)比的基準(zhǔn)，為了提高精確度，中間結(jié)果使用雙精度浮點(diǎn)類型來(lái)存儲(chǔ)；compare_mat函數(shù)的輸入是需要進(jìn)行比較的兩個(gè)矩陣，用來(lái)計(jì)算兩個(gè)矩陣之間的平均相對(duì)誤差以及最大相對(duì)誤差，并打印出比較的結(jié)果；matmultCUDA函數(shù)的輸入是需要相乘的兩個(gè)矩陣，使用GPU實(shí)現(xiàn)矩陣乘法。它首先從顯卡內(nèi)存里申請(qǐng)用于存放矩陣的空間，之后將矩陣數(shù)據(jù)從主內(nèi)存拷貝到顯卡內(nèi)存中。在進(jìn)行內(nèi)存拷貝時(shí)，如果用cudaMemcpy函數(shù)的話，將造成每個(gè)row分開(kāi)進(jìn)行拷貝，這樣的話就需要多次調(diào)用cudaMemcpy函數(shù)，這會(huì)降低程序的效率。所以，在這里使用cudaMemcpy2D

函數(shù)來(lái)拷貝二維數(shù)組，這樣就能夠只執(zhí)行一次函數(shù)調(diào)用就完成拷貝的工作。然后調(diào)用核函數(shù)matMultCUDA，它先通過(guò)bid以及tid計(jì)算出某個(gè)thread需要計(jì)算的行和列，將計(jì)算所得寫(xiě)入結(jié)果中。下面對(duì)程序做一些解釋。1.在使用cuda進(jìn)行計(jì)算時(shí)，用到了下面的代碼：for(j=tid;j<n;j+=blockDim.x){ floatt=0; floaty=0; for(i=0;i<n;i++){ floatr; y-=data[i]*b[i*ldb+j]; r=t-y; y=(r-t)+y; t=r; } c[row*ldc+j]=t;}之所以要這樣進(jìn)行累加是基于下面的考慮：CPU上面的計(jì)算使用了64位浮點(diǎn)數(shù)來(lái)累加中間結(jié)果，而在GPU上卻只能用32位浮點(diǎn)數(shù)進(jìn)行累加，當(dāng)累加的數(shù)據(jù)很大時(shí)，就會(huì)產(chǎn)生舍入誤差。針對(duì)這個(gè)問(wèn)題，CUDA在進(jìn)行加、減、乘法的浮點(diǎn)運(yùn)算時(shí)符合IEEE754規(guī)定的精確度，所以使用Kahan'sSummationFormula來(lái)提高精確度（思想是把每次截?cái)嗟恼`差在一個(gè)較小的數(shù)字里進(jìn)行累加）。2.在調(diào)用核函數(shù)時(shí)使用了這樣的方式：函數(shù)名稱<<<block數(shù)目,thread數(shù)目,sharedmemory大小>>>(參數(shù)...)；因?yàn)橹挥性谕粋€(gè)block中的線程可以共享內(nèi)存，因此一行只能由同一個(gè)block里的線程來(lái)進(jìn)行計(jì)算。另外需要共享內(nèi)存存放整個(gè)row的數(shù)據(jù)。下面的代碼可以把每行的數(shù)據(jù)放到共享內(nèi)存中，便于后面使用共享內(nèi)存進(jìn)行計(jì)算：extern__shared__floatdata[];constinttid=threadIdx.x;constintrow=blockIdx.x;inti,j;for(i=tid;i<n;i+=blockDim.x){ data[i]=a[row*lda+i];}3.GPU在讀取內(nèi)存時(shí)會(huì)選擇一個(gè)固定倍數(shù)的地址開(kāi)始，例如64bytes的整數(shù)倍，這樣可以在最大程度提升效率。但是實(shí)驗(yàn)中使用的矩陣未必是64的整數(shù)倍，這影響了效率。為了消除這個(gè)影響，代碼在內(nèi)存分配時(shí)使用了CUDA提供的cudaMallocPitch函數(shù)，它在配置內(nèi)存時(shí)會(huì)自動(dòng)使用最佳倍數(shù)。使用下面的代碼來(lái)優(yōu)化cudaMalloc部分：size_tpitch_a,pitch_b,pitch_c;

cudaMallocPitch((void**)&ac,&pitch_a,sizeof(float)*n,n);

cudaMallocPitch((void**)&bc,&pitch_b,sizeof(float)*n,n);

cudaMallocPitch((void**)&cc,&pitch_c,sizeof(float)*n,n);cudaMallocPitch函數(shù)會(huì)選擇合適的倍數(shù)來(lái)配置內(nèi)存，并傳回配置寬度。將矩陣拷貝到顯卡內(nèi)存時(shí)要使用傳回的寬度，代碼如下：cudaMemcpy2D(ac,sizeof(float)*n,a,sizeof(float)*lda,sizeof(float)*n,n,cudaMemcpyHostToDevice);cudaMemcpy2D(bc,sizeof(float)*n,b,izeof(float)*ldb,sizeof(float)*n,n,cudaMemcpyHostToDevice);呼叫kernel的部分如下：matMultCUDA<<<n,NUM_THREADS,sizeof(float)*n>>>(ac,pitch_a/sizeof(float),bc,pitch_b/sizeof(float),cc,pitch_c/sizeof(float),n);把結(jié)果拷貝回主內(nèi)存時(shí)，也要用上傳回的寬度：cudaMemcpy2D(c,sizeof(float)*ldc,cc,pitch_c,sizeof(float)*n,n,cudaMemcpyDeviceToHost);實(shí)驗(yàn)內(nèi)容1.打開(kāi)VisualStudio，新建一個(gè)項(xiàng)目，模版要選擇NVIDIA下的CUDA；2.將實(shí)驗(yàn)代碼拷貝進(jìn)去，編譯運(yùn)行，進(jìn)行如下實(shí)驗(yàn)并記錄數(shù)據(jù)（實(shí)驗(yàn)報(bào)告中給出數(shù)據(jù)并繪圖）實(shí)驗(yàn)一：固定矩陣大小為250x250，調(diào)整塊大小NUM_THREADS，記錄程序運(yùn)算時(shí)間及運(yùn)算速度?；鶞?zhǔn)運(yùn)行時(shí)間(s)：基準(zhǔn)運(yùn)算速度(GFLOPS)：NUM_THREADS1632641282565