WiP: New OpenACC features for GPUFORT runtime

Makefile

@@ -0,0 +1,41 @@
+.PHONY: clean_all gpufort_headers gpufort_sources lib/$(LIBGPUFORT) lib/$(LIBGPUFORT_ACC)
+
+SUFFIX         = $(if $(HIP_PLATFORM),$(HIP_PLATFORM),amd)
+LIBGPUFORT     = libgpufort_$(SUFFIX).a
+LIBGPUFORT_ACC = libgpufort_acc_$(SUFFIX).a
+
+GPUFORT_DIR     = $(shell gpufort --path)
+GPUFORT_ACC_DIR = $(GPUFORT_DIR)/runtime/gpufort_acc_runtime
+
+all: | gpufort_headers lib/$(LIBGPUFORT) lib/$(LIBGPUFORT_ACC) make_directories
+
+gpufort_headers:
+	make -C $(GPUFORT_DIR)/include all
+
+gpufort_sources:
+	make -C $(GPUFORT_DIR)/src gpufort_sources
+
+lib/$(LIBGPUFORT): | gpufort_headers gpufort_sources make_directories
+	make -C $(GPUFORT_DIR)/src $(LIBGPUFORT)
+	mv $(GPUFORT_DIR)/src/$(LIBGPUFORT) $(GPUFORT_DIR)/lib
+	mv $(GPUFORT_DIR)/src/*.mod $(GPUFORT_DIR)/include/$(SUFFIX)/
+	make -C $(GPUFORT_DIR)/src clean
+
+lib/$(LIBGPUFORT_ACC): make_directories
+	make -C $(GPUFORT_ACC_DIR)/ lib/$(LIBGPUFORT_ACC)
+	mv $(GPUFORT_ACC_DIR)/lib/$(LIBGPUFORT_ACC)\
+	    $(GPUFORT_DIR)/lib/
+	mv $(GPUFORT_ACC_DIR)/include/*.mod\
+	    $(GPUFORT_DIR)/include/$(SUFFIX)/
+	#-mv $(GPUFORT_ACC_DIR)/include/*.h\
+	#   $(GPUFORT_DIR)/include/$(SUFFIX)
+	make -C $(GPUFORT_ACC_DIR)/ clean
+
+make_directories:
+	mkdir -p $(GPUFORT_DIR)/lib
+	mkdir -p $(GPUFORT_DIR)/include/$(SUFFIX)
+
+clean_all:
+	make -C $(GPUFORT_DIR)/include clean_all
+	make -C $(GPUFORT_DIR)/src     clean_all
+	rm -f lib/$(LIBGPUFORT) lib/$(LIBGPUFORT_ACC)

README.md

@@ -25,6 +25,22 @@ by GPUFORT will in most cases require manual reviewing and fixing.
 
 Please take a look at the (slightly outdated) [user guide](https://rocmsoftwareplatform.github.io/gpufort/).
 
+### Quick start for generating and compiling HIP code form Fortran
+
+Before you can run HIP code generated by GPUFORT, you need to 
+build the GPUFORT static libraries and/or 
+offloading runtime for AMD and/or NVIDIA devices.
+Assuming HIPFORT is installed, you can compile them via:
+
+```bash
+HIP_PLATFORM=amd FC=hipfc make all # use HIP_PLATFORM=nvidia for NVIDIA devices
+```
+
+If the compilation is successful, take a look at the 
+CUDA Fortran and OpenACC examples in the `examples/` subdirectory,
+which all come with a Makefile for generating code and building
+the examples.
+
 ## Implementation details
 
 [This presentation](https://github.com/ROCmSoftwarePlatform/gpufort/blob/main/gpufort_slides.pdf)

examples/gpufort_array/vector-add-hipmalloc/Makefile

@@ -0,0 +1,8 @@
+CFLAGS = -g -ggdb
+
+vector-add:
+	hipcc $(shell gpufort --cpp_config) $(CFLAGS) -c main-kernels.hip.cpp -o main-kernels.hip.o
+	hipfc $(shell gpufort --gfortran_config) main-kernels.hip.o vector-add.f90 -o vector-add $(shell gpufort --ldflags)
+
+clean:
+	rm -f *.mod *.o vector-add

examples/gpufort_array/vector-add-hipmalloc/main-kernels.hip.cpp

@@ -0,0 +1,45 @@
+// This file was generated by gpufort
+#ifndef __KERNELS_HIP_CPP__
+#define __KERNELS_HIP_CPP__
+#include "hip/hip_runtime.h"
+#include "hip/hip_complex.h"
+#include "gpufort.h"
+#include "gpufort_array.h"
+
+
+// BEGIN main_26_e28c45
+/*
+   HIP C++ implementation of the function/loop body of:
+
+     !$cuf kernel do(1) <<<grid, tBlock>>>
+     do i=1,size(y_d,1)
+       y_d(i) = y_d(i) + a*x_d(i)
+     end do
+
+*/
+
+__global__ void  vecadd_kernel(
+    gpufort::array1<float> y_d,
+    const float a,
+    const gpufort::array1<float> x_d
+) {
+  int i = 1 + (1)*(threadIdx.x + blockIdx.x * blockDim.x);
+  if (loop_cond(i,y_d.size(1),1)) {
+    y_d(i)= y_d(i) + a*x_d(i);
+  }
+}
+
+extern "C" void launch_vecadd_kernel_(
+    const dim3& grid,
+    const dim3& block,
+    const int& sharedmem, 
+    hipStream_t& stream,
+    gpufort::array1<float>& y_d,
+    const float& a,
+    const gpufort::array1<float>& x_d
+) {
+  // launch kernel
+  hipLaunchKernelGGL((vecadd_kernel), grid, block, sharedmem, stream, y_d, a, x_d);
+}
+// END vecadd_kernel
+#endif // __KERNELS_HIP_CPP__ 

