cuda.F90

Go to the documentation of this file.

!! Copyright (C) 2019 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"
 
module cuda_oct_m
 
  implicit none
 
  private
  public ::                             &
    cuda_init,                          &
    cuda_end,                           &
    cuda_program_map_init,              &
    cuda_program_map_end,               &
    cuda_program_register,              &
    cuda_program_build,                 &
    cuda_program_get_kernel,            &
    cuda_release_module,                &
    cuda_release_kernel,                &
    cuda_device_max_threads_per_block,  &
    cuda_kernel_max_threads_per_block,  &
    cuda_device_max_grid_dim_x,         &
    cuda_device_max_grid_dim_y,         &
    cuda_device_max_grid_dim_z,         &
    cuda_device_max_block_dim_x,        &
    cuda_device_max_block_dim_y,        &
    cuda_device_max_block_dim_z,        &
    cuda_device_total_memory,           &
    cuda_device_shared_memory,          &
    cuda_mem_alloc,                     &
    cuda_mem_free,                      &
    cuda_alloc_arg_array,               &
    cuda_free_arg_array,                &
    cuda_kernel_set_arg_buffer,         &
    cuda_kernel_set_arg_value,          &
    cuda_context_synchronize,           &
    cuda_synchronize_all_streams,       &
    cuda_launch_kernel,                 &
    cuda_device_name,                   &
    cuda_device_capability,             &
    cuda_driver_version,                &
    cuda_device_get_warpsize,           &
    cuda_deref,                         &
    cuda_set_stream,                    &
    cuda_get_stream,                    &
    cuda_memcpy_htod,                   &
    cuda_memcpy_dtoh,                   &
    cuda_memcpy_dtod,                   &
    cuda_get_pointer_with_offset,       &
    cuda_clean_pointer,                 &
    cuda_mem_free_async,                &
    cuda_mem_set_async,                 &
    cuda_mem_alloc_async,               &
    cublas_init,                        &
    cublas_end,                         &
    cublas_set_stream
 
  integer, parameter, public ::                      &
#ifdef HAVE_HIP
    cublas_diag_non_unit = 131,                      &
    cublas_diag_unit     = 132
#else
  cublas_diag_non_unit = 0,                        &
    cublas_diag_unit     = 1
#endif
 
  integer, parameter, public ::                      &
#ifdef HAVE_HIP
    cublas_op_n = 111,                               &
    cublas_op_t = 112,                               &
    cublas_op_c = 113
#else
  cublas_op_n = 0,                                 &
    cublas_op_t = 1,                                 &
    cublas_op_c = 2
#endif
 
  integer, parameter, public ::                      &
#ifdef HAVE_HIP
    cublas_fill_mode_lower = 122,                    &
    cublas_fill_mode_upper = 121
#else
  cublas_fill_mode_lower = 0,                      &
    cublas_fill_mode_upper = 1
#endif
 
  integer, parameter, public ::                      &
#ifdef HAVE_HIP
    cublas_side_left  = 141,                         &
    cublas_side_right = 142
#else
  cublas_side_left  = 0,                           &
    cublas_side_right = 1
