opencl.F90

!! Copyright (C) 2010 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., 59 Temple Place - Suite 330, Boston, MA
!! 02111-1307, USA.
!!
!! $Id: opencl.F90 3587 2007-11-22 16:43:00Z xavier $

#include "global.h"

module opencl_m
#ifdef HAVE_OPENCL
  use cl
#endif
#ifdef HAVE_CLAMDBLAS
  use clAmdBlas
#endif
#ifdef HAVE_CLAMDFFT
  use clAmdFft
#endif
  use datasets_m
  use global_m
  use io_m
  use loct_m
  use messages_m
  use mpi_m
  use types_m
  use parser_m
  use profiling_m
  use unit_system_m

  implicit none 

  private

  public ::                       &
    opencl_is_enabled,            &
    opencl_t,                     &
    opencl_init,                  &
    opencl_end,                   &
    opencl_padded_size,           &
    opencl_mem_t

#ifdef HAVE_OPENCL
  public ::                       &
    opencl_create_buffer,         &
    opencl_write_buffer,          &
    opencl_read_buffer,           &
    opencl_release_buffer,        &
    opencl_finish,                &
    opencl_set_kernel_arg,        &
    opencl_max_workgroup_size,    &
    opencl_kernel_workgroup_size, &
    opencl_kernel_run,            &
    opencl_build_program,         &
    opencl_release_program,       &
    opencl_release_kernel,        &
    opencl_create_kernel,         &
    opencl_print_error,           &
    clblas_print_error,           &
    clfft_print_error,            &
    opencl_set_buffer_to_zero
#endif

  type opencl_t 
#ifdef HAVE_OPENCL
    type(cl_platform_id)   :: platform_id
    type(cl_context)       :: context
    type(cl_command_queue) :: command_queue
    type(cl_device_id)     :: device
#endif
    integer                :: max_workgroup_size
    integer                :: local_memory_size
    logical                :: enabled
  end type opencl_t

  type opencl_mem_t
#ifdef HAVE_OPENCL
    type(cl_mem)           :: mem
#endif
    integer(SIZEOF_SIZE_T) :: size
    type(type_t)           :: type
  end type opencl_mem_t

  type(opencl_t), public :: opencl

#ifdef HAVE_OPENCL
  ! the kernels
  type(cl_kernel), public :: kernel_vpsi
  type(cl_kernel), public :: kernel_vpsi_spinors
  type(cl_kernel), public :: kernel_daxpy
  type(cl_kernel), public :: kernel_zaxpy
  type(cl_kernel), public :: kernel_copy
  type(cl_kernel), public :: kernel_projector_bra
  type(cl_kernel), public :: kernel_projector_ket
  type(cl_kernel), public :: dpack
  type(cl_kernel), public :: zpack
  type(cl_kernel), public :: dunpack
  type(cl_kernel), public :: zunpack
  type(cl_kernel), public :: kernel_subarray_gather
  type(cl_kernel), public :: kernel_density_real
  type(cl_kernel), public :: kernel_density_complex
  type(cl_kernel), public :: kernel_phase
  type(cl_kernel), public :: dkernel_dot_matrix
  type(cl_kernel), public :: zkernel_dot_matrix
  type(cl_kernel), public :: zkernel_dot_matrix_spinors
  type(cl_kernel), public :: dkernel_dot_vector
  type(cl_kernel), public :: zkernel_dot_vector
  type(cl_kernel), public :: kernel_nrm2_vector
  type(cl_kernel), public :: dzmul
  type(cl_kernel), public :: zzmul

  ! kernels used locally
  type(cl_kernel)         :: set_zero

  interface opencl_create_buffer
    module procedure opencl_create_buffer_4
  end interface opencl_create_buffer

  interface opencl_write_buffer
    module procedure iopencl_write_buffer_1, dopencl_write_buffer_1, zopencl_write_buffer_1
    module procedure iopencl_write_buffer_2, dopencl_write_buffer_2, zopencl_write_buffer_2
    module procedure iopencl_write_buffer_3, dopencl_write_buffer_3, zopencl_write_buffer_3
  end interface opencl_write_buffer

  interface opencl_read_buffer
    module procedure iopencl_read_buffer_1, dopencl_read_buffer_1, zopencl_read_buffer_1
    module procedure iopencl_read_buffer_2, dopencl_read_buffer_2, zopencl_read_buffer_2
    module procedure iopencl_read_buffer_3, dopencl_read_buffer_3, zopencl_read_buffer_3
  end interface opencl_read_buffer

  interface opencl_set_kernel_arg
    module procedure                 &
      opencl_set_kernel_arg_buffer,  &
      iopencl_set_kernel_arg_data,   &
      dopencl_set_kernel_arg_data,   &
      zopencl_set_kernel_arg_data,   &
      opencl_set_kernel_arg_local
  end interface opencl_set_kernel_arg

#endif

  type(profile_t), save :: prof_read, prof_write, prof_kernel_run

  integer, parameter  ::      &
    OPENCL_GPU         = -1,  &
    OPENCL_CPU         = -2,  &
    OPENCL_ACCELERATOR = -3,  &
    OPENCL_DEFAULT     = -4


  integer, parameter  ::      &
    CL_PLAT_INVALID   = -1,   &
    CL_PLAT_AMD       = -2,   &
    CL_PLAT_NVIDIA    = -3,   &
    CL_PLAT_ATI       = -4,   &
    CL_PLAT_INTEL     = -5

  ! a "convenience" public variable
  integer, public :: cl_status

  integer, parameter :: OPENCL_MAX_FILE_LENGTH = 10000

  integer :: buffer_alloc_count
  integer(8) :: allocated_mem

  contains

    pure logical function opencl_is_enabled() result(enabled)
#ifdef HAVE_OPENCL
      enabled = opencl%enabled
#else
      enabled = .false.
#endif
    end function opencl_is_enabled

    ! ------------------------------------------
    
    subroutine opencl_init(base_grp)
      type(mpi_grp_t),  intent(inout) :: base_grp

      logical  :: disable, default, run_benchmark
      integer  :: device_type
      integer  :: idevice, iplatform, ndevices, idev, cl_status, ret_devices, nplatforms, iplat
      character(len=256) :: device_name
#ifdef HAVE_OPENCL
      type(cl_program) :: prog
      type(cl_platform_id), allocatable :: allplatforms(:)
      type(cl_device_id), allocatable :: alldevices(:)