【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个线程。
