万字学习——DCU编程实战补充

参考资料

2.1 详解DCU架构 · DCU 开发与使用文档 (hpccube.com)

DCU架构是什么样的

• 计算单元阵列，如图CU0、CU1等
• 缓存系统（L1一级缓存，L2二级缓存）
• 全局内存(global memory)
• CPU和DCU数据通路（DMA）

我的理解大概是这样的

DCU节点结构

常见的异构计算节点体系结构主要由四个部分组成：主存、多核处理器、I/O Hub和DCU加速器。这种结构在计算机体系结构中被定义为NUMA。

DCU加速器根据其主要功能可以划分为四个主要组件：执行引擎（Execution Engine），一个或多个DMA拷贝引擎（Copy Engine），内存控制器（Memory Controller）和DCU显存（DCU Memory）。

DCU软件栈-HIP

DCU拥有自己的软件栈–HIP软件栈，也叫生态系统或软件层，用来支持基于HIP的异构计算的应用程序。

相关数学库

HIP数学库	CUDA数学库	数学库功能
hipblas	cublas	基础矩阵运算数学库
hiprand	curand	随机数数学库
hipsparse	cusparse	稀疏矩阵数学库
hipfft	cufft	快速傅立叶变换数学库
miopen	cudnn	深度学习基础数学库
hipcub	cub	基础算法库
RCCL	NCCL	通信库
rocThrust	Thrust	并行算法模板库

优化和调试工具

工具名称	功能
rocprofiler	用于程序分析和绘制时间线
roctracer	用于跟踪程序

第一个DCU程序-数组相加

CPU平台C语言版

#include <stdio.h>
#include <stdlib.h>
#define N 10000
int main() {
    //申请数据空间
    float *A = (float *) malloc(N * sizeof(float));
    float *B = (float *) malloc(N * sizeof(float));
    float *C = (float *) malloc(N * sizeof(float));
    //数据初始化
    for (int i = 0; i < N; i++) {
        A[i] = 1;
        B[i] = 1;
        C[i] = 0;
    }
    // 进行数组相加
    for (int i = 0; i < N; i++) {
        C[i] = A[i] + B[i];
    }
    
    printf("%f\n", *A);
    printf("%f\n", *B);
    printf("%f\n", *C);

    //释放数据空间
    free(A);
    free(B);
    free(C);
    return 0;
}

运行

DCU版本

#include <iostream>
#include "hip/hip_runtime.h"
#include <hip/hip_runtime.h>

#define N 10000

__global__ void add(float *d_A, float *d_B, float *d_C) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < N) {
        d_C[tid] = d_A[tid] + d_B[tid];
    }
}

int main() {
    //申请数据空间
    float *A = (float *) malloc(N * sizeof(float));
    float *B = (float *) malloc(N * sizeof(float));
    float *C = (float *) malloc(N * sizeof(float));
    float *d_A = NULL;
    float *d_B = NULL;
    float *d_C = NULL;
    hipMalloc((void **) &d_A, N * sizeof(float));
    hipMalloc((void **) &d_B, N * sizeof(float));
    hipMalloc((void **) &d_C, N * sizeof(float));
    //数据初始化
    for (int i = 0; i < N; i++) {
        A[i] = 1;
        B[i] = 1;
        C[i] = 0;
    }
    hipMemcpy(d_A, A, sizeof(float) * N, hipMemcpyHostToDevice);
    hipMemcpy(d_B, B, sizeof(float) * N, hipMemcpyHostToDevice);
    hipMemcpy(d_C, C, sizeof(float) * N, hipMemcpyHostToDevice);
    dim3 blocksize(256, 1);
    dim3 gridsize(N / 256 + 1, 1);
    // 进行数组相加
    add<<<gridsize, blocksize >>> (d_A, d_B, d_C);
    //结果验证
    hipMemcpy(C, d_C, sizeof(float) * N, hipMemcpyDeviceToHost);
    for (int i = 0; i < N; i++) {
        std::cout << C[i] << std::endl;
    }
    //释放申请空间
    free(A);
    free(B);
    free(C);
    hipFree(d_A);
    hipFree(d_B);
    hipFree(d_C);
}

运行

hipcc vector-DCU.cpp -o vector-DCU
./vector-DCU

rocm-smi命令可以查看DCU负载情况

DCU程序组成

HIP主要API释义

API名称	含义
hipGetDeviceCount	获取机器上的设备个数
hipGetDeviceProperties	获取选定设备的设备属性
hipMalloc	申请DCU内存
hipHostMalloc	在CPU端申请页锁定内存
hipStreamCreate	创建流
hipMemcpyAsync	CPU和DCU内存异步拷贝，拷贝有两个方向，CPU到DCU，DCU到CPU
hipMemcpy	CPU和DCU内存同步拷贝，会造成CPU端程序暂停等待拷贝的完成才会继续下面的指令，同上拷贝有两个方向
hipFree	释放DCU端的内存

HIP核函数

HIP全局内存管理与数据传输

HIP开发执行

HIP设备管理

单进程多CPU编程

HIP性能分析

DCU程序优化