1 概述
1.1 处理器的异构化发展趋势
计算需求飞速增长，在DCU平台上，HIP是一个类似CUDA的显式异构编程模型
2 DCU系统软硬件架构
2.1 详解DCU架构
2.1.1 DCU整体硬件架构
DCU通过PCI-E总线与CPU处理器相连
如图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是满载状态。
- 空间管理
与CPU内存管理类似，通过调用hipMalloc这样函数接口就可以完成DCU内存分配，hipMalloc是最基础的内存分配接口
- 空间管理
-
- 数据拷贝
完成空间申请后，下面一步就需要将CPU端准备好的数据传输给DCU显存。
使用hipMemcpy完成这一过程，需要注意的是hipMemcpy具有方向性，最后一个参数hipMemcpyHostToDevice，指示了方向，代表这一拷贝过程是从主机端到设备端，同样的hipMemcpyDeviceToHost指示的是从设备端到主机端的传输。
- 数据拷贝
- 函数定义
__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-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 线程执行模型
- 线程束和线程块
线程块的大小是通过三尖括号的第二个参数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.2.2 可编程存储结构
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];
