CUDA編程（二） CUDA初始化與核函數

CUDA編程（二）

CUDA初始化與核函數

CUDA初始化

在上一次中已經說過了，CUDA成功安裝之后，新建一個project還是十分簡單的，直接在新建項目的時候選擇NVIDIA CUDA項目就能夠了，我們先新建一個MyCudaTest project。刪掉自帶的演示樣例kernel.cu。然后新建項，新建一個CUDA C/C++ File ,我們首先看一下怎樣初始化CUDA，因此我命名為InitCuda.cu

首先我們要使用CUDA的RunTime API 所以 我們須要include cuda_runtime.h

#include <stdio.h> 

//CUDA RunTime API
#include <cuda_runtime.h>

接下來這個函數會調用 runtime API 中 有關初始化CUDA的內容

//CUDA 初始化
bool InitCUDA()
{
    int count;

    //取得支持Cuda的裝置的數目
    cudaGetDeviceCount(&count);

    //沒有符合的硬件
    if (count == 0) {
        fprintf(stderr, "There is no device.\n");
        return false;
    }

    int i;

    for (i = 0; i < count; i++) {
        cudaDeviceProp prop;
        if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
            if (prop.major >= 1) {
                break;
            }
        }
    }

    if (i == count) {
        fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
        return false;
    }

    cudaSetDevice(i);

    return true;
}

這段程序首先會調用cudaGetDeviceCount 函數。獲得支持 CUDA 的GPU的數量，假設計算機上沒有支持 CUDA 的裝置，則會傳回 1，而這個1是device 0 ，device0 僅僅是一個仿真裝置，可是CUDA的非常多功能都不支持（不支持CUDA1.0以上版本號），因此我們要真正確定系統上是否有支持CUDA的裝置，須要對每一個device調用cudaGetDeviceProperties，來獲得它們的詳細參數，以及所支持的CUDA版本號（prop.major 和 prop.minor 分別代表裝置支持的版本號號碼，比如 6.5 則 prop.major 為 6 而prop.minor 為 5）

cudaGetDeviceProperties除了能夠獲得裝置支持的 CUDA 版本號之外，還有裝置的名稱、內存的大小、最大的 thread 數目、運行單元的頻率等等。詳情可參考NVIDIA 的 CUDA Programming Guide。

在找到支持 CUDA 1.0 以上的裝置之后。就能夠呼叫 cudaSetDevice 函式，把它設為眼下要使用的顯卡。

以下我們在Main函數中調用InitCUDA函數，由於我們使用VS，所以直接ctrl+F5編譯運行就能夠了。運行時假設系統上有支持 CUDA 的裝置。應該會顯示 CUDA initialized。

int main() 
{

    if (!InitCUDA()) 
    { 
        return 0; 
    }

    printf("CUDA initialized.\n"); return 0;

}

完整程序：

#include <stdio.h> 

//CUDA RunTime API
#include <cuda_runtime.h>

//CUDA 初始化
bool InitCUDA()
{
    int count;

    //取得支持Cuda的裝置的數目
    cudaGetDeviceCount(&count);

    //沒有符合的硬件
    if (count == 0) {
        fprintf(stderr, "There is no device.\n");
        return false;
    }

    int i;

    for (i = 0; i < count; i++) {
        cudaDeviceProp prop;
        if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
            if (prop.major >= 1) {
                break;
            }
        }
    }

    if (i == count) {
        fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
        return false;
    }

    cudaSetDevice(i);

    return true;
}

int main() 
{

    if (!InitCUDA()) 
    { 
        return 0; 
    }

    printf("CUDA initialized.\n"); return 0;

}

CUDA核函數

完畢了CUDA的初始化檢查操作，以下我們就能夠使用CUDA完畢一些簡單計算了。這里我們打算計算一系列數字的立方和。

所以我們先寫了一個隨機函數：

#define DATA_SIZE 1048576

