生活随笔
收集整理的這篇文章主要介紹了
CUDA编程指南阅读笔记（六）
小編覺得挺不錯的,現在分享給大家,幫大家做個參考.
4.?CUDA C語言編程接口
? ? ? ? 接上文。
4.3 CUDA C Runtime
4.3.3 共享內存（Shared Memory）
? ? ? ? 共享內存是CUDA設備中非常重要的一個存儲區域，有效地使用共享內存可以充分利用CUDA設備的潛能，極大提升程序性能。那么，共享內存有哪些特點呢？ ? ? ? ? 1、共享內存（shared Memory）是集成在GPU處理器芯片上的（on-chip），因此相比于存在于顯存顆粒中的全局內存（global Memory）和本地內存（local Memory），它具有更高的傳輸帶寬，一般情況下，共享內存的帶寬大約是全局內存帶寬的7-10倍。 ? ? ? ? 2、共享內存的容量很小。根據NVIDIA官方文檔的說法，在計算能力1.x的設備中，每一個流多處理器（Streaming Multiprocessor）上的共享內存容量為16KB。對于計算能力2.x、3.0及3.5的設備該參數為48KB。因此共享內存是稀有資源。 ? ? ? ? 3、共享內存在物理上被劃分為很多塊，每一塊被稱為一個存儲體（bank）。在同一時刻，CUDA設備可以同時訪問多個存儲體。因此，如果一次針對共享內存的訪存操作需要讀取n個地址，而這n個地址恰好分布在n個不同的存儲體（bank）中，那么只需要一個存取周期就可以完成n個地址的訪存任務了。對于計算能力1.x的設備，共享內存被平均劃分為16個存儲體。而對于計算能力2.x、3.0及3.5的設備此參數為32。在共享內存中，相鄰兩塊32bit的數據分別屬于相鄰的兩個存儲體。存儲體每兩個時鐘周期可以傳輸32位數據。 ? ? ? ? 4、共享內存既可以靜態分配，也可以動態分配。 ? ? ? ? 從共享內存的這些特點中我們可以看出，它實際上相當于一個程序員可以操控的緩存（cache），下面，我們使用矩陣乘法的例子來說明如何有效使用共享內存。 ? ? ? ? 首先，我們使用最直觀的方法來完成矩陣乘法C = A x B：讀取A的每一行和B的每一列，順次完成計算任務。矩陣乘法的示意圖如下所示：
下面是矩陣乘法的CUDA C主要實現代碼： [cpp]?view plaincopy
?? ?? typedef?struct?{?? ????int?width;?? ????int?height;?? ????float?*elements;?? }?Matrix;?? ?? ?? #define?BLOCK_SIZE?16?? ?? ?? __global__?void?MatMulKernel(const?Matrix,?const?Matrix,?Matrix);?? ?? ?? ?? void?MatMul(const?Matrix?A,?const?Matrix?B,?Matrix?C)?{?? ?????? ????Matrix?d_A;?? ????d_A.width?=?A.width;?d_A.height?=?A.height;?? ????size_t?size?=?A.width?*?A.height?*?sizeof(float);?? ????cudaMalloc(&d_A.elements,?size);?? ????cudaMemcpy(d_A.elements,?A.elements,?size,?cudaMemcpyHostToDevice);?? ????Matrix?d_B;?? ????d_B.width?=?B.width;?d_B.height?=?B.height;?? ????size?=?B.width?*?B.height?*?sizeof(float);?? ????cudaMalloc(&d_B.elements,?size);?? ????cudaMemcpy(d_B.elements,?B.elements,?size,?cudaMemcpyHostToDevice);?? ?? ?????? ????Matrix?d_C;?? ????d_C.width?=?C.width;?d_C.height?=?C.height;?? ????size?=?C.width?*?C.height?*?sizeof(float);?? ????cudaMalloc(&d_C.elements,?size);?? ?? ?????? ????dim3?dimBlock(BLOCK_SIZE,?BLOCK_SIZE);?? ????dim3?dimGrid(B.width?/?dimBlock.x,?A.height?/?dimBlock.y);?? ????MatMulKernel<<<dimGrid,?dimBlock>>>(d_A,?d_B,?d_C);?? ?? ?????? ????cudaMemcpy(C.elements,?d_c.elements,?size,?cudaMemcpyDeviceToHost);?? ?? ?????? ????cudaFree(d_A.elements);?? ????cudaFree(d_B.elements);?? ????cudaFree(d_C.elements);?? }?? ?? ?? __global__?void?MatMulKernel(Matrix?A,?Matrix?B,?Matrix?C)?{?? ?????? ?????? ????float?Cvalue?=?0;?? ????int?row??=?blockIdx.y?*?blockDim.y?+?threadIdx.y;?? ????int?col?=?blockIdx.x?*?blockDim.x?+?threadIdx.xl?? ????for?(int?e?=?0;?e?<?A.width;?++e)?? ????????Cvalue?+=?A.elements[row?*?A.width?+?e]?*?B.elements[e?*?B.width?+?col];?? ????C.elements[row?*?C.width?+?col]?=?Cvalue;?? }?? 可以看出，為了計算矩陣C的任何一個元素，程序都需要從全局內存（global memory）中獲得矩陣A的一行和矩陣B的一列。因此，完成這一計算矩陣A被讀取了B.width次，矩陣B被讀取了A.height次。
? ? ? ? 現在我們來使用共享內存（shared memory）實現矩陣乘法。假設矩陣C可以被劃分為若干個較小的子方陣C
sub，我們使用一個線程塊（thread block）來負責某一子方陣的計算，線程塊中的每一個線程（thread）正好負責子方陣C
sub中一個元素的計算。這樣劃分后，任何一個結果子方陣C
sub'（尺寸為block_size * block_size）都是與該方陣具有相同行索引的尺寸為A.width * block_size的A的子矩陣A
sub和與該方陣具有相同列索引的尺寸為block_size * B.height的B的子矩陣B
sub相乘所得到。
? ? ? ? 為了匹配設備的計算資源，兩個子矩陣Asub和Bsub被劃分為盡可能多的分離的維度為block_size的子方陣，Csub的值便是這些子矩陣相乘后相加所得到的結果。子矩陣乘法的執行順序都是首先將它們從全局內存（global memory）拷貝到共享內存（shared memory）（線程塊中的每一個線程正好負責方陣一個元素的拷貝），然后由線程自己完成相應元素的計算任務，利用寄存器存儲局部結果，最后將寄存器的內容與新得到的計算結果依此累加起來得到最終運算結果并將其傳輸到全局內存（global memory）中。 ? ? ? ? 通過使用這種分治的計算策略，共享內存得到了很好的利用，采用這種方案計算完成時全局內存中矩陣A被訪問的次數為B.width / block_size，矩陣B被訪問的次數為A.height / block_size，很明顯，這為我們節省了非常多的全局內存帶寬。優化后的矩陣計算示意圖如下所示：
? ? ? ? 為了提升計算效率，我們為類型Matrix增加了一個成員變量stride。__device__函數用來獲得和設置子矩陣的元素。下面是優化后的代碼： [cpp]?view plaincopy
?? ?? typedef?struct?{?? ????int?width;?? ????int?height;?? ????int?stride;?? ????float*?elements;?? }?Matrix;?? ?? ?? __device__?float?GetElement(const?Matrix?A,?int?row,?int?col)?{?? ????return?A.elements[row?*?A.stride?+?col];?? }?? ?? ?? __device__?void?SetElement(Matrix?A,?int?row,?int?col,?float?value)?{?? ????A.elements[row?*?A.stride?+?col]?=?value;?? }?? ?? ?? ?? ?? __device__?Matrix?GetSubMatrix(Matrix?A,?int?row,?int?col)?{?? ????Matrix?Asub;?? ????Asub.width?=?BLOCK_SIZE;?? ????Asub.height?=?BLOCK_SIZE;?? ????Asub.stride?=?A.stride;?? ????Asub.elements?=?&A.elements[A.stride?*?BLOCK_SIZE?*?row?+?BLOCK_SIZE?*?col];?? ????return?Asub;?? }?? ?? ?? #define?BLOCK_SIZE?16?? ?? ?? __global__?void?MatMulKernel(const?Matrix,?const?Matrix,?Matrix);?? ?? ?? ?? void?MatMul(const?Matrix?A,?const?Matrix?B,?Matrix?C)?{?? ?????? ????Matrix?d_A;?? ????d_A.width?=?d_A.stride?=?A.width;?? ????d_A.height?=?A.height;?? ????size_t?size?=?A.width?*?A.height?*?sizeof(float);?? ????cudaMalloc(&d_A.elements,?size);?? ????cudaMemcpy(d_A.elements,?A.elements,?size,?cudaMemcpyHostToDevice);?? ????Matrix?d_B;?? ????d_B.width?=?d_B.stride?=?B.width;?? ????d_B.height?=?B.height;?? ????size?=?B.width?*?B.height?*?sizeof(float);?? ????cudaMalloc(&d_B.elements,?size);?? ????cudaMemcpy(d_B.elements,?B.elements,?size,?cudaMemcpyHostToDevice);?? ?? ?????? ????Matrix?d_C;?? ????d_C.width?=?d_C.stride?=?C.width;?? ????d_C.height?=?C.height;?? ????size?=?C.width?*?C.height?*?sizeof(float);?? ????cudaMalloc(&d_C.elements,?size);?? ?? ?????? ????dim3?dimBlock(BLOCK_SIZE,?BLOCK_SIZE);?? ????dim3?dimGrid(B.width?/?dimBlock.x,?A.height?/?dimBlock.y);?? ????MatMulKernel<<<dimGrid,?dimBlock>>>(d_A,?d_B,?d_C);?? ?? ?????? ????cudaMemcpy(C.elements,?d_C.elements,?size,?cudaMemcpyDeviceToHost);?? ?? ?????? ????cudaFree(d_A.elements);?? ????cudaFree(d_B.elements);?? ????cudaFree(d_C.elements);?? }?? ?? ?? __global__?void?MatMulKernel(Matrix?A,?Matrix?B,?Matrix?C)?{?? ?????? ????int?blockRow?=?blockIdx.y;?? ????int?blockCol?=?blockIdx.x;?? ?? ?????? ????Matrix?Csub?=?GetSubMatrix(C,?blockRow,?blockCol);?? ?? ?????? ?????? ????float?Cvalue?=?0;?? ?? ?????? ????int?row?=?threadIdx.y;?? ????int?col?=?threadIdx.x;?? ?? ?????? ?????? ????for?(int?m?=?0;?m?<?(A.width?/?BLOCK_SIZE);?++m)?{?? ?????????? ????????Matrix?Asub?=?GetSubMatrix(A,?blockRow,?m);?? ?????????? ?????????? ????????Matrix?Bsub?=?GetSubMatrix(B,?m,?blockCol);?? ?? ?????????? ????????__shared__?float?As[BLOCK_SIZE][BLOCK_SIZE];?? ????????__shared__?float?Bs[BLOCK_SIZE][BLOCK_SIZE];?? ?? ?????????? ?????????? ????????As[row][col]?=?GetElement(Asub,?row,?col);?? ????????Bs[row][col]?=?GetElement(Bsub,?row,?col);?? ?? ?????????? ?????????? ????????__syncthreads();?? ?? ?????????? ????????for?(int?e?=?0;?e?<?BLOCK_SIZE;?++e)?? ????????????Cvalue?+=?As[row][e]?*?Bs[e][col];?? ?? ?????????? ?????????? ????????__syncthreads();?? ????}?? ?? ?????? ?????? ????SetElement(Csub,?row,?col,?Cvalue);?? }
同一個block中的線程在95行的for循環中獲取到的Asub，Bsub，Csub是一樣的，每個線程就負責Csub內元素的計算
http://blog.csdn.net/csgxy123/article/details/10018531
總結
以上是生活随笔為你收集整理的CUDA编程指南阅读笔记（六）的全部內容，希望文章能夠幫你解決所遇到的問題。
如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。
