viscosity.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 viscosity kernels
!>  @author Wayne Gaudin
!>  @details Selects the user specified kernel to caluclate the artificial 
!>  viscosity.

MODULE viscosity_module

CONTAINS

  SUBROUTINE viscosity()

    USE clover_module
    USE viscosity_kernel_module
    use cudafor

    IMPLICIT NONE

    INTEGER :: c, xmin, xmax, ymin, ymax
    type(dim3) :: v_grid, v_block

    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
             
             v_block = dim3(32,8,1)
             v_grid = dim3(ceiling((xmax-xmin+1.)/v_block%x), &
                  ceiling((ymax-ymin+1.)/v_block%y), 1)
             
             CALL viscosity_kernel_1<<<v_grid, v_block>>>(xmin,                   &
                  xmax,                     &
                  ymin,                     &
                  ymax,                     &
                  chunks(c)%field%celldx,                    &
                  chunks(c)%field%celldy,                    &
                  chunks(c)%field%density0,                  &
                  chunks(c)%field%pressure,                  &
                  chunks(c)%field%viscosity,                 &
                  chunks(c)%field%xvel0,                     &
                  chunks(c)%field%yvel0                      )
          else
             CALL viscosity_kernel(chunks(c)%field%x_min,                   &
                  chunks(c)%field%x_max,                     &
                  chunks(c)%field%y_min,                     &
                  chunks(c)%field%y_max,                     &
                  chunks(c)%field%celldx,                    &
                  chunks(c)%field%celldy,                    &
                  chunks(c)%field%density0,                  &
                  chunks(c)%field%pressure,                  &
                  chunks(c)%field%viscosity,                 &
                  chunks(c)%field%xvel0,                     &
                  chunks(c)%field%yvel0                      )
          end if
       ENDIF
    ENDDO    
  END SUBROUTINE viscosity
END MODULE viscosity_module