!
! Copyright (C) 2022, MaX CoE
! Distributed under the MIT License 
! (license terms are at http://opensource.org/licenses/MIT).
!
!--
!
! Utility functions to perform memory allocation on the device
! using CUDA-Fortran, OpenACC or OpenMP Offload
!
!==================================================================
!==================================================================
! *DO NOT EDIT*: automatically generated from devxlib_malloc_alloc.jf90
!==================================================================
!==================================================================
!
#include<devxlib_macros.h>
#include<devxlib_defs.h>
!
!=======================================
!
submodule (devxlib_malloc) devxlib_malloc_alloc

   implicit none

   contains
      module subroutine sp_devxlib_malloc_alloc_r1d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          real(real32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:)
          integer(int64), intent(in)           :: dimensions
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          real(real32) :: dummy
          type(c_ptr) :: c_dev_ptr
          real(real32), pointer, contiguous :: fptr(:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions])
                   dev_ptr(lbounds:lbounds+dimensions-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds:lbounds+dimensions-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions), stat = ierr )
#endif
          endif
          !
      end subroutine sp_devxlib_malloc_alloc_r1d
      !
      module subroutine sp_devxlib_malloc_alloc_r2d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          real(real32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:)
          integer(int64), intent(in)           :: dimensions(2)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(2)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          real(real32) :: dummy
          type(c_ptr) :: c_dev_ptr
          real(real32), pointer, contiguous :: fptr(:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2)), stat = ierr )
#endif
          endif
          !
      end subroutine sp_devxlib_malloc_alloc_r2d
      !
      module subroutine sp_devxlib_malloc_alloc_r3d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          real(real32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:)
          integer(int64), intent(in)           :: dimensions(3)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(3)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          real(real32) :: dummy
          type(c_ptr) :: c_dev_ptr
          real(real32), pointer, contiguous :: fptr(:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3)), stat = ierr )
#endif
          endif
          !
      end subroutine sp_devxlib_malloc_alloc_r3d
      !
      module subroutine sp_devxlib_malloc_alloc_r4d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          real(real32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:)
          integer(int64), intent(in)           :: dimensions(4)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(4)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          real(real32) :: dummy
          type(c_ptr) :: c_dev_ptr
          real(real32), pointer, contiguous :: fptr(:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4)), stat = ierr )
#endif
          endif
          !
      end subroutine sp_devxlib_malloc_alloc_r4d
      !
      module subroutine sp_devxlib_malloc_alloc_r5d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          real(real32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(5)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(5)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          real(real32) :: dummy
          type(c_ptr) :: c_dev_ptr
          real(real32), pointer, contiguous :: fptr(:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5)), stat = ierr )
#endif
          endif
          !
      end subroutine sp_devxlib_malloc_alloc_r5d
      !
      module subroutine sp_devxlib_malloc_alloc_r6d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          real(real32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(6)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(6)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          real(real32) :: dummy
          type(c_ptr) :: c_dev_ptr
          real(real32), pointer, contiguous :: fptr(:,:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5), dimensions(6)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                           lbounds(6):lbounds(6)+dimensions(6)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                                                             lbounds(6):lbounds(6)+dimensions(6)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5), dimensions(6)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5),&
                                                     dimensions(6)), stat = ierr )
#endif
          endif
          !
      end subroutine sp_devxlib_malloc_alloc_r6d
      !
      module subroutine dp_devxlib_malloc_alloc_r1d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          real(real64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:)
          integer(int64), intent(in)           :: dimensions
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          real(real64) :: dummy
          type(c_ptr) :: c_dev_ptr
          real(real64), pointer, contiguous :: fptr(:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions])
                   dev_ptr(lbounds:lbounds+dimensions-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds:lbounds+dimensions-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions), stat = ierr )
#endif
          endif
          !
      end subroutine dp_devxlib_malloc_alloc_r1d
      !
      module subroutine dp_devxlib_malloc_alloc_r2d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          real(real64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:)
          integer(int64), intent(in)           :: dimensions(2)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(2)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          real(real64) :: dummy
          type(c_ptr) :: c_dev_ptr
          real(real64), pointer, contiguous :: fptr(:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2)), stat = ierr )
#endif
          endif
          !
      end subroutine dp_devxlib_malloc_alloc_r2d
      !
      module subroutine dp_devxlib_malloc_alloc_r3d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          real(real64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:)
          integer(int64), intent(in)           :: dimensions(3)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(3)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          real(real64) :: dummy
          type(c_ptr) :: c_dev_ptr
          real(real64), pointer, contiguous :: fptr(:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3)), stat = ierr )
