一、为什么线程要分层组织，不能"一大堆线程扁平地摆着"上一节已经知道，一次kernel调用会启动"gridDim × blockDim"个线程， 这里为什么要设计成"Grid包含Block、Block包含Thread"这种 两层嵌套结构，而不是直接给一个线程总数就完事？ 原因主要有两个： ① 硬件层面天然如此 GPU的计算单元本身就是分组的（Streaming Multiprocessor，SM）， 一个Block在执行时，会被整体分配到某一个SM上执行 （不会被拆散分到多个SM），Block内部的线程之间可以 通过共享内存（Shared Memory）互相通信、用__syncthreads() 互相同步——这些"线程间协作"的能力，只在同一个Block内部成立， 不同Block之间的线程默认没有直接协作的手段 → 所以"哪些线程属于同一个Block"这件事， 直接决定了"哪些线程能协作、哪些不能"， 这是一个编程模型上必须显式表达出来的分组 ② 表达数据的多维结构更方便 很多要处理的数据本身是多维的（图像是二维像素矩阵， 体数据是三维体素），Grid和Block都支持用二维、三维的 方式组织线程，这样线程的坐标可以直接对应数据的坐标 （比如像素的行列号），不需要额外做一次"一维下标转二维坐标" 的换算，写代码更直观二、三个层次分别是什么2.1 Thread（线程）最基本的执行单位，一个线程执行一遍kernel函数体里的代码， 处理属于自己的那一份数据。2.2 Block