【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 个线程。
但一个内核可由多个大小相同的线程块执行，因而线程总数应等于每个块的线程数乘以块的数量。这些块将组织为一个一维或二维线程块网格，如图 2-1 所示。该网格的维度由 <<<…>>> 语法的第一个参数指定。网格内的每个块多可由一个一维或二维索引标识，可通过内置的 blockIdx 变量在内核中访问此索引。可以通过内置的 blockDim 变量在内核中访问线程块的维度。此时，之前的示例代码应修改为：
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()
{
// 调用内核
dim3 dimBlock(16, 16);
dim3 dimGrid((N + dimBlock.x – 1) / dimBlock.x,
(N + dimBlock.y – 1) / dimBlock.y);
matAdd<<<dimGrid, dimBlock>>>(A, B, C);
}
我们随机选择了大小为 16x16 的线程块(即包含 256 个线程)，此外创建了一个网格，它具有足够的块，可将每个线程作为一个矩阵元素，这与之前完全相同。
线程块需要独立执行：必须能够以任意顺序执行、能够并行或顺序执行。这种独立性需求允许跨任意数量的核心安排线程块，从而使程序员能够编写出可伸缩的代码。
一个网格内的线程块数量通常是由所处理的数据大小限定的，而不是由系统中的处理器数量决定的，前者可能远远超过后者的数量。
2.3 存储器层次结构
CUDA 线程可在执行过程中访问多个存储器空间的数据，如图 2-2 所示。每个线程都有一个私有的本地存储器。每个线程块都有一个共享存储器，该存储器对于块内的所有线程都是可见的，并且与块具有相同的生命周期。最终，所有线程都可访问相同的全局存储器。
此外还有两个只读的存储器空间，可由所有线程访问，这两个空间是固定存储器空间和纹理存储器空间。全局、固定和纹理存储器空间经过优化，适于不同的存储器用途(参见第 5.1.2.1、5.1.2.3 和 5.1.2.4)。纹理存储器也为某些特殊的数据格式提供了不同的寻址模式以及数据过滤(参见第 4.3.4)。
对于同一个应用程序启动的内核而言，全局、固定和纹理存储器空间都是持久的。
2.4 主机和设备
如图 2-3 所示，CUDA 假设 CUDA 线程可在物理上独立的设备上执行，此类设备作为运行 C 语言程序的主机的协同处理器操作。例如，当内核在 GPU 上执行，而 C 语言程序的其他部分在 CPU 上执行时，就是这样一种情况。
此外，CUDA 还假设主机和设备均维护自己的 DRAM，分别称为主机存储器和设备存储器。因而，一个程序通过调用 CUDA 运行时来管理对内核可见的全局、固定和纹理存储器空间(详见第 4 章)。这包括设备存储器分配和取消分配，还包括主机和设备存储器之间的数据传输。
串行代码在主机上执行，而并行代码在设备上执行。
2.5 计算能力
一个设备的计算能力(compute capability)由主修订号和次修订号定义。
具有相同主修订号的设备属于相同的核心架构。附录 A 中列举的设备均为计算能力是 1.x 的设备(其主要修订号为 1)。
次修订号对应于核心架构的增量式改进，可能包含新特性。
