0. 簡介

最近作者希望系統性的去學習一下CUDA加速的相關知識，正好看到深藍學院有這一門課程。所以這里作者以此課程來作為主線來進行記錄分享，方便能給CUDA網絡加速學習的萌新們去提供一定的幫助。

1. 基礎矩陣乘法

下圖是矩陣乘法的示意圖，下面我們來看一下在CPU和GPU上是怎么表達的。

CPU代碼示意流程：

// Matrix multiplication on the (CPU) host void main(){define A, B, Cfor i= 0 to M-1 dofor j = 0 to N-1 do/* compute element C(i,j) */for k = 0 to K-1 doC(i,j) <= C(i,j) + A(i,k) * B(k,j)endendend }

GPU代碼示意流程：

void main(){define A_cpu, B_cpu, C_cpuin the CPU memorydefine A_gpu, B_gpu, C_gpuin the GPU memorymemcopyA_cputo A_gpumemcopyB_cputo B_gpudim3 dimBlock(16, 16)dim3 dimGrid(N/dimBlock.x, M/dimBlock.y)matrixMul<<<dimGrid, dimBlock>>>(A_gpu,B_gpu,C_gpu,K)memcopyC_gputo C_cpu }__global__ void matrixMul(A_gpu,B_gpu,C_gpu,K){temp <= 0i<= blockIdx.y* blockDim.y+ threadIdx.y// Row iof matrix Cj <= blockIdx.x* blockDim.x+ threadIdx.x// Column j of matrix Cfor k = 0 to K-1 doaccu<= accu+ A_gpu(i,k) * B_gpu(k,j)endC_gpu(i,j) <= accu }

在GPU中，一個線程負責計算C中的一個元素，其中A中的每一行從全局內存中載入N次，B中的每一列從全局內存中載入M次。總共的次數為2mnk的讀取次數（因為在每一個k處都要從A和B中讀取一次，所以要乘二），值得注意的是，可以將多次訪問的數據放到共享內存中，減少重復讀取的次數，并充分利用共享內存的延遲低的優勢

所以我們可以看到共享內存可以被然被內存訪問指令訪問，同時擁有更高的速度訪問(延遲&吞吐)。在共享內存中存在有兩種申請空間方式，靜態申請和動態申請，但是共享內存的大小只有幾十K，過度使用共享內存會降低程序的并行性，共享內存在使用時候需要注意的是：

• 使用__shared__關鍵字，同時有靜態申請和動態申請兩種方式；
• 將每個線程從全局索引位置讀取元素，將它存儲到共享內存之中；
• 注意數據存在著交叉，應該將邊界上的數據拷貝進來；
• 塊內線程同步：__syncthreads()

2. 線程同步函數

上面講到了__syncthreads()函數，該函數是cuda的內建函數，用于塊內線程通信。__syncthreads()函數是對線程進行同步，需要保證的是需要對所有的共享內存需要同步的線程都被同步到，下面是兩個示例。

申請共享內存會存在有兩種申請方式：
靜態申請

__global__ void staticReverse(int *d, int n) {__shared__ int s[64];int t = threadIdx.x;int tr = n-t-1;s[t] = d[t];__syncthreads();d[t] = s[tr]; } staticReverse<<<1,n>>>(d_d, n);

…詳情請參照古月居

總結

以上是生活随笔為你收集整理的CUDA入门和网络加速学习(二)的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。