CUDA編程之快速入門

CUDA（Compute Unified Device Architecture）的中文全稱爲計算統一設備架構。作圖像視覺領域的同窗多多少少都會接觸到CUDA，畢竟要作性能速度優化，CUDA是個很重要的工具，CUDA是作視覺的同窗難以繞過的一個坑，必須踩一踩才踏實。CUDA編程真的是入門容易精通難，具備計算機體系結構和C語言編程知識儲備的同窗上手CUDA編程應該難度不會很大。本文章將經過如下五個方面幫助你們比較全面地瞭解CUDA編程最重要的知識點，作到快速入門：

1. GPU架構特色
2. CUDA線程模型
3. CUDA內存模型
4. CUDA編程模型
5. CUDA應用小例子

1. GPU架構特色

首先咱們先談一談串行計算和並行計算。咱們知道，高性能計算的關鍵利用多核處理器進行並行計算。

當咱們求解一個計算機程序任務時，咱們很天然的想法就是將該任務分解成一系列小任務，把這些小任務一一完成。在串行計算時，咱們的想法就是讓咱們的處理器每次處理一個計算任務，處理完一個計算任務後再計算下一個任務，直到全部小任務都完成了，那麼這個大的程序任務也就完成了。以下圖所示，就是咱們怎麼用串行編程思想求解問題的步驟。

可是串行計算的缺點很是明顯，若是咱們擁有多核處理器，咱們能夠利用多核處理器同時處理多個任務時，並且這些小任務並無關聯關係（不須要相互依賴，好比個人計算任務不須要用到你的計算結果），那咱們爲何還要使用串行編程呢？爲了進一步加快大任務的計算速度，咱們能夠把一些獨立的模塊分配到不一樣的處理器上進行同時計算（這就是並行），最後再將這些結果進行整合，完成一次任務計算。下圖就是將一個大的計算任務分解爲小任務，而後將獨立的小任務分配到不一樣處理器進行並行計算，最後再經過串行程序把結果彙總完成此次的總的計算任務。

因此，一個程序可不能夠進行並行計算，關鍵就在於咱們要分析出該程序能夠拆分出哪幾個執行模塊，這些執行模塊哪些是獨立的，哪些又是強依賴強耦合的，獨立的模塊咱們能夠試着設計並行計算，充分利用多核處理器的優點進一步加速咱們的計算任務，強耦合模塊咱們就使用串行編程，利用串行+並行的編程思路完成一次高性能計算。

接下來咱們談談CPU和GPU有什麼區別，他們倆各自有什麼特色，咱們在談並行、串行計算時屢次談到「多核」的概念，如今咱們先從「核」的角度開始這個話題。首先CPU是專爲順序串行處理而優化的幾個核心組成。而GPU則由數以千計的更小、更高效的核心組成，這些核心專門爲同時處理多任務而設計，可高效地處理並行任務。也就是，CPU雖然每一個核心自身能力極強，處理任務上很是強悍，無奈他核心少，在並行計算上表現不佳；反觀GPU，雖然他的每一個核心的計算能力不算強，但他勝在核心很是多，能夠同時處理多個計算任務，在並行計算的支持上作得很好。

GPU和CPU的不一樣硬件特色決定了他們的應用場景，CPU是計算機的運算和控制的核心，GPU主要用做圖形圖像處理。圖像在計算機呈現的形式就是矩陣，咱們對圖像的處理其實就是操做各類矩陣進行計算，而不少矩陣的運算其實能夠作並行化，這使得圖像處理能夠作得很快，所以GPU在圖形圖像領域也有了大展拳腳的機會。下圖表示的就是一個多GPU計算機硬件系統，能夠看出，一個GPU內存就有不少個SP和各種內存，這些硬件都是GPU進行高效並行計算的基礎。

