upgpkg: hipfort 5.6.0-1

upstream release

Add integration test

5 files changed, 110 insertions, 6 deletions

diff --git a/.SRCINFO b/.SRCINFO
index 1bd2347445ab..812435726ba5 100644
--- a/.SRCINFO
+++ b/.SRCINFO
@@ -1,6 +1,6 @@
 pkgbase = hipfort
 	pkgdesc = Fortran interfaces for ROCm libraries
-	pkgver = 5.4.1
+	pkgver = 5.6.0
 	pkgrel = 1
 	url = https://rocmsoftwareplatform.github.io/hipfort/
 	arch = x86_64
@@ -9,7 +9,7 @@ pkgbase = hipfort
 	depends = hip
 	depends = gcc-fortran
 	options = !strip
-	source = hipfort-5.4.1.tar.gz::https://github.com/ROCmSoftwarePlatform/hipfort/archive/rocm-5.4.1.tar.gz
-	sha256sums = 1e2d4ca623729668fab997dce67365bd5c28570fca7822f17b99f2302f9cea28
+	source = hipfort-5.6.0.tar.gz::https://github.com/ROCmSoftwarePlatform/hipfort/archive/rocm-5.6.0.tar.gz
+	sha256sums = 03176a099bc81e212ad1bf9d86f35561f8f2d21a2f126732d7620e1ea59888d5
 
 pkgname = hipfort
diff --git a/PKGBUILD b/PKGBUILD
index 62259481541e..c203bfff4c86 100644
--- a/PKGBUILD
+++ b/PKGBUILD
@@ -1,7 +1,7 @@
 # Maintainer: Torsten Keßler <t dot kessler at posteo dot de>
 
 pkgname=hipfort
-pkgver=5.4.1
+pkgver=5.6.0
 pkgrel=1
 pkgdesc='Fortran interfaces for ROCm libraries'
 arch=('x86_64')
@@ -11,7 +11,7 @@ depends=('hip' 'gcc-fortran')
 makedepends=('rocm-cmake')
 _git='https://github.com/ROCmSoftwarePlatform/hipfort'
 source=("$pkgname-$pkgver.tar.gz::$_git/archive/rocm-$pkgver.tar.gz")
-sha256sums=('1e2d4ca623729668fab997dce67365bd5c28570fca7822f17b99f2302f9cea28')
+sha256sums=('03176a099bc81e212ad1bf9d86f35561f8f2d21a2f126732d7620e1ea59888d5')
 options=(!strip)
 _dirname="$(basename "$_git")-$(basename "${source[0]}" .tar.gz)"
 
@@ -21,7 +21,7 @@ build() {
         -B build \
         -S "$_dirname" \
         -DCMAKE_BUILD_TYPE=None \
-        -DCMAKE_INSTALL_PREFIX=/opt/rocm/hipfort
+        -DCMAKE_INSTALL_PREFIX=/opt/rocm
     cmake --build build
 }
 
diff --git a/test.cpp b/test.cpp
new file mode 100644
index 000000000000..0c8fbe6074f6
--- /dev/null
+++ b/test.cpp
@@ -0,0 +1,20 @@
+#include <hip/hip_runtime.h>
+#include <cstdio>
+
+__global__ void vector_add(double *out, double *a, double *b, int n)
+{
+  size_t index = blockIdx.x * blockDim.x + threadIdx.x;
+  size_t stride = blockDim.x * gridDim.x;
+
+  for (size_t i = index; i < n; i += stride)
+    out[i] = a[i] + b[i];
+}
+
+extern "C"
+{
+  void launch(dim3* grid, dim3* block, int shmem, hipStream_t stream, double *dout, double *da, double *db, int N)
+  {
+    //printf("launching kernel\n");
+    hipLaunchKernelGGL((vector_add), *grid, *block, shmem, stream, dout, da, db, N);
+  }
+}
diff --git a/test.f03 b/test.f03
new file mode 100644
index 000000000000..8dfbef47d412
--- /dev/null
+++ b/test.f03
@@ -0,0 +1,79 @@
+! Copied from https://github.com/ROCmSoftwarePlatform/hipfort/tree/develop/test/f2008/vecadd
+program fortran_hip
+  use hipfort
+  use hipfort_check
+  
+  implicit none
+
+  interface
+     ! dim3(320), dim3(256), 0, 0
+     subroutine launch(grid,block,shmem,stream,out,a,b,N) bind(c)
+       use iso_c_binding
+       use hipfort_types
+       implicit none
+       type(c_ptr),value :: a, b, out
+       integer(c_int), value :: N, shmem
+       type(dim3) :: grid, block
+       type(c_ptr),value :: stream
+     end subroutine
+  end interface
+
+  integer(c_int), parameter :: N = 1000000
+
+  real(8),allocatable,dimension(:) :: a,b,out
+  real(8),pointer,dimension(:) :: da => null(), db => null(),dout => null()
+
+  !type(dim3) :: grid  = dim3(320,1,1) 
+  !type(dim3) :: block = dim3(256,1,1) 
+  
+  integer :: i
+  type(hipDeviceProp_t),target :: props
+  !
+  call hipCheck(hipGetDeviceProperties(props,0))  
+  write(*,"(a)",advance="no") "-- Running test 'vecadd' (Fortran 2008 interfaces)"
+  write(*,"(a)",advance="no") "- device: "
+  i=1
+  do while ( iachar(props%name(i)) .ne. 0 ) ! print till end char
+    write(*,"(a)",advance="no") props%name(i)
+    i = i+1
+  end do 
+  write(*,"(a)",advance="no") " - "
+
+  ! Allocate host memory
+  allocate(a(N),b(N),out(N))
+
+  ! Initialize host arrays
+  a(:) = 1.0
+  b(:) = 1.0
+
+  ! Allocate array space on the device
+  call hipCheck(hipMalloc(da,N))
+  call hipCheck(hipMalloc(db,N))
+  call hipCheck(hipMalloc(dout,N))
+
+  ! Transfer data from host to device memory
+  call hipCheck(hipMemcpy(da, a, N, hipMemcpyHostToDevice))
+  call hipCheck(hipMemcpy(db, b, N, hipMemcpyHostToDevice))
+
+  ! launch kernel
+  call launch(dim3(320),dim3(256),0,c_null_ptr,c_loc(dout),c_loc(da),c_loc(db),N)
+  !call launch(grid,block,0,c_null_ptr,c_loc(dout),c_loc(da),c_loc(db),N)
+  call hipCheck(hipDeviceSynchronize())
+
+  ! Transfer data back to host memory
+  call hipCheck(hipMemcpy(out, dout, N, hipMemcpyDeviceToHost))
+
+  if ( sum(out) .eq. N*2.0 ) then
+     print *, "PASSED!"
+  else
+     print *, "FAILED!"
+  endif
+
+  call hipCheck(hipFree(da))
+  call hipCheck(hipFree(db))
+  call hipCheck(hipFree(dout))
+
+  ! Deallocate host memory
+  deallocate(a,b,out)
+
+end program fortran_hip
diff --git a/test.sh b/test.sh
new file mode 100755
index 000000000000..bf6fa8a5497b
--- /dev/null
+++ b/test.sh
@@ -0,0 +1,5 @@
+#! /usr/bin/env sh
+
+OUT=$(mktemp -d)
+/opt/rocm/bin/hipfc test.f03 test.cpp -o "$OUT"/test
+"$OUT"/test