Homec4science
Diffusion wireles-gpu (master) Browse

wireles-gpu/src/divstress_gpu.cuf
61fde96d0eeemaster

Last Change
Raw File

divstress_gpu.cuf
View Options

module divergence
use precision
use dimen
contains
subroutine divstress_gpu(divt,tx,ty,tz,stagger_flag,t,plan)
!---------------------------------------------------------------------------
! declaration
!---------------------------------------------------------------------------
use cudafor
use derivative
implicit none
integer :: i,j,k,t,stagger_flag
real(fp),dimension(nx,ny,nz2),intent(inout),device:: divt
real(fp),dimension(nx,ny,nz2),intent(inout),device:: tx,ty,tz
real(fp),dimension(:,:,:),allocatable,device:: tdx,tdy,tdz
real(fp),dimension(nx,ny,nz2):: tdx_host,tdy_host,tdz_host
type(c_ptr),intent(in) :: plan(2)
! Cuda variable
type(dim3) :: grid, tBlock
grid = dim3(nx/16,ny/16,nz2/2)
tBlock = dim3(16,16,2)
save tdx,tdy,tdz
!---------------------------------------------------------------------------
! main code
!---------------------------------------------------------------------------
if (t==1) then
allocate(tdx(nx,ny,nz2))
allocate(tdy(nx,ny,nz2))
allocate(tdz(nx,ny,nz2))
end if
!---------------------------------------------------------------------------
! stress gradients
call ddxy(tdx,tx,1,t,plan)
call ddxy(tdy,ty,2,t,plan)
call ddz_gpu<<<grid, tBlock>>>(tdz,tz,stagger_flag)
!---------------------------------------------------------------------------
! stress divergence
call div_sum<<<grid, tBlock>>>(divt,tdx,tdy,tdz)
! stop
end subroutine divstress_gpu
attributes(global) subroutine div_sum(divt,tdx,tdy,tdz)
!---------------------------------------------------------------------------
! declaration
!---------------------------------------------------------------------------
implicit none
real(fp), dimension(nx,ny,nz2), intent(out) :: divt
real(fp), dimension(nx,ny,nz2), intent(in) :: tdx,tdy,tdz
integer :: i,j,k
i = (blockIdx%x - 1) * blockDim%x + threadIdx%x
j = (blockIdx%y - 1) * blockDim%y + threadIdx%y
k = (blockIdx%z - 1) * blockDim%z + threadIdx%z
! print *,i
!---------------------------------------------------------------------------
! main code
!---------------------------------------------------------------------------
! do k=2,nzb+1
if (k>1 .and. k<nz2) divt(i,j,k)=tdx(i,j,k)+tdy(i,j,k)+tdz(i,j,k)
! end do
end subroutine div_sum
attributes(global) subroutine div_p(rhs_p,dudx,dvdy,dwdz)
!---------------------------------------------------------------------------
! declaration
!---------------------------------------------------------------------------
implicit none
real(fp), dimension(nx,ny,nz2), intent(out) :: rhs_p
real(fp), dimension(nx,ny,nz2), intent(in) :: dudx,dvdy,dwdz
integer :: i,j,k
i = (blockIdx%x - 1) * blockDim%x + threadIdx%x
j = (blockIdx%y - 1) * blockDim%y + threadIdx%y
k = (blockIdx%z - 1) * blockDim%z + threadIdx%z
! print *,i
!---------------------------------------------------------------------------
! main code
!---------------------------------------------------------------------------
! do k=2,nzb+1
if (k>1 .and. k<nz2) rhs_p(i,j,k)=(dudx(i,j,k)+dvdy(i,j,k)+dwdz(i,j,k))/dt
! end do
end subroutine div_p
end module divergence
c4science · Help