| 1268 | } |
| 1269 | |
| 1270 | Real |
| 1271 | MultiFab::sum_unique (int comp, |
| 1272 | bool local, |
| 1273 | const Periodicity& period) const |
| 1274 | { |
| 1275 | BL_PROFILE("MultiFab::sum_unique()"); |
| 1276 | |
| 1277 | // no duplicately distributed points if cell centered |
| 1278 | if (ixType().cellCentered()) { |
| 1279 | return this->sum(comp, local); |
| 1280 | } |
| 1281 | |
| 1282 | // Owner is the grid with the lowest grid number containing the data |
| 1283 | std::unique_ptr<iMultiFab> owner_mask = OwnerMask(period); |
| 1284 | |
| 1285 | Real sm = Real(0.0); |
| 1286 | #ifdef AMREX_USE_GPU |
| 1287 | if (Gpu::inLaunchRegion()) { |
| 1288 | auto const& ma = this->const_arrays(); |
| 1289 | auto const& msk = owner_mask->const_arrays(); |
| 1290 | sm = ParReduce(TypeList<ReduceOpSum>{}, TypeList<Real>{}, *this, IntVect(0), |
| 1291 | [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k) noexcept |
| 1292 | -> GpuTuple<Real> |
| 1293 | { |
| 1294 | return msk[box_no](i,j,k) ? ma[box_no](i,j,k,comp) : 0.0_rt; |
| 1295 | }); |
| 1296 | } else |
| 1297 | #endif |
| 1298 | { |
| 1299 | #ifdef AMREX_USE_OMP |
| 1300 | #pragma omp parallel if (!system::regtest_reduction) reduction(+:sm) |
| 1301 | #endif |
| 1302 | for (MFIter mfi(*this,true); mfi.isValid(); ++mfi) |
| 1303 | { |
| 1304 | Box const& bx = mfi.tilebox(); |
| 1305 | Array4<Real const> const& a = this->const_array(mfi); |
| 1306 | Array4<int const> const& msk = owner_mask->const_array(mfi); |
| 1307 | Real tmp = 0.0_rt; |
| 1308 | AMREX_LOOP_3D(bx, i, j, k, |
| 1309 | { |
| 1310 | tmp += msk(i,j,k) ? a(i,j,k,comp) : 0.0_rt; |
| 1311 | }); |
| 1312 | sm += tmp; // Do it this way so that it does not break regression tests. |
| 1313 | } |
| 1314 | } |
| 1315 | |
| 1316 | if (!local) { |
| 1317 | ParallelAllReduce::Sum(sm, ParallelContext::CommunicatorSub()); |
| 1318 | } |
| 1319 | |
| 1320 | return sm; |
| 1321 | } |
| 1322 | |
| 1323 | Real |
| 1324 | MultiFab::sum_unique (Box const& region, int comp, bool local) const |
no test coverage detected