《并行程序設(shè)計》實驗指導(dǎo)書之五

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

Unified

Device

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

Processing

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

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

——目標內(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)生的行為是沒有定義的。以同步方式執(zhí)行，即當(dāng)函數(shù)返回時，復(fù)制操作就已經(jīng)完成了，并且在輸出緩沖區(qū)中包含了復(fù)制進去的內(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實驗內(nèi)容實驗步驟：打開VisualStudio，新建一個項目，模版要選擇NVIDIA下的CUDA；在生成項目中例子代碼就是向量加法，直接編譯運行，記錄實驗結(jié)果；實驗4.2基于CUDA優(yōu)化矩陣乘法實驗?zāi)康?．掌握基于cuda優(yōu)化科學(xué)計算的方法;2．比較各種cuda優(yōu)化方式的性能提升效果；3．掌握利用加速比、運行時間、效率等測度分析并行程序性能；實驗要求1．熟練掌握C++語言；2、掌握VisualStudio*.NET*集成開發(fā)環(huán)境的使用；3．掌握cuda開發(fā)環(huán)境；實驗原理設(shè)A為

的矩陣，B為

的矩陣，那么稱

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

，其中矩陣C中的第

i行第

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

函數(shù)來拷貝二維數(shù)組，這樣就能夠只執(zhí)行一次函數(shù)調(diào)用就完成拷貝的工作。然后調(diào)用核函數(shù)matMultCUDA，它先通過bid以及tid計算出某個thread需要計算的行和列，將計算所得寫入結(jié)果中。下面對程序做一些解釋。1.在使用cuda進行計算時，用到了下面的代碼：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;}之所以要這樣進行累加是基于下面的考慮：CPU上面的計算使用了64位浮點數(shù)來累加中間結(jié)果，而在GPU上卻只能用32位浮點數(shù)進行累加，當(dāng)累加的數(shù)據(jù)很大時，就會產(chǎn)生舍入誤差。針對這個問題，CUDA在進行加、減、乘法的浮點運算時符合IEEE754規(guī)定的精確度，所以使用Kahan'sSummationFormula來提高精確度（思想是把每次截斷的誤差在一個較小的數(shù)字里進行累加）。2.在調(diào)用核函數(shù)時使用了這樣的方式：函數(shù)名稱<<<block數(shù)目,thread數(shù)目,sharedmemory大小>>>(參數(shù)...)；因為只有在同一個block中的線程可以共享內(nèi)存，因此一行只能由同一個block里的線程來進行計算。另外需要共享內(nèi)存存放整個row的數(shù)據(jù)。下面的代碼可以把每行的數(shù)據(jù)放到共享內(nèi)存中，便于后面使用共享內(nèi)存進行計算：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ù)的地址開始，例如64bytes的整數(shù)倍，這樣可以在最大程度提升效率。但是實驗中使用的矩陣未必是64的整數(shù)倍，這影響了效率。為了消除這個影響，代碼在內(nèi)存分配時使用了CUDA提供的cudaMallocPitch函數(shù)，它在配置內(nèi)存時會自動使用最佳倍數(shù)。使用下面的代碼來優(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ù)會選擇合適的倍數(shù)來配置內(nèi)存，并傳回配置寬度。將矩陣拷貝到顯卡內(nèi)存時要使用傳回的寬度，代碼如下：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)存時，也要用上傳回的寬度：cudaMemcpy2D(c,sizeof(float)*ldc,cc,pitch_c,sizeof(float)*n,n,cudaMemcpyDeviceToHost);實驗內(nèi)容1.打開VisualStudio，新建一個項目，模版要選擇NVIDIA下的CUDA；2.將實驗代碼拷貝進去，編譯運行，進行如下實驗并記錄數(shù)據(jù)（實驗報告中給出數(shù)據(jù)并繪圖）實驗一：固定矩陣大小為250x250，調(diào)整塊大小NUM_THREADS，記錄程序運算時間及運算速度?；鶞蔬\行時間(s)：基準運算速度(GFLOPS)：NUM_THREADS1632641282565