diff --git a/.clang-tidy b/.clang-tidy
index 5ca9d882509..b2ba7d85eed 100644
--- a/.clang-tidy
+++ b/.clang-tidy
@@ -1,6 +1,7 @@
 Checks: >
     '-*,
      bugprone-*,
+       -bugprone-branch-clone,
        -bugprone-easily-swappable-parameters,
        -bugprone-exception-escape,
        -bugprone-implicit-widening-of-multiplication-result,
diff --git a/Src/Base/AMReX_FabArrayBase.H b/Src/Base/AMReX_FabArrayBase.H
index f156f2eadc2..d8395c01711 100644
--- a/Src/Base/AMReX_FabArrayBase.H
+++ b/Src/Base/AMReX_FabArrayBase.H
@@ -468,6 +468,9 @@ public:
         std::unique_ptr<MapOfCopyComTagContainers> m_RcvTags;
     };
 
+    void define_fb_metadata (CommMetaData& cmd, const IntVect& nghost, bool cross,
+                             const Periodicity& period, bool multi_ghost) const;
+
     //
     //! FillBoundary
     struct FB
diff --git a/Src/Base/AMReX_FabArrayBase.cpp b/Src/Base/AMReX_FabArrayBase.cpp
index 71209d83653..274a07cbf91 100644
--- a/Src/Base/AMReX_FabArrayBase.cpp
+++ b/Src/Base/AMReX_FabArrayBase.cpp
@@ -656,26 +656,26 @@ FabArrayBase::FB::FB (const FabArrayBase& fa, const IntVect& nghost,
 }
 
 void
-FabArrayBase::FB::define_fb (const FabArrayBase& fa)
+FabArrayBase::define_fb_metadata (CommMetaData& cmd, const IntVect& nghost,
+                                  bool cross, const Periodicity& period,
+                                  bool multi_ghost) const
 {
-    AMREX_ASSERT(m_multi_ghost ? fa.nGrow() >= 2 : true); // must have >= 2 ghost nodes
-    AMREX_ASSERT(m_multi_ghost ? !m_period.isAnyPeriodic() : true); // this only works for non-periodic
     const int                  MyProc   = ParallelDescriptor::MyProc();
-    const BoxArray&            ba       = fa.boxArray();
-    const DistributionMapping& dm       = fa.DistributionMap();
-    const Vector<int>&         imap     = fa.IndexArray();
+    const BoxArray&            ba       = this->boxArray();
+    const DistributionMapping& dm       = this->DistributionMap();
+    const Vector<int>&         imap     = this->IndexArray();
 
     // For local copy, all workers in the same team will have the identical copy of tags
     // so that they can share work.  But for remote communication, they are all different.
 
     const int nlocal = static_cast<int>(imap.size());
-    const IntVect& ng = m_ngrow;
-    const IntVect ng_ng = m_ngrow - 1;
+    const IntVect& ng = nghost;
+    const IntVect ng_ng =nghost - 1;
     std::vector< std::pair<int,Box> > isects;
 
-    const std::vector<IntVect>& pshifts = m_period.shiftIntVect();
+    const std::vector<IntVect>& pshifts = period.shiftIntVect();
 
-    auto& send_tags = *m_SndTags;
+    auto& send_tags = *cmd.m_SndTags;
 
     for (int i = 0; i < nlocal; ++i)
     {
@@ -697,7 +697,7 @@ FabArrayBase::FB::define_fb (const FabArrayBase& fa)
                     continue;  // local copy will be dealt with later
                 } else if (MyProc == dm[ksnd]) {
                     BoxList bl = amrex::boxDiff(bx, ba[krcv]);
-                    if (m_multi_ghost)
+                    if (multi_ghost)
                     {
                         // In the case where ngrow>1, augment the send/rcv box list
                         // with boxes for overlapping ghost nodes.
@@ -718,7 +718,7 @@ FabArrayBase::FB::define_fb (const FabArrayBase& fa)
         }
     }
 
-    auto& recv_tags = *m_RcvTags;
+    auto& recv_tags = *cmd.m_RcvTags;
 
     bool check_local = false, check_remote = false;
 #if defined(AMREX_USE_GPU)
@@ -735,8 +735,8 @@ FabArrayBase::FB::define_fb (const FabArrayBase& fa)
         check_local = true;
     }
 
