DCU 开发与使用文档

2023-02-14  本文已影响0人  mope

1 概述

1.1 处理器的异构化发展趋势


计算需求飞速增长，在DCU平台上，HIP是一个类似CUDA的显式异构编程模型

2 DCU系统软硬件架构

2.1 详解DCU架构

2.1.1 DCU整体硬件架构

DCU通过PCI-E总线与CPU处理器相连


2-1

如图2-1显示了DCU是个相对完整的系统，由以下几个关键模块组成：
计算单元阵列，如图CU0、CU1等
缓存系统（L1一级缓存，L2二级缓存）
全局内存(global memory)
CPU和DCU数据通路（DMA）

2.1.2 DCU核心架构介绍

3 DCU编程方法

3.1 编写第一个DCU程序

3.1.1 HIP编程实战-数组相加

两数组相加，也即数组A与数组B进行相加，结果赋值给数组C
首先，我们看下CPU平台C语言版数组相加的代码，是一个for循环的简单程序。

#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];
    }
    free(A);
    free(B);
    free(C);
    return 0;
}

malloc内的参数是需要动态分配的字节数

接下来，我们看下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);
}

对DCU的程序需要采用专用的编译器hipcc进行编译

hipcc vector_add_dcu.cpp -o vector_add_dcu

rocm-smi命令:用简单方法确定数组相加这个过程是在DCU上执行了,rocm-smi命令可以查看DCU负载情况.
第一列是DCU的序号，如果你的电脑里有不止一块DCU就会有多行，%DCU列代表了DCU的实时负载，可以看到在运行程序时候第一块DCU是满载状态。


3-1 
__global__ void add(float * d_A,float * d_B,float * d_C)

global是一个函数标识符，代表了定义的函数是在主机端被调用，但是执行是在设备端，这个标识符会被hipcc编译器识别到，然后会按照DCU端指令去翻译下面的代码段。
  被global标识的函数又被称为核函数或者是"kernel"，核函数代表了设备端的入口，一旦进入核函数，函数的执行和控制就交给了设备端。
进入到核函数体中，首先使用threadIdx.x，blockIdx.x ，blockDim.x来定位当前的线程编号，根据线程编号来算出当前线程读取数据的偏移地址，接着再往下需要判断线程索引的数据是否超越数据边界，如果没有超过，运行核函数真正的计算逻辑，进行数据相加。

dim3 blocksize(256,1);
dim3 gridsize(N/256+1,1);
// 进行数组相加
add<<<gridsize, blocksize >>>(d_A,d_B,d_C);

3.1.2 总结

3-2

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

3.2 HIP核函数

3.2.1 线程执行模型

add<<<gridsize, blocksize >>>(d_A,d_B,d_C);

如果需要使用多维线程组织模型，对应于代码中在设置gridsize和blocksize时就需要配置为dim3类型。在下面的例子中可以看到，对于myGpuMatrix这个核函数，我们分配了网格尺寸是dimGrid(90,90),其描述的网格上x轴上线程块90个，y轴线程块90个。每个线程块尺寸是dimBlock(64,1)。

dim3 dimGrid(90,90);
dim3 dimBlock(64,1);
hipLaunchKernelGGL(myGpuMatrix,dimGrid,dimBlock,0,0,dev_A,dev_B,dev_C_t);
3-3

3.2.2 可编程存储结构

3-7
3.2.2.1 存储结构简介
3.2.2.2 使用共享内存

在HIP编程模型中，为了获得高内存带宽，使用共享内存是一种常用的手段。共享内存被分为32个同样大小的内存模型，被称为存储体（Bank），可以被同时访问。
可以把存储体形象的比喻为一串可以穿无数个山楂的糖葫芦串，每一个穿山楂的签子就对应了一个内存通道称为Bank，每个山楂就是存储空间以4字节为单位。
访问共享内存，有串行访问和广播访问两种典型的模式。串行访问是多个线程访问同一个Bank的不同地址。
在编写核函数中使用共享内存可以减少对全局内存的访问。
两种共享内存静态声明如下：

__shared__ int tile[N];
__shared__ int tile[Row][Col];
上一篇 下一篇

猜你喜欢

热点阅读