int data[DATA_SIZE];

//產生大量0-9之間的隨機數
void GenerateNumbers(int *number, int size)
{
    for (int i = 0; i < size; i++) {
        number[i] = rand() % 10;
    }
}

//生成隨機數(main中調用)
//GenerateNumbers(data, DATA_SIZE);

該函數會產生一大堆 0 ~ 9 之間的隨機數，然后我們要對他們進行立方和操作。

那么我們怎樣讓這個工作在顯卡上完畢呢？首先第一件事非常顯而易見，這些數字不能放在內存里了，而是要拷貝到GPU的顯存上。以下我們就來看一下數據復制的部分。

Host&Device架構：

上一次已經講過關於CUDA架構的一些基礎了。這里再略微復習一下。在 CUDA 的架構下，一個程序分為兩個部份：host 端和 device 端。Host 端是指在 CPU 上運行的部份，而 device 端則是在顯示芯片上運行的部份。Device 端的程序又稱為 “kernel”。通常 host 端程序會將數據准備好后，拷貝到顯卡的內存中，再由顯示芯片運行 device 端程序。完畢后再由 host 端程序將結果從顯卡的內存中取回。

我們須要把產生的數據拷貝到Device端的RAM，才干在顯卡上完畢計算。因此我們首先開辟一塊合適的顯存。然后把隨機數從內存復制進去。

    //生成隨機數
    GenerateNumbers(data, DATA_SIZE);

    /*把數據拷貝到顯卡內存中*/

    int* gpudata, *result;

    //cudaMalloc 取得一塊顯卡內存 ( 當中result用來存儲計算結果 )
    cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
    cudaMalloc((void**)&result, sizeof(int));

    //cudaMemcpy 將產生的隨機數拷貝到顯卡內存中 
    //cudaMemcpyHostToDevice - 從內存拷貝到顯卡內存
    //cudaMemcpyDeviceToHost - 從顯卡內存拷貝到內存
    cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE,cudaMemcpyHostToDevice);

凝視已經寫得比較明確了。cudaMalloc 和 cudaMemcpy 的使用方法和一般的 malloc 及 memcpy 相似，只是 cudaMemcpy 則多出一個參數，指示復制內存的方向。

在這里由於是從主內存拷貝到顯卡內存。所以使用 cudaMemcpyHostToDevice。假設是從顯卡內存到主內存，則使用cudaMemcpyDeviceToHost。

完畢了從內存到顯存的數據拷貝之后，我們接下來就要在顯卡上完畢計算了，怎樣讓程序跑在顯卡上？答案是核函數。

CUDA核函數：

要寫在顯示芯片上運行的程序。在 CUDA 中，在函數前面加上__global__ 表示這個函式是要在顯示芯片上運行的,所以我們僅僅要在正常函數之前加上一個__global__即可了：

// __global__ 函數 (GPU上運行) 計算立方和
__global__ static void sumOfSquares(int *num, int* result)
{
    int sum = 0;

    int i;

    for (i = 0; i < DATA_SIZE; i++) {

        sum += num[i] * num[i] * num[i];

    }

    *result = sum;

}

在顯示芯片上運行的程序有一些限制，首先最明顯的一個限制——不能有傳回值。另一些其它的限制，后面會慢慢提到。

運行核函數：

寫好核函數之后須要讓CUDA運行這個函數。

在 CUDA 中，要運行一個核函數，使用以下的語法：

    函數名稱<<<block 數目, thread 數目, shared memory 大小>>>(參數...);

這里我們先不去並行，僅僅是單純地完畢GPU計算，所以我們讓block = 1。thread = 1，share memory = 0

 sumOfSquares<<<1, 1, 0>>>(gpudata, result);

計算完了，千萬別忘了還要把結果從顯示芯片復制回主內存上。然后釋放掉內存~

    int sum;

    //cudaMemcpy 將結果從顯存中復制回內存
    cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);

    //Free
    cudaFree(gpudata);
    cudaFree(result);