#endif
 
  interface
 
    subroutine cuda_init(context, device, stream, device_number, rank) bind(c, name="oct_cuda_init")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: context
      type(c_ptr), intent(inout) :: device
      type(c_ptr), intent(inout) :: stream
      integer(c_int), intent(inout) :: device_number
      integer(c_int), intent(out) :: rank
    end subroutine cuda_init
 
    ! -------------------------------------------------
 
    subroutine cuda_end(context, device) bind(c, name="oct_cuda_end")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: context
      type(c_ptr), intent(inout) :: device
    end subroutine cuda_end
 
    ! -------------------------------------------------
 
    subroutine cuda_program_map_init(program_map) bind(c, name="oct_cuda_program_map_init")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: program_map
    end subroutine cuda_program_map_init
 
    ! -------------------------------------------------
 
    subroutine cuda_program_map_end(program_map) bind(c, name="oct_cuda_program_map_end")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: program_map
    end subroutine cuda_program_map_end
 
    ! -------------------------------------------------
 
    subroutine cuda_program_register(program_map, fname, name) bind(c, name="oct_cuda_program_register")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: program_map
      character(kind=c_char), intent(in) :: fname(*)
      character(kind=c_char), intent(in) :: name(*)
    end subroutine cuda_program_register
 
    ! -------------------------------------------------
 
    subroutine cuda_program_build(program_map, device, fname, flags) bind(c, name="oct_cuda_program_build")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: program_map
      type(c_ptr), intent(inout) :: device
      character(kind=c_char), intent(in) :: fname(*)
      character(kind=c_char), intent(in) :: flags(*)
    end subroutine cuda_program_build
 
    ! -------------------------------------------------
 
    subroutine cuda_program_get_kernel(kernel, program_map, fname, flags, name) &
      bind(c, name="oct_cuda_program_get_kernel")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: kernel
      type(c_ptr), intent(inout) :: program_map
      character(kind=c_char), intent(in) :: fname(*)
      character(kind=c_char), intent(in) :: flags(*)
      character(kind=c_char), intent(in) :: name(*)
    end subroutine cuda_program_get_kernel
    ! -------------------------------------------------
 
    subroutine cuda_release_module(modul) bind(c, name="oct_cuda_release_module")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: modul
    end subroutine cuda_release_module
 
    ! -------------------------------------------------
 
    subroutine cuda_release_kernel(kernel) bind(c, name="oct_cuda_release_kernel")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: kernel
    end subroutine cuda_release_kernel
 
    ! -------------------------------------------------
 
    subroutine cuda_device_max_threads_per_block(device, max_threads) bind(c, name="oct_cuda_device_max_threads_per_block")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: device
      integer(c_int), intent(out) :: max_threads
    end subroutine cuda_device_max_threads_per_block
 
    ! -------------------------------------------------
 
    subroutine cuda_kernel_max_threads_per_block(device, max_threads) bind(c, name="oct_cuda_kernel_max_threads_per_block")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: device
      integer(c_int), intent(out) :: max_threads
    end subroutine cuda_kernel_max_threads_per_block
 
    ! -------------------------------------------------
 
    subroutine cuda_device_max_grid_dim_x(device, max_dim) bind(c, name="oct_cuda_device_max_grid_dim_x")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: device
      integer(c_int), intent(out) :: max_dim
    end subroutine cuda_device_max_grid_dim_x
 
    ! -------------------------------------------------
 
    subroutine cuda_device_max_grid_dim_y(device, max_dim) bind(c, name="oct_cuda_device_max_grid_dim_y")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: device
      integer(c_int), intent(out) :: max_dim
    end subroutine cuda_device_max_grid_dim_y
 
    ! -------------------------------------------------
 
    subroutine cuda_device_max_grid_dim_z(device, max_dim) bind(c, name="oct_cuda_device_max_grid_dim_z")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: device
      integer(c_int), intent(out) :: max_dim
    end subroutine cuda_device_max_grid_dim_z
 
    ! -------------------------------------------------
 
    subroutine cuda_device_max_block_dim_x(device, max_dim) bind(c, name="oct_cuda_device_max_block_dim_x")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: device
      integer(c_int), intent(out) :: max_dim
    end subroutine cuda_device_max_block_dim_x
 
    ! -------------------------------------------------
 
    subroutine cuda_device_max_block_dim_y(device, max_dim) bind(c, name="oct_cuda_device_max_block_dim_y")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: device
      integer(c_int), intent(out) :: max_dim
    end subroutine cuda_device_max_block_dim_y
 
    ! -------------------------------------------------
 
    subroutine cuda_device_max_block_dim_z(device, max_dim) bind(c, name="oct_cuda_device_max_block_dim_z")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: device
      integer(c_int), intent(out) :: max_dim
    end subroutine cuda_device_max_block_dim_z
 
 
    ! -------------------------------------------------
 
    subroutine cuda_device_total_memory(device, total_memory) bind(c, name="oct_cuda_device_total_memory")
      use iso_c_binding
      use, intrinsic :: iso_fortran_env
      implicit none
 
      type(c_ptr), intent(inout) :: device
      integer(c_int64_t), intent(out) :: total_memory
    end subroutine cuda_device_total_memory
 
    ! -------------------------------------------------
 
    subroutine cuda_device_shared_memory(device, shared_memory) bind(c, name="oct_cuda_device_shared_memory")
      use iso_c_binding
      use, intrinsic :: iso_fortran_env
      implicit none
 
      type(c_ptr), intent(inout) :: device
      integer(c_int64_t), intent(out) :: shared_memory
    end subroutine cuda_device_shared_memory
 
    ! -------------------------------------------------
 
    subroutine cuda_mem_alloc(cuda_ptr, size) bind(c, name="oct_cuda_mem_alloc")
      use iso_c_binding
      use, intrinsic :: iso_fortran_env
      implicit none
 
      type(c_ptr), intent(inout) :: cuda_ptr
      integer(c_int64_t), intent(in) :: size
    end subroutine cuda_mem_alloc
 
    ! -------------------------------------------------
 
    subroutine cuda_mem_free(cuda_ptr) bind(c, name="oct_cuda_mem_free")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: cuda_ptr
    end subroutine cuda_mem_free
 
    ! -------------------------------------------------
 
    subroutine cuda_alloc_arg_array(arg_array) bind(c, name="oct_cuda_alloc_arg_array")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: arg_array
    end subroutine cuda_alloc_arg_array
 
    ! -------------------------------------------------
 
    subroutine cuda_free_arg_array(arg_array) bind(c, name="oct_cuda_free_arg_array")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: arg_array
    end subroutine cuda_free_arg_array
 
    ! -------------------------------------------------
 
    subroutine cuda_kernel_set_arg_buffer(arg_array, cuda_ptr, arg_index) bind(c, name="oct_cuda_kernel_set_arg_buffer")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: arg_array
      type(c_ptr), intent(in) :: cuda_ptr
      integer(c_int), intent(in) :: arg_index
    end subroutine cuda_kernel_set_arg_buffer
 
    ! -------------------------------------------------
 
    subroutine cuda_kernel_set_arg_value(arg_array, arg, arg_index, arg_size) bind(c, name="oct_cuda_kernel_set_arg_value")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: arg_array
      type(c_ptr), intent(in) :: arg
      integer(c_int), intent(in) :: arg_index
      integer(c_int), intent(in) :: arg_size
    end subroutine cuda_kernel_set_arg_value
 
    ! -------------------------------------------------
 
    subroutine cuda_context_synchronize() bind(c, name="oct_cuda_context_synchronize")
      implicit none
    end subroutine cuda_context_synchronize
 
    ! -------------------------------------------------
 
    subroutine cuda_synchronize_all_streams() bind(c, name="oct_cuda_synchronize_all_streams")
      implicit none
    end subroutine cuda_synchronize_all_streams
 
    ! -------------------------------------------------
 
    subroutine cuda_launch_kernel(kernel, griddim, blockdim, shared_mem, arg_array) bind(c, name="oct_cuda_launch_kernel")
      use iso_c_binding
      use, intrinsic :: iso_fortran_env
      implicit none
 
      type(c_ptr), intent(inout) :: kernel
      integer(c_int64_t), intent(in) :: griddim
      integer(c_int64_t), intent(in) :: blockdim
      integer(c_int64_t), intent(in) :: shared_mem
      type(c_ptr), intent(inout) :: arg_array
    end subroutine cuda_launch_kernel
 
    ! -------------------------------------------------
 
    subroutine cuda_device_name(device, name) bind(c, name="oct_cuda_device_name")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: device
      character(kind=c_char), intent(inout) :: name(*)
    end subroutine cuda_device_name
 
    ! -------------------------------------------------
 
    subroutine cuda_device_capability(device, major, minor) bind(c, name="oct_cuda_device_capability")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: device
      integer(c_int), intent(out) :: major
      integer(c_int), intent(out) :: minor
    end subroutine cuda_device_capability
 
    ! -------------------------------------------------
 
    subroutine cuda_driver_version(version) bind(c, name="oct_cuda_driver_version")
      use iso_c_binding
      implicit none
 
      integer(c_int),intent(out) :: version
    end subroutine cuda_driver_version
 
    ! -------------------------------------------------
 
    subroutine cuda_device_get_warpsize(device, warpsize) bind(c, name="oct_cuda_device_get_warpsize")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: device
      integer(c_int), intent(out) :: warpsize
    end subroutine cuda_device_get_warpsize
 
    subroutine cuda_deref(cuda_ptr, cuda_deref_ptr) bind(c, name="oct_cuda_deref")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(in)  :: cuda_ptr
      type(c_ptr), intent(out) :: cuda_deref_ptr
    end subroutine cuda_deref
 
    subroutine cuda_set_stream(stream, stream_number) bind(c, name="oct_cuda_set_stream")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: stream
      integer(c_int), intent(in) :: stream_number
    end subroutine cuda_set_stream
 
    subroutine cuda_get_stream(stream_number) bind(c, name="oct_cuda_get_stream")
      use iso_c_binding
      implicit none
 
      integer(c_int), intent(out) :: stream_number
    end subroutine cuda_get_stream
 
    ! -------------------------------------------------
 
    subroutine cuda_memcpy_htod(cuda_ptr, data, size, offset) bind(c, name="oct_cuda_memcpy_htod")
      use iso_c_binding
      use, intrinsic :: iso_fortran_env
      implicit none
 
      type(c_ptr), intent(in) :: cuda_ptr
      type(c_ptr), intent(in) :: data
      integer(c_int64_t), intent(in) :: size
      integer(c_int64_t), intent(in) :: offset
    end subroutine cuda_memcpy_htod
 
    ! -------------------------------------------------
 
    subroutine cuda_memcpy_dtoh(cuda_ptr, data, size, offset) bind(c, name="oct_cuda_memcpy_dtoh")
      use iso_c_binding
      use, intrinsic :: iso_fortran_env
      implicit none
 
      type(c_ptr), intent(in) :: cuda_ptr
      type(c_ptr), intent(inout) :: data
      integer(c_int64_t),  intent(in) :: size
      integer(c_int64_t),  intent(in) :: offset
    end subroutine cuda_memcpy_dtoh
 
    ! -------------------------------------------------
 
    subroutine cuda_memcpy_dtod(cuda_ptr_dest, cuda_ptr_src, size, offset_dest, offset_src) bind(c, name="oct_cuda_memcpy_dtod")
      use iso_c_binding
      use, intrinsic :: iso_fortran_env
      implicit none
 
      type(c_ptr), intent(inout) :: cuda_ptr_dest
      type(c_ptr), intent(in) :: cuda_ptr_src
      integer(c_int64_t), intent(in) :: size
      integer(c_int64_t), intent(in) :: offset_dest
      integer(c_int64_t), intent(in) :: offset_src
    end subroutine cuda_memcpy_dtod
 
    subroutine cuda_get_pointer_with_offset(buffer, offset, buffer_offset) bind(c, name="oct_cuda_get_pointer_with_offset")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(in) :: buffer
      integer(8), intent(in) :: offset
      type(c_ptr), intent(out) :: buffer_offset
    end subroutine cuda_get_pointer_with_offset
 
    subroutine cuda_clean_pointer(buffer) bind(c, name="oct_cuda_clean_pointer")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(in) :: buffer
    end subroutine cuda_clean_pointer
 
    subroutine cuda_mem_free_async(cuda_ptr) bind(c, name="oct_cuda_mem_free_async")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: cuda_ptr
    end subroutine cuda_mem_free_async
 
    subroutine cuda_mem_set_async(cuda_ptr, arg, arg_size, arg_offset) bind(c, name="oct_cuda_mem_set_async")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: cuda_ptr
      integer(c_int8_t), intent(in) :: arg
      integer(c_int64_t) :: arg_size
      integer(c_int64_t) :: arg_offset
    end subroutine cuda_mem_set_async
 
    subroutine cuda_mem_alloc_async(cuda_ptr, arg_size) bind(c, name="oct_cuda_mem_alloc_async")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(inout) :: cuda_ptr
      integer(c_int64_t) :: arg_size
    end subroutine cuda_mem_alloc_async
 
    ! -------------------------------------------------
 
    subroutine cublas_init(stream) bind(c, name="oct_cublas_init")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(in) :: stream
    end subroutine cublas_init
 
    subroutine cublas_end() bind(c, name="oct_cublas_end")
      implicit none
    end subroutine cublas_end
 
    subroutine cublas_set_stream(stream, stream_number) bind(c, name="oct_cublas_set_stream")
      use iso_c_binding
      implicit none
 
      type(c_ptr), intent(in) :: stream
      integer(c_int), intent(in) :: stream_number
    end subroutine cublas_set_stream
 
  end interface
 
end module cuda_oct_m
 
!! Local Variables:
!! mode: f90
!! coding: utf-8
!! End: