历史上遗留下许多优秀的代码中,
不乏用Fortran开发之例。这些程序有待并行化从而在GPU上运行。
一个有效的途径是首先把可并行化的Fortran代码段改成C 代码,其中并行化部分在rocm上实现。
这里介绍一个简单的Fortran + C 混合编程实例, 它通过在Fortran 程序中调用 C程序,进而实现对DCU适配。
该实例包含四个文件,fcode.f , dcu.cpp , main.f,
Makefile 。 把这三个文件拷贝到工作目录,然后再在命令行执行命令 make, 从而生成可执行代码 test。执行test返回PASSED ! 标志实测成功。实测环境中含hipcc模块为dtk22.04 或rocm-4.0.1 。
四个文件及相应代码如下:
fcode.f :
SUBROUTINE fcode()
implicit double precision (a-h, o-z)
parameter (M=64)
dimension x(M),y(M),z(M)
c
do k=1, M
x(k) = 1.0
y(k) = 2.0
z(k) = 0.0
end do
call dcu(x,y,z,M)
do k = 1, M
if ( z(k) .ne. 3.0 ) then
write(6,*) 'u fail !'
return
endif
end do
write(6,*)' PASSED !'
return
end
dcu.cpp :
#include<hip/hip_runtime.h>
#define HIP_ASSERT(status) \
assert(status == hipSuccess)
global void add(double *x, double *y,
double *z, const unsigned int M) {
z[threadIdx.x] = x[threadIdx.x] + y[threadIdx.x];
}
extern "C" void dcu_(double
*h_x,double *h_y,double *h_z, int &M){
HIP_ASSERT(hipSetDevice(0));
double *d_x, *d_y,*d_z;
HIP_ASSERT(hipMalloc((void **)&d_x, M*sizeof(double)));
HIP_ASSERT(hipMalloc((void **)&d_y, M*sizeof(double)));
HIP_ASSERT(hipMalloc((void **)&d_z, M*sizeof(double)));
HIP_ASSERT(hipMemcpy(d_x, h_x, M*sizeof(double),
hipMemcpyHostToDevice));
HIP_ASSERT(hipMemcpy(d_y, h_y, M*sizeof(double), hipMemcpyHostToDevice));
HIP_ASSERT(hipMemcpy(d_z, h_z, M*sizeof(double),
hipMemcpyHostToDevice));
hipLaunchKernelGGL(add, 1, 64, 0, 0, d_x,d_y,d_z,M);
HIP_ASSERT(hipMemcpy(h_z, d_z, M*sizeof(double), hipMemcpyDeviceToHost));
hipFree(d_x);
hipFree(d_y);
hipFree(d_z);
}
main.f :
call fcode()
stop
end
Makefile :
OBJS=main.o fcode.o dcu.o
FC=gfortran
HC=hipcc
FCFLAGS= -c
HCFLAGS= -c
LDFLAGS = -lgfortran
SrcSufF=f
SrcSufCPP=cpp
ObjSuf=o
EXEC=test
all :$(EXEC)
$(EXEC) :$(OBJS)
$(HC) $(OBJS) $(LDFLAGS) -o $(EXEC)
.$(SrcSufF).$(ObjSuf):
$(FC) $(FCFLAGS) $<
.$(SrcSufCPP).$(ObjSuf):
$(HC) $(HCFLAGS) $<
暂无数据
暂无数据