第3章 CUDA程式設計基礎

3.5 “HelloWorld”CUDA 程式設計例項

安裝完Visual Studio 2013軟體並配置好CUDA開發環境。本節我們正式CUDA程式設計。

首先我們開啟vs2013，點選新建工程，選中CUDA 7.5的模板。輸入工程名等資訊，完成工程的建立。

建立完工程後，會預設有一個kernel.cu的檔案，其實現的是矩陣相加，這時候可以直接除錯執行，如果沒有報錯則證明可以進行CUDA程式設計了。

可能會出現，下面這種無法識別<<<符號的情況，但是目測不影響結果的輸出。

執行結果如下：

如果想編寫執行自己的程式，則需要先移除kernel.cu檔案，不然會出現main函式重複宣告。

我們現在開始編寫自己的CUDA程式，首先移除原來的kernel.cu檔案，然後右鍵工程名，選擇新增新item。

選擇CUDA C/C++ File，然後輸入檔名為main，就可以得到一個空的main.cu檔案。在檔案中輸入如下程式碼：

#include <stdio.h>
#include <cuda_runtime.h>
#include <string>
// 要使用 runtime API 的時候，需要 include cuda_runtime.h
bool InitCUDA()
{
    int count;
    // 獲取計算能力>=1.0的裝置數量
 

    cudaGetDeviceCount(&count);
    printf("%d\n", 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) 
        {
            // 輸出裝置名稱
 

            char *c = prop.name;
            printf("%s\n",c);
            // 定義裝置計算能力的主要修訂編號
            printf("%d",prop.major);
            // 定義裝置計算能力的次要修訂編號
            printf("%d\n", prop.minor);
            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__ void kernel(char *str1, char *str2)
{
    while ((*str1) != '\0')
        *str2++ = *str1++;
    *str2 = '\0';
}

int main()
{
    if (!InitCUDA()) {
        return 0;
    }
    char* host_str1;
    host_str1 = "HelloWorld";
    char* host_str2 = new char[11];
    int size = strlen(host_str1) + 1;
    char* dev_str1;
    char* dev_str2;
    cudaMalloc((void**)&dev_str1, size);
    cudaMalloc((void**)&dev_str2, size);

    cudaMemcpy(dev_str1, host_str1, size, cudaMemcpyHostToDevice);
    kernel <<<1, 1 >>>(dev_str1, dev_str2);
    cudaMemcpy(host_str2, dev_str2, size, cudaMemcpyDeviceToHost);
    printf("%s\n",host_str2);
    cudaFree(dev_str1);
    cudaFree(dev_str2);
    getchar();
    return 0;
}

InitCUDA()中會先呼叫 cudaGetDeviceCount 函式，取得支援 CUDA 的裝置的數目。如果系統上沒有支援 CUDA 的裝置，則它會傳回 1，而 device 0 會是一個模擬的裝置，但不支援 CUDA 1.0 以上的功能。所以，要確定系統上是否有支援 CUDA 的裝置，需要對每個 device 呼叫 cudaGetDeviceProperties 函式，取得裝置的各項資料，並判斷裝置支援的 CUDA 版本（prop.major 和 prop.minor 分別代表裝置支援的版本號碼，例如 1.0 則 prop.major 為 1 而 prop.minor 為 0）。

透過 cudaGetDeviceProperties 函式可以取得許多資料，除了裝置支援的 CUDA 版本之外，還有裝置的名稱、記憶體的大小、最大的 thread 數目、執行單元的頻率等等。詳情可參考 NVIDIA 的 CUDA Programming Guide。

在找到支援 CUDA 1.0 以上的裝置之後，就可以呼叫 cudaSetDevice 函式，把它設為目前要使用的裝置。

編譯後輸出如下：

能正確輸出上述結果，則代表第一個完整的CUDA程式到此結束了。但是這個專案的結構看上去一點也不規範化，所有的程式碼都寫到一個原始檔裡，這樣的程式碼不易讀懂，也不好定位程式的錯誤。CUDA程式一般可以進行如下圖所示的檔案結構管理。

如上圖所示，編寫CUDA程式過程中，若某個功能適合序列，則編寫序列程式碼在CPU端執行，若適合並行，則編寫CUDA並行程式碼在GPU端執行。CUDA並行程式碼的機構應進行如下分離：
(1) CUDA程式呼叫介面封裝了和CUDA相關的函式，一般位於.cpp檔案中。

(2) CUDA主機端程式碼的主要功能包括選擇計算裝置，進行GPU端儲存器的分配、主機端與裝置端直接的資料複製及呼叫kernel函式等準備工作，一般位於一個獨立的.cu檔案中。

(3) CUDA裝置端程式碼主要指GPU端執行的核心程式碼kernel函式。kernel函式可能有多個，各個kernel函式各自放在一個單獨的以kernel函式功能命名的.cu檔案下，因此kernel函式相關的檔案可能有多個。

根據以上程式碼管理的方法，對HelloWordl程式改寫如下：在原始檔建立3個檔案(main.cpp、kernel.cu和GPU_HelloWorld.cu)。main.cpp中呼叫核函式GPU_HelloWorld(host_str1, host_str2)相當於上述結構中的“GPU函式”，kernel.cu檔案相當於上述結構中的“kernel函式”。

main.cpp檔案主要包含了一個完整程式執行的主框架，其程式碼改寫如下：

#include<stdio.h>
extern bool GPU_HelloWorld(char* host_str1, char* host_str2);

int main()
{
    char* host_str1;
    host_str1 = "HelloWorld";
    char* host_str2 = new char[11];
    if(!GPU_HelloWorld(host_str1, host_str2))
        return 0;
    printf("%s\n",host_str2);
    return 0;
}

kernel.cu檔案為GPU端執行函式的核心程式碼，其程式碼如下：

#include <cuda_runtime.h>
__global__ void kernel(char *str1, char *str2)
{
    while ((*str1) != '\0')
        *str2++ = *str1++;
    *str2 = '\0';
}

GPU_HelloWorld.cu檔案作用是為kernel函式選擇可用的計算裝置並進行呼叫環境準備，其程式碼如下：

#include "kernel.cu"
#include <cuda_runtime.h>
#include <string>
bool InitCUDA()
{
    int count;
    // 獲取計算能力>=1.0的裝置數量
    cudaGetDeviceCount(&count);
    printf("%d\n", 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) 
        {
            // 輸出裝置名稱
            char *c = prop.name;
            printf("%s\n",c);
            // 定義裝置計算能力的主要修訂編號
            printf("%d",prop.major);
            // 定義裝置計算能力的次要修訂編號
            printf("%d\n", prop.minor);
            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;
}

bool GPU_HelloWorld(char* host_str1, char* host_str2)
{
    if(!InitCUDA())
    {
        return 0;
    }
    int size = strlen(host_str1)+1;
    char* dev_str1;
    char* dev_str2;
    cudaMalloc((void**)&dev_str1, size);
    cudaMalloc((void**)&dev_str2, size);

    cudaMemcpy(dev_str1, host_str1, size, cudaMemcpyHostToDevice);
    kernel <<<1, 1 >>>(dev_str1, dev_str2);
    cudaMemcpy(host_str2, dev_str2, size, cudaMemcpyDeviceToHost);
    cudaFree(dev_str1);
    cudaFree(dev_str2);
    return 1;
}

編譯執行結構修改後的專案，將得到和修改前一樣的執行結果。編譯時需要將kernel.cu檔案從專案生成中排除，否則編譯連線時會出錯。

注：如何從生成專案中移除

(1) 右鍵想要移除的檔案，選擇屬性

(2) 點選右側的從從生成中移除的yes，然後點選確定

(3) 弄好了後看到之前選擇檔案上有一個紅色提示符