4 #include <AMReX_Config.H>
56 template <typename T, std::enable_if_t<!IsBaseFab<T>::value,
int> = 0>
90 template <
typename T,
typename... Ts>
92 tags.emplace_back(std::forward<T>(t));
93 return SetTag(std::forward<Ts>(ts)...);
149 template <
typename T>
159 explicit operator bool() const noexcept {
170 template <
class FAB>
class FabArray;
172 template <
class DFAB,
class SFAB,
173 std::enable_if_t<std::conjunction_v<
176 typename DFAB::value_type>>,
int> BAR = 0>
180 Copy(dst,src,srccomp,dstcomp,numcomp,
IntVect(nghost));
183 template <
class DFAB,
class SFAB,
184 std::enable_if_t<std::conjunction_v<
185 IsBaseFab<DFAB>, IsBaseFab<SFAB>,
187 typename DFAB::value_type>>,
int> BAR = 0>
193 using DT =
typename DFAB::value_type;
198 if constexpr (std::is_same_v<typename SFAB::value_type, typename DFAB::value_type>) {
207 auto const& dstarr = dst.
arrays();
211 dstarr[box_no](i,j,k,dstcomp+n) = DT(srcarr[box_no](i,j,k,srccomp+n));
218 #pragma omp parallel if (Gpu::notInLaunchRegion())
222 const Box& bx = mfi.growntilebox(nghost);
226 auto const& dstFab = dst.
array(mfi);
229 dstFab(i,j,k,dstcomp+n) = DT(srcFab(i,j,k,srccomp+n));
237 class bar = std::enable_if_t<IsBaseFab<FAB>::value> >
241 Add(dst,src,srccomp,dstcomp,numcomp,
IntVect(nghost));
245 class bar = std::enable_if_t<IsBaseFab<FAB>::value> >
253 auto const& dstfa = dst.
arrays();
258 dstfa[box_no](i,j,k,n+dstcomp) += srcfa[box_no](i,j,k,n+srccomp);
267 #pragma omp parallel if (Gpu::notInLaunchRegion())
271 const Box& bx = mfi.growntilebox(nghost);
274 auto const srcFab = src.
array(mfi);
275 auto dstFab = dst.
array(mfi);
278 dstFab(i,j,k,n+dstcomp) += srcFab(i,j,k,n+srccomp);
380 #ifdef AMREX_STRICT_MODE
392 #ifdef AMREX_STRICT_MODE
421 #ifdef AMREX_STRICT_MODE
433 #ifdef AMREX_STRICT_MODE
451 return (
f !=
nullptr);
477 return f->isAllRegular();
538 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
541 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
544 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
547 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
550 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
553 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
556 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
559 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
562 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
565 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
568 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
571 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
574 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
577 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
580 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
583 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
586 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
593 template <class F=
FAB, std::enable_if_t<std::is_move_constructible_v<F>,
int> = 0>
600 template <class F=
FAB, std::enable_if_t<std::is_move_constructible_v<F>,
int> = 0>
628 template <typename SFAB, typename DFAB =
FAB,
629 std::enable_if_t<std::conjunction_v<
648 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
653 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
657 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
665 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
671 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
683 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
690 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
700 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
703 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
711 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
714 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
717 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
718 void abs (
int comp,
int ncomp,
int nghost = 0);
720 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
723 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
726 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
729 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
732 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
735 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
738 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
742 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
746 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
750 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
754 template <class F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
764 template <typename F=
FAB, std::enable_if_t<
IsBaseFab<F>::value,
int> = 0>
766 sum (
int comp,
IntVect const& nghost,
bool local = false) const;
782 [[deprecated(
"Use FabArray::ParallelCopy() instead.")]]
816 {
ParallelCopy(src,src_comp,dest_comp,num_comp,0,0,period,op); }
818 [[deprecated(
"Use FabArray::ParallelCopy() instead.")]]
825 {
ParallelCopy(src,src_comp,dest_comp,num_comp, period, op); }
907 IntVect(dst_nghost),period,op); }
918 bool to_ghost_cells_only =
false);
940 [[deprecated(
"Use FabArray::ParallelCopy() instead.")]]
951 [[deprecated(
"Use FabArray::ParallelCopy() instead.")]]
960 {
ParallelCopy(src,src_comp,dest_comp,num_comp,src_nghost,dst_nghost,period,op); }
983 void copyTo (
FAB& dest,
int scomp,
int dcomp,
int ncomp,
int nghost = 0)
const;
1002 template <
typename BUF=value_type>
1005 template <
typename BUF=value_type>
1008 template <
typename BUF=value_type>
1012 template <
typename BUF=value_type>
1015 template <
typename BUF=value_type>
1018 template <
typename BUF=value_type>
1021 template <
typename BUF=value_type>
1024 template <
typename BUF=value_type>
1027 template <
typename BUF=value_type>
1030 template <
typename BUF=value_type>
1033 template <
typename BUF=value_type>
1036 template <
typename BUF=value_type>
1040 class F=
FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
1149 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
1157 class F=
FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
1160 bool enforce_periodicity_only =
false,
1161 bool override_sync =
false);
1165 int scomp,
int dcomp,
int ncomp,
CpOp op);
1167 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
1170 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
1173 #ifdef AMREX_USE_GPU
1177 int scomp,
int dcomp,
int ncomp,
CpOp op);
1182 #if defined(__CUDACC__)
1184 void FB_local_copy_cuda_graph_1 (
const FB& TheFB,
int scomp,
int ncomp);
1185 void FB_local_copy_cuda_graph_n (
const FB& TheFB,
int scomp,
int ncomp);
1190 #ifdef AMREX_USE_MPI
1192 #ifdef AMREX_USE_GPU
1193 #if defined(__CUDACC__)
1195 void FB_pack_send_buffer_cuda_graph (
const FB& TheFB,
int scomp,
int ncomp,
1200 void FB_unpack_recv_buffer_cuda_graph (
const FB& TheFB,
int dcomp,
int ncomp,
1204 bool is_thread_safe);
1208 template <
typename BUF = value_type>
1214 template <
typename BUF = value_type>
1219 CpOp op,
bool is_thread_safe);
1223 template <
typename BUF = value_type>
1229 template <
typename BUF = value_type>
1234 CpOp op,
bool is_thread_safe);
1247 template <typename F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
1248 typename F::value_type
1250 [[maybe_unused]]
bool ignore_covered =
false)
const;
1261 template <typename IFAB, typename F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
1262 typename F::value_type
1264 bool local =
false)
const;
1280 #ifdef AMREX_USE_GPU
1295 #if defined(BL_USE_MPI3)
1296 if (win != MPI_WIN_NULL) { MPI_Win_free(&win); }
1306 #if defined(BL_USE_MPI3)
1311 #if defined(BL_USE_MPI3)
1312 rhs.win = MPI_WIN_NULL;
1321 #if defined(BL_USE_MPI3)
1323 rhs.win = MPI_WIN_NULL;
1333 #if defined(BL_USE_MPI3)
1334 MPI_Win win = MPI_WIN_NULL;
1346 bool alloc_single_chunk);
1350 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
1360 template <
typename BUF=value_type>
1362 char*& the_recv_data,
1370 template <
typename BUF=value_type>
1380 template <
typename BUF=value_type>
1382 char*& the_send_data,
1390 template <
typename BUF=value_type>
1407 std::unique_ptr<FBData<FAB>>
fbd;
1408 std::unique_ptr<PCData<FAB>>
pcd;
1426 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
1428 int xcomp,
int ycomp,
int ncomp,
IntVect const& nghost);
1441 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
1443 int xcomp,
int ycomp,
int ncomp,
IntVect const& nghost);
1459 template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,
int> = 0>
1463 int dstcomp,
int numcomp,
const IntVect& nghost);
1469 template <
class FAB>
1473 int li = localindex(K);
1474 if (li >= 0 && li <
static_cast<int>(m_fabs_v.size()) && m_fabs_v[li] != 0) {
1482 template <
class FAB>
1486 int li = mfi.LocalIndex();
1487 if (li <
static_cast<int>(m_fabs_v.size()) && m_fabs_v[li] !=
nullptr) {
1495 template <
class FAB>
1501 int li = mfi.LocalIndex();
1502 return m_fabs_v[li];
1505 template <
class FAB>
1511 int li = mfi.LocalIndex();
1512 return m_fabs_v[li];
1515 template <
class FAB>
1519 int li = localindex(K);
1521 return m_fabs_v[li];
1524 template <
class FAB>
1528 int li = localindex(K);
1530 return m_fabs_v[li];
1533 template <
class FAB>
1534 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1538 #ifdef AMREX_USE_CUDA
1539 this->fabPtr(mfi)->prefetchToHost();
1545 template <
class FAB>
1546 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1550 #ifdef AMREX_USE_CUDA
1551 this->fabPtr(mfi)->prefetchToDevice();
1557 template <
class FAB>
1558 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1562 return fabPtr(mfi)->const_array();
1565 template <
class FAB>
1566 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1570 return fabPtr(mfi)->array();
1573 template <
class FAB>
1574 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1578 return fabPtr(K)->const_array();
1581 template <
class FAB>
1582 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1586 return fabPtr(K)->array();
1589 template <
class FAB>
1590 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1594 return fabPtr(mfi)->const_array();
1597 template <
class FAB>
1598 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1602 return fabPtr(K)->const_array();
1605 template <
class FAB>
1606 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1610 return fabPtr(mfi)->const_array(start_comp);
1613 template <
class FAB>
1614 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1618 return fabPtr(mfi)->array(start_comp);
1621 template <
class FAB>
1622 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1626 return fabPtr(K)->const_array(start_comp);
1629 template <
class FAB>
1630 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1634 return fabPtr(K)->array(start_comp);
1637 template <
class FAB>
1638 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1642 return fabPtr(mfi)->const_array(start_comp);
1645 template <
class FAB>
1646 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1650 return fabPtr(K)->const_array(start_comp);
1653 template <
class FAB>
1654 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1662 template <
class FAB>
1663 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1668 return m_const_arrays;
1671 template <
class FAB>
1672 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1677 return m_const_arrays;
1680 template <
class FAB>
1681 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1687 static_assert(
sizeof(A) ==
sizeof(AC),
"sizeof(Array4<T>) != sizeof(Array4<T const>)");
1688 if (!m_hp_arrays && local_size() > 0) {
1689 const int n = local_size();
1690 #ifdef AMREX_USE_GPU
1696 for (
int li = 0; li < n; ++li) {
1698 new ((A*)m_hp_arrays+li) A(m_fabs_v[li]->array());
1699 new ((AC*)m_hp_arrays+li+n) AC(m_fabs_v[li]->const_array());
1701 new ((A*)m_hp_arrays+li) A{};
1702 new ((AC*)m_hp_arrays+li+n) AC{};
1705 m_arrays.
hp = (A*)m_hp_arrays;
1706 m_const_arrays.
hp = (AC*)m_hp_arrays + n;
1707 #ifdef AMREX_USE_GPU
1708 m_arrays.
dp = (A*)m_dp_arrays;
1709 m_const_arrays.
dp = (AC*)m_dp_arrays + n;
1715 template <
class FAB>
1719 #ifdef AMREX_USE_GPU
1722 m_dp_arrays =
nullptr;
1726 m_hp_arrays =
nullptr;
1727 m_arrays.
hp =
nullptr;
1728 m_const_arrays.
hp =
nullptr;
1731 template <
class FAB>
1736 const int li = localindex(K);
1737 if (li >= 0 && li <
static_cast<int>(m_fabs_v.size()) && m_fabs_v[li] !=
nullptr) {
1741 for (
auto const& t : m_tags) {
1742 updateMemUsage(t, -nbytes,
nullptr);
1745 return std::exchange(m_fabs_v[li],
nullptr);
1751 template <
class FAB>
1757 if (li >= 0 && li <
static_cast<int>(m_fabs_v.size()) && m_fabs_v[li] !=
nullptr) {
1761 for (
auto const& t : m_tags) {
1762 updateMemUsage(t, -nbytes,
nullptr);
1765 return std::exchange(m_fabs_v[li],
nullptr);
1771 template <
class FAB>
1775 if (define_function_called)
1777 define_function_called =
false;
1782 for (
auto *
x : m_fabs_v) {
1785 m_factory->destroy(
x);
1791 m_dallocator.
m_arena =
nullptr;
1795 for (
auto const& t : m_tags) {
1796 updateMemUsage(t, -nbytes,
nullptr);
1800 if (m_single_chunk_arena) {
1801 m_single_chunk_arena.reset();
1803 m_single_chunk_size = 0;
1810 template <
class FAB>
1811 template <
typename SFAB,
typename DFAB,
1812 std::enable_if_t<std::conjunction_v<
1815 typename DFAB::value_type>>,
int>>
1820 amrex::Copy(*
this, src, scomp, dcomp, ncomp, nghost);
1823 template <
class FAB>
1824 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1829 amrex::Add(*
this, src, scomp, dcomp, ncomp, nghost);
1832 template <
class FAB>
1833 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1840 template <
class FAB>
1841 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1845 setVal(val,0,n_comp,nghost);
1848 template <
class FAB>
1849 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1856 template <
class FAB>
1857 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
1861 setVal(val,region,0,n_comp,nghost);
1864 template <
class FAB>
1868 m_FA_stats.recordBuild();
1871 template <
class FAB>
1876 m_FA_stats.recordBuild();
1879 template <
class FAB>
1889 template <
class FAB>
1896 : m_factory(factory.clone()),
1903 template <
class FAB>
1905 : m_factory(rhs.Factory().clone()),
1915 auto const& rhsfab = *(rhs.
m_fabs_v[i]);
1925 template <
class FAB>
1928 , m_factory (std::move(rhs.m_factory))
1929 , m_dallocator (std::move(rhs.m_dallocator))
1930 , m_single_chunk_arena(std::move(rhs.m_single_chunk_arena))
1931 , m_single_chunk_size(std::exchange(rhs.m_single_chunk_size,0))
1932 , define_function_called(rhs.define_function_called)
1933 , m_fabs_v (std::move(rhs.m_fabs_v))
1934 #ifdef AMREX_USE_GPU
1935 , m_dp_arrays (std::exchange(rhs.m_dp_arrays,
nullptr))
1937 , m_hp_arrays (std::exchange(rhs.m_hp_arrays,
nullptr))
1938 , m_arrays (rhs.m_arrays)
1939 , m_const_arrays(rhs.m_const_arrays)
1940 , m_tags (std::move(rhs.m_tags))
1941 , shmem (std::move(rhs.shmem))
1944 m_FA_stats.recordBuild();
1945 rhs.define_function_called =
false;
1946 rhs.m_fabs_v.clear();
1950 template <
class FAB>
1959 m_factory = std::move(rhs.m_factory);
1960 m_dallocator = std::move(rhs.m_dallocator);
1961 m_single_chunk_arena = std::move(rhs.m_single_chunk_arena);
1962 std::swap(m_single_chunk_size, rhs.m_single_chunk_size);
1963 define_function_called = rhs.define_function_called;
1965 #ifdef AMREX_USE_GPU
1966 std::swap(m_dp_arrays, rhs.m_dp_arrays);
1968 std::swap(m_hp_arrays, rhs.m_hp_arrays);
1969 m_arrays = rhs.m_arrays;
1970 m_const_arrays = rhs.m_const_arrays;
1972 shmem = std::move(rhs.shmem);
1974 rhs.define_function_called =
false;
1975 rhs.m_fabs_v.clear();
1982 template <
class FAB>
1985 m_FA_stats.recordDelete();
1989 template <
class FAB>
1993 if (!define_function_called) {
return false; }
2001 if (
get(fai).box() != fabbox(fai.index()))
2017 template <
class FAB>
2021 return define_function_called;
2024 template <
class FAB>
2033 define(bxs,dm,nvar,
IntVect(ngrow),info,a_factory);
2036 template <
class FAB>
2045 std::unique_ptr<FabFactory<FAB> > factory(a_factory.
clone());
2047 auto *default_arena = m_dallocator.
m_arena;
2050 m_factory = std::move(factory);
2053 define_function_called =
true;
2069 template <
class FAB>
2074 if (shmem.alloc) { alloc_single_chunk =
false; }
2075 if constexpr (!IsBaseFab_v<FAB>) { alloc_single_chunk =
false; }
2077 const int n = indexArray.size();
2079 shmem.alloc = (nworkers > 1);
2081 bool alloc = !shmem.alloc;
2086 if (alloc_single_chunk) {
2087 m_single_chunk_size = 0L;
2088 for (
int i = 0; i < n; ++i) {
2089 int K = indexArray[i];
2090 const Box& tmpbox = fabbox(K);
2091 m_single_chunk_size += factory.
nBytes(tmpbox, n_comp, K);
2094 m_single_chunk_arena = std::make_unique<detail::SingleChunkArena>(ar, m_single_chunk_size);
2095 fab_info.
SetArena(m_single_chunk_arena.get());
2098 m_fabs_v.reserve(n);
2101 for (
int i = 0; i < n; ++i)
2103 int K = indexArray[i];
2104 const Box& tmpbox = fabbox(K);
2105 m_fabs_v.push_back(factory.
create(tmpbox, n_comp, fab_info, K));
2110 m_tags.emplace_back(
"All");
2111 for (
auto const& t : m_region_tag) {
2112 m_tags.push_back(t);
2114 for (
auto const& t : tags) {
2115 m_tags.push_back(t);
2117 for (
auto const& t: m_tags) {
2118 updateMemUsage(t, nbytes, ar);
2130 for (
int i = 0; i < n; ++i) {
2131 int K = indexArray[i];
2132 int owner = distributionMap[K] - teamlead;
2133 Long s = m_fabs_v[i]->size();
2135 shmem.n_values += s;
2136 shmem.n_points += m_fabs_v[i]->numPts();
2138 if (nextoffset[owner] < 0) {
2140 nextoffset[owner] = s;
2142 offset[i] = nextoffset[owner];
2143 nextoffset[owner] += s;
2147 size_t bytes = shmem.n_values*
sizeof(
value_type);
2152 #if defined (BL_USE_MPI3)
2154 static MPI_Info info = MPI_INFO_NULL;
2155 if (info == MPI_INFO_NULL) {
2156 MPI_Info_create(&info);
2157 MPI_Info_set(info,
"alloc_shared_noncontig",
"true");
2162 BL_MPI_REQUIRE( MPI_Win_allocate_shared(bytes,
sizeof(
value_type),
2163 info, team_comm, &mfp, &shmem.win) );
2165 for (
int w = 0; w < nworkers; ++w) {
2169 BL_MPI_REQUIRE( MPI_Win_shared_query(shmem.win, w, &sz, &disp, &dptr) );
2171 dps.push_back(dptr);
2176 amrex::Abort(
"BaseFab::define: to allocate shared memory, USE_MPI3 must be true");
2180 for (
int i = 0; i < n; ++i) {
2181 int K = indexArray[i];
2182 int owner = distributionMap[K] - teamlead;
2184 m_fabs_v[i]->setPtr(p, m_fabs_v[i]->
size());
2187 for (Long i = 0; i < shmem.n_values; i++, mfp++) {
2196 template <
class FAB>
2208 template <
class FAB>
2213 n_comp = elem->nComp();
2216 setFab_assert(boxno, *elem);
2218 if (m_fabs_v.empty()) {
2219 m_fabs_v.resize(indexArray.size(),
nullptr);
2222 const int li = localindex(boxno);
2224 m_factory->destroy(m_fabs_v[li]);
2226 m_fabs_v[li] = elem.release();
2229 template <
class FAB>
2230 template <
class F, std::enable_if_t<std::is_move_constructible_v<F>,
int> >
2235 n_comp = elem.nComp();
2238 setFab_assert(boxno, elem);
2240 if (m_fabs_v.empty()) {
2241 m_fabs_v.resize(indexArray.size(),
nullptr);
2244 const int li = localindex(boxno);
2246 m_factory->destroy(m_fabs_v[li]);
2248 m_fabs_v[li] =
new FAB(std::move(elem));
2251 template <
class FAB>
2256 n_comp = elem->nComp();
2259 setFab_assert(mfi.
index(), *elem);
2261 if (m_fabs_v.empty()) {
2262 m_fabs_v.resize(indexArray.size(),
nullptr);
2267 m_factory->destroy(m_fabs_v[li]);
2269 m_fabs_v[li] = elem.release();
2272 template <
class FAB>
2273 template <
class F, std::enable_if_t<std::is_move_constructible_v<F>,
int> >
2278 n_comp = elem.nComp();
2281 setFab_assert(mfi.
index(), elem);
2283 if (m_fabs_v.empty()) {
2284 m_fabs_v.resize(indexArray.size(),
nullptr);
2289 m_factory->destroy(m_fabs_v[li]);
2291 m_fabs_v[li] =
new FAB(std::move(elem));
2294 template <
class FAB>
2295 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
2302 template <
class FAB>
2303 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>Z>
2309 if (n_grow.max() > 0)
2311 #ifdef AMREX_USE_GPU
2313 bool use_mfparfor =
true;
2314 const int nboxes = local_size();
2316 if (boxarray[indexArray[0]].numPts() > Long(65*65*65)) {
2317 use_mfparfor =
false;
2320 for (
int i = 0; i < nboxes; ++i) {
2321 const Long npts = boxarray[indexArray[i]].numPts();
2322 if (npts >= Long(64*64*64)) {
2323 use_mfparfor =
false;
2325 }
else if (npts <= Long(17*17*17)) {
2330 const IntVect nghost = n_grow;
2332 auto const& ma = this->arrays();
2336 auto const& a = ma[box_no];
2340 for (
int n = 0; n < ncomp; ++n) {
2341 a(i,j,k,strt_comp+n) = val;
2350 Box const& vbx = mfi.validbox();
2351 auto const& a = this->array(mfi);
2354 #if (AMREX_SPACEDIM == 3)
2355 if (nghost[2] > 0) {
2357 b.setRange(2, vbx.
smallEnd(2)-nghost[2], nghost[2]);
2358 b.grow(
IntVect(nghost[0],nghost[1],0));
2359 tags.emplace_back(Tag{a,
b});
2360 b.shift(2, vbx.
length(2)+nghost[2]);
2361 tags.emplace_back(Tag{a,
b});
2364 #if (AMREX_SPACEDIM >= 2)
2365 if (nghost[1] > 0) {
2367 b.setRange(1, vbx.
smallEnd(1)-nghost[1], nghost[1]);
2368 b.grow(0, nghost[0]);
2369 tags.emplace_back(Tag{a,
b});
2370 b.shift(1, vbx.
length(1)+nghost[1]);
2371 tags.emplace_back(Tag{a,
b});
2374 if (nghost[0] > 0) {
2376 b.setRange(0, vbx.
smallEnd(0)-nghost[0], nghost[0]);
2377 tags.emplace_back(Tag{a,
b});
2378 b.shift(0, vbx.
length(0)+nghost[0]);
2379 tags.emplace_back(Tag{a,
b});
2386 tag.dfab(i,j,k,strt_comp+n) = val;
2392 #ifdef AMREX_USE_OMP
2393 #pragma omp parallel
2397 get(fai).template setComplement<RunOn::Host>(val, fai.validbox(), strt_comp, ncomp);
2403 template <
class FAB>
2404 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
2408 setDomainBndry(val, 0, n_comp, geom);
2411 template <
class FAB>
2412 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
2422 for (
int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
2424 int n = domain_box.
length(idim);
2425 domain_box.
grow(idim, n);
2429 #ifdef AMREX_USE_OMP
2430 #pragma omp parallel if (Gpu::notInLaunchRegion())
2434 const Box& gbx = fai.fabbox();
2437 get(fai).template setComplement<RunOn::Device>(val, domain_box, strt_comp, ncomp);
2442 template <
class FAB>
2443 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int> FOO>
2444 typename F::value_type
2449 using T =
typename FAB::value_type;
2451 #ifdef AMREX_USE_GPU
2453 auto const& ma = this->const_arrays();
2458 return ma[box_no](i,j,k,comp);
2463 #ifdef AMREX_USE_OMP
2464 #pragma omp parallel if (!system::regtest_reduction) reduction(+:sm)
2468 Box const& bx = mfi.growntilebox(nghost);
2469 auto const& a = this->const_array(mfi);
2473 tmp += a(i,j,k,comp);
2486 template <
class FAB>
2490 copyTo(dest, 0, 0, dest.nComp(), nghost);
2493 template <
class FAB>
2494 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
2498 setVal(val,0,n_comp,n_grow);
2501 template <
class FAB>
2502 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
2510 template <
class FAB>
2511 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
2521 template <
class FAB>
2522 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>Z>
2534 #ifdef AMREX_USE_GPU
2536 auto const& fa = this->arrays();
2540 fa[box_no](i,j,k,n+comp) = val;
2548 #ifdef AMREX_USE_OMP
2549 #pragma omp parallel if (Gpu::notInLaunchRegion())
2553 const Box& bx = fai.growntilebox(nghost);
2554 auto fab = this->array(fai);
2557 fab(i,j,k,n+comp) = val;
2563 template <
class FAB>
2564 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
2575 template <
class FAB>
2576 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>Z>
2587 BL_PROFILE(
"FabArray::setVal(val,region,comp,ncomp,nghost)");
2589 #ifdef AMREX_USE_GPU
2591 auto const& fa = this->arrays();
2596 fa[box_no](i,j,k,n+comp) = val;
2605 #ifdef AMREX_USE_OMP
2607 #pragma omp parallel if (Gpu::notInLaunchRegion())
2611 Box b = fai.growntilebox(nghost) & region;
2614 auto fab = this->array(fai);
2617 fab(i,j,k,n+comp) = val;
2624 template <
class FAB>
2625 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
2632 template <
class FAB>
2633 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>Z>
2641 #ifdef AMREX_USE_GPU
2643 auto const& fa = this->arrays();
2647 fa[box_no](i,j,k,n+comp) =
std::abs(fa[box_no](i,j,k,n+comp));
2655 #ifdef AMREX_USE_OMP
2656 #pragma omp parallel if (Gpu::notInLaunchRegion())
2660 const Box& bx = mfi.growntilebox(nghost);
2661 auto fab = this->array(mfi);
2664 fab(i,j,k,n+comp) =
std::abs(fab(i,j,k,n+comp));
2670 template <
class FAB>
2671 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>Z>
2677 #ifdef AMREX_USE_GPU
2679 auto const& fa = this->arrays();
2683 fa[box_no](i,j,k,n+comp) += val;
2691 #ifdef AMREX_USE_OMP
2692 #pragma omp parallel if (Gpu::notInLaunchRegion())
2696 const Box& bx = mfi.growntilebox(nghost);
2697 auto fab = this->array(mfi);
2700 fab(i,j,k,n+comp) += val;
2706 template <
class FAB>
2707 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>Z>
2711 BL_PROFILE(
"FabArray::plus(val, region, comp, num_comp, nghost)");
2713 #ifdef AMREX_USE_GPU
2715 auto const& fa = this->arrays();
2720 fa[box_no](i,j,k,n+comp) += val;
2729 #ifdef AMREX_USE_OMP
2730 #pragma omp parallel if (Gpu::notInLaunchRegion())
2734 const Box& bx = mfi.growntilebox(nghost) & region;
2736 auto fab = this->array(mfi);
2739 fab(i,j,k,n+comp) += val;
2746 template <
class FAB>
2747 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>Z>
2753 #ifdef AMREX_USE_GPU
2755 auto const& fa = this->arrays();
2759 fa[box_no](i,j,k,n+comp) *= val;
2767 #ifdef AMREX_USE_OMP
2768 #pragma omp parallel if (Gpu::notInLaunchRegion())
2772 const Box& bx = mfi.growntilebox(nghost);
2773 auto fab = this->array(mfi);
2776 fab(i,j,k,n+comp) *= val;
2782 template <
class FAB>
2783 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>Z>
2787 BL_PROFILE(
"FabArray::mult(val, region, comp, num_comp, nghost)");
2789 #ifdef AMREX_USE_GPU
2791 auto const& fa = this->arrays();
2796 fa[box_no](i,j,k,n+comp) *= val;
2805 #ifdef AMREX_USE_OMP
2806 #pragma omp parallel if (Gpu::notInLaunchRegion())
2810 const Box& bx = mfi.growntilebox(nghost) & region;
2812 auto fab = this->array(mfi);
2815 fab(i,j,k,n+comp) *= val;
2822 template <
class FAB>
2823 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>Z>
2829 #ifdef AMREX_USE_GPU
2831 auto const& fa = this->arrays();
2835 fa[box_no](i,j,k,n+comp) = numerator / fa[box_no](i,j,k,n+comp);
2843 #ifdef AMREX_USE_OMP
2844 #pragma omp parallel if (Gpu::notInLaunchRegion())
2848 const Box& bx = mfi.growntilebox(nghost);
2849 auto fab = this->array(mfi);
2852 fab(i,j,k,n+comp) = numerator / fab(i,j,k,n+comp);
2858 template <
class FAB>
2859 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>Z>
2863 BL_PROFILE(
"FabArray::invert(numerator, region, comp, num_comp, nghost)");
2865 #ifdef AMREX_USE_GPU
2867 auto const& fa = this->arrays();
2872 fa[box_no](i,j,k,n+comp) = numerator / fa[box_no](i,j,k,n+comp);
2881 #ifdef AMREX_USE_OMP
2882 #pragma omp parallel if (Gpu::notInLaunchRegion())
2886 const Box& bx = mfi.growntilebox(nghost) & region;
2888 auto fab = this->array(mfi);
2891 fab(i,j,k,n+comp) = numerator / fab(i,j,k,n+comp);
2898 template <
class FAB>
2905 #ifdef AMREX_USE_OMP
2906 #pragma omp parallel
2915 template <
class FAB>
2916 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int> FOO>
2918 int xcomp,
int ycomp,
int ncomp,
IntVect const& nghost)
2926 #ifdef AMREX_USE_GPU
2928 auto const& yma = y.
arrays();
2929 auto const& xma =
x.const_arrays();
2933 yma[box_no](i,j,k,ycomp+n) += a * xma[box_no](i,j,k,xcomp+n);
2941 #ifdef AMREX_USE_OMP
2942 #pragma omp parallel if (Gpu::notInLaunchRegion())
2946 const Box& bx = mfi.growntilebox(nghost);
2949 auto const& xfab =
x.const_array(mfi);
2950 auto const& yfab = y.
array(mfi);
2953 yfab(i,j,k,ycomp+n) += a * xfab(i,j,k,xcomp+n);
2960 template <
class FAB>
2961 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int> FOO>
2964 int xcomp,
int ycomp,
int ncomp,
IntVect const& nghost)
2972 #ifdef AMREX_USE_GPU
2974 auto const& yfa = y.
arrays();
2975 auto const& xfa =
x.const_arrays();
2979 yfa[box_no](i,j,k,n+ycomp) = xfa[box_no](i,j,k,n+xcomp)
2980 + a * yfa[box_no](i,j,k,n+ycomp);
2988 #ifdef AMREX_USE_OMP
2989 #pragma omp parallel if (Gpu::notInLaunchRegion())
2993 const Box& bx = mfi.growntilebox(nghost);
2994 auto const& xFab =
x.const_array(mfi);
2995 auto const& yFab = y.
array(mfi);
2998 yFab(i,j,k,n+ycomp) = xFab(i,j,k,n+xcomp)
2999 + a * yFab(i,j,k,n+ycomp);
3005 template <
class FAB>
3006 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int> FOO>
3011 int dstcomp,
int numcomp,
const IntVect& nghost)
3021 #ifdef AMREX_USE_GPU
3023 auto const& dstma = dst.
arrays();
3024 auto const& xma =
x.const_arrays();
3029 dstma[box_no](i,j,k,dstcomp+n) = a*xma[box_no](i,j,k,xcomp+n)
3030 +
b*yma[box_no](i,j,k,ycomp+n);
3038 #ifdef AMREX_USE_OMP
3039 #pragma omp parallel if (Gpu::notInLaunchRegion())
3043 const Box& bx = mfi.growntilebox(nghost);
3044 auto const& xfab =
x.const_array(mfi);
3046 auto const& dfab = dst.
array(mfi);
3049 dfab(i,j,k,dstcomp+n) = a*xfab(i,j,k,xcomp+n) +
b*yfab(i,j,k,ycomp+n);
3055 template <
class FAB>
3056 template <
typename BUF>
3061 if ( n_grow.max() > 0 ) {
3063 FillBoundary_finish<BUF>();
3067 template <
class FAB>
3068 template <
typename BUF>
3073 if ( n_grow.max() > 0 ) {
3074 FillBoundary_nowait<BUF>(0,
nComp(), n_grow, period, cross);
3075 FillBoundary_finish<BUF>();
3079 template <
class FAB>
3080 template <
typename BUF>
3086 "FillBoundary: asked to fill more ghost cells than we have");
3087 if ( nghost.
max() > 0 ) {
3088 FillBoundary_nowait<BUF>(0,
nComp(), nghost, period, cross);
3089 FillBoundary_finish<BUF>();
3093 template <
class FAB>
3094 template <
typename BUF>
3099 if ( n_grow.max() > 0 ) {
3101 FillBoundary_finish<BUF>();
3105 template <
class FAB>
3106 template <
typename BUF>
3111 if ( n_grow.max() > 0 ) {
3112 FillBoundary_nowait<BUF>(scomp, ncomp, n_grow, period, cross);
3113 FillBoundary_finish<BUF>();
3117 template <
class FAB>
3118 template <
typename BUF>
3125 "FillBoundary: asked to fill more ghost cells than we have");
3126 if ( nghost.
max() > 0 ) {
3127 FillBoundary_nowait<BUF>(scomp, ncomp, nghost, period, cross);
3128 FillBoundary_finish<BUF>();
3132 template <
class FAB>
3133 template <
typename BUF>
3140 template <
class FAB>
3141 template <
typename BUF>
3145 FillBoundary_nowait<BUF>(0,
nComp(),
nGrowVect(), period, cross);
3148 template <
class FAB>
3149 template <
typename BUF>
3153 FillBoundary_nowait<BUF>(0,
nComp(), nghost, period, cross);
3156 template <
class FAB>
3157 template <
typename BUF>
3164 template <
class FAB>
3168 BL_PROFILE(
"FabArray::FillBoundaryAndSync()");
3169 if (n_grow.max() > 0 || !is_cell_centered()) {
3170 FillBoundaryAndSync_nowait(0,
nComp(), n_grow, period);
3171 FillBoundaryAndSync_finish();
3175 template <
class FAB>
3180 BL_PROFILE(
"FabArray::FillBoundaryAndSync()");
3181 if (nghost.
max() > 0 || !is_cell_centered()) {
3182 FillBoundaryAndSync_nowait(scomp, ncomp, nghost, period);
3183 FillBoundaryAndSync_finish();
3187 template <
class FAB>
3194 template <
class FAB>
3200 FBEP_nowait(scomp, ncomp, nghost, period,
false,
false,
true);
3203 template <
class FAB>
3211 template <
class FAB>
3216 if (!is_cell_centered()) {
3222 template <
class FAB>
3227 if (!is_cell_centered()) {
3233 template <
class FAB>
3240 template <
class FAB>
3245 FBEP_nowait(scomp, ncomp,
IntVect(0), period,
false,
false,
true);
3248 template <
class FAB>
3256 template <
class FAB>
3260 SumBoundary(0, n_comp,
IntVect(0), period);
3263 template <
class FAB>
3267 SumBoundary(scomp, ncomp,
IntVect(0), period);
3270 template <
class FAB>
3274 SumBoundary(scomp, ncomp, this->
nGrowVect(), nghost, period);
3277 template <
class FAB>
3283 SumBoundary_nowait(scomp, ncomp, src_nghost, dst_nghost, period);
3284 SumBoundary_finish();
3287 template <
class FAB>
3291 SumBoundary_nowait(0, n_comp,
IntVect(0), period);
3294 template <
class FAB>
3298 SumBoundary_nowait(scomp, ncomp,
IntVect(0), period);
3301 template <
class FAB>
3305 SumBoundary_nowait(scomp, ncomp, this->
nGrowVect(), nghost, period);
3308 template <
class FAB>
3312 BL_PROFILE(
"FabArray<FAB>::SumBoundary_nowait()");
3319 amrex::Copy(*tmp, *
this, scomp, 0, ncomp, src_nghost);
3320 this->
setVal(
typename FAB::value_type(0), scomp, ncomp, dst_nghost);
3324 if (!this->pcd) {
delete tmp; }
3327 template <
class FAB>
3331 BL_PROFILE(
"FabArray<FAB>::SumBoundary_finish()");
3341 template <
class FAB>
3352 template <
class FAB>
3358 FBEP_nowait(scomp, ncomp,
nGrowVect(), period,
false,
true);
3363 template <
class FAB>
3370 FBEP_nowait(scomp, ncomp, nghost, period,
false,
true);
3375 template <
class FAB>
3376 template <
typename BUF>
3380 FBEP_nowait<BUF>(scomp, ncomp,
nGrowVect(), period, cross);
3383 template <
class FAB>
3384 template <
typename BUF>
3389 FBEP_nowait<BUF>(scomp, ncomp, nghost, period, cross);
3392 template <
class FAB>
3393 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>Z>
3401 int ncomp = this->
nComp();
3405 for (
int i = 0; i < AMREX_SPACEDIM; ++i) {
3407 domain.
grow(i, ngrow[i]);
3411 #ifdef AMREX_USE_GPU
3413 auto const& fa = this->arrays();
3417 auto const& fab = fa[box_no];
3422 }
else if (domain.
contains(i,j,k)) {
3423 fab(i,j,k,n) = notcovered;
3434 #ifdef AMREX_USE_OMP
3435 #pragma omp parallel if (Gpu::notInLaunchRegion())
3439 auto const& fab = this->array(mfi);
3440 Box const& fbx = mfi.growntilebox();
3441 Box const& gbx = fbx & domain;
3442 Box const& vbx = mfi.validbox();
3446 fab(i,j,k,n) = interior;
3448 fab(i,j,k,n) = notcovered;
3450 fab(i,j,k,n) = physbnd;
3457 setVal(covered, TheFB, 0, ncomp);
3460 template <
class FAB>
3461 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
3465 BL_PROFILE(
"FabArray::setVal(val, thecmd, scomp, ncomp)");
3467 #ifdef AMREX_USE_GPU
3470 CMD_local_setVal_gpu(val, thecmd, scomp, ncomp);
3471 CMD_remote_setVal_gpu(val, thecmd, scomp, ncomp);
3479 auto N_locs =
static_cast<int>(LocTags.size());
3480 #ifdef AMREX_USE_OMP
3481 #pragma omp parallel for if (thecmd.m_threadsafe_loc)
3483 for (
int i = 0; i < N_locs; ++i) {
3485 (*this)[tag.
dstIndex].template setVal<RunOn::Host>(val, tag.
dbox, scomp, ncomp);
3488 for (
const auto & RcvTag : RcvTags) {
3489 auto N =
static_cast<int>(RcvTag.second.size());
3490 #ifdef AMREX_USE_OMP
3491 #pragma omp parallel for if (thecmd.m_threadsafe_rcv)
3493 for (
int i = 0; i < N; ++i) {
3495 (*this)[tag.
dstIndex].template setVal<RunOn::Host>(val, tag.
dbox, scomp, ncomp);
3501 template <
class FAB>
3502 template <class F, std::enable_if_t<IsBaseFab<F>::value,
int>>
3509 #ifdef AMREX_USE_OMP
3510 #pragma omp parallel if (thecmd.m_threadsafe_rcv)
3519 auto N_locs =
static_cast<int>(LocTags.size());
3520 for (
int i = 0; i < N_locs; ++i) {
3525 for (
const auto & RcvTag : RcvTags) {
3526 auto N =
static_cast<int>(RcvTag.second.size());
3527 for (
int i = 0; i < N; ++i) {
3535 template <
class FAB>
3536 template <typename F, std::enable_if_t<IsBaseFab<F>::value,
int> FOO>
3537 typename F::value_type
3539 [[maybe_unused]]
bool ignore_covered)
const
3543 using RT =
typename F::value_type;
3548 if ( this->is_cell_centered() && this->hasEBFabFactory() && ignore_covered )
3552 #ifdef AMREX_USE_GPU
3554 auto const& flagsma = flags.const_arrays();
3555 auto const& ma = this->const_arrays();
3559 if (flagsma[box_no](i,j,k).isCovered()) {
3563 auto const& a = ma[box_no];
3564 for (
int n = 0; n < ncomp; ++n) {
3573 #ifdef AMREX_USE_OMP
3574 #pragma omp parallel reduction(max:nm0)
3577 Box const& bx = mfi.growntilebox(nghost);
3579 auto const& flag = flags.const_array(mfi);
3580 auto const& a = this->const_array(mfi);
3583 if (!flag(i,j,k).isCovered()) {
3594 #ifdef AMREX_USE_GPU
3596 auto const& ma = this->const_arrays();
3600 return std::abs(ma[box_no](i,j,k,comp+n));
3605 #ifdef AMREX_USE_OMP
3606 #pragma omp parallel reduction(max:nm0)
3609 Box const& bx = mfi.growntilebox(nghost);
3610 auto const& a = this->const_array(mfi);
3626 template <
class FAB>
3627 template <typename IFAB, typename F, std::enable_if_t<IsBaseFab<F>::value,
int> FOO>
3628 typename F::value_type
3630 IntVect const& nghost,
bool local)
const
3634 using RT =
typename F::value_type;
3638 #ifdef AMREX_USE_GPU
3640 auto const& ma = this->const_arrays();
3641 auto const& maskma =
mask.const_arrays();
3645 if (maskma[box_no](i,j,k)) {
3647 auto const& a = ma[box_no];
3648 for (
int n = 0; n < ncomp; ++n) {
3659 #ifdef AMREX_USE_OMP
3660 #pragma omp parallel reduction(max:nm0)
3663 Box const& bx = mfi.growntilebox(nghost);
3664 auto const& a = this->const_array(mfi);
3665 auto const& mskfab =
mask.const_array(mfi);
3668 if (mskfab(i,j,k)) {
#define BL_PROFILE(a)
Definition: AMReX_BLProfiler.H:551
#define AMREX_ALWAYS_ASSERT_WITH_MESSAGE(EX, MSG)
Definition: AMReX_BLassert.H:49
#define AMREX_ASSERT(EX)
Definition: AMReX_BLassert.H:38
#define AMREX_ALWAYS_ASSERT(EX)
Definition: AMReX_BLassert.H:50
#define AMREX_NODISCARD
Definition: AMReX_Extension.H:251
#define AMREX_FORCE_INLINE
Definition: AMReX_Extension.H:119
#define AMREX_RESTRICT
Definition: AMReX_Extension.H:37
#define AMREX_HOST_DEVICE_FOR_4D(...)
Definition: AMReX_GpuLaunch.nolint.H:51
#define AMREX_HOST_DEVICE_PARALLEL_FOR_4D(...)
Definition: AMReX_GpuLaunch.nolint.H:55
#define AMREX_IF_ON_DEVICE(CODE)
Definition: AMReX_GpuQualifiers.H:56
#define AMREX_IF_ON_HOST(CODE)
Definition: AMReX_GpuQualifiers.H:58
#define AMREX_GPU_DEVICE
Definition: AMReX_GpuQualifiers.H:18
#define AMREX_GPU_HOST_DEVICE
Definition: AMReX_GpuQualifiers.H:20
Array4< int const > offset
Definition: AMReX_HypreMLABecLap.cpp:1089
Array4< int const > mask
Definition: AMReX_InterpFaceRegister.cpp:93
#define AMREX_LOOP_3D(bx, i, j, k, block)
Definition: AMReX_Loop.nolint.H:4
#define AMREX_LOOP_4D(bx, ncomp, i, j, k, n, block)
Definition: AMReX_Loop.nolint.H:16
int MPI_Comm
Definition: AMReX_ccse-mpi.H:47
if(!(yy_init))
Definition: amrex_iparser.lex.nolint.H:935
A virtual base class for objects that manage their own dynamic memory allocation.
Definition: AMReX_Arena.H:100
virtual void free(void *pt)=0
A pure virtual function for deleting the arena pointed to by pt.
virtual void * alloc(std::size_t sz)=0
A FortranArrayBox(FAB)-like object.
Definition: AMReX_BaseFab.H:183
A collection of Boxes stored in an Array.
Definition: AMReX_BoxArray.H:530
AMREX_GPU_HOST_DEVICE const IntVectND< dim > & smallEnd() const &noexcept
Get the smallend of the BoxND.
Definition: AMReX_Box.H:105
AMREX_GPU_HOST_DEVICE BoxND & grow(int i) noexcept
Definition: AMReX_Box.H:627
AMREX_GPU_HOST_DEVICE IntVectND< dim > length() const noexcept
Return the length of the BoxND.
Definition: AMReX_Box.H:146
AMREX_GPU_HOST_DEVICE bool ok() const noexcept
Checks if it is a proper BoxND (including a valid type).
Definition: AMReX_Box.H:200
AMREX_GPU_HOST_DEVICE bool contains(const IntVectND< dim > &p) const noexcept
Returns true if argument is contained within BoxND.
Definition: AMReX_Box.H:204
Definition: AMReX_FabFactory.H:76
Calculates the distribution of FABs to MPI processes.
Definition: AMReX_DistributionMapping.H:41
Definition: AMReX_EBFabFactory.H:22
const FabArray< EBCellFlagFab > & getMultiEBCellFlagFab() const noexcept
Definition: AMReX_EBFabFactory.H:48
Base class for FabArray.
Definition: AMReX_FabArrayBase.H:41
IntVect nGrowVect() const noexcept
Definition: AMReX_FabArrayBase.H:79
Vector< int > indexArray
Definition: AMReX_FabArrayBase.H:445
FabArrayBase & operator=(const FabArrayBase &rhs)=default
static bool getAllocSingleChunk()
Definition: AMReX_FabArrayBase.H:727
const BoxArray & boxArray() const noexcept
Return a constant reference to the BoxArray that defines the valid region associated with this FabArr...
Definition: AMReX_FabArrayBase.H:94
bool isFusingCandidate() const noexcept
Is this a good candidate for kernel fusing?
void define(const BoxArray &bxs, const DistributionMapping &dm, int nvar, int ngrow)
CopyComTag::CopyComTagsContainer CopyComTagsContainer
Definition: AMReX_FabArrayBase.H:219
CopyComTag::MapOfCopyComTagContainers MapOfCopyComTagContainers
Definition: AMReX_FabArrayBase.H:220
int local_size() const noexcept
Return the number of local FABs in the FabArray.
Definition: AMReX_FabArrayBase.H:112
const DistributionMapping & DistributionMap() const noexcept
Return constant reference to associated DistributionMapping.
Definition: AMReX_FabArrayBase.H:130
CpOp
parallel copy or add
Definition: AMReX_FabArrayBase.H:393
@ ADD
Definition: AMReX_FabArrayBase.H:393
@ COPY
Definition: AMReX_FabArrayBase.H:393
DistributionMapping distributionMap
Definition: AMReX_FabArrayBase.H:444
int nComp() const noexcept
Return number of variables (aka components) associated with each point.
Definition: AMReX_FabArrayBase.H:82
static AMREX_EXPORT FabArrayStats m_FA_stats
Definition: AMReX_FabArrayBase.H:723
An Array of FortranArrayBox(FAB)-like Objects.
Definition: AMReX_FabArray.H:343
void ParallelCopyToGhost_finish()
void setFab(int boxno, std::unique_ptr< FAB > elem)
Explicitly set the Kth FAB in the FabArray to point to elem.
Definition: AMReX_FabArray.H:2210
void copy(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, int src_nghost, int dst_nghost, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:941
F::value_type sum(int comp, IntVect const &nghost, bool local=false) const
Returns the sum of component "comp".
Definition: AMReX_FabArray.H:2445
void EnforcePeriodicity(int scomp, int ncomp, const IntVect &nghost, const Periodicity &period)
Definition: AMReX_FabArray.H:3365
const FAB & get(const MFIter &mfi) const noexcept
Return a constant reference to the FAB associated with mfi.
Definition: AMReX_FabArray.H:508
void abs(int comp, int ncomp, int nghost=0)
Definition: AMReX_FabArray.H:2627
void copy(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const IntVect &src_nghost, const IntVect &dst_nghost, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:952
Array4< typename FabArray< FAB >::value_type const > const_array(const MFIter &mfi) const noexcept
Definition: AMReX_FabArray.H:1592
void * m_dp_arrays
Definition: AMReX_FabArray.H:1281
void FBEP_nowait(int scomp, int ncomp, const IntVect &nghost, const Periodicity &period, bool cross, bool enforce_periodicity_only=false, bool override_sync=false)
static void pack_send_buffer_gpu(FabArray< FAB > const &src, int scomp, int ncomp, Vector< char * > const &send_data, Vector< std::size_t > const &send_size, Vector< const CopyComTagsContainer * > const &send_cctc)
void ParallelCopy(const FabArray< FAB > &src, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:777
typename std::conditional_t< IsBaseFab< FAB >::value, FAB, FABType >::value_type value_type
Definition: AMReX_FabArray.H:354
std::unique_ptr< FabArray< FAB > > os_temp
Definition: AMReX_FabArray.H:1411
void FillBoundary(const IntVect &nghost, const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3082
void prefetchToDevice(const MFIter &mfi) const noexcept
Definition: AMReX_FabArray.H:1548
void FillBoundary_nowait(const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3143
void shift(const IntVect &v)
Shift the boxarray by vector v.
Definition: AMReX_FabArray.H:2900
bool ok() const
Return true if the FabArray is well-defined. That is, the FabArray has a BoxArray and DistributionMap...
Definition: AMReX_FabArray.H:1991
Array4< typename FabArray< FAB >::value_type const > array(const MFIter &mfi, int start_comp) const noexcept
Definition: AMReX_FabArray.H:1608
void CMD_local_setVal_gpu(value_type x, const CommMetaData &thecmd, int scomp, int ncomp)
Definition: AMReX_FBI.H:332
void CMD_remote_setVal_gpu(value_type x, const CommMetaData &thecmd, int scomp, int ncomp)
Definition: AMReX_FBI.H:362
void ParallelAdd_nowait(const FabArray< FAB > &src, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:788
value_type * singleChunkPtr() noexcept
Definition: AMReX_FabArray.H:459
FabArray(FabArray< FAB > &&rhs) noexcept
Definition: AMReX_FabArray.H:1926
Array4< typename FabArray< FAB >::value_type const > array(int K, int start_comp) const noexcept
Definition: AMReX_FabArray.H:1624
FabArray(const BoxArray &bxs, const DistributionMapping &dm, int nvar, const IntVect &ngrow, const MFInfo &info=MFInfo(), const FabFactory< FAB > &factory=DefaultFabFactory< FAB >())
Definition: AMReX_FabArray.H:1890
bool defined(const MFIter &mfi) const noexcept
Definition: AMReX_FabArray.H:1484
void setVal(value_type val, const CommMetaData &thecmd, int scomp, int ncomp)
Definition: AMReX_FabArray.H:3463
void OverrideSync_nowait(int scomp, int ncomp, const Periodicity &period)
Definition: AMReX_FabArray.H:3242
FabArray(const FabArray< FAB > &rhs)=delete
void ParallelCopyToGhost(const FabArray< FAB > &src, int scomp, int dcomp, int ncomp, const IntVect &snghost, const IntVect &dnghost, const Periodicity &period=Periodicity::NonPeriodic())
void FillBoundary_test()
Definition: AMReX_FabArrayCommI.H:831
void ParallelCopy_finish()
const FabFactory< FAB > & Factory() const noexcept
Definition: AMReX_FabArray.H:441
void OverrideSync(const Periodicity &period=Periodicity::NonPeriodic())
Synchronize nodal data. The synchronization will override valid regions by the intersecting valid reg...
Definition: AMReX_FabArray.H:3213
void FillBoundary(bool cross=false)
Copy on intersection within a FabArray. Data is copied from valid regions to intersecting regions of ...
Definition: AMReX_FabArray.H:3058
FAB const * fabPtr(int K) const noexcept
Definition: AMReX_FabArray.H:1526
void clear_arrays()
Definition: AMReX_FabArray.H:1717
void ParallelCopy(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, int src_nghost, int dst_nghost, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:859
void SumBoundary(const Periodicity &period=Periodicity::NonPeriodic())
Sum values in overlapped cells. The destination is limited to valid cells.
Definition: AMReX_FabArray.H:3258
Long m_single_chunk_size
Definition: AMReX_FabArray.H:1271
FAB & get(int K) noexcept
Return a reference to the FAB associated with the Kth element.
Definition: AMReX_FabArray.H:526
static void LinComb(FabArray< FAB > &dst, value_type a, const FabArray< FAB > &x, int xcomp, value_type b, const FabArray< FAB > &y, int ycomp, int dstcomp, int numcomp, const IntVect &nghost)
dst = a*x + b*y
Definition: AMReX_FabArray.H:3008
const FAB & get(int K) const noexcept
Return a constant reference to the FAB associated with the Kth element.
Definition: AMReX_FabArray.H:520
void OverrideSync_nowait(const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3235
AMREX_NODISCARD FAB * release(int K)
Release ownership of the FAB. This function is not thread safe.
Definition: AMReX_FabArray.H:1734
void setDomainBndry(value_type val, const Geometry &geom)
Set all values outside the Geometry domain to val.
Definition: AMReX_FabArray.H:2406
std::unique_ptr< PCData< FAB > > pcd
Definition: AMReX_FabArray.H:1408
void define(const BoxArray &bxs, const DistributionMapping &dm, int nvar, int ngrow, const MFInfo &info=MFInfo(), const FabFactory< FAB > &factory=DefaultFabFactory< FAB >())
Define this FabArray identically to that performed by the constructor having an analogous function si...
Definition: AMReX_FabArray.H:2026
void ParallelAdd_nowait(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const IntVect &src_nghost, const IntVect &dst_nghost, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:888
std::unique_ptr< FBData< FAB > > fbd
Definition: AMReX_FabArray.H:1407
std::unique_ptr< detail::SingleChunkArena > m_single_chunk_arena
Definition: AMReX_FabArray.H:1270
FabArray(const FabArray< FAB > &rhs, MakeType maketype, int scomp, int ncomp)
Definition: AMReX_FabArray.H:1904
void ParallelAdd(const FabArray< FAB > &src, const Periodicity &period=Periodicity::NonPeriodic())
This function copies data from src to this FabArray. Each FAB in fa is intersected with all FABs in t...
Definition: AMReX_FabArray.H:774
FAB fab_type
Definition: AMReX_FabArray.H:356
void BuildMask(const Box &phys_domain, const Periodicity &period, value_type covered, value_type notcovered, value_type physbnd, value_type interior)
Definition: AMReX_FabArray.H:3395
void OverrideSync(int scomp, int ncomp, const Periodicity &period)
Synchronize nodal data. The synchronization will override valid regions by the intersecting valid reg...
Definition: AMReX_FabArray.H:3224
void LocalCopy(FabArray< SFAB > const &src, int scomp, int dcomp, int ncomp, IntVect const &nghost)
Perform local copy of FabArray data.
Definition: AMReX_FabArray.H:1817
bool SharedMemory() const noexcept
Definition: AMReX_FabArray.H:1339
value_type const * singleChunkPtr() const noexcept
Definition: AMReX_FabArray.H:465
LayoutData< int > RecvLayoutMask(const CommMetaData &thecmd)
Definition: AMReX_FabArray.H:3504
void FillBoundary(int scomp, int ncomp, const IntVect &nghost, const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3120
void FillBoundary_nowait(const IntVect &nghost, const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3151
Vector< std::string > m_tags
Definition: AMReX_FabArray.H:1287
void ParallelCopyToGhost_nowait(const FabArray< FAB > &src, int scomp, int dcomp, int ncomp, const IntVect &snghost, const IntVect &dnghost, const Periodicity &period=Periodicity::NonPeriodic())
void FB_local_copy_cpu(const FB &TheFB, int scomp, int ncomp)
Definition: AMReX_FBI.H:211
void invert(value_type numerator, int comp, int num_comp, int nghost=0)
Definition: AMReX_FabArray.H:2825
void FB_local_copy_gpu(const FB &TheFB, int scomp, int ncomp)
Definition: AMReX_FBI.H:272
static void unpack_recv_buffer_cpu(FabArray< FAB > &dst, int dcomp, int ncomp, Vector< char * > const &recv_data, Vector< std::size_t > const &recv_size, Vector< const CopyComTagsContainer * > const &recv_cctc, CpOp op, bool is_thread_safe)
F::value_type norminf(int comp, int ncomp, IntVect const &nghost, bool local=false, [[maybe_unused]] bool ignore_covered=false) const
Return infinity norm.
Definition: AMReX_FabArray.H:3538
void ParallelAdd_nowait(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:827
void PC_local_gpu(const CPC &thecpc, FabArray< FAB > const &src, int scomp, int dcomp, int ncomp, CpOp op)
Definition: AMReX_PCI.H:88
Array4< typename FabArray< FAB >::value_type const > array(const MFIter &mfi) const noexcept
Definition: AMReX_FabArray.H:1560
void setFab_assert(int K, FAB const &fab) const
Definition: AMReX_FabArray.H:2198
Array4< typename FabArray< FAB >::value_type const > const_array(int K) const noexcept
Definition: AMReX_FabArray.H:1600
void plus(value_type val, int comp, int num_comp, int nghost=0)
Definition: AMReX_FabArray.H:2673
DataAllocator m_dallocator
Definition: AMReX_FabArray.H:1269
void FillBoundaryAndSync(int scomp, int ncomp, const IntVect &nghost, const Periodicity &period)
Fill ghost cells and synchronize nodal data. Ghost regions are filled with data from the intersecting...
Definition: AMReX_FabArray.H:3177
void copy(const FabArray< FAB > &src, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:783
void FillBoundaryAndSync_finish()
Definition: AMReX_FabArray.H:3205
void ParallelCopy_nowait(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, int src_nghost, int dst_nghost, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:898
void SumBoundary_nowait(int scomp, int ncomp, IntVect const &src_nghost, IntVect const &dst_nghost, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3310
void SumBoundary_nowait(int scomp, int ncomp, IntVect const &nghost, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3303
void FillBoundary_nowait(bool cross=false)
Definition: AMReX_FabArray.H:3135
void FillBoundary(const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3070
static void Xpay(FabArray< FAB > &y, value_type a, FabArray< FAB > const &x, int xcomp, int ycomp, int ncomp, IntVect const &nghost)
y = x + a*y
Definition: AMReX_FabArray.H:2963
void clear()
Releases FAB memory in the FabArray.
Definition: AMReX_FabArray.H:1773
void FillBoundary_nowait(int scomp, int ncomp, const IntVect &nghost, const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3386
void FillBoundary(int scomp, int ncomp, bool cross=false)
Same as FillBoundary(), but only copies ncomp components starting at scomp.
Definition: AMReX_FabArray.H:3096
void FillBoundary_nowait(int scomp, int ncomp, bool cross=false)
Definition: AMReX_FabArray.H:3159
Array4< typename FabArray< FAB >::value_type > array(const MFIter &mfi) noexcept
Definition: AMReX_FabArray.H:1568
std::vector< FAB * > m_fabs_v
The data.
Definition: AMReX_FabArray.H:1278
void setBndry(value_type val)
Set all values in the boundary region to val.
Definition: AMReX_FabArray.H:2297
void SumBoundary(int scomp, int ncomp, IntVect const &nghost, const Periodicity &period=Periodicity::NonPeriodic())
Sum values in overlapped cells. The destination is limited to valid + ngrow cells.
Definition: AMReX_FabArray.H:3272
void FillBoundaryAndSync(const Periodicity &period=Periodicity::NonPeriodic())
Fill ghost cells and synchronize nodal data. Ghost regions are filled with data from the intersecting...
Definition: AMReX_FabArray.H:3166
void ParallelAdd(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const IntVect &src_nghost, const IntVect &dst_nghost, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:851
bool define_function_called
has define() been called?
Definition: AMReX_FabArray.H:1274
void ParallelAdd(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const Periodicity &period=Periodicity::NonPeriodic())
This function copies data from src to this FabArray. Each FAB in src is intersected with all FABs in ...
Definition: AMReX_FabArray.H:804
FabArray() noexcept
Constructs an empty FabArray<FAB>.
Definition: AMReX_FabArray.H:1865
static void unpack_recv_buffer_gpu(FabArray< FAB > &dst, int dcomp, int ncomp, Vector< char * > const &recv_data, Vector< std::size_t > const &recv_size, Vector< const CopyComTagsContainer * > const &recv_cctc, CpOp op, bool is_thread_safe)
bool defined(int K) const noexcept
Definition: AMReX_FabArray.H:1471
FAB * fabPtr(int K) noexcept
Definition: AMReX_FabArray.H:1517
void SumBoundary(int scomp, int ncomp, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3265
void ParallelAdd(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, int src_nghost, int dst_nghost, const Periodicity &period=Periodicity::NonPeriodic())
Similar to the above function, except that source and destination are grown by src_nghost and dst_ngh...
Definition: AMReX_FabArray.H:842
std::unique_ptr< FabFactory< FAB > > m_factory
Definition: AMReX_FabArray.H:1268
void Redistribute(const FabArray< FAB > &src, int scomp, int dcomp, int ncomp, const IntVect &nghost)
Copy from src to this. this and src have the same BoxArray, but different DistributionMapping.
Definition: AMReX_FabArrayCommI.H:804
void setVal(value_type val)
Set all components in the entire region of each FAB to val.
Definition: AMReX_FabArray.H:2496
typename std::vector< FAB * >::iterator Iterator
Definition: AMReX_FabArray.H:1342
void copy(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:819
MultiArray4< typename FabArray< FAB >::value_type > arrays() noexcept
Definition: AMReX_FabArray.H:1656
void copyTo(FAB &dest, int nghost=0) const
Copy the values contained in the intersection of the valid + nghost region of this FabArray with the ...
Definition: AMReX_FabArray.H:2488
void ParallelAdd_nowait(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, int src_nghost, int dst_nghost, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:878
void SumBoundary(int scomp, int ncomp, IntVect const &src_nghost, IntVect const &dst_nghost, const Periodicity &period=Periodicity::NonPeriodic())
Sum values in overlapped cells. For computing the overlap, the dst is grown by dst_ngrow,...
Definition: AMReX_FabArray.H:3279
void FillBoundary_nowait(int scomp, int ncomp, const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3378
void SumBoundary_finish()
Definition: AMReX_FabArray.H:3329
void SumBoundary_nowait(int scomp, int ncomp, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3296
std::size_t singleChunkSize() const noexcept
Definition: AMReX_FabArray.H:471
void mult(value_type val, int comp, int num_comp, int nghost=0)
Definition: AMReX_FabArray.H:2749
void ParallelCopy_nowait(const FabArray< FAB > &src, int scomp, int dcomp, int ncomp, const IntVect &snghost, const IntVect &dnghost, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY, const FabArrayBase::CPC *a_cpc=nullptr, bool to_ghost_cells_only=false)
static void Saxpy(FabArray< FAB > &y, value_type a, FabArray< FAB > const &x, int xcomp, int ycomp, int ncomp, IntVect const &nghost)
y += a*x
Definition: AMReX_FabArray.H:2917
MultiArray4< value_type > m_arrays
Definition: AMReX_FabArray.H:1284
void AllocFabs(const FabFactory< FAB > &factory, Arena *ar, const Vector< std::string > &tags, bool alloc_single_chunk)
Definition: AMReX_FabArray.H:2071
void FillBoundaryAndSync_nowait(const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3189
Array4< typename FabArray< FAB >::value_type const > const_array(int K, int start_comp) const noexcept
Definition: AMReX_FabArray.H:1648
void FillBoundaryAndSync_nowait(int scomp, int ncomp, const IntVect &nghost, const Periodicity &period)
Definition: AMReX_FabArray.H:3196
Array4< typename FabArray< FAB >::value_type > array(int K) noexcept
Definition: AMReX_FabArray.H:1584
void * m_hp_arrays
Definition: AMReX_FabArray.H:1283
ShMem shmem
Definition: AMReX_FabArray.H:1337
void LocalAdd(FabArray< FAB > const &src, int scomp, int dcomp, int ncomp, IntVect const &nghost)
Perform local addition of FabArray data.
Definition: AMReX_FabArray.H:1826
FabArray< FAB > & operator=(FabArray< FAB > &&rhs) noexcept
Definition: AMReX_FabArray.H:1952
MultiArray4< value_type const > m_const_arrays
Definition: AMReX_FabArray.H:1285
void ParallelCopy(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:810
bool hasEBFabFactory() const noexcept
Definition: AMReX_FabArray.H:448
FAB & atLocalIdx(int L) noexcept
Return a reference to the FAB associated with local index L.
Definition: AMReX_FabArray.H:529
Array4< typename FabArray< FAB >::value_type const > array(int K) const noexcept
Definition: AMReX_FabArray.H:1576
FAB & get(const MFIter &mfi) noexcept
Returns a reference to the FAB associated mfi.
Definition: AMReX_FabArray.H:514
const Vector< std::string > & tags() const noexcept
Definition: AMReX_FabArray.H:446
bool isAllRegular() const noexcept
Definition: AMReX_FabArray.H:473
Array4< typename FabArray< FAB >::value_type const > const_array(const MFIter &mfi, int start_comp) const noexcept
Definition: AMReX_FabArray.H:1640
Arena * arena() const noexcept
Definition: AMReX_FabArray.H:444
static void pack_send_buffer_cpu(FabArray< FAB > const &src, int scomp, int ncomp, Vector< char * > const &send_data, Vector< std::size_t > const &send_size, Vector< const CopyComTagsContainer * > const &send_cctc)
void build_arrays() const
Definition: AMReX_FabArray.H:1683
void SumBoundary_nowait(const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3289
const FAB & atLocalIdx(int L) const noexcept
Definition: AMReX_FabArray.H:530
void ParallelCopy(const FabArray< FAB > &src, int scomp, int dcomp, int ncomp, const IntVect &snghost, const IntVect &dnghost, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY, const FabArrayBase::CPC *a_cpc=nullptr)
FAB * fabPtr(const MFIter &mfi) noexcept
Return pointer to FAB.
Definition: AMReX_FabArray.H:1497
Array4< typename FabArray< FAB >::value_type > array(const MFIter &mfi, int start_comp) noexcept
Definition: AMReX_FabArray.H:1616
void copyTo(FAB &dest, int scomp, int dcomp, int ncomp, int nghost=0) const
Copy the values contained in the intersection of the num_comp component valid + nghost region of this...
void EnforcePeriodicity(const Periodicity &period)
Fill ghost cells with values from their corresponding cells across periodic boundaries,...
Definition: AMReX_FabArray.H:3343
void OverrideSync_finish()
Definition: AMReX_FabArray.H:3250
void FillBoundary(int scomp, int ncomp, const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3108
MultiArray4< typename FabArray< FAB >::value_type const > const_arrays() const noexcept
Definition: AMReX_FabArray.H:1674
void prefetchToHost(const MFIter &mfi) const noexcept
Definition: AMReX_FabArray.H:1536
void define(const BoxArray &bxs, const DistributionMapping &dm, int nvar, const IntVect &ngrow, const MFInfo &info=MFInfo(), const FabFactory< FAB > &factory=DefaultFabFactory< FAB >())
Definition: AMReX_FabArray.H:2038
const FAB & operator[](const MFIter &mfi) const noexcept
Return a constant reference to the FAB associated with mfi.
Definition: AMReX_FabArray.H:505
F::value_type norminf(FabArray< IFAB > const &mask, int comp, int ncomp, IntVect const &nghost, bool local=false) const
Return infinity norm in masked region.
Definition: AMReX_FabArray.H:3629
void ParallelCopy_nowait(const FabArray< FAB > &src, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:791
void EnforcePeriodicity(int scomp, int ncomp, const Periodicity &period)
Definition: AMReX_FabArray.H:3354
void PC_local_cpu(const CPC &thecpc, FabArray< FAB > const &src, int scomp, int dcomp, int ncomp, CpOp op)
Definition: AMReX_PCI.H:6
void FillBoundary_finish()
FAB const * fabPtr(const MFIter &mfi) const noexcept
Definition: AMReX_FabArray.H:1507
void ParallelCopy_nowait(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:833
Array4< typename FabArray< FAB >::value_type > array(int K, int start_comp) noexcept
Definition: AMReX_FabArray.H:1632
bool isDefined() const
Definition: AMReX_FabArray.H:2019
~FabArray()
The destructor – deletes all FABs in the array.
Definition: AMReX_FabArray.H:1983
Definition: AMReX_FabFactory.H:50
virtual AMREX_NODISCARD FAB * create(const Box &box, int ncomps, const FabInfo &info, int box_index) const =0
virtual AMREX_NODISCARD FabFactory< FAB > * clone() const =0
virtual AMREX_NODISCARD Long nBytes(const Box &box, int ncomps, int) const
Definition: AMReX_FabFactory.H:64
Rectangular problem domain geometry.
Definition: AMReX_Geometry.H:73
const Box & Domain() const noexcept
Returns our rectangular domain.
Definition: AMReX_Geometry.H:210
bool isPeriodic(int dir) const noexcept
Is the domain periodic in the specified direction?
Definition: AMReX_Geometry.H:331
Definition: AMReX_Tuple.H:93
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE bool allGE(const IntVectND< dim > &rhs) const noexcept
Returns true if this is greater than or equal to argument for all components. NOTE: This is NOT a str...
Definition: AMReX_IntVect.H:443
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE bool allLE(const IntVectND< dim > &rhs) const noexcept
Returns true if this is less than or equal to argument for all components. NOTE: This is NOT a strict...
Definition: AMReX_IntVect.H:393
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE int max() const noexcept
maximum (no absolute values) value
Definition: AMReX_IntVect.H:214
AMREX_GPU_HOST_DEVICE static constexpr AMREX_FORCE_INLINE IntVectND< dim > TheZeroVector() noexcept
This static member function returns a reference to a constant IntVectND object, all of whose dim argu...
Definition: AMReX_IntVect.H:672
Definition: AMReX_MFIter.H:57
bool isValid() const noexcept
Is the iterator valid i.e. is it associated with a FAB?
Definition: AMReX_MFIter.H:141
int index() const noexcept
The index into the underlying BoxArray of the current FAB.
Definition: AMReX_MFIter.H:144
int LocalIndex() const noexcept
Return local index into the vector of fab pointers, m_fabs_v When AllBoxes is on, local_index_map is ...
Definition: AMReX_MFIter.H:159
This provides length of period for periodic domains. 0 means it is not periodic in that direction....
Definition: AMReX_Periodicity.H:17
static const Periodicity & NonPeriodic() noexcept
Definition: AMReX_Periodicity.cpp:52
bool isAnyPeriodic() const noexcept
Definition: AMReX_Periodicity.H:22
bool isPeriodic(int dir) const noexcept
Definition: AMReX_Periodicity.H:26
Long size() const noexcept
Definition: AMReX_Vector.H:50
#define abs(x)
Definition: complex-type.h:85
@ FAB
Definition: AMReX_AmrvisConstants.H:86
AMREX_GPU_HOST_DEVICE Long size(T const &b) noexcept
integer version
Definition: AMReX_GpuRange.H:26
void streamSynchronize() noexcept
Definition: AMReX_GpuDevice.H:237
bool inLaunchRegion() noexcept
Definition: AMReX_GpuControl.H:86
bool inNoSyncRegion() noexcept
Definition: AMReX_GpuControl.H:146
void htod_memcpy(void *p_d, const void *p_h, const std::size_t sz) noexcept
Definition: AMReX_GpuDevice.H:293
std::enable_if_t< IsBaseFab< FAB >) &&IsDataPacking< DataPacking, FAB >)> ParallelCopy_finish(FabArray< FAB > &dest, CommHandler handler, const FabArrayBase::CommMetaData &cmd, const DataPacking &data_packing)
Definition: AMReX_NonLocalBC.H:793
std::enable_if_t< IsBaseFab< FAB >::value > PrepareSendBuffers(const PackComponents &components, FabArray< FAB > &dest, const FabArray< FAB > &src, CommData &comm, const FabArrayBase::MapOfCopyComTagContainers &cctc)
Calls PrepareComBuffers.
Definition: AMReX_NonLocalBC.H:555
AMREX_NODISCARD CommHandler ParallelCopy_nowait(NoLocalCopy, FabArray< FAB > &dest, const FabArray< FAB > &src, const FabArrayBase::CommMetaData &cmd, const DataPacking &data_packing)
Definition: AMReX_NonLocalBC.H:701
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{})
Finish communication started by FillBoundary_nowait.
void Min(KeyValuePair< K, V > &vi, MPI_Comm comm)
Definition: AMReX_ParallelReduce.H:152
void Sum(T &v, MPI_Comm comm)
Definition: AMReX_ParallelReduce.H:204
void Max(KeyValuePair< K, V > &vi, MPI_Comm comm)
Definition: AMReX_ParallelReduce.H:126
MPI_Comm CommunicatorSub() noexcept
sub-communicator for current frame
Definition: AMReX_ParallelContext.H:70
const ProcessTeam & MyTeam() noexcept
Definition: AMReX_ParallelDescriptor.H:349
int MyProc() noexcept
return the rank number local to the current Parallel Context
Definition: AMReX_ParallelDescriptor.H:125
int MyTeamLead() noexcept
Definition: AMReX_ParallelDescriptor.H:309
int TeamSize() noexcept
Definition: AMReX_ParallelDescriptor.H:294
int SeqNum() noexcept
Returns sequential message sequence numbers, usually used as tags for send/recv.
Definition: AMReX_ParallelDescriptor.H:613
static int f(amrex::Real t, N_Vector y_data, N_Vector y_rhs, void *user_data)
Definition: AMReX_SundialsIntegrator.H:44
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void swap(T &a, T &b) noexcept
Definition: AMReX_algoim_K.H:113
@ max
Definition: AMReX_ParallelReduce.H:17
constexpr bool is_convertible(T)
Definition: AMReX_TypeTraits.H:246
logical function omp_in_parallel()
Definition: AMReX_omp_mod.F90:41
Definition: AMReX_Amr.cpp:49
std::enable_if_t< std::is_integral_v< T > > ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
Definition: AMReX_CTOParallelForImpl.H:200
DistributionMapping const & DistributionMap(FabArrayBase const &fa)
MakeType
Definition: AMReX_MakeType.H:7
@ make_alias
Definition: AMReX_MakeType.H:7
int nComp(FabArrayBase const &fa)
std::unique_ptr< char, TheFaArenaDeleter > TheFaArenaPointer
Definition: AMReX_FabArray.H:103
IntVect nGrowVect(FabArrayBase const &fa)
AMREX_GPU_HOST_DEVICE constexpr AMREX_FORCE_INLINE const T & max(const T &a, const T &b) noexcept
Definition: AMReX_Algorithm.H:35
void Copy(FabArray< DFAB > &dst, FabArray< SFAB > const &src, int srccomp, int dstcomp, int numcomp, int nghost)
Definition: AMReX_FabArray.H:178
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE BoxND< dim > convert(const BoxND< dim > &b, const IntVectND< dim > &typ) noexcept
Returns a BoxND with different type.
Definition: AMReX_Box.H:1435
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE T abs(const GpuComplex< T > &a_z) noexcept
Return the absolute value of a complex number.
Definition: AMReX_GpuComplex.H:356
constexpr AMREX_GPU_HOST_DEVICE GpuTupleElement< I, GpuTuple< Ts... > >::type & get(GpuTuple< Ts... > &tup) noexcept
Definition: AMReX_Tuple.H:179
Long nBytesOwned(T const &) noexcept
Definition: AMReX_FabArray.H:57
BoxArray const & boxArray(FabArrayBase const &fa)
ReduceData< Ts... >::Type ParReduce(TypeList< Ops... > operation_list, TypeList< Ts... > type_list, FabArray< FAB > const &fa, IntVect const &nghost, F &&f)
Parallel reduce for MultiFab/FabArray.
Definition: AMReX_ParReduce.H:47
Arena * The_Comms_Arena()
Definition: AMReX_Arena.cpp:654
IntVectND< AMREX_SPACEDIM > IntVect
Definition: AMReX_BaseFwd.H:30
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void ignore_unused(const Ts &...)
This shuts up the compiler about unused variables.
Definition: AMReX.H:111
void setBndry(MF &dst, typename MF::value_type val, int scomp, int ncomp)
dst = val in ghost cells.
Definition: AMReX_FabArrayUtility.H:1614
bool TilingIfNotGPU() noexcept
Definition: AMReX_MFIter.H:12
void Add(FabArray< FAB > &dst, FabArray< FAB > const &src, int srccomp, int dstcomp, int numcomp, int nghost)
Definition: AMReX_FabArray.H:239
Arena * The_Pinned_Arena()
Definition: AMReX_Arena.cpp:634
void Abort(const std::string &msg)
Print out message to cerr and exit via abort().
Definition: AMReX.cpp:221
void OverrideSync_finish(FabArray< FAB > &fa)
Definition: AMReX_FabArrayUtility.H:1382
void update_fab_stats(Long n, Long s, size_t szt) noexcept
Definition: AMReX_BaseFab.cpp:144
void setVal(MF &dst, typename MF::value_type val)
dst = val
Definition: AMReX_FabArrayUtility.H:1607
Arena * The_Arena()
Definition: AMReX_Arena.cpp:594
void OverrideSync_nowait(FabArray< FAB > &fa, FabArray< IFAB > const &msk, const Periodicity &period)
Definition: AMReX_FabArrayUtility.H:1335
Definition: AMReX_TagParallelFor.H:57
Definition: AMReX_Array4.H:61
Definition: AMReX_DataAllocator.H:9
Arena * arena() const noexcept
Definition: AMReX_DataAllocator.H:24
Arena * m_arena
Definition: AMReX_DataAllocator.H:10
Definition: AMReX_FabArray.H:107
const FabArrayBase::FB * fb
Definition: AMReX_FabArray.H:109
char * the_recv_data
Definition: AMReX_FabArray.H:114
Vector< MPI_Request > recv_reqs
Definition: AMReX_FabArray.H:119
Vector< char * > recv_data
Definition: AMReX_FabArray.H:117
Vector< MPI_Status > recv_stat
Definition: AMReX_FabArray.H:120
int scomp
Definition: AMReX_FabArray.H:110
Vector< int > recv_from
Definition: AMReX_FabArray.H:116
char * the_send_data
Definition: AMReX_FabArray.H:115
Vector< MPI_Request > send_reqs
Definition: AMReX_FabArray.H:123
Vector< char * > send_data
Definition: AMReX_FabArray.H:122
Vector< std::size_t > recv_size
Definition: AMReX_FabArray.H:118
int ncomp
Definition: AMReX_FabArray.H:111
int tag
Definition: AMReX_FabArray.H:124
parallel copy or add
Definition: AMReX_FabArrayBase.H:536
Used by a bunch of routines when communicating via MPI.
Definition: AMReX_FabArrayBase.H:194
Box dbox
Definition: AMReX_FabArrayBase.H:195
int dstIndex
Definition: AMReX_FabArrayBase.H:197
FillBoundary.
Definition: AMReX_FabArrayBase.H:487
void recordBuild() noexcept
Definition: AMReX_FabArrayBase.H:701
Definition: AMReX_FabArray.H:346
FAB value_type
Definition: AMReX_FabArray.H:347
for shared memory
Definition: AMReX_FabArray.H:1290
ShMem(ShMem &&rhs) noexcept
Definition: AMReX_FabArray.H:1304
Long n_values
Definition: AMReX_FabArray.H:1331
Long n_points
Definition: AMReX_FabArray.H:1332
bool alloc
Definition: AMReX_FabArray.H:1330
ShMem(const ShMem &)=delete
ShMem & operator=(ShMem &&rhs) noexcept
Definition: AMReX_FabArray.H:1315
Definition: AMReX_FabFactory.H:27
FabInfo & SetArena(Arena *ar) noexcept
Definition: AMReX_FabFactory.H:42
FabInfo & SetShared(bool s) noexcept
Definition: AMReX_FabFactory.H:37
FabInfo & SetAlloc(bool a) noexcept
Definition: AMReX_FabFactory.H:32
Definition: AMReX_TypeTraits.H:18
FabArray memory allocation information.
Definition: AMReX_FabArray.H:65
MFInfo & SetTag(T &&t, Ts &&... ts) noexcept
Definition: AMReX_FabArray.H:91
Arena * arena
Definition: AMReX_FabArray.H:69
bool alloc
Definition: AMReX_FabArray.H:67
MFInfo & SetTag() noexcept
Definition: AMReX_FabArray.H:78
MFInfo & SetAllocSingleChunk(bool a) noexcept
Definition: AMReX_FabArray.H:74
MFInfo & SetArena(Arena *ar) noexcept
Definition: AMReX_FabArray.H:76
MFInfo & SetAlloc(bool a) noexcept
Definition: AMReX_FabArray.H:72
bool alloc_single_chunk
Definition: AMReX_FabArray.H:68
Vector< std::string > tags
Definition: AMReX_FabArray.H:70
MFInfo & SetTag(const char *t) noexcept
Definition: AMReX_FabArray.H:80
MFInfo & SetTag(const std::string &t) noexcept
Definition: AMReX_FabArray.H:85
Definition: AMReX_FabArray.H:151
Array4< T > const *AMREX_RESTRICT dp
Definition: AMReX_FabArray.H:165
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE Array4< T > const & operator[](int li) const noexcept
Definition: AMReX_FabArray.H:153
Array4< T > const *AMREX_RESTRICT hp
Definition: AMReX_FabArray.H:167
Definition: AMReX_FabArray.H:130
int actual_n_rcvs
Definition: AMReX_FabArray.H:136
Vector< std::size_t > recv_size
Definition: AMReX_FabArray.H:143
int DC
Definition: AMReX_FabArray.H:137
Vector< MPI_Request > send_reqs
Definition: AMReX_FabArray.H:145
int tag
Definition: AMReX_FabArray.H:135
const FabArray< FAB > * src
Definition: AMReX_FabArray.H:133
char * the_recv_data
Definition: AMReX_FabArray.H:139
FabArrayBase::CpOp op
Definition: AMReX_FabArray.H:134
Vector< MPI_Request > recv_reqs
Definition: AMReX_FabArray.H:144
char * the_send_data
Definition: AMReX_FabArray.H:140
const FabArrayBase::CPC * cpc
Definition: AMReX_FabArray.H:132
Vector< int > recv_from
Definition: AMReX_FabArray.H:141
int NC
Definition: AMReX_FabArray.H:137
int SC
Definition: AMReX_FabArray.H:137
Vector< char * > recv_data
Definition: AMReX_FabArray.H:142
const team_t & get() const
Definition: AMReX_ParallelDescriptor.H:185
void MemoryBarrier() const
memory fence
Definition: AMReX_ParallelDescriptor.H:157
Definition: AMReX_FabArray.H:97
char * pointer
Definition: AMReX_FabArray.H:98
void operator()(pointer p) const noexcept
Definition: AMReX_FabArray.H:99
Struct for holding types.
Definition: AMReX_TypeList.H:12