线程层次结构

本文聚焦层次结构、线程块设计、线程束概念、Occupancy 优化。

层次结构回顾

MUSA 采用三层线程层次结构：

网格
├── 线程块 0
│   ├── 线程 0
│   ├── 线程 1
│   └── ...
├── 线程块 1
│   ├── 线程 0
│   └── ...
└── 线程块 N

层级	说明	特点
线程	执行的最小单元	每个线程执行相同代码，处理不同数据
线程块	线程的集合	块内线程可同步、共享内存
网格	所有线程块的集合	一次内核函数启动的所有线程

为什么线程块之间不能同步？ 线程块可被调度到 GPU 上的任意 MP 执行，执行顺序不确定。如果允许跨线程块同步，程序将依赖特定的执行顺序，导致无法在不同 GPU 之间移植。

---

线程块设计

线程块大小选择

线程块大小	适用场景
128	寄存器需求高的内核函数
256	通用场景，推荐起点
512	计算密集型内核函数
1024	最大占用率场景

推荐原则：

• 线程块大小应为线程束大小的倍数（MP31: 32，MP21/MP22: 128）
• 常用值：128、256、512

网格大小计算

// 向上取整计算网格大小
int blockSize = 256;
int gridSize = (n + blockSize - 1) / blockSize;

线程块限制

参数	最大值
每个线程块的线程数	1024
线程块维度 (x)	1024
线程块维度 (y)	1024
线程块维度 (z)	64
共享内存大小	48 KB

网格限制

参数	最大值
网格维度 (x)	2^31 - 1
网格维度 (y)	65535
网格维度 (z)	65535

---

线程束概念

什么是线程束

线程束是 GPU 硬件的基本执行单元，包含固定数量的线程：

GPU 架构	线程束大小
MTT S5000 (MP31)	32 线程
MP21/MP22	128 线程

线程束执行

线程束执行一条指令：

┌─────────────────────────────────────┐
│ 1. 取指 (Fetch)                     │
│ 2. 译码 (Decode)                    │
│ 3. 执行 (Execute) - 所有线程同时    │
│ 4. 写回 (Write-back)                │
└─────────────────────────────────────┘

线程束与线程块的关系

• 一个线程块包含多个线程束
• 线程块大小 = 256 时，包含 8 个线程束（S5000，32 线程/线程束）
• 线程束是硬件调度的基本单位

当线程束内线程执行路径出现分支时，硬件会序列化执行各分支：

// 可能引起分支分化
if (threadIdx.x % 2 == 0) {
    // 路径 A：一半线程执行
    data[idx] = data[idx] * 2.0f;
} else {
    // 路径 B：另一半线程执行
    data[idx] = 0.0f;
}
// 总执行时间 = 路径 A 时间 + 路径 B 时间

更多详情，参见 执行模型。

---

线程同步

线程块内同步

__syncthreads();  // 等待线程块内所有线程到达同步点

使用场景：共享内存读写同步

__shared__ float shared_data[256];

// 步骤 1：写入共享内存
shared_data[threadIdx.x] = input[threadIdx.x];

// 步骤 2：同步，确保所有线程完成写入
__syncthreads();

// 步骤 3：读取其他线程的数据
float neighbor = shared_data[threadIdx.x + 1];

注意事项

• __syncthreads() 只能同步线程块内线程
• 不能跨线程块同步（线程块执行顺序不确定）
• 条件分支中的同步需确保所有线程都执行到同步点

---

占用率（Occupancy）

占用率概念

占用率是指 MP 上活跃的线程束数量与最大可能线程束数量的比率：

占用率 = 活跃线程束数 / 最大线程束数 × 100%

高占用率有助于隐藏内存延迟，提高吞吐量。

影响占用率的因素

因素	说明	影响
寄存器使用	每个线程使用的寄存器数	寄存器越多，活跃线程越少
共享内存	每个线程块使用的共享内存	共享内存越多，活跃线程块越少
线程块大小	每个线程块的线程数	需要与 MP 资源匹配

占用率计算

// 查询推荐线程块大小
int minGridSize, blockSize;
musaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, myKernel, 0, 0);

// 查询最大活跃线程块数
int numBlocks;
musaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, myKernel, blockSize, 0);

// 计算占用率
int maxWarpsPerSM = 32;  // 取决于设备
int activeWarps = numBlocks * (blockSize / 32);
float occupancy = (float)activeWarps / maxWarpsPerSM * 100;

优化策略

// 策略 1：调整线程块大小
// 尝试不同的线程块大小找到最优（128, 256, 512）
int blockSize = 256;  // 通用推荐值

// 策略 2：减少寄存器使用
// 避免大数组，减少临时变量

// 策略 3：减少共享内存使用
// 按需分配，避免浪费

---

最佳实践

1. 选择合适的线程块大小：通常 128-256 个线程
2. 利用线程束对齐：线程块大小应为 32 的倍数
3. 避免分支分化：将相同分支的线程放在一起
4. 使用线程索引：合理计算全局和局部索引
5. 边界检查：始终检查数组边界
6. 资源限制：注意共享内存和寄存器的限制

---

相关文档

• GPU 并行计算基础：GPU 架构基础，线程层次基本概念
• 内存层次结构：内存层次，寄存器/共享内存/全局内存
• 线程索引计算：线程索引，索引计算公式速查
• 执行模型：执行模型，SIMT 并行执行、Warp 调度