(三)线程管理(重新整理,前见基础一)
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)
注意这里有几个限制:
- 只能访问GPU内存
- 必须具有void返回类型
- 参数不支持可变数量的数组
- 不支持静态变量
- 显示异步行为
- 从CPU中调用,在计算能力为3的GPU中,也可以通过GPU调用
除此以外,还有2个限定符:
- __device__ int device_function_name(argument list),此类函数只能通过GPU调用
- __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主程序之间并不处于一个线程,所以