## 一、前置环境 ### 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** 管理系统资源。关键调度参数： ```bash #SBATCH -p kshdmcc2026 # 指定计算分区 #SBATCH -N 2 # 申请 2 个节点 #SBATCH -n 2 # 2 个任务（每节点 1 个） #SBATCH --gres=dcu:4 # 每节点申请 4 张 DCU 卡 #SBATCH --exclusive # 节点独占使用 ``` --- ## 二、申请 DCU 核心的基本指令 ### 2.1 查看可用 DCU 设备 ```bash # 查看 DCU 设备信息（类似 nvidia-smi） rocminfo # 查看设备数量与型号 rocminfo | grep -E "Name|Agent" | head -10 ``` 在 HIP 程序中查看可用设备数量： ```cpp #include #include 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; } ``` 编译与运行： ```bash hipcc -o device_query device_query.cpp ./device_query ``` ### 2.2 通过 SLURM 提交 DCU 任务 **方式一：sbatch 脚本提交** ```bash #!/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`，提交： ```bash sbatch run.sh ``` **方式二：srun 交互式提交** ```bash # 申请 1 个节点 × 2 张 DCU 卡，交互式运行 srun -p kshdmcc2026 -N 1 --gres=dcu:2 --pty bash # 进入节点后直接运行程序 ./my_dcu_program ``` **方式三：srun 直接运行** ```bash srun -N 1 -n 1 --gres=dcu:1 ./my_dcu_program ``` ### 2.3 编译选项详解 ```bash 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 版本。可以通过脚本自动探测： ```bash 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 错误检查宏 ```cpp #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 完整代码 ```cpp #include #include #include #include #include #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<<>>(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(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 编译与运行 ```bash # 编译（指定 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）并发，实现数据传输与计算的流水线重叠： ```cpp hipStream_t s0, s1; hipStreamCreate(&s0); hipStreamCreate(&s1); // 流 0：前半数据 H2D → Kernel → D2H hipMemcpyAsync(d_pool, h_pool, half_bytes, hipMemcpyHostToDevice, s0); kernel<<>>(...); 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<<>>(...); hipMemcpyAsync(h_result + offset, d_result + offset, half_bytes, hipMemcpyDeviceToHost, s1); hipStreamSynchronize(s0); hipStreamSynchronize(s1); ``` ### 5.2 hipHostRegister 固定内存 将 host 内存页锁定（pinned memory），允许 GPU DMA 引擎直接访问，减少一次 CPU 端拷贝： ```cpp 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 性能分析 ```bash # 使用 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 基础认知与核心优势](./01-海光DCU基础认知与核心优势.md) > > 下一篇：[海光 DCU 进阶：高阶实操与性能优化](./03-海光DCU进阶-高阶实操与性能优化.md)