__golobal__修饰的函数为核函数，需要从CPU调用，函数本身在GPU计算。
核函数的调用形式为func <<
核函数中可以调用CUDA的内置变量（gridDim，blockDim，blockIdx，threadIdx）。其中gridDim和blockDim的含义就是DG和DB的含义，blockIdx表示当前线程所属block的序号（即一个Block下的所有线程共享这一个blockIdx），threadIdx表示当前线程在当前Block中的序号。可以得知在这种情况下threadIdx序号是会重复的，例如Block1的thread0线程和Block2的thread0线程，此时需要从局部线程序号换算到全局线程序号上，如下图：
很好理解，blockDim*blockIdx计算当前Block之前的所有线程数目，加上当前的threadIdx即全局唯一的的线程索引。启动线程数可能大于元素个数，可以用判断语句限定线程运算。
事实上上述的DG和DB并不是两个数，而是CUDA中的数据类型dim3，即最多可以创建三维的grid或者block，分别表示为gridDim.x, gridDim.y, gridDim.z, blockDim.x, blockDim.y和blockDim.z，例如指定gridDim(2，2)时表示gridDim.x=2，y=2,z维度为1，此时blockDim(2,4,2)，如下图所示：
需要注意一个block最多包含1024个线程，即blockDim的三个维度乘积不能大于1024。
线程束任意时刻的任务执行都是以线程束warp为单位的，warp_size的大小通常为32，因此block通常设置成32的倍数，warp中的线程有两种状态即活跃（启动的线程）和不活跃（未启动的线程）。
对于gridDim为m，blockDim为n的任务，总共需要的warp数量为m * ceil ( n / warp_size )，ceil表示向上取整。 实际上的物理warp数量有限（68SM，64core per SM，total warp=68*(64/2)=136.因此理论上所有运算并行，但实际上warp调度是有先后顺序的，计算一定程度上也是有先后的。同个block的warp需要一定程度的同步和并行，grid内的block则独立、并行，以下图的按线程赋值打印程序说明：
程序说明：
可以看到void thread()函数中分配了五个线程（1个block * 5个thread）用来执行核函数，此时开启了1个warp，其中5个thread处于active状态，剩余的处于inactivate状态。五个线程同步并行运行核函数。
第一个时钟周期：五个线程同步执行第一行代码，对array数组进行赋值：array[0]=0;array[1]=1...
第二个时钟周期：五个线程同步执行第二行代码，该行代码打印step1开头的语句，打印变量则根据线程变化：step1 0；step1 1；step1 2....
第三和第四时钟周期与第二时钟周期一致，打印前缀变成了step2和step3。因此最后的打印结果如第一张图所示：
线程过多时，可能存在线程轮替执行，因此实际上并不是所有线程都能同步执行的。用__syncthreads可以使线程同步。
warp分歧的逻辑处理：
前面可以知道同一个warp下的线程运行通常都是同步的，但如果遇到分支语句if等，就会出现分歧，warp在这种情况下的处理逻辑就是：首先不满足if的往下执行分支，执行完了就等待；2.执行满足if条件的分支；3.分支都执行完了在同步往下执行。
如上图程序：线程0、1、2、3、4满足if分支%2=0的为0、2、4，不满足的为1、3。因此先执行不满足的分支，打印了step2 1和step2 3，执行完该分支后线程13进行等待；执行满足if分支的0、2、4线程，打印了step1 0，step1 2，step1 4，执行完后和13线程同步执行打印step3的操作，打一局结果如上图所示：
