diff --git a/examples/fortran/fortran_hip_interop/Makefile b/examples/fortran/fortran_hip_interop/Makefile new file mode 100644 index 000000000..8c7f4da62 --- /dev/null +++ b/examples/fortran/fortran_hip_interop/Makefile @@ -0,0 +1,144 @@ +#----------------------------------------------------------------------- +# +# Makefile: demo Makefile for amdgcn target (for flang). +# amdgcn targets begin with gfx. +# +# Run "make help" to see how to use this Makefile +# +#----------------------------------------------------------------------- +# MIT License +# Copyright (c) 2017 Advanced Micro Devices, Inc. All Rights Reserved. +# +# Permission is hereby granted, free of charge, to any person +# obtaining a copy of this software and associated documentation +# files (the "Software"), to deal in the Software without +# restriction, including without limitation the rights to use, copy, +# modify, merge, publish, distribute, sublicense, and/or sell copies +# of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be +# included in all copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, +# EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF +# MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND +# NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS +# BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN +# ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN +# CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +# SOFTWARE. + +TESTNAME =fortran_hip_interop +HIPOBJECT =fortran_callable_init +FILETYPE =f95 + +ifeq ($(AOMP),) +# --- Standard Makefile check for AOMP installation --- +ifeq ("$(wildcard $(AOMP))","") + ifneq ($(AOMP),) + $(warning AOMP not found at $(AOMP)) + endif + AOMP = $(HOME)/rocm/aomp + ifeq ("$(wildcard $(AOMP))","") + $(warning AOMP not found at $(AOMP)) + AOMP = /usr/lib/aomp + ifeq ("$(wildcard $(AOMP))","") + $(warning AOMP not found at $(AOMP)) + $(error Please install AOMP or correctly set env-var AOMP) + endif + endif +endif +# --- End Standard Makefile check for AOMP installation --- +endif +ifeq ($(AOMP_GPU),) +INSTALLED_GPU = $(shell $(AOMP)/bin/mygpu -d gfx900) # Default AOMP_GPU is gfx900 which is vega +AOMP_GPU ?= $(INSTALLED_GPU) +endif +ifeq ($(TARGETS),) +TARGETS =-march=$(AOMP_GPU) +endif + +ifeq (sm_,$(findstring sm_,$(AOMP_GPU))) + AOMP_GPUTARGET = nvptx64-nvidia-cuda +else + AOMP_GPUTARGET = amdgcn-amd-amdhsa +endif + +FORT =$(AOMP)/bin/flang +LFLAGS = -lamdhip64 +FFLAGS = -fopenmp -fopenmp-targets=$(AOMP_GPUTARGET) -Xopenmp-target=$(AOMP_GPUTARGET) + +# ----- Demo compile and link in one step, no object code saved +$(TESTNAME): $(TESTNAME).$(FILETYPE) fortran_callable_init.o + $(FORT) $(FFLAGS) $(TARGETS) $(LFLAGS) $^ -o $@ + +fortran_callable_init.o : fortran_callable_init.hip + $(AOMP)/bin/hipcc -c --offload-arch=$(AOMP_GPU) $^ -o $@ + +run: .FORCE $(TESTNAME) + LIBOMPTARGET_KERNEL_TRACE=1 ./$(TESTNAME) + +# ---- Demo compile and link in two steps, object saved +$(TESTNAME).o: $(TESTNAME).$(FILETYPE) + $(FORT) -c $(FFLAGS) $(TARGETS) $^ + +obin: $(TESTNAME).o fortran_callable_init.o + $(FORT) $(FFLAGS) $(TARGETS) $(LFLAGS) $^ -o obin + +run_obin: obin + LIBOMPTARGET_KERNEL_TRACE=1 ./obin + +# ---- Demo compile to intermediates LLVMIR or assembly +$(TESTNAME).ll: $(TESTNAME).$(FILETYPE) fortran_callable_init.ll + $(FORT) -c -S -emit-llvm $(FFLAGS) $(TARGETS) $(TESTNAME).$(FILETYPE) + +fortran_callable_init.ll : fortran_callable_init.hip + $(AOMP)/bin/hipcc -c -S -emit-llvm --offload-arch=$(AOMP_GPU) $^ + +$(TESTNAME).s: $(TESTNAME).$(FILETYPE) fortran_callable_init.s + $(FORT) -c -S $(FFLAGS) $(TARGETS) $(TESTNAME).$(FILETYPE) + +fortran_callable_init.s : fortran_callable_init.hip + $(AOMP)/bin/hipcc -c -S --offload-arch=$(AOMP_GPU) $^ + +.FORCE: + rm -f $(TESTNAME) + +help: + @echo + @echo "Makefile Help:" + @echo " Source: $(TESTNAME).$(FILETYPE)" + @echo " Compiler: $(FORT)" + @echo " Compiler flags: $(FFLAGS)" + @echo + @echo "Avalable Targets:" + @echo " make // build binary $(TESTNAME)" + @echo " make run // run $(TESTNAME)" + @echo " make $(TESTNAME).o // compile, be, & assemble : -c" + @echo " make obin // link step only" + @echo " make run_obin // run obin " + @echo " make $(TESTNAME).s // compile & backend steps : -c -S" + @echo " make $(TESTNAME).ll // compile step only : -c -S -emit-llvm" + @echo " make clean // cleanup files" + @echo " make help // this help" + @echo + @echo "Environment Variables:" + @echo " AOMP default: $(HOME)/rocm/aomp value: $(AOMP)" + @echo " AOMP_GPU default: gfx900 value: $(AOMP_GPU)" + @echo " CUDA default: /usr/local/cuda value: $(CUDA)" + @echo " TARGETS default: --offload-arch=$(AOMP_GPU)" + @echo " value: $(TARGETS)" + @echo + @echo "Link Flags:" + @echo " Link flags: $(LFLAGS)" + @echo + +# Cleanup anything this makefile can create +clean: + @[ -f ./$(TESTNAME) ] && rm ./$(TESTNAME) ; true + @[ -f ./obin ] && rm ./obin ; true + @[ -f ./$(TESTNAME).ll ] && rm *.ll ; true + @[ -f ./$(TESTNAME).o ] && rm $(TESTNAME).o ; true + @[ -f ./$(TESTNAME).s ] && rm *.s ; true + @[ -f ./$(HIPOBJECT).o ] && rm $(HIPOBJECT).o ; true diff --git a/examples/fortran/fortran_hip_interop/fortran_callable_init.hip b/examples/fortran/fortran_hip_interop/fortran_callable_init.hip new file mode 100644 index 000000000..a6bd393bc --- /dev/null +++ b/examples/fortran/fortran_hip_interop/fortran_callable_init.hip @@ -0,0 +1,11 @@ +#include "hip/hip_runtime.h" +#include +__global__ void init_array(float *arr1){ + int i = hipBlockIdx_x; + arr1[i]=(i)*2.0; +} + +extern "C" int fortran_callable_init(float **a, int N) { + hipLaunchKernelGGL((init_array), dim3(1), dim3(N), 0, 0, *a); + return 0; +} diff --git a/examples/fortran/fortran_hip_interop/fortran_hip_interop.f95 b/examples/fortran/fortran_hip_interop/fortran_hip_interop.f95 new file mode 100644 index 000000000..e664747f5 --- /dev/null +++ b/examples/fortran/fortran_hip_interop/fortran_hip_interop.f95 @@ -0,0 +1,49 @@ +program main + + use omp_lib + implicit none + interface + subroutine fortran_callable_init(a,N) bind(c) + use iso_c_binding + implicit none + type(c_ptr) :: a + integer, value :: N + end subroutine + end interface + integer :: nx,x + integer, parameter :: sp = kind(1.0_4) + real(sp), target, allocatable :: arr1(:), crr1(:) + nx = 16 +!!!!!!!! allocate arrays !!!!!!!! + + allocate(arr1(nx)) + allocate(crr1(nx)) + +!!!!!!!!! Initialise arrays !!!!!!!! + + arr1(:)=0. + + + !$OMP TARGET DATA MAP(tofrom:arr1) MAP(from:crr1) + + !$OMP TARGET DATA USE_DEVICE_PTR(arr1) + call fortran_callable_init(c_loc(arr1),nx) + !$OMP END TARGET DATA + + !$OMP TARGET TEAMS DISTRIBUTE PARALLEL DO PRIVATE(x) & + !$OMP NOWAIT + do x=1,nx + crr1(x)=arr1(x)+1.0 + end do + !$OMP END TARGET TEAMS DISTRIBUTE PARALLEL DO + + !$OMP TASKWAIT + + !$OMP END TARGET DATA + print *,arr1, crr1 + deallocate(arr1) + deallocate(crr1) + + +end +