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

• 1. 空间管理
     与CPU内存管理类似，通过调用hipMalloc这样函数接口就可以完成DCU内存分配，hipMalloc是最基础的内存分配接口
• 2. 数据拷贝
     完成空间申请后，下面一步就需要将CPU端准备好的数据传输给DCU显存。

  使用hipMemcpy完成这一过程，需要注意的是hipMemcpy具有方向性，最后一个参数hipMemcpyHostToDevice，指示了方向，代表这一拷贝过程是从主机端到设备端，同样的hipMemcpyDeviceToHost指示的是从设备端到主机端的传输。

• 3. 函数定义

__global__ void add(float * d_A,float * d_B,float * d_C)

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

• 4. 函数执行

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 线程执行模型

• 1. 线程束和线程块
     线程块的大小是通过三尖括号的第二个参数blocksize来配置的，三尖括号的第一个参数指的是一共启动多少个线程块。例如gridsize为1，blocksize为256，带入下面的代码段，翻译过来就是一共启动1*256个线程来执行add的代码。

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 存储结构简介

• 寄存器
  寄存器是DCU上访问速度最快的内存空间。如果核函数中声明的自变量没有其他修饰符，那么它通常存储在寄存器中。

• 全局内存
  全局内存是DCU中容量最大、延迟最高并且最常使用的内存。

• 共享内存
  共享内存是HIP编程模型中提供的可编程缓存。

• 本地内存
  本地内存，位于全局内存上的一个空间。在逻辑上，它属于线程私有的。

3.2.2.2 使用共享内存

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

__shared__ int tile[N];
__shared__ int tile[Row][Col];