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（下）》。