From 2f13514ba7353e986fba0054170ce331e1a30ba1 Mon Sep 17 00:00:00 2001 From: WeiqunZhang Date: Mon, 13 May 2024 16:23:38 +0000 Subject: [PATCH] =?UTF-8?q?Deploying=20to=20main=20from=20@=20AMReX-Codes/?= =?UTF-8?q?amrex@8eff86d32d17bfbfc6b6bf1091a45bdeb6dd5c86=20=F0=9F=9A=80?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- .../amrex.pdf | Bin 8751143 -> 8751143 bytes .../doxygen/AMReX__FillPatcher_8H_source.html | 4 +- .../doxygen/AMReX__PhysBCFunct_8H_source.html | 949 +++++++++--------- .../doxygen/AMReX__FillPatcher_8H.xml | 4 +- .../doxygen/AMReX__PhysBCFunct_8H.xml | 932 ++++++++--------- .../doxygen/classamrex_1_1GpuBndryFuncFab.xml | 6 +- .../doxygen/classamrex_1_1PhysBCFunct.xml | 12 +- 7 files changed, 957 insertions(+), 950 deletions(-) diff --git a/amrex/docs_html/_downloads/008eb6dbfab802633dff40122ece848c/amrex.pdf b/amrex/docs_html/_downloads/008eb6dbfab802633dff40122ece848c/amrex.pdf index eb8c420c96d507ce04d479675d84804aee0656e2..2c5c8d7a9bfa8b5670485eac0a53f0079799b00f 100644 GIT binary patch delta 685 zcmajaNlzL90EXezidMv3E7}FzEw+^zhGCex6-KOzift8zQoD%azLdJ}mU!@DVj>5< zc=Fbmc;J6{=)H#~{slu!`WkS2Wd5jIG$LxuwioVfk@1`TLL6PnS2w{XD?4XyB? z4e!v74s^ncE_9;@KJ?-}KHwwz&<{TbFo*zx_=FHXV+ddH72oh3!w6#pqljP(R!Ux}m|yD{l(f-H*hn?THXmx%HOW(P+f1SmS>iGc1Ya INbI-iKdu8hX#fBK delta 685 zcmajaHBSQp0EXeSP{QfM585O8R6{=B#T3BF(4RxqT0~*nUX0)IcZD@xb9q2?C zy3vDP^r0UE7{m~UF@jNyVH^{f#1y76gIUbMfq5)o5ldLc3Rba(b!=c0TiC`9cCm+j z9N-W~IK~N1afWkT;1XAEkLvu^`iY`oR9s(!&V;{oCk;8@HFim`N`h=FHJL23C|Q&` ssmve?##&pAVpT+oSgAArI~fTsX_;A>+j8hk2G?6fklwxTa6W2(01uftKmY&$ diff --git a/amrex/docs_html/doxygen/AMReX__FillPatcher_8H_source.html b/amrex/docs_html/doxygen/AMReX__FillPatcher_8H_source.html index becbabcc67..f4dee6565e 100644 --- a/amrex/docs_html/doxygen/AMReX__FillPatcher_8H_source.html +++ b/amrex/docs_html/doxygen/AMReX__FillPatcher_8H_source.html @@ -351,7 +351,7 @@
390  }
391  Gpu::streamSynchronize();
392 
-
393  cbc(*m_cf_crse_data_tmp, 0, ncomp, nghost, time, cbccomp);
+
393  cbc(*m_cf_crse_data_tmp, 0, ncomp, m_cf_crse_data_tmp->nGrowVect(), time, cbccomp);
394 
395  detail::call_interp_hook(pre_interp, *m_cf_crse_data_tmp, 0, ncomp);
396 
@@ -554,7 +554,7 @@
593  }
594  Gpu::streamSynchronize();
595 
-
596  cbc(*m_cf_crse_data_tmp, 0, m_ncomp, m_nghost, time, 0);
+
596  cbc(*m_cf_crse_data_tmp, 0, m_ncomp, m_cf_crse_data_tmp->nGrowVect(), time, 0);
597 
598  if (m_cf_fine_data == nullptr) {
599  m_cf_fine_data = std::make_unique<MF>(detail::make_mf_fine_patch<MF>(fpc, m_ncomp));
diff --git a/amrex/docs_html/doxygen/AMReX__PhysBCFunct_8H_source.html b/amrex/docs_html/doxygen/AMReX__PhysBCFunct_8H_source.html index 17d53c6cf0..3de34f5cc2 100644 --- a/amrex/docs_html/doxygen/AMReX__PhysBCFunct_8H_source.html +++ b/amrex/docs_html/doxygen/AMReX__PhysBCFunct_8H_source.html @@ -247,481 +247,484 @@
153 
154  BL_PROFILE("PhysBCFunct::()");
155 
-
157  const Box& domain = m_geom.Domain();
-
158  Box gdomain = amrex::convert(domain, mf.boxArray().ixType());
-
159  for (int i = 0; i < AMREX_SPACEDIM; ++i) {
-
160  if (m_geom.isPeriodic(i)) {
-
161  gdomain.grow(i, nghost[i]);
-
162  }
-
163  }
-
164 
-
165 #ifdef AMREX_USE_OMP
-
166 #pragma omp parallel if (Gpu::notInLaunchRegion())
-
167 #endif
-
168  {
-
169  Vector<BCRec> bcrs(ncomp);
-
170  for (MFIter mfi(mf); mfi.isValid(); ++mfi)
-
171  {
-
172  FArrayBox& dest = mf[mfi];
-
173  const Box& bx = grow(mfi.validbox(),nghost);
-
174 
-
178  if (!gdomain.contains(bx))
-
179  {
-
181  amrex::setBC(bx, domain, bccomp, 0, ncomp, m_bcr, bcrs);
-
182 
-
184  m_f(bx, dest, icomp, ncomp, m_geom, time, bcrs, 0, bccomp);
-
185  }
-
186  }
-
187  }
-
188  }
-
189 
-
190  // For backward compatibility
-
191  void FillBoundary (MultiFab& mf, int icomp, int ncomp, IntVect const& nghost,
-
192  Real time, int bccomp) {
-
193  // NOLINTNEXTLINE(readability-suspicious-call-argument)
-
194  this->operator()(mf,icomp,ncomp,nghost,time,bccomp);
-
195  }
-
196 
-
197 private:
-
198  Geometry m_geom;
-
199  Vector<BCRec> m_bcr;
-
200  F m_f;
-
201 };
-
202 
-
203 template <class F>
-
204 void
-
205 GpuBndryFuncFab<F>::operator() (Box const& bx, FArrayBox& dest,
-
206  int dcomp, int numcomp,
-
207  Geometry const& geom, Real time,
-
208  const Vector<BCRec>& bcr, int bcomp,
-
209  int orig_comp)
-
210 {
-
211 #if defined(__CUDACC__) && (__CUDACC_VER_MAJOR__ == 11) && (__CUDACC_VER_MINOR__ == 6)
-
212  // The compiler is allowed to always reject static_assert(false,..) even
-
213  // without instantiating this function. The following is to work around
-
214  // the issue. Now the compiler is not allowed to reject the code unless
-
215  // GpuBndryFuncFab::operator() is instantiated.
-
216  static_assert(std::is_same<F,int>::value,
-
217  "CUDA 11.6 bug: https://github.com/AMReX-Codes/amrex/issues/2607");
-
218 #endif
-
219  if (bx.ixType().cellCentered()) {
-
220  ccfcdoit(bx,dest,dcomp,numcomp,geom,time,bcr,bcomp,orig_comp, FilccCell{});
-
221  } else if (bx.ixType().nodeCentered()) {
-
222  nddoit(bx,dest,dcomp,numcomp,geom,time,bcr,bcomp,orig_comp);
-
223  } else {
-
224  ccfcdoit(bx,dest,dcomp,numcomp,geom,time,bcr,bcomp,orig_comp, FilfcFace{});
-
225  }
-
226 }
-
227 
-
228 template <class F>
-
229 void
-
230 GpuBndryFuncFab<F>::nddoit (Box const& bx, FArrayBox& dest,
-
231  int dcomp, int numcomp,
-
232  Geometry const& geom, Real time,
-
233  const Vector<BCRec>& bcr, int bcomp,
-
234  int orig_comp)
-
235 {
-
236  const IntVect& len = bx.length();
-
237 
-
238  Box const& domain = amrex::convert(geom.Domain(),IntVect::TheNodeVector());
-
239  Box gdomain = domain;
-
240  for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
-
241  if (geom.isPeriodic(idim)) {
-
242  gdomain.grow(idim,len[idim]);
-
243  }
-
244  }
-
245 
-
246  if (gdomain.contains(bx)) { return; }
+
156  AMREX_ASSERT(nghost.allLE(mf.nGrowVect()));
+
157 
+
159  const Box& domain = m_geom.Domain();
+
160  Box gdomain = amrex::convert(domain, mf.boxArray().ixType());
+
161  for (int i = 0; i < AMREX_SPACEDIM; ++i) {
+
162  if (m_geom.isPeriodic(i)) {
+
163  gdomain.grow(i, nghost[i]);
+
164  }
+
165  }
+
166 
+
167 #ifdef AMREX_USE_OMP
+
168 #pragma omp parallel if (Gpu::notInLaunchRegion())
+
169 #endif
+
170  {
+
171  Vector<BCRec> bcrs(ncomp);
+
172  for (MFIter mfi(mf); mfi.isValid(); ++mfi)
+
173  {
+
174  FArrayBox& dest = mf[mfi];
+
175  const Box& bx = grow(mfi.validbox(),nghost);
+
176 
+
180  if (!gdomain.contains(bx))
+
181  {
+
183  amrex::setBC(bx, domain, bccomp, 0, ncomp, m_bcr, bcrs);
+
184 
+
186  m_f(bx, dest, icomp, ncomp, m_geom, time, bcrs, 0, bccomp);
+
187  }
+
188  }
+
189  }
+
190  }
+
191 
+
192  // For backward compatibility
+
193  void FillBoundary (MultiFab& mf, int icomp, int ncomp, IntVect const& nghost,
+
194  Real time, int bccomp) {
+
195  // NOLINTNEXTLINE(readability-suspicious-call-argument)
+
196  this->operator()(mf,icomp,ncomp,nghost,time,bccomp);
+
197  }
+
198 
+
199 private:
+
200  Geometry m_geom;
+
201  Vector<BCRec> m_bcr;
+
202  F m_f;
+
203 };
+
204 
+
205 template <class F>
+
206 void
+
207 GpuBndryFuncFab<F>::operator() (Box const& bx, FArrayBox& dest,
+
208  int dcomp, int numcomp,
+
209  Geometry const& geom, Real time,
+
210  const Vector<BCRec>& bcr, int bcomp,
+
211  int orig_comp)
+
212 {
+
213 #if defined(__CUDACC__) && (__CUDACC_VER_MAJOR__ == 11) && (__CUDACC_VER_MINOR__ == 6)
+
214  // The compiler is allowed to always reject static_assert(false,..) even
+
215  // without instantiating this function. The following is to work around
+
216  // the issue. Now the compiler is not allowed to reject the code unless
+
217  // GpuBndryFuncFab::operator() is instantiated.
+
218  static_assert(std::is_same<F,int>::value,
+
219  "CUDA 11.6 bug: https://github.com/AMReX-Codes/amrex/issues/2607");
+
220 #endif
+
221  if (bx.ixType().cellCentered()) {
+
222  ccfcdoit(bx,dest,dcomp,numcomp,geom,time,bcr,bcomp,orig_comp, FilccCell{});
+
223  } else if (bx.ixType().nodeCentered()) {
+
224  nddoit(bx,dest,dcomp,numcomp,geom,time,bcr,bcomp,orig_comp);
+
225  } else {
+
226  ccfcdoit(bx,dest,dcomp,numcomp,geom,time,bcr,bcomp,orig_comp, FilfcFace{});
+
227  }
+
228 }
+
229 
+
230 template <class F>
+
231 void
+
232 GpuBndryFuncFab<F>::nddoit (Box const& bx, FArrayBox& dest,
+
233  int dcomp, int numcomp,
+
234  Geometry const& geom, Real time,
+
235  const Vector<BCRec>& bcr, int bcomp,
+
236  int orig_comp)
+
237 {
+
238  const IntVect& len = bx.length();
+
239 
+
240  Box const& domain = amrex::convert(geom.Domain(),IntVect::TheNodeVector());
+
241  Box gdomain = domain;
+
242  for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
+
243  if (geom.isPeriodic(idim)) {
+
244  gdomain.grow(idim,len[idim]);
+
245  }
+
246  }
247 
-
248  Array4<Real> const& fab = dest.array();
-
249  const auto geomdata = geom.data();
-
250 
-
251  AsyncArray<BCRec> bcr_aa(bcr.data()+bcomp, numcomp);
-
252  BCRec* bcr_p = bcr_aa.data();
-
253 
-
254  const auto f_user = m_user_f;
+
248  if (gdomain.contains(bx)) { return; }
+
249 
+
250  Array4<Real> const& fab = dest.array();
+
251  const auto geomdata = geom.data();
+
252 
+
253  AsyncArray<BCRec> bcr_aa(bcr.data()+bcomp, numcomp);
+
254  BCRec* bcr_p = bcr_aa.data();
255 
-
256  // xlo
-
257  if (!geom.isPeriodic(0) && bx.smallEnd(0) < domain.smallEnd(0)) {
-
258  Box bndry = bx;
-
259  int dxlo = domain.smallEnd(0);
-
260  bndry.setBig(0,dxlo-1);
-
261  amrex::ParallelFor(bndry, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
-
262  {
-
263  for (int n = dcomp; n < dcomp+numcomp; ++n) {
-
264  fab(i,j,k,n) = fab(dxlo,j,k,n);
-
265  }
-
266  f_user(IntVect(AMREX_D_DECL(i,j,k)), fab, dcomp, numcomp, geomdata, time,
-
267  bcr_p, 0, orig_comp);
-
268  });
-
269  }
-
270 
-
271  // xhi
-
272  if (!geom.isPeriodic(0) && bx.bigEnd(0) > domain.bigEnd(0)) {
-
273  Box bndry = bx;
-
274  int dxhi = domain.bigEnd(0);
-
275  bndry.setSmall(0,dxhi+1);
-
276  amrex::ParallelFor(bndry, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
-
277  {
-
278  for (int n = dcomp; n < dcomp+numcomp; ++n) {
-
279  fab(i,j,k,n) = fab(dxhi,j,k,n);
-
280  }
-
281  f_user(IntVect(AMREX_D_DECL(i,j,k)), fab, dcomp, numcomp, geomdata, time,
-
282  bcr_p, 0, orig_comp);
-
283  });
-
284  }
-
285 
-
286 #if (AMREX_SPACEDIM >= 2)
-
287  // ylo
-
288  if (!geom.isPeriodic(1) && bx.smallEnd(1) < domain.smallEnd(1)) {
-
289  Box bndry = bx;
-
290  int dylo = domain.smallEnd(1);
-
291  bndry.setBig(1,dylo-1);
-
292  amrex::ParallelFor(bndry, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
-
293  {
-
294  for (int n = dcomp; n < dcomp+numcomp; ++n) {
-
295  fab(i,j,k,n) = fab(i,dylo,k,n);
-
296  }
-
297  f_user(IntVect(AMREX_D_DECL(i,j,k)), fab, dcomp, numcomp, geomdata, time,
-
298  bcr_p, 0, orig_comp);
-
299  });
-
300  }
-
301 
-
302  // yhi
-
303  if (!geom.isPeriodic(1) && bx.bigEnd(1) > domain.bigEnd(1)) {
-
304  Box bndry = bx;
-
305  int dyhi = domain.bigEnd(1);
-
306  bndry.setSmall(1,dyhi+1);
-
307  amrex::ParallelFor(bndry, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
-
308  {
-
309  for (int n = dcomp; n < dcomp+numcomp; ++n) {
-
310  fab(i,j,k,n) = fab(i,dyhi,k,n);
-
311  }
-
312  f_user(IntVect(AMREX_D_DECL(i,j,k)), fab, dcomp, numcomp, geomdata, time,
-
313  bcr_p, 0, orig_comp);
-
314  });
-
315  }
-
316 #endif
-
317 
-
318 #if (AMREX_SPACEDIM == 3)
-
319  // zlo
-
320  if (!geom.isPeriodic(2) && bx.smallEnd(2) < domain.smallEnd(2)) {
-
321  Box bndry = bx;
-
322  int dzlo = domain.smallEnd(2);
-
323  bndry.setBig(2,dzlo-1);
-
324  amrex::ParallelFor(bndry, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
-
325  {
-
326  for (int n = dcomp; n < dcomp+numcomp; ++n) {
-
327  fab(i,j,k,n) = fab(i,j,dzlo,n);
-
328  }
-
329  f_user(IntVect(AMREX_D_DECL(i,j,k)), fab, dcomp, numcomp, geomdata, time,
-
330  bcr_p, 0, orig_comp);
-
331  });
-
332  }
-
333 
-
334  // zhi
-
335  if (!geom.isPeriodic(2) && bx.bigEnd(2) > domain.bigEnd(2)) {
-
336  Box bndry = bx;
-
337  int dzhi = domain.bigEnd(2);
-
338  bndry.setSmall(2,dzhi+1);
-
339  amrex::ParallelFor(bndry, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
-
340  {
-
341  for (int n = dcomp; n < dcomp+numcomp; ++n) {
-
342  fab(i,j,k,n) = fab(i,j,dzhi,n);
-
343  }
-
344  f_user(IntVect(AMREX_D_DECL(i,j,k)), fab, dcomp, numcomp, geomdata, time,
-
345  bcr_p, 0, orig_comp);
-
346  });
-
347  }
-
348 #endif
-
349 }
-
350 
-
351 template <class F>
-
352 template <class FF>
-
353 void
-
354 GpuBndryFuncFab<F>::ccfcdoit (Box const& bx, FArrayBox& dest,
-
355  int dcomp, int numcomp,
-
356  Geometry const& geom, Real time,
-
357  const Vector<BCRec>& bcr, int bcomp,
-
358  int orig_comp, FF const& fillfunc)
-
359 {
-
360  const IntVect& len = bx.length();
-
361 
-
362  IndexType idxType = bx.ixType();
-
363  Box const& domain = amrex::convert(geom.Domain(),idxType);
-
364  Box gdomain = domain;
-
365  for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
-
366  if (geom.isPeriodic(idim)) {
-
367  gdomain.grow(idim,len[idim]);
-
368  }
-
369  }
-
370 
-
371  if (gdomain.contains(bx)) { return; }
+
256  const auto f_user = m_user_f;
+
257 
+
258  // xlo
+
259  if (!geom.isPeriodic(0) && bx.smallEnd(0) < domain.smallEnd(0)) {
+
260  Box bndry = bx;
+
261  int dxlo = domain.smallEnd(0);
+
262  bndry.setBig(0,dxlo-1);
+
263  amrex::ParallelFor(bndry, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
+
264  {
+
265  for (int n = dcomp; n < dcomp+numcomp; ++n) {
+
266  fab(i,j,k,n) = fab(dxlo,j,k,n);
+
267  }
+
268  f_user(IntVect(AMREX_D_DECL(i,j,k)), fab, dcomp, numcomp, geomdata, time,
+
269  bcr_p, 0, orig_comp);
+
270  });
+
271  }
+
272 
+
273  // xhi
+
274  if (!geom.isPeriodic(0) && bx.bigEnd(0) > domain.bigEnd(0)) {
+
275  Box bndry = bx;
+
276  int dxhi = domain.bigEnd(0);
+
277  bndry.setSmall(0,dxhi+1);
+
278  amrex::ParallelFor(bndry, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
+
279  {
+
280  for (int n = dcomp; n < dcomp+numcomp; ++n) {
+
281  fab(i,j,k,n) = fab(dxhi,j,k,n);
+
282  }
+
283  f_user(IntVect(AMREX_D_DECL(i,j,k)), fab, dcomp, numcomp, geomdata, time,
+
284  bcr_p, 0, orig_comp);
+
285  });
+
286  }
+
287 
+
288 #if (AMREX_SPACEDIM >= 2)
+
289  // ylo
+
290  if (!geom.isPeriodic(1) && bx.smallEnd(1) < domain.smallEnd(1)) {
+
291  Box bndry = bx;
+
292  int dylo = domain.smallEnd(1);
+
293  bndry.setBig(1,dylo-1);
+
294  amrex::ParallelFor(bndry, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
+
295  {
+
296  for (int n = dcomp; n < dcomp+numcomp; ++n) {
+
297  fab(i,j,k,n) = fab(i,dylo,k,n);
+
298  }
+
299  f_user(IntVect(AMREX_D_DECL(i,j,k)), fab, dcomp, numcomp, geomdata, time,
+
300  bcr_p, 0, orig_comp);
+
301  });
+
302  }
+
303 
+
304  // yhi
+
305  if (!geom.isPeriodic(1) && bx.bigEnd(1) > domain.bigEnd(1)) {
+
306  Box bndry = bx;
+
307  int dyhi = domain.bigEnd(1);
+
308  bndry.setSmall(1,dyhi+1);
+
309  amrex::ParallelFor(bndry, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
+
310  {
+
311  for (int n = dcomp; n < dcomp+numcomp; ++n) {
+
312  fab(i,j,k,n) = fab(i,dyhi,k,n);
+
313  }
+
314  f_user(IntVect(AMREX_D_DECL(i,j,k)), fab, dcomp, numcomp, geomdata, time,
+
315  bcr_p, 0, orig_comp);
+
316  });
+
317  }
+
318 #endif
+
319 
+
320 #if (AMREX_SPACEDIM == 3)
+
321  // zlo
+
322  if (!geom.isPeriodic(2) && bx.smallEnd(2) < domain.smallEnd(2)) {
+
323  Box bndry = bx;
+
324  int dzlo = domain.smallEnd(2);
+
325  bndry.setBig(2,dzlo-1);
+
326  amrex::ParallelFor(bndry, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
+
327  {
+
328  for (int n = dcomp; n < dcomp+numcomp; ++n) {
+
329  fab(i,j,k,n) = fab(i,j,dzlo,n);
+
330  }
+
331  f_user(IntVect(AMREX_D_DECL(i,j,k)), fab, dcomp, numcomp, geomdata, time,
+
332  bcr_p, 0, orig_comp);
+
333  });
+
334  }
+
335 
+
336  // zhi
+
337  if (!geom.isPeriodic(2) && bx.bigEnd(2) > domain.bigEnd(2)) {
+
338  Box bndry = bx;
+
339  int dzhi = domain.bigEnd(2);
+
340  bndry.setSmall(2,dzhi+1);
+
341  amrex::ParallelFor(bndry, [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
+
342  {
+
343  for (int n = dcomp; n < dcomp+numcomp; ++n) {
+
344  fab(i,j,k,n) = fab(i,j,dzhi,n);
+
345  }
+
346  f_user(IntVect(AMREX_D_DECL(i,j,k)), fab, dcomp, numcomp, geomdata, time,
+
347  bcr_p, 0, orig_comp);
+
348  });
+
349  }
+
350 #endif
+
351 }
+
352 
+
353 template <class F>
+
354 template <class FF>
+
355 void
+
356 GpuBndryFuncFab<F>::ccfcdoit (Box const& bx, FArrayBox& dest,
+
357  int dcomp, int numcomp,
+
358  Geometry const& geom, Real time,
+
359  const Vector<BCRec>& bcr, int bcomp,
+
360  int orig_comp, FF const& fillfunc)
+
361 {
+
362  const IntVect& len = bx.length();
+
363 
+
364  IndexType idxType = bx.ixType();
+
365  Box const& domain = amrex::convert(geom.Domain(),idxType);
+
366  Box gdomain = domain;
+
367  for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
+
368  if (geom.isPeriodic(idim)) {
+
369  gdomain.grow(idim,len[idim]);
+
370  }
+
371  }
372 
-
373  Array4<Real> const& fab = dest.array();
-
374  const auto geomdata = geom.data();
-
375 
-
376 #ifdef AMREX_USE_GPU
-
377  AsyncArray<BCRec> bcr_aa(bcr.data()+bcomp, numcomp);
-
378  BCRec* bcr_p = bcr_aa.data();
-
379 
-
380  const auto f_user = m_user_f;
+
373  if (gdomain.contains(bx)) { return; }
+
374 
+
375  Array4<Real> const& fab = dest.array();
+
376  const auto geomdata = geom.data();
+
377 
+
378 #ifdef AMREX_USE_GPU
+
379  AsyncArray<BCRec> bcr_aa(bcr.data()+bcomp, numcomp);
+
380  BCRec* bcr_p = bcr_aa.data();
381 
-
382  // fill on the faces first
-
383  {
-
384  Array<Box,2*AMREX_SPACEDIM> dom_face_boxes
-
385  = { AMREX_D_DECL(amrex::convert(amrex::adjCellLo(gdomain, 0, len[0]),idxType),
-
386  amrex::convert(amrex::adjCellLo(gdomain, 1, len[1]),idxType),
-
387  amrex::convert(amrex::adjCellLo(gdomain, 2, len[2]),idxType)),
-
388  AMREX_D_DECL(amrex::convert(amrex::adjCellHi(gdomain, 0, len[0]),idxType),
-
389  amrex::convert(amrex::adjCellHi(gdomain, 1, len[1]),idxType),
-
390  amrex::convert(amrex::adjCellHi(gdomain, 2, len[2]),idxType)) };
-
391 
-
392  Vector<Box> face_boxes;
-
393  for (const Box& b : dom_face_boxes) {
-
394  Box tmp = b & bx;
-
395  if (tmp.ok()) { face_boxes.push_back(tmp); }
-
396  }
-
397  const int n_face_boxes = face_boxes.size();
-
398  if (n_face_boxes == 1) {
-
399  amrex::ParallelFor(face_boxes[0],
-
400  [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
-
401  {
-
402  AMREX_D_PICK(amrex::ignore_unused(j,k),amrex::ignore_unused(k),(void)0);
-
403  IntVect const idx(AMREX_D_DECL(i,j,k));
-
404  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
-
405  f_user(idx, fab, dcomp, numcomp, geomdata, time,
-
406  bcr_p, 0, orig_comp);
-
407  });
-
408  } else if (n_face_boxes > 1) {
-
409  AsyncArray<Box> face_boxes_aa(face_boxes.data(), n_face_boxes);
-
410  Box* boxes_p = face_boxes_aa.data();
-
411  Long ncounts = 0;
-
412  for (const auto& b : face_boxes) {
-
413  ncounts += b.numPts();
-
414  }
-
415  amrex::ParallelFor(ncounts,
-
416  [=] AMREX_GPU_DEVICE (Long icount) noexcept
-
417  {
-
418  const auto& idx = getCell(boxes_p, n_face_boxes, icount);
-
419  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
-
420  f_user(idx, fab, dcomp, numcomp, geomdata, time,
-
421  bcr_p, 0, orig_comp);
-
422  });
-
423  }
-
424  }
-
425 
-
426 #if (AMREX_SPACEDIM >= 2)
-
427  // fill on the box edges
-
428  {
-
429 #if (AMREX_SPACEDIM == 2)
-
430  Array<Box,4> dom_edge_boxes
-
431  = {{ amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
-
432  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType),
-
433  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
-
434  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType) }};
-
435 #else
-
436  Array<Box,12> dom_edge_boxes
-
437  = {{ amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
-
438  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType),
-
439  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
-
440  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType),
-
441  //
-
442  amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType),
-
443  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType),
-
444  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType),
-
445  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType),
-
446  //
-
447  amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType),
-
448  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType),
-
449  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType),
-
450  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType) }};
-
451 #endif
-
452  Vector<Box> edge_boxes;
-
453  for (const Box& b : dom_edge_boxes) {
-
454  Box tmp = b & bx;
-
455  if (tmp.ok()) { edge_boxes.push_back(tmp); }
-
456  }
-
457  const int n_edge_boxes = edge_boxes.size();
-
458  if (n_edge_boxes == 1) {
-
459  amrex::ParallelFor(edge_boxes[0],
-
460  [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
-
461  {
-
462  AMREX_D_PICK(amrex::ignore_unused(j,k),amrex::ignore_unused(k),(void)0);
-
463  IntVect const idx(AMREX_D_DECL(i,j,k));
-
464  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
-
465  f_user(idx, fab, dcomp, numcomp, geomdata, time,
-
466  bcr_p, 0, orig_comp);
-
467  });
-
468  } else if (n_edge_boxes > 1) {
-
469  AsyncArray<Box> edge_boxes_aa(edge_boxes.data(), n_edge_boxes);
-
470  Box* boxes_p = edge_boxes_aa.data();
-
471  Long ncounts = 0;
-
472  for (const auto& b : edge_boxes) {
-
473  ncounts += b.numPts();
-
474  }
-
475  amrex::ParallelFor(ncounts,
-
476  [=] AMREX_GPU_DEVICE (Long icount) noexcept
-
477  {
-
478  const auto& idx = getCell(boxes_p, n_edge_boxes, icount);
-
479  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
-
480  f_user(idx, fab, dcomp, numcomp, geomdata, time,
-
481  bcr_p, 0, orig_comp);
-
482  });
-
483  }
-
484  }
-
485 #endif
-
486 
-
487 #if (AMREX_SPACEDIM == 3)
-
488  // fill on corners
-
489  {
-
490  Array<Box,8> dom_corner_boxes
-
491  = {{ amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
492  amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
493  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
494  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
495  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
496  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
497  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
498  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType) }};
-
499 
-
500  Vector<Box> corner_boxes;
-
501  for (const Box& b : dom_corner_boxes) {
-
502  Box tmp = b & bx;
-
503  if (tmp.ok()) { corner_boxes.push_back(tmp); }
-
504  }
-
505  const int n_corner_boxes = corner_boxes.size();
-
506  if (n_corner_boxes == 1) {
-
507  amrex::ParallelFor(corner_boxes[0],
-
508  [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
-
509  {
-
510  IntVect const idx(AMREX_D_DECL(i,j,k));
-
511  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
-
512  f_user(idx, fab, dcomp, numcomp, geomdata, time,
-
513  bcr_p, 0, orig_comp);
-
514  });
-
515  } else if (n_corner_boxes > 1) {
-
516  AsyncArray<Box> corner_boxes_aa(corner_boxes.data(), n_corner_boxes);
-
517  Box* boxes_p = corner_boxes_aa.data();
-
518  Long ncounts = 0;
-
519  for (const auto& b : corner_boxes) {
-
520  ncounts += b.numPts();
-
521  }
-
522  amrex::ParallelFor(ncounts,
-
523  [=] AMREX_GPU_DEVICE (Long icount) noexcept
-
524  {
-
525  const auto& idx = getCell(boxes_p, n_corner_boxes, icount);
-
526  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
-
527  f_user(idx, fab, dcomp, numcomp, geomdata, time,
-
528  bcr_p, 0, orig_comp);
-
529  });
-
530  }
-
531  }
-
532 #endif
-
533 
-
534 #else
-
535  BCRec const* bcr_p = bcr.data()+bcomp;
-
536 
-
537  const auto& f_user = m_user_f;
+
382  const auto f_user = m_user_f;
+
383 
+
384  // fill on the faces first
+
385  {
+
386  Array<Box,2*AMREX_SPACEDIM> dom_face_boxes
+
387  = { AMREX_D_DECL(amrex::convert(amrex::adjCellLo(gdomain, 0, len[0]),idxType),
+
388  amrex::convert(amrex::adjCellLo(gdomain, 1, len[1]),idxType),
+
389  amrex::convert(amrex::adjCellLo(gdomain, 2, len[2]),idxType)),
+
390  AMREX_D_DECL(amrex::convert(amrex::adjCellHi(gdomain, 0, len[0]),idxType),
+
391  amrex::convert(amrex::adjCellHi(gdomain, 1, len[1]),idxType),
+
392  amrex::convert(amrex::adjCellHi(gdomain, 2, len[2]),idxType)) };
+
393 
+
394  Vector<Box> face_boxes;
+
395  for (const Box& b : dom_face_boxes) {
+
396  Box tmp = b & bx;
+
397  if (tmp.ok()) { face_boxes.push_back(tmp); }
+
398  }
+
399  const int n_face_boxes = face_boxes.size();
+
400  if (n_face_boxes == 1) {
+
401  amrex::ParallelFor(face_boxes[0],
+
402  [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
+
403  {
+
404  AMREX_D_PICK(amrex::ignore_unused(j,k),amrex::ignore_unused(k),(void)0);
+
405  IntVect const idx(AMREX_D_DECL(i,j,k));
+
406  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
+
407  f_user(idx, fab, dcomp, numcomp, geomdata, time,
+
408  bcr_p, 0, orig_comp);
+
409  });
+
410  } else if (n_face_boxes > 1) {
+
411  AsyncArray<Box> face_boxes_aa(face_boxes.data(), n_face_boxes);
+
412  Box* boxes_p = face_boxes_aa.data();
+
413  Long ncounts = 0;
+
414  for (const auto& b : face_boxes) {
+
415  ncounts += b.numPts();
+
416  }
+
417  amrex::ParallelFor(ncounts,
+
418  [=] AMREX_GPU_DEVICE (Long icount) noexcept
+
419  {
+
420  const auto& idx = getCell(boxes_p, n_face_boxes, icount);
+
421  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
+
422  f_user(idx, fab, dcomp, numcomp, geomdata, time,
+
423  bcr_p, 0, orig_comp);
+
424  });
+
425  }
+
426  }
+
427 
+
428 #if (AMREX_SPACEDIM >= 2)
+
429  // fill on the box edges
+
430  {
+
431 #if (AMREX_SPACEDIM == 2)
+
432  Array<Box,4> dom_edge_boxes
+
433  = {{ amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
+
434  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType),
+
435  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
+
436  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType) }};
+
437 #else
+
438  Array<Box,12> dom_edge_boxes
+
439  = {{ amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
+
440  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType),
+
441  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
+
442  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType),
+
443  //
+
444  amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType),
+
445  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType),
+
446  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType),
+
447  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType),
+
448  //
+
449  amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType),
+
450  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType),
+
451  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType),
+
452  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType) }};
+
453 #endif
+
454  Vector<Box> edge_boxes;
+
455  for (const Box& b : dom_edge_boxes) {
+
456  Box tmp = b & bx;
+
457  if (tmp.ok()) { edge_boxes.push_back(tmp); }
+
458  }
+
459  const int n_edge_boxes = edge_boxes.size();
+
460  if (n_edge_boxes == 1) {
+
461  amrex::ParallelFor(edge_boxes[0],
+
462  [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
+
463  {
+
464  AMREX_D_PICK(amrex::ignore_unused(j,k),amrex::ignore_unused(k),(void)0);
+
465  IntVect const idx(AMREX_D_DECL(i,j,k));
+
466  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
+
467  f_user(idx, fab, dcomp, numcomp, geomdata, time,
+
468  bcr_p, 0, orig_comp);
+
469  });
+
470  } else if (n_edge_boxes > 1) {
+
471  AsyncArray<Box> edge_boxes_aa(edge_boxes.data(), n_edge_boxes);
+
472  Box* boxes_p = edge_boxes_aa.data();
+
473  Long ncounts = 0;
+
474  for (const auto& b : edge_boxes) {
+
475  ncounts += b.numPts();
+
476  }
+
477  amrex::ParallelFor(ncounts,
+
478  [=] AMREX_GPU_DEVICE (Long icount) noexcept
+
479  {
+
480  const auto& idx = getCell(boxes_p, n_edge_boxes, icount);
+
481  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
+
482  f_user(idx, fab, dcomp, numcomp, geomdata, time,
+
483  bcr_p, 0, orig_comp);
+
484  });
+
485  }
+
486  }
+
487 #endif
+
488 
+
489 #if (AMREX_SPACEDIM == 3)
+
490  // fill on corners
+
491  {
+
492  Array<Box,8> dom_corner_boxes
+
493  = {{ amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
494  amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
495  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
496  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
497  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
498  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
499  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
500  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType) }};
+
501 
+
502  Vector<Box> corner_boxes;
+
503  for (const Box& b : dom_corner_boxes) {
+
504  Box tmp = b & bx;
+
505  if (tmp.ok()) { corner_boxes.push_back(tmp); }
+
506  }
+
507  const int n_corner_boxes = corner_boxes.size();
+
508  if (n_corner_boxes == 1) {
+
509  amrex::ParallelFor(corner_boxes[0],
+
510  [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
+
511  {
+
512  IntVect const idx(AMREX_D_DECL(i,j,k));
+
513  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
+
514  f_user(idx, fab, dcomp, numcomp, geomdata, time,
+
515  bcr_p, 0, orig_comp);
+
516  });
+
517  } else if (n_corner_boxes > 1) {
+
518  AsyncArray<Box> corner_boxes_aa(corner_boxes.data(), n_corner_boxes);
+
519  Box* boxes_p = corner_boxes_aa.data();
+
520  Long ncounts = 0;
+
521  for (const auto& b : corner_boxes) {
+
522  ncounts += b.numPts();
+
523  }
+
524  amrex::ParallelFor(ncounts,
+
525  [=] AMREX_GPU_DEVICE (Long icount) noexcept
+
526  {
+
527  const auto& idx = getCell(boxes_p, n_corner_boxes, icount);
+
528  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
+
529  f_user(idx, fab, dcomp, numcomp, geomdata, time,
+
530  bcr_p, 0, orig_comp);
+
531  });
+
532  }
+
533  }
+
534 #endif
+
535 
+
536 #else
+
537  BCRec const* bcr_p = bcr.data()+bcomp;
538 
-
539  // fill on the box faces first
-
540  {
-
541  Array<Box,2*AMREX_SPACEDIM> dom_face_boxes
-
542  = {{ AMREX_D_DECL(amrex::convert(amrex::adjCellLo(gdomain, 0, len[0]),idxType),
-
543  amrex::convert(amrex::adjCellLo(gdomain, 1, len[1]),idxType),
-
544  amrex::convert(amrex::adjCellLo(gdomain, 2, len[2]),idxType)),
-
545  AMREX_D_DECL(amrex::convert(amrex::adjCellHi(gdomain, 0, len[0]),idxType),
-
546  amrex::convert(amrex::adjCellHi(gdomain, 1, len[1]),idxType),
-
547  amrex::convert(amrex::adjCellHi(gdomain, 2, len[2]),idxType)) }};
-
548 
-
549  for (const Box& b : dom_face_boxes) {
-
550  Box tmp = b & bx;
-
551  amrex::For(tmp, [=] (int i, int j, int k) noexcept
-
552  {
-
553  amrex::ignore_unused(j,k);
-
554  IntVect const idx(AMREX_D_DECL(i,j,k));
-
555  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
-
556  f_user(idx, fab, dcomp, numcomp, geomdata, time,
-
557  bcr_p, 0, orig_comp);
-
558  });
-
559  }
-
560  }
-
561 
-
562 #if (AMREX_SPACEDIM >= 2)
-
563  // fill on the box edges
-
564  {
-
565 #if (AMREX_SPACEDIM == 2)
-
566  Array<Box,4> dom_edge_boxes
-
567  = {{ amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
-
568  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType),
-
569  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
-
570  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType) }};
-
571 #else
-
572  Array<Box,12> dom_edge_boxes
-
573  = {{ amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
-
574  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType),
-
575  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
-
576  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType),
-
577  //
-
578  amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType),
-
579  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType),
-
580  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType),
-
581  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType),
-
582  //
-
583  amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType),
-
584  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType),
-
585  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType),
-
586  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType) }};
-
587 #endif
-
588 
-
589  for (const Box& b : dom_edge_boxes) {
-
590  Box tmp = b & bx;
-
591  amrex::For(tmp, [=] (int i, int j, int k) noexcept
-
592  {
-
593  amrex::ignore_unused(j,k);
-
594  IntVect const idx(AMREX_D_DECL(i,j,k));
-
595  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
-
596  f_user(idx, fab, dcomp, numcomp, geomdata, time,
-
597  bcr_p, 0, orig_comp);
-
598  });
-
599  }
-
600  }
-
601 #endif
-
602 
-
603 #if (AMREX_SPACEDIM == 3)
-
604  // fill on box corners
-
605  {
-
606  Array<Box,8> dom_corner_boxes
-
607  = {{ amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
608  amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
609  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
610  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
611  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
612  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
613  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
-
614  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType) }};
-
615 
-
616  for (const Box& b : dom_corner_boxes) {
-
617  Box tmp = b & bx;
-
618  amrex::For(tmp, [=] (int i, int j, int k) noexcept
-
619  {
-
620  IntVect const idx(AMREX_D_DECL(i,j,k));
-
621  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
-
622  f_user(idx, fab, dcomp, numcomp, geomdata, time,
-
623  bcr_p, 0, orig_comp);
-
624  });
-
625  }
-
626  }
-
627 #endif
-
628 
+
539  const auto& f_user = m_user_f;
+
540 
+
541  // fill on the box faces first
+
542  {
+
543  Array<Box,2*AMREX_SPACEDIM> dom_face_boxes
+
544  = {{ AMREX_D_DECL(amrex::convert(amrex::adjCellLo(gdomain, 0, len[0]),idxType),
+
545  amrex::convert(amrex::adjCellLo(gdomain, 1, len[1]),idxType),
+
546  amrex::convert(amrex::adjCellLo(gdomain, 2, len[2]),idxType)),
+
547  AMREX_D_DECL(amrex::convert(amrex::adjCellHi(gdomain, 0, len[0]),idxType),
+
548  amrex::convert(amrex::adjCellHi(gdomain, 1, len[1]),idxType),
+
549  amrex::convert(amrex::adjCellHi(gdomain, 2, len[2]),idxType)) }};
+
550 
+
551  for (const Box& b : dom_face_boxes) {
+
552  Box tmp = b & bx;
+
553  amrex::For(tmp, [=] (int i, int j, int k) noexcept
+
554  {
+
555  amrex::ignore_unused(j,k);
+
556  IntVect const idx(AMREX_D_DECL(i,j,k));
+
557  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
+
558  f_user(idx, fab, dcomp, numcomp, geomdata, time,
+
559  bcr_p, 0, orig_comp);
+
560  });
+
561  }
+
562  }
+
563 
+
564 #if (AMREX_SPACEDIM >= 2)
+
565  // fill on the box edges
+
566  {
+
567 #if (AMREX_SPACEDIM == 2)
+
568  Array<Box,4> dom_edge_boxes
+
569  = {{ amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
+
570  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType),
+
571  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
+
572  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType) }};
+
573 #else
+
574  Array<Box,12> dom_edge_boxes
+
575  = {{ amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
+
576  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType),
+
577  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType),
+
578  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType),
+
579  //
+
580  amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType),
+
581  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType),
+
582  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType),
+
583  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType),
+
584  //
+
585  amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType),
+
586  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType),
+
587  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType),
+
588  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType) }};
+
589 #endif
+
590 
+
591  for (const Box& b : dom_edge_boxes) {
+
592  Box tmp = b & bx;
+
593  amrex::For(tmp, [=] (int i, int j, int k) noexcept
+
594  {
+
595  amrex::ignore_unused(j,k);
+
596  IntVect const idx(AMREX_D_DECL(i,j,k));
+
597  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
+
598  f_user(idx, fab, dcomp, numcomp, geomdata, time,
+
599  bcr_p, 0, orig_comp);
+
600  });
+
601  }
+
602  }
+
603 #endif
+
604 
+
605 #if (AMREX_SPACEDIM == 3)
+
606  // fill on box corners
+
607  {
+
608  Array<Box,8> dom_corner_boxes
+
609  = {{ amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
610  amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
611  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
612  amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
613  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
614  amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
615  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType),
+
616  amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType) }};
+
617 
+
618  for (const Box& b : dom_corner_boxes) {
+
619  Box tmp = b & bx;
+
620  amrex::For(tmp, [=] (int i, int j, int k) noexcept
+
621  {
+
622  IntVect const idx(AMREX_D_DECL(i,j,k));
+
623  fillfunc(idx, fab, dcomp, numcomp, domain, bcr_p, 0);
+
624  f_user(idx, fab, dcomp, numcomp, geomdata, time,
+
625  bcr_p, 0, orig_comp);
+
626  });
+
627  }
+
628  }
629 #endif
-
630 }
-
631 
+
630 
+
631 #endif
632 }
-
633 #endif
+
633 
+
634 }
+
635 #endif
AMReX_ArrayLim.H
AMReX_BCRec.H
BL_PROFILE
#define BL_PROFILE(a)
Definition: AMReX_BLProfiler.H:551
+
AMREX_ASSERT
#define AMREX_ASSERT(EX)
Definition: AMReX_BLassert.H:38
AMReX_FilCC_C.H
AMReX_FilFC_C.H
AMReX_FilND_C.H
@@ -758,6 +761,7 @@
amrex::CpuBndryFuncFab::CpuBndryFuncFab
CpuBndryFuncFab(UserFillBox a_f)
Definition: AMReX_PhysBCFunct.H:107
amrex::CpuBndryFuncFab::operator()
void operator()(Box const &bx, FArrayBox &dest, int dcomp, int numcomp, Geometry const &geom, Real time, const Vector< BCRec > &bcr, int bcomp, int orig_comp)
Definition: AMReX_PhysBCFunct.cpp:48
amrex::FArrayBox
A Fortran Array of REALs.
Definition: AMReX_FArrayBox.H:229
+
amrex::FabArrayBase::nGrowVect
IntVect nGrowVect() const noexcept
Definition: AMReX_FabArrayBase.H:79
amrex::FabArrayBase::boxArray
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
amrex::Geometry
Rectangular problem domain geometry.
Definition: AMReX_Geometry.H:73
amrex::Geometry::Domain
const Box & Domain() const noexcept
Returns our rectangular domain.
Definition: AMReX_Geometry.H:210
@@ -766,11 +770,11 @@
amrex::Geometry::isPeriodic
bool isPeriodic(int dir) const noexcept
Is the domain periodic in the specified direction?
Definition: AMReX_Geometry.H:318
amrex::GpuBndryFuncFab
Definition: AMReX_PhysBCFunct.H:64
amrex::GpuBndryFuncFab::GpuBndryFuncFab
GpuBndryFuncFab(F const &a_f)
Definition: AMReX_PhysBCFunct.H:67
-
amrex::GpuBndryFuncFab::nddoit
void nddoit(Box const &bx, FArrayBox &dest, int dcomp, int numcomp, Geometry const &geom, Real time, const Vector< BCRec > &bcr, int bcomp, int orig_comp)
Definition: AMReX_PhysBCFunct.H:230
+
amrex::GpuBndryFuncFab::nddoit
void nddoit(Box const &bx, FArrayBox &dest, int dcomp, int numcomp, Geometry const &geom, Real time, const Vector< BCRec > &bcr, int bcomp, int orig_comp)
Definition: AMReX_PhysBCFunct.H:232
amrex::GpuBndryFuncFab::GpuBndryFuncFab
GpuBndryFuncFab(F &&a_f)
Definition: AMReX_PhysBCFunct.H:68
-
amrex::GpuBndryFuncFab::ccfcdoit
void ccfcdoit(Box const &bx, FArrayBox &dest, int dcomp, int numcomp, Geometry const &geom, Real time, const Vector< BCRec > &bcr, int bcomp, int orig_comp, FF const &fillfunc)
Definition: AMReX_PhysBCFunct.H:354
-
amrex::GpuBndryFuncFab::operator()
void operator()(Box const &bx, FArrayBox &dest, int dcomp, int numcomp, Geometry const &geom, Real time, const Vector< BCRec > &bcr, int bcomp, int orig_comp)
Definition: AMReX_PhysBCFunct.H:205
+
amrex::GpuBndryFuncFab::ccfcdoit
void ccfcdoit(Box const &bx, FArrayBox &dest, int dcomp, int numcomp, Geometry const &geom, Real time, const Vector< BCRec > &bcr, int bcomp, int orig_comp, FF const &fillfunc)
Definition: AMReX_PhysBCFunct.H:356
+
amrex::GpuBndryFuncFab::operator()
void operator()(Box const &bx, FArrayBox &dest, int dcomp, int numcomp, Geometry const &geom, Real time, const Vector< BCRec > &bcr, int bcomp, int orig_comp)
Definition: AMReX_PhysBCFunct.H:207
amrex::GpuBndryFuncFab::m_user_f
F m_user_f
Definition: AMReX_PhysBCFunct.H:90
amrex::Gpu::AsyncArray
Definition: AMReX_GpuAsyncArray.H:29
amrex::Gpu::AsyncArray::data
T const * data() const noexcept
Definition: AMReX_GpuAsyncArray.H:69
@@ -779,6 +783,7 @@
amrex::IndexType::nodeCentered
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE bool nodeCentered() const noexcept
True if the IndexType is NODE based in all directions.
Definition: AMReX_IndexType.H:88
amrex::IntVect
Definition: AMReX_IntVect.H:46
amrex::IntVect::TheNodeVector
AMREX_GPU_HOST_DEVICE static constexpr AMREX_FORCE_INLINE IntVect TheNodeVector() noexcept
This static member function returns a reference to a constant IntVect object, all of whose AMREX_SPAC...
Definition: AMReX_IntVect.H:553
+
amrex::IntVect::allLE
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE bool allLE(const IntVect &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:314
amrex::MFIter
Definition: AMReX_MFIter.H:57
amrex::MFIter::isValid
bool isValid() const noexcept
Is the iterator valid i.e. is it associated with a FAB?
Definition: AMReX_MFIter.H:141
amrex::MultiFab
A collection (stored as an array) of FArrayBox objects.
Definition: AMReX_MultiFab.H:38
@@ -788,13 +793,13 @@
amrex::PhysBCFunct::define
void define(const Geometry &geom, const Vector< BCRec > &bcr, F const &f)
Definition: AMReX_PhysBCFunct.H:141
amrex::PhysBCFunct::define
void define(const Geometry &geom, const Vector< BCRec > &bcr, F &&f)
Definition: AMReX_PhysBCFunct.H:145
amrex::PhysBCFunct::PhysBCFunct
PhysBCFunct(const Geometry &geom, const Vector< BCRec > &bcr, F const &f)
Definition: AMReX_PhysBCFunct.H:133
-
amrex::PhysBCFunct::m_bcr
Vector< BCRec > m_bcr
Definition: AMReX_PhysBCFunct.H:199
+
amrex::PhysBCFunct::m_bcr
Vector< BCRec > m_bcr
Definition: AMReX_PhysBCFunct.H:201
amrex::PhysBCFunct::operator()
void operator()(MultiFab &mf, int icomp, int ncomp, IntVect const &nghost, Real time, int bccomp)
Definition: AMReX_PhysBCFunct.H:149
amrex::PhysBCFunct::PhysBCFunct
PhysBCFunct(const Geometry &geom, const Vector< BCRec > &bcr, F &&f)
Definition: AMReX_PhysBCFunct.H:137
amrex::PhysBCFunct::PhysBCFunct
PhysBCFunct()=default
-
amrex::PhysBCFunct::m_geom
Geometry m_geom
Definition: AMReX_PhysBCFunct.H:198
-
amrex::PhysBCFunct::m_f
F m_f
Definition: AMReX_PhysBCFunct.H:200
-
amrex::PhysBCFunct::FillBoundary
void FillBoundary(MultiFab &mf, int icomp, int ncomp, IntVect const &nghost, Real time, int bccomp)
Definition: AMReX_PhysBCFunct.H:191
+
amrex::PhysBCFunct::m_geom
Geometry m_geom
Definition: AMReX_PhysBCFunct.H:200
+
amrex::PhysBCFunct::m_f
F m_f
Definition: AMReX_PhysBCFunct.H:202
+
amrex::PhysBCFunct::FillBoundary
void FillBoundary(MultiFab &mf, int icomp, int ncomp, IntVect const &nghost, Real time, int bccomp)
Definition: AMReX_PhysBCFunct.H:193
amrex::Vector
This class is a thin wrapper around std::vector. Unlike vector, Vector::operator[] provides bound che...
Definition: AMReX_Vector.H:27
amrex::Vector::size
Long size() const noexcept
Definition: AMReX_Vector.H:50
amrex::SundialsUserFun::f
static int f(sunrealtype t, N_Vector y_data, N_Vector y_rhs, void *user_data)
Definition: AMReX_SundialsIntegrator.H:43
diff --git a/amrex/docs_xml/doxygen/AMReX__FillPatcher_8H.xml b/amrex/docs_xml/doxygen/AMReX__FillPatcher_8H.xml index b398915418..1b05b68511 100644 --- a/amrex/docs_xml/doxygen/AMReX__FillPatcher_8H.xml +++ b/amrex/docs_xml/doxygen/AMReX__FillPatcher_8H.xml @@ -2466,7 +2466,7 @@ } Gpu::streamSynchronize(); -cbc(*m_cf_crse_data_tmp,0,ncomp,nghost,time,cbccomp); +cbc(*m_cf_crse_data_tmp,0,ncomp,m_cf_crse_data_tmp->nGrowVect(),time,cbccomp); detail::call_interp_hook(pre_interp,*m_cf_crse_data_tmp,0,ncomp); @@ -2669,7 +2669,7 @@ } Gpu::streamSynchronize(); -cbc(*m_cf_crse_data_tmp,0,m_ncomp,m_nghost,time,0); +cbc(*m_cf_crse_data_tmp,0,m_ncomp,m_cf_crse_data_tmp->nGrowVect(),time,0); if(m_cf_fine_data==nullptr){ m_cf_fine_data=std::make_unique<MF>(detail::make_mf_fine_patch<MF>(fpc,m_ncomp)); diff --git a/amrex/docs_xml/doxygen/AMReX__PhysBCFunct_8H.xml b/amrex/docs_xml/doxygen/AMReX__PhysBCFunct_8H.xml index 3a83d867f2..b70f520bae 100644 --- a/amrex/docs_xml/doxygen/AMReX__PhysBCFunct_8H.xml +++ b/amrex/docs_xml/doxygen/AMReX__PhysBCFunct_8H.xml @@ -2347,478 +2347,480 @@ BL_PROFILE("PhysBCFunct::()"); -constBox&domain=m_geom.Domain(); -Boxgdomain=amrex::convert(domain,mf.boxArray().ixType()); -for(inti=0;i<AMREX_SPACEDIM;++i){ -if(m_geom.isPeriodic(i)){ -gdomain.grow(i,nghost[i]); -} -} - -#ifdefAMREX_USE_OMP -#pragmaompparallelif(Gpu::notInLaunchRegion()) -#endif -{ -Vector<BCRec>bcrs(ncomp); -for(MFItermfi(mf);mfi.isValid();++mfi) -{ -FArrayBox&dest=mf[mfi]; -constBox&bx=grow(mfi.validbox(),nghost); - -if(!gdomain.contains(bx)) -{ -amrex::setBC(bx,domain,bccomp,0,ncomp,m_bcr,bcrs); - -m_f(bx,dest,icomp,ncomp,m_geom,time,bcrs,0,bccomp); -} -} -} -} - -//Forbackwardcompatibility -voidFillBoundary(MultiFab&mf,inticomp,intncomp,IntVectconst&nghost, -Realtime,intbccomp){ -//NOLINTNEXTLINE(readability-suspicious-call-argument) -this->operator()(mf,icomp,ncomp,nghost,time,bccomp); -} - -private: -Geometrym_geom; -Vector<BCRec>m_bcr; -Fm_f; -}; - -template<classF> -void -GpuBndryFuncFab<F>::operator() (Boxconst&bx,FArrayBox&dest, -intdcomp,intnumcomp, -Geometryconst&geom,Realtime, -constVector<BCRec>&bcr,intbcomp, -intorig_comp) -{ -#ifdefined(__CUDACC__)&&(__CUDACC_VER_MAJOR__==11)&&(__CUDACC_VER_MINOR__==6) -//Thecompilerisallowedtoalwaysrejectstatic_assert(false,..)even -//withoutinstantiatingthisfunction.Thefollowingistoworkaround -//theissue.Nowthecompilerisnotallowedtorejectthecodeunless -//GpuBndryFuncFab::operator()isinstantiated. -static_assert(std::is_same<F,int>::value, -"CUDA11.6bug:https://github.com/AMReX-Codes/amrex/issues/2607"); -#endif -if(bx.ixType().cellCentered()){ -ccfcdoit(bx,dest,dcomp,numcomp,geom,time,bcr,bcomp,orig_comp,FilccCell{}); -}elseif(bx.ixType().nodeCentered()){ -nddoit(bx,dest,dcomp,numcomp,geom,time,bcr,bcomp,orig_comp); -}else{ -ccfcdoit(bx,dest,dcomp,numcomp,geom,time,bcr,bcomp,orig_comp,FilfcFace{}); -} -} - -template<classF> -void -GpuBndryFuncFab<F>::nddoit(Boxconst&bx,FArrayBox&dest, -intdcomp,intnumcomp, -Geometryconst&geom,Realtime, -constVector<BCRec>&bcr,intbcomp, -intorig_comp) -{ -constIntVect&len=bx.length(); - -Boxconst&domain=amrex::convert(geom.Domain(),IntVect::TheNodeVector()); -Boxgdomain=domain; -for(intidim=0;idim<AMREX_SPACEDIM;++idim){ -if(geom.isPeriodic(idim)){ -gdomain.grow(idim,len[idim]); -} -} - -if(gdomain.contains(bx)){return;} +AMREX_ASSERT(nghost.allLE(mf.nGrowVect())); + +constBox&domain=m_geom.Domain(); +Boxgdomain=amrex::convert(domain,mf.boxArray().ixType()); +for(inti=0;i<AMREX_SPACEDIM;++i){ +if(m_geom.isPeriodic(i)){ +gdomain.grow(i,nghost[i]); +} +} + +#ifdefAMREX_USE_OMP +#pragmaompparallelif(Gpu::notInLaunchRegion()) +#endif +{ +Vector<BCRec>bcrs(ncomp); +for(MFItermfi(mf);mfi.isValid();++mfi) +{ +FArrayBox&dest=mf[mfi]; +constBox&bx=grow(mfi.validbox(),nghost); + +if(!gdomain.contains(bx)) +{ +amrex::setBC(bx,domain,bccomp,0,ncomp,m_bcr,bcrs); + +m_f(bx,dest,icomp,ncomp,m_geom,time,bcrs,0,bccomp); +} +} +} +} + +//Forbackwardcompatibility +voidFillBoundary(MultiFab&mf,inticomp,intncomp,IntVectconst&nghost, +Realtime,intbccomp){ +//NOLINTNEXTLINE(readability-suspicious-call-argument) +this->operator()(mf,icomp,ncomp,nghost,time,bccomp); +} + +private: +Geometrym_geom; +Vector<BCRec>m_bcr; +Fm_f; +}; + +template<classF> +void +GpuBndryFuncFab<F>::operator() (Boxconst&bx,FArrayBox&dest, +intdcomp,intnumcomp, +Geometryconst&geom,Realtime, +constVector<BCRec>&bcr,intbcomp, +intorig_comp) +{ +#ifdefined(__CUDACC__)&&(__CUDACC_VER_MAJOR__==11)&&(__CUDACC_VER_MINOR__==6) +//Thecompilerisallowedtoalwaysrejectstatic_assert(false,..)even +//withoutinstantiatingthisfunction.Thefollowingistoworkaround +//theissue.Nowthecompilerisnotallowedtorejectthecodeunless +//GpuBndryFuncFab::operator()isinstantiated. +static_assert(std::is_same<F,int>::value, +"CUDA11.6bug:https://github.com/AMReX-Codes/amrex/issues/2607"); +#endif +if(bx.ixType().cellCentered()){ +ccfcdoit(bx,dest,dcomp,numcomp,geom,time,bcr,bcomp,orig_comp,FilccCell{}); +}elseif(bx.ixType().nodeCentered()){ +nddoit(bx,dest,dcomp,numcomp,geom,time,bcr,bcomp,orig_comp); +}else{ +ccfcdoit(bx,dest,dcomp,numcomp,geom,time,bcr,bcomp,orig_comp,FilfcFace{}); +} +} + +template<classF> +void +GpuBndryFuncFab<F>::nddoit(Boxconst&bx,FArrayBox&dest, +intdcomp,intnumcomp, +Geometryconst&geom,Realtime, +constVector<BCRec>&bcr,intbcomp, +intorig_comp) +{ +constIntVect&len=bx.length(); + +Boxconst&domain=amrex::convert(geom.Domain(),IntVect::TheNodeVector()); +Boxgdomain=domain; +for(intidim=0;idim<AMREX_SPACEDIM;++idim){ +if(geom.isPeriodic(idim)){ +gdomain.grow(idim,len[idim]); +} +} -Array4<Real>const&fab=dest.array(); -constautogeomdata=geom.data(); - -AsyncArray<BCRec>bcr_aa(bcr.data()+bcomp,numcomp); -BCRec*bcr_p=bcr_aa.data(); - -constautof_user=m_user_f; +if(gdomain.contains(bx)){return;} + +Array4<Real>const&fab=dest.array(); +constautogeomdata=geom.data(); + +AsyncArray<BCRec>bcr_aa(bcr.data()+bcomp,numcomp); +BCRec*bcr_p=bcr_aa.data(); -//xlo -if(!geom.isPeriodic(0)&&bx.smallEnd(0)<domain.smallEnd(0)){ -Boxbndry=bx; -intdxlo=domain.smallEnd(0); -bndry.setBig(0,dxlo-1); -amrex::ParallelFor(bndry,[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept -{ -for(intn=dcomp;n<dcomp+numcomp;++n){ -fab(i,j,k,n)=fab(dxlo,j,k,n); -} -f_user(IntVect(AMREX_D_DECL(i,j,k)),fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -} - -//xhi -if(!geom.isPeriodic(0)&&bx.bigEnd(0)>domain.bigEnd(0)){ -Boxbndry=bx; -intdxhi=domain.bigEnd(0); -bndry.setSmall(0,dxhi+1); -amrex::ParallelFor(bndry,[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept -{ -for(intn=dcomp;n<dcomp+numcomp;++n){ -fab(i,j,k,n)=fab(dxhi,j,k,n); -} -f_user(IntVect(AMREX_D_DECL(i,j,k)),fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -} - -#if(AMREX_SPACEDIM>=2) -//ylo -if(!geom.isPeriodic(1)&&bx.smallEnd(1)<domain.smallEnd(1)){ -Boxbndry=bx; -intdylo=domain.smallEnd(1); -bndry.setBig(1,dylo-1); -amrex::ParallelFor(bndry,[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept -{ -for(intn=dcomp;n<dcomp+numcomp;++n){ -fab(i,j,k,n)=fab(i,dylo,k,n); -} -f_user(IntVect(AMREX_D_DECL(i,j,k)),fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -} - -//yhi -if(!geom.isPeriodic(1)&&bx.bigEnd(1)>domain.bigEnd(1)){ -Boxbndry=bx; -intdyhi=domain.bigEnd(1); -bndry.setSmall(1,dyhi+1); -amrex::ParallelFor(bndry,[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept -{ -for(intn=dcomp;n<dcomp+numcomp;++n){ -fab(i,j,k,n)=fab(i,dyhi,k,n); -} -f_user(IntVect(AMREX_D_DECL(i,j,k)),fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -} -#endif - -#if(AMREX_SPACEDIM==3) -//zlo -if(!geom.isPeriodic(2)&&bx.smallEnd(2)<domain.smallEnd(2)){ -Boxbndry=bx; -intdzlo=domain.smallEnd(2); -bndry.setBig(2,dzlo-1); -amrex::ParallelFor(bndry,[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept -{ -for(intn=dcomp;n<dcomp+numcomp;++n){ -fab(i,j,k,n)=fab(i,j,dzlo,n); -} -f_user(IntVect(AMREX_D_DECL(i,j,k)),fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -} - -//zhi -if(!geom.isPeriodic(2)&&bx.bigEnd(2)>domain.bigEnd(2)){ -Boxbndry=bx; -intdzhi=domain.bigEnd(2); -bndry.setSmall(2,dzhi+1); -amrex::ParallelFor(bndry,[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept -{ -for(intn=dcomp;n<dcomp+numcomp;++n){ -fab(i,j,k,n)=fab(i,j,dzhi,n); -} -f_user(IntVect(AMREX_D_DECL(i,j,k)),fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -} -#endif -} - -template<classF> -template<classFF> -void -GpuBndryFuncFab<F>::ccfcdoit(Boxconst&bx,FArrayBox&dest, -intdcomp,intnumcomp, -Geometryconst&geom,Realtime, -constVector<BCRec>&bcr,intbcomp, -intorig_comp,FFconst&fillfunc) -{ -constIntVect&len=bx.length(); - -IndexTypeidxType=bx.ixType(); -Boxconst&domain=amrex::convert(geom.Domain(),idxType); -Boxgdomain=domain; -for(intidim=0;idim<AMREX_SPACEDIM;++idim){ -if(geom.isPeriodic(idim)){ -gdomain.grow(idim,len[idim]); -} -} - -if(gdomain.contains(bx)){return;} +constautof_user=m_user_f; + +//xlo +if(!geom.isPeriodic(0)&&bx.smallEnd(0)<domain.smallEnd(0)){ +Boxbndry=bx; +intdxlo=domain.smallEnd(0); +bndry.setBig(0,dxlo-1); +amrex::ParallelFor(bndry,[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept +{ +for(intn=dcomp;n<dcomp+numcomp;++n){ +fab(i,j,k,n)=fab(dxlo,j,k,n); +} +f_user(IntVect(AMREX_D_DECL(i,j,k)),fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +} + +//xhi +if(!geom.isPeriodic(0)&&bx.bigEnd(0)>domain.bigEnd(0)){ +Boxbndry=bx; +intdxhi=domain.bigEnd(0); +bndry.setSmall(0,dxhi+1); +amrex::ParallelFor(bndry,[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept +{ +for(intn=dcomp;n<dcomp+numcomp;++n){ +fab(i,j,k,n)=fab(dxhi,j,k,n); +} +f_user(IntVect(AMREX_D_DECL(i,j,k)),fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +} + +#if(AMREX_SPACEDIM>=2) +//ylo +if(!geom.isPeriodic(1)&&bx.smallEnd(1)<domain.smallEnd(1)){ +Boxbndry=bx; +intdylo=domain.smallEnd(1); +bndry.setBig(1,dylo-1); +amrex::ParallelFor(bndry,[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept +{ +for(intn=dcomp;n<dcomp+numcomp;++n){ +fab(i,j,k,n)=fab(i,dylo,k,n); +} +f_user(IntVect(AMREX_D_DECL(i,j,k)),fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +} + +//yhi +if(!geom.isPeriodic(1)&&bx.bigEnd(1)>domain.bigEnd(1)){ +Boxbndry=bx; +intdyhi=domain.bigEnd(1); +bndry.setSmall(1,dyhi+1); +amrex::ParallelFor(bndry,[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept +{ +for(intn=dcomp;n<dcomp+numcomp;++n){ +fab(i,j,k,n)=fab(i,dyhi,k,n); +} +f_user(IntVect(AMREX_D_DECL(i,j,k)),fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +} +#endif + +#if(AMREX_SPACEDIM==3) +//zlo +if(!geom.isPeriodic(2)&&bx.smallEnd(2)<domain.smallEnd(2)){ +Boxbndry=bx; +intdzlo=domain.smallEnd(2); +bndry.setBig(2,dzlo-1); +amrex::ParallelFor(bndry,[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept +{ +for(intn=dcomp;n<dcomp+numcomp;++n){ +fab(i,j,k,n)=fab(i,j,dzlo,n); +} +f_user(IntVect(AMREX_D_DECL(i,j,k)),fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +} + +//zhi +if(!geom.isPeriodic(2)&&bx.bigEnd(2)>domain.bigEnd(2)){ +Boxbndry=bx; +intdzhi=domain.bigEnd(2); +bndry.setSmall(2,dzhi+1); +amrex::ParallelFor(bndry,[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept +{ +for(intn=dcomp;n<dcomp+numcomp;++n){ +fab(i,j,k,n)=fab(i,j,dzhi,n); +} +f_user(IntVect(AMREX_D_DECL(i,j,k)),fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +} +#endif +} + +template<classF> +template<classFF> +void +GpuBndryFuncFab<F>::ccfcdoit(Boxconst&bx,FArrayBox&dest, +intdcomp,intnumcomp, +Geometryconst&geom,Realtime, +constVector<BCRec>&bcr,intbcomp, +intorig_comp,FFconst&fillfunc) +{ +constIntVect&len=bx.length(); + +IndexTypeidxType=bx.ixType(); +Boxconst&domain=amrex::convert(geom.Domain(),idxType); +Boxgdomain=domain; +for(intidim=0;idim<AMREX_SPACEDIM;++idim){ +if(geom.isPeriodic(idim)){ +gdomain.grow(idim,len[idim]); +} +} -Array4<Real>const&fab=dest.array(); -constautogeomdata=geom.data(); - -#ifdefAMREX_USE_GPU -AsyncArray<BCRec>bcr_aa(bcr.data()+bcomp,numcomp); -BCRec*bcr_p=bcr_aa.data(); - -constautof_user=m_user_f; +if(gdomain.contains(bx)){return;} + +Array4<Real>const&fab=dest.array(); +constautogeomdata=geom.data(); + +#ifdefAMREX_USE_GPU +AsyncArray<BCRec>bcr_aa(bcr.data()+bcomp,numcomp); +BCRec*bcr_p=bcr_aa.data(); -//fillonthefacesfirst -{ -Array<Box,2*AMREX_SPACEDIM>dom_face_boxes -={AMREX_D_DECL(amrex::convert(amrex::adjCellLo(gdomain,0,len[0]),idxType), -amrex::convert(amrex::adjCellLo(gdomain,1,len[1]),idxType), -amrex::convert(amrex::adjCellLo(gdomain,2,len[2]),idxType)), -AMREX_D_DECL(amrex::convert(amrex::adjCellHi(gdomain,0,len[0]),idxType), -amrex::convert(amrex::adjCellHi(gdomain,1,len[1]),idxType), -amrex::convert(amrex::adjCellHi(gdomain,2,len[2]),idxType))}; - -Vector<Box>face_boxes; -for(constBox&b:dom_face_boxes){ -Boxtmp=b&bx; -if(tmp.ok()){face_boxes.push_back(tmp);} -} -constintn_face_boxes=face_boxes.size(); -if(n_face_boxes==1){ -amrex::ParallelFor(face_boxes[0], -[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept -{ -AMREX_D_PICK(amrex::ignore_unused(j,k),amrex::ignore_unused(k),(void)0); -IntVectconstidx(AMREX_D_DECL(i,j,k)); -fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); -f_user(idx,fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -}elseif(n_face_boxes>1){ -AsyncArray<Box>face_boxes_aa(face_boxes.data(),n_face_boxes); -Box*boxes_p=face_boxes_aa.data(); -Longncounts=0; -for(constauto&b:face_boxes){ -ncounts+=b.numPts(); -} -amrex::ParallelFor(ncounts, -[=]AMREX_GPU_DEVICE(Longicount)noexcept -{ -constauto&idx=getCell(boxes_p,n_face_boxes,icount); -fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); -f_user(idx,fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -} -} - -#if(AMREX_SPACEDIM>=2) -//fillontheboxedges -{ -#if(AMREX_SPACEDIM==2) -Array<Box,4>dom_edge_boxes -={{amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType)}}; -#else -Array<Box,12>dom_edge_boxes -={{amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType), -// -amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType), -// -amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType)}}; -#endif -Vector<Box>edge_boxes; -for(constBox&b:dom_edge_boxes){ -Boxtmp=b&bx; -if(tmp.ok()){edge_boxes.push_back(tmp);} -} -constintn_edge_boxes=edge_boxes.size(); -if(n_edge_boxes==1){ -amrex::ParallelFor(edge_boxes[0], -[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept -{ -AMREX_D_PICK(amrex::ignore_unused(j,k),amrex::ignore_unused(k),(void)0); -IntVectconstidx(AMREX_D_DECL(i,j,k)); -fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); -f_user(idx,fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -}elseif(n_edge_boxes>1){ -AsyncArray<Box>edge_boxes_aa(edge_boxes.data(),n_edge_boxes); -Box*boxes_p=edge_boxes_aa.data(); -Longncounts=0; -for(constauto&b:edge_boxes){ -ncounts+=b.numPts(); -} -amrex::ParallelFor(ncounts, -[=]AMREX_GPU_DEVICE(Longicount)noexcept -{ -constauto&idx=getCell(boxes_p,n_edge_boxes,icount); -fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); -f_user(idx,fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -} -} -#endif - -#if(AMREX_SPACEDIM==3) -//filloncorners -{ -Array<Box,8>dom_corner_boxes -={{amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType)}}; - -Vector<Box>corner_boxes; -for(constBox&b:dom_corner_boxes){ -Boxtmp=b&bx; -if(tmp.ok()){corner_boxes.push_back(tmp);} -} -constintn_corner_boxes=corner_boxes.size(); -if(n_corner_boxes==1){ -amrex::ParallelFor(corner_boxes[0], -[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept -{ -IntVectconstidx(AMREX_D_DECL(i,j,k)); -fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); -f_user(idx,fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -}elseif(n_corner_boxes>1){ -AsyncArray<Box>corner_boxes_aa(corner_boxes.data(),n_corner_boxes); -Box*boxes_p=corner_boxes_aa.data(); -Longncounts=0; -for(constauto&b:corner_boxes){ -ncounts+=b.numPts(); -} -amrex::ParallelFor(ncounts, -[=]AMREX_GPU_DEVICE(Longicount)noexcept -{ -constauto&idx=getCell(boxes_p,n_corner_boxes,icount); -fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); -f_user(idx,fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -} -} -#endif - -#else -BCRecconst*bcr_p=bcr.data()+bcomp; - -constauto&f_user=m_user_f; +constautof_user=m_user_f; + +//fillonthefacesfirst +{ +Array<Box,2*AMREX_SPACEDIM>dom_face_boxes +={AMREX_D_DECL(amrex::convert(amrex::adjCellLo(gdomain,0,len[0]),idxType), +amrex::convert(amrex::adjCellLo(gdomain,1,len[1]),idxType), +amrex::convert(amrex::adjCellLo(gdomain,2,len[2]),idxType)), +AMREX_D_DECL(amrex::convert(amrex::adjCellHi(gdomain,0,len[0]),idxType), +amrex::convert(amrex::adjCellHi(gdomain,1,len[1]),idxType), +amrex::convert(amrex::adjCellHi(gdomain,2,len[2]),idxType))}; + +Vector<Box>face_boxes; +for(constBox&b:dom_face_boxes){ +Boxtmp=b&bx; +if(tmp.ok()){face_boxes.push_back(tmp);} +} +constintn_face_boxes=face_boxes.size(); +if(n_face_boxes==1){ +amrex::ParallelFor(face_boxes[0], +[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept +{ +AMREX_D_PICK(amrex::ignore_unused(j,k),amrex::ignore_unused(k),(void)0); +IntVectconstidx(AMREX_D_DECL(i,j,k)); +fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); +f_user(idx,fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +}elseif(n_face_boxes>1){ +AsyncArray<Box>face_boxes_aa(face_boxes.data(),n_face_boxes); +Box*boxes_p=face_boxes_aa.data(); +Longncounts=0; +for(constauto&b:face_boxes){ +ncounts+=b.numPts(); +} +amrex::ParallelFor(ncounts, +[=]AMREX_GPU_DEVICE(Longicount)noexcept +{ +constauto&idx=getCell(boxes_p,n_face_boxes,icount); +fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); +f_user(idx,fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +} +} + +#if(AMREX_SPACEDIM>=2) +//fillontheboxedges +{ +#if(AMREX_SPACEDIM==2) +Array<Box,4>dom_edge_boxes +={{amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType)}}; +#else +Array<Box,12>dom_edge_boxes +={{amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType), +// +amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType), +// +amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType)}}; +#endif +Vector<Box>edge_boxes; +for(constBox&b:dom_edge_boxes){ +Boxtmp=b&bx; +if(tmp.ok()){edge_boxes.push_back(tmp);} +} +constintn_edge_boxes=edge_boxes.size(); +if(n_edge_boxes==1){ +amrex::ParallelFor(edge_boxes[0], +[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept +{ +AMREX_D_PICK(amrex::ignore_unused(j,k),amrex::ignore_unused(k),(void)0); +IntVectconstidx(AMREX_D_DECL(i,j,k)); +fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); +f_user(idx,fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +}elseif(n_edge_boxes>1){ +AsyncArray<Box>edge_boxes_aa(edge_boxes.data(),n_edge_boxes); +Box*boxes_p=edge_boxes_aa.data(); +Longncounts=0; +for(constauto&b:edge_boxes){ +ncounts+=b.numPts(); +} +amrex::ParallelFor(ncounts, +[=]AMREX_GPU_DEVICE(Longicount)noexcept +{ +constauto&idx=getCell(boxes_p,n_edge_boxes,icount); +fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); +f_user(idx,fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +} +} +#endif + +#if(AMREX_SPACEDIM==3) +//filloncorners +{ +Array<Box,8>dom_corner_boxes +={{amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType)}}; + +Vector<Box>corner_boxes; +for(constBox&b:dom_corner_boxes){ +Boxtmp=b&bx; +if(tmp.ok()){corner_boxes.push_back(tmp);} +} +constintn_corner_boxes=corner_boxes.size(); +if(n_corner_boxes==1){ +amrex::ParallelFor(corner_boxes[0], +[=]AMREX_GPU_DEVICE(inti,intj,intk)noexcept +{ +IntVectconstidx(AMREX_D_DECL(i,j,k)); +fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); +f_user(idx,fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +}elseif(n_corner_boxes>1){ +AsyncArray<Box>corner_boxes_aa(corner_boxes.data(),n_corner_boxes); +Box*boxes_p=corner_boxes_aa.data(); +Longncounts=0; +for(constauto&b:corner_boxes){ +ncounts+=b.numPts(); +} +amrex::ParallelFor(ncounts, +[=]AMREX_GPU_DEVICE(Longicount)noexcept +{ +constauto&idx=getCell(boxes_p,n_corner_boxes,icount); +fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); +f_user(idx,fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +} +} +#endif + +#else +BCRecconst*bcr_p=bcr.data()+bcomp; -//fillontheboxfacesfirst -{ -Array<Box,2*AMREX_SPACEDIM>dom_face_boxes -={{AMREX_D_DECL(amrex::convert(amrex::adjCellLo(gdomain,0,len[0]),idxType), -amrex::convert(amrex::adjCellLo(gdomain,1,len[1]),idxType), -amrex::convert(amrex::adjCellLo(gdomain,2,len[2]),idxType)), -AMREX_D_DECL(amrex::convert(amrex::adjCellHi(gdomain,0,len[0]),idxType), -amrex::convert(amrex::adjCellHi(gdomain,1,len[1]),idxType), -amrex::convert(amrex::adjCellHi(gdomain,2,len[2]),idxType))}}; - -for(constBox&b:dom_face_boxes){ -Boxtmp=b&bx; -amrex::For(tmp,[=](inti,intj,intk)noexcept -{ -amrex::ignore_unused(j,k); -IntVectconstidx(AMREX_D_DECL(i,j,k)); -fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); -f_user(idx,fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -} -} - -#if(AMREX_SPACEDIM>=2) -//fillontheboxedges -{ -#if(AMREX_SPACEDIM==2) -Array<Box,4>dom_edge_boxes -={{amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType)}}; -#else -Array<Box,12>dom_edge_boxes -={{amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType), -// -amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType), -// -amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType)}}; -#endif - -for(constBox&b:dom_edge_boxes){ -Boxtmp=b&bx; -amrex::For(tmp,[=](inti,intj,intk)noexcept -{ -amrex::ignore_unused(j,k); -IntVectconstidx(AMREX_D_DECL(i,j,k)); -fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); -f_user(idx,fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -} -} -#endif - -#if(AMREX_SPACEDIM==3) -//fillonboxcorners -{ -Array<Box,8>dom_corner_boxes -={{amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), -amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType)}}; - -for(constBox&b:dom_corner_boxes){ -Boxtmp=b&bx; -amrex::For(tmp,[=](inti,intj,intk)noexcept -{ -IntVectconstidx(AMREX_D_DECL(i,j,k)); -fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); -f_user(idx,fab,dcomp,numcomp,geomdata,time, -bcr_p,0,orig_comp); -}); -} -} -#endif - +constauto&f_user=m_user_f; + +//fillontheboxfacesfirst +{ +Array<Box,2*AMREX_SPACEDIM>dom_face_boxes +={{AMREX_D_DECL(amrex::convert(amrex::adjCellLo(gdomain,0,len[0]),idxType), +amrex::convert(amrex::adjCellLo(gdomain,1,len[1]),idxType), +amrex::convert(amrex::adjCellLo(gdomain,2,len[2]),idxType)), +AMREX_D_DECL(amrex::convert(amrex::adjCellHi(gdomain,0,len[0]),idxType), +amrex::convert(amrex::adjCellHi(gdomain,1,len[1]),idxType), +amrex::convert(amrex::adjCellHi(gdomain,2,len[2]),idxType))}}; + +for(constBox&b:dom_face_boxes){ +Boxtmp=b&bx; +amrex::For(tmp,[=](inti,intj,intk)noexcept +{ +amrex::ignore_unused(j,k); +IntVectconstidx(AMREX_D_DECL(i,j,k)); +fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); +f_user(idx,fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +} +} + +#if(AMREX_SPACEDIM>=2) +//fillontheboxedges +{ +#if(AMREX_SPACEDIM==2) +Array<Box,4>dom_edge_boxes +={{amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType)}}; +#else +Array<Box,12>dom_edge_boxes +={{amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),idxType), +// +amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),2,len[2]),idxType), +// +amrex::convert(amrex::adjCellLo(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellLo(gdomain,1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellHi(gdomain,1,len[1]),2,len[2]),idxType)}}; +#endif + +for(constBox&b:dom_edge_boxes){ +Boxtmp=b&bx; +amrex::For(tmp,[=](inti,intj,intk)noexcept +{ +amrex::ignore_unused(j,k); +IntVectconstidx(AMREX_D_DECL(i,j,k)); +fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); +f_user(idx,fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +} +} +#endif + +#if(AMREX_SPACEDIM==3) +//fillonboxcorners +{ +Array<Box,8>dom_corner_boxes +={{amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellLo(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellLo(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellLo(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType), +amrex::convert(amrex::adjCellHi(amrex::adjCellHi(amrex::adjCellHi(gdomain,0,len[0]),1,len[1]),2,len[2]),idxType)}}; + +for(constBox&b:dom_corner_boxes){ +Boxtmp=b&bx; +amrex::For(tmp,[=](inti,intj,intk)noexcept +{ +IntVectconstidx(AMREX_D_DECL(i,j,k)); +fillfunc(idx,fab,dcomp,numcomp,domain,bcr_p,0); +f_user(idx,fab,dcomp,numcomp,geomdata,time, +bcr_p,0,orig_comp); +}); +} +} #endif -} - + +#endif } -#endif + +} +#endif diff --git a/amrex/docs_xml/doxygen/classamrex_1_1GpuBndryFuncFab.xml b/amrex/docs_xml/doxygen/classamrex_1_1GpuBndryFuncFab.xml index fa802a09bc..e567962603 100644 --- a/amrex/docs_xml/doxygen/classamrex_1_1GpuBndryFuncFab.xml +++ b/amrex/docs_xml/doxygen/classamrex_1_1GpuBndryFuncFab.xml @@ -118,7 +118,7 @@ - + @@ -176,7 +176,7 @@ - + void @@ -225,7 +225,7 @@ - + diff --git a/amrex/docs_xml/doxygen/classamrex_1_1PhysBCFunct.xml b/amrex/docs_xml/doxygen/classamrex_1_1PhysBCFunct.xml index 16a6e150b5..d9b96677b5 100644 --- a/amrex/docs_xml/doxygen/classamrex_1_1PhysBCFunct.xml +++ b/amrex/docs_xml/doxygen/classamrex_1_1PhysBCFunct.xml @@ -20,7 +20,7 @@ - + Vector< BCRec > @@ -33,7 +33,7 @@ - + F @@ -46,7 +46,7 @@ - + @@ -202,7 +202,7 @@ Based on BCRec for the domain, we need to make BCRec for this Box Note that we pass 0 as starting component of bcrs. - + void @@ -239,7 +239,7 @@ - + @@ -336,7 +336,7 @@ - + amrex::PhysBCFunctdefine amrex::PhysBCFunctdefine