如今再從數據處理的角度來對比CPU和GPU的特色。CPU須要很強的通用性來處理各類不一樣的數據類型，好比整型、浮點數等，同時它又必須擅長處理邏輯判斷所致使的大量分支跳轉和中斷處理，因此CPU其實就是一個能力很強的夥計，他能把不少事處理得妥穩當當，固然啦咱們須要給他不少資源供他使用（各類硬件），這也致使了CPU不可能有太多核心（核心總數不超過16）。而GPU面對的則是類型高度統一的、相互無依賴的大規模數據和不須要被打斷的純淨的計算環境，GPU有很是多核心（費米架構就有512核），雖然其核心的能力遠沒有CPU的核心強，可是勝在多，
在處理簡單計算任務時呈現出「人多力量大」的優點，這就是並行計算的魅力。

整理一下二者特色就是：

• CPU：擅長流程控制和邏輯處理，不規則數據結構，不可預測存儲結構，單線程程序，分支密集型算法
• GPU：擅長數據並行計算，規則數據結構，可預測存儲模式

如今的計算機體系架構中，要完成CUDA並行計算，單靠GPU一人之力是不能完成計算任務的，必須藉助CPU來協同配合完成一次高性能的並行計算任務。

通常而言，並行部分在GPU上運行，串行部分在CPU運行，這就是異構計算。具體一點，異構計算的意思就是不一樣體系結構的處理器相互協做完成計算任務。CPU負責整體的程序流程，而GPU負責具體的計算任務，當GPU各個線程完成計算任務後，咱們就將GPU那邊計算獲得的結果拷貝到CPU端，完成一次計算任務。

因此應用程序利用GPU實現加速的整體分工就是：密集計算代碼（約佔5%的代碼量）由GPU負責完成，剩餘串行代碼由CPU負責執行。

2. CUDA線程模型

下面咱們介紹CUDA的線程組織結構。首先咱們都知道，線程是程序執行的最基本單元，CUDA的並行計算就是經過成千上萬個線程的並行執行來實現的。下面的機構圖說明了GPU的不一樣層次的結構。

CUDA的線程模型從小往大來總結就是：

1. Thread：線程，並行的基本單位
2. Thread Block：線程塊，互相合做的線程組，線程塊有以下幾個特色：

• 容許彼此同步
• 能夠經過共享內存快速交換數據
• 以1維、2維或3維組織

1. Grid：一組線程塊

• 以1維、2維組織
• 共享全局內存

Kernel：在GPU上執行的核心程序，這個kernel函數是運行在某個Grid上的。

• One kernel <-> One Grid

每個block和每一個thread都有本身的ID，咱們經過相應的索引找到相應的線程和線程塊。

• threadIdx，blockIdx
• Block ID: 1D or 2D
• Thread ID: 1D, 2D or 3D

理解kernel，必需要對kernel的線程層次結構有一個清晰的認識。首先GPU上不少並行化的輕量級線程。kernel在device上執行時其實是啓動不少線程，一個kernel所啓動的全部線程稱爲一個網格（grid），同一個網格上的線程共享相同的全局內存空間，grid是線程結構的第一層次，而網格又能夠分爲不少線程塊（block），一個線程塊裏面包含不少線程，這是第二個層次。線程兩層組織結構如上圖所示，這是一個gird和block均爲2-dim的線程組織。grid和block都是定義爲dim3類型的變量，dim3能夠當作是包含三個無符號整數（x，y，z）成員的結構體變量，在定義時，缺省值初始化爲1。所以grid和block能夠靈活地定義爲1-dim，2-dim以及3-dim結構，kernel調用時也必須經過執行配置<<<grid, block>>>來指定kernel所使用的網格維度和線程塊維度。舉個例子，咱們以上圖爲例，分析怎麼經過<<<grid,block>>>>這種標記方式索引到咱們想要的那個線程。CUDA的這種<<<grid,block>>>其實就是一個多級索引的方法，第一級索引是(grid.xIdx, grid.yIdy)，對應上圖例子就是(1, 1)，經過它咱們就能找到了這個線程塊的位置，而後咱們啓動二級索引(block.xIdx, block.yIdx, block.zIdx)來定位到指定的線程。這就是咱們CUDA的線程組織結構。

這裏想談談SP和SM（流處理器），不少人會被這兩個專業名詞搞得暈頭轉向。

