海光 DCU 实战：基本指令与操作

一、前置环境

1.1 硬件要求

• 海光 DCU 加速卡（如 Z100L / K100 系列）
• 或配备 DCU 的异构计算节点

1.2 软件栈

DCU 的软件栈称为 DTK（Deep Computing Toolkit），核心组件包括：

组件	说明	类比（NVIDIA）
hipcc	HIP C/C++ 编译器（基于 Clang/LLVM）	nvcc
HIP Runtime	运行时 API，管理设备与内存	CUDA Runtime
rocBLAS	高性能 BLAS 库	cuBLAS
MIOpen	深度学习算子库	cuDNN
rocprof	硬件性能分析器	nvprof / Nsight
rocminfo	查看 DCU 设备信息	nvidia-smi

1.3 作业调度系统

DCU 集群通常通过 SLURM 管理系统资源。关键调度参数：

#SBATCH -p kshdmcc2026       # 指定计算分区
#SBATCH -N 2                 # 申请 2 个节点
#SBATCH -n 2                 # 2 个任务（每节点 1 个）
#SBATCH --gres=dcu:4         # 每节点申请 4 张 DCU 卡
#SBATCH --exclusive          # 节点独占使用

---

二、申请 DCU 核心的基本指令

2.1 查看可用 DCU 设备

# 查看 DCU 设备信息（类似 nvidia-smi）
rocminfo

# 查看设备数量与型号
rocminfo | grep -E "Name|Agent" | head -10

在 HIP 程序中查看可用设备数量：

#include <hip/hip_runtime.h>
#include <cstdio>

int main() {
    int ndevices;
    hipGetDeviceCount(&ndevices);
    printf("可用 DCU 设备数: %d\n", ndevices);

    for (int i = 0; i < ndevices; i++) {
        hipDeviceProp_t prop;
        hipGetDeviceProperties(&prop, i);
        printf("设备 %d: %s\n", i, prop.name);
        printf("  显存: %.2f GB\n", prop.totalGlobalMem / (1024.0 * 1024.0 * 1024.0));
        printf("  计算单元: %d\n", prop.multiProcessorCount);
    }
    return 0;
}

编译与运行：

hipcc -o device_query device_query.cpp
./device_query

2.2 通过 SLURM 提交 DCU 任务

方式一：sbatch 脚本提交

#!/bin/bash
#SBATCH -p kshdmcc2026
#SBATCH -N 1
#SBATCH -n 1
#SBATCH --gres=dcu:1
#SBATCH -J my_dcu_job

./my_dcu_program

保存为 run.sh，提交：

sbatch run.sh

方式二：srun 交互式提交

# 申请 1 个节点 × 2 张 DCU 卡，交互式运行
srun -p kshdmcc2026 -N 1 --gres=dcu:2 --pty bash

# 进入节点后直接运行程序
./my_dcu_program

方式三：srun 直接运行

srun -N 1 -n 1 --gres=dcu:1 ./my_dcu_program

2.3 编译选项详解

hipcc -std=c++14 -O3 -march=native \
    --offload-arch=gfx906 \          # ⬅ 关键：指定 DCU 架构
    -o output_binary \
    source.cpp

选项	含义
--offload-arch=gfx906	指定面向 DCU Z100L 架构编译
-O3	最高级别代码优化
-march=native	针对当前 CPU 架构优化宿主代码
-std=c++14	C++ 标准版本

2.4 DTK 多版本兼容探测

在实际集群中，不同节点可能安装了不同的 DTK 版本。可以通过脚本自动探测：

DTK_VER=""
for try_ver in "dtk-25.04.4" "dtk-24.04.1" "dtk-22.10.1"; do
    candidate="/public/software/compiler/rocm/${try_ver}/hip/bin/hipcc"
    if [ -f "$candidate" ]; then
        DTK_VER="$try_ver"
        break
    fi
done

ROCM_PATH="/public/software/compiler/rocm/${DTK_VER}"
HIPCC="${ROCM_PATH}/hip/bin/hipcc"
export LD_LIBRARY_PATH="${ROCM_PATH}/lib:${ROCM_PATH}/hip/lib:$LD_LIBRARY_PATH"

