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> 接下來是一個 InitCUDA 函式,會呼叫 runtime API 中,有關初始化 CUDA 的功能: bool 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 函式,把它設為目前要使用的裝置。 最後是 main 函式。在 main 函式中我們直接呼叫剛才的 InitCUDA 函式,並顯示適當的訊息: int main() 這樣就可以利用 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> 並加入一個新函式 GenerateNumbers: void GenerateNumbers(int *number, int size) 這個函式會產生一大堆 0 ~ 9 之間的亂數。 要利用 CUDA 進行計算之前,要先把資料複製到顯示記憶體中,才能讓顯示晶片使用。因此,需要取得一塊適當大小的顯示記憶體,再把產生好的資料複製進去。在 main 函式中加入: GenerateNumbers(data, DATA_SIZE); 上面這段程式會先呼叫 GenerateNumbers 產生亂數,並呼叫 cudaMalloc 取得一塊顯示記憶體(result 則是用來存取計算結果,在稍後會用到),並透過 cudaMemcpy 將產生的亂數複製到顯示記憶體中。cudaMalloc 和 cudaMemcpy 的用法和一般的 malloc 及 memcpy 類似,不過 cudaMemcpy 則多出一個參數,指示複製記憶體的方向。在這裡因為是從主記憶體複製到顯示記憶體,所以使用 cudaMemcpyHostToDevice。如果是從顯示記憶體到主記憶體,則使用 cudaMemcpyDeviceToHost。這在之後會用到。 接下來是要寫在顯示晶片上執行的程式。在 CUDA 中,在函式前面加上 __global__ 表示這個函式是要在顯示晶片上執行的。因此,加入以下的函式: __global__ static void sumOfSquares(int *num, int* result) 在顯示晶片上執行的程式有一些限制,例如它不能有傳回值。其它的限制會在之後提到。 接下來是要讓 CUDA 執行這個函式。在 CUDA 中,要執行一個函式,使用以下的語法: 函式名稱<<<block 數目, thread 數目, shared memory 大小>>>(參數...); 呼叫完後,還要把結果從顯示晶片複製回主記憶體上。在 main 函式中加入以下的程式: sumOfSquares<<<1, 1, 0>>>(gpudata, result); 因為這個程式只使用一個 thread,所以 block 數目、thread 數目都是 1。我們也沒有使用到任何 shared memory,所以設為 0。編譯後執行,應該可以看到執行的結果。 為了確定執行的結果正確,我們可以加上一段以 CPU 執行的程式碼,來驗證結果: sum = 0; 編譯後執行,確認兩個結果相同。 計算執行時間CUDA 提供了一個 clock 函式,可以取得目前的 timestamp,很適合用來判斷一段程式執行所花費的時間(單位為 GPU 執行單元的時脈)。這對程式的最佳化也相當有用。要在我們的程式中記錄時間,把 sumOfSquares 函式改成: __global__ static void sumOfSquares(int *num, int* result, 把 main 函式中間部份改成: int* gpudata, *result; sumOfSquares<<<1, 1, 0>>>(gpudata, result, time); 編譯後執行,就可以看到執行所花費的時間了。 如果計算實際執行時間的話,可能會注意到它的執行效率並不好。這是因為我們的程式並沒有利用到 CUDA 的主要的優勢,即平行化執行。在下一段文章中,會討論如何進行最佳化的動作。 |
Copyright(c) 2007 Ping-Che Chen