CUDA 编程模型采用了一个三层的线程层次结构，旨在映射到 GPU 硬件的多级架构，实现最大的并行性和数据局部性。
1. 线程 (Thread)
线程是 CUDA 并行计算的基本执行单元。
定义：在 Kernel 函数中，每个并行计算的实例就是一个线程。例如，在向量加法中，一个线程负责计算C[i]=A[i]+B[i]C[i] = A[i] + B[i]C[i]=A[i]+B[i]。
标识：每个线程都有一个内置的唯一标识符
threadIdx。threadIdx.x、threadIdx.y、threadIdx.z：线程在当前线程块内的坐标。
执行：线程执行 Kernel 代码中定义的指令，访问私有寄存器（Registers）和线程块共享内存（Shared Memory）。
独立性：理想情况下，每个线程应独立执行其任务，减少相互依赖，以实现最大并行效率。
2. 线程块 (Block)
线程块是线程的分组容器，是线程间协作的基本单位。
定义：一组协作的线程集合，它们共同执行 Kernel 的一个子任务。
标识：每个线程块都有一个内置的唯一标识符
blockIdx。blockIdx.x、blockIdx.y、blockIdx.z：线程块在网格内的坐标。
维度：线程块可以定义为一维、二维或三维结构（
dim3 blockDim;）。blockDim.x、blockDim.y、blockDim.z：定义了线程块中线程的数量和排列方式。
协作与通信：
共享内存：块内的线程可以通过共享内存（Shared Memory）进行高速数据交换。
同步屏障：块内的线程可以通过
__syncthreads()函数进行同步，确保所有线程都到达某个执行点后才能继续，这对于数据依赖的算法至关重要。
硬件映射：一个线程块内的所有线程保证会被调度到同一个SM（流多处理器）上执行。这保证了块内线程对共享内存的访问非常快。
3. 网格 (Grid)
网格是线程块的最高级容器，代表一次完整的 Kernel 启动。
定义：由所有线程块组成的集合，是整个并行计算任务的总和。
标识：网格没有内置的索引变量，它的维度由主机在 Kernel 启动时通过
gridDim参数指定。gridDim.x、gridDim.y、gridDim.z：定义了网格中线程块的数量和排列方式。
协作与通信：网格中的不同线程块原则上是相互独立的。
无法直接同步：不同线程块之间不能使用
__syncthreads()进行同步。全局内存通信：线程块之间只能通过访问速度较慢的全局内存（Global Memory）进行通信。如果需要全局同步，必须终止当前 Kernel，在主机端同步后，再次启动一个新的 Kernel。
硬件映射：网格中的线程块由 GPU 调度器分发到不同的 SM 上并行执行。
4. 如何计算全局唯一索引
理解线程层次结构的关键在于知道如何将线程的局部坐标 (blockIdx和threadIdx) 转换为它在整个网格中的全局唯一索引，即它应该处理的数据元素的索引iii。
4.1 一维索引计算
对于大多数简单的一维数据结构（如向量），全局索引iii的计算公式如下：
i=blockIdx.x×blockDim.x+threadIdx.xi = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}i=blockIdx.x×blockDim.x+threadIdx.x
| 变量 | 含义 |
|---|---|
| blockIdx.x\text{blockIdx.x}blockIdx.x | 当前块的索引 |
| blockDim.x\text{blockDim.x}blockDim.x | 每个块的线程数 |
| threadIdx.x\text{threadIdx.x}threadIdx.x | 线程在块内的索引 |
4.2 二维索引计算（例如矩阵）
对于二维数据结构（如矩阵），我们通常需要计算两个索引rrr（行）和ccc（列）：
r=blockIdx.y×blockDim.y+threadIdx.yc=blockIdx.x×blockDim.x+threadIdx.xr = \text{blockIdx.y} \times \text{blockDim.y} + \text{threadIdx.y} \\ c = \text{blockIdx.x} \times \text{blockDim.x} + \text{threadIdx.x}r=blockIdx.y×blockDim.y+threadIdx.yc=blockIdx.x×blockDim.x+threadIdx.x
| 变量 | 含义 |
|---|---|
| blockIdx.y\text{blockIdx.y}blockIdx.y | 块的行索引 |
| threadIdx.y\text{threadIdx.y}threadIdx.y | 线程在块内的行索引 |
| blockDim.y\text{blockDim.y}blockDim.y | 每个块的线程行数 |
5. 层次结构与硬件的映射关系
CUDA 层次结构的设计直接反映了 NVIDIA GPU 的硬件结构，这是高效性能的关键：
Grid→\to→GPU：整个网格映射到整个 GPU。
Block→\to→SM：每个线程块被调度到一个 SM 上执行。多个块可以按时间片轮转的方式在同一个 SM 上执行，或者同时在多个 SM 上执行。
Thread→\to→Core：块内的线程最终映射到 SM 内部的 CUDA 核心。
- Warp：SM 实际上是以Warp（32 个连续线程）为单位进行指令调度和执行的。一个线程块会被分解成一个或多个 Warp。
性能考虑：分块大小
选择合适的线程块大小 (blockDim) 至关重要：
太小：无法充分利用 SM 的资源（如寄存器、共享内存），浪费硬件并行潜力。
太大：可能超出 SM 能够提供的资源限制（如最大线程数、共享内存大小），导致 Kernel 启动失败或运行效率降低。
通常，
blockDim选择 128、256 或 512，并且应为 32（Warp 大小）的倍数，以避免线程分化带来的性能损失。
