技術標籤：矩陣的加法

這一篇我們來使用CUDA寫一個矩陣加法的小程式，實際上手一下。

環境搭建

筆者的測試環境是Ubuntu 16.04系統 + Titan XP顯示卡。CUDA Toolkit和顯示卡驅動的安裝我就不贅述了，相信玩深度學習的各位都安裝過。如果你不知道自己安裝沒安裝過，但你用GPU跑過pytorch、tensorflow等框架的程式碼，那你也其實是安裝過這些環境的，可以直接編譯CUDA程式。

CUDA程式碼檔案字尾名是.cu。編譯單檔案CUDA程式非常方便，執行下面這條指令即可。

nvcc -o test test.cu

之後會生成一個可執行檔案test，使用下面的命令即可執行。

./test

是不是非常簡單，目前這種方式已經足夠我們進行CUDA程式設計的入門了，之後我們再學習使用其他工具進行程式的除錯。

從矩陣加法入手

既然是第一個程式，我們從最經典也最適合GPU的矩陣加法入手，學習一下標準的CUDA程式會由哪些部分組成。我們會實現一個矩陣求和的程式，然後統計執行時間，看看相比於CPU序列程式設計，GPU到底達到了多高的加速比。

先貼程式碼。

#include #include #include "cudastart.h"//CPU對照組，用於對比加速比void sumMatrix2DonCPU(float * MatA,float * MatB,float * MatC,int nx,int ny){    float* a = MatA;    float* b = MatB;    float* c = MatC;    for(int j=0; j    {        for(int i=0; i        {          c[i] = a[i]+b[i];        }        c += nx;        b += nx;        a += nx;    }}//核函式，每一個執行緒計算矩陣中的一個元素。__global__ void sumMatrix(float * MatA,float * MatB,float * MatC,int nx,int ny){    int ix = threadIdx.x+blockDim.x*blockIdx.x;    int iy = threadIdx.y+blockDim.y*blockIdx.y;    int idx = ix+iy*ny;    if (ix    {        MatC[idx] = MatA[idx]+MatB[idx];    }}//主函式int main(int argc,char** argv){    //裝置初始化    printf("strating...\n");    initDevice(0);    //輸入二維矩陣，4096*4096，單精度浮點型。    int nx = 1<<12;    int ny = 1<<12;    int nBytes = nx*ny*sizeof(float);    //Malloc，開闢主機記憶體    float* A_host = (float*)malloc(nBytes);    float* B_host = (float*)malloc(nBytes);    float* C_host = (float*)malloc(nBytes);    float* C_from_gpu = (float*)malloc(nBytes);    initialData(A_host, nx*ny);    initialData(B_host, nx*ny);    //cudaMalloc，開闢裝置記憶體    float* A_dev = NULL;    float* B_dev = NULL;    float* C_dev = NULL;    CHECK(cudaMalloc((void**)&A_dev, nBytes));    CHECK(cudaMalloc((void**)&B_dev, nBytes));    CHECK(cudaMalloc((void**)&C_dev, nBytes));    //輸入資料從主機記憶體拷貝到裝置記憶體    CHECK(cudaMemcpy(A_dev, A_host, nBytes, cudaMemcpyHostToDevice));    CHECK(cudaMemcpy(B_dev, B_host, nBytes, cudaMemcpyHostToDevice));    //二維執行緒塊，32×32    dim3 block(32, 32);    //二維執行緒網格，128×128    dim3 grid((nx-1)/block.x+1, (ny-1)/block.y+1);    //測試GPU執行時間    double gpuStart = cpuSecond();    //將核函式放線上程網格中執行    sumMatrix<<>>(A_dev, B_dev, C_dev, nx, ny);    CHECK(cudaDeviceSynchronize());    double gpuTime = cpuSecond() - gpuStart;    printf("GPU Execution Time: %f sec\n", gpuTime);    //在CPU上完成相同的任務    cudaMemcpy(C_from_gpu, C_dev, nBytes, cudaMemcpyDeviceToHost);    double cpuStart=cpuSecond();    sumMatrix2DonCPU(A_host, B_host, C_host, nx, ny);    double cpuTime = cpuSecond() - cpuStart;    printf("CPU Execution Time: %f sec\n", cpuTime);    //檢查GPU與CPU計算結果是否相同    CHECK(cudaMemcpy(C_from_gpu, C_dev, nBytes, cudaMemcpyDeviceToHost));    checkResult(C_host, C_from_gpu, nx*ny);    //釋放記憶體    cudaFree(A_dev);    cudaFree(B_dev);    cudaFree(C_dev);    free(A_host);    free(B_host);    free(C_host);    free(C_from_gpu);    cudaDeviceReset();    return 0;}
 

程式碼思路比較直觀，看註釋應該就能懂。我稍微解釋一下。首先對裝置進行初始化，這個我包裝起來放在標頭檔案裡了。之後初始化需要的主機記憶體和裝置記憶體。接著在主機端隨機初始化兩個矩陣A_host、B_host，把這兩個陣列拷貝到裝置記憶體A_dev、B_dev中。我們初始化128×128的執行緒網格，和32×32的執行緒塊。匯流排程數為4096×4096，和我們輸入矩陣的大小相同，每個執行緒處理矩陣中的一個元素。處理完後把裝置記憶體上的計算結果拷貝回主機，並且與CPU對照組的結果進行比較，看是否相同。

程式碼我傳在github上，需要的同學可以clone下來跑一下。

https://github.com/ZihaoZhao/CUDA_study/tree/master/Sum_Matrix

執行結果

strating...Using device 0: TITAN X (Pascal)GPU Execution Time: 0.000634 secCPU Execution Time: 0.149086 secCheck result success!

這樣測量出的加速比為0.149086/0.000634=235倍。

在測量執行時間時，我們在cpuSecond()函式裡面實際使用的是gettimeofday()函式，這樣對於測量GPU時間來說其實是有一定誤差的，不過也能一定程度反映GPU的加速比。學習一門語言，上手程式設計是必須的，建議想要深入學習CUDA的讀者也親手編寫一下程式碼，至少run一下，會有更多收穫。

來源：https://zhuanlan.zhihu.com/p/97044592

End

宣告：部分內容來源於網路，僅供讀者學習、交流之目的。文章版權歸原作者所有。如有不妥，請聯絡刪除。