# 你可以学习到的内容 学完本文后可以掌握： - DCU 编程的基本流程； - HIP 程序的基本结构 - kernel、grid、block、thread 的含义； - 主机端和设备端数据如何传输； - 如何编译和运行一个简单的 DCU 程序； # 1、DCU、DTK概述 **DCU（Deep Computing Unit），深度计算单元**，是**海光信息**研发、面向数据中心的**国产 GPGPU 通用并行加速协处理器。其目前发布的产品如图1-1 所示，BW1000**采用 GPGPU 通用架构，支持全精度 AI 算力与Switch形式卡间高速互联，兼容主流 GPGPU 生态，源码平滑迁移等特性**。K100-AI 是**国产化高性能GPGPU加速卡,采用通用架构设计，支持FP32、TF32、BF16、FP16.INT8等主流AI精度。提供成熟完善的异构软件栈兼容主流GPGPU开源生态。随着云端训练、 大模型推理以及 HPC+AI 等多个领域软硬件国产化的推进，**如何使用国产异构计算芯片进行 AI 领域的相关研发和产业化显得愈发重要**。国产AI 芯片的应用场景将会越来越广泛，国产化的AI生态也将会越来越完善。  海光发布的DTK（Deep Computing Toolkit）深度计算工具栈**，底层兼容 AMD ROCm 计算生态，实现**类 CUDA 兼容。完整配套编译器、数学库、深度学习框架适配（PyTorch/TensorFlow/MindSpore 等），同时支持大模型离线训练、线上高并发推理。下表1-1 描述了 DTK 中的库与 CUDA 平台中的库的对应关系与功能介绍。  DTK计算架构建立在在DCU加速器上，其架构自底向上可以划分为基础计算层，计算执行层，计算工具层，计算服务层，异构计算语言层。**本次介绍主要围绕异构计算语言层中的显示并行语言部分，异构计算语言层也是对于刚入门异构编程的开发者比较友好的一层，接下来开始我们的第一个实践案例吧**。  # 2、环境准备 对于DCU异构编程的环境搭建有多种方式，主要有以下几种。 + 方案1：容器化方式 + 方案2：物理机方式 + 方案3：使用搭建了DCU服务器的云计算平台方式 方案1 和方案2 要求具备一台搭载了DCU加速卡和Linux操作系统的物理机。这里使用方案3 进行后续章节介绍。由于开发者社区案例生态已经非常的丰富了，如图2-1所示，这里通过海光社区官方链接给出环境搭建指南（[DCU上手指南](https://developer.sourcefind.cn/document)）。  # 3、编写第一个异构并行程序 现对程序的运行流程进行一个简单的介绍，如图3-1所示。该流程图对比展示了 CPU 串行程序与 DCU 异构并行程序的执行过程。左侧 CPU 程序主要按照“申请内存、初始化数据、逐个计算、释放内存”的顺序串行完成任务；右侧 DCU 程序需要先在主机端和设备端分别申请内存，将数据从 CPU 内存拷贝到 DCU 显存，再启动核函数进行并行计算，并通过同步、事件记录等方式统计执行时间，最后将计算结果拷贝回主机并释放相关资源。DCU 编程相比 CPU 编程多了设备内存管理、数据传输、核函数启动和同步控制等步骤，但因此能够利用大量并行线程显著提升计算效率。  ## 3.1、CPU串行程序：向量加法 ```cpp #include #include #include // 定义向量长度：一千万个int元素 #define N 10000000 int main() { // 1. 堆内存动态申请三个向量数组空间 int *A = (int *) malloc(N * sizeof(int)); int *B = (int *) malloc(N * sizeof(int)); int *C = (int *) malloc(N * sizeof(int)); // 2. 向量初始化 for (int i = 0; i < N; i++) { A[i] = 1; B[i] = 1; C[i] = 0; } // 记录向量加法开始CPU时钟 clock_t start = clock(); // 3. 核心：向量逐元素加法计算 for (int i = 0; i < N; i++) { C[i] = A[i] + B[i]; } // 记录向量加法结束CPU时钟 clock_t end = clock(); // 计算耗时：(结束时钟-开始时钟)/每秒时钟数，单位秒 double cost_time = (double)(end - start) / CLOCKS_PER_SEC * 1000; printf("CPU串行程序-向量加法计算总耗时：%.4f 毫秒\n", cost_time); // 释放堆内存，防止内存泄漏 free(A); free(B); free(C); return 0; } ``` 编译：g++ vectoradd_cpu.cpp -o vectoradd_cpu 运行 ./vectoradd_cpu  ## 3.2、DCU异构并行程序：向量加法 ```cpp #include #include // HIP 编程头文件，对应海光DCU DTK #include // 向量长度，和之前CPU串行代码保持一致 #define N 10000000 // GPU核函数：向量加法，__global__ 代表在设备端并行执行 __global__ void vector_add(int *d_A, int *d_B, int *d_C) { // 计算当前线程全局索引，对应数组下标 int idx = blockIdx.x * blockDim.x + threadIdx.x; // 边界判断，防止数组越界 if (idx < N) { d_C[idx] = d_A[idx] + d_B[idx]; } } int main() { // 1. 主机(CPU)内存申请 // 在CPU堆上分配三个向量数组 int *A = (int *) malloc(N * sizeof(int)); int *B = (int *) malloc(N * sizeof(int)); int *C = (int *) malloc(N * sizeof(int)); // 设备(GPU)端指针，初始置空 int *d_A = NULL; int *d_B = NULL; int *d_C = NULL; // 在GPU显存中分配三块向量空间 hipMalloc((void **) &d_A, N * sizeof(int)); hipMalloc((void **) &d_B, N * sizeof(int)); hipMalloc((void **) &d_C, N * sizeof(int)); // 2. CPU端向量初始化 for (int i = 0; i < N; i++) { A[i] = 1; B[i] = 1; C[i] = 0; } // 3. 数据从CPU内存拷贝至GPU显存 hipMemcpy(d_A, A, sizeof(int) * N, hipMemcpyHostToDevice); hipMemcpy(d_B, B, sizeof(int) * N, hipMemcpyHostToDevice); hipMemcpy(d_C, C, sizeof(int) * N, hipMemcpyHostToDevice); // 4. 创建GPU事件，用于统计核函数运行耗时 hipEvent_t beg, end; hipEventCreate(&beg); hipEventCreate(&end); float time_ms; // 存储计算耗时，单位毫秒 // 设置线程块、网格大小 dim3 blocksize(256, 1); // 每个线程块包含256个线程 dim3 gridsize(N / 256 + 1, 1); // 网格总块数，向上取整避免覆盖不到数据 // 同步GPU，确保前面数据拷贝全部完成 hipDeviceSynchronize(); // 记录GPU计算起始事件 hipEventRecord(beg, NULL); // 5. 启动GPU并行核函数（核心向量加法计算） vector_add<<>>(d_A, d_B, d_C); // 记录GPU计算结束事件 hipEventRecord(end, NULL); // 等待GPU计算完成 hipEventSynchronize(end); // 计算beg到end之间的执行时间，结果存入time_ms(单位ms) hipEventElapsedTime(&time_ms, beg, end); // 6. GPU计算结果拷贝回CPU内存 hipMemcpy(C, d_C, sizeof(int) * N, hipMemcpyDeviceToHost); // 7. 校验计算结果正确性 for (int i = 0; i < N; i++) { if (C[i] != 2) { std::cout << "计算结果错误 : " << C[i] << std::endl; } } printf("计算结果正确, DCU异构并行程序-向量加法计算总耗时：%.4f 毫秒\n", time_ms); // 8. 释放所有申请的资源 // 释放CPU堆内存 free(A); free(B); free(C); // 释放GPU显存 hipFree(d_A); hipFree(d_B); hipFree(d_C); // 销毁计时事件对象 hipEventDestroy(beg); hipEventDestroy(end); return 0; } ``` 编译：hipcc vectoradd_dcu.cpp -o vectoradd_dcu 运行：./vectoradd_dcu  ## 3.3、对比分析 向量长度 N=10000000（一千万个 int 元素） + CPU 串行纯加法耗时：81.6060 ms + DCU 并行核函数纯计算耗时：1.7859 ms 异构编程核心收益来自 DCU 海量轻量级线程并发计算，非常适合大数据并行数值运算。 # 4、异构并行编程模型 在第3节中编写了一个简单的DCU异构并行程序，其中使用到了HIP模型。那么接下来介绍什么是HIP。HIP (Heterogeneous-compute Interface for Portability)是一种显式异构并行编程模型，是对C/C++语言扩展的异构并行编程模型。使用HIP可以实现异构加速卡程序编写。在CPU上运行的部分称为主机端（host），在加速卡上运行的部分称为设备端（device）。主机端代码使用HIP Runtime API来完成对 设备端的状态控制，并与之交互。这些操作通常为设备端存储控制，双端数据传输，设备端核函数加载和执行等。 HIP 基于标准 C/C++ 做了三类专用扩展，用于异构并行编程，分别是核函数修饰符、线程层次结构内置变量、设备存储限定符。 ## 4.1、核函数 核函数区分主机 CPU 函数**和**设备 DCU 并行函数**，定义能被 CPU 调度、在 DCU 上千线程并行执行的计算函数，通过专属关键字扩展标准 C/C++。核函数返回类型必须为 void，表 4-1 中总结了 HIP 程序中的**函数类型限定符。函数类型限定符指定一个函数在主机上执行还是在设备上执行以及可被主机调用还是被设备调用  CPU 调用核函数的特有语法，用于向硬件分配并行线程资源。例如：\_\_global\_\_ void vector_add(int *d_A, int *d_B, int *d_C)** ## 4.2、**线程层次结构** 标准 C/C++ 没有线程网格、线程块的硬件分层概念，HIP 新增 **全局内置只读变量**，供核函数内获取当前线程的坐标、分层大小，用来划分数组计算下标。 + Grid（网格）：由多个 Block 组成，整个核函数的总任务容器 + Block（线程块）：网格内的基础分组，一个 Block 内的线程可同步、共享共享内存 + Thread（线程）：最小并行执行单元 grid、block、thread 是 kernel 启动时组织并行任务的三层结构。一次 kernel 启动会产生一个 grid，一个 grid 里面有很多 block，一个 block 里面有很多 thread。thread 是 kernel 中真正执行代码的单位。比如这个vector_add kernel。每一个 thread 都会执行一遍这段代码，但是**不同 thread** 。 ```cpp __global__ void vector_add(int *d_A, int *d_B, int *d_C) { // 计算当前线程全局索引，对应数组下标 int idx = blockIdx.x * blockDim.x + threadIdx.x; // 边界判断，防止数组越界 if (idx < N) { d_C[idx] = d_A[idx] + d_B[idx]; } } ``` 1D线程网格实例如图4-1所示。  在计算线程在全局中的位置时可使用：int idx = blockDim.x * blockIdx.x + threadIdx.x 2D线程网格实例如图4-2所示  不同的颜色代表不同block，每个小方块代表thread，每个block和thread都有2维的索引，当需要计算线程在全局中的位置时，需要计算x和y两个维度信息：int idx_x = blockDim.x * blockIdx.x + threadIdx.x 和 int idx_y = blockDim.y * blockIdx.y + threadIdx.y ## 4.3、存储结构 对于开发人员来说，一般有两种类型的存储器。可编程的：需要显式地控制哪些数据存放在可编程内存中，例如全局内存和共享内存。 不可编程的：程序将自动生成存放位置，例如 L1 和 L2。存储结构分布如图4-3所示。  **全局内存**是 DCU中容量最大、延迟最高并且最常使用的内存。 全局内存常驻于 DCU 内存中，可通过内存事务进行访问，这些内存事务最好是 4 字节对齐，这样可以利用全局内存合并访存的特性来减少对全局内存的访问次数 **共享内存**在 DCU 上，每个 CU 都有一块共享内存，大小为 64KB，其带宽和延迟远远优于全局内存。共享内存在核函数的范围内声明，是线程之间相互通信的基本方式。其生命周期伴随着整个线程块，一个块内的线程通过使用共享内存中的数据可以相互合作。 # 5、参考文献 [DCU异构编程与优化--基础](https://developer.sourcefind.cn/study/courses/coursesneiye?post_id=7953) [DTK-光合开发者社区](https://developer.sourcefind.cn/dtk)