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

admin2024-07-10  2

参考资料

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

DCU架构是什么样的

万字学习——DCU编程实战补充,image-20240706142835011,第1张

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

万字学习——DCU编程实战补充,image-20240706143119392,第2张

我的理解大概是这样的

万字学习——DCU编程实战补充,image-20240706145851060,第3张

DCU节点结构

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

万字学习——DCU编程实战补充,image-20240706150247449,第4张

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

万字学习——DCU编程实战补充,image-20240706150330616,第5张

DCU软件栈-HIP

DCU拥有自己的软件栈–HIP软件栈,也叫生态系统或软件层,用来支持基于HIP的异构计算的应用程序。万字学习——DCU编程实战补充,image-20240706150500537,第6张

相关数学库

HIP数学库CUDA数学库数学库功能
hipblascublas基础矩阵运算数学库
hiprandcurand随机数数学库
hipsparsecusparse稀疏矩阵数学库
hipfftcufft快速傅立叶变换数学库
miopencudnn深度学习基础数学库
hipcubcub基础算法库
RCCLNCCL通信库
rocThrustThrust并行算法模板库

优化和调试工具

工具名称功能
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编程实战补充,image-20240706152759044,第7张

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

万字学习——DCU编程实战补充,image-20240707085756455,第8张

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

万字学习——DCU编程实战补充,image-20240707090016671,第9张

DCU程序组成
万字学习——DCU编程实战补充,image-20240707091544599,第10张

HIP主要API释义

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

万字学习——DCU编程实战补充,a60fbc68e524516776f62f278e8bc25,第11张

HIP核函数

万字学习——DCU编程实战补充,f99e0b3ed236f186627e9d00114e513,第12张

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

万字学习——DCU编程实战补充,image-20240710091030117,第13张

HIP开发执行

万字学习——DCU编程实战补充,image-20240710091113231,第14张

HIP设备管理

万字学习——DCU编程实战补充,image-20240710091134872,第15张

单进程多CPU编程

万字学习——DCU编程实战补充,第16张

HIP性能分析

万字学习——DCU编程实战补充,image-20240710091240966,第17张

DCU程序优化

万字学习——DCU编程实战补充,image-20240710091337734,第18张

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明原文出处。如若内容造成侵权/违法违规/事实不符,请联系SD编程学习网:675289112@qq.com进行投诉反馈,一经查实,立即删除!