【IT168 技术文档】本章引入了CUDA编程模型背后的主要概念，方式是概述它们是怎样使用C语言表示的。更多的关于CUDA C的描述在3.2节。
本章使用的向量相加例子的完整代码和下一个例子可在SDK中的vectorAdd代码样本中找到。
2.1 内核
CUDA通过允许程序员定义称为内核的C函数扩展了C，内核调用时会被N个CUDA线程执行N次(注：这句话要好好理解，其实每个线程只执行了一次)，这和普通的C函数只执行一次不同。
内核使用__global__声明符定义，使用一种新<<<...>>>执行配置语法指定执行某一指定内核调用的线程数(参看附录B.13)。每个执行内核的线程拥有一个少有的线程ID，可以通过内置的threadIdx变量在内核中访问。
下面的样本代码将两个长度为N的向量A和B相加，并将结果存入向量C中。
__global__ void VecAdd(float* A, float* B, float* C)
{
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main()
{
...
// Kernel invocation with N threads
这里，N个线程中的每一个执行VecAdd()的一次成对加法。
2.2 线程层次
为简便起见，threadIdx是一个有3个部件的向量，所以线程可以使用一维，二维，三维索引标识，形成一维，二维，三维的线程块。这提供了一种自然的方式来调用作用在域内元素上的计算，如向量，矩阵，体元(volume)。
线程索引和线程ID直接相关：对于一维的块，它们相同;对于二维长度为(Dx,Dy)的块，线程索引为(x,y)的线程ID是(x+yDx);对于三维长度为(Dx,Dy,Dz)的块，索引为(x,y,z)的线程ID为(x+yDx+zDxDy)。
下面的例子代码将两个长度为N*N的矩阵A和B相加，然后将结果写入矩阵C。
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
int i = threadIdx.x;
int j = threadIdx.y;
C[i][j] = A[i][j] + B[i][j];
}
int main() {
...
// Kernel invocation with one block of N * N * 1 threads
int numBlocks = 1;
dim3 threadsPerBlock(N, N);
由于块内的所有线程必须存在于同一个处理器核心中且共享该核心有限的存储器资源，因此，一个块内的线程数目是有限的。在目前的GPU上，一个线程块可以包含多达512个线程。
然而，一个内核可被多个同样大小的线程块执行，所以总的线程数等于每个块内的线程数乘以线程块数。
线程块被组织成一维或二维的线程网格，如图2-1所示。一个网格内的线程块数往往由被处理的数据量而不是系统的处理器数决定，前者往往远超后者。
线程块内线程数和网格内线程块数由<<<...>>>语法确定，参数可以是整形或者dim3类型。二维的块或网格的尺寸可以以和上一个例子相同的方式指定。
网格内的每个块可以通过可在内核中访问的一维或二维索引唯一确定，此索引可通过内置的blockIdx变量获得。块的尺寸(dimension)可以在内核中通过内置变量blockDim访问。
为了处理多个块，扩展前面的MatAdd()例子后，代码成了下面的样子。
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N)
C[i][j] = A[i][j] + B[i][j];
}
int main()
{
...
// Kernel invocation
dim3 threadsPerBlock(16, 16);
一个长度为16*16(256线程)的块，虽然是特意指定，但是常见。像以前一样，创建了内有足够的块的网格，使得一个线程处理一个矩阵元素。为简便起见，此例假设网格每一维上的线程数可被块内对应维上的线程数整除，尽管这并不常见。
线程块必须独立执行。而且能够以任意顺序，串行或者并行执行。这种独立性要求使得线程块可以以任何顺序在任意数目核心上调度，保证程序员能够写出能够随核心数目扩展的代码(enabling programmers to write code that scales with the number of cores)。
块内线程可通过共享存储器和同步执行协作，共享存储器可以共享数据，同步执行可以协调存储器访问。更精确一点说，可以在内核中调用__syncthreads()内置函数指明同步点;__syncthreads()起栅栏的作用，在其调用点，块内线程必须等待，直到所以线程都到达此点才能向前执行。节3.2.2给出了一个使用共享存储器的例子。
为了能有效协作，共享存储器要求是靠近每个处理器核心的低延迟存储器(更像L1缓存)，而且__syncthreads()要是轻量级的。
2.3 存储器层次
在执行期间，CUDA线程可能访问来自多个存储器空间的数据，如图2-2所示。每个线程有私有的本地存储器。每个块有对块内所有线程可见的共享存储器，共享存储器的生命期和块相同。所有的线程可访问同一全局存储器。
另外还有两种可被所有线程访问的只读存储器：常量和纹理存储器空间。全局，常量和纹理存储器空间为不同的存储器用途作了优化(参看5.3.2.1节,5.3.2.4节和5.3.2.5节)。纹理存储器还为一些特殊数据格式提供了不同的寻址模式和数据滤波(参看3.2.4节)。
在同一应用发射的内核之间，全局，常量和纹理存储器空间是持久的。
2.4 异构编程
如图2-3所示，CUDA编程模型假设CUDA线程在物理上独立的设备上执行，设备作为主机的协处理器，主机运行C程序。例如，内核在GPU上执行，而C程序的其它部分在CPU上执行就是这种模式。
CUDA编程模型同时假设主机和设备各自都维护着自己独立的DRAM存储器空间，各自被称为主机存储器空间和设备存储器空间。因此，程序通过调用CUDA 运行时，来管理对内核可见的全局、常量和纹理存储器空间(参看第三章)。这包括设备存储器分配和释放，也包括在主机和设备间的数据传输。
2.5 计算能力
设备的计算能力由主修订号和次修订号定义。
主修订号相同的设备基于相同的核心架构。Fermi架构的主修订号为2。以前的设备的计算能力都是1.x(它们的主修订号为1)。
次修订号对应着对核心架构的增量提升，也可能包含了新特性。
附录A列出了所有支持CUDA的设备，包括它们的计算能力。附录G给出了各计算能力设备的技术规范。