examples/gpufort_array/vector-add-hipmalloc/vector-add.f90

@@ -0,0 +1,30 @@
+program main
+  use hipfort
+  use hipfort_check
+  use gpufort_array
+  implicit none
+  integer, parameter :: N = 40000
+  real :: x(N), y(N), a
+  real,pointer :: x_d(:), y_d(:)
+  type(dim3)  :: grid, tBlock
+  type(c_ptr) :: stream = c_null_ptr
+  !
+  tBlock = dim3(256,1,1)
+  grid   = dim3(ceiling(real(N)/tBlock%x),1,1)
+
+  x = 1.0; y = 2.0; a = 2.0
+
+  call hipCheck(hipMalloc(x_d,source=x))
+  call hipCheck(hipMalloc(y_d,source=y))
+
+  call launch_vecadd_kernel(grid,tBlock,0,stream,&
+    gpufort_array_wrap_device_ptr(y_d,lbound(y)),a,&
+    gpufort_array_wrap_device_ptr(x_d))
+  call hipCheck(hipStreamSynchronize(stream))
+
+  call hipCheck(hipMemcpy(y,y_d,hipMemcpyDeviceToHost))
+  call hipCheck(hipFree(x_d))
+  call hipCheck(hipFree(y_d))
+
+  write(*,*) 'Max error: ', maxval(abs(y-4.0))
+end program main

examples/gpufort_array/vector-add-interoperable-types/Makefile

@@ -0,0 +1,8 @@
+CFLAGS = -g -ggdb
+
+vector-add:
+	hipcc $(shell gpufort --cpp_config) $(CFLAGS) -c main-kernels.hip.cpp -o main-kernels.hip.o
+	hipfc $(shell gpufort --gfortran_config) main-kernels.hip.o vector-add.f90 -o vector-add $(shell gpufort --ldflags)
+
+clean:
+	rm -f *.mod *.o vector-add

examples/gpufort_array/vector-add-interoperable-types/main-kernels.hip.cpp

