CUDA 编程指南 第5章：性能指南

5.1 指令性能

要处理一个 warp 的线程的指令，多处理器必须：
(1) 读取 warp 的每个线程的指令操作数
(2) 执行指令
(3) 写入 warp 的每个线程的结果

因此，有效的指令吞吐量取决于名义指令吞吐量以及内存延迟和带宽。它通过下列方式最大化：

• 最小化具有低吞吐量的指令的使用（参见 5.1.1）
• 最大化每种内存的可用内存带宽（参见 5.1.2）
• 允许线程调度器尽可能地将内存事务与数学计算重叠，这需要：
  • 程序具有高的算术密度（每个内存操作对应高的算术操作数）
  • 每个多处理器具有许多活动线程（参见 5.2）

5.1.1 指令吞吐量

5.1.1.1 算术指令

要执行 warp 的一个指令，多处理器需要：

注意：

• 整数除法和模数操作特别昂贵，应尽量替换为位操作：
  若 n 是 2 的幂，则 (i/n) 等价于 (i>>log2(n))，(i%n) 等价于 (i&(n-1))
• __[u]mul24 在将来架构中可能比 32 位整数乘法慢，建议提供两个内核版本
• 强烈建议使用单精度浮点和单精度数学函数（sinf、logf、expf 等）
• 浮点常量使用 f 后缀（如 3.141592653589793f, 1.0f, 0.5f）

5.1.1.2 控制流指令

任何流控制指令 (if, switch, do, for, while) 通过导致同一 warp 的线程分散
（按不同的执行路径执行），可以显著影响有效的指令吞吐量。

• 不同的执行路径必须序列化，增加 warp 执行的指令总数
• 当控制流取决于线程 ID 时，应最小化分散的 warp 数
• 控制条件仅取决于 (threadIdx / WSIZE) 时，warp 不会分散

编译器优化：

• 分支预测：将分支指令替换为谓词指令，每个指令与每线程条件代码相关
• 仅当分支指令数 <= 临界值（分散 warp 多时为 7，少时为 4）时才做此优化
• #pragma unroll 指令可控制循环展开（参见 4.2.5.2）

5.1.1.3 内存指令

内存指令包括从共享或全局内存中读取/写入的任何指令。

• 多处理器使用 4 个时钟周期执行 warp 的一个内存指令
• 访问全局内存时，还有 400 到 600 个时钟周期的内存延迟

如果等待全局内存访问时可以执行足够的独立算术指令，则大部分延迟可由
线程调度器隐藏。

5.1.1.4 同步指令

__syncthreads 使用 4 个时钟周期执行（如果没有任何线程需要等待）。

5.1.2 内存带宽

每个内存空间的有效带宽主要取决于内存访问模式。

典型的编程模式（最小化设备内存访问）：

1. 将设备内存中的数据加载到共享内存中
2. 与块的所有其他线程同步（__syncthreads）
3. 处理共享内存中的数据
4. 如有必要，重新同步
5. 将结果写回设备内存

5.1.2.1 全局内存

全局内存空间没有高速缓存，最重要的是按照正确的访问模式获得最大内存带宽。

(1) 字大小和对齐
设备能在单个指令中读取 32 位、64 位或 128 位字。
type 必须使得 sizeof(type) 等于 4、8 或 16，且地址必须对齐到 sizeof(type)。
内置类型（如 float2、float4）自动满足对齐要求。
结构体使用 align(8) 或 align(16) 强制对齐。

(2) 内存合并（Memory Coalescing）
半 warp 的每个线程同时访问的全局内存地址应排列为单个邻近的、对齐的内存访问。

合并条件：

• 半 warp 中线程号 N 应访问地址 HalfWarpBaseAddress + N
• HalfWarpBaseAddress 对齐为 16*sizeof(type) 字节

建议履行整个 warp 的合并要求（将来设备需要）。

带宽对比：

• 已合并 64 位访问 < 已合并 32 位访问（稍低）
• 已合并 128 位访问 << 已合并 32 位访问（低很多）
• 未合并 vs 已合并：32 位约差一个数量级，64 位约差 4 倍，128 位约差 2 倍

图 5-1: 已合并全局内存访问模式示例

• 左：已合并的 float 内存访问
• 右：已合并的 float 内存访问（分散 warp）

图 5-2: 未合并全局内存访问模式示例

• 左：非顺序的 float 内存访问
• 右：未对齐的开始地址

图 5-3: 未合并全局内存访问模式示例

• 左：不相邻的 float 内存访问
• 右：未合并的 float3 内存访问

常见访问模式：
(a) 线程 ID 为 tid 的线程访问 BaseAddress[tid]
→ type 需满足大小和对齐要求
(b) 索引为 (tx,ty) 的线程访问 2D 数组 BaseAddress[ty*width + tx]
→ 合并条件：块宽度是半 warp 大小的倍数，且 width 是 16 的倍数
→ 使用 cudaMallocPitch() / cuMemAllocPitch() 分配

5.1.2.2 常量内存

常量内存空间具有高速缓存。

• 半 warp 所有线程读取同一地址时，速度与寄存器一样快
• 成本随不同地址数线性扩展
• 建议整个 warp 所有线程读取同一地址（将来设备需要）

5.1.2.3 纹理内存

纹理内存空间具有高速缓存。

• 针对 2D 空间局部性优化
• 读取紧密相邻纹理地址的 warp 线程达到最佳性能
• 设计用于流水化具有恒定延迟的拾取
• 高速缓存命中降低 DRAM 带宽需求，但不降低拾取延迟

5.1.2.4 共享内存

