万字学习——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的异构计算的应用程序。

相关数学库

优化和调试工具

第一个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;
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29

运行

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);
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50

运行

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

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

DCU程序组成

HIP主要API释义

HIP核函数

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

HIP开发执行

HIP设备管理

单进程多CPU编程

HIP性能分析

DCU程序优化