第三篇：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 的是引數本身，還是引數的地址，這很關鍵。