#endif
          endif
          !
      end subroutine dp_devxlib_malloc_alloc_r3d
      !
      module subroutine dp_devxlib_malloc_alloc_r4d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          real(real64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:)
          integer(int64), intent(in)           :: dimensions(4)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(4)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          real(real64) :: dummy
          type(c_ptr) :: c_dev_ptr
          real(real64), pointer, contiguous :: fptr(:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4)), stat = ierr )
#endif
          endif
          !
      end subroutine dp_devxlib_malloc_alloc_r4d
      !
      module subroutine dp_devxlib_malloc_alloc_r5d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          real(real64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(5)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(5)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          real(real64) :: dummy
          type(c_ptr) :: c_dev_ptr
          real(real64), pointer, contiguous :: fptr(:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5)), stat = ierr )
#endif
          endif
          !
      end subroutine dp_devxlib_malloc_alloc_r5d
      !
      module subroutine dp_devxlib_malloc_alloc_r6d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          real(real64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(6)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(6)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          real(real64) :: dummy
          type(c_ptr) :: c_dev_ptr
          real(real64), pointer, contiguous :: fptr(:,:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5), dimensions(6)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                           lbounds(6):lbounds(6)+dimensions(6)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                                                             lbounds(6):lbounds(6)+dimensions(6)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5), dimensions(6)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5),&
                                                     dimensions(6)), stat = ierr )
#endif
          endif
          !
      end subroutine dp_devxlib_malloc_alloc_r6d
      !
      module subroutine sp_devxlib_malloc_alloc_c1d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          complex(real32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:)
          integer(int64), intent(in)           :: dimensions
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          complex(real32) :: dummy
          type(c_ptr) :: c_dev_ptr
          complex(real32), pointer, contiguous :: fptr(:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions])
                   dev_ptr(lbounds:lbounds+dimensions-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds:lbounds+dimensions-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions), stat = ierr )
#endif
          endif
          !
      end subroutine sp_devxlib_malloc_alloc_c1d
      !
      module subroutine sp_devxlib_malloc_alloc_c2d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          complex(real32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:)
          integer(int64), intent(in)           :: dimensions(2)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(2)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          complex(real32) :: dummy
          type(c_ptr) :: c_dev_ptr
          complex(real32), pointer, contiguous :: fptr(:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2)), stat = ierr )
#endif
          endif
          !
      end subroutine sp_devxlib_malloc_alloc_c2d
      !
      module subroutine sp_devxlib_malloc_alloc_c3d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          complex(real32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:)
          integer(int64), intent(in)           :: dimensions(3)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(3)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          complex(real32) :: dummy
          type(c_ptr) :: c_dev_ptr
          complex(real32), pointer, contiguous :: fptr(:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3)), stat = ierr )
#endif
          endif
          !
      end subroutine sp_devxlib_malloc_alloc_c3d
      !
      module subroutine sp_devxlib_malloc_alloc_c4d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          complex(real32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:)
          integer(int64), intent(in)           :: dimensions(4)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(4)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          complex(real32) :: dummy
          type(c_ptr) :: c_dev_ptr
          complex(real32), pointer, contiguous :: fptr(:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4)), stat = ierr )
#endif
          endif
          !
      end subroutine sp_devxlib_malloc_alloc_c4d
      !
      module subroutine sp_devxlib_malloc_alloc_c5d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          complex(real32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(5)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(5)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          complex(real32) :: dummy
          type(c_ptr) :: c_dev_ptr
          complex(real32), pointer, contiguous :: fptr(:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5)), stat = ierr )
#endif
          endif
          !
      end subroutine sp_devxlib_malloc_alloc_c5d
      !
      module subroutine sp_devxlib_malloc_alloc_c6d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          complex(real32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(6)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(6)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          complex(real32) :: dummy
          type(c_ptr) :: c_dev_ptr
          complex(real32), pointer, contiguous :: fptr(:,:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5), dimensions(6)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                           lbounds(6):lbounds(6)+dimensions(6)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                                                             lbounds(6):lbounds(6)+dimensions(6)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5), dimensions(6)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5),&
                                                     dimensions(6)), stat = ierr )
#endif
          endif
          !
      end subroutine sp_devxlib_malloc_alloc_c6d
      !
      module subroutine dp_devxlib_malloc_alloc_c1d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          complex(real64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:)
          integer(int64), intent(in)           :: dimensions
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          complex(real64) :: dummy
          type(c_ptr) :: c_dev_ptr
          complex(real64), pointer, contiguous :: fptr(:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions])
                   dev_ptr(lbounds:lbounds+dimensions-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds:lbounds+dimensions-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions), stat = ierr )
