Homec4science
Diffusion miniapp_gbs_gradients (master) Browse

R11986/gradz_n2n_fd4.F90
1bfd752d5fccmaster

Last Change
Raw File

gradz_n2n_fd4.F90
View Options

program main
use space_grid
use gradients
use prec_const
#ifdef _OPENMP
use omp_lib, ONLY : omp_get_num_devices, omp_get_initial_device, omp_get_default_device
#elif _OPENACC
use openacc
#endif
#ifdef _CUDA
use cudafor
#endif
implicit none
! INPUT: Number of cells in each direction
integer :: nx = 10000, ny = 10000, nz = 10
!
real(dp), dimension(:, :, :), allocatable :: fin,fout ! array in input and array in output (computed from different gradients)
real(dp), dimension(:, :, :), allocatable :: fin_omp, fout_omp ! same but computed with openmp
#ifdef _CUDA
real(dp), dimension(:, :, :), device, allocatable :: fin_d, fout_d
real(dp), DIMENSION(1:4), device :: coef_int_d, coef_der1_stag_d, coef_der1_n2n_d, coef_der2_stag_d
integer, device :: xs, xe, ys, ye, zs, ze
real :: simTime
type(cudaEvent) :: startEvent, stopEvent
integer :: istat
type(dim3) :: grid_sp, block_sp
#endif
integer :: i, j, k, ix, iy, iz
real :: start, finish, t0s, t0e, t1s, t1e, t2s, t2e, t3s, t3e
integer :: ii, it, h = 0, t = 0, nd = 0
integer :: asize
integer, dimension(1:3) :: block_size
real(dp), dimension(1:4) :: coef_der
!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!
! Compute mesh
!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!
!
#ifdef _OPENMP
!$omp parallel
h = omp_get_initial_device()
t = omp_get_default_device()
nd = omp_get_num_devices()
!$omp end parallel
#elif _OPENACC
!!$acc parallel
!h = acc_get_initial_device()
!t = acc_get_default_device()
nd = acc_get_num_devices(acc_device_nvidia)
!!$acc end parallel
#endif
!
write(6,*) ' - Initial device = ', h, 'default device =', t, 'number of devices ', nd
!
write(*,*) ' - Computing mesh'
call compute_mesh(nx,ny,nz)
asize = (ixeg - ixsg)*(iyeg - iysg)*(izeg - izsg)
write(*,*) ' size = ', asize, ixeg, ixsg, iyeg, iysg, izeg, izsg
!
! Initialize coefficient for gradients (as in gbs)
!
write(*,*) ' - Initializing coefficients'
call intialize_coefficients
!
! allocate arrays (as in gbs)
!
write(*,*) ' - Allocating the arrays'
call gbs_allocate(fin, iysg, iyeg, ixsg, ixeg, izsg, izeg)
call gbs_allocate(fout, iysg, iyeg, ixsg, ixeg, izsg, izeg)
!
! initialize input and output arrays
!
write(*,*) ' - Initializing the arrays'
call init_array(fin)
call init_array(fout)
write(*,*) ' L2_norm = ', norm2(fin(iys:iye,ixs:ixe,izs:ize)), 'Linf norm = ', &
& maxval(fin(iys:iye,ixs:ixe,izs:ize))
!
! Benchmark is starting now
!
#ifdef _CUDA
!
allocate(fin_d (iysg:iyeg, ixsg:ixeg, izsg:izeg))
allocate(fout_d(iysg:iyeg, ixsg:ixeg, izsg:izeg))
call cpu_time(t0s)
!call cudamalloc(fout_d, asize)
fin_d = fin
!fout_d = fout
call cpu_time(t0e)
write(*,*) '0 - data transfer: ', asize*8/(t0e - t0s)/1e9, ' GiB/s'
!
coef_der1_n2n_d = coef_der1_n2n
!
block_size = [16, 16, 4]
block_sp = dim3(block_size(1), block_size(2), block_size(3))
grid_sp = dim3((ixeg - ixsg)/block_sp%x, (iyeg - iysg)/block_sp%y, (izeg - izsg)/block_sp%z)
!
call cpu_time(t1s)
call gradz_n2n_fd4_cuda<<< grid_sp, block_sp >>>(fin_d(:,:,:), fout_d(:,:,:), ixsg, ixeg, iysg, iyeg, izsg, izeg, coef_der1_n2n_d(1:4))
istat = cudaDeviceSynchronize()
call cpu_time(t1e)
write(*,*) '1 - kernel time: ', t1e - t1s
!
call cpu_time(t2s)
fout = fout_d
call cpu_time(t2e)
write(*,*) '2 - data transfer: ', asize*8/(t2e - t2s)/1e9, ' GiB/s'
#else
!
coef_der(:) = deltazi*coef_der1_n2n(:)
!
!
! Moving to the device
!
call cpu_time(start)
!$omp target enter data map(to:fin(iysg:iyeg, ixsg:ixeg, izsg:izeg)) map(alloc:fout(iysg:iyeg, ixsg:ixeg, izsg:izeg))
!$acc data copyin(fin, coef_der(1:4)) copyout(fout)
!!$omp target data map(tofrom: fout, fin)
! Done moving to the device
call cpu_time(finish)
write(*,*) '0 - data trasnfer: ', asize*8_dp/1e9/(finish - start), ' GiB/s, ', finish - start, ' s'
write(*,*) '1 - Starting the benchmarks'
!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!
! Test gradients
!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!
call cpu_time(start)
!!$omp target map(tofrom: fin, fout)
!!$omp target teams distribute parallel do simd collapse(3)
!!$omp target data map(to: fin) map(alloc: fout)
!!$omp target data map(to: fin)
#if 1
call gradz_n2n_fd4_omp(fin(:,:,:), fout(:,:,:), coef_der(1:4)) !, coef_der1_n2n(:))
#else
!$omp target teams distribute parallel do simd collapse(3)
do iz=izs,ize
do ix=ixsg,ixeg
do iy=iysg,iyeg
fout(iy, ix, iz) = coef_der(1)*fin(iy, ix, iz-2) &
+ coef_der(2)*fin(iy, ix, iz-1) &
+ coef_der(3)*fin(iy, ix, iz+1) &
+ coef_der(4)*fin(iy, ix, iz+2)
! fout(iy, ix, iz) = 1_dp
end do
end do
end do
#endif
call cpu_time(finish)
write(*,*) '2 - gradz_n2n_fd4: ', finish - start, ' s'
call cpu_time(start)
#ifndef __GNUC__
!!$acc end kernels
!!$omp end target data
!!$omp end target teams distribute parallel do
!$omp target update from(fout)
!!$omp target end data from(fout)
!$omp target exit data map(from:fout) map(release:fin)
#endif
!$acc end data
call cpu_time(finish)
#endif
write(*,*) ' gradz_n2n_fd4 : ', finish - start, ' seconds, BW = ', 2*nx*ny*nz/(finish - start)/1024/1024/1024, ' GB/s'
!write(*,*) 'gradz_n2n_fd4: ', finish - start, ' seconds'
write(*,*) 'gradz_n2n_fd4: L2_norm = ', norm2(fout(iys:iye,ixs:ixe,izs:ize)), 'Linf norm = ', &
& maxval(fout(iys:iye,ixs:ixe,izs:ize))
write(*,*) 'Exiting...'
end program main
c4science · Help