[Feature]: inputs to hipMemcyp2D with rank greater than two

[marsdeno]

Suggestion Description

cudaMemcpy2D accepts fortran pointers to arrays with ranks greater than 2., while hip equivalent does not. Is this a conscious choice? I have code which uses memcpy2d on 3- and 4-d arrays. I can obviously rewrite, but was hoping to keep changes with respect to cuda version minimal.

Operating System

No response

GPU

No response

ROCm Component

No response

[domcharrier]

Hi @marsdeno,

We are currently providing hipmemcpy2D interface members for up to dimension 2 (as you have observed) via assumed-shape array arguments (ex: real :: A(:,:)).

Example of such an interface member procedure (from : https://github.com/ROCm/hipfort/blob/develop/lib/hipfort/hipfort_hipmemcpy.f):

function hipMemcpy_l_2_c_int(dest,src,length,myKind)
      use iso_c_binding
#ifdef USE_CUDA_NAMES
      use hipfort_cuda_errors
#endif
      use hipfort_enums
      use hipfort_types
      implicit none
      logical(c_bool),target,dimension(:,:),intent(inout) :: dest
      logical(c_bool),target,dimension(:,:),intent(in)    :: src
      integer(c_int),intent(in) :: length
      integer(kind(hipMemcpyHostToHost)) :: myKind
#ifdef USE_CUDA_NAMES
      integer(kind(cudaSuccess)) :: hipMemcpy_l_2_c_int 
#else
      integer(kind(hipSuccess)) :: hipMemcpy_l_2_c_int
#endif
      !
      hipMemcpy_l_2_c_int = hipMemcpy_(c_loc(dest),c_loc(src),length*1_8,myKind)
end function

On the short term, you could create an overloaded interface for hipmemcpy2D yourself; see the next section.

Short term: Workaround?

As a workaround, you can create a similar procedure yourself that agrees with your application's datatypes and ranks and then overload the interface hipmemcpy2D yourself to support what you want to to do.

Below is a minimal example that showcases the steps to do (developed and tested via https://onlinegdb.com/2elbMJizz):

module hipfort
   interface hipmemcpy2D 
      module procedure :: foo
   end interface
contains
   subroutine foo(bar)
     integer :: bar
     print *,"a scalar"
   end subroutine
end module

Program Hello
  use hipfort

  ! user-side overloading of interface
  interface hipmemcpy2D
    procedure :: foo2
  end interface

  integer :: a, b(2)

  ! call already existing hipfort interface member
  call hipmemcpy2D(a) ! output: "a scalar"
  ! call our custom interface member
  call hipmemcpy2D(b) ! output: "an array"

contains
  subroutine foo2(bar)
     integer :: bar(:)
     print *, "an array"
  end subroutine 
End Program Hello

Should we use Fortran 2018 features?

Out of curiosity: What Fortran standard do you currently assume for your application?

With Fortran 2018, one could use assumed-rank array arguments (ex: real :: A(..)). This would allow to construct interface members such as the routine sub1 in the below snippet that can take arrays of any rank as argument:

REAL :: a0
REAL :: a1(10)
REAL :: a2(10, 20)
REAL, POINTER :: a3(:,:,:)

CALL sub1(a0)
CALL sub1(a1)
CALL sub1(a2)
CALL sub1(a3)

CONTAINS
  SUBROUTINE sub1(a)
    REAL :: a(..)
    PRINT *, RANK(a)
  END

END

There is ongoing work to revise the interfaces to support more ranks/provide better interfaces for users of more modern Fortran compilers. So feedback from your side on this is highly appreciated!