在CUDA中,host和device是两个重要的概念,我们用host指代CPU及其内存,而用device指代GPU及其内存。
典型的CUDA程序的执行流程如下:
- 分配host内存,并进行数据初始化;
- 分配device内存,并从host将数据拷贝到device上;
- 调用CUDA的核函数在device上完成指定的运算;
- 将device上的运算结果拷贝到host上;
- 释放device和host上分配的内存。
上面流程中最重要的一个过程是调用CUDA的核函数来执行并行计算,kernel是CUDA中一个重要的概念,kernel是在device上线程中并行执行的函数,核函数用__global__
符号声明,在调用时需要用<<<grid, block>>>
来指定kernel要执行的线程数量,在CUDA中,每一个线程都要执行核函数,并且每个线程会分配一个唯一的线程号thread ID,这个ID值可以通过核函数的内置变量threadIdx
来获得。
由于GPU实际上是异构模型,所以需要区分host和device上的代码,在CUDA中是通过函数类型限定词开区别host和device上的函数,主要的三个函数类型限定词如下:
__global__
:在device上执行,从host中调用(一些特定的GPU也可以从device上调用),返回类型必须是void
,不支持可变参数参数,不能成为类成员函数。注意用__global__
定义的kernel是异步的,这意味着host不会等待kernel执行完就执行下一步。__device__
:在device上执行,单仅可以从device中调用,不可以和__global__
同时用。__host__
:在host上执行,仅可以从host上调用,一般省略不写,不可以和__global__
同时用,但可和__device__
,此时函数会在device和host都编译。
kernel在device上执行时实际上是启动很多线程,一个kernel所启动的所有线程称为一个网格(grid),同一个网格上的线程共享相同的全局内存空间,grid是线程结构的第一层次,而网格又可以分为很多线程块(block),一个线程块里面包含很多线程,这是第二个层次。