一、线程束（Warp）：GPU 的最小执行单元，避免 “分化” 是关键

线程束是 GPU 的最小调度和执行单元（固定 32 个线程，部分新架构如 Hopper 支持 64 线程束，但主流仍为 32），同一线程束内的线程必须执行完全相同的指令（SIMT 架构），但操作不同数据。线程束的组织优化核心是避免 “线程束分化”（Warp Divergence）。

1. 线程束分化：性能杀手

当线程束内的线程因分支语句（如if-else、switch）执行不同路径时，会触发 “分化”：

• 硬件会序列化执行不同分支（先执行if路径的线程，再执行else路径的线程，未执行的线程暂时闲置）；
• 分支覆盖率越低，效率损失越大（如 50% 线程走if，性能可能减半）。

例：

// 低效：线程束内分化（假设threadIdx.x为0~31）
if (threadIdx.x % 2 == 0) {
    result[threadIdx.x] = a[threadIdx.x] + b[threadIdx.x];
} else {
    result[threadIdx.x] = a[threadIdx.x] * b[threadIdx.x];
}

上述代码中，线程束内一半线程执行加法，一半执行乘法，需分两次执行，效率降低 50%。

2. 优化技巧：消除或减少分化

技巧 1：用算术操作替代分支
将条件判断转化为数学运算（如(condition) ? x : y可改为condition * x + (1 - condition) * y），避免分支。
优化后代码：

int condition = (threadIdx.x % 2 == 0) ? 1 : 0;
result[threadIdx.x] = condition * (a[threadIdx.x] + b[threadIdx.x]) + 
                     (1 - condition) * (a[threadIdx.x] * b[threadIdx.x]);

技巧 2：确保分支在 “线程束粒度” 上一致
若必须分支，让整个线程束执行同一分支（而非束内部分线程）。例如，按线程束 ID（threadIdx.x / 32）划分分支：

// 高效：同一线程束内无分化（所有32线程执行同一分支）
if ((threadIdx.x / 32) % 2 == 0) { // 按线程束ID分支
    // 整个线程束执行此路径
} else {
    // 整个线程束执行此路径
}

技巧 3：重构算法，将分支逻辑提到线程块外
对数据预处理，让线程块内的线程执行相同逻辑（如按数据特征划分任务，避免块内分支）。

二、线程块（Block）：资源分配与并发的基本单位，大小与资源匹配是核心

线程块由 1 个或多个线程束组成（如 256 线程的块含 8 个 32 线程束），是 SM（流式多处理器）上资源分配的基本单位（共享内存、寄存器等按块分配）。线程块的组织需解决两个核心问题：块大小如何选择，以及如何最大化 SM 上的块驻留数。

1. 线程块大小：128~256 线程为黄金区间，需是 32 的倍数

线程块大小（blockDim）直接影响 SM 的利用率和资源消耗，需满足：

• 必须是 32 的倍数：避免 “部分填充的线程束”（如 255 线程的块会产生 8 个完整束 + 1 个仅 31 线程的束，浪费 1 个线程的计算资源）；
• 不宜过小（如 < 64 线程）：SM 上的计算单元（如 CUDA 核心）无法被充分利用（SM 通常有 64 + 核心，小线程块会导致大量核心空闲）；
• 不宜过大（如 > 1024 线程）：可能因共享内存或寄存器不足，导致 SM 无法同时驻留多个块（块驻留数 = SM 资源总量 / 单块资源需求，块过大则驻留数下降，影响并发切换）。

推荐值：128、256、512 线程 / 块（需结合具体代码测试，例如计算密集型任务可用较大块，内存密集型任务可用较小块）。

2. 最大化 SM 上的 “块驻留数”

块驻留数（一个 SM 同时容纳的线程块数量）决定了 SM 的并发能力 —— 驻留块越多，调度器可切换的线程束越多，越能隐藏内存延迟。驻留数受以下资源限制（取最小值）：

• 寄存器限制：SM总寄存器数 / 单块寄存器消耗（单块寄存器消耗 = 块内线程数 × 单线程寄存器数）；
• 共享内存限制：SM共享内存总量 / 单块共享内存消耗；
• 线程数限制：SM最大线程数（如2048） / 块内线程数；
• 块数上限：SM 硬件支持的最大驻留块数（如 32）。

优化技巧：

• 控制单块资源消耗：减少共享内存申请（如用寄存器存储临时数据）、避免单线程使用过多寄存器（如减少大数组），提升驻留数。
  例：若 SM 有 65536 寄存器，单线程用 32 寄存器，256 线程块的单块消耗为256×32=8192寄存器，则寄存器限制的驻留数为65536 / 8192 = 8；若单线程寄存器降至 16，则驻留数提升至 16。
• 平衡块大小与驻留数：例如，128 线程块的驻留数可能是 256 线程块的 2 倍（若资源允许），更适合内存延迟高的场景（通过更多切换隐藏延迟）。

3. 线程块内的线程映射：匹配数据布局，优化内存访问

线程块内的线程索引（threadIdx.x/y/z）需与数据的内存布局对齐，确保全局内存访问合并（连续地址）和共享内存无冲突。

