《并行程序設計》實驗指導書之五實驗5.1運行CUDA向量加法的樣例代碼實驗目的1．掌握cuda開發環境的使用與配置；2．掌握利用cuda調試和運行；3．掌握cuda并行計算的原理。實驗要求1．熟練掌握C++語言；2．掌握VisualStudio*.NET*集成開發環境的使用；3．掌握cuda開發環境。實驗原理CUDA(Compute

Unified

Device

Architecture)是由NVIDIA公司創立的基于他們公司生產的圖形處理器GPUs（Graphics

Processing

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

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

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

的矩陣，B為

的矩陣，那么稱

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

，其中矩陣C中的第

i行第

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

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