cuda并行编程基础（三）

（三）线程管理（重新整理，前见基础一）

1.前面我们提到过网格，一个内核函数启动产生的所有线程称为网格，同一网格内线程可共享相同的全局内存空间。

2.多个线程块则组成网格，同一线程块内的的线程可以通过同步或共享内存来实现协作。

3.最后就是线程了，每个线程具有唯一的标识，通过blockIdx和threadIdx来区分。

gridIdx.x/gridIdx.y/gridIdx.z 理解为某线程块在网络的三维坐标

blockIdx.x/blockIdx.y/blockIdx.z 理解为某线程在线程块的三维坐标

有了他们，就可以在网格当中准确定位某线程了。

如果有必要的话，可以获取网格的维度gridDim和线程块的维度blockDim，那么容易知道blockIdx肯定要小于或等于gridDim,同样threadIdx要小于或等于blockDim。

一般情况下，我们会将网格组织为二数数组，而线程块且组织为三维数组，以上均为dim3类型，默认情况下，各维索引均为1.

（四）编写核函数

核函数定义：__global__ void kernel_name(argument list)

注意这里有几个限制：

1. 只能访问GPU内存
2. 必须具有void返回类型
3. 参数不支持可变数量的数组
4. 不支持静态变量
5. 显示异步行为
6. 从CPU中调用，在计算能力为3的GPU中，也可以通过GPU调用

除此以外，还有2个限定符：

1. __device__ int device_function_name(argument list)，此类函数只能通过GPU调用
2. __host__ int host_function_name(argument list)，此类函数只能通过CPU调用，跟平时程序一样，也可以忽略__host__限定符。

这里限定符__device__和__host__可以一起使用，这样的话，函数既可以用CPU调用，也可以在GPU里调用。

（五）错误处理

一般CUDA函数都会返回一个错误值，如函数：

cudaError_t cudaMemcpy(void *dst,const void *src,size_t count,cudaMemcpyKind kind);

它返回的是一个错误码：cudaError_t error，错误码如果不熟悉的话，我们可以通过以下函数获得具体错误信息：

Char *cudaGetErrorString(cudaError_t error);

其中error==cudaSucess或error==0代表运行成功。

（六）确定核函数性能

核函数运行后，我们怎么知道它要跑多长时间呢？前面我们知道核函数与CPU主程序之间并不处于一个线程，所以