例：矩阵转置中的线程映射
矩阵在内存中按 “行优先” 存储，转置时需按列访问，易导致非连续全局内存访问。通过线程块内的二维映射（threadIdx.x对应行，threadIdx.y对应列），结合共享内存分块，可优化访问模式：

__global__ void transpose(float *out, float *in, int rows, int cols) {
    __shared__ float tile[32][33]; // 32×33共享内存（填充1列避免冲突）
    
    // 线程映射：(x,y)对应输入矩阵的(row, col)
    int row = blockIdx.y * 32 + threadIdx.y;
    int col = blockIdx.x * 32 + threadIdx.x;
    
    // 行访问输入（连续，合并），存入共享内存
    tile[threadIdx.y][threadIdx.x] = in[row * cols + col];
    __syncthreads();
    
    // 列访问共享内存（无冲突），行访问输出（连续，合并）
    out[col * rows + row] = tile[threadIdx.x][threadIdx.y];
}

此处线程块大小为32×32，线程索引与矩阵分块对齐，既保证全局内存合并，又通过共享内存 Padding 避免列访问冲突。

三、网格（Grid）：任务划分的顶层单位，需覆盖所有 SM 并匹配问题维度

网格由多个线程块组成（gridDim），其设计需确保所有 SM 被充分利用，并与问题的维度（如 1D 向量、2D 图像、3D 体积数据）匹配。

1. 网格大小：至少为 SM 数量的 2~4 倍，避免 SM 空闲

GPU 的 SM 数量通常为数十到数百（如 RTX 3090 有 82 个 SM，A100 有 108 个）。若网格中的线程块数远少于 SM 数量（如 82 个 SM 仅分配 40 个块），会导致部分 SM 空闲，浪费计算资源。

推荐值：网格中的块数 = SM 数量 × 2~4（例如 82 个 SM，网格块数取 200~300），确保 SM 可持续调度新块（即使部分块已完成）。

2. 网格维度：匹配问题的空间维度，简化索引计算

网格可声明为 1D、2D 或 3D（gridDim.x/y/z），维度选择应与问题的自然维度一致，减少索引计算的复杂度和错误：

• 1D 网格：适合向量、序列等 1D 数据（如数组求和）；
• 2D 网格：适合图像、矩阵等 2D 数据（如卷积、转置）；
• 3D 网格：适合 3D 体积数据（如 CT/MRI 图像重建）。

例：2D 网格处理图像

// 图像尺寸为W×H，线程块大小32×32
dim3 block(32, 32);
dim3 grid((W + 31) / 32, (H + 31) / 32); // 向上取整，覆盖所有像素
kernel<<<grid, block>>>(image, W, H);

网格维度与图像维度一致，线程索引(blockIdx.x*blockDim.x + threadIdx.x, blockIdx.y*blockDim.y + threadIdx.y)可直接对应像素坐标，计算直观且不易出错。

3. 避免跨块同步，用 “分阶段计算” 替代

线程块间无法直接同步（__syncthreads()仅作用于块内），跨块同步需通过全局内存 + 原子操作实现，开销极大（数百到数千时钟周期）。

优化技巧：将任务拆分为 “块内独立计算”+“全局汇总” 的分阶段流程，避免跨块依赖。例如，全局求和可分为：

1. 每个线程块计算部分和（块内同步）；
2. 单个线程块汇总所有部分和（仅一次全局同步）。

四、特殊场景的线程组织优化

1. 计算密集型任务（如矩阵乘法）

• 线程块大小可适当增大（如 256~512 线程），利用更多寄存器存储中间结果，减少内存访问；
• 采用 “2D 线程块 + 共享内存分块”（如 16×16 块），最大化数据复用（每个元素被多次访问）。

2. 内存密集型任务（如数据拷贝、滤波）

• 线程块大小可适当减小（如 128 线程），增加 SM 上的块驻留数，通过更多线程束切换隐藏内存延迟；
• 确保线程索引与内存地址严格对齐（如按 128 字节边界访问），提升全局内存合并效率。

3. 利用硬件特性的线程组织（如 Tensor Core）

• Tensor Core（用于混合精度矩阵乘法）要求线程块按特定尺寸组织（如 128×128 线程块，对应 16×16×16 的矩阵分块），需严格匹配硬件计算单元的并行粒度；
• 通过wmma API（CUDA 的矩阵乘累加接口）强制线程组织与 Tensor Core 对齐，否则无法利用硬件加速。

五、总结：线程组织的核心原则

1. 线程束层面：避免分化，确保同一束内线程执行相同指令，用算术操作替代分支；
2. 线程块层面：大小选择 128~256（32 的倍数），控制资源消耗以提升 SM 驻留数，线程索引匹配数据布局；
3. 网格层面：块数为 SM 数量的 2~4 倍，维度匹配问题空间，避免跨块同步；
4. 硬件匹配：根据任务类型（计算 / 内存密集）和特殊硬件（Tensor Core）调整组织方式，最大化硬件利用率。

线程组织的本质是 “将问题的并行性映射到 GPU 的硬件并行结构”—— 通过合理划分任务、匹配资源、减少低效操作，让计算单元和内存带宽被充分利用，最终实现性能最大化。