因为位于芯片上，共享内存比本地和全局内存快得多。访问共享内存与访问寄存器
一样快（只要没有库冲突）。

内存库（Memory Bank）：

• 共享内存划分为相同大小的内存模块（库），共 16 个库
• 连续 32 位字分配到连续的库中
• 每两个时钟周期每个库有 32 位带宽
• warp 的共享内存请求分为第一半 warp 和第二半 warp 各一个请求
• 同一半 warp 内不同线程访问不同库 → 可同时服务
• 同一半 warp 内不同线程访问同一库 → 库冲突，需序列化

库冲突避免（计算能力 1.x，库数 m=16）：

• 线程使用跨度 s 访问 32 位字时，若 s 是奇数则无冲突（因为 m 是 2 的幂）
• d = gcd(m, s)，仅当半 warp 大小 <= m/d 时无冲突

图 5-4: 无库冲突的共享内存访问模式示例

• 左：跨度为一个 32 位字的线性寻址
• 右：随机排列

[图片 fig5-5.png] 图 5-5: 无库冲突的共享内存访问模式示例

• 跨度为三个 32 位字的线性寻址
  

[图片 fig5-6.png] 图 5-6: 有库冲突的共享内存访问模式示例

• 左：跨度为两个 32 位字的线性寻址 → 2 路库冲突
• 右：跨度为八个 32 位字的线性寻址 → 8 路库冲突

广播机制：

• 当半 warp 的多个线程从同一 32 位字内的地址读取时，可减少库冲突
• 通过多步服务：每步选择一个广播字，包含位于广播字内的所有地址
  以及由剩余地址指向的每个库的一个地址

[图片 fig5-7.png] 图 5-7: 有广播的共享内存读取访问模式示例

• 左：所有线程从同一 32 位字中的地址读取 → 无冲突
• 右：如果第一步选择库 5 中的字作为广播字，则无冲突；否则 2 路冲突

结构体访问：

• 结构体赋值编译为与成员数一样多的内存请求
• 成员使用较大跨度访问可避免冲突

5.1.2.5 寄存器

访问寄存器对于每条指令需要零个额外时钟周期。

• 读后写依赖关系导致的延迟：每多处理器至少 192 个活动线程可隐藏
• 寄存器内存库冲突：编译器自动优化，每块线程数为 64 的倍数时效果最佳
• 应用程序对寄存器库冲突无直接控制

5.2 每块的线程数

选择每块线程数时应最大化可用计算资源的利用率。

基本原则：
(1) 块数至少 >= 多处理器数目
(2) 每个多处理器允许两个或多个活动块（重叠等待与运行）
→ 块数 >= 多处理器数目 × 2
→ 每块共享内存量 <= 多处理器共享内存总量的一半
(3) 每块线程数为 warp 大小的倍数（避免未充满的 warp 浪费资源）
→ 更好是 64 的倍数
(4) 平衡：线程越多 → 时间分片越好，但每线程可用寄存器越少

对于计算能力 1.x 的设备：
每线程可用寄存器数 = R / (B × ceil(T, 32))
其中 R = 每多处理器寄存器总数，B = 活动块数，T = 每块线程数

推荐：

• 每块 64 个线程：最小，仅多个活动块时有意义
• 每块 192 或 256 个线程：比较好，通常允许足够的寄存器
• 每个网格至少 100 个块；1000 个块可扩展到几代

多处理器占有率 = 活动 warp 数 / 活动 warp 最大数
→ 最大化占有率：编译器最小化寄存器使用 + 程序员小心选择执行配置
→ CUDA SDK 提供电子表格辅助选择线程块大小

5.3 主机和设备之间的数据传送

• 设备↔设备内存带宽 >> 设备↔主机内存带宽
• 最小化主机和设备间数据传送：
  • 将更多代码从主机移动到设备（即使低并行计算）
  • 中间数据结构在设备内存中创建、操作、销毁
• 将许多小的传送分批为一次大的传送（减少每传送开销）
• 使用页面锁定内存获得较高带宽（参见 4.5.1.2）

5.4 纹理拾取 vs 全局或常量内存读取

通过纹理拾取读取设备内存的优点：

对于 CUDA 数组的纹理，硬件额外提供：

注意：纹理高速缓存与同一内核调用中的全局内存写不保持一致。
线程只能通过纹理安全读取先前内核调用或内存复制已更新的内存位置。

5.5 整体性能优化策略

性能优化围绕三个基本策略：

┌─────────────────────────────────────────────────────┐
│ (1) 最大化并行执行 │
│ - 暴露尽可能多的数据并行 │
│ - 同块线程同步：__syncthreads() + 共享内存 │
│ - 跨块线程通信：两个内核调用 + 全局内存 │
│ - 仔细选择执行配置（参见 5.2） │
│ - 使用流实现设备上并发 + 主机设备间并发 │
├─────────────────────────────────────────────────────┤
│ (2) 优化内存使用以获得最大内存带宽 │
│ - 最小化主机↔设备数据传送（参见 5.3） │
│ - 最大化共享内存使用，最小化设备↔全局内存传送 │
│ - 按最佳访问模式组织内存访问（参见 5.1.2.1-5.1.2.4）│
│ - 全局内存访问优化最重要（低带宽+高延迟） │
│ - 有时重新计算数据比传送数据更好 │
├─────────────────────────────────────────────────────┤
│ (3) 优化指令使用以获得最大指令吞吐量 │
│ - 最小化低吞吐量算术指令 │
│ - 用精度换速度（固有函数、单精度） │
│ - 注意控制流指令导致 warp 分散（参见 5.1.1.2） │
└─────────────────────────────────────────────────────┘

================================================================================
关键数字速查表