[Hotball's Hive]CUDA 教学-3

 

2008/3/25

[Hotball's Hive]第一個 CUDA 程式

CUDA 目前有兩種不同的 API：Runtime API 和 Driver API，兩種 API 各有其適用的範圍。由於 runtime API 較容易使用，一開始我們會以 runetime API 為主。

CUDA 的初始化

首先，先建立一個檔案 first_cuda.cu。如果是使用 Visual Studio 的話，則請先按照這裡的設定方式設定 project。

要使用 runtime API 的時候，需要 include cuda_runtime.h。所以，在程式的最前面，加上

#include <stdio.h>
#include <cuda_runtime.h>

接下來是一個 InitCUDA 函式，會呼叫 runtime API 中，有關初始化 CUDA 的功能：

bool InitCUDA()
{
    int count;
    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 的裝置的數目。如果系統上沒有支援 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 函式，把它設為目前要使用的裝置。

最後是 main 函式。在 main 函式中我們直接呼叫剛才的 InitCUDA 函式，並顯示適當的訊息：

int main()
{
    if(!InitCUDA()) {
        return 0;
    }
    printf("CUDA initialized./n");
    return 0;
}

這樣就可以利用 nvcc 來 compile 這個程式了。使用 Visual Studio 的話，若按照先前的設定方式，可以直接 Build Project 並執行。

nvcc 是 CUDA 的 compile 工具，它會將 .cu 檔拆解出在 GPU 上執行的部份，及在 host 上執行的部份，並呼叫適當的程式進行 compile 動作。在 GPU 執行的部份會透過 NVIDIA 提供的 compiler 編譯成中介碼，而 host 執行的部份則會透過系統上的 C++ compiler 編譯（在 Windows 上使用 Visual C++ 而在 Linux 上使用 gcc）。

編譯後的程式，執行時如果系統上有支援 CUDA 的裝置，應該會顯示 CUDA initialized. 的訊息，否則會顯示相關的錯誤訊息。

利用 CUDA 進行運算

到目前為止，我們的程式並沒有做什麼有用的工作。所以，現在我們加入一個簡單的動作，就是把一大堆數字，計算出它的平方和。

首先，把程式最前面的 include 部份改成：

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#define DATA_SIZE 1048576
int data[DATA_SIZE];

並加入一個新函式 GenerateNumbers：

void GenerateNumbers(int *number, int size)
{
    for(int i = 0; i < size; i++) {
        number[i] = rand() % 10;
    }
}

這個函式會產生一大堆 0 ~ 9 之間的亂數。

要利用 CUDA 進行計算之前，要先把資料複製到顯示記憶體中，才能讓顯示晶片使用。因此，需要取得一塊適當大小的顯示記憶體，再把產生好的資料複製進去。在 main 函式中加入：

    GenerateNumbers(data, DATA_SIZE);
    int* gpudata, *result;
    cudaMalloc((void**) &gpudata, sizeof(int) * DATA_SIZE);
    cudaMalloc((void**) &result, sizeof(int));
    cudaMemcpy(gpudata, data, sizeof(int) * DATA_SIZE,
        cudaMemcpyHostToDevice);

上面這段程式會先呼叫 GenerateNumbers 產生亂數，並呼叫 cudaMalloc 取得一塊顯示記憶體（result 則是用來存取計算結果，在稍後會用到），並透過 cudaMemcpy 將產生的亂數複製到顯示記憶體中。cudaMalloc 和 cudaMemcpy 的用法和一般的 malloc 及 memcpy 類似，不過 cudaMemcpy 則多出一個參數，指示複製記憶體的方向。在這裡因為是從主記憶體複製到顯示記憶體，所以使用 cudaMemcpyHostToDevice。如果是從顯示記憶體到主記憶體，則使用 cudaMemcpyDeviceToHost。這在之後會用到。

接下來是要寫在顯示晶片上執行的程式。在 CUDA 中，在函式前面加上 __global__ 表示這個函式是要在顯示晶片上執行的。因此，加入以下的函式：

__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];
    }
    *result = sum;
}

在顯示晶片上執行的程式有一些限制，例如它不能有傳回值。其它的限制會在之後提到。

接下來是要讓 CUDA 執行這個函式。在 CUDA 中，要執行一個函式，使用以下的語法：

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

呼叫完後，還要把結果從顯示晶片複製回主記憶體上。在 main 函式中加入以下的程式：

    sumOfSquares<<<1, 1, 0>>>(gpudata, result);
    int sum;
    cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);
    cudaFree(gpudata);
    cudaFree(result);
    printf("sum: %d/n", sum);

因為這個程式只使用一個 thread，所以 block 數目、thread 數目都是 1。我們也沒有使用到任何 shared memory，所以設為 0。編譯後執行，應該可以看到執行的結果。

為了確定執行的結果正確，我們可以加上一段以 CPU 執行的程式碼，來驗證結果：

    sum = 0;
    for(int i = 0; i < DATA_SIZE; i++) {
        sum += data[i] * data[i];
    }
    printf("sum (CPU): %d/n", sum);

編譯後執行，確認兩個結果相同。

計算執行時間

CUDA 提供了一個 clock 函式，可以取得目前的 timestamp，很適合用來判斷一段程式執行所花費的時間（單位為 GPU 執行單元的時脈）。這對程式的最佳化也相當有用。要在我們的程式中記錄時間，把 sumOfSquares 函式改成：

__global__ static void sumOfSquares(int *num, int* result,
    clock_t* time)
{
    int sum = 0;
    int i;
    clock_t start = clock();
    for(i = 0; i < DATA_SIZE; i++) {
        sum += num[i] * num[i];
    }
    *result = sum;
    *time = clock() - start;
}

把 main 函式中間部份改成：

    int* gpudata, *result;
    clock_t* time;
    cudaMalloc((void**) &gpudata, sizeof(int) * DATA_SIZE);
    cudaMalloc((void**) &result, sizeof(int));
    cudaMalloc((void**) &time, sizeof(clock_t));
    cudaMemcpy(gpudata, data, sizeof(int) * DATA_SIZE,
        cudaMemcpyHostToDevice);

sumOfSquares<<<1, 1, 0>>>(gpudata, result, time);
    int sum;
    clock_t time_used;
    cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);
    cudaMemcpy(&time_used, time, sizeof(clock_t),
        cudaMemcpyDeviceToHost);
    cudaFree(gpudata);
    cudaFree(result);
    printf("sum: %d time: %d/n", sum, time_used);

編譯後執行，就可以看到執行所花費的時間了。

如果計算實際執行時間的話，可能會注意到它的執行效率並不好。這是因為我們的程式並沒有利用到 CUDA 的主要的優勢，即平行化執行。在下一段文章中，會討論如何進行最佳化的動作。

---

本文經作者授權轉載，原文網址為：http://www.kimicat.com/%E7%AC%AC%E4%B8%80%E5%80%8Bcuda%E7%A8%8B%E5%BC%8F