但一个内核可由多个大小相同的线程块执行，因而线程总数应等于每个块的线程数乘以块的数量。这些块将组织为一个一维或二维线程块网格，如图 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)。
次修订号对应于核心架构的增量式改进，可能包含新特性。
