Permit INV_TRANS with OpenMP offload

src/trans/gpu/algor/buffered_allocator_mod.F90

@@ -146,8 +146,6 @@ SUBROUTINE ASSIGN_PTR_FLOAT(DST, SRC, START_IN_BYTES, LENGTH_IN_BYTES, SET_VALUE
     ELSE
 #ifdef ACCGPU
         SET_STREAM_EFF = ACC_ASYNC_SYNC
-#endif
-#ifdef OMPGPU
 #endif
     ENDIF
     IF (SET_VALUE_EFF .AND. LENGTH_IN_BYTES > 0) THEN
@@ -157,14 +155,13 @@ SUBROUTINE ASSIGN_PTR_FLOAT(DST, SRC, START_IN_BYTES, LENGTH_IN_BYTES, SET_VALUE
       !$ACC PARALLEL PRESENT(SRC) ASYNC(SET_STREAM_EFF)
 #endif
 #ifdef OMPGPU
+      !$OMP TARGET TEAMS DISTRIBUTE PARALLEL DO DEFAULT(NONE) SHARED(LENGTH_IN_BYTES,SRC)
 #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, &
@@ -196,8 +193,6 @@ SUBROUTINE ASSIGN_PTR_DOUBLE(DST, SRC, START_IN_BYTES, LENGTH_IN_BYTES, SET_VALU
     ELSE
 #ifdef ACCGPU
         SET_STREAM_EFF = ACC_ASYNC_SYNC
-#endif
-#ifdef OMPGPU
 #endif
     ENDIF
     IF (SET_VALUE_EFF .AND. LENGTH_IN_BYTES > 0) THEN
@@ -208,14 +203,13 @@ SUBROUTINE ASSIGN_PTR_DOUBLE(DST, SRC, START_IN_BYTES, LENGTH_IN_BYTES, SET_VALU
       !$ACC PARALLEL PRESENT(SRC) ASYNC(SET_STREAM_EFF)
 #endif
 #ifdef OMPGPU
+      !$OMP TARGET TEAMS DISTRIBUTE PARALLEL DO DEFAULT(NONE) SHARED(LENGTH_IN_BYTES,SRC)
 #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, &

src/trans/gpu/algor/ext_acc.F90

@@ -290,6 +290,7 @@ subroutine ext_acc_create(ptrs, stream)
       !$acc enter data create(pp) async(stream_act)
 #endif
 #ifdef OMPGPU
+      !$omp target enter data map(alloc:pp)
 #endif
     enddo
   end subroutine
@@ -327,6 +328,7 @@ subroutine ext_acc_copyin(ptrs, stream)
       !$acc enter data copyin(pp) async(stream_act)
 #endif
 #ifdef OMPGPU
+      !$omp target enter data map(to:pp)
 #endif
     enddo
   end subroutine
@@ -363,6 +365,7 @@ subroutine ext_acc_copyout(ptrs, stream)
       !$acc exit data copyout(pp) async(stream_act)
 #endif
 #ifdef OMPGPU
+      !$omp target exit data map(from:pp)
 #endif
     enddo
   end subroutine
@@ -399,6 +402,7 @@ subroutine ext_acc_delete(ptrs, stream)
       !$acc exit data delete(pp) async(stream_act)
 #endif
 #ifdef OMPGPU
+      !$omp target exit data map(delete:pp)
 #endif
     enddo
   end subroutine

src/trans/gpu/algor/hicblas_gemm.hip.cpp

@@ -25,11 +25,7 @@
 bool hip_alreadyAllocated_sgemm = false;
 bool hip_alreadyAllocated_sgemm_handle = false;
 
-bool hip_alreadyAllocated_dsgemm = false;
-bool hip_alreadyAllocated_dgemm_handle = false;
-
 hipblasHandle_t handle_hip_sgemm;
-hipblasHandle_t handle_hip_dgemm;
 
 namespace {
 struct cache_key {
@@ -199,15 +195,25 @@ template <typename Real> struct hipblas_gemm_grouped {
   void operator()(hipStream_t stream, int m, int n, int k, Real alpha,
                   const Real *A, int lda, const Real *B, int ldb, Real beta,
                   Real *C, int ldc) const {
+    // TODO: sort out this nonsense
+#ifdef OMPGPU
+    hipblasHandle_t handle;
+    HICBLAS_CHECK(hipblasCreate(&handle));
+#endif
+#ifdef ACCGPU
     hipblasHandle_t handle = get_hipblas_handle();
     HICBLAS_CHECK(hipblasSetStream(handle, stream));
+#endif
 
     if constexpr (std::is_same<Real, float>::value)
       HICBLAS_CHECK(hipblasSgemm(handle, transa_, transb_, m, n, k, &alpha, A,
                                  lda, B, ldb, &beta, C, ldc));
     if constexpr (std::is_same<Real, double>::value)
       HICBLAS_CHECK(hipblasDgemm(handle, transa_, transb_, m, n, k, &alpha, A,
                                  lda, B, ldb, &beta, C, ldc));
+#ifdef OMPGPU
+    HICBLAS_CHECK(hipblasDestroy(handle));
+#endif
   }
 
 private:
@@ -278,10 +284,6 @@ void hipblas_dgemm_wrapper(char transa, char transb, int m, int n, int k,
   if (transb == 'T' || transb == 't')
     op_t2 = HIPBLAS_OP_T;
 
-  if (!hip_alreadyAllocated_dgemm_handle) {
-    HICBLAS_CHECK(hipblasCreate(&handle_hip_dgemm));
-    hip_alreadyAllocated_dgemm_handle = true;
-  }
   hipblasHandle_t handle = get_hipblas_handle();
   HICBLAS_CHECK(hipblasSetStream(handle, *(hipStream_t *)stream));
 

src/trans/gpu/algor/hicblas_mod.F90

@@ -7,21 +7,8 @@
 ! nor does it submit to any jurisdiction.
 !
 
-#if defined CUDAGPU
-#define ACC_GET_HIP_STREAM ACC_GET_CUDA_STREAM
-#define OPENACC_LIB OPENACC
-#endif
-
 MODULE HICBLAS_MOD
 
-USE EC_PARKIND, ONLY: JPIM, JPRM, JPRD, JPIB
-USE GROWING_ALLOCATOR_MOD, ONLY: GROWING_ALLOCATION_TYPE
-#ifdef ACCGPU
-USE OPENACC_LIB, ONLY: ACC_GET_HIP_STREAM
-#endif
-#ifdef OMPGPU
-#endif
-
 IMPLICIT NONE
 
 INTERFACE
@@ -118,224 +105,4 @@ SUBROUTINE HIP_SGEMM_GROUPED( &
   END SUBROUTINE HIP_SGEMM_GROUPED
 END INTERFACE
 
-CONTAINS
-
-SUBROUTINE HIP_DGEMM_BATCHED_OVERLOAD( &
-  & TRANSA, TRANSB, &
-  & M, N, K, &
-  & ALPHA, &
-  & AARRAY, LDA, STRIDEA, &
-  & BARRAY, LDB, STRIDEB, &
-  & BETA, &
-  & CARRAY, LDC, STRIDEC, &
-  & BATCHCOUNT, STREAM, ALLOC)
-  USE ISO_C_BINDING, ONLY: C_CHAR, C_INT, C_LONG, C_LOC
-  CHARACTER(1,C_CHAR), VALUE :: TRANSA, TRANSB
-  INTEGER(KIND=JPIM) :: M
-  INTEGER(KIND=JPIM) :: N
-  INTEGER(KIND=JPIM) :: K
-  REAL(KIND=JPRD) :: ALPHA
-  REAL(KIND=JPRD), DIMENSION(:) :: AARRAY
-  INTEGER(KIND=JPIM) :: LDA
-  INTEGER(KIND=JPIM) :: STRIDEA
-  REAL(KIND=JPRD), DIMENSION(:,:) :: BARRAY
-  INTEGER(KIND=JPIM) :: LDB
-  INTEGER(KIND=JPIM) :: STRIDEB
-  REAL(KIND=JPRD) :: BETA
-  REAL(KIND=JPRD), DIMENSION(:) :: CARRAY
-  INTEGER(KIND=JPIM) :: LDC
-  INTEGER(KIND=JPIM) :: STRIDEC
-  INTEGER(KIND=JPIM) :: BATCHCOUNT
-  INTEGER(KIND=C_INT) :: STREAM
-  TYPE(GROWING_ALLOCATION_TYPE), INTENT(IN), POINTER :: ALLOC
-
-  INTEGER(KIND=C_LONG) :: HIP_STREAM
-
-#ifdef ACCGPU
-  HIP_STREAM = INT(ACC_GET_HIP_STREAM(STREAM), C_LONG)
-#endif
-#ifdef OMPGPU
-#endif
-
-#if defined(_CRAYFTN)
-#ifdef ACCGPU
-  !$ACC HOST_DATA USE_DEVICE(AARRAY,BARRAY,CARRAY)
-#endif
-#endif
-  CALL HIP_DGEMM_BATCHED( &
-    & TRANSA, TRANSB, &
-    & M, N, K, &
-    & ALPHA, &
-    & AARRAY, LDA, STRIDEA, &
-    & BARRAY, LDB, STRIDEB, &
-    & BETA, &
-    & CARRAY, LDC, STRIDEC, &
-    & BATCHCOUNT, HIP_STREAM, C_LOC(ALLOC))
-#if defined(_CRAYFTN)
-#ifdef ACCGPU
-  !$ACC END HOST_DATA
-#endif
-#endif
-END SUBROUTINE HIP_DGEMM_BATCHED_OVERLOAD
-
-SUBROUTINE HIP_SGEMM_BATCHED_OVERLOAD( &
-  & TRANSA, TRANSB, &
-  & M, N, K, &
-  & ALPHA, &
-  & AARRAY, LDA, STRIDEA, &
-  & BARRAY, LDB, STRIDEB, &
-  & BETA, &
-  & CARRAY, LDC, STRIDEC, &
-  & BATCHCOUNT, STREAM, ALLOC)
-  USE ISO_C_BINDING, ONLY: C_CHAR, C_INT, C_LONG, C_LOC
-  CHARACTER(1,C_CHAR), VALUE :: TRANSA, TRANSB
-  INTEGER(KIND=JPIM) :: M
-  INTEGER(KIND=JPIM) :: N
-  INTEGER(KIND=JPIM) :: K
-  REAL(KIND=JPRM) :: ALPHA
-  REAL(KIND=JPRM), DIMENSION(:) :: AARRAY
-  INTEGER(KIND=JPIM) :: LDA
-  INTEGER(KIND=JPIM) :: STRIDEA
-  REAL(KIND=JPRM), DIMENSION(*) :: BARRAY
-  INTEGER(KIND=JPIM) :: LDB
-  INTEGER(KIND=JPIM) :: STRIDEB
-  REAL(KIND=JPRM) :: BETA
-  REAL(KIND=JPRM), DIMENSION(:) :: CARRAY
-  INTEGER(KIND=JPIM) :: LDC
-  INTEGER(KIND=JPIM) :: STRIDEC
-  INTEGER(KIND=JPIM) :: BATCHCOUNT
-  INTEGER(KIND=C_INT) :: STREAM
-  TYPE(GROWING_ALLOCATION_TYPE), INTENT(IN), POINTER :: ALLOC
-
-  INTEGER(KIND=C_LONG) :: HIP_STREAM
-
-#ifdef ACCGPU
-  HIP_STREAM = INT(ACC_GET_HIP_STREAM(STREAM), C_LONG)
-#endif
-#ifdef OMPGPU
-#endif
-
-  CALL HIP_SGEMM_BATCHED( &
-    & TRANSA, TRANSB, &
-    & M, N, K, &
-    & ALPHA, &
-    & AARRAY, LDA, STRIDEA, &
-    & BARRAY, LDB, STRIDEB, &
-    & BETA, &
-    & CARRAY, LDC, STRIDEC, &
-    & BATCHCOUNT, HIP_STREAM, C_LOC(ALLOC))
-END SUBROUTINE HIP_SGEMM_BATCHED_OVERLOAD
-
-SUBROUTINE HIP_DGEMM_GROUPED_OVERLOAD( &
-  & RESOL_ID, BLAS_ID, TRANSA, TRANSB, &
-  & M, N, K, &
-  & ALPHA, &
-  & AARRAY, LDA, OFFSETA, &
-  & BARRAY, LDB, OFFSETB, &
-  & BETA, &
-  & CARRAY, LDC, OFFSETC, &
-  & BATCHCOUNT, STREAM, ALLOC)
-  USE ISO_C_BINDING, ONLY: C_INT, C_CHAR, C_LONG, C_LOC
-  INTEGER(KIND=C_INT), INTENT(IN) :: RESOL_ID
-  INTEGER(KIND=C_INT), INTENT(IN) :: BLAS_ID
-  CHARACTER(1,C_CHAR), VALUE :: TRANSA, TRANSB
-  INTEGER(KIND=JPIM) :: M
-  INTEGER(KIND=JPIM) :: N(:)
-  INTEGER(KIND=JPIM) :: K(:)
-  REAL(KIND=JPRD) :: ALPHA
-  REAL(KIND=JPRD), DIMENSION(:) :: AARRAY
-  INTEGER(KIND=JPIM) :: LDA
-  INTEGER(KIND=JPIB) :: OFFSETA(:)
-  REAL(KIND=JPRD), DIMENSION(*) :: BARRAY
-  INTEGER(KIND=JPIM) :: LDB(:)
-  INTEGER(KIND=JPIB) :: OFFSETB(:)
-  REAL(KIND=JPRD) :: BETA
-  REAL(KIND=JPRD), DIMENSION(:) :: CARRAY
-  INTEGER(KIND=JPIM) :: LDC
-  INTEGER(KIND=JPIB) :: OFFSETC(:)
-  INTEGER(KIND=JPIM) :: BATCHCOUNT
-  INTEGER(KIND=C_INT) :: STREAM
-  TYPE(GROWING_ALLOCATION_TYPE), INTENT(IN), POINTER :: ALLOC
-
-  INTEGER(KIND=C_LONG) :: HIP_STREAM
-
-#ifdef ACCGPU
-  HIP_STREAM = INT(ACC_GET_HIP_STREAM(STREAM), C_LONG)
-#endif
-#ifdef OMPGPU
-#endif
-
-  CALL HIP_DGEMM_GROUPED( &
-    & RESOL_ID, BLAS_ID, TRANSA, TRANSB, &
-    & M, N, K, &
-    & ALPHA, &
-    & AARRAY, LDA, OFFSETA, &
-    & BARRAY, LDB, OFFSETB, &
-    & BETA, &
-    & CARRAY, LDC, OFFSETC, &
-    & BATCHCOUNT, HIP_STREAM, C_LOC(ALLOC))
-
-END SUBROUTINE HIP_DGEMM_GROUPED_OVERLOAD
-
-SUBROUTINE HIP_SGEMM_GROUPED_OVERLOAD(&
-  & RESOL_ID, BLAS_ID, TRANSA, TRANSB, &
-  & M, N, K, &
-  & ALPHA, &
-  & AARRAY, LDA, OFFSETA, &
-  & BARRAY, LDB, OFFSETB, &
-  & BETA, &
-  & CARRAY, LDC, OFFSETC, &
-  & BATCHCOUNT, STREAM, ALLOC)
-  USE ISO_C_BINDING, ONLY: C_INT, C_CHAR, C_LONG, C_LOC
-  INTEGER(KIND=C_INT), INTENT(IN) :: RESOL_ID
-  INTEGER(KIND=C_INT), INTENT(IN) :: BLAS_ID
-  CHARACTER(1,C_CHAR), VALUE :: TRANSA, TRANSB
-  INTEGER(KIND=JPIM) :: M
-  INTEGER(KIND=JPIM) :: N(:)
-  INTEGER(KIND=JPIM) :: K(:)
-  REAL(KIND=JPRM) :: ALPHA
-  REAL(KIND=JPRM), DIMENSION(:) :: AARRAY
-  INTEGER(KIND=JPIM) :: LDA
-  INTEGER(KIND=JPIB) :: OFFSETA(:)
-  REAL(KIND=JPRM), DIMENSION(*) :: BARRAY
-  INTEGER(KIND=JPIM) :: LDB(:)
-  INTEGER(KIND=JPIB) :: OFFSETB(:)
-  REAL(KIND=JPRM) :: BETA
-  REAL(KIND=JPRM), DIMENSION(:) :: CARRAY
-  INTEGER(KIND=JPIM) :: LDC
-  INTEGER(KIND=JPIB) :: OFFSETC(:)
-  INTEGER(KIND=JPIM) :: BATCHCOUNT
-  INTEGER(KIND=C_INT) :: STREAM
-  TYPE(GROWING_ALLOCATION_TYPE), INTENT(IN), POINTER :: ALLOC
-
-  INTEGER(KIND=C_LONG) :: HIP_STREAM
-
-#ifdef ACCGPU
-  HIP_STREAM = INT(ACC_GET_HIP_STREAM(STREAM), C_LONG)
-#endif
-#ifdef OMPGPU
-#endif
-
-#if defined(_CRAYFTN)
-#ifdef ACCGPU
-  !$ACC HOST_DATA USE_DEVICE(AARRAY,BARRAY,CARRAY)
-#endif
-#endif
-  CALL HIP_SGEMM_GROUPED( &
-    & RESOL_ID, BLAS_ID, TRANSA, TRANSB, &
-    & M, N, K, &
-    & ALPHA, &
-    & AARRAY, LDA, OFFSETA, &
-    & BARRAY, LDB, OFFSETB, &
-    & BETA, &
-    & CARRAY, LDC, OFFSETC, &
-    & BATCHCOUNT, HIP_STREAM, C_LOC(ALLOC))
-#if defined(_CRAYFTN)
-#ifdef ACCGPU
-  !$ACC END HOST_DATA
-#endif
-#endif
-
-END SUBROUTINE HIP_SGEMM_GROUPED_OVERLOAD
-
 END MODULE HICBLAS_MOD

src/trans/gpu/external/setup_trans.F90

@@ -562,13 +562,13 @@ SUBROUTINE SETUP_TRANS(KSMAX,KDGL,KDLON,KLOEN,LDSPLIT,PSTRET,&
   !$ACC WAIT(1)
 #endif
 #ifdef OMPGPU
-  !$OMP TARGET ENTER DATA MAP(TO:R)
-  !$OMP TARGET ENTER DATA MAP(TO:F%RLAPIN,F%RACTHE,F%RW)
-  !$OMP TARGET ENTER DATA MAP(TO:FG%ZAA,FG%ZAS,FG%ZEPSNM)
-  !$OMP TARGET ENTER DATA MAP(TO:D%MYMS,D%NPNTGTB0,D%NPNTGTB1,D%NSTAGT0B,D%NSTAGT1B,D%NSTAGTF,&
+  !$OMP TARGET ENTER DATA MAP(TO:R,R%NSMAX,R%NTMAX,R%NDGL,R%NDGNH)
+  !$OMP TARGET ENTER DATA MAP(TO:F,F%RLAPIN,F%RACTHE,F%RW)
+  !$OMP TARGET ENTER DATA MAP(TO:FG,FG%ZAA,FG%ZAS,FG%ZEPSNM)
+  !$OMP TARGET ENTER DATA MAP(TO:D,D%NUMP,D%MYMS,D%NPNTGTB0,D%NPNTGTB1,D%NSTAGT0B,D%NSTAGT1B,D%NSTAGTF,&
   !$OMP&                         D%NPROCM,D%NPROCL,D%NPTRLS,D%MSTABF,D%NASM0,D%OFFSETS_GEMM1,&
-  !$OMP&                         D%OFFSETS_GEMM2)
-  !$OMP TARGET ENTER DATA MAP(TO:G%NDGLU,G%NMEN,G%NLOEN)
+  !$OMP&                         D%OFFSETS_GEMM2,D%NDGL_FS)
+  !$OMP TARGET ENTER DATA MAP(TO:G,G%NDGLU,G%NMEN,G%NLOEN)
 #endif
 
   WRITE(NOUT,*) '===GPU arrays successfully allocated'