nVidia CUDA API（下）

在《nVidia CUDA API（上）》的部分，已經大致說明了 extension 的部分；這邊要來講的則是 runtime library 的部分。

CUDA 的 runtime library 分成下面三部分：

• common component
  提供 CUDA 的 vector 等型別，以及可以同時在 host 及 device 上可以執行的函式。主要包括：
  • 內建的 vector 型別
  • dim3 型別（用於指定 kernel 參數的型別）
  • texture 型別
  • 數學函式
• device component
  在 device 上執行的部分，提供 device 上特殊的函式。主要有：
  • 數學函式。這邊的數學函式是 device 的特殊版本，精確度較低，但是速度較快。
  • 同步函式
  • Atomic 函式
  • 型別轉換函式（Conversion/Casting）
  • texture 函式
• host component
  在 host 上執行的部分，負責控制 device。主要是處理：
  • Device 管理
  • Context 管理（所謂 Context 大致相當於 CPU？）
  • 記憶體管理
  • Code module 管理（Code module 似乎相當於 dynamic library）
  • Execution control
  • Texture reference 管理
  • Interoperability with OpenGL and Direct3D.

不過，這邊應該就不會列舉了∼個別的說明，請參考《CUDA Programming Guide 1.0》這份文件。而這邊，就只大概介紹一些 Heresy 覺得比較會用到的部分了。

記憶體管理

要說最重要、一定會用到的部分，應該就是記憶體管理的部分了！因為在 CUDA 中，要在 device 的程式中存取的資料，一定要先存在 device 上，所以在記憶體管理的部分，CUDA 提供了很類似 C 的一些函式：cudaMalloc()、cudaFree()。原則上，cudaMalloc 大致上就是 C 語言中的 malloc，也就像是 C 的 new；而 cudaFree 則就相當於 C 的 free，以及 C 的 delete。

而要使用的方法，則是要先宣告 pointer，在 allocate 記憶體給他；下面就是一個簡單的範例：

float	*dev_Array; cudaMalloc( (void**)&dev_Array, ArraySize * sizeof( float ) );

這樣，就可以在 device 上配置一塊大小是 ArraySize 的 float 陣列了∼

如果要把已經存在一般程式中的資料複製到 device 上呢？這時候就要用對應到 C 語言中 memcpy 的函式－cudaMemcpy 了∼以 CUDA SDK 範例裡的寫法，會是：

// define, allocate on host
 float	*host_Array = new float[100]; 
for( int i = 0; i < 100;    i )
 	host_Array[i] = i;
  
 // define, allocate on device
 float	*dev_Array;
 cudaMalloc( (void**)&dev_Array, 100 * sizeof( float ) );
  
 // copy from host to device
 cudaMemcpy>( dev_Array, host_Array, 100 * sizeof( float ),
 	cudaMemcpyHostToDevice );

cudaMemcpy 這個函式的四個參數依序為：目標、來源、大小、方法。最後一項方法的型別是 CUDA 的一個列舉型別 cudaMemcpyKind，他有四種值：cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice；相當直覺的，就是代表由 host/device 複製到 device/host。所以，如果是要反過來把 device 上的資料複製回 host 的話，就只要把最後一項參數改成 cudaMemcpyDeviceToHost 就可以了！

而要釋放掉 allocate 出來的記憶體空間的話，就直接使用 cudaFree 就可以了。

cudaFree( dev_Array );

實際上 CUDA 的記憶體管理除了這邊提到的基本類型，還有 cudaMallocPitch(), cudaMemset(), cudaMemcpy2D(), cudaMemcpyToArray(), cudaMemcpyToSymbol() 等其他的記憶體管理函式，在這邊就不一一提出來了。

裝置管理

一般情況下，可能用不到這一部分的 API 吧∼但是如果有兩個以上的 CUDA Device 的話，就須要透過這些函式來選擇要用的裝置了∼此外，也可以先透過這些函式，來確認機器上的一些基本規格。

這部分的函式不多，有下面幾個：

• cudaGetDeviceCount( int* count )
  

  取得目前機器上，支援 CUDA 的 device 數目，會把結果存在 count 中。函式則會 return 一個 cudaError_t，用來表示 CUDA 的錯誤。

• cudaGetDeviceProperties( cudaDeviceProp* prop, int dev )
  

  取得目前機器上第 dev 個 device 的屬性；prop 的結構會是：

  struct cudaDeviceProp
  {
    char name[256];		//裝置名稱
     size_t totalGlobalMem;	//Device 的 global memory 總量
     size_t sharedMemPerBlock;	//Device 的 shared memory 總量
      int regsPerBlock;		//每一個 block 的 registers 數量
     int warpSize;		//warp size
    size_t memPitch;		//最大的 memory pitch
    int maxThreadsPerBlock;	//Block 的最大 thread 數 
     int maxThreadsDim[3];	//Block 中 thread 最大的維度
     int maxGridSize[3];	//Grid 中 block 的最大維度
     size_t totalConstMem;	//Device 的 constant memory 總量
     int major;		//主要版本編號
     int minor;		//次要版本編號
     int clockRate;		//時脈
     size_t textureAlignment;	//alignment requirement mentioned 
  }; 

• cudaChooseDevice( int* dev, const cudaDeviceProp* prop)
  

  在目前的 device 中，找一個最符合給定的 prop 的，並把他的索引直存在 dev 裡。

• cudaSetDevice(int dev)
  

  指定現在要使用第 dev 個 device。

• cudaGetDevice(int* dev)
  

  取得目前正在使用的 device 的索引編號，儲存在 dev 中。

快速數學計算

從上面的列表可以發現，在 common component 和 device component 中，都有數學函式的功能。其中，common component 中的應該算是一般的數學計算函式，而 device component 所提供的，則是精確度較低，但是速度較快的版本！

也就是說 CUDA 在 device component 的部分，提供了部分數學計算的函式，是以 GPU 來做計算的快速版本。這些函式在 CUDA 中統一以「__」來做開頭，例如 __sinf(x)、__fadd_rz(x,y)、__logf(x) 等等。這些函式在《CUDA Programming Guide 1.0》的 Appendix B 中，可以找到比較完整的列表可以參考。

而在 nvcc 的編譯參數中，也還提供了一個「-use_fast_math」的參數，可以強制所有可以的數學計算函數，都使用叫快速的版本來計算。這在精確度比較不重要的情況下，應該可以對效能有些增益。

上面講的部分，都是 runtime API 的部分；大部分的函式，都是以 cuda 來做為開頭。而實際上，CUDA 還有提供 Driver API；這一部分的函式會是以 cu 來做為開頭。而許多函式，都可以和 runtime API 來做一一對應。而以文件裡的說法，Driver API 應該是比 runtime API 來的低階，但是實際的差異，Heresy 也不是很清楚就是了。

本文主要參考nVidia 的《CUDA Programming Guide 1.0》（PDF 文件）的第四章。

原始發表： nVidia CUDA API（下）