曙光超算Fortran混编C并启动DCU计算

简介

将在Fortran程序中初始化的矩阵数组传递给C程序，由C程序启动DCU，每个线程根据自己的线程ID号定位到矩阵的对应行上，将该行每个元素的值加上该线程ID号。

特别注意与错误定位

（1）Fortran数组为列优先存储，两层循环初始化数组时注意外层循环变量在矩阵索引的第二维度上。否则可能出现内存读写错误。

（2）传递到C程序的参数皆为指针，即在C中要以指针类型接收数据。

（3）C数组为行优先存储，在传递之后从软件层来看矩阵自动进行了转置，然而在内存中的存储位置并没有发生任何改变，因此只需将矩阵的行列号交换即可获得转置后的矩阵。

（4）参与混编的C函数必须是小写（Fortran不区分大小写但C区分），函数名后加‘_’（下划线），函数名前用extern “C”声明。否则编译时找不到该混编函数。

（5）调用DCU需要指定编译指令，编译指令与所在环境息息相关。编译指令与编译程序时所在的环境与运行程序时所在的环境需统一。

（6）目前DCU核函数内不支持C++库（以后可能会支持），也不支持malloc动态内存分配。DCU所需要的内存需要在主机端使用hipMalloc等提前开辟，再将显存指针传入DCU中。否则报DCU显存页不存在或不可用错误。

Fortran源代码

文件名：fMain.F90

program mainimplicit none! 进行混合编译的c函数名(computefordcu)必须是小写，原因是Fortran语言不区分大小写，而混合编译时对其是区分大小写的external :: computefordcu! 测试矩阵，对矩阵进行每一列的元素加上该列号的操作，Fortran存储为列优先real*8, allocatable, dimension(:,:) :: matrix! 测试矩阵的行数，列数integer :: row, col! 测试结果real*8, allocatable, dimension(:) :: relvectorinteger :: i, jrow = 10col = 8! 为测试矩阵分配内存allocate(matrix(row, col))! 为测试结果分配内存allocate(relvector(col))! 初始化测试矩阵，Fortran数组索引从1开始，同时计算测试结果即每列元素之和do i = 1, colrelvector(i) = 0do j = 1, row! 列优先初始化，颠倒顺序极易出错matrix(j, i) = jrelvector(i) = relvector(i) + matrix(j, i)end do! 打印测试结果print *, "Sum of matrix column(", i, ")elements :", relvector(i)end do! C数组的存储为行优先，因此传入的矩阵会自动转置，所以在此颠倒行列数call computefordcu(matrix, col, row)! 打印、验证结果do i = 1, colrelvector(i) = 0do j = 1, rowrelvector(i) = relvector(i) + matrix(j, i)end do! 打印测试结果print *, "Sum of matrix column(", i, ")elements :", relvector(i)end do
end program

C源代码

文件名：cMain.cpp

