第三篇：CUDA 標准編程模式

前言

       本文將介紹 CUDA 編程的基本模式，所有 CUDA 程序都基於此模式編寫，即使是調用庫，庫的底層也是這個模式實現的。

模式描述

       1. 定義需要在 device 端執行的核函數。( 函數聲明前加 _golbal_ 關鍵字 )

       2. 在顯存中為待運算的數據以及需要存放結果的變量開辟顯存空間。( cudaMalloc 函數實現 )

       3. 將待運算的數據傳輸進顯存。( cudaMemcpy，cublasSetVector 等函數實現 )

       4. 調用 device 端函數，同時要將需要為 device 端函數創建的塊數線程數等參數傳遞進 <<<>>>。( 注: <<<>>>下方編譯器可能顯示語法錯誤，不用管 )

       5. 從顯存中獲取結果變量。( cudaMemcpy，cublasGetVector 等函數實現 )

       6. 釋放申請的顯存空間。( cudaFree 實現 )

       PS：每個 device 端函數在被調用時都能獲取到調用它的具體塊號，線程號，從而實現並行( 獲取方法請參考下面的編程規范說明以及代碼示例 )。

編程規范說明

　　在 CUDA 標准編程模式中，增加了一些編程規范，在這里簡要說明：

　　函數聲明關鍵字：

　　　　1. __device__

　　　　表明此函數只能在 GPU 中被調用，在 GPU 中執行。這類函數只能被 __global__ 類型函數或 __device__ 類型函數調用。

　　　　2. __global__

　　　　表明此函數在 CPU 上調用，在 GPU 中執行。這也是以后會常提到的 "內核函數"，有時為了便於理解也稱 "device" 端函數。

　　　　3. __host__

　　　　表明此函數在 CPU 上調用和執行，這也是默認情況。
　　內核函數配置運算符 <<<>>> - 這個運算符在調用內核函數的時候使用，一般情況下傳遞進三個參數：

　　　　1. 塊數

　　　　2. 線程數

　　　　3. 共享內存大小 (此參數默認為0 )

　　內核函數中的幾個系統變量 - 這幾個變量可以在內核函數中使用，從而控制塊與線程的工作：

　　　　1. gridDim：塊數

　　　　2. blockDim：塊中線程數

　　　　3. blockIdx：塊編號 (0 - gridDim-1)

　　　　4. threadIdx：線程編號 (0 - blockDim-1)

　　知道這些已經足夠編寫 CUDA 程序了，更多的編程說明將在以后的文章中介紹。

代碼示例

　　該程序采用 CUDA 並行化思想來對數組進行求和 (代碼下方如果出現紅色波浪線無視之)：

// 相關 CUDA 庫
#include "cuda_runtime.h"
#include "cuda.h"
#include "device_launch_parameters.h"

#include <iostream>
#include <cstdlib>

using namespace std;

const int N = 100;

// 塊數
const int BLOCK_data = 3; 
// 各塊中的線程數
const int THREAD_data = 10; 

// CUDA初始化函數
bool InitCUDA()
{
    int deviceCount; 

    // 獲取顯示設備數
    cudaGetDeviceCount (&deviceCount);

    if (deviceCount == 0) 
    {
        cout << "找不到設備" << endl;
        return EXIT_FAILURE;
    }

    int i;
    for (i=0; i<deviceCount; i++)
    {
        cudaDeviceProp prop;
        if (cudaGetDeviceProperties(&prop,i)==cudaSuccess) // 獲取設備屬性
        {
            if (prop.major>=1) //cuda計算能力
            {
                break;
            }
        }
    }

    if (i==deviceCount)
    {
        cout << "找不到支持 CUDA 計算的設備" << endl;
        return EXIT_FAILURE;
    }

    cudaSetDevice(i); // 選定使用的顯示設備

    return EXIT_SUCCESS;
}

// 此函數在主機端調用，設備端執行。
__global__ 
static void Sum (int *data,int *result)
{
    // 取得線程號
    const int tid = threadIdx.x; 
    // 獲得塊號
    const int bid = blockIdx.x; 
    
    int sum = 0;

    // 有點像網格計算的思路
    for (int i=bid*THREAD_data+tid; i<N; i+=BLOCK_data*THREAD_data)
    {
        sum += data[i];
    }
    
    // result 數組存放各個線程的計算結果
    result[bid*THREAD_data+tid] = sum; 
}

int main ()
{
    // 初始化 CUDA 編譯環境
    if (InitCUDA()) {
        return EXIT_FAILURE;
    }
    cout << "成功建立 CUDA 計算環境" << endl << endl;

    // 建立，初始化，打印測試數組
    int *data = new int [N];
    cout << "測試矩陣: " << endl;
    for (int i=0; i<N; i++)
    {
        data[i] = rand()%10;
        cout << data[i] << " ";
        if ((i+1)%10 == 0) cout << endl;
    }
    cout << endl;

    int *gpudata, *result; 
    
    // 在顯存中為計算對象開辟空間
    cudaMalloc ((void**)&gpudata, sizeof(int)*N); 
    // 在顯存中為結果對象開辟空間
    cudaMalloc ((void**)&result, sizeof(int)*BLOCK_data*THREAD_data);
    
    // 將數組數據傳輸進顯存
    cudaMemcpy (gpudata, data, sizeof(int)*N, cudaMemcpyHostToDevice); 
    // 調用 kernel 函數 - 此函數可以根據顯存地址以及自身的塊號，線程號處理數據。
    Sum<<<BLOCK_data,THREAD_data,0>>> (gpudata,result);
    
    // 在內存中為計算對象開辟空間
    int *sumArray = new int[THREAD_data*BLOCK_data];
    // 從顯存獲取處理的結果
    cudaMemcpy (sumArray, result, sizeof(int)*THREAD_data*BLOCK_data, cudaMemcpyDeviceToHost);
    
    // 釋放顯存
    cudaFree (gpudata); 
    cudaFree (result);

    // 計算 GPU 每個線程計算出來和的總和
    int final_sum=0;
    for (int i=0; i<THREAD_data*BLOCK_data; i++)
    {
        final_sum += sumArray[i];
    }

    cout << "GPU 求和結果為: " << final_sum << endl;

    // 使用 CPU 對矩陣進行求和並將結果對照
    final_sum = 0;
    for (int i=0; i<N; i++)
    {
        final_sum += data[i];
    }
    cout << "CPU 求和結果為: " << final_sum << endl;

    getchar();

    return 0;
}

運行測試

      

       PS：矩陣元素是隨機生成的

小結

       1. 掌握本節知識的關鍵除了要掌握各個API，還要深刻理解內核函數中的塊及線程變量的控制，或者說施展 :) 

       2. 一定要明確傳遞進 API 的是參數本身，還是參數的地址，這很關鍵。