kernel函数即扩展的C函数，通常一个内核函数会被每个定义的CUDA线程调用，所以总共可以执行线程个次数。kernel函数一定要用类型定义符__global__定义，调用时用<<<。。。>>>说明执行参数，如调用的线程数、块数等。kernel的每个thread拥有一个唯一标识的thread ID，每个block拥有一个唯一标识的block ID，用户通过内置的threadIdx、blockIdx变量在kernel中访问。

如图2。1[15]所示，thread可以是一维，二维或三维的，用3维的向量threadIdx索引标识，如图2。1。若是一维的block，thread索引和thread ID相同；对于二维的（Dx,Dy）的block，坐标为（x,y)，则thread的ID为(x+yDx)；在三维 (Dx,Dy,Dz)的block中，坐标（x,y,z)的thread ID为（x+yDx+zDxDy）。每个block通过一维或二维索引内置变量blockIdx访问。block的大小可在kernel中通过内置变量blockDim访问得到。

图2。1  CUDA编程模型

一个block包含多个最小执行单位thread，只有同一个block中的线程能够共享同一块存储区域， thread若在两个不同block中则不能直接传递数据、消息。但是我们仍然需要调用多个block共同完成任务，如顺序执行，而非同时执行各个block中的thread，这样可以克服一个block中的thread数目不足的缺陷。在kernel中调用__syncthreads( )函数，起线程同步的功能，即调用时必须所有block内thread都执行到这儿才能继续执行下一步。

2。2。2  CUDA存储器模型

图2。2  CUDA存储器模型

CUDA的存储器模型[15]如图2。2，开发者应该根据它们的特点与局限性，有选择性地使用。根据它们的访问权限，大致可以分为三种类型的存取器。

寄存器（Registers）和本地存储器（local memory）是私有的，只能被自己的线程访问，属于第一层存储器。寄存器的访问速度很快，延迟极低，但数量有限，当寄存器不足时，就会调用本地存储器存储数据，但本地存储器的存取速度有很大延迟。

共享存储器（Shared Memory）属于第二层，可以被同一block所有thread共享，位于GPU片内的高速存储器[13]。

全局存储器（Global Memory）、常量存储器（Constant Memory）、纹理存取器（Texture Memory）[16]属于第三层，可被所有thread访问，它们成本低但容量大，访问速度远低于Registers和local memory。全局存储器的容量是GPU中最大的，是GPU与CPU交互的桥梁。

2。2。3  CUDA工具库

CUDA提供了丰富的工具库[17]，主要包括：

基于CUDA实现的线代库cuBLAS[18]，以column-major的方式存储，利用GPU加速向量、矩阵的线性运算；

库CUFFT[18]提供了一系列帮助开发者进行Fourier变换的运算的函数，CUFFT库分别由CUFFT和CUFFTW两个构成。CUFFTW库为研究人员准备了一些接口，可让用户在CUDA GPU上编写使用FFTW库的程序，CUFFT是纯CUDA接口的傅里叶变换库；

能生成高质量伪随机数的库cuRAND[18]，它通过创建一组伪随机数并每次简单地从列表中选择下一个元素，从而实现随机数的生成，任何指导生成随机数所使用的种子的人都可以准确的预测出给定序列中的下一个随机数；文献综述

库cuSPARSE[18]，是一个线性代数库，对稀疏矩阵之类的操作有独到的用法，使用很宽泛，对稠密和稀疏的数据格式都支持[13]。

2。3  CUDA软硬件架构

2。3。1  CUDA软件体系

CUDA的软件体系大致由CUDA Library（CUDA库函数）、CUDA runtime API（CUDA运行时API）、CUDA driver API（CUDA驱动API）三层组成[15]。CUDA C语言是对C语言的最小扩展集，是CUDA的核心，必须通过nvcc编译器编译才能使用这些扩展和运行时库的源文件。