accel.F90

Go to the documentation of this file.

!! Copyright (C) 2010-2016 X. Andrade
!!
!! This program is free software; you can redistribute it and/or modify
!! it under the terms of the GNU General Public License as published by
!! the Free Software Foundation; either version 2, or (at your option)
!! any later version.
!!
!! This program is distributed in the hope that it will be useful,
!! but WITHOUT ANY WARRANTY; without even the implied warranty of
!! MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
!! GNU General Public License for more details.
!!
!! You should have received a copy of the GNU General Public License
!! along with this program; if not, write to the Free Software
!! Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA
!! 02110-1301, USA.
!!
 
#include "global.h"
 
#if defined(HAVE_CUDA)
#define HAVE_ACCEL 1
#endif
 
module accel_oct_m
  use alloc_cache_oct_m
  use cuda_oct_m
  use debug_oct_m
  use global_oct_m
  use iso_c_binding, only: c_size_t
  use, intrinsic :: iso_fortran_env
  use loct_oct_m
  use math_oct_m
  use messages_oct_m
  use mpi_oct_m
  use namespace_oct_m
  use types_oct_m
  use parser_oct_m
  use profiling_oct_m
  use unit_system_oct_m
  use varinfo_oct_m
  use string_oct_m
 
  implicit none
 
  private
 
  public ::                       &
    accel_context_t,              &
    accel_device_t,               &
    accel_mem_t,                  &
    accel_kernel_t,               &
    accel_t,                      &
    accel_is_enabled,             &
    accel_allow_cpu_only,         &
    accel_init,                   &
    accel_end,                    &
    accel_padded_size,            &
    accel_kernel_start_call,      &
    accel_kernel_build,           &
    accel_create_buffer,          &
    accel_write_buffer,           &
    accel_read_buffer,            &
    accel_release_buffer,         &
    accel_ensure_buffer_size,     &
    accel_buffer_is_allocated,    &
    accel_finish,                 &
    accel_set_kernel_arg,         &
    accel_max_workgroup_size,     &
    accel_kernel_workgroup_size,  &
    accel_kernel_run,             &
    accel_set_buffer_to_zero,     &
    accel_use_shared_mem,         &
    accel_local_memory_size,      &
    accel_global_memory_size,     &
    accel_max_size_per_dim,       &
    accel_get_device_pointer,     &
    daccel_get_pointer_with_offset,&
    zaccel_get_pointer_with_offset,&
    accel_clean_pointer,          &
    accel_set_stream,             &
    accel_get_stream,             &
    accel_synchronize_all_streams,&
    accel_get_unfolded_size,      &
    accel_create_blas_alpha_beta_buffer, &
    accel_release_blas_alpha_beta_buffer
 
  integer, public, parameter ::                 &
    ACCEL_MEM_READ_ONLY  = 0,                   &
    accel_mem_read_write = 1,                   &
    accel_mem_write_only = 2
 
  type accel_context_t
    ! Components are public by default
#if defined(HAVE_CUDA)
    type(c_ptr)      :: cuda_context
#else
    integer          :: dummy
#endif
  end type accel_context_t
 
  type accel_device_t
    ! Components are public by default
#if defined(HAVE_CUDA)
    type(c_ptr)      :: cuda_device
#else
    integer         :: dummy
#endif
  end type accel_device_t
 
  type accel_t
    ! Components are public by default
    type(accel_context_t)  :: context
    type(accel_device_t)   :: device
    type(c_ptr)            :: cublas_handle
    type(c_ptr)            :: cuda_stream
    type(c_ptr)            :: module_map
    integer                :: max_workgroup_size
    integer(int64)         :: local_memory_size
    integer(int64)         :: global_memory_size
    logical                :: enabled
    logical                :: allow_CPU_only
    logical                :: shared_mem
    logical                :: cuda_mpi
    integer                :: warp_size
    integer(int64)         :: initialize_buffers
    character(len=32)      :: debug_flag
    integer(int64)         :: max_block_dim(3)
    integer(int64)         :: max_grid_dim(3)
  end type accel_t
 
  type accel_mem_t
    ! Components are public by default
    type(c_ptr)            :: mem
    integer(c_size_t)      :: size = 0
    type(type_t)           :: type
    integer                :: flags = 0
    logical                :: allocated = .false.
  end type accel_mem_t
 
  type accel_kernel_t
    ! Components are public by default
#ifdef HAVE_CUDA
    type(c_ptr)                   :: cuda_kernel
    type(c_ptr)                   :: cuda_module
    type(c_ptr)                   :: arguments
