核函数与线程层级
- CUDA核函数
- 线程层级
- 线程排布
- blockId和threadId的计算
CUDA核函数
在GPU上并行执行的函数称为CUDA核函数（Kernel Function)，它属于CUDA编程中最为重要且核心的一个环节。
核函数用__global__符号声明，在devie（GPU）上执行，在host（CPU）上调用，返回类型必须时void，不支持可变参数，不支持静态变量，不支持函数指针，核函数相对于CPU是异步的，在核函数执行完之前就会返回，这样CPU可以不用等待核函数的完成，继续执行后续代码。
核函数在调用时需要用<<<Dg, Db>>>来指定kernel要执行的线程数量，不同的GPU架构，grid和block的维度有限度。在host端核函数的调用方式为：
kernel_function<<<Dg, Db, Ns, S>>>(param list);
其中，
- Dg：int型或者dim3类型(x,y,z)，用于定义一个Grid中Block是如何组织的，如果是int型，则表示一维组织结构。
- Db：int型或者dim3类型(x,y,z)，用于定义一个Block中Thread是如何组织的，如果是int型，则表示一维组织结构。
- Ns：size_t类型，可缺省，默认为0；用于设置每个block除了静态分配的共享内存外，最多能动态分配的共享内存大小，单位为byte。0表示不需要动态分配。
- S：cudaStream_t类型，可缺省，默认为0。表示该核函数位于哪个流。
在CUDA中，每一个线程都要执行核函数，并且每个线程会分配一个唯一的线程号thread ID，这个ID值可以在CUDA上可以使用以下列举的内置变量来获取Thread ID和Block ID：
threadIdx.[x, y, z]表示Block内Thread的x或y或z三者其中一个维度的编号。blockIdx.[x, y, z]表示Gird内Block的x或y或z三者其中一个维度的编号。blockDim.[x, y, z]表示Block的维度，也就是Block中x或y或z三者其中一个维度上的Thread的数目。gridDim.[x, y, z]表示Gird的维度，也就是Grid中x或y或z三者其中一个维度上Block的数目。
代码示例：
// 下面来定义一个核函数，用来做向量加法，这里同时做了declaration和definition
// 正常的C++语法定义一个void函数，接受三个参数A、B、C，均为浮点数的指针
// 请注意这个核函数将会并行化地运行，其中每个thread负责向量中的一个位置，做逐位加法
__global__ void VecAdd(float* A, float* B, float* C) // kernel函数需要在最前面加上__global__关键字
{
int i = threadIdx.x; // threadIdx是一个内置的变量，告诉我们当前运行这个函数的thread的ID
C[i] = A[i] + B[i]; // 对于第i个线程，它的任务就是把A向量的第i个元素与B的第i个元素相加，写到C的第i个元素
}
int main()
{
int gs = 1;
int bs = 4;
...
// 调用核函数，这里声明我们想使用gs个grid，bs个thread
VecAdd<<<gs, bs>>>(A, B, C);
...
}
线程层级
线程排布
Grid/Block/Thread都是软件的组织结构，并不是硬件的，因此理论上我们可以以任意的维度（一维、二维、三维）去排列Thread。
当使用int类型时表示一维排布时（注：也可以使用dim3表示一维排布）：
int grid = 4; // dim3 grid(4);
int block = 8;// dim3 block(8);
kernel_name<<<grid,block>>>(param list);
表示一个Grid中有4个Block，在(x,y,z)三个方向上的排布方式分别是4、1、1；一个Block中有8个Thread，在(x,y,z)三个方向上的排布方式分别是8、1、1。
当使用dim3类型表示二维排布时：
dim3 grid(3,2), block(4,3);
kernel_name<<<grid, block>>>(param list);
表示一个Grid中有3x2x1=6个Block，在(x,y,z)三个方向上的排布方式分别是3、2、1；一个Block中有4x3x1=12个Thread，在(x,y,z)三个方向上的排布方式分别是4、3、1。
当使用dim3类型表示三维排布时：
dim3 grid(3,2,2), block(4,3,1);
kernel_name<<<grid, block>>>(param list);
表示一个Grid中有3x2x2=12个Block，在(x,y,z)三个方向上的排布方式分别是3、2、1；一个Block中有4x3x1=12个Thread，在(x,y,z)三个方向上的排布方式分别是4、3、1。
blockId和threadId的计算
若Grid排布为一维，Block排布为一维：
int blockId = blockIdx.x
int threadId = blockIdx.x *blockDim.x + threadIdx.x
若Grid排布为二维，Block排布为二维：
int blockId = blockIdx.x + blockId.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y *blockDim.x) + threadIdx.x;
若Grid排布为三维，Block排布为三维：
int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;
int threadIc = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x) + threadIdx.x;
