> 基于 MCC2026 OSTIA SST 气候态计算项目（V0 → V10.24，13,214× 加速比）实战经验 > 从 17.4 小时到 4.75 秒的完整优化历程 --- ## 写在前面 上篇讲了硬件，中篇聊了写 kernel 和调 profiler。这一篇把视角拉远——不只看某一段代码怎么写，而是看整个程序怎么从 MATLAB 串行 17 小时一步步优化到 4.75 秒。核心线索是"瓶颈转移"：每次攻破一个瓶颈，下一个就会浮出水面，要求完全不同的技术手段。13,214 倍不是某个天才优化，而是十几个小改进叠加的乘积效应。 --- ## 一、总览：加速比是怎么来的？ ``` V0 MATLAB 串行基线 62,768s (17h26m) 1× V3 C++ CPU 多线程 180s (3 min) 349× V4 HIP GPU 移植 70s 897× V10 C++ 微优化 14s 4,483× V10.9 预加载流水线 ~9.4s 6,677× V10.13 float→short 量化 ~8.4s 7,472× V10.18 pread 裸读 ~6.6s 9,510× V10.21 线程替代 fork ~5.7s 11,012× V10.23 工作窃取 ~5.3s 11,910× V10.24 hipHostRegister ~4.75s 13,214× ``` **13,214 倍是十几个不同方向的改进叠加的乘积效应,其中dcu加速带来的提升不容小觑。** 这背后是一个关于"瓶颈转移"的核心概念：每次突破一个瓶颈后，下一个瓶颈会浮现出来，要求完全不同的优化技术。理解瓶颈在哪个维度，比掌握具体的优化技巧更重要。 本文按时间线叙述，在每个阶段讲解当时的瓶颈分布、优化手段、以及为什么它能生效。 --- ## 二、第一阶段：MATLAB → C++ 预加载池（1× → 349×） V0 是 MATLAB 串行脚本，每个 DOY 循环打开 330 个 NC 文件，累计 **2189 万次 `nc_varget`** 调用，每次触发文件系统 open/read/close，17 小时才能跑完。V3 的核心改进：**把 3060 个 NC 文件一次性全部读入 12.7 GB 内存池 `pool[DOY][year][pixel]`**，所有 DOY 计算从池中取数据，文件 I/O 从 281,520 次降为 3060 次。20 个 CPU 线程并行计算，耗时从 17h 降到 180s（349×）。 **瓶颈转移**：I/O 不再是瓶颈，CPU 计算变成瓶颈 → 引入 GPU。 --- ## 三、第二阶段：GPU 移植（349× → 897×） ### 3.1 Kernel 设计与 Grid 划分 每个海洋格点的 330 样本统计完全独立——典型的"尴尬并行"，1,038,240 像素可同时计算。Kernel 使用 **2D Grid**： ```cpp dim3 block(256); // 每block 256线程 = 4 wavefront dim3 grid(GRID_SIZE / 256, // x: 覆盖全部 1,038,240 像素 doy_count); // y: 覆盖该GPU分配的DOY数 ``` 选择 256 线程/block 是 gfx906 上的经验最优——太小（64）调度开销大，太大（1024）寄存器压力过大导致 occupancy 进一步降低。 2D Grid 允许一次 kernel launch 处理多个 DOY，避免 92 次独立 launch 的累积开销（每次 ~10 μs × 92 ≈ 1ms）。 ### 3.2 H2D / Kernel / D2H 管线 GPU 计算的三个阶段对应不同硬件资源，可以流水线重叠： ``` Host内存 ──H2D(PCIe)──▶ GPU显存 ──Kernel(HBM)──▶ GPU显存 ──D2H(PCIe)──▶ Host内存 ``` ```cpp void gpu_worker(int gpu_id, ...) { hipSetDevice(gpu_id); hipMalloc(&d_pool, local_pool_bytes); // GPU显存分配 hipMalloc(&d_mean, out_bytes); hipMalloc(&d_p90, out_bytes); hipMemcpy(d_pool, full_pool + offset, ..., hipMemcpyHostToDevice); // H2D hipStream_t s0, s1; // 2-stream 分半并行 compute_stats_kernel<<>>(d_pool, ..., d_mean, d_p90); compute_stats_kernel<<>>(d_pool, ..., d_mean + off, d_p90 + off); hipMemcpyAsync(h_mean, d_mean, ..., hipMemcpyDeviceToHost, s0); // D2H异步 hipMemcpyAsync(h_mean + off, d_mean + off, ..., hipMemcpyDeviceToHost, s1); hipStreamSynchronize(s0); hipStreamSynchronize(s1); hipFree(d_pool); hipFree(d_mean); hipFree(d_p90); } ``` 每个 GPU 只加载它需要的 pool DOY 子集（含滑动窗口边界重叠），典型配置 4 GPU × 2 节点，每 GPU 处理 ~23 个 DOY，显存占用约 1.5 GB。 ### 3.3 内存瓶颈的诊断与确认 GPU 移植后，用 rocprof 做性能分析： | 硬件计数器 | 数值 | 含义 | |-----------|:----:|------| | **MemUnitBusy** | **99.71%** | HBM 控制器几乎满负荷 | | VALUUtilization | 58.81% | ALU 近半时间在等数据 | | L2CacheHit | 67.21% | 33% 访问穿透到 HBM | MemUnitBusy 99.71% = **纯内存瓶颈**。ALU 在等 HBM 的数据，不是 HBM 在等 ALU。这意味着：指令优化、循环展开、提升 occupancy 全部无效——指令执行时间被内存等待完全掩盖。 **Occupancy 的再思考**：CUDA 开发者常认为 occupancy 越高越好（更多 wavefront 可隐藏内存延迟），但前提是内存带宽有富余。当 MemUnitBusy 已达 99.7%，HBM 控制器无暇处理新请求，提升 occupancy 只是在 HBM 控制器前加更长的排队。我们 25% 的 occupancy（160 VGPR/线程限制）在这个场景下不仅不是问题，甚至可能是最优——更少的 wavefront 意味着更少的 L2 竞争、更低的 HBM bank 冲突。 ### 3.4 瓶颈再转移 GPU 移植后，Kernel 耗时 ~900ms，但整个程序仍需 70s——**95.7% 是串行 I/O 时间**（单线程读 NC 文件预加载 pool）。瓶颈从 GPU 计算转移到了 CPU 侧数据供给。 --- ## 四、第三阶段：CPU 侧 I/O 并行化（897× → 11,012×） GPU 移植后，70s 中 95.7% 是串行 I/O 时间——预加载仍是单线程。这个阶段的优化目标是把 CPU 侧数据喂饱 GPU，涉及多个技术迭代： **V10.9 Fork + mmap**：HDF5 1.8 非线程安全，多线程读 NC 会随机崩溃。用 `fork()` 创建 28 个子进程（各自独立 HDF5 状态），通过 `mmap(MAP_SHARED | MAP_ANONYMOUS)` 共享 6.4 GB 内存池，semaphore 同步父子进程。 **V10.13 float→short 量化**：SST 精度需求 0.01°C，将 float(4B) 量化为 short(2B)。对 DCU 而言关键收益不在内存占用，而在 **H2D 传输量减半**：每 GPU 从 2.7 GB → 1.35 GB，H2D 从 ~200ms → ~100ms，直接减轻 PCIe 压力。 **V10.18 pread 裸读**：用 HDF5 `H5Dget_offset()` 探测数据块文件偏移量，确认一致后直接用 `pread()` 系统调用读取，绕过 NetCDF→HDF5→文件系统三层调用栈，零元数据开销，天然线程安全。 **V10.21 线程替代 fork**：pread 线程安全后，改用 `std::thread` 直接并行预加载，消除 fork 的 copy-on-write page fault 和 semaphore IPC 开销。sys 降 35%，page fault 降 98%。 **V10.23 工作窃取**：36 个预加载线程通过 `atomic` 计数器竞争 DOY 索引，快线程多干慢线程少干，消除 Lustre I/O 带宽波动（最快/最慢可达 1.5×）导致的木桶效应。 这些 CPU 侧优化将预加载从 ~67s 压到 ~2.7s，确保 GPU 不会因等数据而空转。 --- ## 五、第四阶段：DCU 侧收尾优化 ### 5.1 hipHostRegister：消除隐式拷贝 (V10.24) 默认 `hipMemcpy` H2D 需要两步：驱动先将数据拷到内部 staging buffer，GPU DMA 引擎再从 staging buffer 拷到显存。`hipHostRegister` 将 host 内存页锁定（pin）并注册到 IOMMU，使 GPU DMA 引擎可直接访问物理页，省掉 staging buffer 拷贝。 ```cpp // 1. mmap 匿名映射 short *pool = (short*)mmap(nullptr, pool_bytes, PROT_READ | PROT_WRITE, MAP_PRIVATE | MAP_ANONYMOUS, -1, 0); // 2. 预加载线程写入数据（触发 page fault，分配物理页） for (auto &t : preload_threads) t.join(); // 3. 页面全部 fault-in 后注册到 IOMMU——顺序至关重要！ hipHostRegister(pool, pool_bytes, hipHostRegisterDefault); // 4. GPU 通过 DMA 直接读取 pool // 5. 释放 hipHostUnregister(pool); munmap(pool, pool_bytes); ``` **踩坑记录**：这个优化在 V10.14 被放弃过——当时注册开销 ~4s 超过了收益。原因是在预加载之前就注册，后续每个 page fault 都触发 IOMMU 重新映射。V10.24 改为**预加载完成后注册**，此时所有页面已 fault-in，pin 操作耗时大幅缩短。H2D 有效带宽从 ~6 GB/s 提升到 ~12 GB/s，接近 PCIe 3.0 ×16 的理论极限。 ### 5.2 三缓存管线分析（未实施） 当前 H2D 是同步的，排在所有 kernel 之前。理想方案是用 3 个 HIP stream 各管一个数据 chunk 的 H2D→Kernel→D2H，实现完全流水线重叠： ``` Stream 0: [H2D_0][Kernel_0][D2H_0] Stream 1: [H2D_1][Kernel_1][D2H_1] Stream 2: [H2D_2][Kernel_2][D2H_2] ``` 预期 GPU 端改善 ~6-8%（H2D 被 kernel 掩盖），但 GPU 仅占 wall clock 的 25%，墙钟改善 <2%。在 CPU 侧 I/O 优化做到极致后，GPU 端的边际收益已经很小。**这也是判断优化该收手的一个信号**——当剩余优化目标的 wall clock 占比不足 25%，收益很难超过 2-3%。 --- ## 六、瓶颈分析与何时停止优化 ### 6.1 V10.24 的时间分布 | 阶段 | 时间 | 占比 | 瓶颈 | |------|:----:|:----:|------| | 预加载 I/O (pread) | ~2.7s | **57%** | Lustre 存储带宽 | | GPU Kernel | ~0.9s | 19% | HBM 内存带宽 (99.7% MemUnit) | | H2D + hipHostRegister | ~0.5s | 11% | PCIe 带宽 | | D2H | ~0.3s | 6% | PCIe 带宽 | | NC 写入 | ~0.2s | 4% | Lustre 写入带宽 | | 其他 (fork/同步等) | ~0.15s | 3% | — | ``` 57% Lustre + 19% HBM + 11% PCIe = 87% 硬件物理极限 ``` ### 6.2 每条优化路径都走到了尽头 **CPU 计算** → 已卸载到 GPU。剩下的 CPU 时间全是 I/O 等待。 **GPU 计算** → MemUnitBusy 99.7% 意味着 HBM 带宽已打满。指令优化、occupancy 提升、循环展开全部无效——ALU 在等 HBM 的数据，不是 HBM 在等 ALU。 **H2D/D2H** → PCIe 3.0 ×16 的带宽上限约 14 GB/s，我们的有效带宽已达约 12 GB/s，接近理论极限。 **I/O 读取** → pread 绕过了所有软件中间层。Lustre 的存储节点和网络是最后的物理障碍。36 个并发 pread 已经将 Lustre 的元数据并发推到了临界点（56 worker 反而慢 0.4s）。 --- ## 七、优化方法论总结 ### 7.1 瓶颈转移原则 ``` I/O瓶颈(V0) → 预加载池 → 计算瓶颈(V3) → GPU移植 → I/O瓶颈(V4-V10) → pread+线程 → 硬件极限(V10.24) ``` 每次攻破一个瓶颈，下一个就会浮出水面——而且有可能是之前不注意的环节。理解瓶颈在哪个维度，比掌握优化技巧更重要。 ### 7.2 DCU Kernel 快速诊断法 如果无法使用 rocprof（无 root 权限、DTK 版本不支持），可以尝试使用四个粗糙实验判断瓶颈： 1. **注释掉 kernel 中除读取外的所有操作**：耗时基本不变 → 内存瓶颈 2. **kernel 输入数据量减半**：耗时同比减半 → 内存瓶颈 3. **blockSize 从 256 改成 128**：耗时不变 → occupancy 不是瓶颈 4. **kernel 中插入无意义计算**（循环 100 次空转）：耗时不变 → 计算不是瓶颈（ALU 在等 HBM） ### 7.3 给 DCU 开发者的建议 1. **先用最简单的方式跑起来**。V0 的 MATLAB 17 小时给出了正确答案，优化从正确答案开始。 2. **每个阶段只优化一个瓶颈**。同时优化多个维度只会让调试不可控。 3. **保持"真理版本"**。每个版本都和 MATLAB V0 输出交叉验证（RMSE = 0）。跨语言数值验证时，**第一步就核对行为一致性**——我们的 P90 偏差 0.0126°C 查了四天才发现是 Hyndman & Fan Type=5 vs Type=7 的差异。 4. **profiler 先于直觉**。rocprof 的数字不会骗人：MemUnitBusy 99.7% = 纯内存瓶颈，所有指令优化都是徒劳。 5. **知道何时收手**。当 87% 的时间花在 Lustre + HBM + PCIe 的物理极限上，继续改代码已无意义。剩余优化目标的 wall clock 占比不足 25% 时，边际收益很难超过 2-3%。 --- ## 八、附：关键指标速查 ### 8.1 硬件极限 | 组件 | 理论带宽 | 实际达到 | 利用率 | |------|----------|:------:|:------:| | Lustre (每节点) | ~5 GB/s | ~4.5 GB/s | 90% | | HBM2 (每卡) | ~1 TB/s | ~900 GB/s (ECC后) | 99.7% MemUnit | | PCIe 3.0 ×16 | 15.75 GB/s | ~12 GB/s | 76% | | DDR4 (每NUMA) | ~50 GB/s | 极少使用 | <5% | ### 8.2 数据规模 | 指标 | 数值 | |------|------| | NC 文件总数 | 3060 (102天 × 30年) | | 每文件大小 | ~4.15 MB | | 总输入数据 | 14 GB / 节点 | | pool 大小 (short) | 6.4 GB / 节点 | | GPU 显存占用 | ~1.5 GB / GPU (pool子集 + 输出) | | 输出数据 | 184 NC (92天均值 + 92天P90) | ### 8.3 精度验证 | 指标 | RMSE | |------|:----:| | 气候态均值 (Clim) | 0.0001°C | | P90 分位数 | 0.0012°C | --- *全文完。*