• SP：最基本的處理單元，streaming processor，也稱爲CUDA core。最後具體的指令和任務都是在SP上處理的。GPU進行並行計算，也就是不少個SP同時作處理。
• SM：多個SP加上其餘的一些資源組成一個streaming multiprocessor。也叫GPU大核，其餘資源如：warp scheduler，register，shared memory等。SM能夠看作GPU的心臟（對比CPU核心），register和shared memory是SM的稀缺資源。CUDA將這些資源分配給全部駐留在SM中的threads。所以，這些有限的資源就使每一個SM中active warps有很是嚴格的限制，也就限制了並行能力。

須要指出，每一個SM包含的SP數量依據GPU架構而不一樣，Fermi架構GF100是32個，GF10X是48個，Kepler架構都是192個，Maxwell都是128個。

簡而言之，SP是線程執行的硬件單位，SM中包含多個SP，一個GPU能夠有多個SM（好比16個），最終一個GPU可能包含有上千個SP。這麼多核心「同時運行」，速度可想而知，這個引號只是想代表實際上，軟件邏輯上是全部SP是並行的，可是物理上並非全部SP都能同時執行計算（好比咱們只有8個SM卻有1024個線程塊須要調度處理），由於有些會處於掛起，就緒等其餘狀態，這有關GPU的線程調度。

下面這個圖將從硬件角度和軟件角度解釋CUDA的線程模型。

• 每一個線程由每一個線程處理器（SP）執行
• 線程塊由多核處理器（SM）執行
• 一個kernel其實由一個grid來執行，一個kernel一次只能在一個GPU上執行

block是軟件概念，一個block只會由一個sm調度，程序員在開發時，經過設定block的屬性，告訴GPU硬件，我有多少個線程，線程怎麼組織。而具體怎麼調度由sm的warps scheduler負責，block一旦被分配好SM，該block就會一直駐留在該SM中，直到執行結束。一個SM能夠同時擁有多個blocks，但須要序列執行。下圖顯示了GPU內部的硬件架構：

3. CUDA內存模型

CUDA中的內存模型分爲如下幾個層次：

• 每一個線程都用本身的registers（寄存器）
• 每一個線程都有本身的local memory（局部內存）
• 每一個線程塊內都有本身的shared memory（共享內存），全部線程塊內的全部線程共享這段內存資源
• 每一個grid都有本身的global memory（全局內存），不一樣線程塊的線程均可使用
• 每一個grid都有本身的constant memory（常量內存）和texture memory（紋理內存），），不一樣線程塊的線程均可使用

線程訪問這幾類存儲器的速度是register > local memory >shared memory > global memory

下面這幅圖表示就是這些內存在計算機架構中的所在層次。

4. CUDA編程模型

上面講了這麼多硬件相關的知識點，如今終於能夠開始說說CUDA是怎麼寫程序的了。

咱們先捋一捋常見的CUDA術語：

第一個要掌握的編程要點：咱們怎麼寫一個能在GPU跑的程序或函數呢？

經過關鍵字就能夠表示某個程序在CPU上跑仍是在GPU上跑！以下表所示，好比咱們用__global__定義一個kernel函數，就是CPU上調用，GPU上執行，注意__global__函數的返回值必須設置爲void。

第二個編程要點：CPU和GPU間的數據傳輸怎麼寫？

首先介紹在GPU內存分配回收內存的函數接口：

• cudaMalloc(): 在設備端分配global memory
• cudaFree(): 釋放存儲空間

CPU的數據和GPU端數據作數據傳輸的函數接口是同樣的，他們經過傳遞的函數實參（枚舉類型）來表示傳輸方向：

cudaMemcpy(void dst, void src, size_t nbytes,
enum cudaMemcpyKind direction)

enum cudaMemcpyKind:

• cudaMemcpyHostToDevice（CPU到GPU）
• cudaMemcpyDeviceToHost（GPU到CPU）
• cudaMemcpyDeviceToDevice（GPU到GPU）

第三個編程要點是：怎麼用代碼表示線程組織模型？
咱們能夠用dim3類來表示網格和線程塊的組織方式，網格grid能夠表示爲一維和二維格式，線程塊block能夠表示爲一維、二維和三維的數據格式。

