CUDA核函数

在GPU上调用的函数成为CUDA核函数(Kernel function)，核函数会被GPU上的多个线程执行。每个线程都会执行核函数里的代码，当然由于线程编号的不同，执行的代码路径可能会有所不同。下面的几行代码是向量加计算的CUDA核函数：

__global__ void vectoradd (int *a, int *b, int *c, int n){

int i = blockDim.x * blockIdx.x + threadIdx.x;

if (i < n){

c[i] = a[i] + b[i];

}

}

vectoradd<<>>(d_a, d_b, d_c, N);

从这段代码可以看出CUDA核函数的几个特征：

-函数的最前面是声明标识符__global__，该标识符表示这个函数可以在GPU上执行。需要指出的是尽管是在GPU上执行，但是仍然是由CPU端发起调用的

-核函数调用时需要用<<<...>>>符号来指定线程配置

-在核函数内部可以调用CUDA内置变量，比如threadIdx，blockDim等

-核函数相对于CPU代码是异步的，也就是控制会在核函数执行完成之前就返回，这样CPU就可以不用等待核函数的完成而继续执行后面的CPU代码

关于线程的配置以及内置变量将会在后面详细介绍。CUDA核函数除了上面提到的几个特征之外，还有一些限制：

-核函数内部只能访问device内存。因为核函数是执行在设备端，所以只能访问设备端内存。

-必须返回void类型。我们知道核函数是由CPU端发起的并执行在GPU上的函数。在核函数内部的数据均是位于GPU上的，假设核函数有返回值，那么返回值是位于GPU上的数据，CPU去直接接收这个数据是不被允许的。所以，核函数没有返回值。

-核函数不支持可变参数

-核函数不支持静态变量

-核函数不支持函数指针

在CUDA编程中，除了__global__外，常用的标识符还有：

__device__

-有标识符__device__的函数只能在GPU段执行

-只能在GPU段调用，比如可以在__global__以及__device__函数中调用

-__global__与__device__不能同时使用

另外一个常用的标识符是__host__：

-只能在host端执行

-只能在host端调用

单独使用__host__的情况时，该函数与普通的CPU函数的性质及使用方法没有任何差别。那既然这样为什么还要引入这个标识符呢？我们可以想象有这样一种情况，一个函数我们希望它既可以在CPU上调用也可以在GPU上调用，那么我们这样声明这个函数：__host__ __device__ funForCPUandGPU(args), 则这个函数既可以在CPU上执行也可以在GPU上执行。

线程配置

前面提到，在调用核函数时需要通过<<<...>>>指定线程配置，在具体介绍之前，我们先来了解CUDA编程中几个基本的概念。

线程(Thread)是CUDA程序的基本执行单元，每个线程内的执行都会顺序执行。所有的线程都会执行相同的代码，当然有可能会执行相同代码的不同分支。所有的线程之间是并行执行的，没有先后之分。

线程块(Thread Block)是由一组线程组成。每个线程块内部的线程之间可以进行协作，有可以共同访问的内存-共享内存。每个线程块会在GPU上的某一个流处理器(Streaming Multiprocessor, SM)中执行。

线程网格(Thread Grid)是一组线程块的集合。线程网格里的线程块会被调度到GPU的多个SM上去执行。线程块之间并没有同步机制，线程块被执行的先后顺序是不确定的。线程块之间的通讯比较昂贵，需要通过全局内存(global memory)来实现。

CUDA线程层次。每个核函数对应一个线程网格，每一个线程网格包含多个线程块，每个线程块包含多个线程。线程网格以及线程块可以是1维、2维或者3维的。

在调用核函数时需要指定的线程配置就是需要给定每个线程网格中有多少线程块，每个线程块有多少线程，并且他们的排列方式是怎样的。一个线程配置的例子如下：

dim3 grid(3,2,1), block(5,3,1)

kernel_name<<>>(…)

线程网格以及线程块的数据类型是dim3，实质上是一个结构体，有三个变量分别用来描述x、y、z三个方向的长度。<<<...>>>中的第一个参数用来指定线程网格的结构，也就是每个线程网格中有多少线程块，上面的例子中每个线程网格中有321=6个线程块，排布方式是三个方向上分别是3、2、1。第二个参数是用来指定线程块的结构，也就是每个线程块中有多少个线程，上面的例子中每个线程块中有531=15个线程，排布方式是三个方向上分别是5、3、1。<<<...>>>也可以接受整型变量，比如<<<6, 32>>>代表一个线程网格中有6个线程块，一维排布，一个线程网格块内有32个线程，同样一维排布。这样整个核函数内的总的线程数就是6*32=192。

另外核函数内部可以使用CUDA的内置变量来获取线程号以及线程块号：

threadIdx.[x y z]指的是线程块内线程的编号

blockIdx.[x y z]指的是线程网格内线程块的编号

blockDim.[x y z]指的是线程块的维度，也就是线程块中每个方向上线程的数目

gridDim.[x y z]指的是线程网格的维度，也就是线程网格中每个方向上线程块的数目

下面我们来看一个简单的例子，线程网格有4个线程块，每个线程块内有8个线程，并且都是一维排布：

kernel_name<<<4, 8>>> (argument list)

具体的线程配置以及相应的内置变量的值如下图所示

内置变量均从0开始编号

从上图可以看出，我们可以很轻易的获取一个线程在线程块的位置。在核函数中，我们经常需要得知一个线程在一个线程网格中的位置，那么该怎么计算呢？同样来看一个简单的例子：

dim3 grid(4,1,1), block(4,1,1)

上面的例子中有4个线程块，每个线程块中4个线程，假设我们需要计算红色标记的线程在线程网格中的位置。观察上图，我们可以分成两个部分进程计算，首先计算该线程所在线程块前面总共有多少线程，然后在加上该线程在当前线程块的位置就可以获取在整个线程网格中的位置。该线程所在的线程块编号是blockIdx.x，每个线程块内的线程数是blockDim.x，那总的线程数是blockIdx.x * blockDim.x. 再加上该线程在当前线程块中的位置threadIdx.x，则有：

int idx = blockIdx.x * blockDim.x + threadIdx.x;

二维与三维的情况会稍微复杂些，但计算方法是一样的。下面是一段打印二维线程编号的核函数的例子，自己可以尝试编译运行，相信会有助于对线程位置的计算加深理解。

#include

#include

__global__ void printThreadIndex() {

int ix = threadIdx.x + blockIdx.x * blockDim.x;

int iy = threadIdx.y + blockIdx.y * blockDim.y;

unsigned int idx = iy*blockDim.x * gridDim.x + ix;

printf("thread_id (%d,%d) block_id (%d,%d) coordinate (%d, %d), global index %2d \n",

threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y, ix, iy, idx);

}

int main(void) {

dim3 grid(2, 3, 1), block(4, 8, 1);

printThreadIndex<<>>();

cudaResetDevice();

return 0;

}