【IT168 文档】以下内容摘自CUDA编程手册 版本2.1 并参考并引用了2.0中文版编程手册中的内容，并修正了一些术语和错误。
本章节介绍CUDA编程模型的主要的概念并勾画出其是如何以C的形式进行表述。关于对应CUDA的C的描述的更加详细的内容将会在第四章给出。
2.1 内核
CUDA 允许程序员定义称为内核(kernel)的 C 语言函数，从而扩展了 C 语言，在调用此类函数时，它将由 N 个不同的 CUDA 线程并行执行 N 次，这与普通的 C 语言函数只执行一次的方式不同。
在定义内核时，需要使用 _global_ 声明定义符，使用一种全新的 <<<…>>> 语法指定每次调用的 CUDA 线程数：
__global__ void vecAdd(float* A, float* B, float* C)
{
}
int main()
{
// 内核调用，N这里是线程的数目
vecAdd<<<1, N>>>(A, B, C);
}
执行内核的每个线程都会被分配一个独特的线程 ID，可通过内置的 threadIdx 变量在内核中访问此 ID。以下示例代码将大小为 N 的向量 A 和向量 B 相加，并将结果存储在向量 C 中：
{
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
int main()
{
// 内核调用
vecAdd<<<1, N>>>(A, B, C);
}
执行 vecAdd( ) 的每个线程都会执行一次成对的加法运算。
2.2 线程层次结构
为方便起见，我们将 threadIdx 设置为一个包含 3 个组件的向量，因而可使用一维、二维或三维索引标识线程，构成一维、二维或三维线程块。这提供了一种自然的方法，可为一个域中的各元素调用计算，如向量、矩阵或字段。下面的示例代码将大小为 N x N 的矩阵 A 和矩阵 B 相加，并将结果存储在矩阵 C 中：
{
int i = threadIdx.x;
int j = threadIdx.y;
C[i][j] = A[i][j] + B[i][j];
}
int main()
{
// Kernel invocation
dim3 dimBlock(N, N);
matAdd<<<1, dimBlock>>>(A, B, C);
}
线程的索引及其线程 ID 有着直接的关系：对于一维块来说，两者是相同的;对于大小为 (Dx，Dy) 的二维块来说，索引为 (x,y) 的线程的ID 是 (x + yDx);对于大小为 (Dx，Dy， Dz) 的三维块来说，索引为 (x, y, z) 的线程的ID 是 (x + y Dx +z DxDy)。
一个块内的线程可彼此协作，通过一些共享存储器来共享数据，并同步其执行来协调存储器访问。更具体地说，可以通过调用 __syncthreads()__ 内蕴函数在内核中指定同步点;__syncthreads()__ 起到栅栏的作用，块中的所有线程都必须在这里等待进一步的处理。
为实现有效的协作，共享存储器应该是接近各处理器核心的低延迟存储器，最好是类似 L1 缓存这样的，__syncthreads()__ 应是轻量级的，一个块中的所有线程都必须驻留在同一个处理器核心中。因而，一个处理器核心的有限存储器资源制约了每个块的线程数量。在 NVIDIA Tesla 架构中，一个线程块最多可以包含 512 个线程。