dim3 DimGrid(100, 50);  //5000個線程塊，維度是100*50
dim3 DimBlock(4, 8, 8);  //每一個線層塊內包含256個線程，線程塊內的維度是4*8*8

接下來介紹一個很是重要又很難懂的一個知識點，咱們怎麼計算線程號呢？

1.使用N個線程塊，每個線程塊只有一個線程，即

dim3 dimGrid(N);
dim3 dimBlock(1);

此時的線程號的計算方式就是

threadId = blockIdx.x;

其中threadId的取值範圍爲0到N-1。對於這種狀況，咱們能夠將其看做是一個列向量，列向量中的每一行對應一個線程塊。列向量中每一行只有1個元素，對應一個線程。

2.使用M×N個線程塊，每一個線程塊1個線程

因爲線程塊是2維的，故能夠看作是一個M*N的2維矩陣，其線程號有兩個維度，即：

dim3 dimGrid(M,N);
dim3 dimBlock(1);

其中

blockIdx.x 取值0到M-1
blcokIdx.y 取值0到N-1

這種狀況通常用於處理2維數據結構，好比2維圖像。每個像素用一個線程來處理，此時須要線程號來映射圖像像素的對應位置，如

pos = blockIdx.y * blcokDim.x + blockIdx.x; //其中gridDim.x等於M

3.使用一個線程塊，該線程具備N個線程，即

dim3 dimGrid(1);
dim3 dimBlock(N);

此時線程號的計算方式爲

threadId = threadIdx.x;

其中threadId的範圍是0到N-1，對於這種狀況，能夠看作是一個行向量，行向量中的每個元素的每個元素對應着一個線程。

4.使用M個線程塊，每一個線程塊內含有N個線程，即

dim3 dimGrid(M);
dim3 dimBlock(N);

這種狀況，能夠把它想象成二維矩陣，矩陣的行與線程塊對應，矩陣的列與線程編號對應，那線程號的計算方式爲

threadId = threadIdx.x + blcokIdx*blockDim.x;

上面其實就是把二維的索引空間轉換爲一維索引空間的過程。

5.使用M×N的二維線程塊，每個線程塊具備P×Q個線程，即

dim3 dimGrid(M, N);
dim3 dimBlock(P, Q);

這種狀況實際上是咱們遇到的最多狀況，特別適用於處理具備二維數據結構的算法，好比圖像處理領域。

其索引有兩個維度

threadId.x = blockIdx.x*blockDim.x+threadIdx.x;
threadId.y = blockIdx.y*blockDim.y+threadIdx.y;

上述公式就是把線程和線程塊的索引映射爲圖像像素座標的計算方法。

CUDA應用例子

咱們已經掌握了CUDA編程的基本語法，如今咱們開始以一些小例子來真正上手CUDA。

首先咱們編寫一個程序，查看咱們GPU的一些硬件配置狀況。

#include "device_launch_parameters.h"
#include <iostream>

int main()
{
    int deviceCount;
    cudaGetDeviceCount(&deviceCount);
    for(int i=0;i<deviceCount;i++)
    {
        cudaDeviceProp devProp;
        cudaGetDeviceProperties(&devProp, i);
        std::cout << "使用GPU device " << i << ": " << devProp.name << std::endl;
        std::cout << "設備全局內存總量： " << devProp.totalGlobalMem / 1024 / 1024 << "MB" << std::endl;
        std::cout << "SM的數量：" << devProp.multiProcessorCount << std::endl;
        std::cout << "每一個線程塊的共享內存大小：" << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
        std::cout << "每一個線程塊的最大線程數：" << devProp.maxThreadsPerBlock << std::endl;
        std::cout << "設備上一個線程塊（Block）種可用的32位寄存器數量： " << devProp.regsPerBlock << std::endl;
        std::cout << "每一個EM的最大線程數：" << devProp.maxThreadsPerMultiProcessor << std::endl;
        std::cout << "每一個EM的最大線程束數：" << devProp.maxThreadsPerMultiProcessor / 32 << std::endl;
        std::cout << "設備上多處理器的數量： " << devProp.multiProcessorCount << std::endl;
        std::cout << "======================================================" << std::endl;     
        
    }
    return 0;
}

咱們利用nvcc來編譯程序。

