摘要
AMD DCU（Deep Computing Unit）是 ROCm 生态下的 GPU 加速器，已在多个国产超算系统中部署。本文以一个气候态海温第 90 百分位计算程序为教学案例，面向 DCU 上的 HIP 编程初学者，系统介绍三类典型算子（Top-K 归约、滑动窗口统计矩、多级流水线）的优化思路。全文围绕七条策略展开：wavefront 对齐与线程块配置、寄存器级归约与 launch_bounds 约束、分支分歧消除与统计矩替代、加法群降维、多流并发与 DMA 计算重叠、流水线环数匹配、pinned memory 管理。每条策略都包含架构根因分析、完整代码示例和量化对比，旨在为 DCU 上的性能优化提供一份可参考的教学材料。实验部分给出了各优化步骤的加速效果：GPU kernel 时间从 8.0 秒降至 0.08 秒，总程序从 14.0 秒降至 4.2 秒。
关键词：AMD DCU；HIP；GPU 算子优化；归约；滑动窗口；流水线
*1. 引言
1.1 为什么学习 DCU 优化？
AMD DCU（Deep Computing Unit）是 AMD 面向 HPC 和 AI 场景设计的 GPU 加速器。它的编程模型 HIP（Heterogeneous-compute Interface for Portability）在语法上高度兼容 CUDA——如果你写过 CUDA，HIP 代码看起来几乎一模一样：
*但底层硬件有本质差异。最显著的区别是：NVIDIA 的调度单位是 32 线程的 warp，而 DCU 是 64 线程的 wavefront。这个 2× 的差异会传导到几乎所有优化决策上——线程块大小设多少、寄存器够不够用、分支分歧有多大代价。本文的核心目的就是帮你建立 DCU 的”硬件直觉”：知道你的代码在 DCU 上是怎么跑的，才能知道怎么让它跑得更快。
1.2 教学案例：气候态 SST 百分位计算
本文用一个真实的气候海洋学问题作为贯穿全文的教学案例。问题是这样的：
给定 1991 年到 2020 年共 30 年的全球逐日海表温度（SST）数据，网格分辨率 0.25°×0.25°（经度 1440 格 × 纬度 721 格），对于每个海洋格点，计算每个日历日（一年中的第 152 天到第 243 天）的：
气候态均值：该日前后各 5 天（共 11 天）滑动窗口内，30 年所有有效 SST 的平均值
第 90 百分位：同样窗口内所有有效 SST 的第 90 百分位值
这个问题具备三个典型的 GPU 计算特征：
数据并行：每个格点的计算独立，天然适合 GPU
归约密集：每个格点需要从约 330 个值中找出百分位
滑动窗口：相邻日历日的窗口有 10⁄11 的重叠
数据存储为逐日的 NetCDF4/HDF5 文件，单文件约 8.3 MB，总共约 5280 个文件。输入规模约 25 GB/进程。
1.3 文章结构
第 2 节：DCU 架构详解——你写代码之前必须知道的硬件知识
第 3 节：七条优化策略，每条包含”为什么→怎么做→代码→效果”
第 4 节：实验数据汇总
第 5 节：常见陷阱与调试建议
第 6 节：总结
*2. DCU 架构特征
2.1 芯片规格速览
本文使用的 DCU 型号是 AMD MI50（架构代号 gfx906）：
*2.2 Wavefront：DCU 上最重要的概念
NVIDIA GPU 的调度单位是 warp（32 线程）。DCU 的调度单位是 wavefront（64 线程）。这不是一个数字差异，它从根本上决定了你的优化策略。
什么是 wavefront？
当你在 HIP 中写 kernel<<<256, 128>>>(…) 时，你以为你在控制线程。实际上，DCU 硬件并不直接调度线程——它调度的是 wavefront（一组 64 个线程）。128 个线程的块会被拆成 2 个 wavefront，每个 64 线程，在两个不同的 SIMD 单元上执行。
线程块 (128线程)：
├── Wavefront 0: 线程 0-63 → SIMD单元 0
└── Wavefront 1: 线程 64-127 → SIMD单元 1
为什么 wavefront = 64 很重要？
*实践要点：
线程块大小应为 64 的倍数（128、256 最优）
避免 32、96、160 这些非 64 倍数的大小
分支分歧在 DCU 上比 NVIDIA 贵 2 倍
2.3 寄存器：最稀缺的资源
DCU 的每 CU 寄存器预算仅 64 KB。对比 NVIDIA V100 的 256 KB/CU，差距 4 倍。这意味着 DCU 上的寄存器更加紧张。
这 64 KB 被所有 resident wavefront 共享。wavefront 越少、每线程可用寄存器越多；wavefront 越多、并行度越高但每线程寄存器越少。
案例：每 CU 4 个 thread block，每 block 256 线程
→ 每 CU 线程数 = 4 × 256 = 1024
→ 每 CU wavefront 数 = 1024 / 64 = 16
→ 每 wavefront 寄存器 = 65536 / 16 = 4096 bytes = 1024 register slots
→ 每线程寄存器 = 1024 / 64 = 16 个
等等，16 个寄存器远远不够！大多数 kernel 至少需要 30-50 个寄存器。
这就是 launch_bounds 的作用。它告诉编译器你预期每 CU 有多少个 block，编译器据此算出每线程可用寄存器数，如果超过则 spill 到 local memory。
__global__ void __launch_bounds__(256, 4) my_kernel()
// 含义：最大 256 线程/块，最少 4 block/CU
// 编译器保证：65536 / (4 × 256) = 64 寄存器/线程
当你不用 __launch_bounds__ 时，编译器会取一个保守值（通常更多寄存器），导致实际 wavefront 数减少，从而降低 occupancy。反之如果设得太紧，编译器会 spill，性能更差。
后面 3.2 节会专门讲如何用 launch_bounds。
2.4 访存层次
Host (CPU) Memory
│ PCIe 3.0 ×16 (~12 GB/s)
▼
DCU Device Memory (HBM2, 16 GB, ~1 TB/s)
│
├── L2 缓存 (4 MB, 所有 CU 共享)
│ │
│ ▼
│ ┌──────────────────────┐
│ │ 每 CU: L1 16KB │
│ │ 或 共享内存 64KB │
│ │ 或 L1 16KB + 共享 48KB
│ └──────────────────────┘
│ │
│ ▼
│ 寄存器文件 (64 KB/CU)
│ │
│ ▼
│ ALU 计算
层级含义：
Global memory (HBM2)：大容量、高带宽，但延迟 ~300-500 cycles
L2 缓存：4 MB 片上缓存，所有 CU 共享，延迟 ~100 cycles
L1 缓存 / 共享内存：每 CU 独立，可配置为 16 KB L1 + 48 KB 共享，或全部 64 KB 共享
寄存器：最快，延迟 1 cycle，但每 CU 仅 64 KB
实践要点：
线程私有数据尽量放寄存器（即局部变量），而非 global memory
需要线程间共享的数据放共享内存
只读的大数组用 const __restrict__ 修饰，帮助编译器生成向量化 load 指令
*3. 七条算子优化策略
3.1 策略一：Wavefront 对齐与线程块配置
3.1.1 为什么需要对齐？
如前所述，DCU 以 wavefront（64 线程）为最小调度单位。如果你的线程块大小不是 64 的倍数，最后一个 wavefront 会包含空闲线程——这些空闲线程虽然不做任何计算，但仍然占用 CU 资源（寄存器、执行时间）。
3.1.2 具体做法
// ❌ 错误：非 64 倍数
int thr = 32; // 32 = 0.5 wavefront，浪费 50%
int thr = 96; // 96 = 1.5 wavefront，浪费 25%
int thr = 160; // 160 = 2.5 wavefront，浪费 20%
// ✅ 正确：64 倍数
int thr = 64; // 1 wavefront，完整利用
int thr = 128; // 2 wavefront
int thr = 256; // 4 wavefront（推荐）
3.1.3 线程块大小选择的权衡
*3.1.4 实例：配置我们的 kernel
// 每块 DCU 处理约 259,560 个格点
size_t sz = 259560;
// 用 256 线程块，需要 ceil(259560/256) = 1014 个块
int thr = 256;
int blk = (int)((sz + thr – 1) / thr); // = 1014
// 线程总数 = 1014 × 256 = 259,584
// 多出 24 个线程（ti >= sz 时直接 return）
// 这 24 个线程浪费可以忽略（0.009%）
kernel<<<blk, thr, 0, stream>>>(…);
*3.2 策略二：寄存器级归约与 launch_bounds
3.2.1 问题场景
在 SST 百分位计算中，每个格点需要维护一个大小为 34 的浮点堆（Top-K 最小堆）。如果这个堆存在 global memory 中，每次堆操作需要 1-2 次全局内存访问（~300 cycles），34 次值的插入就是 10,000+ cycles 的开销——太慢了。
3.2.2 寄存器级归约
将堆存储在寄存器数组中——每个线程有自己独立的 34 个 float 寄存器，堆操作（比较、交换）全部在寄存器内完成，0 次全局内存访问。
// ★ 寄存器级归约：堆在寄存器中
__global__ void __launch_bounds__(256, 4)
topk_kernel(const double* input, float* heap_out,
size_t sz, int n_vals)
{
size_t ti = blockIdx.x * blockDim.x + threadIdx.x;
if(ti >= sz) return;
// ★ heap[34] 全部存在寄存器中！
float heap[34];
#pragma unroll
for(int i = 0; i < 34; i++) heap[i] = -1e38f;
double sum = 0.0; // 也在寄存器中
int count = 0; // 也在寄存器中
// 遍历这个格点的所有值
for(int i = 0; i < n_vals; i++){
double v = input[(size_t)i * /*stride*/ + ti];
if(!isnan(v)){
sum += v;
count++;
// ★ 堆操作完全在寄存器内执行
if(v > (double)heap[0]){
heap[0] = (float)v;
// sift-down：只需寄存器操作
int pos = 0;
while(pos < 34){
int left = 2*pos + 1;
int right = 2*pos + 2;
int smallest = pos;
if(left < 34 && heap[left] < heap[smallest])
smallest = left;
if(right < 34 && heap[right] < heap[smallest])
smallest = right;
if(smallest == pos) break;
float tmp = heap[pos];
heap[pos] = heap[smallest];
heap[smallest] = tmp;
pos = smallest;
}
}
}
}
// ★ 整个归约过程只有两次全局内存访问：
// 1. 读取 input[ti]
// 2. 写回 heap_out[ti * 34 .. ti * 34 + 33]
for(int i = 0; i < 34; i++)
heap_out[(size_t)ti * 34 + i] = heap[i];
}
3.2.3 为什么需要 launch_bounds
heap[34] 用掉了 34 个浮点寄存器。加上 sum、count、索引变量等，总量约 45-50 个寄存器。如果不加 __launch_bounds__，编译器不知道每 CU 会有多少线程，会按最坏情况分配——可能只允许 2 block/CU，导致 occupancy 从 4 block 降到 2 block。
// 不加 __launch_bounds__:
// 编译器假设：未知线程数 → 保守分配 ~50 reg/thread
// 每 CU 线程数 = 65536/50 ≈ 1310 → 约 5 wavefront
// 每 block 256 线程 = 4 wavefront → 只能 1 block/CU
// occupancy ≈ 25% ← 差！
// 加 __launch_bounds__(256, 4):
// 编译器知道：256 thread/block，至少 4 block/CU
// 每线程寄存器上限 = 65536/(256×4) = 64
// 实际用 50 < 64 → 不 spill
// occupancy ≈ 100% ← 优
验证寄存器使用量：
# 用 hipcc 编译时加上 –save-temps 查看汇编
hipcc -c kernel.cpp –save-temps -o kernel.o
# 查看寄存器使用（在 .hsaco 或 .o 文件中）
# 搜索 "vgpr" 或 "sgpr"
# 例如：vgpr_count: 48 表示用了 48 个向量寄存器
3.2.4 寄存器 vs 共享内存 vs 全局内存
*结论：对于线程私有的归约操作，寄存器永远是第一选择。
*3.3 策略三：分支分歧消除与统计矩替代
3.3.1 为什么分支在 DCU 上特别贵？
在 wavefront 内，所有 64 个线程共享一个程序计数器（PC）。这意味着它们在同一时刻执行同一条指令。当遇到条件分支时：
if(condition){
// path A: 线程 0-31 走这里
} else {
// path B: 线程 32-63 走这里
}
DCU 的处理方式是：
先执行 path A，同时掩蔽（mask off）不走 path A 的 32 个线程
再执行 path B，同时掩蔽不走 path B 的 32 个线程
浪费：50% 的执行资源被掩蔽
在 NVIDIA 上（32 线程 warp）：
16 线程 + 16 线程 = 50% 浪费
在 DCU 上（64 线程 wavefront）：
32 线程 + 32 线程 = 50% 浪费（比例相同）
但 64 个线程被掩蔽 vs 32 个线程被掩蔽，总时间片浪费翻倍
3.3.2 堆插入的分歧问题
Top-K 堆方案中，每个值都需要判断 if(v > heap[0])——大约只有 34⁄330 ≈ 10% 的值会进入堆。这意味着 90% 的线程走 else 路径，10% 走 if 路径，分歧严重。
// ★ 分歧在这里
if(v > (double)heap[0]){ // ~10% true, ~90% false
heap[0] = (float)v; // 仅 10% 线程执行
// sift-down 循环（更多分支）
while(…){ … }
}
// else: 什么都不做（90% 线程空闲）
3.3.3 统计矩：完全消除分支
统计矩方案的核心思想是：不维护精确堆，只维护 4 个统计矩。矩的累加不需要任何条件判断，所有线程执行完全相同的 4 条 FMA 指令：
// ★ 无分支、无分歧、所有线程执行相同代码
float fv = (float)v;
m1 += fv; // 一阶矩：∑x
m2 += fv*fv; // 二阶矩：∑x²
m3 += fv*fv*fv; // 三阶矩：∑x³
m4 += fv*fv*fv*fv; // 四阶矩：∑x⁴
count++;
DCU 上的执行效果：
没有 if-else → 没有 wavefront 分歧
所有线程 ALU 100% 忙碌
DCU 利用率从 ~40% 提升到 ~95%
3.3.4 从矩到 P90：Cornish-Fisher 展开
有了 4 个统计矩之后，如何得到第 90 百分位？Cornish-Fisher 展开提供了从矩到分位数的近似公式：
Step 1: 计算均值和方差
float n = (float)count;
float mean = m1 / n;
float variance = m2 / n – mean * mean;
float std = sqrtf(variance);
Step 2: 计算偏度和峰度
// 偏度 Skewness：衡量分布不对称程度
float skewness = (m3/n – 3*mean*variance – mean*mean*mean)
/ (variance * std);
// 峰度 Kurtosis：衡量分布尾部厚度
float kurtosis = (m4/n – 4*mean*m3/n + 6*mean*mean*variance
+ mean*mean*mean*mean) / (variance*variance) – 3;
Step 3: Cornish-Fisher 展开求 P90
// 标准正态分布的 P90 分位数
float z = 1.2816f;
// Cornish-Fisher 展开（取前三阶修正）
float cf = z
+ (z*z – 1) * skewness / 6
+ (z*z*z – 3*z) * kurtosis / 24
– (2*z*z*z – 5*z) * skewness * skewness / 36;
float p90 = mean + std * cf;
精度说明：对于 SST 数据，如果分布接近正态，CF 展开的误差 < 0.1°C。对于有偏分布（如洋流交汇区），误差最大约 0.5°C。在气候监测应用中，这个精度是可接受的。
3.3.5 性能对比
**3.4 策略四：加法群降维
3.4.1 问题：滑动窗口的冗余计算
SST 百分位计算中，每个 DOY 使用前后各 5 天的数据（11 天窗口）。相邻 DOY 的窗口有 10 天是重叠的：
DOY 152: [147, 148, 149, 150, 151, 152, 153, 154, 155, 156, 157]
DOY 153: [148, 149, 150, 151, 152, 153, 154, 155, 156, 157, 158]
↑ 这10天是重复计算的！
DOY 154: [149, 150, 151, 152, 153, 154, 155, 156, 157, 158, 159]
↑ 又有10天重复
传统方法中，每个 DOY 独立计算 30年 × 11天 = 330 次 FMA。12 个 DOY 就是 12 × 330 = 3960 次 FMA/格点。其中约 91% 是重复的！
3.4.2 核心洞察：矩是加法群
统计矩具有完美的加法性质：
*
也就是说，窗口的矩 = 窗口内单天矩的和。这意味着我们不需要为每个窗口独立计算——先算好每个”单天”的矩，然后用”加法 + 减法”来滑动窗口。
3.4.3 两阶段 kernel 实现
第一阶段：day_moment kernel——只算单天，不做窗口
这个 kernel 只做一件事：对每个日历日，累加 30 年数据的 4 阶矩。不涉及任何滑动窗口。
// ★ 第一阶段：单天矩累加
// 输入：d_ring（BATCH年 × ns天的数据）
// 输出：day_mom（ns天 × 每格点 × 4矩）
__global__ void
day_moment(const double* r, // 输入数据
float* dm, // day_mom输出：ns × sz × 4
int* dc, // day_cnt输出：ns × sz
size_t sz, int nd, int ny)
{
size_t ti = blockIdx.x * blockDim.x + threadIdx.x;
if(ti >= sz) return;
// 遍历所有年份和天数
for(int y = 0; y < ny; y++) // 年份循环
for(int d = 0; d < nd; d++){ // 天数循环
double v = r[((size_t)y * nd + d) * sz + ti];
if(!isnan(v)){
float fv = (float)v;
// ★ 累加到第 d 天、第 ti 格点的矩中
float* m = &dm[((size_t)d * sz + ti) * 4];
m[0] += fv; // 矩1
m[1] += fv*fv; // 矩2
m[2] += fv*fv*fv; // 矩3
m[3] += fv*fv*fv*fv; // 矩4
dc[(size_t)d * sz + ti]++; // 计数
}
}
}
第二阶段：slide_moment kernel——滑动窗口合成 DOY 矩
这个 kernel 利用已算好的单天矩，用前缀和方法快速合成每个 DOY 的 11 天窗口矩。
// ★ 第二阶段：滑动窗口合成
// 输入：day_mom（单天矩）
// 输出：doy_mom（每个DOY的窗口矩）
__global__ void
slide_moment(
const float* dm, // day_mom: 单天矩
const int* dc, // day_cnt: 单天计数
float* ym, // doy_mom: 输出DOY矩
int* yc, // doy_cnt: 输出DOY计数
size_t sz, int ndo, int ns, int nd)
{
size_t ti = blockIdx.x * blockDim.x + threadIdx.x;
if(ti >= sz) return;
// 跨 DOY 的滑动累加器（在寄存器中持续更新）
float pm1 = 0, pm2 = 0, pm3 = 0, pm4 = 0;
int pvc = 0;
for(int di = 0; di < ndo; di++){
if(di == 0){
// ★ 第一个 DOY：累加 nd 天（11天）
for(int d = 0; d < nd; d++){
const float* m = &dm[((size_t)d * sz + ti) * 4];
pm1 += m[0]; pm2 += m[1];
pm3 += m[2]; pm4 += m[3];
pvc += dc[(size_t)d * sz + ti];
}
} else {
// ★ 后续 DOY：只做 加新天 – 减旧天
int L = di – 1; // 滑出的天
int R = di + nd – 1; // 滑入的天
const float* mL = &dm[((size_t)L * sz + ti) * 4];
const float* mR = &dm[((size_t)R * sz + ti) * 4];
pm1 += -mL[0] + mR[0];
pm2 += -mL[1] + mR[1];
pm3 += -mL[2] + mR[2];
pm4 += -mL[3] + mR[3];
pvc += -dc[(size_t)L * sz + ti] + dc[(size_t)R * sz + ti];
}
// 写回当前 DOY 的结果
float* ym2 = &ym[((size_t)di * sz + ti) * 4];
ym2[0] = pm1; ym2[1] = pm2;
ym2[2] = pm3; ym2[3] = pm4;
yc[(size_t)di * sz + ti] = pvc;
}
}
3.4.4 计算量对比
传统方法 ka_moment（每DOY遍历11天×30年）：
12 DOY × 30年 × 11天 × 4 FMA = 15,840 FMA
降维法 day_moment（每个单天遍历30年）：
22天 × 30年 × 4 FMA = 2,640 FMA
降维法 slide_moment（滑动窗口）：
第1 DOY: 11天 × 8 add/sub = 88 ops
第2~12 DOY: 各 8 add/sub = 88 ops
合计: 176 ops
降维法总操作: 2,640 + 176 = 2,816 ops（含FMA + 加减）
不对，让我们回到 kernel 的实际循环。
ka_moment 的双重循环：
for di(12) for y(30) for d(11) {
m1 += v; m2 += v²; m3 += v³; m4 += v⁴;
}
= 12 × 30 × 11 × 4 = 15,840 FMA
day_moment + slide_moment：
day: for y(30) for d(22) { 4 FMA } = 30×22×4 = 2,640 FMA
slide: 第1DOY = 11×4 = 44 加减 + 后续11×8 = 88 = 132 加减
总 = 2,640 + 132 = 2,772 ops
节省 = (15,840 – 2,772) / 15,840 = 82.5%
GPU kernel 时间从 0.35s（ka_moment）降至 0.08s（day_moment + slide_moment），降幅 77%。
3.4.5 可推广性
加法群降维不仅适用于 SST 百分位计算。任何具有以下特征的滑动窗口计算都可以用：
窗口高度重叠（相邻窗口重叠率 > 50%）
统计量满足可加性（矩、直方图、计数、求和）
可推广场景：
气象数据的滑动平均、滑动方差
金融时间序列的移动窗口统计
信号处理中的短时傅里叶变换
图像处理的滑动窗口滤波
*3.5 策略五：多流并发与 DMA 计算重叠
3.5.1 背景：DCU 的 DMA 引擎
DCU 内部有两个独立工作的硬件单元：
DMA 引擎：负责 host↔device 数据传输
计算引擎：负责执行 kernel 指令
这两个引擎可以同时运行。这意味着 DCU 可以在传输数据的同时进行计算——前提是你用对了编程模式。
3.5.2 默认行为 vs 多流并发
默认（单流）：
hipMemcpy(d_ring, host_ring, size, hipMemcpyHostToDevice); // 同步传输
kernel<<<…, stream>>>(d_ring, …); // 然后计算
hipStreamSynchronize(stream);
时间线：[H2D][H2D]….[Compute][Compute] → 串行执行
多流并发：
// 创建两个流
hipStream_t st_copy, st_cmpt;
hipStreamCreate(&st_copy);
hipStreamCreate(&st_cmpt);
// H2D 传输在 st_copy 上
hipMemcpy2DAsync(d_ring[1], dst_pitch,
host_ring[next], src_pitch, width, height,
hipMemcpyHostToDevice, st_copy);
// 计算在 st_cmpt 上（使用另一份数据）
kernel<<<grid, block, 0, st_cmpt>>>(d_ring[0], …);
时间线：[H2D…H2D] 与 [Compute…Compute] 并行 → 重叠执行
3.5.3 双缓冲 + 双流实战
// ★ 双缓冲 + 双流：H2D 和 kernel 在不同流上重叠
// 流 0：负责 H2D 传输（未来批次的预加载）
hipStream_t st_copy[4];
// 流 1：负责 kernel 计算（当前批次的处理）
hipStream_t st_cmpt[4];
// d_ring 也分两份，避免数据竞争
double* d_ring[4][2];
for(int b = 0; b < nb; b++){
int cur_buf = b % 2; // 当前用哪份 buffer
int next_buf = 1 – cur_buf; // 下一份 buffer
// ★ 在当前流上启动当前批次的 kernel
// （kernel 读 d_ring[g][cur_buf]）
ka_kernel<<<blk, thr, 0, st_cmpt[g]>>>(
d_ring[g][cur_buf], …);
// ★ 在拷贝流上启动下一批数据的 H2D 传输
// （DMA 写 d_ring[g][next_buf]）
// ★ 和上方的 kernel 同时执行！
HIP_CHECK(hipMemcpy2DAsync(
d_ring[g][next_buf], dpitch,
host_ring[next], spitch,
width, height,
hipMemcpyHostToDevice,
st_copy[g])); // ← 不同流
// ★ 只需要同步计算流
HIP_CHECK(hipStreamSynchronize(st_cmpt[g]));
// 拷贝流不需要在这里同步——下次用到 next_buf 时自然保证就绪
}
3.5.4 收益分析
*当 H2D 时间 ≈ GPU 时间时，双流重叠的加速比接近 2×。
*3.6 策略六：流水线环数匹配
3.6.1 核心矛盾
在 SST 计算程序中，文件 I/O 是比 GPU 计算慢得多的环节：
*GPU 只用了 0.15s 就处理完一批数据，但下一批数据还要 2.3s 才能读完。这就意味着 GPU 大部分时间在等数据——即我们前面测量的 HIDDEN_WAIT。
3.6.2 环数的数学推导
流水线中的每个”环”（ring buffer）经过三个阶段：
Ring 生命周期：[文件读取 (T_io)] → [H2D+GPU (T_gpu)] → [空闲]
如果有 R 个环交替使用，那么一个新批次的读取可以提前 R-1 个批次开始。只要提前开始的时间足够覆盖读取时间，GPU 就不用等。
条件：提前开始的时间 ≥ 读取时间
*
最小环数：
*
代入我们的数据：*，最小 17 环。
但 17 个环 × 915 MB = 15.5 GB 的 pinned memory，超出了每节点的 /dev/shm 容量。所以只能取折中。
3.6.3 实际效果
*等待时间的可视化（4环 vs 2环）：
2环:
GPU: [批0][____等____][批1][____等____][批2][____等____]…
读: [读0] [读1] [读2]
← 每批GPU只做0.15s，等2.15s，利用率6.5%
4环:
GPU: [批0][批1][批2][批3][__等__][批4][批5]…
读: [读0][读1][读2][读3] [读4][读5]
← 前3批零等待，后3批少量等待，利用率约40%
3.6.4 环数实现要点
// 4 环流水线主循环
// 预读第 1,2,3 批到 ring[1], ring[2], ring[3]
for(int i = 1; i < 4; i++) start_read(i, i);
for(int b = 0; b < nb; b++){
int ri = b % 4; // 当前使用的环
// 等待当前环的数据就绪
if(rd_batch[ri] >= 0) rd_fut[ri].get();
// 用这个刚释放的环预读后面的批次
start_read(b + 4, ri);
// H2D + GPU 计算
hipMemcpy2DAsync(d_ring, host_ring[ri], …);
kernel<<<…>>>(d_ring, …);
hipStreamSynchronize(stream);
}
*3.7 策略七：Pinned Memory 管理
3.7.1 为什么需要 pinned memory？
DCU 通过 PCIe DMA 引擎访问 host 内存。DMA 要求源/目标地址的物理页是锁定在内存中的（不会被交换出去）。malloc 分配的内存物理页可以被操作系统换出，因此 DMA 无法直接访问——需要 ROCm 驱动做一次中转拷贝，速度减半。
// ❌ malloc：DMA 无法直接访问
double* buf = (double*)malloc(size); // 虚拟内存，页可能被换出
hipMemcpy(d_ring, buf, size, …); // 实际走：buf→临时pinned buffer→device
// ✅ hipHostMalloc：DMA 可直接访问
double* buf;
hipHostMalloc(&buf, size); // /dev/shm 上的 pinned memory
hipMemcpy(d_ring, buf, size, …); // DMA 直传，无中转
3.7.2 性能差异
*对于 915 MB 的 H2D 传输：
malloc：915/4 = 228 ms
hipHostMalloc：915/12 = 76 ms
3× 差距。
3.7.3 注意事项
1. /dev/shm 容量限制
hipHostMalloc 分配的内存来自 /dev/shm（tmpfs）。总量受 /dev/shm 大小限制：
$ df -h /dev/shm
tmpfs 62G 14G 49G 22% /dev/shm
每节点 49 GB 可用。4 ring × 915 MB × 4 rank = 14.6 GB，在安全范围内。
2. 多进程共享
DCU 上使用 hipHostMallocPortable 标记可以让分配的内存被所有进程访问（适用于 MPI 场景）：
hipHostMalloc(&buf, size, hipHostMallocPortable);
3. 清理顺序
// ★ 正确的清理顺序
hipHostFree(host_ring[0]); // 立即释放 /dev/shm
hipHostFree(host_ring[1]);
hipHostFree(host_ring[2]);
hipHostFree(host_ring[3]);
// ★ 跳过 hipFree（GPU 显存释放不是 /dev/shm 的问题）
// ★ 必须调用 MPI_Finalize（清理 PMIX 小文件残留）
MPI_Finalize();
// ★ 最后兜底清扫
system("rm -rf /dev/shm/* 2>/dev/null");
*4. 实验数据汇总
4.1 硬件环境
*4.2 各版本性能
*4.3 各策略加速因子
**5. 常见陷阱与调试建议
5.1 陷阱一：hipHostMalloc 静默失败
hipHostMalloc(&buf, huge_size); // 如果 /dev/shm 不够，返回非NULL但性能极差
// 应该：
hipError_t e = hipHostMalloc(&buf, huge_size);
if(e != hipSuccess) printf("pinned alloc failed, will be slow!\\n");
5.2 陷阱二：kernel 编译成功但不执行
kernel<<<grid, block, 0, stream>>>(…);
// ★ 没有返回值检查！kernel launch 可能失败。
// 应该：
kernel<<<grid, block, 0, stream>>>(…);
hipError_t e = hipGetLastError(); // launch 错误
hipStreamSynchronize(stream); // 执行错误
5.3 陷阱三：核函数访问越界
用 hipMemcheck 工具检测：
hipcc -g -O0 kernel.cpp -o kernel
hipMemcheck ./kernel
5.4 陷阱四：launch_bounds 设得太紧
__launch_bounds__(256, 4) // 承诺最多 256 线程/块，至少 4 block/CU
// 但如果 kernel 实际需要 >64 寄存器，编译器只能 spill
// 症状：kernel 编译成功但性能比不加还差
// 解决：调低 N（min blocks），让编译器分配更多寄存器
__launch_bounds__(256, 2) // 允许 2 block/CU，每线程 128 寄存器
5.5 陷阱五：忽略 device 切换
hipSetDevice(0);
hipMalloc(&ptr0, size); // 在 device 0 上分配
// ★ 忘了 hipSetDevice(1)！
hipMalloc(&ptr1, size); // 仍然在 device 0 上分配！
// 然后：
kernel<<<…, stream_on_dev1>>>(ptr1, …);
// HIP_CHECK(hipGetLastError()) → hipErrorInvalidValue
// 因为 ptr1 在 device 0 上，但 stream 在 device 1 上
*6. 总结
本文以气候态 SST 百分位计算为教学案例，系统介绍了 AMD DCU 上的七条算子优化策略。核心要点：
DCU 不是 CUDA：wavefront 64、寄存器 64K/CU、L1 16KB——这些数字决定了你的优化方向
分歧在 DCU 上最贵：64 线程 wavefront 使分支分歧的代价是 NVIDIA 的 2 倍
寄存器是命根子：善用 __launch_bounds__ + 寄存器级归约，避免 spill
数学可以替代计算：统计矩替代精确堆（-60% 计算量），加法群降维（-83% 计算量）
系统优化比 kernel 优化更重要：当 GPU 时间降到 0.08s（仅占 2%），I/O 才是真瓶颈
最终 GPU kernel 从 8.0s 降到 0.08s（99% 降幅），整体程序从 14.0s 降到 4.2s（3.3× 加速）。
*参考文献
[1] AMD, Inc. “AMD ROCm Documentation: HIP Programming Guide v5.x.” https://rocm.docs.amd.com/projects/HIP/en/latest/
[2] AMD, Inc. “AMD Instinct MI50 Accelerator Specifications.” https://www.amd.com/en/products/accelerators/instinct/mi50.html
[3] AMD, Inc. “AMD GPU Architecture: CDNA 2.” Whitepaper, 2022.
[4] Hobday, A.J., et al. “A hierarchical approach to defining marine heatwaves.” Progress in Oceanography, 141:227-238, 2016. DOI: 10.1016/j.pocean.2015.12.014
[5] Cornish, E.A. & Fisher, R.A. “Moments and cumulants in the specification of distributions.” Revue de l’Institut International de Statistique, 5(4):307-320, 1937.
[6] NVIDIA Corporation. “CUDA C++ Programming Guide.” https://docs.nvidia.com/cuda/cuda-c-programming-guide/
[7] Unidata/UCAR. “NetCCF Documentation.” NetCDF | NSF Unidata
[8] The HDF Group. “HDF5 Documentation.” https://portal.hdfgroup.org/display/HDF5/HDF5
https://zhuanlan.zhihu.com/p/2052383479850078572https://zhuanlan.zhihu.com/p/2052383479850078572（这个是本人zh发布同一文章）