#endif
    integer(int64)                :: cuda_shared_mem
    logical                       :: initialized = .false.
    type(accel_kernel_t), pointer :: next
    integer                       :: arg_count
    character(len=128)            :: kernel_name
  end type accel_kernel_t
 
  type(accel_t), public :: accel
 
  ! Global variables defined on device
  type(accel_mem_t), public, save :: zM_0_buffer, zM_1_buffer
  type(accel_mem_t), public, save :: dM_0_buffer, dM_1_buffer
 
  ! the kernels
  type(accel_kernel_t), public, target, save :: kernel_vpsi
  type(accel_kernel_t), public, target, save :: kernel_vpsi_complex
  type(accel_kernel_t), public, target, save :: kernel_vpsi_spinors
  type(accel_kernel_t), public, target, save :: kernel_vpsi_spinors_complex
  type(accel_kernel_t), public, target, save :: kernel_daxpy
  type(accel_kernel_t), public, target, save :: kernel_zaxpy
  type(accel_kernel_t), public, target, save :: kernel_copy
  type(accel_kernel_t), public, target, save :: kernel_copy_complex_to_real
  type(accel_kernel_t), public, target, save :: kernel_copy_real_to_complex
  type(accel_kernel_t), public, target, save :: dpack
  type(accel_kernel_t), public, target, save :: zpack
  type(accel_kernel_t), public, target, save :: dunpack
  type(accel_kernel_t), public, target, save :: zunpack
  type(accel_kernel_t), public, target, save :: kernel_ghost_reorder
  type(accel_kernel_t), public, target, save :: kernel_density_real
  type(accel_kernel_t), public, target, save :: kernel_density_complex
  type(accel_kernel_t), public, target, save :: kernel_density_spinors
  type(accel_kernel_t), public, target, save :: kernel_phase
  type(accel_kernel_t), public, target, save :: kernel_phase_spiral
  type(accel_kernel_t), public, target, save :: dkernel_dot_matrix
  type(accel_kernel_t), public, target, save :: zkernel_dot_matrix
  type(accel_kernel_t), public, target, save :: zkernel_dot_matrix_spinors
  type(accel_kernel_t), public, target, save :: dkernel_batch_axpy
  type(accel_kernel_t), public, target, save :: zkernel_batch_axpy
  type(accel_kernel_t), public, target, save :: dkernel_ax_function_py
  type(accel_kernel_t), public, target, save :: zkernel_ax_function_py
  type(accel_kernel_t), public, target, save :: dkernel_batch_dotp
  type(accel_kernel_t), public, target, save :: zkernel_batch_dotp
  type(accel_kernel_t), public, target, save :: dzmul
  type(accel_kernel_t), public, target, save :: zzmul
 
  interface accel_padded_size
    module procedure accel_padded_size_i8, accel_padded_size_i4
  end interface accel_padded_size
 
  interface accel_create_buffer
    module procedure accel_create_buffer_4, accel_create_buffer_8
  end interface accel_create_buffer
 
  interface accel_kernel_run
    module procedure accel_kernel_run_4, accel_kernel_run_8
  end interface accel_kernel_run
 
  interface accel_set_buffer_to_zero
    module procedure accel_set_buffer_to_zero_i8, accel_set_buffer_to_zero_i4
  end interface accel_set_buffer_to_zero
 
  interface accel_write_buffer
    module procedure iaccel_write_buffer_single, laccel_write_buffer_single, daccel_write_buffer_single, zaccel_write_buffer_single
    module procedure iaccel_write_buffer_0, laccel_write_buffer_0, daccel_write_buffer_0, zaccel_write_buffer_0
    module procedure iaccel_write_buffer_1, laccel_write_buffer_1, daccel_write_buffer_1, zaccel_write_buffer_1
    module procedure iaccel_write_buffer_2, laccel_write_buffer_2, daccel_write_buffer_2, zaccel_write_buffer_2
    module procedure iaccel_write_buffer_3, laccel_write_buffer_3, daccel_write_buffer_3, zaccel_write_buffer_3
    module procedure iaccel_write_buffer_4, laccel_write_buffer_4, daccel_write_buffer_4, zaccel_write_buffer_4
    module procedure iaccel_write_buffer_5, laccel_write_buffer_5, daccel_write_buffer_5, zaccel_write_buffer_5
    module procedure iaccel_write_buffer_6, laccel_write_buffer_6, daccel_write_buffer_6, zaccel_write_buffer_6
    module procedure iaccel_write_buffer_0_int32, laccel_write_buffer_0_int32, daccel_write_buffer_0_int32, &
      zaccel_write_buffer_0_int32
    module procedure iaccel_write_buffer_1_int32, laccel_write_buffer_1_int32, daccel_write_buffer_1_int32, &
      zaccel_write_buffer_1_int32
    module procedure iaccel_write_buffer_2_int32, laccel_write_buffer_2_int32, daccel_write_buffer_2_int32, &
      zaccel_write_buffer_2_int32
    module procedure iaccel_write_buffer_3_int32, laccel_write_buffer_3_int32, daccel_write_buffer_3_int32, &
      zaccel_write_buffer_3_int32
    module procedure iaccel_write_buffer_4_int32, laccel_write_buffer_4_int32, daccel_write_buffer_4_int32, &
      zaccel_write_buffer_4_int32
    module procedure iaccel_write_buffer_5_int32, laccel_write_buffer_5_int32, daccel_write_buffer_5_int32, &
      zaccel_write_buffer_5_int32
    module procedure iaccel_write_buffer_6_int32, laccel_write_buffer_6_int32, daccel_write_buffer_6_int32, &
      zaccel_write_buffer_6_int32
  end interface accel_write_buffer
 
  interface accel_read_buffer
    module procedure iaccel_read_buffer_0, laccel_read_buffer_0, daccel_read_buffer_0, zaccel_read_buffer_0
    module procedure iaccel_read_buffer_1, laccel_read_buffer_1, daccel_read_buffer_1, zaccel_read_buffer_1
    module procedure iaccel_read_buffer_2, laccel_read_buffer_2, daccel_read_buffer_2, zaccel_read_buffer_2
    module procedure iaccel_read_buffer_3, laccel_read_buffer_3, daccel_read_buffer_3, zaccel_read_buffer_3
    module procedure iaccel_read_buffer_4, laccel_read_buffer_4, daccel_read_buffer_4, zaccel_read_buffer_4
    module procedure iaccel_read_buffer_5, laccel_read_buffer_5, daccel_read_buffer_5, zaccel_read_buffer_5
    module procedure iaccel_read_buffer_6, laccel_read_buffer_6, daccel_read_buffer_6, zaccel_read_buffer_6
    module procedure iaccel_read_buffer_0_int32, laccel_read_buffer_0_int32, daccel_read_buffer_0_int32, zaccel_read_buffer_0_int32
    module procedure iaccel_read_buffer_1_int32, laccel_read_buffer_1_int32, daccel_read_buffer_1_int32, zaccel_read_buffer_1_int32
    module procedure iaccel_read_buffer_2_int32, laccel_read_buffer_2_int32, daccel_read_buffer_2_int32, zaccel_read_buffer_2_int32
    module procedure iaccel_read_buffer_3_int32, laccel_read_buffer_3_int32, daccel_read_buffer_3_int32, zaccel_read_buffer_3_int32
    module procedure iaccel_read_buffer_4_int32, laccel_read_buffer_4_int32, daccel_read_buffer_4_int32, zaccel_read_buffer_4_int32
    module procedure iaccel_read_buffer_5_int32, laccel_read_buffer_5_int32, daccel_read_buffer_5_int32, zaccel_read_buffer_5_int32
    module procedure iaccel_read_buffer_6_int32, laccel_read_buffer_6_int32, daccel_read_buffer_6_int32, zaccel_read_buffer_6_int32
  end interface accel_read_buffer
 
  interface accel_set_kernel_arg
    module procedure                       &
      accel_set_kernel_arg_buffer,  &
      iaccel_set_kernel_arg_data,   &
      laccel_set_kernel_arg_data,   &
      daccel_set_kernel_arg_data,   &
      zaccel_set_kernel_arg_data,   &
      accel_set_kernel_arg_local
  end interface accel_set_kernel_arg
 
  interface accel_get_device_pointer
    module procedure iaccel_get_device_pointer_1, laccel_get_device_pointer_1
    module procedure iaccel_get_device_pointer_2, laccel_get_device_pointer_2
    module procedure iaccel_get_device_pointer_3, laccel_get_device_pointer_3
    module procedure daccel_get_device_pointer_1, zaccel_get_device_pointer_1
    module procedure daccel_get_device_pointer_2, zaccel_get_device_pointer_2
    module procedure daccel_get_device_pointer_3, zaccel_get_device_pointer_3
    module procedure iaccel_get_device_pointer_1l, laccel_get_device_pointer_1l
    module procedure iaccel_get_device_pointer_2l, laccel_get_device_pointer_2l
    module procedure iaccel_get_device_pointer_3l, laccel_get_device_pointer_3l
    module procedure daccel_get_device_pointer_1l, zaccel_get_device_pointer_1l
    module procedure daccel_get_device_pointer_2l, zaccel_get_device_pointer_2l
    module procedure daccel_get_device_pointer_3l, zaccel_get_device_pointer_3l
  end interface accel_get_device_pointer
 
  interface accel_create_blas_alpha_beta_buffer
    module procedure                                       &
      daccel_create_blas_alpha_beta_buffer,                &
      zaccel_create_blas_alpha_beta_buffer,                &
      iaccel_create_blas_alpha_beta_buffer,                &
      laccel_create_blas_alpha_beta_buffer
  end interface accel_create_blas_alpha_beta_buffer
 
  interface accel_release_blas_alpha_beta_buffer
    module procedure                                       &
      daccel_release_blas_alpha_beta_buffer,               &
      zaccel_release_blas_alpha_beta_buffer,               &
      iaccel_release_blas_alpha_beta_buffer,               &
      laccel_release_blas_alpha_beta_buffer
  end interface accel_release_blas_alpha_beta_buffer
 
  integer :: buffer_alloc_count
  integer(int64) :: allocated_mem
  type(accel_kernel_t), pointer :: head
  type(alloc_cache_t) :: memcache
 