#endif
          endif
          !
      end subroutine dp_devxlib_malloc_alloc_c1d
      !
      module subroutine dp_devxlib_malloc_alloc_c2d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          complex(real64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:)
          integer(int64), intent(in)           :: dimensions(2)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(2)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          complex(real64) :: dummy
          type(c_ptr) :: c_dev_ptr
          complex(real64), pointer, contiguous :: fptr(:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2)), stat = ierr )
#endif
          endif
          !
      end subroutine dp_devxlib_malloc_alloc_c2d
      !
      module subroutine dp_devxlib_malloc_alloc_c3d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          complex(real64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:)
          integer(int64), intent(in)           :: dimensions(3)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(3)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          complex(real64) :: dummy
          type(c_ptr) :: c_dev_ptr
          complex(real64), pointer, contiguous :: fptr(:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3)), stat = ierr )
#endif
          endif
          !
      end subroutine dp_devxlib_malloc_alloc_c3d
      !
      module subroutine dp_devxlib_malloc_alloc_c4d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          complex(real64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:)
          integer(int64), intent(in)           :: dimensions(4)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(4)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          complex(real64) :: dummy
          type(c_ptr) :: c_dev_ptr
          complex(real64), pointer, contiguous :: fptr(:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4)), stat = ierr )
#endif
          endif
          !
      end subroutine dp_devxlib_malloc_alloc_c4d
      !
      module subroutine dp_devxlib_malloc_alloc_c5d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          complex(real64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(5)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(5)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          complex(real64) :: dummy
          type(c_ptr) :: c_dev_ptr
          complex(real64), pointer, contiguous :: fptr(:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5)), stat = ierr )
#endif
          endif
          !
      end subroutine dp_devxlib_malloc_alloc_c5d
      !
      module subroutine dp_devxlib_malloc_alloc_c6d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          complex(real64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(6)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(6)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          complex(real64) :: dummy
          type(c_ptr) :: c_dev_ptr
          complex(real64), pointer, contiguous :: fptr(:,:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5), dimensions(6)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                           lbounds(6):lbounds(6)+dimensions(6)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                                                             lbounds(6):lbounds(6)+dimensions(6)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5), dimensions(6)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5),&
                                                     dimensions(6)), stat = ierr )
#endif
          endif
          !
      end subroutine dp_devxlib_malloc_alloc_c6d
      !
      module subroutine i4_devxlib_malloc_alloc_i1d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          integer(int32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:)
          integer(int64), intent(in)           :: dimensions
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          integer(int32) :: dummy
          type(c_ptr) :: c_dev_ptr
          integer(int32), pointer, contiguous :: fptr(:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions])
                   dev_ptr(lbounds:lbounds+dimensions-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds:lbounds+dimensions-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions), stat = ierr )
#endif
          endif
          !
      end subroutine i4_devxlib_malloc_alloc_i1d
      !
      module subroutine i4_devxlib_malloc_alloc_i2d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          integer(int32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:)
          integer(int64), intent(in)           :: dimensions(2)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(2)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          integer(int32) :: dummy
          type(c_ptr) :: c_dev_ptr
          integer(int32), pointer, contiguous :: fptr(:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2)), stat = ierr )
#endif
          endif
          !
      end subroutine i4_devxlib_malloc_alloc_i2d
      !
      module subroutine i4_devxlib_malloc_alloc_i3d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          integer(int32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:)
          integer(int64), intent(in)           :: dimensions(3)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(3)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          integer(int32) :: dummy
          type(c_ptr) :: c_dev_ptr
          integer(int32), pointer, contiguous :: fptr(:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3)), stat = ierr )
#endif
          endif
          !
      end subroutine i4_devxlib_malloc_alloc_i3d
      !
      module subroutine i4_devxlib_malloc_alloc_i4d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          integer(int32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:)
          integer(int64), intent(in)           :: dimensions(4)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(4)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          integer(int32) :: dummy
          type(c_ptr) :: c_dev_ptr
          integer(int32), pointer, contiguous :: fptr(:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4)), stat = ierr )
#endif
          endif
          !
      end subroutine i4_devxlib_malloc_alloc_i4d
      !
      module subroutine i4_devxlib_malloc_alloc_i5d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          integer(int32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(5)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(5)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          integer(int32) :: dummy
          type(c_ptr) :: c_dev_ptr
          integer(int32), pointer, contiguous :: fptr(:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5)), stat = ierr )
