## 摘要 **海光 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<<>>(d_ptr); // HIP 版本（只有前缀不同） hipMalloc(&d_ptr, size); hipMemcpy(d_ptr, h_ptr, size, hipMemcpyHostToDevice); kernel<<>>(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<<>>(...); ``` **正确做法**： ``` int thr = 256;  // 256 / 64 = 4 个完整 wavefront int blk = (N + thr - 1) / thr; // blk × thr 可能略多于 N，多出的线程通过 if(idx >= N) return; 处理 kernel<<>>(...); ``` **速查表**： | **线程块大小** | **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 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 上的七条算子优化策略。核心要点：** 1. **DCU 不是 CUDA**：wavefront 64、寄存器 64K/CU、L1 16KB——这些数字决定了你的优化方向与 NVIDIA 不同 2. **分歧在 DCU 上最贵**：64 线程 wavefront 使分支分歧的代价是 NVIDIA 的 2 倍 3. **寄存器是命根子**：善用 `__launch_bounds__` + 寄存器级归约，避免 spill（spill 到 L1 只有 16KB） 4. **数学可以替代计算**：统计矩替代精确堆（-60% 计算量），加法群降维（-83% 计算量） 5. **系统优化比 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](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/](https://docs.nvidia.com/cuda/) **[5] NetCDF Documentation. **[https://www.unidata.ucar.edu/software/netcdf/](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](https://portal.hdfgroup.org/)