accelerate.cuf

!Crown Copyright 2012 AWE.
!
! This file is part of CloverLeaf.
!
! CloverLeaf 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 3 of the License, or (at your option) 
! any later version.
!
! CloverLeaf 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 
! CloverLeaf. If not, see http://www.gnu.org/licenses/.

!>  @brief Driver for the acceleration kernels
!>  @author Wayne Gaudin
!>  @details Calls user requested kernel

MODULE accelerate_module
  
CONTAINS
  
  SUBROUTINE accelerate()
    
    USE clover_module
    USE accelerate_kernel_module
    use cudafor 

    IMPLICIT NONE
    
    type(dim3) :: acc_grid, acc_block
    integer :: xmin, xmax, ymin, ymax, istat

    INTEGER :: c

    REAL(KIND=8) :: kernel_time,timer
    
    IF(profiler_on) kernel_time=timer()
    DO c=1,number_of_chunks       
       IF(chunks(c)%task.EQ.parallel%task) THEN          
          if (use_cuda_fortran) then
             xmin = chunks(c)%field%x_min
             xmax = chunks(c)%field%x_max
             ymin = chunks(c)%field%y_min
             ymax = chunks(c)%field%y_max
             
             acc_block = dim3(32,8,1)
             acc_grid = dim3(ceiling((xmax-xmin+2.)/(acc_block%x)), ceiling((ymax-ymin+2.)/acc_block%y), 1)             
             
             CALL accelerate_kernel_1<<<acc_grid, acc_block>>>(xmin,                &
                  xmax,                  &
                  ymin,                  &
                  ymax,                  &
                  dt,                                     &
                  chunks(c)%field%xarea,                  &
                  chunks(c)%field%yarea,                  &
                  chunks(c)%field%volume,                 &
                  chunks(c)%field%density0,               &
                  chunks(c)%field%pressure,               &
                  chunks(c)%field%viscosity,              &
                  chunks(c)%field%xvel0,                  &
                  chunks(c)%field%yvel0,                  &
                  chunks(c)%field%xvel1,                  &
                  chunks(c)%field%yvel1,                  &
                  chunks(c)%field%work_array1             )

          else
             CALL accelerate_kernel(chunks(c)%field%x_min,                &
                  chunks(c)%field%x_max,                  &
                  chunks(c)%field%y_min,                  &
                  chunks(c)%field%y_max,                  &
                  dt,                                     &
                  chunks(c)%field%xarea,                  &
                  chunks(c)%field%yarea,                  &
                  chunks(c)%field%volume,                 &
                  chunks(c)%field%density0,               &
                  chunks(c)%field%pressure,               &
                  chunks(c)%field%viscosity,              &
                  chunks(c)%field%xvel0,                  &
                  chunks(c)%field%yvel0,                  &
                  chunks(c)%field%xvel1,                  &
                  chunks(c)%field%yvel1,                  &
                  chunks(c)%field%work_array1             )
          end if
       ENDIF          
    ENDDO
    IF(profiler_on) profiler%acceleration=profiler%acceleration+(timer()-kernel_time)    
  END SUBROUTINE accelerate
END MODULE accelerate_module