Gpu 2decomp dev

.gitmodules

@@ -1,10 +1,9 @@
-[submodule "2decomp-fft"]
-	path = 2decomp-fft
-	url = https://github.com/pbartholomew08/2decomp-fft.git
-	branch = update
-
 [submodule "tools/View3D"]
 	path = tools/View3D
 	url = https://github.com/uDALES/View3D.git
 	branch = master
 
+[submodule "2decomp-fft"]
+	path = 2decomp-fft
+	url = https://github.com/pbartholomew08/2decomp-fft.git
+	branch = feat--halo-exchange--udales

CMakeLists.txt

@@ -55,7 +55,7 @@ elseif(CMAKE_Fortran_COMPILER_ID MATCHES "Cray")
     set(CMAKE_Fortran_FLAGS_DEBUG "-G0 -O0 -ei -R bcdps -m 0")
     set(CMAKE_Fortran_FLAGS_RELEASE "-O3")
 elseif(CMAKE_Fortran_COMPILER_ID MATCHES "NVHPC")
-    set(CMAKE_Fortran_FLAGS "-gpu=cc86,ccall -cuda -acc -mp -Mpreprocess -cpp -r8 -Mextend -Ktrap=divz,inv,ovf,fp")
+    set(CMAKE_Fortran_FLAGS "-gpu=cc86,ccall,lineinfo -cuda -cudalib=cufft -acc -mp -Minfo=accel -Mpreprocess -cpp -r8 -Mextend -Ktrap=divz,inv,ovf,fp")
     set(CMAKE_Fortran_FLAGS_DEBUG "-g -O0 -Mincoming=nan -Mbounds -Mwarn -Mchkstk -traceback -Minfo=all")
     set(CMAKE_Fortran_FLAGS_RELEASE "-fast -Kieee")
     add_definitions(-D_GPU)
@@ -128,9 +128,7 @@ set_property(CACHE FFT_Choice PROPERTY STRINGS generic fftw3)
 add_executable(u-dales ${src_f90})
 #include_directories(2decomp-fft)
 include_directories(${CMAKE_SOURCE_DIR}/2decomp-fft/build/include)
-if(CMAKE_Fortran_COMPILER_ID MATCHES "NVHPC")
-    include_directories(/opt/nvidia/hpc_sdk/Linux_x86_64/24.7/math_libs/12.5/targets/x86_64-linux/include)
-endif()
+
 target_link_libraries(u-dales ${CMLIB_CMAKE_LIBRARY})
 target_link_libraries(u-dales MPI::MPI_Fortran)
 target_link_libraries(u-dales ${NETCDF_F90_LIBRARIES})
@@ -143,9 +141,6 @@ message("FFTW_LIBRARIES :  (${FFTW_LIBRARIES}).")
 target_link_libraries(u-dales ${FFTW_LIBRARIES})
 #target_link_libraries(u-dales 2decomp-fft)
 target_link_libraries(u-dales ${CMAKE_SOURCE_DIR}/2decomp-fft/build/lib/libdecomp2d.a)
-if(CMAKE_Fortran_COMPILER_ID MATCHES "NVHPC")
-    target_link_libraries(u-dales /opt/nvidia/hpc_sdk/Linux_x86_64/24.7/math_libs/12.5/targets/x86_64-linux/lib/libcufft.so)
-endif()
 
 #target_include_directories(u-dales PUBLIC $<TARGET_PROPERTY:2decomp-fft,BINARY_DIR>)
 target_include_directories(u-dales PUBLIC ${FFTW_INCLUDE_DIRS})

docs/udales-geometry-tutorial.md

@@ -329,21 +329,17 @@ help udgeom.createRealistic
                      that contains the buildings. The function adds the
                      ground surface.  
 
-     geom = createRealistic(stlfile, xsize_og, ysize_og, shift, edgelength)
+     geom = createRealistic(stlfile, xsize, ysize, shift, edgelength)
      returns a geom instance that can be saved to an stl file. 
          stlfile:    the STL file that contains the buildings (NOT the
                      ground)