---

三、DCU 编程模型入门

3.1 基本概念对照表

HIP (DCU)	CUDA (NVIDIA)	含义
__global__	__global__	内核函数，从 host 调用，在 device 执行
__device__	__device__	设备端函数，仅可在 device 调用
hipMalloc	cudaMalloc	在 device 上分配显存
hipMemcpy	cudaMemcpy	host ↔ device 数据传输
hipFree	cudaFree	释放 device 显存
hipLaunchKernelGGL	<<< >>>	启动 kernel
threadIdx.x	threadIdx.x	线程在线程块内的索引
blockIdx.x	blockIdx.x	线程块在网格中的索引
blockDim.x	blockDim.x	每块的线程数
hipStream_t	cudaStream_t	异步执行流

3.2 HIP 错误检查宏

#define HIP_CHECK(cmd) do { \
    hipError_t err = cmd; \
    if (err != hipSuccess) { \
        fprintf(stderr, "HIP error at %s:%d: %s (%d)\n", \
                __FILE__, __LINE__, hipGetErrorString(err), err); \
        exit(1); \
    } \
} while(0)

#define HIP_KERNEL_CHECK() do { \
    hipError_t err = hipGetLastError(); \
    if (err != hipSuccess) { \
        fprintf(stderr, "HIP kernel error at %s:%d: %s (%d)\n", \
                __FILE__, __LINE__, hipGetErrorString(err), err); \
        exit(1); \
    } \
} while(0)

---

四、实战示例：向量加法

4.1 完整代码

#include <hip/hip_runtime.h>
#include <cstdio>
#include <cstdlib>
#include <cmath>
#include <chrono>

#define HIP_CHECK(cmd) do { \
    hipError_t err = cmd; \
    if (err != hipSuccess) { \
        fprintf(stderr, "HIP error at %s:%d: %s (%d)\n", \
                __FILE__, __LINE__, hipGetErrorString(err), err); \
        exit(1); \
    } \
} while(0)

// ── GPU Kernel：向量加法 ──────────────────────────────────
__global__ void vector_add(const float *__restrict__ A,
                           const float *__restrict__ B,
                           float       *__restrict__ C,
                           int N)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

int main() {
    const int N = 1 << 24;  // 16M 元素 ≈ 64 MB × 3 = 192 MB
    const size_t bytes = N * sizeof(float);

    // ── CPU 端分配 ──────────────────────────────────────
    float *h_A = new float[N];
    float *h_B = new float[N];
    float *h_C = new float[N];

    for (int i = 0; i < N; i++) {
        h_A[i] = (float)i;
        h_B[i] = (float)(i * 2);
    }

    // ── DCU 端分配 ──────────────────────────────────────
    float *d_A, *d_B, *d_C;
    HIP_CHECK(hipMalloc(&d_A, bytes));
    HIP_CHECK(hipMalloc(&d_B, bytes));
    HIP_CHECK(hipMalloc(&d_C, bytes));

    // ── H2D 数据传输 ────────────────────────────────────
    HIP_CHECK(hipMemcpy(d_A, h_A, bytes, hipMemcpyHostToDevice));
    HIP_CHECK(hipMemcpy(d_B, h_B, bytes, hipMemcpyHostToDevice));

    // ── 启动 Kernel ─────────────────────────────────────
    const int block_size = 256;
    const int grid_size  = (N + block_size - 1) / block_size;

    auto t0 = std::chrono::steady_clock::now();

    vector_add<<<dim3(grid_size), dim3(block_size), 0, 0>>>(d_A, d_B, d_C, N);
    HIP_KERNEL_CHECK();
    HIP_CHECK(hipDeviceSynchronize());

    auto t1 = std::chrono::steady_clock::now();
    double ms = std::chrono::duration<double, std::milli>(t1 - t0).count();

    // ── D2H 数据回传 ────────────────────────────────────
    HIP_CHECK(hipMemcpy(h_C, d_C, bytes, hipMemcpyDeviceToHost));

    // ── 验证结果 ────────────────────────────────────────
    int errors = 0;
    for (int i = 0; i < N; i++) {
        float expected = h_A[i] + h_B[i];
        if (fabsf(h_C[i] - expected) > 1e-5f) {
            if (++errors <= 5) {
                printf("Mismatch at %d: got %.2f, expected %.2f\n",
                       i, h_C[i], expected);
            }
        }
    }

    printf("N = %d (%.1fM)\n", N, N / 1e6);
    printf("Kernel time: %.3f ms\n", ms);
    printf("Bandwidth:   %.2f GB/s\n",
           (3.0 * bytes) / (ms / 1000.0) / 1e9);
    printf("Errors: %d\n", errors);

    // ── 清理 ────────────────────────────────────────────
    HIP_CHECK(hipFree(d_A));
    HIP_CHECK(hipFree(d_B));
    HIP_CHECK(hipFree(d_C));
    delete[] h_A;
    delete[] h_B;
    delete[] h_C;

    return errors == 0 ? 0 : 1;
}