contains
 
  pure logical function accel_is_enabled() result(enabled)
#ifdef HAVE_ACCEL
    enabled = accel%enabled
#else
    enabled = .false.
#endif
  end function accel_is_enabled
 
  ! ------------------------------------------
 
  pure logical function accel_allow_cpu_only() result(allow)
#ifdef HAVE_ACCEL
    allow = accel%allow_CPU_only
#else
    allow = .true.
#endif
  end function accel_allow_cpu_only
 
  ! ------------------------------------------
 
  subroutine accel_init(base_grp, namespace)
    type(mpi_grp_t),     intent(inout) :: base_grp
    type(namespace_t),   intent(in)    :: namespace
 
    logical  :: disable, default, run_benchmark
    integer  :: idevice
#ifdef HAVE_CUDA
    integer :: dim
#ifdef HAVE_MPI
    character(len=256) :: sys_name
#endif
#endif
 
    push_sub(accel_init)
 
    buffer_alloc_count = 0
 
    !%Variable DisableAccel
    !%Type logical
    !%Default yes
    !%Section Execution::Accel
    !%Description
    !% If Octopus was compiled with CUDA support, it will
    !% try to initialize and use an accelerator device. By setting this
    !% variable to <tt>yes</tt> you force Octopus not to use an accelerator even it is available.
    !%End
    call messages_obsolete_variable(namespace, 'DisableOpenCL', 'DisableAccel')
#ifdef HAVE_ACCEL
    default = .false.
#else
    default = .true.
#endif
    call parse_variable(namespace, 'DisableAccel', default, disable)
    accel%enabled = .not. disable
 
#ifndef HAVE_ACCEL
    if (accel%enabled) then
      message(1) = 'Octopus was compiled without Cuda support.'
      call messages_fatal(1)
    end if
#endif
 
    if (.not. accel_is_enabled()) then
      pop_sub(accel_init)
      return
    end if
 
    call messages_obsolete_variable(namespace, 'AccelPlatform')
    call messages_obsolete_variable(namespace, 'OpenCLPlatform', 'AccelPlatform')
 
    !%Variable AccelDevice
    !%Type integer
    !%Default 0
    !%Section Execution::Accel
    !%Description
    !% This variable selects the GPU that Octopus will use. You can specify a
    !% numerical id to select a specific device.
    !%
    !% In case of MPI enabled runs devices are distributed in a round robin fashion,
    !% starting at this value.
    !%End
    call parse_variable(namespace, 'AccelDevice', 0, idevice)
 
    call messages_obsolete_variable(namespace, 'OpenCLDevice', 'AccelDevice')
 
    if (idevice < 0) then
      call messages_write('Invalid AccelDevice')
      call messages_fatal()
    end if
 
    call messages_print_with_emphasis(msg="GPU acceleration", namespace=namespace)
 
#ifdef HAVE_CUDA
    if (idevice<0) idevice = 0
    call cuda_init(accel%context%cuda_context, accel%device%cuda_device, accel%cuda_stream, &
      idevice, base_grp%rank)
#ifdef HAVE_MPI
    call loct_sysname(sys_name)
    write(message(1), '(A,I5,A,I5,2A)') "Rank ", base_grp%rank, " uses device number ", idevice, &
      " on ", trim(sys_name)
    call messages_info(1, all_nodes = .true.)
#endif
 
    accel%shared_mem = .true.
 
    call cublas_init(accel%cublas_handle, accel%cuda_stream)
#endif
 
 
    ! Get some device information that we will need later
#ifdef HAVE_CUDA
    call cuda_device_total_memory(accel%device%cuda_device, accel%global_memory_size)
    call cuda_device_shared_memory(accel%device%cuda_device, accel%local_memory_size)
    call cuda_device_max_threads_per_block(accel%device%cuda_device, accel%max_workgroup_size)
    call cuda_device_get_warpsize(accel%device%cuda_device, accel%warp_size)
    call cuda_device_max_block_dim_x(accel%device%cuda_device, dim)
    accel%max_block_dim(1) = int(dim, int64)
    call cuda_device_max_block_dim_y(accel%device%cuda_device, dim)
    accel%max_block_dim(2) = int(dim, int64)
    call cuda_device_max_block_dim_z(accel%device%cuda_device, dim)
    accel%max_block_dim(3) = int(dim, int64)
    call cuda_device_max_grid_dim_x(accel%device%cuda_device, dim)
    accel%max_grid_dim(1) = int(dim, int64)
    call cuda_device_max_grid_dim_y(accel%device%cuda_device, dim)
    accel%max_grid_dim(2) = int(dim, int64)
    call cuda_device_max_grid_dim_z(accel%device%cuda_device, dim)
    accel%max_grid_dim(3) = int(dim, int64)
#endif
 
    if (base_grp%is_root()) call device_info()
 
    ! initialize the cache used to speed up allocations
    call alloc_cache_init(memcache, nint(0.25_real64*accel%global_memory_size, int64))
 
    ! now initialize the kernels
    call accel_kernel_global_init()
 
#if defined(HAVE_HIP)
    accel%debug_flag = "-g"
#elif defined(HAVE_CUDA)
    accel%debug_flag = "-lineinfo"