@@ -0,0 +1,51 @@
+// This file was generated by gpufort
+#ifndef __KERNELS_HIP_CPP__
+#define __KERNELS_HIP_CPP__
+#include "hip/hip_runtime.h"
+#include "hip/hip_complex.h"
+#include "gpufort.h"
+#include "gpufort_array.h"
+
+// C++ structs with same layout as Fortran interop types
+struct node_t {
+  float val;
+};
+
+struct mesh_t {
+  float a;
+  gpufort::array1<node_t> x; 
+  gpufort::array1<float> y;
+};
+
+/*
+   HIP C++ implementation of the function/loop body of:
+
+     !$cuf kernel do(1) <<<grid, tBlock>>>
+     do i=1,size(mesh%y,1)
+       mesh%y(i) = mesh%y(i) + mesh%a*mesh%x(i)%val
+     end do
+*/
+
+__global__ void  vecadd_kernel(mesh_t mesh) {
+  int i = 1 + (1)*(threadIdx.x + blockIdx.x * blockDim.x);
+  if (loop_cond(i,mesh.y.size(1),1)) {
+    mesh.y(i)= mesh.y(i) + mesh.a*mesh.x(i).val;
+  }
+}
+
+extern "C" void launch_vecadd_kernel_auto_(
+    const int& sharedmem, 
+    hipStream_t& stream,
+    mesh_t& mesh
+) {
+  const int vecadd_kernel_blockX = 128;
+  dim3 block(vecadd_kernel_blockX);
+  const int vecadd_kernel_NX = (1 + ((mesh.y.size(1)) - (1)));
+  const int vecadd_kernel_gridX = divideAndRoundUp( vecadd_kernel_NX, vecadd_kernel_blockX );
+  dim3 grid(vecadd_kernel_gridX);
+
+  // launch kernel
+  hipLaunchKernelGGL((vecadd_kernel), grid, block, sharedmem, stream, mesh);
+  //mesh.y.print_device_data(std::cout,"prefix",gpufort::PrintMode::PrintCValues);
+}
+#endif // __KERNELS_HIP_CPP__ 

examples/gpufort_array/vector-add-interoperable-types/vector-add.f90

