Allow GPU version to compile with OpenMP

src/programs/CMakeLists.txt

@@ -35,7 +35,8 @@ foreach( program ectrans-benchmark )
                                   fiat
                                   parkind_${prec}
                                   trans_gpu_${prec}
-                                  OpenACC::OpenACC_Fortran
+                                  $<${HAVE_ACC}:OpenACC::OpenACC_Fortran>
+                                  $<${HAVE_OMP}:OpenMP::OpenMP_Fortran>
                                 DEFINITIONS
                                   VERSION="gpu"
                                   $<$<COMPILE_LANG_AND_ID:Fortran,NVHPC>:USE_PINNED>

src/trans/gpu/CMakeLists.txt

@@ -58,6 +58,7 @@ ecbuild_add_library(
   PUBLIC_LIBS          fiat ectrans_common
   PRIVATE_LIBS         ${ECTRANS_GPU_HIP_LIBRARIES}
                        $<${HAVE_ACC}:OpenACC::OpenACC_Fortran>
+                       $<${HAVE_OMP}:OpenMP::OpenMP_Fortran>
                        $<${HAVE_CUTLASS}:nvidia::cutlass::cutlass>
   PRIVATE_DEFINITIONS  ${GPU_RUNTIME}GPU ${GPU_OFFLOAD}GPU
                        $<${HAVE_CUTLASS}:USE_CUTLASS>

src/trans/gpu/algor/buffered_allocator_mod.F90

@@ -12,7 +12,9 @@ MODULE BUFFERED_ALLOCATOR_MOD
   USE ABORT_TRANS_MOD,       ONLY: ABORT_TRANS
   USE ISO_C_BINDING,         ONLY: C_INT8_T, C_SIZE_T, C_LOC, C_F_POINTER
   USE GROWING_ALLOCATOR_MOD, ONLY: GROWING_ALLOCATION_TYPE
+#ifdef ACCGPU
   USE OPENACC,               ONLY: ACC_ASYNC_SYNC
+#endif
 
   IMPLICIT NONE
 
@@ -142,16 +144,28 @@ SUBROUTINE ASSIGN_PTR_FLOAT(DST, SRC, START_IN_BYTES, LENGTH_IN_BYTES, SET_VALUE
     IF (PRESENT(SET_STREAM)) THEN
         SET_STREAM_EFF = SET_STREAM
     ELSE
+#ifdef ACCGPU
         SET_STREAM_EFF = ACC_ASYNC_SYNC
+#endif
+#ifdef OMPGPU
+#endif
     ENDIF
     IF (SET_VALUE_EFF .AND. LENGTH_IN_BYTES > 0) THEN
       ! This option is turned off by default, but for experimentation we can turn it on. This is
       ! setting all bits to 1 (meaning NaN in floating point)
+#ifdef ACCGPU
       !$ACC PARALLEL PRESENT(SRC) ASYNC(SET_STREAM_EFF)
+#endif
+#ifdef OMPGPU
+#endif
       DO J=1_C_SIZE_T,LENGTH_IN_BYTES
         SRC(J) = -1
       ENDDO
+#ifdef ACCGPU
       !$ACC END PARALLEL
+#endif
+#ifdef OMPGPU
+#endif
     ENDIF
     CALL C_F_POINTER(C_LOC(SRC(START_IN_BYTES:START_IN_BYTES+LENGTH_IN_BYTES-1)), DST, &
         & [C_SIZEOF(SRC(START_IN_BYTES:START_IN_BYTES+LENGTH_IN_BYTES-1))/C_SIZEOF(DST(0))])
@@ -180,17 +194,29 @@ SUBROUTINE ASSIGN_PTR_DOUBLE(DST, SRC, START_IN_BYTES, LENGTH_IN_BYTES, SET_VALU
     IF (PRESENT(SET_STREAM)) THEN
         SET_STREAM_EFF = SET_STREAM
     ELSE
+#ifdef ACCGPU
         SET_STREAM_EFF = ACC_ASYNC_SYNC
+#endif
+#ifdef OMPGPU
+#endif
     ENDIF
     IF (SET_VALUE_EFF .AND. LENGTH_IN_BYTES > 0) THEN
       ! This option is turned off by default, but for experimentation we can turn it on. This is
       ! setting all bits to 1 (meaning NaN in floating point)
       END_IN_BYTES=START_IN_BYTES+LENGTH_IN_BYTES-1
+#ifdef ACCGPU
       !$ACC PARALLEL PRESENT(SRC) ASYNC(SET_STREAM_EFF)
+#endif
+#ifdef OMPGPU
+#endif
       DO J=1_C_SIZE_T,LENGTH_IN_BYTES
         SRC(J) = -1
       ENDDO
+#ifdef ACCGPU
       !$ACC END PARALLEL
+#endif
+#ifdef OMPGPU
+#endif
     ENDIF
     CALL C_F_POINTER(C_LOC(SRC(START_IN_BYTES:START_IN_BYTES+LENGTH_IN_BYTES-1)), DST, &
         & [C_SIZEOF(SRC(START_IN_BYTES:START_IN_BYTES+LENGTH_IN_BYTES-1))/C_SIZEOF(DST(0))])

src/trans/gpu/algor/ext_acc.F90

@@ -18,13 +18,22 @@ module openacc_ext_type
 end module
 module openacc_ext
   use iso_c_binding, only: c_ptr, c_size_t, c_loc, c_sizeof
+#ifdef ACCGPU
   use openacc, only: acc_handle_kind
+#endif
+#ifdef OMPGPU
+#endif
   use openacc_ext_type, only: ext_acc_arr_desc
   implicit none
 
   private
   public :: ext_acc_pass, ext_acc_create, ext_acc_copyin, ext_acc_copyout, &
-    & ext_acc_delete, ext_acc_arr_desc, acc_handle_kind
+#ifdef ACCGPU
+       & ext_acc_delete, ext_acc_arr_desc, acc_handle_kind
+#endif
+#ifdef OMPGPU
+       & ext_acc_delete, ext_acc_arr_desc
+#endif
 
   type common_pointer_descr
     type(c_ptr) :: ptr
@@ -247,107 +256,150 @@ function get_common_pointers(in_ptrs, out_ptrs) result(num_ranges)
     enddo
   end function
   subroutine ext_acc_create(ptrs, stream)
+#ifdef ACCGPU
     use openacc, only: acc_async_sync
+#endif
     use iso_fortran_env, only: int32
     implicit none
     type(ext_acc_arr_desc), intent(in) :: ptrs(:)
+#ifdef ACCGPU
     integer(acc_handle_kind), optional :: stream
+#endif
+#ifdef OMPGPU
+    integer(4), optional :: stream
+#endif
 
     type(common_pointer_descr), allocatable :: common_ptrs(:)
 
     integer :: i, num_ranges
     integer(kind=int32), pointer :: pp(:)
+#ifdef ACCGPU
     integer(acc_handle_kind) :: stream_act
-
     if (present(stream)) then
       stream_act = stream
     else
       stream_act = acc_async_sync
     endif
+#endif
     allocate(common_ptrs(size(ptrs)))
     num_ranges = get_common_pointers(ptrs, common_ptrs)
 
     do i = 1, num_ranges
-      call c_f_pointer(common_ptrs(i)%ptr, pp, shape=[common_ptrs(i)%sz/c_sizeof(pp(1))])
+      call c_f_pointer(common_ptrs(i)%ptr, pp, [common_ptrs(i)%sz/c_sizeof(pp(1))])
+#ifdef ACCGPU
       !$acc enter data create(pp) async(stream_act)
+#endif
+#ifdef OMPGPU
+#endif
     enddo
   end subroutine
   subroutine ext_acc_copyin(ptrs, stream)
+#ifdef ACCGPU
     use openacc, only: acc_async_sync
+#endif
     implicit none
     type(ext_acc_arr_desc), intent(in) :: ptrs(:)
+#ifdef ACCGPU
     integer(acc_handle_kind), optional :: stream
+#endif
+#ifdef OMPGPU
+    integer(4), optional :: stream
+#endif
 
     type(common_pointer_descr), allocatable :: common_ptrs(:)
 
     integer :: i, num_ranges
     integer(4), pointer :: pp(:)
-
+#ifdef ACCGPU
     integer(acc_handle_kind) :: stream_act
-
     if (present(stream)) then
       stream_act = stream
     else
       stream_act = acc_async_sync
     endif
+#endif
     allocate(common_ptrs(size(ptrs)))
     num_ranges = get_common_pointers(ptrs, common_ptrs)
 
     do i = 1, num_ranges
-      call c_f_pointer(common_ptrs(i)%ptr, pp, shape=[common_ptrs(i)%sz/c_sizeof(pp(1))])
+      call c_f_pointer(common_ptrs(i)%ptr, pp, [common_ptrs(i)%sz/c_sizeof(pp(1))])
+#ifdef ACCGPU
       !$acc enter data copyin(pp) async(stream_act)
+#endif
+#ifdef OMPGPU
+#endif
     enddo
   end subroutine
   subroutine ext_acc_copyout(ptrs, stream)
+#ifdef ACCGPU
     use openacc, only: acc_async_sync
+#endif
     implicit none
     type(ext_acc_arr_desc), intent(in) :: ptrs(:)
+#ifdef ACCGPU
     integer(acc_handle_kind), optional :: stream
-
+#endif
+#ifdef OMPGPU
+    integer(4), optional :: stream
+#endif
     type(common_pointer_descr), allocatable :: common_ptrs(:)
 
     integer :: i, num_ranges
     integer(4), pointer :: pp(:)
-
+#ifdef ACCGPU
     integer(acc_handle_kind) :: stream_act
-
     if (present(stream)) then
       stream_act = stream
     else
       stream_act = acc_async_sync
     endif
+#endif
     allocate(common_ptrs(size(ptrs)))
     num_ranges = get_common_pointers(ptrs, common_ptrs)
 
     do i = 1, num_ranges
-      call c_f_pointer(common_ptrs(i)%ptr, pp, shape=[common_ptrs(i)%sz/c_sizeof(pp(1))])
+      call c_f_pointer(common_ptrs(i)%ptr, pp, [common_ptrs(i)%sz/c_sizeof(pp(1))])
+#ifdef ACCGPU
       !$acc exit data copyout(pp) async(stream_act)
+#endif
+#ifdef OMPGPU
+#endif
     enddo
   end subroutine
   subroutine ext_acc_delete(ptrs, stream)
+#ifdef ACCGPU
     use openacc, only: acc_async_sync
+#endif
     implicit none
     type(ext_acc_arr_desc), intent(in) :: ptrs(:)
+#ifdef ACCGPU
     integer(acc_handle_kind), optional :: stream
-
+#else
+    integer(4), optional :: stream
+#endif
     type(common_pointer_descr), allocatable :: common_ptrs(:)
 
     integer :: i, num_ranges
     integer(4), pointer :: pp(:)
-
+#ifdef ACCGPU
     integer(acc_handle_kind) :: stream_act
 
     if (present(stream)) then
       stream_act = stream
     else
       stream_act = acc_async_sync
     endif
+#endif
     allocate(common_ptrs(size(ptrs)))
     num_ranges = get_common_pointers(ptrs, common_ptrs)
 
     do i = 1, num_ranges
-      call c_f_pointer(common_ptrs(i)%ptr, pp, shape=[common_ptrs(i)%sz/c_sizeof(pp(1))])
+      call c_f_pointer(common_ptrs(i)%ptr, pp, [common_ptrs(i)%sz/c_sizeof(pp(1))])
+#ifdef ACCGPU
       !$acc exit data delete(pp) async(stream_act)
+#endif
+#ifdef OMPGPU
+#endif
     enddo
   end subroutine
 end module

src/trans/gpu/algor/growing_allocator_mod.F90

@@ -43,7 +43,12 @@ SUBROUTINE REALLOCATE_GROWING_ALLOCATION(ALLOC, SZ)
 
     IF (.NOT. ASSOCIATED(ALLOC%PTR)) THEN
       ALLOCATE(ALLOC%PTR(SZ))
+#ifdef OMPGPU
+      !$OMP TARGET ENTER DATA MAP(ALLOC:ALLOC%PTR)
+#endif
+#ifdef ACCGPU
       !$ACC ENTER DATA CREATE(ALLOC%PTR)
+#endif
       ALLOC%FREE_FUNCS_SZ = 0
     ENDIF
   END SUBROUTINE
@@ -93,7 +98,12 @@ SUBROUTINE DESTROY_GROWING_ALLOCATOR(ALLOC)
         CALL ALLOC%FREE_FUNCS(I)%FUNC(ALLOC%PTR, &
                                       SIZE(ALLOC%PTR, 1, C_SIZE_T))
       ENDDO
+#ifdef OMPGPU
+      !$OMP TARGET EXIT DATA MAP(DELETE:ALLOC%PTR)
+#endif
+#ifdef ACCGPU
       !$ACC EXIT DATA DELETE(ALLOC%PTR)
+#endif
       DEALLOCATE(ALLOC%PTR)
       NULLIFY(ALLOC%PTR)
     ENDIF