概述
简介
将在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 main
implicit none
! 进行混合编译的c函数名(computefordcu)必须是小写,原因是Fortran语言不区分大小写,而混合编译时对其是区分大小写的
external :: computefordcu
! 测试矩阵,对矩阵进行每一列的元素加上该列号的操作,Fortran存储为列优先
real*8, allocatable, dimension(:,:) :: matrix
! 测试矩阵的行数,列数
integer :: row, col
! 测试结果
real*8, allocatable, dimension(:) :: relvector
integer :: i, j
row = 10
col = 8
! 为测试矩阵分配内存
allocate(matrix(row, col))
! 为测试结果分配内存
allocate(relvector(col))
! 初始化测试矩阵,Fortran数组索引从1开始,同时计算测试结果即每列元素之和
do i = 1, col
relvector(i) = 0
do j = 1, row
! 列优先初始化,颠倒顺序极易出错
matrix(j, i) = j
relvector(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, col
relvector(i) = 0
do j = 1, row
relvector(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.out
all:
$(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
最后
以上就是真实大象为你收集整理的曙光超算Fortran混编C并启动DCU计算的全部内容,希望文章能够帮你解决曙光超算Fortran混编C并启动DCU计算所遇到的程序开发问题。
如果觉得靠谱客网站的内容还不错,欢迎将靠谱客网站推荐给程序员好友。
本图文内容来源于网友提供,作为学习参考使用,或来自网络收集整理,版权属于原作者所有。
发表评论 取消回复