然而，一个内核可被多个同样大小的线程块执行，所以总的线程数等于每个块内的线程数乘以线程块数。
线程块被组织成一维或二维的线程网格，如图2-1所示。一个网格内的线程块数往往由被处理的数据量而不是系统的处理器数决定，前者往往远超后者。
线程块内线程数和网格内线程块数由<<<...>>>语法确定，参数可以是整形或者dim3类型。二维的块或网格的尺寸可以以和上一个例子相同的方式指定。
网格内的每个块可以通过可在内核中访问的一维或二维索引唯一确定，此索引可通过内置的blockIdx变量获得。块的尺寸(dimension)可以在内核中通过内置变量blockDim访问。
为了处理多个块，扩展前面的MatAdd()例子后，代码成了下面的样子。
__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 < N && j < N)
C[i][j] = A[i][j] + B[i][j];
}
int main()
{
...
// Kernel invocation
dim3 threadsPerBlock(16, 16);
一个长度为16*16(256线程)的块，虽然是特意指定，但是常见。像以前一样，创建了内有足够的块的网格，使得一个线程处理一个矩阵元素。为简便起见，此例假设网格每一维上的线程数可被块内对应维上的线程数整除，尽管这并不常见。
线程块必须独立执行。而且能够以任意顺序，串行或者并行执行。这种独立性要求使得线程块可以以任何顺序在任意数目核心上调度，保证程序员能够写出能够随核心数目扩展的代码(enabling programmers to write code that scales with the number of cores)。
块内线程可通过共享存储器和同步执行协作，共享存储器可以共享数据，同步执行可以协调存储器访问。更精确一点说，可以在内核中调用__syncthreads()内置函数指明同步点;__syncthreads()起栅栏的作用，在其调用点，块内线程必须等待，直到所以线程都到达此点才能向前执行。节3.2.2给出了一个使用共享存储器的例子。
为了能有效协作，共享存储器要求是靠近每个处理器核心的低延迟存储器(更像L1缓存)，而且__syncthreads()要是轻量级的。
