Add fortran_hip_interop example which is based on use_device_ptr smok…

…e test.

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

examples/fortran/fortran_hip_interop/fortran_callable_init.hip

@@ -0,0 +1,11 @@
+#include "hip/hip_runtime.h"
+#include <stdio.h>
+__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;
+}

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
+