nvcc test1.cu -o test1

輸出結果：由於個人服務器是8個TITAN GPU，爲了省略重複信息，下面只顯示兩個GPU結果

使用GPU device 0: TITAN X (Pascal)
設備全局內存總量： 12189MB
SM的數量：28
每一個線程塊的共享內存大小：48 KB
每一個線程塊的最大線程數：1024
設備上一個線程塊（Block）種可用的32位寄存器數量： 65536
每一個EM的最大線程數：2048
每一個EM的最大線程束數：64
設備上多處理器的數量： 28
======================================================
使用GPU device 1: TITAN X (Pascal)
設備全局內存總量： 12189MB
SM的數量：28
每一個線程塊的共享內存大小：48 KB
每一個線程塊的最大線程數：1024
設備上一個線程塊（Block）種可用的32位寄存器數量： 65536
每一個EM的最大線程數：2048
每一個EM的最大線程束數：64
設備上多處理器的數量： 28
======================================================

.......

第一個計算任務：將兩個元素數目爲1024×1024的float數組相加。

首先咱們思考一下若是隻用CPU咱們怎麼串行完成這個任務。

#include <iostream>
#include <stdlib.h>
#include <sys/time.h>
#include <math.h>

using namespace std;

int main()
{
    struct timeval start, end;
    gettimeofday( &start, NULL );
    float*A, *B, *C;
    int n = 1024 * 1024;
    int size = n * sizeof(float);
    A = (float*)malloc(size);
    B = (float*)malloc(size);
    C = (float*)malloc(size);

    for(int i=0;i<n;i++)
    {
        A[i] = 90.0;
        B[i] = 10.0;
    }
    
    for(int i=0;i<n;i++)
    {
        C[i] = A[i] + B[i];
    }

    float max_error = 0.0;
    for(int i=0;i<n;i++)
    {
        max_error += fabs(100.0-C[i]);
    }
    cout << "max_error is " << max_error << endl;
    gettimeofday( &end, NULL );
    int timeuse = 1000000 * ( end.tv_sec - start.tv_sec ) + end.tv_usec - start.tv_usec;
    cout << "total time is " << timeuse/1000 << "ms" <<endl;
    return 0;
}

CPU方式輸出結果

max_error is 0
total time is 22ms

若是咱們使用GPU來作並行計算，速度將會如何呢？

編程要點：

1. 每一個Block中的Thread數最大不超過512；
2. 爲了充分利用SM，Block數儘量多，>100。

#include "cuda_runtime.h"
#include <stdlib.h>
#include <iostream>
#include <sys/time.h>

using namespace std;

__global__ void Plus(float A[], float B[], float C[], int n)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    C[i] = A[i] + B[i];
}

