nVidia CUDA API（上）

最新推荐文章于 2023-04-11 00:02:00 发布

nicole_yaoyao

最新推荐文章于 2023-04-11 00:02:00 发布

阅读量2k

点赞数 1

分类专栏： cuda路漫漫文章标签： nvidia cuda api extension thread function library

cuda路漫漫专栏收录该内容

4 篇文章 0 订阅

订阅专栏

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 )。