4.2 编译与运行

# 编译（指定 DCU gfx906 架构）
hipcc -std=c++14 -O3 --offload-arch=gfx906 -o vec_add vec_add.cpp

# 直接运行
./vec_add

# 或提交到 SLURM 队列
srun -N 1 -n 1 --gres=dcu:1 ./vec_add

4.3 输出示例

N = 16777216 (16.8M)
Kernel time: 0.856 ms
Bandwidth:   225.32 GB/s
Errors: 0

---

五、进阶概念速览

5.1 HIP Stream 异步执行

DCU 支持多流（Stream）并发，实现数据传输与计算的流水线重叠：

hipStream_t s0, s1;
hipStreamCreate(&s0);
hipStreamCreate(&s1);

// 流 0：前半数据 H2D → Kernel → D2H
hipMemcpyAsync(d_pool, h_pool, half_bytes, hipMemcpyHostToDevice, s0);
kernel<<<grid0, block, 0, s0>>>(...);
hipMemcpyAsync(h_result, d_result, half_bytes, hipMemcpyDeviceToHost, s0);

// 流 1：后半数据 H2D → Kernel → D2H（与流 0 并发）
hipMemcpyAsync(d_pool + offset, h_pool + offset, half_bytes, hipMemcpyHostToDevice, s1);
kernel<<<grid1, block, 0, s1>>>(...);
hipMemcpyAsync(h_result + offset, d_result + offset, half_bytes, hipMemcpyDeviceToHost, s1);

hipStreamSynchronize(s0);
hipStreamSynchronize(s1);

5.2 hipHostRegister 固定内存

将 host 内存页锁定（pinned memory），允许 GPU DMA 引擎直接访问，减少一次 CPU 端拷贝：

short *pool = (short*)mmap(nullptr, pool_bytes, PROT_READ | PROT_WRITE,
                            MAP_PRIVATE | MAP_ANONYMOUS, -1, 0);

// 锁定内存页，GPU 可直接 DMA 读取
hipHostRegister(pool, pool_bytes, hipHostRegisterDefault);

// ... GPU 计算 ...

hipHostUnregister(pool);
munmap(pool, pool_bytes);

5.3 性能分析

# 使用 rocprof 采集硬件计数器
rocprof --stats ./your_dcu_program

# 输出包括：
#   MemUnitBusy    - 内存单元占用率
#   VALUUtilization - 向量 ALU 利用率
#   L2CacheHit     - L2 缓存命中率
#   FETCH_SIZE     - 从 HBM 读取数据量

---

六、小结

本文覆盖了 DCU 编程的三个核心环节：

1. 资源申请：通过 SLURM 的 --gres=dcu:N 申请 DCU 卡
2. 编译：使用 hipcc --offload-arch=gfx906 编译 HIP 程序
3. 编程模式：内存分配（hipMalloc）→ 数据传输（hipMemcpy）→ 内核启动（kernel<<<>>>）→ 结果回传 → 释放资源

从向量加法的简单例子出发，你已经掌握了 DCU 编程的基本骨架。接下来可以进入更复杂的实战场景。

---

上一篇：海光 DCU 基础认知与核心优势

下一篇：海光 DCU 进阶：高阶实操与性能优化