#include <iostream>
#include <stdio.h>
#include <hip/hip_runtime.h>using namespace std;// 测试标志
#define TESTON true
// 错误码
#define MATRIX_SIZE_ERROR        0x100001  // 矩阵规模错误
#define MATRIX_POINTER_NULL      0x100002  // 矩阵空指针
#define DCUMEM_ALLOC_ERROR       0x100003  // 内存分配错误/* DCU核函数，对矩阵的每一行元素加上其行号* pfMatrix: 需要处理的矩阵首地址* row: 矩阵的行数* col: 矩阵的列数* */
__global__ void MatrixProcess(double *pfMatrix, int row, int col)
{// 测试if (TESTON){// DCU编程不支持C++库printf("DCU Program Started! \n");}// 错误检查if (row < 1 || col < 1){printf("ERROR OCCURRE! ERROR CODE : %x", MATRIX_SIZE_ERROR);return;}if (pfMatrix == NULL){printf("ERROR OCCURRE! ERROR CODE : %x", MATRIX_POINTER_NULL);return;}int tid;double *pfMatRow;tid = blockDim.x * blockIdx.x + threadIdx.x;// 依据该线程ID找到该线程需要处理的数据的首地址pfMatRow = (double*)((char*)pfMatrix + col * tid * sizeof(double));for (int i = 0; i < col; i++){pfMatRow[i] = pfMatRow[i] + tid;}return;
}/* Fortran调用C的入口函数，进行混编的c函数名需要加'_'(下划线)，传入函数的参数均为指针* matrix: 矩阵首地址* row: 矩阵的行数* col: 矩阵的列数* */
extern "C" void computefordcu_(double *matrix, int *row, int *col)
{// 测试if (TESTON){cout << "C Program Started!" << endl;}// 错误检查if (*row < 1 || *col < 1){printf("ERROR OCCURRE! ERROR CODE : %x", MATRIX_SIZE_ERROR);return;}int iErr, iSize;// DCU显存指针double *pfDevMatrix;// 统计时间float fTimeElapsed;// 任务流hipStream_t hipStream;// 记录时间hipEvent_t start, stop;//设置当前设备hipSetDevice(0);// 创建事件hipEventCreate(&start);hipEventCreate(&stop);// 创建任务流hipStreamCreate(&hipStream);// 为待处理矩阵指针分配DCU显存iSize = (*row) * (*col) * sizeof(double);hipMalloc((void**)&pfDevMatrix, iSize);if (pfDevMatrix == nullptr){printf("ERROR OCCURRE! ERROR CODE : %x", DCUMEM_ALLOC_ERROR);return;}hipEventRecord(start, hipStream);// 内存数据拷贝到显存hipMemcpyAsync(pfDevMatrix, matrix, iSize, hipMemcpyHostToDevice, hipStream);// 启动核函数，每个线程处理矩阵的一行数据hipLaunchKernelGGL(MatrixProcess, dim3(1), dim3(*row), 0, hipStream, pfDevMatrix, *row, *col);// 显存数据拷贝回内存hipMemcpyAsync(matrix, pfDevMatrix, iSize, hipMemcpyDeviceToHost, hipStream);hipEventRecord(stop, hipStream);// 同步hipEventSynchronize(stop);hipEventElapsedTime(&fTimeElapsed, start, stop);cout << "DCU TIME : " << fTimeElapsed << " ms" << endl;//清除设备显存hipFree(pfDevMatrix);//清除事件hipEventDestroy(start);hipEventDestroy(stop);//清除流hipStreamDestroy(hipStream);return;
}

MakeFile文件

文件名：makefile

HIP_INSTALL_PATH = /public/software/compiler/rocm/rocm-4.0.1# 编译器
CC = hipcc
FC = gfortran# 链接器
LD = gfortran# 编译C程序的指令
CFLAGS = --offload-arch=gfx906
# 编译Fortran程序的指令
FFLAGS =# 链接指令
LIBS = -lstdc++ -lamdhip64 -lhsa-runtime64 -lamd_comgr -lhsakmt -L$(HIP_INSTALL_PATH)/lib64/ -L$(HIP_INSTALL_PATH)/hip/lib -L$(HIP_INSTALL_PATH)/lib/ -L$(HIP_INSTALL_PATH)/hipfort/lib/# Fortran源文件
FFILES = fMain.F90# C/HIP源文件
CFILES = cMain.cpp# 中间文件
OBJECTS = fMain.o cMain.o# 目标文件
TARGET = main.outall:$(FC) $(FFLAGS) -c $(FFILES)$(CC) $(CFLAGS) -c $(CFILES)$(LD) $(LIBS) $(OBJECTS) -o $(TARGET)clean:rm -f *.o $(TARGET)

申请DCU节点

salloc -p kshdnormal -N 1 --gres=dcu:1

切换运行环境及编译环境

module switch compiler/rocm/4.0.1

编译及运行程序

make