@@ -0,0 +1,81 @@
+module types
+  use iso_c_binding
+  use gpufort_array
+
+  ! interoperable types
+  type,bind(c) :: node_t
+    real(c_float) :: val
+  end type
+
+  type,bind(c) :: mesh_t
+    real(c_float)        :: a
+    type(gpufort_array1) :: x ! array of type `node_t`
+    type(gpufort_array1) :: y ! array of type `real(c_float)`
+  end type
+end module
+
+program main
+  use types
+  use hipfort
+  use hipfort_check
+  implicit none
+  integer, parameter :: N = 40000
+  !
+  type(mesh_t)          :: mesh
+  real(c_float),pointer :: mesh_t_y(:)
+  !
+  call init_mesh()
+  call gpufort_array_hostptr(mesh%y,mesh_t_y)
+
+  call launch_vecadd_kernel_auto(0,c_null_ptr,mesh)
+  call hipCheck(gpufort_array_copy_to_host(mesh%y))
+
+  write(*,*) 'Max error: ', maxval(abs(mesh_t_y-4.0))
+
+  call destroy_mesh()
+contains
+  subroutine init_mesh()
+    use gpufort_array
+    use iso_c_binding
+    implicit none
+    integer i
+    type(node_t)          :: node_t_dummy
+    type(node_t),pointer  :: mesh_t_x(:)
+    real(c_float),pointer :: mesh_t_y(:)
+
+    ! Init scalar
+    mesh%a = 2.0
+
+    ! Init basic datatype array member
+    call hipCheck(gpufort_array_init(mesh%y,&
+      c_float,sizes=[N],lbounds=[1],&
+      alloc_mode=gpufort_array_alloc_pinned_host_alloc_device)) 
+    call gpufort_array_hostptr(mesh%y,mesh_t_y) ! Get Fortran pointer to host data (overloaded routine exists for basic datatypes.
+    ! 
+    mesh_t_y(:) = 2.0                            ! Also sets lower bounds.
+    call hipCheck(gpufort_array_copy_to_device(mesh%y))
+
+    ! Init struct member.  
+    call hipCheck(gpufort_array_init(mesh%x,&
+      int(c_sizeof(node_t_dummy),c_int),&
+      sizes=[N],lbounds=[1],&
+      alloc_mode=gpufort_array_alloc_pinned_host_alloc_device))
+    !    
+    ! Get Fortran pointer to host data. No overloaded routine exists (must be type-specific).
+    ! In this case, you must/can set lower bounds manually if they are different from 1 and 
+    ! the initial data is depending on the index.
+    call c_f_pointer(mesh%x%data%data_host,mesh_t_x,shape=[N])
+    !mesh_t_x(1:) => mesh_t_x ! Set lower bounds
+    mesh_t_x(:)%val = 1
+    call hipCheck(gpufort_array_copy_to_device(mesh%x))
+  end subroutine
+
+  ! (will be) autogenerated by GPUFORT
+  subroutine destroy_mesh()
+    use gpufort_array
+    use iso_c_binding
+    implicit none
+    call hipCheck(gpufort_array_destroy(mesh%x))
+    call hipCheck(gpufort_array_destroy(mesh%y))
+  end subroutine
+end program main

examples/gpufort_array/vector-add-non-interoperable-types/Makefile

@@ -0,0 +1,8 @@
+CFLAGS = -g -ggdb
+
+vector-add:
+	hipcc $(shell gpufort --cpp_config) $(CFLAGS) -c main-kernels.hip.cpp -o main-kernels.hip.o
+	hipfc $(shell gpufort --gfortran_config) main-kernels.hip.o vector-add.f90 -o vector-add $(shell gpufort --ldflags)
+
+clean:
+	rm -f *.mod *.o vector-add

examples/gpufort_array/vector-add-non-interoperable-types/main-kernels.hip.cpp

@@ -0,0 +1,50 @@
+// This file was generated by gpufort
+#ifndef __KERNELS_HIP_CPP__
+#define __KERNELS_HIP_CPP__
+#include "hip/hip_runtime.h"
+#include "hip/hip_complex.h"
+#include "gpufort.h"
+#include "gpufort_array.h"
+
+// C++ structs with same layout as Fortran interop types
+struct node_t {
+  float val;
+};
+
+struct mesh_t {
+  float a;
+  gpufort::array1<node_t> x; 
+  gpufort::array1<float> y;
+};
+
+/*
+   HIP C++ implementation of the function/loop body of:
+
+     !$cuf kernel do(1) <<<grid, tBlock>>>
+     do i=1,size(mesh%y,1)
+       mesh%y(i) = mesh%y(i) + mesh%a*mesh%x(i)%val
+     end do
+*/
+
+__global__ void  vecadd_kernel(mesh_t mesh) {
+  int i = 1 + (1)*(threadIdx.x + blockIdx.x * blockDim.x);
+  if (loop_cond(i,mesh.y.size(1),1)) {
+    mesh.y(i)= mesh.y(i) + mesh.a*mesh.x(i).val;
+  }
+}
+
+extern "C" void launch_vecadd_kernel_auto_(
+    const int& sharedmem, 
+    hipStream_t& stream,
+    mesh_t& mesh
+) {
+  const int vecadd_kernel_blockX = 128;
+  dim3 block(vecadd_kernel_blockX);
+  const int vecadd_kernel_NX = (1 + ((mesh.y.size(1)) - (1)));
+  const int vecadd_kernel_gridX = divideAndRoundUp( vecadd_kernel_NX, vecadd_kernel_blockX );
+  dim3 grid(vecadd_kernel_gridX);
+
+  // launch kernel
+  hipLaunchKernelGGL((vecadd_kernel), grid, block, sharedmem, stream, mesh);
+}
+#endif // __KERNELS_HIP_CPP__