CUDA之所以能实现并行计算，主要得益于开发人员能够对GPU计算资源有效分配。
CUDA的粒度有三种，从大到小以此为: grid、block、thread
三者关系如下：

Grid好比一个小区，block是这个小区的每一栋房子，而thread则是每栋房子里的小隔间。我们要进行工作的最小单位就是一个小隔间，即thread。

咱们可以从一个简单的cuda程序入手，了解这几个概念：

#include <cuda_runtime.h>
#include <stdio.h>
__global__ void checkIndex(void)
{
  printf("threadIdx:(%d,%d,%d) blockIdx:(%d,%d,%d) blockDim:(%d,%d,%d)\
  gridDim(%d,%d,%d)\n",
  threadIdx.x,threadIdx.y,threadIdx.z,
  blockIdx.x,blockIdx.y,blockIdx.z,
  blockDim.x,blockDim.y,blockDim.z,
  gridDim.x,gridDim.y,gridDim.z);
}
int main(int argc,char **argv)
{
  int nElem=6;
  dim3 thread(3);
  dim3 block((nElem + thread.x - 1) / thread.x);
  printf("block.x %d block.y %d block.z %d\n", block.x, block.y, block.z);
  printf("thread.x %d thread.y %d thread.z %d\n", thread.x, thread.y, thread.z);
  checkIndex<<<block, thread>>>();
  cudaDeviceReset();
  return 0;
}

将代码保存为grain.cu，然后编译运行：

nvcc -o grain grain.cu
./grain

上面代码表示，将6个元素安排给1个grid中的2个block中的2*3个thread。用更白话的方式说：咱们在GPU上分配了一个维度是(2, 1, 1)的grid，表示grid含有2x1x1=2个block，其中每个block的维度是(3, 1, 1)，表示block含有3x1x1=3个thread。

而在kernel函数（即被__global__修饰，可以定义每个thread上运行的程序。）使用threadID、blockID、gridID等可以直接获取当前的线程号、区块号和网格号，从而使当前thread知道自己该计算哪一块内容，就可以避免每个thread做一样的事情了。