最后我們把結果打印出來就大功告成了：

    printf("GPUsum: %d \n", sum);

之后我們再用CPU計算一下來驗證一下上面的過程是否有錯，這一步還是十分必要的：

    sum = 0;

    for (int i = 0; i < DATA_SIZE; i++) {
        sum += data[i] * data[i] * data[i];
    }

    printf("CPUsum: %d \n", sum);

完整程序：

程序代碼：

#include <stdio.h>
#include <stdlib.h>

//CUDA RunTime API
#include <cuda_runtime.h>

#define DATA_SIZE 1048576

int data[DATA_SIZE];

//產生大量0-9之間的隨機數
void GenerateNumbers(int *number, int size)
{
    for (int i = 0; i < size; i++) {
        number[i] = rand() % 10;
    }
}

//CUDA 初始化
bool InitCUDA()
{
    int count;

    //取得支持Cuda的裝置的數目
    cudaGetDeviceCount(&count);

    if (count == 0) {
        fprintf(stderr, "There is no device.\n");
        return false;
    }

    int i;

    for (i = 0; i < count; i++) {
        cudaDeviceProp prop;
        if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
            if (prop.major >= 1) {
                break;
            }
        }
    }

    if (i == count) {
        fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
        return false;
    }

    cudaSetDevice(i);

    return true;
}


// __global__ 函數 (GPU上運行) 計算立方和
__global__ static void sumOfSquares(int *num, int* result)
{
    int sum = 0;

    int i;

    for (i = 0; i < DATA_SIZE; i++) {

        sum += num[i] * num[i] * num[i];

    }

    *result = sum;

}





int main()
{

    //CUDA 初始化
    if (!InitCUDA()) {
        return 0;
    }

    //生成隨機數
    GenerateNumbers(data, DATA_SIZE);

    /*把數據拷貝到顯卡內存中*/

    int* gpudata, *result;

    //cudaMalloc 取得一塊顯卡內存 ( 當中result用來存儲計算結果 )
    cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
    cudaMalloc((void**)&result, sizeof(int));

    //cudaMemcpy 將產生的隨機數拷貝到顯卡內存中 
    //cudaMemcpyHostToDevice - 從內存拷貝到顯卡內存
    //cudaMemcpyDeviceToHost - 從顯卡內存拷貝到內存
    cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);

    // 在CUDA 中運行函數 語法：函數名稱<<<block 數目, thread 數目, shared memory 大小>>>(參數...);
    sumOfSquares << <1, 1, 0 >> >(gpudata, result);


    /*把結果從顯示芯片復制回主內存*/

    int sum;

    //cudaMemcpy 將結果從顯存中復制回內存
    cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);

    //Free
    cudaFree(gpudata);
    cudaFree(result);

    printf("GPUsum: %d \n", sum);

    sum = 0;

    for (int i = 0; i < DATA_SIZE; i++) {
        sum += data[i] * data[i] * data[i];
    }

    printf("CPUsum: %d \n", sum);

    return 0;
}

運行結果：

總結：

這次給大家介紹了CUDA的初始化和怎樣在顯卡上運行程序。即先將數據從內存拷貝到顯存。再寫好運算的核函數，之后用CUDA調用核函數，完畢GPU上的計算。之后當然不要忘記將結果復制回內存，釋放掉顯存。

總的來說一個CUDA程序的骨架已經搭建起來了，而GPU計算的重中之重即並行加速還沒有進行介紹，只是在加速之前我們另一件非常重要的事情須要考慮，那就是我們的程序究竟有沒有加速，也就是我們要輸出程序的運行時間，這個時間我們須要使用CUDA提供的一個Clock函數，能夠取得GPU運行單元的頻率，所以下一篇博客我將主要解說這個函數~希望能給大家的學習帶來幫助~

參考資料：《深入淺出談CUDA》