| 686 | } |
| 687 | |
| 688 | void RadSolve::levelFlux(int level, |
| 689 | Array<MultiFab, AMREX_SPACEDIM>& Flux, |
| 690 | MultiFab& Er, int igroup) |
| 691 | { |
| 692 | BL_PROFILE("RadSolve::levelFlux"); |
| 693 | const BoxArray& grids = parent->boxArray(level); |
| 694 | const DistributionMapping& dmap = parent->DistributionMap(level); |
| 695 | |
| 696 | // grow a larger MultiFab to hold Er so we can difference across faces |
| 697 | MultiFab Erborder(grids, dmap, 1, 1); |
| 698 | Erborder.setVal(0.0); |
| 699 | MultiFab::Copy(Erborder, Er, igroup, 0, 1, 0); |
| 700 | |
| 701 | Erborder.FillBoundary(parent->Geom(level).periodicity()); // zeroes left in off-level boundaries |
| 702 | |
| 703 | auto dx = parent->Geom(level).CellSizeArray(); |
| 704 | |
| 705 | for (int n = 0; n < AMREX_SPACEDIM; n++) { |
| 706 | |
| 707 | const MultiFab *bp; |
| 708 | |
| 709 | if (hd) { |
| 710 | bp = &hd->bCoefficients(n); |
| 711 | } |
| 712 | else if (hm) { |
| 713 | bp = &hm->bCoefficients(level, n); |
| 714 | } |
| 715 | else if (hem) { |
| 716 | bp = &hem->bCoefficients(level, n); |
| 717 | } |
| 718 | |
| 719 | MultiFab &bcoef = *(MultiFab*)bp; |
| 720 | |
| 721 | #ifdef _OPENMP |
| 722 | #pragma omp parallel |
| 723 | #endif |
| 724 | for (MFIter mfi(Flux[n], TilingIfNotGPU()); mfi.isValid(); ++mfi) { |
| 725 | const Box& bx = mfi.tilebox(); |
| 726 | |
| 727 | auto Erborder_arr = Erborder[mfi].array(); |
| 728 | auto bcoef_arr = bcoef[mfi].array(); |
| 729 | auto Flux_arr = Flux[n][mfi].array(); |
| 730 | |
| 731 | Real beta = radsolve::beta; |
| 732 | |
| 733 | amrex::ParallelFor(bx, |
| 734 | [=] AMREX_GPU_HOST_DEVICE (int i, int j, int k) |
| 735 | { |
| 736 | if (n == 0) { |
| 737 | |
| 738 | const Real fac = -beta / dx[0]; |
| 739 | |
| 740 | Flux_arr(i,j,k) = bcoef_arr(i,j,k) * (Erborder_arr(i,j,k) - Erborder_arr(i-1,j,k)) * fac; |
| 741 | |
| 742 | } |
| 743 | else if (n == 1) { |
| 744 | |
| 745 | const Real fac = -beta / dx[1]; |
no test coverage detected