-         xsize_og:   length of the original domain in x-direction
-         ysize_og:   length of the original domain in y-direction
+         xsize:      length of the new domain in x-direction
+         ysize:      length of the new domain in y-direction
          shift:      array that shifts the geometry. shift[1], shift[2] and
                      shift[3] represent, respectively, the shift in x-, y-
                      and z-direction.
          edgelength: the length of individual facets. Best taken as xsize
                      (or ysize) divided by an integer number.
-
-  Note that the output domain size is
-      xsize = xsize_og + shift[1]
-      ysize = ysize_og + shift[2]
 ```
 
 
@@ -357,16 +353,16 @@ See below for an example how to use this function.
 stlfile = 'uDALES.stl';
 
 % original domain size
-xsize_og = 256; 
-ysize_og = 128; 
+xsize = 256; 
+ysize = 128; 
 
 % example translation to make domain larger in x direction
 shift = [20 0 0];
 
 % ground facet size
 edgelength = 16;
 
-geom = udgeom.createRealistic(stlfile, xsize_og, ysize_og, shift, edgelength);
+geom = udgeom.createRealistic(stlfile, xsize, ysize, shift, edgelength);
 geom.show; 
 xlabel('x [m]'); 
 ylabel('y [m]')

src/heatpump.f90

@@ -77,7 +77,7 @@ subroutine heatpump
         if (lfan_hp) then ! Heat pump fan is on
           wm(i,j,k+1) = w_hp_exhaust ! Set exhaust velocity at heat pump point [m/s], at input 'w' cell face k+1
           w0(i,j,k+1) = w_hp_exhaust
-          !wp(i,j,k+1) = 0.
+          wp(i,j,k+1) = 0.
         end if
 
         thlp(i,j,k) = thlp(i,j,k) - thl_dot_hp * dxi * dyi * dzfi(k)  ! [K/s], at cell center k

src/modboundary.f90

@@ -72,26 +72,28 @@ subroutine halos
                             BCym_periodic, BCyT_periodic, BCyq_periodic, BCys_periodic, &
                             ibrank, ierank, jbrank, jerank
       use modfields, only : u0, v0, w0, um, vm, wm, thl0, thlm, qt0, qtm, sv0, svm, thl0c
-      use decomp_2d, only : exchange_halo_z
+      use m_halo, only : halo_exchange
+      use decomp_2d, only: decomp_main
+
       implicit none
       integer i, k, n
 
 !$acc data create(u0, v0, w0, um, vm, wm, thl0, thlm, thl0c, qt0, qtm, sv0, svm)
 !$acc update device(u0, v0, w0, um, vm, wm, thl0, thlm, thl0c, qt0, qtm, sv0, svm)
-      call exchange_halo_z(u0)
-      call exchange_halo_z(v0)
-      call exchange_halo_z(w0)
-      call exchange_halo_z(um)
-      call exchange_halo_z(vm)
-      call exchange_halo_z(wm)
-      call exchange_halo_z(thl0)
-      call exchange_halo_z(thlm)
-      call exchange_halo_z(thl0c, opt_zlevel=(/ihc,jhc,khc/))
-      call exchange_halo_z(qt0)
-      call exchange_halo_z(qtm)
+      call halo_exchange(u0, 3)
+      call halo_exchange(v0, 3)
+      call halo_exchange(w0, 3)
+      call halo_exchange(um, 3)
+      call halo_exchange(vm, 3)
+      call halo_exchange(wm, 3)
+      call halo_exchange(thl0, 3)
+      call halo_exchange(thlm, 3)
+      call halo_exchange(thl0c, 3, opt_levels=(/ihc,jhc,khc/))
+      call halo_exchange(qt0, 3)
+      call halo_exchange(qtm, 3)
       do n = 1, nsv
-         call exchange_halo_z(sv0(:, :, :, n), opt_zlevel=(/ihc,jhc,khc/))
-         call exchange_halo_z(svm(:, :, :, n), opt_zlevel=(/ihc,jhc,khc/))
+         call halo_exchange(sv0(:, :, :, n), 3, opt_levels=(/ihc,jhc,khc/))
+         call halo_exchange(svm(:, :, :, n), 3, opt_levels=(/ihc,jhc,khc/))
       enddo
 !$acc update host(u0, v0, w0, um, vm, wm, thl0, thlm, thl0c, qt0, qtm, sv0, svm)
 !$acc end data
@@ -538,7 +540,8 @@ subroutine closurebc
                                 ibrank, ierank, jbrank, jerank, BCtopm, BCxm, BCym, &
                                 BCtopm_freeslip, BCtopm_noslip, BCtopm_pressure, &
                                 BCxm_periodic, BCym_periodic
-     use decomp_2d,      only : exchange_halo_z, zstart, zend
+     use decomp_2d,      only : zstart, zend
+     use m_halo,         only : halo_exchange
 #if defined(_GPU)
      use cudafor
      use modcuda,        only : ekm_d, ekh_d, griddim, blockdim, checkCUDA
@@ -547,11 +550,11 @@ subroutine closurebc
 
 #if defined(_GPU)
      call checkCUDA( cudaDeviceSynchronize(), 'cudaDeviceSynchronize in closurebc' )
-     call exchange_halo_z(ekm_d)
-     call exchange_halo_z(ekh_d)
+     call halo_exchange(ekm_d, 3)
+     call halo_exchange(ekh_d, 3)
 #else
-     call exchange_halo_z(ekm)
-     call exchange_halo_z(ekh)
+     call halo_exchange(ekm, 3)
+     call halo_exchange(ekh, 3)
 #endif
 
 #if defined(_GPU)
@@ -1322,7 +1325,7 @@ subroutine bcpup(pup, pvp, pwp, rk3coef)
      use modfields,    only : pres0, up, vp, wp, um, vm, wm, w0, u0, v0, uouttot, vouttot, uinit, vinit, uprof, vprof, pres0, IIc, IIcs
      use modmpi,       only : excjs, excis, myid, avexy_ibm
      use modinletdata, only : u0driver
-     use decomp_2d,    only : exchange_halo_z
+     use m_halo,       only : halo_exchange
 
      real, dimension(ib - ih:ie + ih, jb - jh:je + jh, kb:ke + kh), intent(inout) :: pup
      real, dimension(ib - ih:ie + ih, jb - jh:je + jh, kb:ke + kh), intent(inout) :: pvp
@@ -1342,9 +1345,9 @@ subroutine bcpup(pup, pvp, pwp, rk3coef)
      ! maybe safer to just resize to kb-kh:ke+kh
 !$acc data create(pup, pvp, pwp)
 !$acc update device(pup, pvp, pwp)
-     call exchange_halo_z(pup, opt_zlevel=(/ih,jh,0/))
-     call exchange_halo_z(pvp, opt_zlevel=(/ih,jh,0/))
-     call exchange_halo_z(pwp, opt_zlevel=(/ih,jh,0/))
+     call halo_exchange(pup, 3, opt_levels=(/ih,jh,0/))
+     call halo_exchange(pvp, 3, opt_levels=(/ih,jh,0/))
+     call halo_exchange(pwp, 3, opt_levels=(/ih,jh,0/))
 !$acc update host(pup, pvp, pwp)
 !$acc end data
      ! if (jbrank) write(*,*) "jb after exhange_halo ", pvp(ie/2,jb,ke)
@@ -1474,7 +1477,7 @@ subroutine bcp(p)
      use modglobal, only : ib, ie, jb, je, ih, jh, kb, ke, kh, dyi, rk3step, dt, &
                            ibrank, ierank, jbrank, jerank, BCxm, BCym, BCxm_periodic, BCym_periodic
      use modfields, only : pres0, up, u0, um, uouttot, vp, v0
-     use decomp_2d, only : exchange_halo_z
+     use m_halo,    only : halo_exchange
 
      real, dimension(ib - ih:ie + ih, jb - jh:je + jh, kb - kh:ke + kh), intent(inout) :: p !< pressure
      integer i, j, k
@@ -1489,8 +1492,8 @@ subroutine bcp(p)
 
 !$acc data create(p, pres0)
 !$acc update device(p, pres0)
-     call exchange_halo_z(p)
-     call exchange_halo_z(pres0)
+     call halo_exchange(p, 3)
+     call halo_exchange(pres0, 3)
 !$acc update host(p, pres0)
 !$acc end data
 

src/modfields.f90

@@ -587,10 +587,10 @@ subroutine initfields
     ! call alloc_z(wp); wp = 0.
     call alloc_z(pres0); pres0 = 0.
 
-    call alloc_z(div, opt_zlevel=(/0,0,0/))
-    call alloc_z(dudx, opt_zlevel=(/0,0,0/))
-    call alloc_z(dvdy, opt_zlevel=(/0,0,0/))
-    call alloc_z(dwdz, opt_zlevel=(/0,0,0/))
+    call alloc_z(div, opt_levels=(/0,0,0/))
+    call alloc_z(dudx, opt_levels=(/0,0,0/))
+    call alloc_z(dvdy, opt_levels=(/0,0,0/))
+    call alloc_z(dwdz, opt_levels=(/0,0,0/))
 
     ! Always have to allocate these, even if ltempeq/lmoist = .false.
     call alloc_z(thlm); thlm = 0.
@@ -610,7 +610,7 @@ subroutine initfields
     call alloc_z(e12m); e12m = 0.
     call alloc_z(e120); e120 = 0.
     ! call alloc_z(e12p); e12p = 0.
-    call alloc_z(mindist, opt_zlevel=(/0,0,0/)); mindist = 0.
+    call alloc_z(mindist, opt_levels=(/0,0,0/)); mindist = 0.
 
     ! These always have to be allocated differently?
     !if (nsv>0) then

src/modibm.f90

@@ -128,7 +128,7 @@ subroutine initibm
      use modglobal, only : libm, xh, xf, yh, yf, zh, zf, xhat, yhat, zhat, vec0, &
                            ib, ie, ih, ihc, jb, je, jh, jhc, kb, ke, kh, khc, nsv, &
                            iwallmom, lmoist, ltempeq, cexpnr, nfcts, lwritefac
-     use decomp_2d, only : exchange_halo_z
+     use m_halo, only : halo_exchange
      use modmpi,    only : myid
      use modstat_nc,only: open_nc, define_nc, ncinfo, writestat_dims_nc
 
@@ -160,9 +160,9 @@ subroutine initibm
      call solid(solid_info_w, mask_w, rhs, 0., ih, jh, kh)
 !$acc data create(mask_u, mask_v, mask_w)
 !$acc update device(mask_u, mask_v, mask_w)
-     call exchange_halo_z(mask_u)!, opt_zlevel=(/ih,jh,0/))
-     call exchange_halo_z(mask_v)!, opt_zlevel=(/ih,jh,0/))
-     call exchange_halo_z(mask_w)!, opt_zlevel=(/ih,jh,0/))
+     call halo_exchange(mask_u, 3)!, opt_zlevel=(/ih,jh,0/))
+     call halo_exchange(mask_v, 3)!, opt_zlevel=(/ih,jh,0/))
+     call halo_exchange(mask_w, 3)!, opt_zlevel=(/ih,jh,0/))
 !$acc update host(mask_u, mask_v, mask_w)
 !$acc end data
 
@@ -191,7 +191,7 @@ subroutine initibm
        call solid(solid_info_c, mask_c, rhs, 0., ih, jh, kh)
 !$acc data create(mask_c)
 !$acc update device(mask_c)
-       call exchange_halo_z(mask_c)!, opt_zlevel=(/ih,jh,0/))
+       call halo_exchange(mask_c, 3)!, opt_zlevel=(/ih,jh,0/))
 !$acc update host(mask_c)
 !$acc end data
      end if
@@ -2320,7 +2320,8 @@ subroutine createmasks
                             IIcs, IIus, IIvs, IIws, IIuws, IIvws, IIuvs, &
                             IIct, IIut, IIvt, IIwt, IIuwt, um, u0, vm, v0, wm, w0
       use modmpi,    only : myid, comm3d, mpierr, MY_REAL, nprocs
-      use decomp_2d, only : zstart, exchange_halo_z
+      use decomp_2d, only : zstart
+      use m_halo,    only : halo_exchange
 
       integer :: IIcl(kb:ke + khc), IIul(kb:ke + khc), IIvl(kb:ke + khc), IIwl(kb:ke + khc), IIuwl(kb:ke + khc), IIvwl(kb:ke + khc), IIuvl(kb:ke + khc)
       integer :: IIcd(ib:ie, kb:ke)
@@ -2404,9 +2405,9 @@ subroutine createmasks
       end do
 
       ! Can't do this because no interface for integers
-      ! call exchange_halo_z(IIuv, opt_zlevel=(/ihc,jhc,0/))
-      ! call exchange_halo_z(IIuv, opt_zlevel=(/ihc,jhc,0/))
-      ! call exchange_halo_z(IIvw, opt_zlevel=(/ihc,jhc,0/))
+      ! call halo_exchange(IIuv, 3, opt_zlevel=(/ihc,jhc,0/))
+      ! call halo_exchange(IIuv, 3, opt_zlevel=(/ihc,jhc,0/))
+      ! call halo_exchange(IIvw, 3, opt_zlevel=(/ihc,jhc,0/))
 
       do k = kb, ke + khc
          IIcl(k) = sum(IIc(ib:ie, jb:je, k))

src/modpois.f90

@@ -81,12 +81,12 @@ subroutine initpois
     allocate(pvp(ib-ih:ie+ih,jb-jh:je+jh,kb:ke+kh))
     allocate(pwp(ib-ih:ie+ih,jb-jh:je+jh,kb:ke+kh))
     call alloc_z(p)
-    call alloc_z(rhs, opt_zlevel=(/0,0,0/))
+    call alloc_z(rhs, opt_levels=(/0,0,0/))
     call alloc_z(dpupdx)
     call alloc_z(dpvpdy)
     call alloc_z(dpwpdz)
-    call alloc_z(Fxy, opt_zlevel=(/0,0,0/))
-    call alloc_z(Fxyz, opt_zlevel=(/0,0,0/))
+    call alloc_z(Fxy, opt_levels=(/0,0,0/))
+    call alloc_z(Fxyz, opt_levels=(/0,0,0/))
 
     if (ipoiss == POISS_FFT2D) then
       allocate(xrt(itot))
@@ -440,9 +440,9 @@ subroutine poisson
 
     select case (ipoiss)
     case (POISS_FFT2D)
-      call alloc_x(px, opt_xlevel=(/0,0,0/))
-      call alloc_y(py, opt_ylevel=(/0,0,0/))
-      call alloc_z(pz, opt_zlevel=(/0,0,0/))
+      call alloc_x(px, opt_levels=(/0,0,0/))
+      call alloc_y(py, opt_levels=(/0,0,0/))
+      call alloc_z(pz, opt_levels=(/0,0,0/))
 
       pz = p(ib:ie,jb:je,kb:ke)
 
@@ -763,9 +763,9 @@ subroutine poisson
 #else
       write(*,*) "POISS_FFT2D_2DECOMP on CPU."
 #endif
-      call alloc_x(px, opt_xlevel=(/0,0,0/))
-      call alloc_y(py, opt_ylevel=(/0,0,0/))
-      call alloc_z(pz, opt_zlevel=(/0,0,0/))
+      call alloc_x(px, opt_levels=(/0,0,0/))
+      call alloc_y(py, opt_levels=(/0,0,0/))
+      call alloc_z(pz, opt_levels=(/0,0,0/))
       allocate(Fx(sp%xsz(1),sp%xsz(2),sp%xsz(3)))
       allocate(Fy(sp%ysz(1),sp%ysz(2),sp%ysz(3)))
       allocate(Fz(sp%zsz(1),sp%zsz(2),sp%zsz(3)))
@@ -939,7 +939,7 @@ subroutine poisson
       allocate(Fz(sp%zsz(1),sp%zsz(2),sp%zsz(3)))
       allocate(Fzr(imax,jmax,ktot))
 
-      call alloc_z(pz, opt_zlevel=(/0,0,0/))
+      call alloc_z(pz, opt_levels=(/0,0,0/))
       pz = p(ib:ie,jb:je,kb:ke)
 
       call decomp_2d_fft_3d(pz,Fx) ! start in z-pencil in physical space, end in x-pencil in Fourier space

src/modstartup.f90

@@ -651,6 +651,7 @@ end subroutine readnamelists
 
    subroutine init2decomp
      use decomp_2d
+     use decomp_2d_mpi, only : DECOMP_2D_COMM_CART_Z
      use modglobal, only : itot, jtot, ktot, BCxm, BCym, BCxm_periodic, Bcym_periodic
      use modmpi,    only : comm3d, myid, myidx, myidy, cmyidx, cmyidy, nprocx, nprocy, &
                            nbreast, nbrwest, nbrnorth, nbrsouth, mpierr
@@ -953,7 +954,7 @@ subroutine readinitfiles
          nstepreaddriver
       use modinlet, only:readinletfile
       use moddriver, only: readdriverfile,initdriver,drivergen,readdriverfile_chunk
-      use decomp_2d, only : exchange_halo_z, update_halo, decomp_main
+      use decomp_2d, only : decomp_main
 
       integer i, j, k, n
 

src/modstatsdump.f90

@@ -1688,7 +1688,7 @@ subroutine tkestatsdump
   use modmpi,           only : myid,cmyid,my_real,mpi_sum,avey_ibm,mpierr,comm3d,excjs,avexy_ibm
   use modsurfdata,      only : thls
   use modsubgrid,       only : ekh
-  use decomp_2d,        only : exchange_halo_z
+  use m_halo,           only : halo_exchange
   implicit none
 
   real, dimension(ib:ie,jb:je,kb:ke)  :: disssgsfl     ! average subgrid visc. * average rate of strain squared : 2*<nu_t>*<Sij>*<Sij>
@@ -1815,17 +1815,17 @@ subroutine tkestatsdump
 
 !$acc data create(tvmx, tsgsmx1, tsgsmx2, dummyx, ttmx, tvmy, tsgsmy1, tsgsmy2, dummyy, ttmy)
 !$acc update device(tvmx, tsgsmx1, tsgsmx2, dummyx, ttmx, tvmy, tsgsmy1, tsgsmy2, dummyy, ttmy)
-      call exchange_halo_z(tvmx, opt_zlevel=(/ih,jh,0/))
-      call exchange_halo_z(tsgsmx1, opt_zlevel=(/ih,jh,0/))
-      call exchange_halo_z(tsgsmx2, opt_zlevel=(/ih,jh,0/))
-      call exchange_halo_z(dummyx, opt_zlevel=(/ih,jh,0/))
-      call exchange_halo_z(ttmx, opt_zlevel=(/ih,jh,0/))
-
-      call exchange_halo_z(tvmy, opt_zlevel=(/ih,jh,0/))
-      call exchange_halo_z(tsgsmy1, opt_zlevel=(/ih,jh,0/))
-      call exchange_halo_z(tsgsmy2, opt_zlevel=(/ih,jh,0/))
-      call exchange_halo_z(dummyy, opt_zlevel=(/ih,jh,0/))
-      call exchange_halo_z(ttmy, opt_zlevel=(/ih,jh,0/))
+      call halo_exchange(tvmx, 3, opt_levels=(/ih,jh,0/))
+      call halo_exchange(tsgsmx1, 3, opt_levels=(/ih,jh,0/))
+      call halo_exchange(tsgsmx2, 3, opt_levels=(/ih,jh,0/))
+      call halo_exchange(dummyx, 3, opt_levels=(/ih,jh,0/))
+      call halo_exchange(ttmx, 3, opt_levels=(/ih,jh,0/))
+
+      call halo_exchange(tvmy, 3, opt_levels=(/ih,jh,0/))
+      call halo_exchange(tsgsmy1, 3, opt_levels=(/ih,jh,0/))
+      call halo_exchange(tsgsmy2, 3, opt_levels=(/ih,jh,0/))
+      call halo_exchange(dummyy, 3, opt_levels=(/ih,jh,0/))
+      call halo_exchange(ttmy, 3, opt_levels=(/ih,jh,0/))
 !$acc update host(tvmx, tsgsmx1, tsgsmx2, dummyx, ttmx, tvmy, tsgsmy1, tsgsmy2, dummyy, ttmy)
 !$acc end data