线程层次结构：Thread， Block， Grid

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 的硬件结构，这是高效性能的关键：

1. Grid→\to→GPU：整个网格映射到整个 GPU。

2. Block→\to→SM：每个线程块被调度到一个 SM 上执行。多个块可以按时间片轮转的方式在同一个 SM 上执行，或者同时在多个 SM 上执行。

3. Thread→\to→Core：块内的线程最终映射到 SM 内部的 CUDA 核心。

   • Warp：SM 实际上是以Warp（32 个连续线程）为单位进行指令调度和执行的。一个线程块会被分解成一个或多个 Warp。

性能考虑：分块大小

选择合适的线程块大小 (blockDim) 至关重要：

• 太小：无法充分利用 SM 的资源（如寄存器、共享内存），浪费硬件并行潜力。

• 太大：可能超出 SM 能够提供的资源限制（如最大线程数、共享内存大小），导致 Kernel 启动失败或运行效率降低。

通常，blockDim选择 128、256 或 512，并且应为 32（Warp 大小）的倍数，以避免线程分化带来的性能损失。