CUDA线程布局和内存层次
Caution
TODO : CUDA 中 CUDA Core 硬件结构(SM/SP等)与软件层面布局对应关系及介绍
CUDA线程布局：
如图所示，CUDA线程布局分为三层：网格(Grid),线程块(Block)以及线程(thread)
Important
在计算机中，内存的访问是一维的，线程的访问实质上也是一维的。CUDA对硬件进行了抽象，grid和block是为了方便进行线程管理而抽象出来的线程组织模式。每个核函数能够调度的所有线程在一个核函数唯一对应的grid中，一个grid中包含了多个Block，每个Block包含了相同数量的线程。
网格：同一线程网格中的所有线程共享全局内存空间，一个网格有多个线程块（Block）构成。一个网格有效地代表了一个kernel启动，即它包含了要为一个特定内核启动运行的所有块(以及线程)。
线程块：一个线程块包含一组线程，同一线程块内的线程协同可以通过“同步”和“共享内存”的方式来实现。线程块之间彼此独立，不同线程块内的线程不能相互影响，是"物理隔离"的。每个block会由一个专门的SM进行运算。由于warp是SM的基本执行单元，当前一个warp都是32个线程，所以block被分配给SM后还会根据warp的大小将一个block分成多次执行。
线程：线程是CUDA架构下GPU芯片执行时的最小单位，一个block中的thread能存取同一块共享内存，可以快速进行同步和通信操作。
- 另外：为了便于定位thread和block，需要对thread和block建立索引作为它们的唯一标识。在一个网格中，我们通过blockIdx和threadIdx两个坐标变量来定位一个线程，通过gridDim和blockDim确定grid和block维度大小。通过这些索引，可以构成对每个线程在一维、二维或三维的唯一标识。如一维下的线程标识为：int globalThreadId = blockIdx.x * blockDim.x + threadIdx.x;
CUDA内存层次：
如图所示，CUDA内存包括以下几个层次：
全局内存（Global Memory）
- 全局内存是设备内存，可以被Grid中所有线程访问。全局内存的访问速度较慢，但容量较大，适合存储大量数据。在Host端或Device端使用__device__声明。
常量内存（Constant Memory）
- 常量内存是一种只读内存，用于存储常量数据。常量内存的访问速度较快，但容量较小，在Host端使用__constant__声明，不可以在核函数内声明。
纹理内存（Texture Memory）
- 纹理内存是一种只读内存，主要用于存储图像数据。纹理内存具有缓存特性，可以加速对图像数据的访问。
共享内存（Shared Memory）
- 共享内存是线程块内的所有线程共享的内存。共享内存的访问速度较快，可以用于线程间的数据交换和临时数据存储。每个线程块的共享内存容量有限，在核函数中使用__shared__可以声明共享内存变量。
本地内存（Local Memory）
- 本地内存是每个线程私有的内存，主要用于存储寄存器溢出数据和局部变量。本地内存实际上位于全局内存中，访问速度较慢。
寄存器（Registers）
- 寄存器是每个线程私有的高速存储器，用于存储局部变量和临时数据。寄存器的访问速度最快，但数量有限。