#endif
 
    call accel_kernel_start_call(kernel_vpsi, 'vpsi.cu', "vpsi")
    call accel_kernel_start_call(kernel_vpsi_complex, 'vpsi.cu', "vpsi_complex")
    call accel_kernel_start_call(kernel_vpsi_spinors, 'vpsi.cu', "vpsi_spinors")
    call accel_kernel_start_call(kernel_vpsi_spinors_complex, 'vpsi.cu', "vpsi_spinors_complex")
    call accel_kernel_start_call(kernel_daxpy, 'axpy.cu', "daxpy", flags = '-DRTYPE_DOUBLE')
    call accel_kernel_start_call(kernel_zaxpy, 'axpy.cu', "zaxpy", flags = '-DRTYPE_COMPLEX')
    call accel_kernel_start_call(dkernel_batch_axpy, 'axpy.cu', "dbatch_axpy_function", &
      flags = ' -DRTYPE_DOUBLE')
    call accel_kernel_start_call(zkernel_batch_axpy, 'axpy.cu', "zbatch_axpy_function", &
      flags = '-DRTYPE_COMPLEX')
    call accel_kernel_start_call(dkernel_ax_function_py, 'axpy.cu', "dbatch_ax_function_py", &
      flags = '-DRTYPE_DOUBLE')
    call accel_kernel_start_call(zkernel_ax_function_py, 'axpy.cu', "zbatch_ax_function_py", &
      flags = '-DRTYPE_COMPLEX')
    call accel_kernel_start_call(dkernel_batch_dotp, 'mesh_batch_single.cu', "dbatch_mf_dotp")
    call accel_kernel_start_call(zkernel_batch_dotp, 'mesh_batch_single.cu', "zbatch_mf_dotp")
    call accel_kernel_start_call(dpack, 'pack.cu', "dpack")
    call accel_kernel_start_call(zpack, 'pack.cu', "zpack")
    call accel_kernel_start_call(dunpack, 'pack.cu', "dunpack")
    call accel_kernel_start_call(zunpack, 'pack.cu', "zunpack")
    call accel_kernel_start_call(kernel_copy, 'copy.cu', "copy")
    call accel_kernel_start_call(kernel_copy_complex_to_real, 'copy.cu', "copy_complex_to_real")
    call accel_kernel_start_call(kernel_copy_real_to_complex, 'copy.cu', "copy_real_to_complex")
    call accel_kernel_start_call(kernel_ghost_reorder, 'ghost.cu', "ghost_reorder")
    call accel_kernel_start_call(kernel_density_real, 'density.cu', "density_real")
    call accel_kernel_start_call(kernel_density_complex, 'density.cu', "density_complex")
    call accel_kernel_start_call(kernel_density_spinors, 'density.cu', "density_spinors")
    call accel_kernel_start_call(kernel_phase, 'phase.cu', "phase")
    call accel_kernel_start_call(dkernel_dot_matrix, 'mesh_batch.cu', "ddot_matrix")
    call accel_kernel_start_call(zkernel_dot_matrix, 'mesh_batch.cu', "zdot_matrix")
    call accel_kernel_start_call(zkernel_dot_matrix_spinors, 'mesh_batch.cu', "zdot_matrix_spinors")
 
 
    call accel_kernel_start_call(dzmul, 'mul.cu', "dzmul", flags = '-DRTYPE_DOUBLE')
    call accel_kernel_start_call(zzmul, 'mul.cu', "zzmul", flags = '-DRTYPE_COMPLEX')
 
    ! Define global buffers
    if(.not. accel_buffer_is_allocated(zm_0_buffer)) then
      call accel_create_buffer(zm_0_buffer, accel_mem_read_only, type_cmplx, 1)
      call accel_write_buffer(zm_0_buffer, m_z0)
    end if
    if(.not. accel_buffer_is_allocated(zm_1_buffer)) then
      call accel_create_buffer(zm_1_buffer, accel_mem_read_only, type_cmplx, 1)
      call accel_write_buffer(zm_1_buffer, m_z1)
    end if
    if(.not. accel_buffer_is_allocated(dm_0_buffer)) then
      call accel_create_buffer(dm_0_buffer, accel_mem_read_only, type_float, 1)
      call accel_write_buffer(dm_0_buffer, m_zero)
    end if
    if(.not. accel_buffer_is_allocated(dm_1_buffer)) then
      call accel_create_buffer(dm_1_buffer, accel_mem_read_only, type_float, 1)
      call accel_write_buffer(dm_1_buffer, m_one)
    end if
 
 
    !%Variable AccelBenchmark
    !%Type logical
    !%Default no
    !%Section Execution::Accel
    !%Description
    !% If this variable is set to yes, Octopus will run some
    !% routines to benchmark the performance of the accelerator device.
    !%End
    call parse_variable(namespace, 'AccelBenchmark', .false., run_benchmark)
 
    call messages_obsolete_variable(namespace, 'OpenCLBenchmark', 'AccelBenchmark')
 
    if (run_benchmark) then
      call accel_check_bandwidth()
    end if
 
    !%Variable GPUAwareMPI
    !%Type logical
    !%Section Execution::Accel
    !%Description
    !% If Octopus was compiled with GPU support and MPI support and if the MPI
    !% implementation is GPU-aware (i.e., it supports communication using device pointers),
    !% this switch can be set to true to use the GPU-aware MPI features. The advantage
    !% of this approach is that it can do, e.g., peer-to-peer copies between devices without
    !% going through the host memory.
    !% The default is false, except when the configure switch --enable-cudampi is set, in which
    !% case this variable is set to true.
    !%End
#ifdef HAVE_CUDA_MPI
    default = .true.
#else
    default = .false.
#endif
    call parse_variable(namespace, 'GPUAwareMPI', default, accel%cuda_mpi)
    if (accel%cuda_mpi) then
#ifndef HAVE_CUDA_MPI
      call messages_write("Warning: trying to use GPU-aware MPI, but we have not detected support in the linked MPI library.")
      call messages_warning()
#endif
      call messages_write("Using GPU-aware MPI.")
      call messages_info()
    end if
 
 
    !%Variable AllowCPUonly
    !%Type logical
    !%Section Execution::Accel
    !%Description
    !% In order to prevent waste of resources, the code will normally stop when the GPU is disabled due to
    !% incomplete implementations or incompatibilities. AllowCPUonly = yes overrides this and allows the
    !% code execution also in these cases.
    !%End
#if defined (HAVE_ACCEL)
    default = .false.
#else
    default = .true.
#endif
    call parse_variable(namespace, 'AllowCPUonly', default, accel%allow_CPU_only)
 
 
    !%Variable InitializeGPUBuffers
    !%Type integer
    !%Default no
    !%Section Execution::Accel
    !%Description
    !% Initialize new GPU buffers to zero on creation (use only for debugging, as it has a performance impact!).
    !%Option no 0
    !% Do not initialize GPU buffers.
    !%Option yes 1
    !% Initialize GPU buffers to zero.
    !%Option nan 2
    !% Initialize GPU buffers to nan.
    !%End
    call parse_variable(namespace, 'InitializeGPUBuffers', option__initializegpubuffers__no, accel%initialize_buffers)
    if (.not. varinfo_valid_option('InitializeGPUBuffers', accel%initialize_buffers)) then
      call messages_input_error(namespace, 'InitializeGPUBuffers')
    end if
 
 
    call messages_print_with_emphasis(namespace=namespace)
 
    pop_sub(accel_init)
 
  contains
 
    subroutine device_info()
#ifdef HAVE_CUDA
      integer :: version
      character(kind=c_char) :: cval_str(257)
#endif
      integer :: major, minor
      character(len=256) :: val_str
 
      push_sub(accel_init.device_info)
 
      call messages_new_line()
      call messages_write('Selected device:')
      call messages_new_line()
 
#ifdef HAVE_CUDA
#ifdef __HIP_PLATFORM_AMD__
      call messages_write('      Framework              : ROCm')
#else
      call messages_write('      Framework              : CUDA')
#endif
#endif
      call messages_info()
 
#ifdef HAVE_CUDA
      call messages_write('      Device type            : GPU', new_line = .true.)
#ifdef __HIP_PLATFORM_AMD__
      call messages_write('      Device vendor          : AMD Corporation', new_line = .true.)
#else
      call messages_write('      Device vendor          : NVIDIA Corporation', new_line = .true.)
#endif
#endif
 
#ifdef HAVE_CUDA
      cval_str = c_null_char
      call cuda_device_name(accel%device%cuda_device, cval_str)
      call string_c_to_f(cval_str, val_str)