int main()
{
    struct timeval start, end;
    gettimeofday( &start, NULL );
    float*A, *Ad, *B, *Bd, *C, *Cd;
    int n = 1024 * 1024;
    int size = n * sizeof(float);

    // CPU端分配內存
    A = (float*)malloc(size);
    B = (float*)malloc(size);
    C = (float*)malloc(size);

    // 初始化數組
    for(int i=0;i<n;i++)
    {
        A[i] = 90.0;
        B[i] = 10.0;
    }

    // GPU端分配內存
    cudaMalloc((void**)&Ad, size);
    cudaMalloc((void**)&Bd, size);
    cudaMalloc((void**)&Cd, size);

    // CPU的數據拷貝到GPU端
    cudaMemcpy(Ad, A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(Bd, B, size, cudaMemcpyHostToDevice);
    cudaMemcpy(Bd, B, size, cudaMemcpyHostToDevice);

    // 定義kernel執行配置，（1024*1024/512）個block，每一個block裏面有512個線程
    dim3 dimBlock(512);
    dim3 dimGrid(n/512);

    // 執行kernel
    Plus<<<dimGrid, dimBlock>>>(Ad, Bd, Cd, n);

    // 將在GPU端計算好的結果拷貝回CPU端
    cudaMemcpy(C, Cd, size, cudaMemcpyDeviceToHost);

    // 校驗偏差
    float max_error = 0.0;
    for(int i=0;i<n;i++)
    {
        max_error += fabs(100.0 - C[i]);
    }

    cout << "max error is " << max_error << endl;

    // 釋放CPU端、GPU端的內存
    free(A);
    free(B);
    free(C);
    cudaFree(Ad);
    cudaFree(Bd);
    cudaFree(Cd);
    gettimeofday( &end, NULL );
    int timeuse = 1000000 * ( end.tv_sec - start.tv_sec ) + end.tv_usec - start.tv_usec;
    cout << "total time is " << timeuse/1000 << "ms" <<endl;
    return 0;
}

GPU方式輸出結果

max error is 0
total time is 1278ms

由上面的例子看出，使用CUDA編程時咱們看不到for循環了，由於CPU編程的循環已經被分散到各個thread上作了，因此咱們也就看到不到for一類的語句。從結果上看，CPU的循環計算的速度比GPU計算快多了，緣由就在於CUDA中有大量的內存拷貝操做（數據傳輸花費了大量時間，而計算時間卻很是少），若是計算量比較小的話，CPU計算會更合適一些。

下面計算一個稍微複雜的例子，矩陣加法，即對兩個矩陣對應座標的元素相加後的結果存儲在第三個的對應位置的元素上。

值得注意的是，這個計算任務我採用了二維數組的計算方式，注意一下二維數組在CUDA編程中的寫法。

CPU版本

#include <stdlib.h>
#include <iostream>
#include <sys/time.h>
#include <math.h>

#define ROWS 1024
#define COLS 1024

using namespace std;

int main()
{
    struct timeval start, end;
    gettimeofday( &start, NULL );
    int *A, **A_ptr, *B, **B_ptr, *C, **C_ptr;
    int total_size = ROWS*COLS*sizeof(int);
    A = (int*)malloc(total_size);
    B = (int*)malloc(total_size);
    C = (int*)malloc(total_size);
    A_ptr = (int**)malloc(ROWS*sizeof(int*));
    B_ptr = (int**)malloc(ROWS*sizeof(int*));
    C_ptr = (int**)malloc(ROWS*sizeof(int*));
    
    //CPU一維數組初始化
    for(int i=0;i<ROWS*COLS;i++)
    {
        A[i] = 80;
        B[i] = 20;
    }
    
    for(int i=0;i<ROWS;i++)
    {
        A_ptr[i] = A + COLS*i;
        B_ptr[i] = B + COLS*i;
        C_ptr[i] = C + COLS*i;
    }
    
    for(int i=0;i<ROWS;i++)
        for(int j=0;j<COLS;j++)
        {
            C_ptr[i][j] = A_ptr[i][j] + B_ptr[i][j];
        }
        
    //檢查結果
    int max_error = 0;
    for(int i=0;i<ROWS*COLS;i++)
    {
        //cout << C[i] << endl;
        max_error += abs(100-C[i]);
    }
    
    cout << "max_error is " << max_error <<endl;     
    gettimeofday( &end, NULL );
    int timeuse = 1000000 * ( end.tv_sec - start.tv_sec ) + end.tv_usec - start.tv_usec;
    cout << "total time is " << timeuse/1000 << "ms" <<endl;
    
    return 0;
}

CPU方式輸出

max_error is 0
total time is 29ms

GPU版本

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <sys/time.h> 
#include <stdio.h>
#include <math.h>
#define Row  1024
#define Col 1024
 
 
__global__ void addKernel(int **C,  int **A, int ** B)
{
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    int idy = threadIdx.y + blockDim.y * blockIdx.y;
    if (idx < Col && idy < Row) {
        C[idy][idx] = A[idy][idx] + B[idy][idx];
    }
}
 
int main()
{
    struct timeval start, end;
    gettimeofday( &start, NULL );

    int **A = (int **)malloc(sizeof(int*) * Row);
    int **B = (int **)malloc(sizeof(int*) * Row);
    int **C = (int **)malloc(sizeof(int*) * Row);
    int *dataA = (int *)malloc(sizeof(int) * Row * Col);
    int *dataB = (int *)malloc(sizeof(int) * Row * Col);
    int *dataC = (int *)malloc(sizeof(int) * Row * Col);
    int **d_A;
    int **d_B;
    int **d_C;
    int *d_dataA;
    int *d_dataB;
    int *d_dataC;
    //malloc device memory
    cudaMalloc((void**)&d_A, sizeof(int **) * Row);
    cudaMalloc((void**)&d_B, sizeof(int **) * Row);
    cudaMalloc((void**)&d_C, sizeof(int **) * Row);
    cudaMalloc((void**)&d_dataA, sizeof(int) *Row*Col);
    cudaMalloc((void**)&d_dataB, sizeof(int) *Row*Col);
    cudaMalloc((void**)&d_dataC, sizeof(int) *Row*Col);
    //set value
    for (int i = 0; i < Row*Col; i++) {
        dataA[i] = 90;
        dataB[i] = 10;
    }
    //將主機指針A指向設備數據位置，目的是讓設備二級指針可以指向設備數據一級指針
    //A 和  dataA 都傳到了設備上，可是兩者尚未創建對應關係
    for (int i = 0; i < Row; i++) {
        A[i] = d_dataA + Col * i;
        B[i] = d_dataB + Col * i;
        C[i] = d_dataC + Col * i;
    }
                                                                
    cudaMemcpy(d_A, A, sizeof(int*) * Row, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, B, sizeof(int*) * Row, cudaMemcpyHostToDevice);
    cudaMemcpy(d_C, C, sizeof(int*) * Row, cudaMemcpyHostToDevice);
    cudaMemcpy(d_dataA, dataA, sizeof(int) * Row * Col, cudaMemcpyHostToDevice);
    cudaMemcpy(d_dataB, dataB, sizeof(int) * Row * Col, cudaMemcpyHostToDevice);
    dim3 threadPerBlock(16, 16);
    dim3 blockNumber( (Col + threadPerBlock.x - 1)/ threadPerBlock.x, (Row + threadPerBlock.y - 1) / threadPerBlock.y );
    printf("Block(%d,%d)   Grid(%d,%d).\n", threadPerBlock.x, threadPerBlock.y, blockNumber.x, blockNumber.y);
    addKernel << <blockNumber, threadPerBlock >> > (d_C, d_A, d_B);
    //拷貝計算數據-一級數據指針
    cudaMemcpy(dataC, d_dataC, sizeof(int) * Row * Col, cudaMemcpyDeviceToHost);
                                                                                             
    int max_error = 0;
    for(int i=0;i<Row*Col;i++)
    {
        //printf("%d\n", dataC[i]);
        max_error += abs(100-dataC[i]);
    }

    //釋放內存
    free(A);
    free(B);
    free(C);
    free(dataA);
    free(dataB);
    free(dataC);
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    cudaFree(d_dataA);
    cudaFree(d_dataB);
    cudaFree(d_dataC);

    printf("max_error is %d\n", max_error);
    gettimeofday( &end, NULL );
    int timeuse = 1000000 * ( end.tv_sec - start.tv_sec ) + end.tv_usec - start.tv_usec;
    printf("total time is %d ms\n", timeuse/1000);

    return 0;
}

GPU輸出

Block(16,16)   Grid(64,64).
max_error is 0
total time is 442 ms

從結果看出，CPU計算時間仍是比GPU的計算時間短。這裏須要指出的是，這種二維數組的程序寫法的效率並不高（雖然比較符合咱們的思惟方式），由於咱們作了兩次訪存操做。因此通常而言，作高性能計算通常不會採起這種編程方式。

最後一個例子咱們將計算一個更加複雜的任務，矩陣乘法

回顧一下矩陣乘法：兩矩陣相乘，左矩陣第一行乘以右矩陣第一列（分別相乘，第一個數乘第一個數），乘完以後相加，即爲結果的第一行第一列的數，依次往下算，直到計算完全部矩陣元素。

CPU版本

#include <iostream>
#include <stdlib.h>
#include <sys/time.h>

#define ROWS 1024
#define COLS 1024

using namespace std;

void matrix_mul_cpu(float* M, float* N, float* P, int width)
{
    for(int i=0;i<width;i++)
        for(int j=0;j<width;j++)
        {
            float sum = 0.0;
            for(int k=0;k<width;k++)
            {
                float a = M[i*width+k];
                float b = N[k*width+j];
                sum += a*b;
            }
            P[i*width+j] = sum;
        }
}

int main()
{
    struct timeval start, end;
    gettimeofday( &start, NULL );
    float *A, *B, *C;
    int total_size = ROWS*COLS*sizeof(float);
    A = (float*)malloc(total_size);
    B = (float*)malloc(total_size);
    C = (float*)malloc(total_size);

    //CPU一維數組初始化
    for(int i=0;i<ROWS*COLS;i++)
    {
        A[i] = 80.0;
        B[i] = 20.0;
    }

    matrix_mul_cpu(A, B, C, COLS);

    gettimeofday( &end, NULL );
    int timeuse = 1000000 * ( end.tv_sec - start.tv_sec ) + end.tv_usec - start.tv_usec;
    cout << "total time is " << timeuse/1000 << "ms" <<endl;

    return 0;
}

CPU輸出

total time is 7617ms

梳理一下CUDA求解矩陣乘法的思路：由於C=A×B，咱們利用每一個線程求解C矩陣每一個(x, y)的元素，每一個線程載入A的一行和B的一列，遍歷各自行列元素，對A、B對應的元素作一次乘法和一次加法。

GPU版本

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <sys/time.h> 
#include <stdio.h>
#include <math.h>
#define Row  1024
#define Col 1024

 
__global__ void matrix_mul_gpu(int *M, int* N, int* P, int width)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;
                
    int sum = 0;
    for(int k=0;k<width;k++)
    {
        int a = M[j*width+k];
        int b = N[k*width+i];
        sum += a*b;
    }
    P[j*width+i] = sum;
}
 
