为方便起见，我们将 threadIdx 设置为一个包含 3 个组件的向量，因而可使用一维、二维或三维缩影标识线程，构成一维、二维或三维线程块。这提供了一种自然的方法，可为一个域中的各元素调用计算，如向量、矩阵或字段。下面的示例代码将大小为 NxN 的矩阵 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[j] = A[j] + B[j];
}
int main()
{
// Kernel invocation
dim3 dimBlock(N, N);
matAdd>>(A, B, C);
}
线程的索引及其线程 ID 有着直接的关系：对于一维块来说，两者是相同的；对于大小为 (Dx，Dy) 的二维块来说，索引为 (x,y) 的线程的ID 是 （x + yDx）；对于大小为 （Dx，Dy， Dz） 的三维块来说，索引为 （x, y, z） 的线程的ID 是 （x + yDx + ZDxDy）。
一个块内的线程可彼此协作，通过一些共享存储器来共享数据，并同步其执行来协调存储器访问。更具体地说，可以通过调用 _syncthreads()_ 内函数在内核中指定同步点；_syncthreads()_ 起到屏障的作用，块中的所有线程都必须在这里等待处理。
为实现有效的协作，共享存储器应该是接近各处理器核心的低延迟存储器，如 L1 缓存，_syncthreads()_ 应是轻量级的，一个块中的所有线程都必须位于同一个处理器核心中。因而，一个处理器核心的有限存储器资源制约了每个块的线程数量。在 NVIDIA Tesla 架构中，一个线程块最多可以包含 512 个线程。
但一个内核可能由多个大小相同的线程块执行，因而线程总数应等于每个块的线程数乘以块的数量。这些块将组织为一个一维或二维线程块网格，如图 2-1 所示。该网格的维度由 >> 语法的第一个参数指定。网格内的每个块多可由一个一维或二维索引标识，可通过内置的 blockIdx 变量在内核中访问此索引。可以通过内置的 blockDim 变量在内核中访问线程块的维度。此时，之前的示例代码应修改为：
__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 >>(A, B, C);
}
我们随机选择了大小为 16x16 的线程块（即包含 256 个线程），此外创建了一个网格，它具有足够的块，可将每个线程作为一个矩阵元素，这与之前完全相同。
线程块需要独立执行：必须能够以任意顺序执行、能够并行或顺序执行。这种独立性需求允许跨任意数量的核心安排线程块，从而使程序员能够编写出可伸缩的代码。
一个网格内的线程块数量通常是由所处理的数据大小限定的，而不是由系统中的处理器数量决定的，前者可能远远超过后者的数量。
|