-    m_threadsafe_loc = true;
-    m_threadsafe_rcv = true;
+    cmd.m_threadsafe_loc = true;
+    cmd.m_threadsafe_rcv = true;
     for (int i = 0; i < nlocal; ++i)
     {
         BoxList bl_local(ba.ixType());
@@ -759,7 +759,7 @@ FabArrayBase::FB::define_fb (const FabArrayBase& fa)
 
                 BoxList bl = amrex::boxDiff(dst_bx, vbx);
 
-                if (m_multi_ghost)
+                if (multi_ghost)
                 {
                     // In the case where ngrow>1, augment the send/rcv box list
                     // with boxes for overlapping ghost nodes.
@@ -779,7 +779,7 @@ FabArrayBase::FB::define_fb (const FabArrayBase& fa)
                         const BoxList tilelist(blbx, FabArrayBase::comm_tile_size);
                         for (auto const& it_tile : tilelist)
                         {
-                            m_LocTags->emplace_back(it_tile, it_tile+pit, krcv, ksnd);
+                            cmd.m_LocTags->emplace_back(it_tile, it_tile+pit, krcv, ksnd);
                         }
                         if (check_local) {
                             bl_local.push_back(blbx);
@@ -794,20 +794,20 @@ FabArrayBase::FB::define_fb (const FabArrayBase& fa)
             }
         }
 
-        if (m_threadsafe_loc) {
+        if (cmd.m_threadsafe_loc) {
             if ((bl_local.size() > 1)
                 && ! BoxArray(std::move(bl_local)).isDisjoint())
             {
-                m_threadsafe_loc = false;
+                cmd.m_threadsafe_loc = false;
                 check_local = false; // No need to check anymore
             }
         }
 
-        if (m_threadsafe_rcv) {
+        if (cmd.m_threadsafe_rcv) {
             if ((bl_remote.size() > 1)
                 && ! BoxArray(std::move(bl_remote)).isDisjoint())
             {
-                m_threadsafe_rcv = false;
+                cmd.m_threadsafe_rcv = false;
                 check_remote = false; // No need to check anymore
             }
         }
@@ -815,7 +815,7 @@ FabArrayBase::FB::define_fb (const FabArrayBase& fa)
 
     for (int ipass = 0; ipass < 2; ++ipass) // pass 0: send; pass 1: recv
     {
-        CopyComTag::MapOfCopyComTagContainers & Tags = (ipass == 0) ? *m_SndTags : *m_RcvTags;
+        CopyComTag::MapOfCopyComTagContainers & Tags = (ipass == 0) ? *cmd.m_SndTags : *cmd.m_RcvTags;
 
         for (auto& kv : Tags)
         {
@@ -833,7 +833,7 @@ FabArrayBase::FB::define_fb (const FabArrayBase& fa)
                 const IntVect& d2s = tag.sbox.smallEnd() - tag.dbox.smallEnd();
 
                 std::vector<Box> boxes;
-                if (m_cross) {
+                if (cross) {
                     const Box& dstvbx = ba[tag.dstIndex];
                     for (int dir = 0; dir < AMREX_SPACEDIM; dir++)
                     {
@@ -861,7 +861,7 @@ FabArrayBase::FB::define_fb (const FabArrayBase& fa)
                 {
                     for (auto const& cross_box : boxes)
                     {
-                        if (m_cross)
+                        if (cross)
                         {
                             cctv_tags_cross.emplace_back(cross_box, cross_box+d2s,
                                                          tag.dstIndex, tag.srcIndex);
@@ -877,6 +877,15 @@ FabArrayBase::FB::define_fb (const FabArrayBase& fa)
     }
 }
 
+void
+FabArrayBase::FB::define_fb (const FabArrayBase& fa)
+{
+    AMREX_ASSERT(m_multi_ghost ? fa.nGrow() >= 2 : true); // must have >= 2 ghost nodes
+    AMREX_ASSERT(m_multi_ghost ? !m_period.isAnyPeriodic() : true); // this only works for non-periodic
+
+    fa.define_fb_metadata(*this, m_ngrow, m_cross, m_period, m_multi_ghost);
+}
+
 void
 FabArrayBase::FB::define_epo (const FabArrayBase& fa)
 {
diff --git a/Src/Base/AMReX_IntVect.H b/Src/Base/AMReX_IntVect.H
index 13ecfa43fe1..5762c50de5a 100644
--- a/Src/Base/AMReX_IntVect.H
+++ b/Src/Base/AMReX_IntVect.H
@@ -122,6 +122,9 @@ public:
     explicit IntVect (const Array<int,AMREX_SPACEDIM>& a) noexcept
         : vect{AMREX_D_DECL(a[0],a[1],a[2])} {}
 
+    explicit constexpr IntVect (Dim3 const& a) noexcept
+        : vect{AMREX_D_DECL(a.x,a.y,a.z)} {}
+
     // dtor, copy-ctor, copy-op=, move-ctor, and move-op= are compiler generated.
 
     AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
diff --git a/Src/Base/AMReX_NonLocalBC.H b/Src/Base/AMReX_NonLocalBC.H
index 189a0164437..fc7f4cb6583 100644
--- a/Src/Base/AMReX_NonLocalBC.H
+++ b/Src/Base/AMReX_NonLocalBC.H
@@ -99,7 +99,7 @@ struct MultiBlockIndexMapping {
 //! \return Returns IntVect{dtos(Dim3{iv})}
 template <typename DTOS>
 std::enable_if_t<IsCallableR<Dim3, DTOS, Dim3>::value, IntVect>
-Apply(DTOS dtos, const IntVect& iv)
+Apply(DTOS const& dtos, const IntVect& iv)
 {
     Dim3 i = dtos(iv.dim3());
     return IntVect{AMREX_D_DECL(i.x, i.y, i.z)};
@@ -115,7 +115,7 @@ Apply(DTOS dtos, const IntVect& iv)
 //! Apply(dtos, box.smallEnd()) and Apply(dtos, box.bigEnd()).
 template <typename DTOS>
 std::enable_if_t<IsCallableR<Dim3, DTOS, Dim3>::value && !IsCallableR<IndexType, DTOS, IndexType>::value, Box>
-Image(DTOS dtos, const Box& box)
+Image (DTOS const& dtos, const Box& box)
 {
     IntVect mapped_smallEnd = Apply(dtos, box.smallEnd());
     IntVect mapped_bigEnd = Apply(dtos, box.bigEnd());
@@ -138,7 +138,7 @@ Image(DTOS dtos, const Box& box)
 //! Apply(dtos, box.smallEnd()) and Apply(dtos, box.bigEnd()).
 template <typename DTOS>
 std::enable_if_t<IsCallableR<Dim3, DTOS, Dim3>::value && IsCallableR<IndexType, DTOS, IndexType>::value, Box>
-Image(DTOS dtos, const Box& box)
+Image (DTOS const& dtos, const Box& box)
 {
     // "Forget" the index type mapping and invoke Image without changing the index type.
     Box srcbox = Image([&dtos](Dim3 d) { return dtos(d); }, box);
@@ -160,7 +160,7 @@ Image(DTOS dtos, const Box& box)
 //! \return Returns IntVect{dtos.Inverse(Dim3{iv})}
 template <typename DTOS>
 std::enable_if_t<HasInverseMemFn<DTOS>::value, IntVect>
-ApplyInverse(DTOS dtos, const IntVect& iv)
+ApplyInverse(DTOS const& dtos, const IntVect& iv)
 {
     return Apply([&dtos](Dim3 i) { return dtos.Inverse(i); }, iv);
 }
@@ -175,7 +175,7 @@ ApplyInverse(DTOS dtos, const IntVect& iv)
 //! ApplyInverse(box.smallEnd()) and ApplyInverse(box.bigEnd()).
 template <typename DTOS>
 std::enable_if_t<HasInverseMemFn<DTOS>::value && !IsCallableR<IndexType, DTOS, IndexType>::value, Box>
-InverseImage(DTOS dtos, const Box& box)
+InverseImage (DTOS const& dtos, const Box& box)
 {
     return Image([&dtos](Dim3 i) { return dtos.Inverse(i); }, box);
 }
@@ -190,7 +190,7 @@ InverseImage(DTOS dtos, const Box& box)
 //! ApplyInverse(box.smallEnd()) and ApplyInverse(box.bigEnd()).
 template <typename DTOS>
 std::enable_if_t<HasInverseMemFn<DTOS>::value && IsCallableR<IndexType, DTOS, IndexType>::value, Box>
-InverseImage(DTOS dtos, const Box& box)
+InverseImage (DTOS const& dtos, const Box& box)
 {
     return Image([&dtos](auto&& i) { return dtos.Inverse(i); }, box);
 }
@@ -216,7 +216,7 @@ struct MultiBlockCommMetaData : FabArrayBase::CommMetaData {
     template <typename DTOS,
               typename = std::enable_if_t<IsIndexMapping<DTOS>::value>>
     MultiBlockCommMetaData(const FabArrayBase& dst, const Box& dstbox, const FabArrayBase& src,
-                           const IntVect& ngrow, DTOS dtos);
+                           const IntVect& ngrow, DTOS const& dtos);
 
     //! \brief Build global meta data by calling the define() member function.
     //!
@@ -225,7 +225,7 @@ struct MultiBlockCommMetaData : FabArrayBase::CommMetaData {
               typename = std::enable_if_t<IsIndexMapping<DTOS>::value>>
     MultiBlockCommMetaData(const BoxArray& dstba, const DistributionMapping& dstdm,
                            const Box& dstbox, const BoxArray& srcba,
-                           const DistributionMapping& srcdm, const IntVect& ngrow, DTOS dtos);
+                           const DistributionMapping& srcdm, const IntVect& ngrow, DTOS const& dtos);
 
     //! \name Manipulators
 
@@ -257,7 +257,7 @@ struct MultiBlockCommMetaData : FabArrayBase::CommMetaData {
     std::enable_if_t<IsIndexMapping<DTOS>::value>
     define(const BoxArray& dstba, const DistributionMapping& dstdm, const Box& dstbox,
            const BoxArray& srcba, const DistributionMapping& srcdm, const IntVect& ngrow,
-           DTOS dtos);
+           DTOS const& dtos);
 };
 
 ////////////////////////////////////////////////////////////////////////////////////
@@ -491,22 +491,22 @@ struct IsDataPacking :
 template <class FAB, class DTOS = Identity, class Proj = Identity>
 std::enable_if_t<IsBaseFab<FAB>() && IsCallableR<Dim3, DTOS, Dim3>() && IsFabProjection<Proj, FAB>()>
 local_copy_cpu (FabArray<FAB>& dest, const FabArray<FAB>& src, int dcomp, int scomp, int ncomp,
-                FabArrayBase::CopyComTagsContainer const& local_tags, DTOS dtos = DTOS{},
-                Proj proj = Proj{}) noexcept;
+                FabArrayBase::CopyComTagsContainer const& local_tags, DTOS const& dtos = DTOS{},
+                Proj const& proj = Proj{}) noexcept;
 
 template <class FAB, class DTOS = Identity, class Proj = Identity>
 std::enable_if_t<IsBaseFab<FAB>() && IsCallableR<Dim3, DTOS, Dim3>() && IsFabProjection<Proj, FAB>()>
 unpack_recv_buffer_cpu (FabArray<FAB>& mf, int dcomp, int ncomp, Vector<char*> const& recv_data,
                         Vector<std::size_t> const& recv_size,
                         Vector<FabArrayBase::CopyComTagsContainer const*> const& recv_cctc,
-                        DTOS dtos = DTOS{}, Proj proj = Proj{}) noexcept;
+                        DTOS const& dtos = DTOS{}, Proj const& proj = Proj{}) noexcept;
 
 #ifdef AMREX_USE_GPU
 template <class FAB, class DTOS = Identity, class Proj = Identity>
 std::enable_if_t<IsBaseFab<FAB>() && IsCallableR<Dim3, DTOS, Dim3>() && IsFabProjection<Proj, FAB>()>
 local_copy_gpu (FabArray<FAB>& dest, const FabArray<FAB>& src, int dcomp, int scomp, int ncomp,
-                FabArrayBase::CopyComTagsContainer const& local_tags, DTOS dtos = DTOS{},
-                Proj proj = Proj{}) noexcept;
+                FabArrayBase::CopyComTagsContainer const& local_tags, DTOS const& dtos = DTOS{},
+                Proj const& proj = Proj{}) noexcept;
 
 template <class FAB, class DTOS = Identity, class Proj = Identity>
 std::enable_if_t<IsBaseFab<FAB>() && IsCallableR<Dim3, DTOS, Dim3>() && IsFabProjection<Proj, FAB>()>
@@ -514,7 +514,7 @@ unpack_recv_buffer_gpu (FabArray<FAB>& mf, int scomp, int ncomp,
                         Vector<char*> const& recv_data,
                         Vector<std::size_t> const& recv_size,
                         Vector<FabArrayBase::CopyComTagsContainer const*> const& recv_cctc,
-                        DTOS dtos = DTOS{}, Proj proj = Proj{});
+                        DTOS const& dtos = DTOS{}, Proj const& proj = Proj{});
 #endif
 
 ////////////////////////////////////////////////////////////////////////////////////
@@ -877,12 +877,12 @@ ParallelCopy_finish (DoLocalCopy, FabArray<FAB>& dest, const FabArray<FAB>& src,
 template <typename FAB, typename DTOS = Identity, typename Proj = Identity>
 std::enable_if_t<IsBaseFab<FAB>() && IsCallableR<Dim3, DTOS, Dim3>() && IsFabProjection<Proj, FAB>()>
 ParallelCopy (FabArray<FAB>& dest, const FabArray<FAB>& src, const FabArrayBase::CommMetaData& cmd,
-              SrcComp srccomp, DestComp destcomp, NumComps numcomp, DTOS dtos = DTOS{}, Proj proj = Proj{}) {
+              SrcComp srccomp, DestComp destcomp, NumComps numcomp, DTOS const& dtos = DTOS{}, Proj const& proj = Proj{}) {
     PackComponents components{};
     components.dest_component = destcomp.i;
     components.src_component = srccomp.i;
     components.n_components = numcomp.n;
-    ApplyDtosAndProjectionOnReciever<DTOS, Proj> packing{components, std::move(dtos), std::move(proj)};
+    ApplyDtosAndProjectionOnReciever<DTOS, Proj> packing{components, dtos, proj};
     CommHandler handler = ParallelCopy_nowait(dest, src, cmd, packing);
     ParallelCopy_finish(dest, std::move(handler), cmd, packing); // NOLINT
 }
@@ -914,8 +914,8 @@ ParallelCopy (FabArray<FAB>& dest, const FabArray<FAB>& src, const FabArrayBase:
 template <typename FAB, typename DTOS = Identity, typename Proj = Identity>
 std::enable_if_t<IsBaseFab<FAB>() && IsCallableR<Dim3, DTOS, Dim3>() && IsFabProjection<Proj, FAB>()>
 ParallelCopy (FabArray<FAB>& dest, const FabArray<FAB>& src, const FabArrayBase::CommMetaData& cmd,
-              int srccomp, int destcomp, int numcomp, DTOS dtos = DTOS{}, Proj proj = Proj{}) {
-    ParallelCopy(dest, src, cmd, SrcComp(srccomp), DestComp(destcomp), NumComps(numcomp), std::move(dtos), std::move(proj));
+              int srccomp, int destcomp, int numcomp, DTOS const& dtos = DTOS{}, Proj const& proj = Proj{}) {
+    ParallelCopy(dest, src, cmd, SrcComp(srccomp), DestComp(destcomp), NumComps(numcomp), dtos, proj);
 }
 
 //! \brief Call ParallelCopy_nowait followed by ParallelCopy_finish, strong typed version.
@@ -951,9 +951,9 @@ template <typename FAB, typename DTOS = Identity, typename Proj = Identity>
 std::enable_if_t<IsBaseFab<FAB>() && IsIndexMapping<DTOS>() && IsFabProjection<Proj, FAB>(),
 MultiBlockCommMetaData>
 ParallelCopy (FabArray<FAB>& dest, const Box& destbox, const FabArray<FAB>& src, SrcComp srccomp,
-              DestComp destcomp, NumComps numcomp, const IntVect& ngrow, DTOS dtos = DTOS{}, Proj proj = Proj{}) {
+              DestComp destcomp, NumComps numcomp, const IntVect& ngrow, DTOS const& dtos = DTOS{}, Proj const& proj = Proj{}) {
     MultiBlockCommMetaData cmd(dest, destbox, src, ngrow, dtos);
-    ParallelCopy(dest, src, cmd, srccomp, destcomp, numcomp, std::move(dtos), std::move(proj));
+    ParallelCopy(dest, src, cmd, srccomp, destcomp, numcomp, dtos, proj);
     return cmd;
 }
 
@@ -990,8 +990,8 @@ template <typename FAB, typename DTOS = Identity, typename Proj = Identity>
 std::enable_if_t<IsBaseFab<FAB>() && IsIndexMapping<DTOS>() && IsFabProjection<Proj, FAB>(),
 MultiBlockCommMetaData>
 ParallelCopy (FabArray<FAB>& dest, const Box& destbox, const FabArray<FAB>& src, int srccomp,
-              int destcomp, int numcomp, const IntVect& ngrow, DTOS dtos = DTOS{}, Proj proj = Proj{}) {
-    return ParallelCopy(dest, destbox, src, SrcComp(srccomp), DestComp(destcomp), NumComps(numcomp), ngrow, std::move(dtos), std::move(proj));
+              int destcomp, int numcomp, const IntVect& ngrow, DTOS const& dtos = DTOS{}, Proj const& proj = Proj{}) {
+    return ParallelCopy(dest, destbox, src, SrcComp(srccomp), DestComp(destcomp), NumComps(numcomp), ngrow, dtos, proj);
 }
 
 // Rotate90 fills the lo-x and lo-y boundary regions by rotating the data
@@ -1030,11 +1030,171 @@ template <class FAB>
 std::enable_if_t<IsBaseFab<FAB>::value>
 FillPolar (FabArray<FAB>& mf, Box const& domain);
 
+/**
+ * \brief Start communication to fill boundary
+ *
+ * This starts communication to fill ghost cells of a FabArray. This
+ * function is supposed to be used together with FillBoundary_finish and
+ * makeFillBoundaryMetaData as follows.
+\verbatim
+    auto cmd = makeFillBoundaryMetaData(mf, mf.nGrowVect, geom, dtos);
+    // The metadata cmd can be cached and reused on a MultiFab/FabArray with
+    // the same BoxArray and DistributionMapping.
+    auto handler = FillBoundary_nowait(mf, cmd, scomp, ncomp, dtos, proj);
+    // Independent computation can be performed.
+    FillBoundary_finish(std::move(handler), mf, cmd, scomp, ncomp, dtos, proj);
+\endverbatim
+ *
+ * The FillBoundary capability here is more flexible than FabArray's
+ * FillBoundary member functions, which only fill ghost cells inside the
+ * domain and ghost cells at periodic boundaries. The FillBoundary here can
+ * be used to fill non-local domain boundaries (e.g., in spherical and
+ * cylindrical coordinates) given appropriate index and component mappings.
+ *
+ * \tparam FAB  MultiFab/FabArray type
+ * \tparam DTOS Index mapping from destination from source. See
+ *              SphThetaPhiRIndexMapping for an example.
+ * \tparam Proj Component mapping from source to destination. See
+ *              SphThetaPhiRComponentMapping for an example.
+ *
+ * \param mf    FabArray/MultiFab whose ghost cells need to be filled.
+ * \param cmd   communication metadata.
+ * \param scomp starting component.
+ * \param ncomp number of components.
+ * \param dtos  index mapping.
+ * \param proj  component mapping.
+ *
+ * \return      a CommHandler object needed for calling FillBoundary_finish.
+ */
+template <typename FAB, typename DTOS, typename Proj = Identity>
+[[nodiscard]]
+std::enable_if_t<IsBaseFab<FAB>() &&
+                 IsCallableR<Dim3,DTOS,Dim3>() &&
+                 IsFabProjection<Proj,FAB>(),
+                 CommHandler>
+FillBoundary_nowait (FabArray<FAB>& mf, const FabArrayBase::CommMetaData& cmd,
+                     int scomp, int ncomp, DTOS const& dtos,
+                     Proj const& proj = Proj{});
+
+/**
+ * \brief Finish communication started by FillBoundary_nowait
+ *
+ * This finishes the communication to fill ghost cells of a FabArray. This
+ * function is supposed to be used together with FillBoundary_nowait and
+ * makeFillBoundaryMetaData as follows.
+\verbatim
+    auto cmd = makeFillBoundaryMetaData(mf, mf.nGrowVect, geom, dtos);
+    // The metadata cmd can be cached and reused on a MultiFab/FabArray with
+    // the same BoxArray and DistributionMapping.
+    auto handler = FillBoundary_nowait(mf, cmd, scomp, ncomp, dtos, proj);
+    // Independent computation can be performed.
+    FillBoundary_finish(std::move(handler), mf, cmd, scomp, ncomp, dtos, proj);
+\endverbatim
+ *
+ * The FillBoundary capability here is more flexible than FabArray's
+ * FillBoundary member functions, which only fill ghost cells inside the
+ * domain and ghost cells at periodic boundaries. The FillBoundary here can
+ * be used to fill non-local domain boundaries (e.g., in spherical and
+ * cylindrical coordinates) given appropriate index and component mappings.
+ *
+ * \tparam FAB  MultiFab/FabArray type
+ * \tparam DTOS Index mapping from destination from source. See
+ *              SphThetaPhiRIndexMapping for an example.
+ * \tparam Proj Component mapping from source to destination. See
+ *              SphThetaPhiRComponentMapping for an example.
+ *
+ * \param handler CommHandler returned by FillBoundary_nowait.
+ * \param mf      FabArray/MultiFab whose ghost cells need to be filled.
+ * \param cmd     communication metadata.
+ * \param scomp   starting component.
+ * \param ncomp   number of components.
+ * \param dtos    index mapping
+ * \param proj    component mapping
+ */
+template <typename FAB, typename DTOS, typename Proj = Identity>
+std::enable_if_t<IsBaseFab<FAB>() &&
+                 IsCallableR<Dim3,DTOS,Dim3>() &&
+                 IsFabProjection<Proj,FAB>()>
+FillBoundary_finish (CommHandler handler,
+                     FabArray<FAB>& mf, const FabArrayBase::CommMetaData& cmd,
+                     int scomp, int ncomp, DTOS const& dtos,
+                     Proj const& proj = Proj{});
+
+/**
+ * \brief Fill ghost cells for FabArray/MultiFab
+ *
+ * This fills ghost cells of a FabArray. This function is supposed to be
+ * used together with makeFillBoundaryMetaData as follows.
+\verbatim
+    auto cmd = makeFillBoundaryMetaData(mf, mf.nGrowVect, geom, dtos);
+    // The metadata cmd can be cached and reused on a MultiFab/FabArray with
+    // the same BoxArray and DistributionMapping.
+    FillBoundary_finish(mf, cmd, scomp, ncomp, dtos, proj);
+\endverbatim
+ *
+ * The FillBoundary capability here is more flexible than FabArray's
+ * FillBoundary member functions, which only fill ghost cells inside the
+ * domain and ghost cells at periodic boundaries. The FillBoundary here can
+ * be used to fill non-local domain boundaries (e.g., in spherical and
+ * cylindrical coordinates) given appropriate index and component mappings.
+ *
+ * \tparam FAB  MultiFab/FabArray type
+ * \tparam DTOS Index mapping from destination from source. See
+ *              SphThetaPhiRIndexMapping for an example.
+ * \tparam Proj Component mapping from source to destination. See
+ *              SphThetaPhiRComponentMapping for an example.
+ *
+ * \param mf    FabArray/MultiFab whose ghost cells need to be filled.
+ * \param cmd   communication metadata.
+ * \param scomp starting component.
+ * \param ncomp number of components.
+ * \param dtos  index mapping.
+ * \param proj  component mapping.
+ */
+template <typename FAB, typename DTOS, typename Proj = Identity>
+std::enable_if_t<IsBaseFab<FAB>() &&
+                 IsCallableR<Dim3,DTOS,Dim3>() &&
+                 IsFabProjection<Proj,FAB>()>
+FillBoundary (FabArray<FAB>& mf, const FabArrayBase::CommMetaData& cmd,
+              int scomp, int ncomp, DTOS const& dtos, Proj const& proj = Proj{})
+{
+    BL_PROFILE("FillBoundary(cmd)");
+    auto handler = FillBoundary_nowait(mf, cmd, scomp, ncomp, dtos, proj);
+    FillBoundary_finish(std::move(handler), mf, cmd, scomp, ncomp, dtos, proj);
+}
+
+/**
+ * \brief Make metadata for FillBoundary
+ *
+ * \tparam FAB  MultiFab/FabArray type
+ * \tparam DTOS Index mapping from destination from source. See
+ *              SphThetaPhiRIndexMapping for an example.
+ *
+ * \param mf     FabArray/MultiFab whose ghost cells need to be filled.
+ * \param nghost number of ghost cells to be filled.
+ * \param geom   a Geometry object that contains the domain information.
+ * \param dtos   index mapping.
+ *
+ * \return      communication metadata
+ */
+template <typename FAB, typename DTOS>
+[[nodiscard]]
+std::enable_if_t<IsBaseFab<FAB>() && IsCallableR<Dim3,DTOS,Dim3>(),
+                 FabArrayBase::CommMetaData>
+makeFillBoundaryMetaData (FabArray<FAB>& mf, IntVect const& nghost,
+                          Geometry const& geom, DTOS const& dtos);
+
 }
 
 #include <AMReX_NonLocalBCImpl.H>
 
 namespace amrex {
+    using NonLocalBC::FillBoundary;
+    using NonLocalBC::FillBoundary_nowait;
+    using NonLocalBC::FillBoundary_finish;
+    using NonLocalBC::makeFillBoundaryMetaData;
+    using NonLocalBC::SphThetaPhiRIndexMapping;
+    using NonLocalBC::SphThetaPhiRComponentMapping;
     using NonLocalBC::ParallelCopy;
     using NonLocalBC::ParallelCopy_nowait;
     using NonLocalBC::ParallelCopy_finish;
diff --git a/Src/Base/AMReX_NonLocalBC.cpp b/Src/Base/AMReX_NonLocalBC.cpp
index 691d46f90f2..3c148b19362 100644
--- a/Src/Base/AMReX_NonLocalBC.cpp
+++ b/Src/Base/AMReX_NonLocalBC.cpp
@@ -1,5 +1,39 @@
 #include "AMReX_NonLocalBC.H"
 
+namespace amrex::NonLocalBC::detail {
+void split_boxes (BoxList& bl, Box const& domain)
+{
+    BoxList bltmp(bl.ixType());
+    for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
+        for (auto& b : bl) {
+            if (b.smallEnd(idim) < domain.smallEnd(idim) &&
+                b.bigEnd(idim) >= domain.smallEnd(idim))
+            {
+                auto btmp = b;
+                b.setBig(idim,domain.smallEnd(idim)-1);
+                btmp.setSmall(idim,domain.smallEnd(idim));
+                bltmp.push_back(btmp);
+            }
+        }
+        bl.join(bltmp);
+        bltmp.clear();
+
+        for (auto& b : bl) {
+            if (b.smallEnd(idim) <= domain.bigEnd(idim) &&
+                b.bigEnd(idim) > domain.bigEnd(idim))
+            {
+                auto btmp = b;
+                b.setBig(idim,domain.bigEnd(idim));
+                btmp.setSmall(idim,domain.bigEnd(idim)+1);
+                bltmp.push_back(btmp);
+            }
+        }
+        bl.join(bltmp);
+        bltmp.clear();
+    }
+}
+}
+
 namespace amrex::NonLocalBC {
 
 #ifdef AMREX_USE_MPI
@@ -104,6 +138,6 @@ void PostSends(CommData& send, int mpi_tag) {
 template MultiBlockCommMetaData ParallelCopy(FabArray<FArrayBox>& dest, const Box& destbox,
                                              const FabArray<FArrayBox>& src, int destcomp,
                                              int srccomp, int numcomp, const IntVect& ngrow,
-                                             MultiBlockIndexMapping, Identity);
+                                             MultiBlockIndexMapping const&, Identity const&);
 
 }
diff --git a/Src/Base/AMReX_NonLocalBCImpl.H b/Src/Base/AMReX_NonLocalBCImpl.H
index 13eb69c68d6..ab2a0dfb978 100644
--- a/Src/Base/AMReX_NonLocalBCImpl.H
+++ b/Src/Base/AMReX_NonLocalBCImpl.H
@@ -186,7 +186,7 @@ struct PolarFn2 { // for the x-y corners
 template <class FAB, class DTOS, class Proj>
 std::enable_if_t<IsBaseFab<FAB>() && IsCallableR<Dim3, DTOS, Dim3>() && IsFabProjection<Proj, FAB>()>
 local_copy_cpu (FabArray<FAB>& dest, const FabArray<FAB>& src, int dcomp, int scomp, int ncomp,
-                FabArrayBase::CopyComTagsContainer const& local_tags, DTOS dtos, Proj proj) noexcept {
+                FabArrayBase::CopyComTagsContainer const& local_tags, DTOS const& dtos, Proj const& proj) noexcept {
     const auto N_locs = static_cast<int>(local_tags.size());
     if (N_locs == 0) return;
 #ifdef AMREX_USE_OMP
@@ -214,7 +214,7 @@ std::enable_if_t<IsBaseFab<FAB>() && IsCallableR<Dim3, DTOS, Dim3>() && IsFabPro
 unpack_recv_buffer_cpu (FabArray<FAB>& mf, int dcomp, int ncomp, Vector<char*> const& recv_data,
                         Vector<std::size_t> const& recv_size,
                         Vector<FabArrayBase::CopyComTagsContainer const*> const& recv_cctc,
-                        DTOS dtos, Proj proj) noexcept {
+                        DTOS const& dtos, Proj const& proj) noexcept {
     amrex::ignore_unused(recv_size);
 
     const auto N_rcvs = static_cast<int>(recv_cctc.size());
@@ -254,7 +254,7 @@ struct Array4Array4Box {
 template <class FAB, class DTOS, class Proj>
 std::enable_if_t<IsBaseFab<FAB>() && IsCallableR<Dim3, DTOS, Dim3>() && IsFabProjection<Proj, FAB>()>
 local_copy_gpu (FabArray<FAB>& dest, const FabArray<FAB>& src, int dcomp, int scomp, int ncomp,
-                FabArrayBase::CopyComTagsContainer const& local_tags, DTOS dtos, Proj proj) noexcept {
+                FabArrayBase::CopyComTagsContainer const& local_tags, DTOS const& dtos, Proj const& proj) noexcept {
     int N_locs = local_tags.size();
     if (N_locs == 0) return;
 
@@ -279,7 +279,7 @@ unpack_recv_buffer_gpu (FabArray<FAB>& mf, int scomp, int ncomp,
                         Vector<char*> const& recv_data,
                         Vector<std::size_t> const& recv_size,
                         Vector<FabArrayBase::CopyComTagsContainer const*> const& recv_cctc,
-                        DTOS dtos, Proj proj)
+                        DTOS const& dtos, Proj const& proj)
 {
     amrex::ignore_unused(recv_size);
 
@@ -337,14 +337,14 @@ unpack_recv_buffer_gpu (FabArray<FAB>& mf, int scomp, int ncomp,
 
 template <typename DTOS, typename>
 MultiBlockCommMetaData::MultiBlockCommMetaData (const FabArrayBase& dst, const Box& dstbox, const FabArrayBase& src,
-                                                const IntVect& ngrow, DTOS dtos)
+                                                const IntVect& ngrow, DTOS const& dtos)
     : MultiBlockCommMetaData(dst.boxArray(), dst.DistributionMap(), dstbox, src.boxArray(),
                                 src.DistributionMap(), ngrow, dtos) {}
 
 template <typename DTOS, typename>
 MultiBlockCommMetaData::MultiBlockCommMetaData (const BoxArray& dstba, const DistributionMapping& dstdm,
                                                 const Box& dstbox, const BoxArray& srcba,
-                                                const DistributionMapping& srcdm, const IntVect& ngrow, DTOS dtos) {
+                                                const DistributionMapping& srcdm, const IntVect& ngrow, DTOS const& dtos) {
     define(dstba, dstdm, dstbox, srcba, srcdm, ngrow, dtos);
 }
 
@@ -352,7 +352,7 @@ template <typename DTOS>
 std::enable_if_t<IsIndexMapping<DTOS>::value>
 MultiBlockCommMetaData::define (const BoxArray& dstba, const DistributionMapping& dstdm, const Box& dstbox,
                                 const BoxArray& srcba, const DistributionMapping& srcdm, const IntVect& ngrow,
-                                DTOS dtos) {
+                                DTOS const& dtos) {
     m_LocTags = std::make_unique<FabArrayBase::CopyComTagsContainer>();
     m_SndTags = std::make_unique<FabArrayBase::MapOfCopyComTagContainers>();
     m_RcvTags = std::make_unique<FabArrayBase::MapOfCopyComTagContainers>();
@@ -395,13 +395,13 @@ MultiBlockCommMetaData::define (const BoxArray& dstba, const DistributionMapping
     }
 }
 
-template <class FAB, class DTOS>
+template <class FAB, class DTOS, class Proj>
 #ifdef AMREX_USE_MPI
 AMREX_NODISCARD
 #endif
 CommHandler
 Comm_nowait (FabArray<FAB>& mf, int scomp, int ncomp, FabArrayBase::CommMetaData const& cmd,
-             DTOS dtos)
+             DTOS const& dtos, Proj const& proj)
 {
 #ifdef AMREX_USE_MPI
     if (ParallelContext::NProcsSub() == 1)
@@ -410,11 +410,11 @@ Comm_nowait (FabArray<FAB>& mf, int scomp, int ncomp, FabArrayBase::CommMetaData
         if (cmd.m_LocTags->empty()) { return CommHandler{}; }
 #ifdef AMREX_USE_GPU
         if (Gpu::inLaunchRegion()) {
-            local_copy_gpu(mf, mf, scomp, scomp, ncomp, *cmd.m_LocTags, dtos);
+            local_copy_gpu(mf, mf, scomp, scomp, ncomp, *cmd.m_LocTags, dtos, proj);
         } else
 #endif
         {
-            local_copy_cpu(mf, mf, scomp, scomp, ncomp, *cmd.m_LocTags, dtos);
+            local_copy_cpu(mf, mf, scomp, scomp, ncomp, *cmd.m_LocTags, dtos, proj);
         }
         return CommHandler{};
     }
@@ -466,11 +466,11 @@ Comm_nowait (FabArray<FAB>& mf, int scomp, int ncomp, FabArrayBase::CommMetaData
     {
 #ifdef AMREX_USE_GPU
         if (Gpu::inLaunchRegion()) {
-            local_copy_gpu(mf, mf, scomp, scomp, ncomp, *cmd.m_LocTags, dtos);
+            local_copy_gpu(mf, mf, scomp, scomp, ncomp, *cmd.m_LocTags, dtos, proj);
         } else
 #endif
         {
-            local_copy_cpu(mf, mf, scomp, scomp, ncomp, *cmd.m_LocTags, dtos);
+            local_copy_cpu(mf, mf, scomp, scomp, ncomp, *cmd.m_LocTags, dtos, proj);
         }
     }
 
@@ -479,10 +479,10 @@ Comm_nowait (FabArray<FAB>& mf, int scomp, int ncomp, FabArrayBase::CommMetaData
 }
 
 #ifdef AMREX_USE_MPI
-template <class FAB, class DTOS>
+template <class FAB, class DTOS, class Proj>
 void
 Comm_finish (FabArray<FAB>& mf, int scomp, int ncomp, FabArrayBase::CommMetaData const& cmd,
-             CommHandler handler, DTOS dtos)
+             CommHandler handler, DTOS const& dtos, Proj const& proj)
 {
     if (ParallelContext::NProcsSub() == 1) return;
 
@@ -506,12 +506,12 @@ Comm_finish (FabArray<FAB>& mf, int scomp, int ncomp, FabArrayBase::CommMetaData
         if (Gpu::inLaunchRegion())
         {
             unpack_recv_buffer_gpu(mf, scomp, ncomp, handler.recv.data,
-                                   handler.recv.size, handler.recv.cctc, dtos);
+                                   handler.recv.size, handler.recv.cctc, dtos, proj);
         } else
 #endif
         {
             unpack_recv_buffer_cpu(mf, scomp, ncomp, handler.recv.data,
-                                   handler.recv.size, handler.recv.cctc, dtos);
+                                   handler.recv.size, handler.recv.cctc, dtos, proj);
         }
     }
 
@@ -539,7 +539,8 @@ Rotate90 (FabArray<FAB>& mf, int scomp, int ncomp, IntVect const& nghost, Box co
 
     const FabArrayBase::RB90& TheRB90 = mf.getRB90(nghost, domain);
 
-    auto handler = Comm_nowait(mf, scomp, ncomp, TheRB90,Rotate90DstToSrc{});
+    auto handler = Comm_nowait(mf, scomp, ncomp, TheRB90,Rotate90DstToSrc{},
+                               Identity{});
 
     Box corner(-nghost, IntVect{AMREX_D_DECL(-1,-1,domain.bigEnd(2)+nghost[2])});
 #ifdef AMREX_USE_OMP
@@ -557,7 +558,8 @@ Rotate90 (FabArray<FAB>& mf, int scomp, int ncomp, IntVect const& nghost, Box co
     }
 
 #ifdef AMREX_USE_MPI
-    Comm_finish(mf, scomp, ncomp, TheRB90, std::move(handler), Rotate90DstToSrc{});
+    Comm_finish(mf, scomp, ncomp, TheRB90, std::move(handler), Rotate90DstToSrc{},
+                Identity{});
 #else
     amrex::ignore_unused(handler);
 #endif
@@ -587,10 +589,12 @@ Rotate180 (FabArray<FAB>& mf, int scomp, int ncomp, IntVect const& nghost, Box c
 
     const FabArrayBase::RB180& TheRB180 = mf.getRB180(nghost, domain);
 
-    auto handler = Comm_nowait(mf, scomp, ncomp, TheRB180, Rotate180Fn{domain.length(1)});
+    auto handler = Comm_nowait(mf, scomp, ncomp, TheRB180,
+                               Rotate180Fn{domain.length(1)}, Identity{});
 
 #ifdef AMREX_USE_MPI
-    Comm_finish(mf, scomp, ncomp, TheRB180, std::move(handler), Rotate180Fn{domain.length(1)});
+    Comm_finish(mf, scomp, ncomp, TheRB180, std::move(handler),
+                Rotate180Fn{domain.length(1)}, Identity{});
 #else
     amrex::ignore_unused(handler);
 #endif
@@ -621,11 +625,12 @@ FillPolar (FabArray<FAB>& mf, int scomp, int ncomp, IntVect const& nghost, Box c
     const FabArrayBase::PolarB& ThePolarB = mf.getPolarB(nghost, domain);
 
     auto handler = Comm_nowait(mf, scomp, ncomp, ThePolarB,
-                               PolarFn{domain.length(0), domain.length(1)});
+                               PolarFn{domain.length(0), domain.length(1)},
+                               Identity{});
 
 #ifdef AMREX_USE_MPI
     Comm_finish(mf, scomp, ncomp, ThePolarB, std::move(handler),
-                PolarFn{domain.length(0), domain.length(1)});
+                PolarFn{domain.length(0), domain.length(1)}, Identity{});
 #else
     amrex::ignore_unused(handler);
 #endif
@@ -638,10 +643,348 @@ FillPolar (FabArray<FAB>& mf, Box const& domain)
     FillPolar(mf, 0, mf.nComp(), mf.nGrowVect(), domain);
 }
 
+template <typename FAB, typename DTOS, typename Proj>
+std::enable_if_t<IsBaseFab<FAB>() &&
+                 IsCallableR<Dim3,DTOS,Dim3>() &&
+                 IsFabProjection<Proj,FAB>(),
+                 CommHandler>
+FillBoundary_nowait (FabArray<FAB>& mf, const FabArrayBase::CommMetaData& cmd,
+                     int scomp, int ncomp, DTOS const& dtos, Proj const& proj)
+{
+    BL_PROFILE("FillBoundary_nowait(cmd)");
+    AMREX_ASSERT(scomp < mf.nComp() && scomp+ncomp <= mf.nComp());
+    return Comm_nowait(mf, scomp, ncomp, cmd, dtos, proj);
+}
+
+template <typename FAB, typename DTOS, typename Proj>
+std::enable_if_t<IsBaseFab<FAB>() &&
+                 IsCallableR<Dim3,DTOS,Dim3>() &&
+                 IsFabProjection<Proj,FAB>()>
+FillBoundary_finish (CommHandler handler,
+                     FabArray<FAB>& mf, const FabArrayBase::CommMetaData& cmd,
+                     int scomp, int ncomp, DTOS const& dtos, Proj const& proj)
+{
+#ifdef AMREX_USE_MPI
+    BL_PROFILE("FillBoundary_finish(cmd)");
+    Comm_finish(mf, scomp, ncomp, cmd, std::move(handler), dtos, proj);
+#else
+    amrex::ignore_unused(handler,mf,cmd,scomp,ncomp,dtos,proj);
+#endif
+}
+
+template <typename DTOS>
+Vector<std::pair<Box,Box>>
+get_src_dst_boxes (DTOS const& dtos, Box const& dstbox, Box const& domain)
+{
+    Vector<std::pair<Box,Box>> r;
+    IntVect mapped_smallend(dtos(amrex::lbound(dstbox)));
+    IntVect mapped_bigend  (dtos(amrex::ubound(dstbox)));
+    if (!domain.contains(mapped_smallend) || !domain.contains(mapped_bigend)) {
+        return r;
+    }
+
+    auto sign = dtos.sign(amrex::lbound(dstbox));
+    auto perm = dtos.permutation(amrex::lbound(dstbox));
+    auto dtype = dstbox.type();
+    IntVect stype{AMREX_D_DECL(dtype[perm[0]],
+                               dtype[perm[1]],
+                               dtype[perm[2]])};
+    Array<Array<std::pair<int,int>,2>,AMREX_SPACEDIM> ends;
+    Array<Array<std::pair<int,int>,2>,AMREX_SPACEDIM> dst_ends;
+    Array<int,AMREX_SPACEDIM> nboxes;
+    for (int ddim = 0; ddim < AMREX_SPACEDIM; ++ddim) {
+        int sdim = perm[ddim];
+        auto mm = std::minmax(mapped_smallend[sdim],mapped_bigend[sdim]);
+        if (((sign[ddim] > 0) && (mapped_smallend[sdim] <= mapped_bigend[sdim])) ||
+            ((sign[ddim] < 0) && (mapped_bigend[sdim] <= mapped_smallend[sdim])))
+        {
+            nboxes[sdim] = 1;
+            ends[sdim][0] = mm;
+            dst_ends[ddim][0] = std::make_pair(dstbox.smallEnd(ddim),
+                                               dstbox.bigEnd(ddim));
+        } else {
+            nboxes[sdim] = 2;
+            ends[sdim][0].first  = domain.smallEnd(sdim);
+            ends[sdim][0].second = mm.first;
+            ends[sdim][1].first  = mm.second;
+            ends[sdim][1].second = domain.bigEnd(sdim);
+            int n0 = ends[sdim][0].second - ends[sdim][0].first;
+            int n1 = ends[sdim][1].second - ends[sdim][1].first;
+            if (mm.first == mapped_smallend[sdim]) {
+                dst_ends[ddim][0] = std::make_pair(dstbox.smallEnd(ddim),
+                                                   dstbox.smallEnd(ddim)+n0);
+                dst_ends[ddim][1] = std::make_pair(dstbox.bigEnd(ddim)-n1,
+                                                   dstbox.bigEnd(ddim));
+            } else {
+                dst_ends[ddim][0] = std::make_pair(dstbox.bigEnd(ddim)-n0,
+                                                   dstbox.bigEnd(ddim));
+                dst_ends[ddim][1] = std::make_pair(dstbox.smallEnd(ddim),
+                                                   dstbox.smallEnd(ddim)+n1);
+            }
+        }
+    }
+
+    r.reserve(AMREX_D_TERM(nboxes[0],*nboxes[1],*nboxes[2]));
+
+#if (AMREX_SPACEDIM == 3)
+    for (int kbox = 0; kbox < nboxes[2]; ++kbox)
+#endif
+#if (AMREX_SPACEDIM >=2 )
+    for (int jbox = 0; jbox < nboxes[1]; ++jbox)
+#endif
+    for (int ibox = 0; ibox < nboxes[0]; ++ibox)
+    {
+        IntVect siv(AMREX_D_DECL(ibox,jbox,kbox));
+        IntVect div(AMREX_D_DECL(siv[perm[0]],siv[perm[1]],siv[perm[2]]));
+        r.emplace_back(Box(IntVect(AMREX_D_DECL(ends[0][ibox].first,
+                                                ends[1][jbox].first,
+                                                ends[2][kbox].first)),
+                           IntVect(AMREX_D_DECL(ends[0][ibox].second,
+                                                ends[1][jbox].second,
+                                                ends[2][kbox].second)),
+                           stype),
+                       Box(IntVect(AMREX_D_DECL(dst_ends[0][div[0]].first,
+                                                dst_ends[1][div[1]].first,
+                                                dst_ends[2][div[2]].first)),
+                           IntVect(AMREX_D_DECL(dst_ends[0][div[0]].second,
+                                                dst_ends[1][div[1]].second,
+                                                dst_ends[2][div[2]].second)),
+                           dtype));
+    }
+
+    return r; // NOLINT(readability-misleading-indentation,-warnings-as-errors)
+}
+
+template <typename DTOS>
+Box get_dst_subbox (DTOS const& dtos, std::pair<Box,Box> const& sdboxes,
+                    Box const& srcsubbox)
+{
+    Box const& srcbox = sdboxes.first;
+    Box const& dstbox = sdboxes.second;
+    if (srcbox == srcsubbox) {
+        return dstbox;
+    } else {
+        auto sign = dtos.sign(amrex::lbound(dstbox));
+        auto perm = dtos.permutation(amrex::lbound(dstbox));
+        Box dstsubbox = dstbox;
+        for (int ddim = 0; ddim < AMREX_SPACEDIM; ++ddim) {
+            int sdim = perm[ddim];
+            if (sign[ddim] > 0) {
+                dstsubbox.growLo(ddim, srcbox.smallEnd(sdim)-srcsubbox.smallEnd(sdim));
+                dstsubbox.growHi(ddim, srcsubbox.bigEnd(sdim)-srcbox.bigEnd(sdim));
+            } else {
+                dstsubbox.growLo(ddim, srcsubbox.bigEnd(sdim)-srcbox.bigEnd(sdim));
+                dstsubbox.growHi(ddim, srcbox.smallEnd(sdim)-srcsubbox.smallEnd(sdim));
+            }
+        }
+        return dstsubbox;
+    }
+}
+
+namespace detail {
+    void split_boxes (BoxList& bl, Box const& domain);
+}
+
+template <typename FAB, typename DTOS>
+std::enable_if_t<IsBaseFab<FAB>() && IsCallableR<Dim3,DTOS,Dim3>(),
+                 FabArrayBase::CommMetaData>
+makeFillBoundaryMetaData (FabArray<FAB>& mf, IntVect const& nghost,
+                          Geometry const& geom, DTOS const& dtos)
+{
+    FabArrayBase::CommMetaData cmd;
+    cmd.m_LocTags = std::make_unique<FabArrayBase::CopyComTagsContainer>();
+    cmd.m_SndTags = std::make_unique<FabArrayBase::MapOfCopyComTagContainers>();
+    cmd.m_RcvTags = std::make_unique<FabArrayBase::MapOfCopyComTagContainers>();
+
+    // Normal FillBoundary part
+    mf.define_fb_metadata(cmd, nghost, false, geom.periodicity(), false);
+
+    BoxArray const& ba = mf.boxArray();
+    DistributionMapping const& dm = mf.DistributionMap();
+    Box dombox = amrex::convert(geom.Domain(), ba.ixType());
+    Box pdombox = amrex::convert(geom.growPeriodicDomain(nghost), ba.ixType());
+
+    const int myproc = ParallelDescriptor::MyProc();
+    const auto nboxes = static_cast<int>(ba.size());
+    std::vector<std::pair<int,Box> > isects;
+
+    for (int i = 0; i < nboxes; ++i) {
+        Box const& gbx = amrex::grow(ba[i], nghost);
+        BoxList bl = amrex::boxDiff(gbx, pdombox);
+        if (bl.isEmpty()) { continue; }
+
+        detail::split_boxes(bl, dombox);
+
+        const int dst_owner = dm[i];
+        for (auto const& dst_box : bl) {
+            auto const& src_dst_boxes = get_src_dst_boxes(dtos, dst_box, dombox);
+            for (auto const& sd_box_pair : src_dst_boxes) {
+                ba.intersections(sd_box_pair.first, isects);
+                for (auto const& is : isects) {
+                    int const     k = is.first;
+                    Box const src_b = is.second;
+                    int const src_owner = dm[k];
+                    if (dst_owner == myproc || src_owner == myproc) {
+                        Box const& dst_b = get_dst_subbox(dtos, sd_box_pair, src_b);
+                        if (src_owner == dst_owner) {
+                            cmd.m_LocTags->emplace_back(dst_b, src_b, i, k);
+                        } else {
+                            auto& tags = (dst_owner == myproc) ?
+                                (*cmd.m_RcvTags)[src_owner] :
+                                (*cmd.m_SndTags)[dst_owner];
+                            tags.emplace_back(dst_b, src_b, i, k);
+                        }
+                    }
+                }
+            }
+        }
+    }
+
+    return cmd;
+}
+
+struct SphThetaPhiRIndexMapping
+{
+    SphThetaPhiRIndexMapping (Box const& a_domain)
+        : nx(a_domain.length(0)),
+          ny(a_domain.length(1)),
+          nz(a_domain.length(2))
+    {
+        AMREX_ASSERT(a_domain.smallEnd() == 0);
+    }
+
+    [[nodiscard]] AMREX_GPU_HOST_DEVICE
+    Dim3 operator() (Dim3 const& ijk) const noexcept
+    {
+        const int i = ijk.x;
+        const int j = ijk.y;
+        const int k = ijk.z;
+        bool ilo = i < 0;
+        bool ihi = i >= nx;
+        bool imd = i >= 0 && i < nx;
+        bool jlo = j < 0;
+        bool jhi = j >= ny;
+        bool jmd = j >= 0 && j < ny;
+        bool klo = k < 0;
+        bool kmd = k >= 0 && k < nz;
+        // We do not need to do anything at the theta-lo/r-lo edge,
+        // theta-hi/r-lo edge, and r > r-hi.
+        if (ilo && jmd && kmd)
+        {
+            return {-1-i, (j+ny/2)%ny, k};
+        }
+        else if (ihi && jmd && kmd)
+        {
+            return {2*nx-1-i, (j+ny/2)%ny, k};
+        }
+        else if (imd && jlo && kmd)
+        {
+            return {i, j+ny, k};
+        }
+        else if (imd && jhi && kmd)
+        {
+            return {i, j-ny, k};
+        }
+        else if (imd && jmd & klo)
+        {
+            return {nx-1-i, (j+ny/2)%ny, -1-k};
+        }
+        else if (ilo && jlo && kmd)
+        {
+            return {-1-i, (j+ny/2)%ny, k};
+        }
+        else if (ihi && jlo && kmd)
+        {
+            return {2*nx-1-i, (j+ny/2)%ny, k};
+        }
+        else if (ilo && jhi && kmd)
+        {
+            return {-1-i, (j+ny/2)%ny, k};
+        }
+        else if (ihi && jhi && kmd)
+        {
+            return {2*nx-1-i, (j+ny/2)%ny, k};
+        }
+        else if (imd && jlo && klo)
+        {
+            return {nx-1-i, (j+ny/2)%ny, -1-k};
+        }
+        else if (imd && jhi && klo)
+        {
+            return {nx-1-i, (j+ny/2)%ny, -1-k};
+        }
+        else
+        {
+            return ijk;
+        }
+    }
+
+    [[nodiscard]] IntVect sign (Dim3 const& ijk) const noexcept
+    {
+        if (ijk.z < 0) {
+            return IntVect{AMREX_D_DECL(-1, 1,-1)};
+        } else if (ijk.z >=0 && ijk.z < nz &&
+                   (ijk.x < 0 || ijk.x >= nx)) {
+            return IntVect{AMREX_D_DECL(-1, 1, 1)};
+        } else {
+            return IntVect{AMREX_D_DECL( 1, 1, 1)};
+        }
+    }
+
+    [[nodiscard]] IntVect permutation (Dim3 const& /*ijk*/) const noexcept // NOLINT(readability-convert-member-functions-to-static)
+    {
+        return IntVect(AMREX_D_DECL(0,1,2));
+    }
+
+private:
+    int nx, ny, nz;
+};
+
+struct SphThetaPhiRComponentMapping
+{
+    SphThetaPhiRComponentMapping (Box const& a_domain, int a_start_index)
+        : nx(a_domain.length(0)),
+          ny(a_domain.length(1)),
+          nz(a_domain.length(2)),
+          scomp(a_start_index) {}
+
+    template <typename T>
+    [[nodiscard]] AMREX_GPU_HOST_DEVICE
+    T operator()(Array4<const T> const& a, Dim3 const& ijk, int n) const noexcept
+    {
+        const int i = ijk.x;
+        const int j = ijk.y;
+        const int k = ijk.z;
+        auto r = a(i,j,k,n);
+        if (n == scomp) {
+            if ((i >= 0 && i <  nx) &&
+                (j <  0 || j >= ny) &&
+                (k >= 0 && k <  nz)) {
+                return r;
+            } else {
+                // We do not need to worry about the theta-lo/r-lo edge,
+                // theta-hi/r-lo edge, and r > r-hi.
+                return -r;
+            }
+        } else if (n == scomp+2) {
+            if (k < 0) {
+                return -r;
+            } else {
+                return r;
+            }
+        } else {
+            return r;
+        }
+    }
+private:
+    int nx, ny, nz;
+    int scomp;
+};
+
 extern template MultiBlockCommMetaData ParallelCopy(FabArray<FArrayBox>& dest, const Box& destbox,
                                                     const FabArray<FArrayBox>& src, int destcomp,
                                                     int srccomp, int numcomp, const IntVect& ngrow,
-                                                    MultiBlockIndexMapping, Identity);
+                                                    MultiBlockIndexMapping const&, Identity const&);
 }
 
 #endif
diff --git a/Src/EB/AMReX_EB_STL_utils.cpp b/Src/EB/AMReX_EB_STL_utils.cpp
index 29ec723b2b1..ba729c12ce3 100644
--- a/Src/EB/AMReX_EB_STL_utils.cpp
+++ b/Src/EB/AMReX_EB_STL_utils.cpp
@@ -51,7 +51,7 @@ namespace {
                                               Real dlevset)
     {
         if ((dlevset > 0._rt && norm.x > 0._rt) || (dlevset < 0._rt && norm.x < 0._rt))
-        { // This triangle has the wrong direction // NOLINT(bugprone-branch-clone)
+        { // This triangle has the wrong direction
             return std::make_pair(false,0.0_rt);
         }
         else if (x1 > amrex::max(v1.x,v2.x,v3.x) ||
diff --git a/Src/LinearSolvers/MLMG/AMReX_MLNodeLaplacian_misc.cpp b/Src/LinearSolvers/MLMG/AMReX_MLNodeLaplacian_misc.cpp
index c8b530927e8..907b0483400 100644
--- a/Src/LinearSolvers/MLMG/AMReX_MLNodeLaplacian_misc.cpp
+++ b/Src/LinearSolvers/MLMG/AMReX_MLNodeLaplacian_misc.cpp
@@ -584,7 +584,7 @@ MLNodeLaplacian::Fsmooth (int amrlev, int mglev, MultiFab& sol, const MultiFab&
                 }
             }
             else if ( (m_use_harmonic_average && mglev > 0) || m_use_mapped )
-            { // NOLINT(bugprone-branch-clone)
+            {
 #ifdef AMREX_USE_OMP
 #pragma omp parallel
 #endif
diff --git a/Tests/Amr/Advection_AmrCore/Source/AdvancePhiAllLevels.cpp b/Tests/Amr/Advection_AmrCore/Source/AdvancePhiAllLevels.cpp
index 8f79d63f6ab..6d9fc8125d6 100644
--- a/Tests/Amr/Advection_AmrCore/Source/AdvancePhiAllLevels.cpp
+++ b/Tests/Amr/Advection_AmrCore/Source/AdvancePhiAllLevels.cpp
@@ -100,7 +100,7 @@ AmrCoreAdv::AdvancePhiAllLevels (Real time, Real dt_lev, int /*iteration*/)
                 // y -------------------------
                 Array4<Real> phiy = tmpfab.array(itmp);
                 Array4<Real const> phiy_c = phiy;
-                itmp += 1;
+                itmp += 1; // NOLINT(clang-analyzer-deadcode.DeadStores)
 
                 amrex::launch(amrex::grow(gbx,Direction::y,1),
                 [=] AMREX_GPU_DEVICE (const Box& tbx)
diff --git a/Tests/MultiBlock/Advection/main.cpp b/Tests/MultiBlock/Advection/main.cpp
index ce0b148c3b8..74418ef9efc 100644
--- a/Tests/MultiBlock/Advection/main.cpp
+++ b/Tests/MultiBlock/Advection/main.cpp
@@ -74,7 +74,7 @@ class AdvectionAmrCore : public AmrCore {
         // Perform first order accurate upwinding with velocity 1 in the stored direction.
         const double dx = Geom(0).CellSize(0);
         const double a_dt_over_dx = dt / dx * (velocity == dir);
-        if (dir == Direction::x) { // NOLINT(bugprone-branch-clone)
+        if (dir == Direction::x) {
 #ifdef AMREX_USE_OMP
 #pragma omp parallel if (Gpu::notInLaunchRegion())
 #endif
diff --git a/Tools/Plotfile/fcompare.cpp b/Tools/Plotfile/fcompare.cpp
index 88edc1981c1..7dc5ed361ad 100644
--- a/Tools/Plotfile/fcompare.cpp
+++ b/Tools/Plotfile/fcompare.cpp
@@ -369,7 +369,7 @@ int main_main()
         if (abort_if_not_all_found) return EXIT_FAILURE;
     }
 
-    if (any_nans) { // NOLINT(bugprone-branch-clone)
+    if (any_nans) {
         return EXIT_FAILURE;
     } else if (global_error == 0.0) {
         amrex::Print() << " PLOTFILE AGREE" << std::endl;