int main()
{
    struct timeval start, end;
    gettimeofday( &start, NULL );

    int *A = (int *)malloc(sizeof(int) * Row * Col);
    int *B = (int *)malloc(sizeof(int) * Row * Col);
    int *C = (int *)malloc(sizeof(int) * Row * Col);
    //malloc device memory
    int *d_dataA, *d_dataB, *d_dataC;
    cudaMalloc((void**)&d_dataA, sizeof(int) *Row*Col);
    cudaMalloc((void**)&d_dataB, sizeof(int) *Row*Col);
    cudaMalloc((void**)&d_dataC, sizeof(int) *Row*Col);
    //set value
    for (int i = 0; i < Row*Col; i++) {
        A[i] = 90;
        B[i] = 10;
    }
                                                                
    cudaMemcpy(d_dataA, A, sizeof(int) * Row * Col, cudaMemcpyHostToDevice);
    cudaMemcpy(d_dataB, B, sizeof(int) * Row * Col, cudaMemcpyHostToDevice);
    dim3 threadPerBlock(16, 16);
    dim3 blockNumber((Col+threadPerBlock.x-1)/ threadPerBlock.x, (Row+threadPerBlock.y-1)/ threadPerBlock.y );
    printf("Block(%d,%d)   Grid(%d,%d).\n", threadPerBlock.x, threadPerBlock.y, blockNumber.x, blockNumber.y);
    matrix_mul_gpu << <blockNumber, threadPerBlock >> > (d_dataA, d_dataB, d_dataC, Col);
    //拷貝計算數據-一級數據指針
    cudaMemcpy(C, d_dataC, sizeof(int) * Row * Col, cudaMemcpyDeviceToHost);
                                                                                             
    //釋放內存
    free(A);
    free(B);
    free(C);
    cudaFree(d_dataA);
    cudaFree(d_dataB);
    cudaFree(d_dataC);

    gettimeofday( &end, NULL );
    int timeuse = 1000000 * ( end.tv_sec - start.tv_sec ) + end.tv_usec - start.tv_usec;
    printf("total time is %d ms\n", timeuse/1000);

    return 0;
}

GPU輸出

Block(16,16)   Grid(64,64).
total time is 506 ms

從這個矩陣乘法任務能夠看出，咱們經過GPU進行並行計算的方式僅花費了0.5秒，可是CPU串行計算方式卻花費了7.6秒，計算速度提高了十多倍，可見並行計算的威力！