#endif
      call messages_write('      Device name            : '//trim(val_str))
      call messages_new_line()
 
#ifdef HAVE_CUDA
      call cuda_device_capability(accel%device%cuda_device, major, minor)
#endif
      call messages_write('      Cuda capabilities      :')
      call messages_write(major, fmt = '(i2)')
      call messages_write('.')
      call messages_write(minor, fmt = '(i1)')
      call messages_new_line()
 
      ! VERSION
#ifdef HAVE_CUDA
      call cuda_driver_version(version)
      call messages_write('      Driver version         : ')
      call messages_write(version)
#endif
      call messages_new_line()
 
 
      call messages_write('      Device memory          :')
      call messages_write(accel%global_memory_size, units=unit_megabytes)
      call messages_new_line()
 
      call messages_write('      Local/shared memory    :')
      call messages_write(accel%local_memory_size, units=unit_kilobytes)
      call messages_new_line()
 
      call messages_write('      Max. group/block size  :')
      call messages_write(accel%max_workgroup_size)
      call messages_new_line()
 
      call messages_info()
 
      pop_sub(accel_init.device_info)
    end subroutine device_info
 
  end subroutine accel_init
 
  ! ------------------------------------------
  subroutine accel_end(namespace)
    type(namespace_t), intent(in) :: namespace
 
    integer(int64) :: hits, misses
    real(real64) :: volume_hits, volume_misses
    logical :: found
    type(accel_mem_t) :: tmp
 
    push_sub(accel_end)
 
    if (accel_is_enabled()) then
 
      ! Release global buffers
      call accel_release_buffer(zm_0_buffer)
      call accel_release_buffer(zm_1_buffer)
      call accel_release_buffer(dm_0_buffer)
      call accel_release_buffer(dm_1_buffer)
 
      do
        call alloc_cache_get(memcache, alloc_cache_any_size, found, tmp%mem)
        if (.not. found) exit
 
#ifdef HAVE_CUDA
        call cuda_mem_free(tmp%mem)
#endif
      end do
 
      call alloc_cache_end(memcache, hits, misses, volume_hits, volume_misses)
 
      call messages_print_with_emphasis(msg="Acceleration-device allocation cache", namespace=namespace)
 
      call messages_new_line()
      call messages_write('    Number of allocations    =')
      call messages_write(hits + misses, new_line = .true.)
      call messages_write('    Volume of allocations    =')
      call messages_write(volume_hits + volume_misses, fmt = 'f18.1', units = unit_gigabytes, align_left = .true., &
        new_line = .true.)
      call messages_write('    Hit ratio                =')
      if (hits + misses > 0) then
        call messages_write(hits/real(hits + misses, real64)*100, fmt='(f6.1)', align_left = .true.)
      else
        call messages_write(m_zero, fmt='(f6.1)', align_left = .true.)
      end if
      call messages_write('%', new_line = .true.)
      call messages_write('    Volume hit ratio         =')
      if (volume_hits + volume_misses > 0) then
        call messages_write(volume_hits/(volume_hits + volume_misses)*100, fmt='(f6.1)', align_left = .true.)
      else
        call messages_write(m_zero, fmt='(f6.1)', align_left = .true.)
      end if
      call messages_write('%')
      call messages_new_line()
      call messages_info()
 
      call messages_print_with_emphasis(namespace=namespace)
    end if
 
    call accel_kernel_global_end()
 
    if (accel_is_enabled()) then
#ifdef HAVE_CUDA
      call cublas_end(accel%cublas_handle)
      if (.not. accel%cuda_mpi) then ! CUDA aware MPI finalize will do the cleanup
        call cuda_end(accel%context%cuda_context, accel%device%cuda_device)
      end if
#endif
 
      if (buffer_alloc_count /= 0) then
        call messages_write('Accel:')
        call messages_write(real(allocated_mem, real64) , fmt = 'f12.1', units = unit_megabytes, align_left = .true.)
        call messages_write(' in ')
        call messages_write(buffer_alloc_count)
        call messages_write(' buffers were not deallocated.')
        call messages_fatal()
      end if
 
    end if
 
    pop_sub(accel_end)
  end subroutine accel_end
 
  ! ------------------------------------------
 
  integer(int64) function accel_padded_size_i8(nn) result(psize)
    integer(int64), intent(in) :: nn
 
    integer(int64) :: modnn, bsize
 
    psize = nn
 
    if (accel_is_enabled()) then
 
      bsize = accel_max_workgroup_size()
 
      psize = nn
      modnn = mod(nn, bsize)
      if (modnn /= 0) psize = psize + bsize - modnn
 
    end if
 
  end function accel_padded_size_i8
 
  ! ------------------------------------------
 
  integer(int32) function accel_padded_size_i4(nn) result(psize)
    integer(int32), intent(in) :: nn
 
    psize = int(accel_padded_size_i8(int(nn, int64)), int32)
 
  end function accel_padded_size_i4
 
  ! ------------------------------------------
 
  subroutine accel_create_buffer_4(this, flags, type, size, set_zero, async)
    type(accel_mem_t),  intent(inout) :: this
    integer,            intent(in)    :: flags
    type(type_t),       intent(in)    :: type
    integer,            intent(in)    :: size
    logical,  optional, intent(in)    :: set_zero
    logical,  optional, intent(in)    :: async
 
    call accel_create_buffer_8(this, flags, type, int(size, int64), set_zero, async)
  end subroutine accel_create_buffer_4
 
  ! ------------------------------------------
 
  subroutine accel_create_buffer_8(this, flags, type, size, set_zero, async)
    type(accel_mem_t),  intent(inout) :: this
    integer,            intent(in)    :: flags
    type(type_t),       intent(in)    :: type
    integer(int64),     intent(in)    :: size
    logical,  optional, intent(in)    :: set_zero
    logical,  optional, intent(in)    :: async
 
    integer(int64) :: fsize
    logical    :: found
    integer(int64) :: initialize_buffers
 
    push_sub(accel_create_buffer_8)
 
    this%type = type
    this%size = size
    this%flags = flags
    fsize = int(size, int64)*types_get_size(type)
    this%allocated = .true.
 
    if (fsize > 0) then
 
      call alloc_cache_get(memcache, fsize, found, this%mem)
 
      if (.not. found) then
#ifdef HAVE_CUDA
        if(optional_default(async, .false.)) then
          call cuda_mem_alloc_async(this%mem, fsize)
        else
          call cuda_mem_alloc(this%mem, fsize)
        end if
#endif
      end if
 
      buffer_alloc_count = buffer_alloc_count + 1
      allocated_mem = allocated_mem + fsize
 
    end if
 
    if (present(set_zero)) then
      initialize_buffers = merge(option__initializegpubuffers__yes, option__initializegpubuffers__no, set_zero)
    else
      initialize_buffers = accel%initialize_buffers
    end if
    select case (initialize_buffers)
    case (option__initializegpubuffers__yes)
      call accel_set_buffer_to(this, type, int(z'00', int8), size)
    case (option__initializegpubuffers__nan)
      call accel_set_buffer_to(this, type, int(z'FF', int8), size)
    end select
 
    pop_sub(accel_create_buffer_8)
  end subroutine accel_create_buffer_8
 
  ! ------------------------------------------
 
  subroutine accel_release_buffer(this, async)
    type(accel_mem_t), intent(inout) :: this
    logical, optional, intent(in) :: async
 
    logical :: put
    integer(int64) :: fsize
 
    push_sub(accel_release_buffer)
 
    if (this%size > 0) then
 
      fsize = int(this%size, int64)*types_get_size(this%type)
 
      call alloc_cache_put(memcache, fsize, this%mem, put)
 
      if (.not. put) then
#ifdef HAVE_CUDA
        if (optional_default(async, .false.)) then
          call cuda_mem_free_async(this%mem)
        else
          call cuda_mem_free(this%mem)
        end if
#endif
      end if
 
      buffer_alloc_count = buffer_alloc_count - 1
      allocated_mem = allocated_mem + fsize
 
    end if
 
    this%size = 0
    this%flags = 0
 
    this%allocated = .false.
 
    pop_sub(accel_release_buffer)
  end subroutine accel_release_buffer
 
  ! ------------------------------------------------------
 
  ! Check if the temporary buffers are the right size, if not reallocate them
  subroutine accel_ensure_buffer_size(buffer, flags, type, required_size, set_zero, async)
    type(accel_mem_t), intent(inout) :: buffer
    integer,           intent(in)    :: flags
    type(type_t),      intent(in)    :: type
    integer,           intent(in)    :: required_size
    logical,           intent(in)    :: set_zero
    logical, optional, intent(in)    :: async
 
    push_sub(accel_ensure_buffer_size)
 
 
    if (accel_buffer_is_allocated(buffer) .and. buffer%size < required_size) then
      call accel_release_buffer(buffer, async=optional_default(async, .false.))
    end if
 
    if (.not. accel_buffer_is_allocated(buffer)) then
      call accel_create_buffer(buffer, flags, type, required_size, set_zero=set_zero, async=optional_default(async, .false.))
    end if
 
    pop_sub(accel_ensure_buffer_size)
  end subroutine accel_ensure_buffer_size
 
  ! ------------------------------------------
 
  logical pure function accel_buffer_is_allocated(this) result(allocated)
    type(accel_mem_t), intent(in) :: this
 
    allocated = this%allocated
  end function accel_buffer_is_allocated
 
  ! -----------------------------------------
 
  subroutine accel_finish()
    ! no push_sub, called too frequently
 
    if (accel_is_enabled()) then
#ifdef HAVE_CUDA
      call cuda_context_synchronize()
#endif
    end if
  end subroutine accel_finish
 
  ! ------------------------------------------
 
  subroutine accel_set_kernel_arg_buffer(kernel, narg, buffer)
    type(accel_kernel_t), intent(inout) :: kernel
    integer,              intent(in)    :: narg
    type(accel_mem_t),    intent(in)    :: buffer
 
    assert(accel_buffer_is_allocated(buffer))
 
    ! no push_sub, called too frequently
#ifdef HAVE_CUDA
    call cuda_kernel_set_arg_buffer(kernel%arguments, buffer%mem, narg)
#endif
 
  end subroutine accel_set_kernel_arg_buffer
 
  ! ------------------------------------------
 
  subroutine accel_set_kernel_arg_local(kernel, narg, type, size)
    type(accel_kernel_t), intent(inout) :: kernel
    integer,              intent(in)    :: narg
    type(type_t),         intent(in)    :: type
    integer,              intent(in)    :: size
 
    integer(int64) :: size_in_bytes
 
    push_sub(accel_set_kernel_arg_local)
 
 
    size_in_bytes = int(size, int64)*types_get_size(type)
 
    if (size_in_bytes > accel%local_memory_size) then
      write(message(1), '(a,f12.6,a)') "CL Error: requested local memory: ", real(size_in_bytes, real64) /1024.0, " Kb"
      write(message(2), '(a,f12.6,a)') "          available local memory: ", real(accel%local_memory_size, real64) /1024.0, " Kb"
      call messages_fatal(2)
    else if (size_in_bytes <= 0) then
      write(message(1), '(a,i10)') "CL Error: invalid local memory size: ", size_in_bytes
      call messages_fatal(1)
    end if
 
#ifdef HAVE_CUDA
    kernel%cuda_shared_mem = size_in_bytes
#endif
 
    pop_sub(accel_set_kernel_arg_local)
  end subroutine accel_set_kernel_arg_local
 
  ! ------------------------------------------
 
  subroutine accel_kernel_run_8(kernel, globalsizes, localsizes)
    type(accel_kernel_t), intent(inout) :: kernel
    integer(int64),          intent(in) :: globalsizes(:)
    integer(int64),          intent(in) :: localsizes(:)
 
    integer :: dim
    integer(int64) :: gsizes(1:3)
    integer(int64) :: lsizes(1:3)
 
    ! no push_sub, called too frequently
 
    ! cuda needs all dimensions
    gsizes = 1
    lsizes = 1
 
    dim = ubound(globalsizes, dim=1)
 
    assert(dim == ubound(localsizes, dim=1))
 
    ! if one size is zero, there is nothing to do
    if (any(globalsizes == 0)) return
 
    assert(all(localsizes > 0))
    assert(all(localsizes <= accel_max_workgroup_size()))
    assert(all(mod(globalsizes, localsizes) == 0))
 
    gsizes(1:dim) = globalsizes(1:dim)
    lsizes(1:dim) = localsizes(1:dim)
 
#ifdef HAVE_CUDA
    ! Maximum dimension of a block
    if (any(lsizes(1:3) > accel%max_block_dim(1:3))) then
      message(1) = "Maximum dimension of a block too large in kernel "//trim(kernel%kernel_name)
      message(2) = "The following conditions should be fulfilled:"
      write(message(3), "(A, I8, A, I8)") "Dim 1: ", lsizes(1), " <= ", accel%max_block_dim(1)
      write(message(4), "(A, I8, A, I8)") "Dim 2: ", lsizes(2), " <= ", accel%max_block_dim(2)
      write(message(5), "(A, I8, A, I8)") "Dim 3: ", lsizes(3), " <= ", accel%max_block_dim(3)
      message(6) = "This is an internal error, please contact the developers."
      call messages_fatal(6)
    end if
 
 
    ! Maximum number of threads per block
    if (product(lsizes) > accel_max_workgroup_size()) then
      message(1) = "Maximum number of threads per block too large in kernel "//trim(kernel%kernel_name)
      message(2) = "The following condition should be fulfilled:"
      write(message(3), "(I8, A, I8)") product(lsizes), " <= ", accel_max_workgroup_size()
      message(4) = "This is an internal error, please contact the developers."
      call messages_fatal(4)
    end if
 
    gsizes(1:3) = gsizes(1:3)/lsizes(1:3)
 
    ! Maximum dimensions of the grid of thread block
    if (any(gsizes(1:3) > accel%max_grid_dim(1:3))) then
      message(1) = "Maximum dimension of grid too large in kernel "//trim(kernel%kernel_name)
      message(2) = "The following conditions should be fulfilled:"
      write(message(3), "(A, I8, A, I10)") "Dim 1: ", gsizes(1), " <= ", accel%max_grid_dim(1)
      write(message(4), "(A, I8, A, I10)") "Dim 2: ", gsizes(2), " <= ", accel%max_grid_dim(2)
      write(message(5), "(A, I8, A, I10)") "Dim 3: ", gsizes(3), " <= ", accel%max_grid_dim(3)
      message(6) = "This is an internal error, please contact the developers."
      call messages_fatal(6)
    end if
 
    call cuda_launch_kernel(kernel%cuda_kernel, gsizes(1), lsizes(1), kernel%cuda_shared_mem, kernel%arguments)
 
    kernel%cuda_shared_mem = 0
#endif
 
  end subroutine accel_kernel_run_8
 
  ! -----------------------------------------------
 
  subroutine accel_kernel_run_4(kernel, globalsizes, localsizes)
    type(accel_kernel_t), intent(inout) :: kernel
    integer,              intent(in)    :: globalsizes(:)
    integer,              intent(in)    :: localsizes(:)
 
    call accel_kernel_run_8(kernel, int(globalsizes, int64), int(localsizes, int64))
 
  end subroutine accel_kernel_run_4
 
  ! -----------------------------------------------
 
  integer pure function accel_max_workgroup_size() result(max_workgroup_size)
    max_workgroup_size = accel%max_workgroup_size
  end function accel_max_workgroup_size
 
  ! -----------------------------------------------
 
  integer function accel_kernel_workgroup_size(kernel) result(workgroup_size)
    type(accel_kernel_t), intent(inout) :: kernel
 
#ifdef HAVE_CUDA
    integer :: max_workgroup_size
#endif
 
    workgroup_size = 0
 
#ifdef HAVE_CUDA
    call cuda_kernel_max_threads_per_block(kernel%cuda_kernel, max_workgroup_size)
    if (debug%info .and. max_workgroup_size /= accel%max_workgroup_size) then
      write(message(1), "(A, I5, A)") "A kernel can use only less threads per block (", workgroup_size, ")", &
        "than available on the device (", accel%max_workgroup_size, ")"
      call messages_info(1)
    end if
    ! recommended number of threads per block is 256 according to the CUDA best practice guide
    ! see https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#thread-and-block-heuristics
    workgroup_size = 256
    ! make sure we do not use more threads per block than available for this kernel
    workgroup_size = min(workgroup_size, max_workgroup_size)
#endif
 
  end function accel_kernel_workgroup_size
 
  ! ----------------------------------------------------
 
  subroutine accel_set_buffer_to(buffer, type, val, nval, offset, async)
    type(accel_mem_t),        intent(inout) :: buffer
    type(type_t),             intent(in)    :: type
    integer(int8),            intent(in)    :: val
    integer(int64),           intent(in)    :: nval
    integer(int64), optional, intent(in)    :: offset
    logical,        optional, intent(in)    :: async
 
    integer(int64) :: nval_, offset_, type_size
 
    push_sub(accel_set_buffer_to)
 
    if (nval == 0) then
      pop_sub(accel_set_buffer_to)
      return
    end if
    assert(nval > 0)
 
    if (present(offset)) then
      assert(offset >= 0)
      if(offset > buffer%size) then
        pop_sub(accel_set_buffer_to)
        return
      end if
    end if
 
    type_size = types_get_size(type)
 
    nval_ = nval*type_size
 
    offset_ = 0_int64
    if (present(offset)) offset_ = offset*type_size
 
    call cuda_mem_set_async(buffer%mem, val, nval_, offset_)
    if(.not. optional_default(async, .false.)) call accel_finish()
 
    pop_sub(accel_set_buffer_to)
  end subroutine accel_set_buffer_to
 
  ! ----------------------------------------------------
 
  subroutine accel_set_buffer_to_zero_i8(buffer, type, nval, offset, async)
    type(accel_mem_t),        intent(inout) :: buffer
    type(type_t),             intent(in)    :: type
    integer(int64),           intent(in)    :: nval
    integer(int64), optional, intent(in)    :: offset
    logical,        optional, intent(in)    :: async
 
    push_sub(accel_set_buffer_to_zero_i8)
 
    call accel_set_buffer_to(buffer, type, int(z'00', int8), nval, offset, async)
 
    pop_sub(accel_set_buffer_to_zero_i8)
  end subroutine accel_set_buffer_to_zero_i8
 
  ! ----------------------------------------------------
 
  subroutine accel_set_buffer_to_zero_i4(buffer, type, nval, offset, async)
    type(accel_mem_t),        intent(inout) :: buffer
    type(type_t),             intent(in)    :: type
    integer(int32),           intent(in)    :: nval
    integer(int32), optional, intent(in)    :: offset
    logical,        optional, intent(in)    :: async
 
    push_sub(accel_set_buffer_to_zero_i4)
 
    if (present(offset)) then
      call accel_set_buffer_to_zero_i8(buffer, type, int(nval, int64), int(offset, int64), async=async)
    else
      call accel_set_buffer_to_zero_i8(buffer, type, int(nval, int64), async=async)
    end if
 
    pop_sub(accel_set_buffer_to_zero_i4)
  end subroutine accel_set_buffer_to_zero_i4
 
  ! ----------------------------------------------------
 
  subroutine accel_check_bandwidth()
    integer :: itime
    integer, parameter :: times = 10
    integer :: size
    real(real64)   :: time, stime
    real(real64)   :: read_bw, write_bw
    type(accel_mem_t) :: buff
    real(real64), allocatable :: data(:)
 
    call messages_new_line()
    call messages_write('Info: Benchmarking the bandwidth between main memory and device memory')
    call messages_new_line()
    call messages_info()
 
    call messages_write(' Buffer size   Read bw  Write bw')
    call messages_new_line()
    call messages_write('       [MiB]   [MiB/s]   [MiB/s]')
    call messages_info()
 
    size = 15000
    do
      safe_allocate(data(1:size))
      call accel_create_buffer(buff, accel_mem_read_write, type_float, size)
 
      stime = loct_clock()
      do itime = 1, times
        call accel_write_buffer(buff, size, data)
        call accel_finish()
      end do
      time = (loct_clock() - stime)/real(times, real64)
 
      write_bw = real(size, real64) *8.0_real64/time
 
      stime = loct_clock()
      do itime = 1, times
        call accel_read_buffer(buff, size, data)
      end do
      call accel_finish()
 
      time = (loct_clock() - stime)/real(times, real64)
      read_bw = real(size, real64) *8.0_real64/time
 
      call messages_write(size*8.0_real64/1024.0_real64**2)
      call messages_write(write_bw/1024.0_real64**2, fmt = '(f10.1)')
      call messages_write(read_bw/1024.0_real64**2, fmt = '(f10.1)')
      call messages_info()
 
      call accel_release_buffer(buff)
 
      safe_deallocate_a(data)
 
      size = int(size*2.0)
 
      if (size > 50000000) exit
    end do
  end subroutine accel_check_bandwidth
 
  ! ----------------------------------------------------
 
  logical pure function accel_use_shared_mem() result(use_shared_mem)
 
    use_shared_mem = accel%shared_mem
 
  end function accel_use_shared_mem
 
  !------------------------------------------------------------
 
  subroutine accel_kernel_global_init()
 
    push_sub(accel_kernel_global_init)
 
    nullify(head)
 
    call cuda_module_map_init(accel%module_map)
 
    pop_sub(accel_kernel_global_init)
  end subroutine accel_kernel_global_init
 
  !------------------------------------------------------------
 
  subroutine accel_kernel_global_end()
    type(accel_kernel_t), pointer :: next_head
 
    push_sub(accel_kernel_global_end)
 
    do
      if (.not. associated(head)) exit
      next_head => head%next
      call accel_kernel_end(head)
      head => next_head
    end do
 
    if (accel_is_enabled()) then
      call cuda_module_map_end(accel%module_map)
    end if
 
    pop_sub(accel_kernel_global_end)
  end subroutine accel_kernel_global_end
 
  !------------------------------------------------------------
 
  subroutine accel_kernel_build(this, file_name, kernel_name, flags)
    type(accel_kernel_t),        intent(inout) :: this
    character(len=*),            intent(in)    :: file_name
    character(len=*),            intent(in)    :: kernel_name
    character(len=*), optional,  intent(in)    :: flags
 
#ifdef HAVE_CUDA
    character(len=1000) :: all_flags
#endif
 
    push_sub(accel_kernel_build)
 
    call profiling_in("ACCEL_COMPILE", exclude = .true.)
 
#ifdef HAVE_CUDA
    all_flags = '-I'//trim(conf%share)//'/kernels/'//" "//trim(accel%debug_flag)
 
    if (accel_use_shared_mem()) then
      all_flags = trim(all_flags)//' -DSHARED_MEM'
    end if
 
    if (present(flags)) then
      all_flags = trim(all_flags)//' '//trim(flags)
    end if
 
    call cuda_build_program(accel%module_map, this%cuda_module, accel%device%cuda_device, &
      string_f_to_c(trim(file_name)), string_f_to_c(trim(all_flags)))
 
    call cuda_create_kernel(this%cuda_kernel, this%cuda_module, string_f_to_c(trim(kernel_name)))
    call cuda_alloc_arg_array(this%arguments)
 
    this%cuda_shared_mem = 0
#endif
 
    this%initialized = .true.
    this%kernel_name = trim(kernel_name)
 
    call profiling_out("ACCEL_COMPILE")
 
    pop_sub(accel_kernel_build)
  end subroutine accel_kernel_build
 
  !------------------------------------------------------------
 
  subroutine accel_kernel_end(this)
    type(accel_kernel_t), intent(inout) :: this
 
    push_sub(accel_kernel_end)
 
#ifdef HAVE_CUDA
    call cuda_free_arg_array(this%arguments)
    call cuda_release_kernel(this%cuda_kernel)
    ! modules are not released here, since they are not associated to a kernel
#endif
 
    this%initialized = .false.
 
    pop_sub(accel_kernel_end)
  end subroutine accel_kernel_end
 
  !------------------------------------------------------------
 
  subroutine accel_kernel_start_call(this, file_name, kernel_name, flags)
    type(accel_kernel_t), target, intent(inout) :: this
    character(len=*),             intent(in)    :: file_name
    character(len=*),             intent(in)    :: kernel_name
    character(len=*), optional,   intent(in)    :: flags
 
    push_sub(accel_kernel_start_call)
 
    if (.not. this%initialized) then
      call accel_kernel_build(this, file_name, kernel_name, flags)
      this%next => head
      head => this
    end if
 
    pop_sub(accel_kernel_start_call)
  end subroutine accel_kernel_start_call
 
  !--------------------------------------------------------------
 
  integer(int64) pure function accel_global_memory_size() result(size)
 
    size = accel%global_memory_size
 
  end function accel_global_memory_size
 
  !--------------------------------------------------------------
 
  integer(int64) pure function accel_local_memory_size() result(size)
 
    size = accel%local_memory_size
 
  end function accel_local_memory_size
 
  !--------------------------------------------------------------
 
  integer pure function accel_max_size_per_dim(dim) result(size)
    integer, intent(in) :: dim
 
    size = 0
#ifdef HAVE_CUDA
    size = 32768
    if (dim == 1) size = 2**30
#endif
  end function accel_max_size_per_dim
 
  ! ------------------------------------------------------
 
  subroutine accel_set_stream(stream_number)
    integer, intent(in) :: stream_number
 
    push_sub(accel_set_stream)
 
    if (accel_is_enabled()) then
#ifdef HAVE_CUDA
      call cuda_set_stream(accel%cuda_stream, stream_number)
      call cublas_set_stream(accel%cublas_handle, accel%cuda_stream)
#endif
    end if
 
    pop_sub(accel_set_stream)
  end subroutine accel_set_stream
 
  ! ------------------------------------------------------
 
  subroutine accel_get_stream(stream_number)
    integer, intent(inout) :: stream_number
 
    push_sub(accel_get_stream)
 
    if (accel_is_enabled()) then
#ifdef HAVE_CUDA
      call cuda_get_stream(stream_number)
#endif
    end if
 
    pop_sub(accel_get_stream)
  end subroutine accel_get_stream
 
  ! ------------------------------------------------------
 
  subroutine accel_synchronize_all_streams()
    push_sub(accel_synchronize_all_streams)
 
    if (accel_is_enabled()) then
#ifdef HAVE_CUDA
      call cuda_synchronize_all_streams()
#endif
    end if
 
    pop_sub(accel_synchronize_all_streams)
  end subroutine accel_synchronize_all_streams
 
  function daccel_get_pointer_with_offset(buffer, offset) result(buffer_offset)
    type(c_ptr), intent(in) :: buffer
    integer(int64), intent(in) :: offset
    type(c_ptr) :: buffer_offset
 
    push_sub(daccel_get_pointer_with_offset)
#ifdef HAVE_CUDA
    call cuda_get_pointer_with_offset(buffer, offset, buffer_offset)
#else
    ! this is needed to make the compiler happy for non-GPU compilations
    buffer_offset = buffer
#endif
    pop_sub(daccel_get_pointer_with_offset)
  end function daccel_get_pointer_with_offset
 
  function zaccel_get_pointer_with_offset(buffer, offset) result(buffer_offset)
    type(c_ptr), intent(in) :: buffer
    integer(int64), intent(in) :: offset
    type(c_ptr) :: buffer_offset
 
    push_sub(zaccel_get_pointer_with_offset)
#ifdef HAVE_CUDA
    call cuda_get_pointer_with_offset(buffer, 2_int64*offset, buffer_offset)
#else
    ! this is needed to make the compiler happy for non-GPU compilations
    buffer_offset = buffer
#endif
    pop_sub(zaccel_get_pointer_with_offset)
  end function zaccel_get_pointer_with_offset
 
  subroutine accel_clean_pointer(buffer)
    type(c_ptr), intent(in) :: buffer
 
    push_sub(accel_clean_pointer)
#ifdef HAVE_CUDA
    call cuda_clean_pointer(buffer)
#endif
    pop_sub(accel_clean_pointer)
  end subroutine accel_clean_pointer
 
  subroutine accel_get_unfolded_size(size, grid_size, thread_block_size)
    integer(int64), intent(in)  :: size
    integer(int64), intent(out) :: grid_size
    integer(int64), intent(out) :: thread_block_size
 
    push_sub(accel_get_unfolded_size)
#ifdef __HIP_PLATFORM_AMD__
    ! not benefitial for AMD chips
    grid_size = size
    thread_block_size = size
#else
    grid_size = size * accel%warp_size
    thread_block_size = accel%warp_size
#endif
    pop_sub(accel_get_unfolded_size)
  end subroutine accel_get_unfolded_size
 
#include "undef.F90"
#include "real.F90"
#include "accel_inc.F90"
 
#include "undef.F90"
#include "complex.F90"
#include "accel_inc.F90"
 
#include "undef.F90"
#include "integer.F90"
#include "accel_inc.F90"
 
#include "undef.F90"
#include "integer8.F90"
#include "accel_inc.F90"
 
end module accel_oct_m
 
!! Local Variables:
!! mode: f90
!! coding: utf-8
!! End: