Identify locations

Exawind · Jul 8, 2024 · b898303 · b898303
1 parent 0230d3a
commit b898303
Show file tree

Hide file tree

Showing 100 changed files with 225 additions and 116 deletions.
diff --git a/amr-wind/CMakeLists.txt b/amr-wind/CMakeLists.txt
@@ -20,7 +20,6 @@ target_include_directories(${amr_wind_lib_name} PUBLIC
 add_subdirectory(core)
 add_subdirectory(boundary_conditions)
 add_subdirectory(convection)
-add_subdirectory(derive)
 add_subdirectory(diffusion)
 add_subdirectory(projection)
 add_subdirectory(setup)

diff --git a/amr-wind/core/Field.cpp b/amr-wind/core/Field.cpp
@@ -391,22 +391,17 @@ void Field::to_uniform_space() noexcept
 
     // scale velocity to accommodate for mesh mapping -> U^bar = U * J/fac
     for (int lev = 0; lev < m_repo.num_active_levels(); ++lev) {
-        for (amrex::MFIter mfi(mesh_fac(lev)); mfi.isValid(); ++mfi) {
-
-            amrex::Array4<amrex::Real> const& field = operator()(lev).array(
-                mfi);
-            amrex::Array4<amrex::Real const> const& fac =
-                mesh_fac(lev).const_array(mfi);
-            amrex::Array4<amrex::Real const> const& detJ =
-                mesh_detJ(lev).const_array(mfi);
-
-            amrex::ParallelFor(
-                mfi.growntilebox(), AMREX_SPACEDIM,
-                [=] AMREX_GPU_DEVICE(int i, int j, int k, int n) noexcept {
-                    field(i, j, k, n) *= detJ(i, j, k) / fac(i, j, k, n);
-                });
-        }
+        const auto& fac = mesh_fac(lev).const_arrays();
+        const auto& detJ = mesh_detJ(lev).const_arrays();
+        const auto& field = operator()(lev).arrays();
+        amrex::ParallelFor(
+            mesh_fac(lev), amrex::IntVect(0), operator()(lev).nComp(),
+            [=] AMREX_GPU_DEVICE(int nbx, int i, int j, int k, int n) noexcept {
+                field[nbx](i, j, k, n) *=
+                    detJ[nbx](i, j, k) / fac[nbx](i, j, k, n);
+            });
     }
+    amrex::Gpu::synchronize();
     m_mesh_mapped = true;
 }
 
@@ -426,21 +421,18 @@ void Field::to_stretched_space() noexcept
 
     // scale field back to stretched mesh -> U = U^bar * fac/J
     for (int lev = 0; lev < m_repo.num_active_levels(); ++lev) {
-        for (amrex::MFIter mfi(mesh_fac(lev)); mfi.isValid(); ++mfi) {
-            amrex::Array4<amrex::Real> const& field = operator()(lev).array(
-                mfi);
-            amrex::Array4<amrex::Real const> const& fac =
-                mesh_fac(lev).const_array(mfi);
-            amrex::Array4<amrex::Real const> const& detJ =
-                mesh_detJ(lev).const_array(mfi);
-
-            amrex::ParallelFor(
-                mfi.growntilebox(), AMREX_SPACEDIM,
-                [=] AMREX_GPU_DEVICE(int i, int j, int k, int n) noexcept {
-                    field(i, j, k, n) *= fac(i, j, k, n) / detJ(i, j, k);
-                });
-        }
+
+        const auto& fac = mesh_fac(lev).const_arrays();
+        const auto& detJ = mesh_detJ(lev).const_arrays();
+        const auto& field = operator()(lev).arrays();
+        amrex::ParallelFor(
+            mesh_fac(lev), num_grow(), operator()(lev).nComp(),
+            [=] AMREX_GPU_DEVICE(int nbx, int i, int j, int k, int n) noexcept {
+                field[nbx](i, j, k, n) *=
+                    fac[nbx](i, j, k, n) / detJ[nbx](i, j, k);
+            });
     }
+    amrex::Gpu::synchronize();
     m_mesh_mapped = false;
 }
 

diff --git a/amr-wind/core/FieldFillPatchOps.H b/amr-wind/core/FieldFillPatchOps.H
@@ -373,7 +373,7 @@ public:
             const auto& dbx =
                 ori.isLow() ? amrex::adjCellLo(domain, idir, nghost[idir])
                             : amrex::adjCellHi(domain, idir, nghost[idir]);
-
+            // needs openmp pragma?
             for (amrex::MFIter mfi(mfab); mfi.isValid(); ++mfi) {
                 const auto& gbx = amrex::grow(mfi.validbox(), nghost);
                 const auto& bx = gbx & dbx;

diff --git a/amr-wind/core/field_ops.H b/amr-wind/core/field_ops.H
@@ -310,18 +310,15 @@ lower_bound(FType& field, const amrex::Real min_value, const int icomp = 0)
     const auto& repo = field.repo();
     const int nlevels = repo.num_active_levels();
     for (int lev = 0; lev < nlevels; ++lev) {
-
-        for (amrex::MFIter mfi(field(lev)); mfi.isValid(); ++mfi) {
-            const auto& bx = mfi.tilebox();
-            const auto& field_arr = field(lev).array(mfi);
-
-            amrex::ParallelFor(
-                bx, [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
-                    field_arr(i, j, k, icomp) =
-                        amrex::max(min_value, field_arr(i, j, k, icomp));
-                });
-        }
+        const auto& farrs = field(lev).arrays();
+        amrex::ParallelFor(
+            field(lev),
+            [=] AMREX_GPU_DEVICE(int nbx, int i, int j, int k) noexcept {
+                farrs[nbx](i, j, k, icomp) =
+                    amrex::max(min_value, farrs[nbx](i, j, k, icomp));
+            });
     }
+    amrex::Gpu::synchronize();
 }
 
 /** Computes the global maximum of a field from all levels
@@ -378,28 +375,25 @@ inline void normalize(FType& field)
 
     const int nlevels = repo.num_active_levels();
     for (int lev = 0; lev < nlevels; ++lev) {
-
-        for (amrex::MFIter mfi(field(lev)); mfi.isValid(); ++mfi) {
-            const auto& bx = mfi.tilebox();
-            const auto& field_arr = field(lev).array(mfi);
-
-            amrex::ParallelFor(
-                bx, [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
-                    // Compute magnitude
-                    amrex::Real mag = 0.;
+        const auto& farrs = field(lev).arrays();
+        amrex::ParallelFor(
+            field(lev),
+            [=] AMREX_GPU_DEVICE(int nbx, int i, int j, int k) noexcept {
+                auto farr = farrs[nbx];
+                // Compute magnitude
+                amrex::Real mag = 0.;
+                for (int icomp = 0; icomp < ncomp; ++icomp) {
+                    mag = mag + farr(i, j, k, icomp) * farr(i, j, k, icomp);
+                }
+                if (mag > eps) {
                     for (int icomp = 0; icomp < ncomp; ++icomp) {
-                        mag = mag + field_arr(i, j, k, icomp) *
-                                        field_arr(i, j, k, icomp);
-                    }
-                    if (mag > eps) {
-                        for (int icomp = 0; icomp < ncomp; ++icomp) {
-                            field_arr(i, j, k, icomp) =
-                                field_arr(i, j, k, icomp) / std::sqrt(mag);
-                        }
+                        farr(i, j, k, icomp) =
+                            farr(i, j, k, icomp) / std::sqrt(mag);
                     }
-                });
-        }
+                }
+            });
     }
+    amrex::Gpu::synchronize();
 }
 
 } // namespace amr_wind::field_ops

diff --git a/amr-wind/derive/CMakeLists.txt b/amr-wind/derive/CMakeLists.txt
diff --git a/amr-wind/derive/field_algebra.cpp b/amr-wind/derive/field_algebra.cpp
diff --git a/amr-wind/diffusion/incflo_diffusion.cpp b/amr-wind/diffusion/incflo_diffusion.cpp
@@ -221,6 +221,7 @@ void viscosity_to_uniform_space(
         repo.get_mesh_mapping_detJ(amr_wind::FieldLoc::ZFACE);
 
     // beta accounted for mesh mapping (x-face) = J/fac^2 * mu
+    // needs openmp pragma?
     for (amrex::MFIter mfi(b[0]); mfi.isValid(); ++mfi) {
         amrex::Array4<amrex::Real> const& mu = b[0].array(mfi);
         amrex::Array4<amrex::Real const> const& fac =
@@ -235,6 +236,7 @@ void viscosity_to_uniform_space(
             });
     }
     // beta accounted for mesh mapping (y-face) = J/fac^2 * mu
+    // needs openmp pragma?
     for (amrex::MFIter mfi(b[1]); mfi.isValid(); ++mfi) {
         amrex::Array4<amrex::Real> const& mu = b[1].array(mfi);
         amrex::Array4<amrex::Real const> const& fac =
@@ -249,6 +251,7 @@ void viscosity_to_uniform_space(
             });
     }
     // beta accounted for mesh mapping (z-face) = J/fac^2 * mu
+    // needs openmp pragma?
     for (amrex::MFIter mfi(b[2]); mfi.isValid(); ++mfi) {
         amrex::Array4<amrex::Real> const& mu = b[2].array(mfi);
         amrex::Array4<amrex::Real const> const& fac =

diff --git a/amr-wind/equation_systems/CompRHSOps.H b/amr-wind/equation_systems/CompRHSOps.H
@@ -78,6 +78,7 @@ struct ComputeRHSOp
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+            // this one
             for (amrex::MFIter mfi(field(lev)); mfi.isValid(); ++mfi) {
                 const auto& bx = mfi.tilebox();
                 auto fld = field(lev).array(mfi);
@@ -197,6 +198,7 @@ struct ComputeRHSOp
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+            // this one
             for (amrex::MFIter mfi(field(lev)); mfi.isValid(); ++mfi) {
                 const auto& bx = mfi.tilebox();
                 auto fld = field(lev).array(mfi);

diff --git a/amr-wind/equation_systems/DiffusionOps.cpp b/amr-wind/equation_systems/DiffusionOps.cpp
@@ -128,6 +128,7 @@ void DiffSolverIface<LinOp>::linsys_solve_impl()
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+        // this one
         for (amrex::MFIter mfi(rhs, amrex::TilingIfNotGPU()); mfi.isValid();
              ++mfi) {
             const auto& bx = mfi.tilebox();

diff --git a/amr-wind/equation_systems/PDEOps.H b/amr-wind/equation_systems/PDEOps.H
@@ -100,6 +100,7 @@ struct SrcTermOpBase
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+            // this one
             for (amrex::MFIter mfi(src_term, amrex::TilingIfNotGPU());
                  mfi.isValid(); ++mfi) {
                 const auto& bx = mfi.tilebox();

diff --git a/amr-wind/equation_systems/density/density_ops.H b/amr-wind/equation_systems/density/density_ops.H
@@ -30,6 +30,7 @@ struct ComputeRHSOp<Density, Scheme>
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+            // this one
             for (amrex::MFIter mfi(field(lev)); mfi.isValid(); ++mfi) {
                 const auto& bx = mfi.tilebox();
                 auto rho = field(lev).array(mfi);
@@ -62,6 +63,7 @@ struct ComputeRHSOp<Density, Scheme>
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+            // this one
             for (amrex::MFIter mfi(field(lev)); mfi.isValid(); ++mfi) {
                 const auto& bx = mfi.tilebox();
                 auto rho = field(lev).array(mfi);

diff --git a/amr-wind/equation_systems/icns/icns_advection.cpp b/amr-wind/equation_systems/icns/icns_advection.cpp
@@ -287,6 +287,8 @@ void MacProjOp::mac_proj_to_uniform_space(
     // scale U^mac to accommodate for mesh mapping -> U^bar = J/fac *
     // U^mac beta accounted for mesh mapping = J/fac^2 * 1/rho construct
     // rho and mesh map u_mac on x-face
+    // this one
+    // needs openmp pragma?
     for (amrex::MFIter mfi(*(rho_face[0])); mfi.isValid(); ++mfi) {
         amrex::Array4<amrex::Real> const& u = u_mac(lev).array(mfi);
         amrex::Array4<amrex::Real> const& rho = rho_face[0]->array(mfi);
@@ -303,6 +305,8 @@ void MacProjOp::mac_proj_to_uniform_space(
             });
     }
     // construct rho on y-face
+    // this one
+    // needs openmp pragma?
     for (amrex::MFIter mfi(*(rho_face[1])); mfi.isValid(); ++mfi) {
         amrex::Array4<amrex::Real> const& v = v_mac(lev).array(mfi);
         amrex::Array4<amrex::Real> const& rho = rho_face[1]->array(mfi);
@@ -319,6 +323,8 @@ void MacProjOp::mac_proj_to_uniform_space(
             });
     }
     // construct rho on z-face
+    // this one
+    // needs openmp pragma?
     for (amrex::MFIter mfi(*(rho_face[2])); mfi.isValid(); ++mfi) {
         amrex::Array4<amrex::Real> const& w = w_mac(lev).array(mfi);
         amrex::Array4<amrex::Real> const& rho = rho_face[2]->array(mfi);

diff --git a/amr-wind/equation_systems/icns/icns_diffusion.H b/amr-wind/equation_systems/icns/icns_diffusion.H
@@ -56,6 +56,7 @@ public:
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+                // this one
                 for (amrex::MFIter mfi(divtau(lev), amrex::TilingIfNotGPU());
                      mfi.isValid(); ++mfi) {
                     const auto& bx = mfi.tilebox();
@@ -198,6 +199,7 @@ public:
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+                // this one
                 for (amrex::MFIter mfi(divtau(lev), amrex::TilingIfNotGPU());
                      mfi.isValid(); ++mfi) {
                     const auto& bx = mfi.tilebox();
@@ -269,6 +271,7 @@ public:
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+            // this one
             for (amrex::MFIter mfi(rhs, amrex::TilingIfNotGPU()); mfi.isValid();
                  ++mfi) {
                 const auto& bx = mfi.tilebox();
@@ -430,6 +433,7 @@ public:
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+                // this one
                 for (amrex::MFIter mfi(divtau(lev), amrex::TilingIfNotGPU());
                      mfi.isValid(); ++mfi) {
                     const auto& bx = mfi.tilebox();
@@ -512,6 +516,7 @@ public:
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+            // this one
             for (amrex::MFIter mfi(rhs, amrex::TilingIfNotGPU()); mfi.isValid();
                  ++mfi) {
                 const auto& bx = mfi.tilebox();

diff --git a/amr-wind/equation_systems/icns/icns_ops.H b/amr-wind/equation_systems/icns/icns_ops.H
@@ -94,6 +94,7 @@ struct SrcTermOp<ICNS> : SrcTermOpBase<ICNS>
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+            // this one
             for (amrex::MFIter mfi(src_term, amrex::TilingIfNotGPU());
                  mfi.isValid(); ++mfi) {
                 const auto& bx = mfi.tilebox();

diff --git a/amr-wind/equation_systems/levelset/levelset_ops.H b/amr-wind/equation_systems/levelset/levelset_ops.H
@@ -69,6 +69,7 @@ struct ComputeRHSOp<Levelset, Scheme>
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+            // this one
             for (amrex::MFIter mfi(field(lev)); mfi.isValid(); ++mfi) {
                 const auto& bx = mfi.tilebox();
                 auto phi = field(lev).array(mfi);
@@ -101,6 +102,7 @@ struct ComputeRHSOp<Levelset, Scheme>
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
 #endif
+            // this one
             for (amrex::MFIter mfi(field(lev)); mfi.isValid(); ++mfi) {
                 const auto& bx = mfi.tilebox();
                 auto phi = field(lev).array(mfi);

diff --git a/amr-wind/equation_systems/tke/tke_ops.H b/amr-wind/equation_systems/tke/tke_ops.H
@@ -42,6 +42,7 @@ struct PostSolveOp<TKE>
         const int nlevels = repo.num_active_levels();
         const auto clip_value = m_clip_value;
         for (int lev = 0; lev < nlevels; ++lev) {
+            // this one
             for (amrex::MFIter mfi(field(lev)); mfi.isValid(); ++mfi) {
                 const auto& bx = mfi.tilebox();
                 const auto& field_arr = field(lev).array(mfi);

diff --git a/amr-wind/equation_systems/vof/vof_hybridsolver_ops.H b/amr-wind/equation_systems/vof/vof_hybridsolver_ops.H
@@ -19,6 +19,7 @@ static void replace_masked_vof(
         auto& vof = f_vof(lev);
         const auto& vof_new = f_vof_new(lev);
 
+        // this one
         for (amrex::MFIter mfi(iblank); mfi.isValid(); ++mfi) {
             const auto& gbx = mfi.growntilebox();
             const amrex::Array4<const int>& native_flag =

diff --git a/amr-wind/immersed_boundary/bluff_body/bluff_body_ops.cpp b/amr-wind/immersed_boundary/bluff_body/bluff_body_ops.cpp
@@ -47,6 +47,7 @@ void apply_mms_vel(CFDSim& sim)
         const auto& dx = geom[lev].CellSizeArray();
         const auto& problo = geom[lev].ProbLoArray();
 
+        // this one
         for (amrex::MFIter mfi(levelset(lev)); mfi.isValid(); ++mfi) {
             const auto& bx = mfi.growntilebox();
             const auto& phi = levelset(lev).const_array(mfi);
@@ -90,6 +91,7 @@ void apply_dirichlet_vel(CFDSim& sim, const amrex::Vector<amrex::Real>& vel_bc)
         // Defining the "ghost-cell" band distance
         amrex::Real phi_b = std::cbrt(dx[0] * dx[1] * dx[2]);
 
+        // this one
         for (amrex::MFIter mfi(levelset(lev)); mfi.isValid(); ++mfi) {
             const auto& bx = mfi.tilebox();
             const auto& varr = velocity(lev).array(mfi);