1 核函数
先看一个cuda版本的hello world
#include <stdio.h>
__global__ void helloworld()
{
printf("hello world\n");
}
int main()
{
helloworld()<<<1, 1>>>();
cudaDeviceSynchronize();
return 0;
}
这里helloworld()<<<1, 1>>>();就是核函数，关于核函数：
1 核函数（kernel function）在CPU调用，GPU执行
2 核函数必须用__global__修饰，返回值必须是void
3 核函数有两种写法，__global__ void xxx()或者 void __global__ xxx()
4 核函数的几个限制
4.1 核函数只能访问GPU内存
4.2 核函数不能使用变长参数
4.3 核函数不能使用static变量
4.4 核函数不能使用函数指针
4.5 核函数具有异步性
helloworld()<<<1, 1>>>();这行代码是什么意思，我们还不知道，不着急，接着往下看。
2 线程管理
当内核函数开始执行，如何组织GPU的线程就变成了最主要的问题了，我们必须明确，一个核函数只能有一个grid，一个grid可以有很多个块，每个块可以有很多的线程，这种分层的组织结构使得我们的并行过程更加自如灵活：
一个线程块block中的线程可以完成下述协作：
- 同步
- 共享内存
不同块内线程不能相互影响！他们是物理隔离的！
3 线程索引计算方式
3.1 内建变量gridDim / blockDim / blockIdx / threadIdx
内建变量只在核函数有限，且无需定义。
threadIdx是一个uint3类型，表示一个线程的索引。
blockIdx是一个uint3类型，表示一个线程块的索引，一个线程块中通常有多个线程。
blockDim是一个dim3类型，表示线程块的大小。
gridDim是一个dim3类型，表示网格的大小，一个网格中通常有多个线程块。
gridDim和blockDim
gridDim和blockDim是dim3类型(基于uint3定义的数据结构)的变量，都包含三个字段x,y,z，这两个内建变量三个字段的值分别等于初始化时grid_size和block_size的值。如：
dim3 grid_size(2,1,1);
dim3 block_size(3,2,1);
kernel_xxx<<<<gird_size, block_size>>>();
/*
gridDim.x=2, gridDim.y=1, gridDim.z=1
blockDim.x=3, blockDim.y=2, blockDim.z=1
*/
Note: gridDim和blockDim没有指定的维度默认为1
blockIdx和threadIdx
blockIdx和threadIdx是类型为uint3的变量，该类型是一个结构体，包含x,y,z三个字段，其取值范围如下：
blockIdx.x 取值范围[0, gridDim.x -1]
blockIdx.y 取值范围[0, gridDim.y -1]
blockIdx.z 取值范围[0, gridDim.y -1]
threadIdx.x 取值范围[0, blockDim.x -1]
threadIdx.y 取值范围[0, blockDim.y -1]
threadIdx.z 取值范围[0, blockDim.z -1]
3.2 线程索引计算方式
一个Grid可以包含多个Blocks，Blocks的组织方式可以是一维的，二维或者三维的。block包含多个Threads，这些Threads的组织方式也可以是一维，二维或者三维的。所以共有九种方式。
CUDA中每一个线程都有一个唯一的标识ID—ThreadIdx，这个ID随着Grid和Block的划分方式的不同而变化，这里给出Grid和Block不同划分方式下线程索引ID的计算公式。
第一种：一维grid，一维block
int threadId = blockIdx.x *blockDim.x + threadIdx.x;
第二种：一维grid，二维block
int threadId = blockIdx.x * blockDim.x * blockDim.y+ threadIdx.y * blockDim.x + threadIdx.x;
第三种：一维grid，三维block
int threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z
+ threadIdx.z * blockDim.y * blockDim.x
+ threadIdx.y * blockDim.x + threadIdx.x;
第四种：二维grid，一维block
int blockId = blockIdx.y * gridDim.x + blockIdx.x;
int threadId = blockId * blockDim.x + threadIdx.x;
第五种：二维grid，二维block
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y)
第六种：二维grid，三维block
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (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 threadId = blockId * blockDim.x + threadIdx.x;
第八种：三维grid，二维block
int blockId = blockIdx.x + blockIdx.y * gridDim.x
+ gridDim.x * gridDim.y * blockIdx.z;
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 threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x) + threadIdx.x;
3.3 网格大小限制
gridDim.x 最大值 2^31-1
gridDim.y 最大值 2^16-1
gridDim.z 最大值 2^16-1
blockDim.x 最大值 1024
blockDim.y 最大值 1024
blockDim.z 最大值 64