#endif
          endif
          !
      end subroutine i4_devxlib_malloc_alloc_i5d
      !
      module subroutine i4_devxlib_malloc_alloc_i6d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          integer(int32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(6)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(6)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          integer(int32) :: dummy
          type(c_ptr) :: c_dev_ptr
          integer(int32), pointer, contiguous :: fptr(:,:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5), dimensions(6)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                           lbounds(6):lbounds(6)+dimensions(6)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                                                             lbounds(6):lbounds(6)+dimensions(6)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5), dimensions(6)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5),&
                                                     dimensions(6)), stat = ierr )
#endif
          endif
          !
      end subroutine i4_devxlib_malloc_alloc_i6d
      !
      module subroutine i8_devxlib_malloc_alloc_i1d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          integer(int64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:)
          integer(int64), intent(in)           :: dimensions
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          integer(int64) :: dummy
          type(c_ptr) :: c_dev_ptr
          integer(int64), pointer, contiguous :: fptr(:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions])
                   dev_ptr(lbounds:lbounds+dimensions-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds:lbounds+dimensions-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions), stat = ierr )
#endif
          endif
          !
      end subroutine i8_devxlib_malloc_alloc_i1d
      !
      module subroutine i8_devxlib_malloc_alloc_i2d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          integer(int64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:)
          integer(int64), intent(in)           :: dimensions(2)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(2)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          integer(int64) :: dummy
          type(c_ptr) :: c_dev_ptr
          integer(int64), pointer, contiguous :: fptr(:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2)), stat = ierr )
#endif
          endif
          !
      end subroutine i8_devxlib_malloc_alloc_i2d
      !
      module subroutine i8_devxlib_malloc_alloc_i3d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          integer(int64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:)
          integer(int64), intent(in)           :: dimensions(3)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(3)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          integer(int64) :: dummy
          type(c_ptr) :: c_dev_ptr
          integer(int64), pointer, contiguous :: fptr(:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3)), stat = ierr )
#endif
          endif
          !
      end subroutine i8_devxlib_malloc_alloc_i3d
      !
      module subroutine i8_devxlib_malloc_alloc_i4d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          integer(int64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:)
          integer(int64), intent(in)           :: dimensions(4)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(4)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          integer(int64) :: dummy
          type(c_ptr) :: c_dev_ptr
          integer(int64), pointer, contiguous :: fptr(:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4)), stat = ierr )
#endif
          endif
          !
      end subroutine i8_devxlib_malloc_alloc_i4d
      !
      module subroutine i8_devxlib_malloc_alloc_i5d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          integer(int64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(5)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(5)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          integer(int64) :: dummy
          type(c_ptr) :: c_dev_ptr
          integer(int64), pointer, contiguous :: fptr(:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5)), stat = ierr )
#endif
          endif
          !
      end subroutine i8_devxlib_malloc_alloc_i5d
      !
      module subroutine i8_devxlib_malloc_alloc_i6d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          integer(int64), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(6)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(6)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          integer(int64) :: dummy
          type(c_ptr) :: c_dev_ptr
          integer(int64), pointer, contiguous :: fptr(:,:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5), dimensions(6)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                           lbounds(6):lbounds(6)+dimensions(6)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                                                             lbounds(6):lbounds(6)+dimensions(6)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5), dimensions(6)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5),&
                                                     dimensions(6)), stat = ierr )
#endif
          endif
          !
      end subroutine i8_devxlib_malloc_alloc_i6d
      !
      module subroutine l4_devxlib_malloc_alloc_l1d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          logical(int32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:)
          integer(int64), intent(in)           :: dimensions
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          logical(int32) :: dummy
          type(c_ptr) :: c_dev_ptr
          logical(int32), pointer, contiguous :: fptr(:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions])
                   dev_ptr(lbounds:lbounds+dimensions-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds:lbounds+dimensions-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * dimensions / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * dimensions / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions), stat = ierr )
#endif
          endif
          !
      end subroutine l4_devxlib_malloc_alloc_l1d
      !
      module subroutine l4_devxlib_malloc_alloc_l2d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          logical(int32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:)
          integer(int64), intent(in)           :: dimensions(2)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(2)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          logical(int32) :: dummy
          type(c_ptr) :: c_dev_ptr
          logical(int32), pointer, contiguous :: fptr(:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2)), stat = ierr )
