nVidia CUDA API(上)

CUDA 的 API 架構,大概分成「C 語言的 extension」、「runtime library」兩部分。extension 的部分,提供了 C 的一些延伸,來訂一 CUDA 的變數、程式等等;在最簡單的 case 裡,應該只要用到 extension 的部分就夠了∼而 runtime library 的部分,則又提供了一些 CUDA device 的控制函式,以及一些針對 GPU 編寫的特殊的函式。

接下來,先講一些一定要知道的 extension 吧。

首先,CUDA 在 C 的 extension 分成四種:Function type qualifiers、Variable type qualifiers、Directive to specify how a kernel is executed、Built-in variables;下面來各自做一些簡短的介紹。

  • Function type qualifiers

用來指定 function 是要在 host 或 device 上執行,以及是用來被 host 或 device 呼叫。他的類別有三種:

  • __device__
    在 device 上執行,且只能被 device 呼叫。
    同時,他永遠是 inline function。
  • __global__
    將 function 宣告成一個 kernel,在 device 上執行,只能被 host 呼叫。
    他的 return type 必須要是 void;傳入的參數會是透過 shared memory 給 device,大小不能超過 256byte(這個 Heresy 比較不清楚)。而在乎叫時,必須要指定執行的參數(請參考下面的「Directive to specify how a kernel is executed」)
  • __host__
    在 host 上執行,且只能被 host 呼叫。(相當於一般的 function)

如果沒有指定的話,自然就是一般的 function,和 __host__ 一樣了。而 __host____device__ 可以同時使用,這樣 function 會編譯成同時可以在 device 和 host 執行。

而在 device 上執行的function(__device__ __global__)有一些基本的限制:

  • 不支援遞迴
  • 不能有 static 變數
  • 不能使用 variable number of arguments
  • Variable type qualifiers

在變數類型方面,是用來指定記憶體的類型。分成三種:

  • __device__

    宣告變數存在 device 上;可以和下面兩者同時使用,來做更進一步的設定。如果沒有額外指定的話,那這個變數會

    • 存在 global memory 空間
    • 生命週期和程式相同
    • 可以被 grid 中的所有 thread 透過 runtime library 存取。
  • __constant__

    可和 __device__ 同時使用,會將變數宣告成:

    • 存在 constant memory 空間
    • 生命週期和程式相同
    • 可以被 grid 中的所有 thread 透過 runtime library 存取。
  • __shared__

    可和 __device__ 同時使用,會將變數宣告成:

    • 存在 thread block 的 shared memory 空間
    • 生命週期和 thread block 相同
    • 只能被 block 中的 thread 存取
  • Directive to specify how a kernel is executed

指定 kernel 在 device 上執行的設定參數,主要就是指定這份 kernel 要用多大的 block grid(也就是多少個 block)、每個 block 多大(也就是每個 block 有多少 thread)。所以呼叫 __global__ 的地方,都要指定 execution configuration;他的形式式在 function name 和參數之間,加入「<<< Dg, Db, Ns >>>」。

其中,三個值的意義如下:

  • Dg 的型別是 dim3(屬於 common runtime component 的部分,一種簡單的資料結構),用來指定 grid 的維度和大小;Dg.x * Dg.y 就是 grid 中會被執行的 block 數目。
  • Db 的型別是 dim3,用來指定 block 的維度和大小;Db.x * Db.y * Db.z 就是每個 block 中的 thread 數目。
  • Ns 的型別是 size_t,用來指定每個 block 在 shared memory 中動態分配的變數的位元數。這個值可以不用指定,預設值是 0。

如果 function 是宣告成「__global__ void Func(float* parameter);」,那呼叫的方法就是「Func<<< Dg, Db, Ns >>>(parameter);」。而在 device 上總共會被產生的執行序數目,就會是 ( Dg.x * Dg.y ) * ( Db.x * Db.y * Db.z )

  • Built-in variables

指定 grid 和 block 的維度,以及 block 和 thread 的索引,有下面這些:

  • gridDim
    資料型別是 dim3,儲存 grid 的維度資料。
  • blockIdx
    資料型別是 uint3,儲存 grid 中 block 的索引值。
  • blockDim
    資料型別是 dim3,儲存 block 的維度資料。
  • threadIdx
    資料型別是 uint3,儲存 block 中 thread 的索引值。

而這些變數都是唯獨的,不能去修改他們的值;此外,也不能去用他們的位址。而他們主要的用處,是讓 device 上的程式,可以知道自己是哪一個 block 的 哪一個 thread,進而知道自己在陣列或 texture 中該取的值;某種程度上,就相當於迴圈中不斷累加、用來計數的 index。

Extension 的部分大概就是上面這些了∼而 runtime library 的部分,請參考《nVidia CUDA API(下)》。


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


原始發表: nVidia CUDA API(上)

發佈留言

發佈留言必須填寫的電子郵件地址不會公開。 必填欄位標示為 *