海光 DCU 加速器上的算子优化策略教学——以气候态海温百分位计算为例
摘要
海光 DCU（Deep Computing Unit）是 ROCm 生态下的 GPU 加速器，已在多个国产超算系统中部署。本文以一个气候态海温第 90 百分位计算程序为教学案例，面向 DCU 上的 HIP 编程初学者，系统介绍三类典型算子——Top-K 归约、滑动窗口统计矩、多级流水线——的优化思路。全文围绕七条策略展开：wavefront 对齐与线程块配置、寄存器级归约与 launch_bounds 约束、分支分歧消除与统计矩替代、加法群降维、多流并发与 DMA 计算重叠、流水线环数匹配、pinned memory 管理。每条策略都包含架构根因分析、完整代码示例和量化对比。实验部分给出了各优化步骤的加速效果：GPU kernel 时间从 8.0 秒降至 0.08 秒，总程序从 14.0 秒降至 4.2 秒。
关键词：海光 DCU；HIP；GPU 算子优化；归约；滑动窗口；流水线
1. 引言
1.1 DCU 是什么
DCU（Deep Computing Unit）是海光面向 HPC 和 AI 场景设计的 GPU 加速器。它的编程模型 HIP（Heterogeneous-compute Interface for Portability）在语法上高度兼容 CUDA——如果你写过 CUDA，HIP 代码看起来几乎一模一样：
// CUDA 版本 cudaMalloc(&d_ptr, size); cudaMemcpy(d_ptr, h_ptr, size, cudaMemcpyHostToDevice); kernel<<<grid, block>>>(d_ptr); // HIP 版本（只有前缀不同） hipMalloc(&d_ptr, size); hipMemcpy(d_ptr, h_ptr, size, hipMemcpyHostToDevice); kernel<<<grid, block>>>(d_ptr);
但 DCU 的底层架构与 NVIDIA GPU 有显著差异。本文从架构差异出发，解释每条优化策略背后的 WHY——为什么这样做在 DCU 上有效。
1.2 计算问题
本文使用的教学案例是一个气候态海温（SST）第 90 百分位计算程序。输入数据为 1991—2020 年（30 年）全球逐日海温，空间分辨率 0.25 度 x 0.25 度，网格 1440 x 721。需要计算每个海洋格点每个日历日（DOY 152-243）的：
-
气候态均值：该格点 30 年同一天附近 11 天窗口的平均值
-
第 90 百分位 SST：该格点 30 年同一天附近 11 天窗口的 P90
数据以 NetCDF4/HDF5 格式存储，单文件解压后约 8.3 MB。程序使用 MPI + OpenMP + HIP 三级混合并行，部署在 2 节点 x 4 DCU/节点的超算集群上。
1.3 三类典型算子
| 算子类型 | 对应计算 | 优化难度 |
|---|---|---|
| Top-K 归约 | 从 330 个值中找第 90 百分位 | 中等——寄存器管理 + 分支分歧 |
| 滑动窗口统计矩 | 相邻 DOY 的 11 天窗口矩计算 | 高——需利用数学性质降维 |
| 多级流水线 | 文件读 -> H2D -> GPU -> D2H -> 写文件 | 低——但收益最大 |
2. DCU 架构关键参数
2.1 架构规格
| 参数 | 海光 DCU (gfx906) | NVIDIA V100 | 对优化策略的影响 |
|---|---|---|---|
| 计算单元 (CU/SM) | 64 | 80 | 需足够线程填满所有 CU |
| 每 CU 最大 resident 线程 | 2560 (40 wave x 64) | 2048 (64 warp x 32) | — |
| Wavefront / Warp 大小 | 64 | 32 | 分歧代价翻倍 |
| 每 CU 寄存器文件 | 64 KB | 256 KB | 寄存器更紧张，需精确控制 |
| L1 缓存 / CU | 16 KB | 128 KB | 更依赖寄存器而非缓存 |
| 共享内存 / CU | 64 KB | 96 KB | 共享内存受限 |
2.2 三个核心原则
原则 1：线程块应为 64 的倍数
DCU 的最小调度单位是 wavefront（64 线程）。线程块大小不是 64 的倍数时，最后一个 wavefront 中的空闲线程仍占用 CU 资源。
-
推荐：64、128、256
-
不推荐：32、96、160、224
原则 2：每线程寄存器预算约 64 个
每 CU 64K 寄存器被所有 resident wavefronts 共享：
__launch_bounds__(256, 4) → 65536 / (4 × 256) = 64 reg/thread __launch_bounds__(256, 2) → 65536 / (2 × 256) = 128 reg/thread __launch_bounds__(128, 8) → 65536 / (8 × 128) = 64 reg/thread
超过预算的寄存器会 spill 到 local memory（L1 仅 16 KB，spill 代价极高）。
原则 3：分歧代价与 wavefront 大小成正比
64 线程 wavefront 中若出现二分叉，最多 50% 线程闲置。消除一个分支在 DCU 上的收益是 NVIDIA 的 2 倍（因为 wavefront 是 warp 的 2 倍大）。
3. 七条算子优化策略
3.1 策略一：Wavefront 对齐与线程块配置
问题：线程块大小不是 64 的倍数时，最后一个 wavefront 利用不充分。
错误示范：
int thr = 160; // 160 / 64 = 2.5 wavefront // 第 3 个 wavefront 只有 32 个有效线程 // 32 个线程槽位被浪费 int blk = (N + thr - 1) / thr; kernel<<<blk, thr>>>(...);
正确做法：
int thr = 256; // 256 / 64 = 4 个完整 wavefront int blk = (N + thr - 1) / thr; // blk × thr 可能略多于 N，多出的线程通过 if(idx >= N) return; 处理 kernel<<<blk, thr>>>(...);
速查表：
| 线程块大小 | wavefront 数 | 利用率 | 适用场景 |
|---|---|---|---|
| 64 | 1 | 100% | 寄存器压力大时 |
| 128 | 2 | 100% | 通用 |
| 256 | 4 | 100% | 推荐（寄存器充裕时最优） |
| 32 | 0.5 | 50% | 不推荐 |
| 512 | 8 | 100% | 可能超寄存器预算 |
3.2 策略二：寄存器级归约与 launch_bounds 约束
问题：DCU 的 L1 缓存仅 16 KB/CU，依赖全局内存或共享内存的归约会受带宽限制。
架构根因：寄存器是 DCU 上最快的存储层次——比共享内存快 5-10 倍，比全局内存快 100 倍。对于线程内独立的归约（每个线程处理不同格点），寄存器是最优选择。
__launch_bounds__ 用法详解：
// 形式：__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
// 告诉编译器：这个 kernel 最多用 M 线程/块，每 CU 至少 N 个块
__global__ void __launch_bounds__(256, 4) my_kernel(...)
{
// 编译器会根据 256*4=1024 thread/CU 来分配寄存器
// 每线程最多 65536/1024 = 64 个寄存器
// 如果实际需要更多，编译器选择 spill（伤性能）或报错
}
寄存器级 Top-K 堆的完整实现：
这个 kernel 维护 34 个 float 的最小堆，全部在寄存器中：
// 34 float × 4 byte = 136 byte ≈ 34 个寄存器
// 加上索引变量：~45 个寄存器
// __launch_bounds__(256,4) → 64 寄存器预算 ✓
__global__ void __launch_bounds__(256, 4)
topk_reduce(const double* data, float* heap_out, int cnt_per_thread, int stride)
{
size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
// ★ 34 元素 float 堆，全部在寄存器中
float heap[34];
#pragma unroll
for(int i = 0; i < 34; i++) heap[i] = -1e38f;
// 遍历当前线程负责的所有值
for(int i = 0; i < cnt_per_thread; i++){
float v = (float)data[i * stride + tid];
if(v > heap[0]){ // 大于堆顶（堆中最小值）才插入
heap[0] = v;
// ★ 寄存器内 sift-down，无显存访问
int pos = 0;
while(pos < 34){
int child = 2 * pos + 1;
if(child + 1 < 34 && heap[child+1] < heap[child])
child++;
if(child >= 34 || heap[pos] <= heap[child])
break;
float tmp = heap[pos];
heap[pos] = heap[child];
heap[child] = tmp;
pos = child;
}
}
}
// ★ 只写一次全局内存
for(int i = 0; i < 34; i++)
heap_out[tid * 34 + i] = heap[i];
}
寄存器分配诊断：编译时添加 --save-temps 选项，查看生成的 .s 文件中寄存器使用量：
hipcc -O3 --save-temps kernel.cpp -o kernel grep "sgpr_count\|vgpr_count" kernel.s # sgpr_count: 42 (标量寄存器) # vgpr_count: 12 (向量寄存器) # 总寄存器 = 42 + 12 = 54 < 64 ✓
效果：GPU 计算时间从约 8 秒（未优化基线）降至 0.9 秒（寄存器级归约 + device ring），加速约 8.9 倍。
3.3 策略三：分支分歧消除与统计矩替代
问题：Top-K 堆插入中的条件分支在 DCU 上产生严重的 wavefront 分歧。
分歧代价量化分析：
考虑一个 wavefront 中 64 个线程执行：
if(v > heap[0]){ // 分歧点：部分 true，部分 false
// true 路径：堆插入
while(1){ ... // 内部分歧（不同循环次数） }
} else {
// false 路径：什么都不做
}
DCU 的执行方式：wavefront 中的每个线程都有一个状态掩码（active mask），只有掩码为 1 的线程执行当前指令。当遇到分支时：
if(v > heap[0]) → 64 线程中约一半 true，一半 false 第 1 阶段：mask = 0b00001111（假设低 4 位为 true） 执行 true 路径的指令 false 路径的线程被掩蔽 第 2 阶段：mask = 0b11110000（高 4 位为 false） 执行 false 路径的指令 true 路径的线程被掩蔽 总指令数 = true 路径指令数 + false 路径指令数 理想情况（无分歧）= max(true, false) 路径指令数
堆插入的分歧尤为严重：64 个线程中，每个线程可能执行不同次数的 sift-down 循环（有些 0 次、有些 1 次、有些 2 次...），导致 wavefront 内严重分歧。
无分支替代方案：统计矩法
// --- 原始：有分歧 ---
if(v > heap[0]){ ← 分歧 1
heap[0] = v;
while(1){
// sift-down ← 分歧 2（循环次数不同）
...
}
}
// --- 优化后：无分歧 ---
// 所有线程执行完全相同的 4 条 FMA 指令
float fv = (float)v;
if(!isnan(v)){ ← 唯一分歧，但概率低（海洋数据 ~70% 有效）
m1 += fv; // FMA，无分歧
m2 += fv * fv; // FMA，无分歧
m3 += fv * fv * fv; // FMA，无分歧
m4 += fv * fv * fv * fv; // FMA，无分歧
}
Cornish-Fisher 展开的完整推导：
要从 4 阶矩求 P90，使用 Cornish-Fisher 展开：
Step 1：从矩计算统计量 均值 μ = M1 / n 方差 σ² = M2/n - μ² 标准差 σ = sqrt(σ²) 偏度 S = (M3/n - 3μσ² - μ³) / σ³ 峰度 K = (M4/n - 4μ·M3/n + 6μ²σ² + μ⁴) / (σ²)² - 3 Step 2：Cornish-Fisher 展开求分位数 z = 1.2816 (标准正态分布的 P90 分位数) CF = z + (z²-1)·S/6 + (z³-3z)·K/24 - (2z³-5z)·S²/36 Step 3：P90 = μ + σ · CF
代码实现：
__global__ void kf_cf(const float* mom, const int* cnt,
double* cl, double* p, size_t sz, int ndo)
{
size_t ti = blockIdx.x * blockDim.x + threadIdx.x;
if(ti >= sz) return;
for(int di = 0; di < ndo; di++){
int vc = cnt[di * sz + ti];
if(vc == 0){ // 陆地，无效
cl[di * sz + ti] = NAN;
p[di * sz + ti] = NAN;
continue;
}
const float* m = &mom[((size_t)di * sz + ti) * 4];
float n = (float)vc, mu = m[0] / n;
float s2 = m[1] / n - mu * mu, s = sqrtf(s2);
if(s < 1e-6f){ // 零方差：P90 = 均值
cl[di * sz + ti] = mu;
p[di * sz + ti] = mu;
continue;
}
float s3 = s2 * s, s4 = s2 * s2;
float sk = (m[2]/n - 3*mu*s2 - mu*mu*mu) / s3;
float ku = (m[3]/n - 4*mu*m[2]/n + 6*mu*mu*s2 + mu*mu*mu*mu) / s4 - 3.0f;
float z = 1.2816f;
float cf = z + (z*z - 1)*sk/6 + (z*z*z - 3*z)*ku/24
- (2*z*z*z - 5*z)*sk*sk/36;
cl[di * sz + ti] = mu;
p[di * sz + ti] = mu + s * cf;
}
}
精度分析：统计矩法 P90 与精确 Top-K 堆的偏差通常 < 0.5°C。在全球 SST 数据验证中，约 90% 格点的偏差 < 0.3°C，满足气候监测需求。
DCU 利用率提升：
Top-K 堆：warp 分歧率约 60%，DCU 利用率 ~40% 统计矩法：分歧率 0%，DCU 利用率 ~95% GPU 时间：0.9s → 0.35s（-61%）
3.4 策略四：加法群降维
数学原理：统计矩是加法群同态。对于任意两个不相交的集合 A 和 B：
M_k(A ∪ B) = M_k(A) + M_k(B)
其中 M_k 是 k 阶矩（M₁ = Σx, M₂ = Σx², M₃ = Σx³, M₄ = Σx⁴）。
这一性质意味着：滑动窗口的矩可以通过"窗口起点"到"窗口终点"的前缀和相减得到——无需重新计算窗口内的所有元素。
三阶段演进：
阶段 1——原始 ka_moment（每 DOY 窗口独立计算）：
for each DOY: for each year: for each day in window (11 days): m1 += val; m2 += val*val; ... // 22 FMA // 每个 DOY 做 330 次 FMA // 12 个 DOY = 3960 FMA/格点
阶段 2——day_moment（只算单天跨年矩）：
for each year: for each day (22 days): m1 += val; m2 += val*val; ... // 4 FMA // 30 年 x 22 天 = 660 FMA/格点
阶段 3——slide_moment（滑动窗口合成）：
// DOY[0] = day[0] + day[1] + ... + day[10] (11 个) // DOY[1] = DOY[0] - day[0] + day[11] (-1 +1 = 2 个) // DOY[2] = DOY[1] - day[1] + day[12] (-1 +1 = 2 个) // ... // 1 + 11 x 2 = 23 次加减/格点 // 换算为 FMA ≈ 12 FMA/格点
三个 kernel 的完整代码：
// ★ day_moment：单天跨年矩累加
__global__ void day_moment(const double* r, float* dm, int* dc,
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[(y * nd + d) * sz + ti];
if(!isnan(v)){
float fv = (float)v;
float* m = &dm[((size_t)d * sz + ti) * 4];
m[0] += fv; m[1] += fv * fv;
m[2] += fv * fv * fv;
m[3] += fv * fv * fv * fv;
dc[(size_t)d * sz + ti]++;
}
}
}
// ★ slide_moment：滑动窗口合成（前缀和）
__global__ void slide_moment(const float* dm, const int* dc,
float* ym, int* yc, size_t sz, int ndo, int ns, int nd)
{
size_t ti = blockIdx.x * blockDim.x + threadIdx.x;
if(ti >= sz) return;
float pm1=0, pm2=0, pm3=0, pm4=0;
int pvc = 0;
for(int di = 0; di < ndo; di++){
if(di == 0){
// 第一个 DOY：累加 nd 天
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, 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];
}
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;
}
}
效果对比：
| kernel | 每格点 FMA | 相对原始 | 绝对节省 |
|---|---|---|---|
| ka_moment | 3960 | 100% | — |
| day_moment | 660 | 17% | -3300 FMA |
| slide_moment | 12 | 0.3% | -3876 FMA |
| 合计 | 672 | 17% | -83% |
GPU 时间从 0.35s 降至 0.08s（-77%）。
3.5 策略五：多流并发与 DMA 计算重叠
架构根因：DCU 的 DMA 引擎（负责 PCIe 传输）与计算引擎（执行 kernel）是独立的硬件单元，可同时操作。
错误做法——串行传输和计算：
hipMemcpy(d_ring, h_data, size, hipMemcpyHostToDevice); // 等传输完 kernel<<<...>>>(d_ring, ...); // 再算 hipStreamSynchronize(st); // DMA 引擎空闲等待 kernel 执行 // 计算引擎空闲等待 DMA // 利用率低下
正确做法——双流重叠：
hipStream_t st_copy, st_cmpt; hipStreamCreate(&st_copy); hipStreamCreate(&st_cmpt); // 流 1 (st_copy)：负责 H2D 传输 hipMemcpy2DAsync(d_ring[next], ..., h_data[next], ..., hipMemcpyHostToDevice, st_copy); // 流 2 (st_cmpt)：负责计算（使用上次已传好的数据） kernel<<<..., st_cmpt>>>(d_ring[cur], ...); // 两个流同时进行：DMA 和 kernel 并发 // 同步 hipStreamSynchronize(st_cmpt); hipStreamSynchronize(st_copy);
双缓冲 d_ring 的完整流水线：
// 两块 d_ring 交替：当前批计算时，下一批已经在传
int buf = 0;
for(int b = 0; b < nb; b++){
int next_buf = 1 - buf;
// 启动下一批 H2D（在 stream_copy 上，与计算并行）
if(b + 1 < nb){
hipMemcpy2DAsync(d_ring[next_buf], ...,
hipMemcpyHostToDevice, st_copy);
}
// 当前批计算（在 stream_compute 上）
kernel<<<..., st_cmpt>>>(d_ring[buf], ...);
hipStreamSynchronize(st_cmpt);
buf = next_buf;
}
如何验证重叠生效：
# 在程序运行时监控 rocm-smi -P # 查看功耗（计算+传输时功耗更高） rocm-smi -b # 查看 PCIe 带宽使用
效果：H2D 传输时间从 0.2s 降至几乎完全隐藏（与 kernel 执行重叠），有效减少流水线气泡。
3.6 策略六：流水线环数匹配
问题：文件读取时间（~2.3s）远大于 GPU 计算时间（~0.2s）。在双缓冲流水线中，GPU 大部分时间在等待异步读完成。
时序分析：
2 环流水线： 读[0] │████████████████████│ │ 2.3s 算[0] │ │███│ │ 0.2s 读[1] │ │████████████████████│ 2.3s（读[0]之后发起） ↑ GPU 空闲 2.1s ← HIDDEN_WAIT
理论分析：
设一个流水线周期为 $T{cycle} = \max(T{io}, T{gpu})$。对于 $R$ 个环，首张读完后，GPU 以 $T{gpu}$ 的节奏处理。只要读能在 GPU 需要之前完成，就不产生等待。
读的提前量：第 $i$ 批数据的读在 GPU 处理第 $i-R+1$ 批时发起（因为只有 $R$ 个环，读只能写在已用过的环上）。这意味着读的提前窗口为 $(R-1) \times T_{gpu}$。
无等待条件： $$(R-1) \times T{gpu} \geq T{io}$$ $$R \geq \frac{T{io}}{T{gpu}} + 1$$
不同环数的 HIDDEN_WAIT 计算：
| 环数 R | 提前窗口 | 读用时 | 每批等待 | 总等待 (5批) |
|---|---|---|---|---|
| 2 | 0.2s | 2.3s | 2.1s | 10.5s |
| 3 | 0.4s | 2.3s | 1.9s | 7.6s |
| 4 | 0.6s | 2.3s | 1.7s | 5.1s |
| 7 | 1.2s | 2.3s | 1.1s | 2.2s |
| 12 | 2.2s | 2.3s | 0.1s | 0.1s |
实践中选 4 环的考量：
-
4 环 x 915 MB = 3.6 GB pinned memory，可接受
-
7 环 x 915 MB = 6.4 GB，/dev/shm 空间可能紧张
-
12 环理论上最优，但 pinned memory 太大
代码实现（4 环异步流水线）：
// 4 环：提前 3 批启动异步读
std::future<void> rd_fut[4];
int rd_batch[4] = {-1, -1, -1, -1};
auto start_read = [&](int bi, int ri){
if(bi >= nb) return;
rd_batch[ri] = bi;
rd_fut[ri] = std::async([&, bi, ri](){
read_files(host_ring[ri], ...);
});
};
// 预读前 3 批
for(int i = 1; i < 4; i++) start_read(i, i);
for(int b = 0; b < nb; b++){
int ri = b % 4;
// 等本批数据就绪（如果还没读完）
if(b >= 1 && rd_batch[ri] >= 0) rd_fut[ri].get();
// 提前启动 b+4 批的读（用刚释放的环）
start_read(b + 4, ri);
// H2D + GPU 计算
hipMemcpy2DAsync(d_ring[g], ..., host_ring[ri], ..., st);
kernel<<<..., st>>>(d_ring[g], ...);
hipStreamSynchronize(st);
}
3.7 策略七：Pinned Memory 管理
问题：malloc 分配的主机内存，DCU 的 DMA 引擎无法直接访问，需 CPU 中转拷贝，带宽下降约 3 倍。
架构根因：DCU 的 PCIe DMA 引擎只能访问物理地址连续且固定的内存。普通 malloc 分配的页面可能被换出，且物理地址不连续，DMA 无法直接访问。
三种内存分配方式的对比：
// 1. 普通 malloc：DCU 不可直接访问 double* h = (double*)malloc(size); hipMemcpy(d, h, size, hipMemcpyHostToDevice); // → 运行时自动分配一个临时的 pinned buffer // → 从 h 拷贝到临时 buffer（CPU 中转） // → 从临时 buffer DMA 到 DCU // → 总带宽：~4 GB/s // 2. hipHostMalloc：DCU DMA 直传 double* h; hipHostMalloc(&h, size); hipMemcpy(d, h, size, hipMemcpyHostToDevice); // → DMA 直接从 h 读取 // → 总带宽：~12 GB/s // 3. hipMalloc + 寄存器缓存：最快 double* d; hipMalloc(&d, size); kernel<<<...>>>(d, ...); // 数据直接在 DCU 上分配 // → 零传输开销（如果数据在 DCU 上生成）
hipHostMalloc 注意事项：
// 分配标志 hipHostMalloc(&p, size); // 默认 hipHostMalloc(&p, size, hipHostMallocPortable); // 多进程可访问 hipHostMalloc(&p, size, hipHostMallocWriteCombined); // 写合并（读慢写快） // 检查是否真的分配到 pinned memory hipPointerAttribute_t attr; hipPointerGetAttributes(&attr, p); // attr.isManaged 为 false 表示不是 pinned
/dev/shm 管理与清理：
// 分配：每环 ~915 MB，4 环 = 3.6 GB
for(int i = 0; i < 4; i++) hipHostMalloc(&host_ring[i], bB);
// 使用完毕后及时释放
for(int i = 0; i < 4; i++) hipHostFree(host_ring[i]);
// 程序退出前兜底清理
system("rm -rf /dev/shm/* 2>/dev/null");
4. 实验与分析
4.1 实验环境
| 组件 | 配置 |
|---|---|
| DCU | 4 x 海光 DCU (gfx906), 16 GB HBM2/卡 |
| CPU | 2 x x86_64, 32 cores/node |
| 内存 | 256 GB DDR4/node |
| 互联 | InfiniBand |
| ROCm | dtk-26.04 |
| MPI | OpenMPI 4.0.2 |
| 编译 | hipcc -O3 -fopenmp -mavx2 -mfma |
4.2 各优化阶段性能
| 阶段 | 核心改动 | GPU 时间 | 总时间 | 加速比 | 关键策略 |
|---|---|---|---|---|---|
| 基线 | Top-K堆 + 2环 | ~8.0s | 14.0s | 1.0x | — |
| A | 寄存器堆 + device ring | 0.9s | 8.6s | 1.6x | 策略2 |
| B | 统计矩 (零分支) | 0.35s | 10.8s | 1.3x | 策略3 |
| C | 4环流水线 | 0.6s | 6.8s | 2.1x | 策略6 |
| D | 滑动窗口矩 | 0.08s | 4.2s | 3.3x | 策略4 |
注：阶段 B 总时间高于 A 是因为流水线未优化（2 环），HIDDEN_WAIT 抵消了 GPU 收益。阶段 C 修复了流水线。
4.3 各阶段 GPU 时间变化
GPU 时间 (s) 8.0 | ████████████████████████ | 0.9 | ███ 0.35| █ 0.08| ▏ └─────────────────────────▶ 基线 A B C+D
4.4 瓶颈转移分析
阶段 D 的时间构成 (4.2s)： ┌────────────────────────────────────────────┐ │ ████████████████████████████████ READ 1.2s │ 28% │ ████████████████████████████████████████ │ │ ████████████████████ HIDDEN_WAIT 2.5s │ 60% │ ████████████████████████████████████████ │ │ ██ GPU 0.08s │ 2% │ ██████ FINAL+WRITE 0.4s │ 10% └────────────────────────────────────────────┘
GPU 从最初的 57% 占比降至 2%，瓶颈完全转移到文件 I/O。
4.5 各策略加速因子汇总
| 策略 | 贡献加速比 | 累积加速比 (理想) |
|---|---|---|
| wavefront 对齐 | 1.2x | 1.2x |
| 寄存器级归约 | 2.0x | 2.4x |
| 消除分支分歧 | 2.6x | 6.2x |
| 加法群降维 | 4.4x | 27.3x |
| 多流并发 | 1.5x | 41.0x |
| 环数匹配 | 2.5x | 102.5x |
| pinned memory | 2.0x | 205x |
5. 总结
本文以气候态 SST 百分位计算为教学案例，系统介绍了海光 DCU 上的七条算子优化策略。核心要点：
-
DCU 不是 CUDA：wavefront 64、寄存器 64K/CU、L1 16KB——这些数字决定了你的优化方向与 NVIDIA 不同
-
分歧在 DCU 上最贵：64 线程 wavefront 使分支分歧的代价是 NVIDIA 的 2 倍
-
寄存器是命根子：善用
__launch_bounds__+ 寄存器级归约，避免 spill（spill 到 L1 只有 16KB） -
数学可以替代计算：统计矩替代精确堆（-60% 计算量），加法群降维（-83% 计算量）
-
系统优化比 kernel 优化更重要：当 GPU 时间降到 0.08s（仅占 2%），I/O 才是真瓶颈
优化路径总结：
新手优化顺序建议： 第 1 步：线程块配 64/128/256 → 免费提速 20% 第 2 步：用统计矩替代分支归约 → 提速 2-3x（DCU 上收益最大） 第 3 步：加法群降维滑动窗口 → 提速 2-4x（如有窗口运算） 第 4 步：检查流水线环数 → 看 T_io/T_gpu 比例 第 5 步：检查 pinned memory → 确认 hipMemcpy 带宽
最终 GPU kernel 从 8.0s 降到 0.08s（99% 降幅），总程序从 14.0s 降到 4.2s（3.3x 加速）。
参考文献
[1] ROCm Documentation: HIP Programming Guide. https://rocm.docs.amd.com [2] Hobday, A.J., et al. "A hierarchical approach to defining marine heatwaves." Progress in Oceanography, 141:227-238, 2016. [3] 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. [4] NVIDIA CUDA C++ Programming Guide. https://docs.nvidia.com/cuda/ [5] NetCDF Documentation. https://www.unidata.ucar.edu/software/netcdf/ [6] AMD GPU Architecture CDNA 2. Whitepaper, 2022. [7] The HDF Group. HDF5 Documentation. https://portal.hdfgroup.org
免费领 150 小时云算力，进群参与显卡、AI PC 幸运抽奖
更多推荐
26
0- 0
已为社区贡献2条内容
所有评论(0)