#endif
          endif
          !
      end subroutine l4_devxlib_malloc_alloc_l2d
      !
      module subroutine l4_devxlib_malloc_alloc_l3d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          logical(int32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:)
          integer(int64), intent(in)           :: dimensions(3)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(3)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          logical(int32) :: dummy
          type(c_ptr) :: c_dev_ptr
          logical(int32), pointer, contiguous :: fptr(:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3)), stat = ierr )
#endif
          endif
          !
      end subroutine l4_devxlib_malloc_alloc_l3d
      !
      module subroutine l4_devxlib_malloc_alloc_l4d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          logical(int32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:)
          integer(int64), intent(in)           :: dimensions(4)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(4)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          logical(int32) :: dummy
          type(c_ptr) :: c_dev_ptr
          logical(int32), pointer, contiguous :: fptr(:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4)), stat = ierr )
#endif
          endif
          !
      end subroutine l4_devxlib_malloc_alloc_l4d
      !
      module subroutine l4_devxlib_malloc_alloc_l5d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          logical(int32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(5)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(5)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          logical(int32) :: dummy
          type(c_ptr) :: c_dev_ptr
          logical(int32), pointer, contiguous :: fptr(:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5)), stat = ierr )
#endif
          endif
          !
      end subroutine l4_devxlib_malloc_alloc_l5d
      !
      module subroutine l4_devxlib_malloc_alloc_l6d(dev_ptr, dimensions, ierr, device_id, lbounds)
          implicit none
          !
          logical(int32), pointer, contiguous, intent(out) DEV_ATTR :: dev_ptr(:,:,:,:,:,:)
          integer(int64), intent(in)           :: dimensions(6)
          integer(int32), intent(out)          :: ierr
          integer(int32), intent(in), optional :: device_id
          integer(int64), intent(in), optional :: lbounds(6)
#if defined __DXL_CUDAF
          integer(kind=c_int) :: info
#endif
          logical(int32) :: dummy
          type(c_ptr) :: c_dev_ptr
          logical(int32), pointer, contiguous :: fptr(:,:,:,:,:,:)
#if defined __DXL_OPENMP_GPU
          integer(int32) :: device_id_local
#endif
          !
#if defined __DXL_CUDAF
          if (present(device_id)) ierr = cudaSetDevice(device_id)
#elif defined __DXL_OPENMP_GPU
          if (present(device_id)) then
             device_id_local = device_id
          else
             device_id_local = omp_get_default_device()
          endif
#endif
          !
          if (present(lbounds)) then
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, fptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                     dimensions(5), dimensions(6)])
                   dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                           lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                           lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                           lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                           lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                           lbounds(6):lbounds(6)+dimensions(6)-1_int64) => fptr
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(lbounds(1):lbounds(1)+dimensions(1)-1_int64,&
                                                             lbounds(2):lbounds(2)+dimensions(2)-1_int64,&
                                                             lbounds(3):lbounds(3)+dimensions(3)-1_int64,&
                                                             lbounds(4):lbounds(4)+dimensions(4)-1_int64,&
                                                             lbounds(5):lbounds(5)+dimensions(5)-1_int64,&
                                                             lbounds(6):lbounds(6)+dimensions(6)-1_int64), stat = ierr )
#endif
          else
             !
#if defined __DXL_CUDAF || defined __DXL_OPENACC || defined __DXL_OPENMP_GPU
             if (.not. associated(dev_ptr)) then
#  if defined __DXL_CUDAF
                info = cudaMalloc_f(c_dev_ptr,int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENACC
                c_dev_ptr = acc_malloc_f(int(storage_size(dummy) * product(dimensions) / 8, c_size_t))
#  elif defined __DXL_OPENMP_GPU
                c_dev_ptr = omp_target_alloc(int(storage_size(dummy) * product(dimensions) / 8, c_size_t), device_id_local)
#  endif
                if (c_associated(c_dev_ptr)) then
                   call c_f_pointer(c_dev_ptr, dev_ptr, [dimensions(1), dimensions(2), dimensions(3), dimensions(4), &
                                                         dimensions(5), dimensions(6)])
                   ierr = 0
                else
                   ierr = 1000
                endif
             endif
#else
             if (.not.associated(dev_ptr)) allocate( dev_ptr(dimensions(1),&
                                                     dimensions(2),&
                                                     dimensions(3),&
                                                     dimensions(4),&
                                                     dimensions(5),&
                                                     dimensions(6)), stat = ierr )
#endif
          endif
          !
      end subroutine l4_devxlib_malloc_alloc_l6d
      !

endsubmodule devxlib_malloc_alloc