149 Vector<std::string> varnames;
152 const int ncomp_mf = varnames.size();
153 const auto ngrow_vars = IntVect(
NGROW-1,
NGROW-1,0);
161 for (
int lev = 0; lev <= finest_level; ++lev) {
168 Real fill_value = 0.0_rt;
169 for (
int lev = 0; lev <= finest_level; ++lev) {
174 Vector<MultiFab> mf(finest_level+1);
175 for (
int lev = 0; lev <= finest_level; ++lev) {
176 mf[lev].define(grids[lev], dmap[lev], ncomp_mf, ngrow_vars);
180 Vector<MultiFab> mf_nd(finest_level+1);
181 for (
int lev = 0; lev <= finest_level; ++lev) {
182 BoxArray nodal_grids(grids[lev]); nodal_grids.surroundingNodes();
183 mf_nd[lev].define(nodal_grids, dmap[lev], AMREX_SPACEDIM, 0);
184 mf_nd[lev].setVal(0.);
188 Vector<MultiFab> mf_u(finest_level+1);
189 Vector<MultiFab> mf_v(finest_level+1);
190 Vector<MultiFab> mf_w(finest_level+1);
192 for (
int lev = 0; lev <= finest_level; ++lev) {
193 BoxArray grid_stag_u(grids[lev]); grid_stag_u.surroundingNodes(0);
194 BoxArray grid_stag_v(grids[lev]); grid_stag_v.surroundingNodes(1);
195 BoxArray grid_stag_w(grids[lev]); grid_stag_w.surroundingNodes(2);
196 mf_u[lev].define(grid_stag_u, dmap[lev], 1, 0);
197 mf_v[lev].define(grid_stag_v, dmap[lev], 1, 0);
198 mf_w[lev].define(grid_stag_w, dmap[lev], 1, 0);
199 MultiFab::Copy(mf_u[lev],*
xvel_new[lev],0,0,1,0);
200 MultiFab::Copy(mf_v[lev],*
yvel_new[lev],0,0,1,0);
201 MultiFab::Copy(mf_w[lev],*
zvel_new[lev],0,0,1,0);
206 Vector<MultiFab> mf_cc_vel(finest_level+1);
213 for (
int lev = 0; lev <= finest_level; ++lev) {
214 mf_cc_vel[lev].define(grids[lev], dmap[lev], AMREX_SPACEDIM, IntVect(1,1,0));
215 mf_cc_vel[lev].setVal(0.0_rt);
216 average_face_to_cellcenter(mf_cc_vel[lev],0,
218 mf_cc_vel[lev].FillBoundary(geom[lev].periodicity());
222 amrex::Interpolater* mapper = &cell_cons_interp;
224 for (
int lev = 1; lev <= finest_level; ++lev) {
225 Vector<MultiFab*> fmf = {&(mf_cc_vel[lev]), &(mf_cc_vel[lev])};
226 Vector<Real> ftime = {
t_new[lev],
t_new[lev]};
227 Vector<MultiFab*> cmf = {&mf_cc_vel[lev-1], &mf_cc_vel[lev-1]};
228 Vector<Real> ctime = {
t_new[lev],
t_new[lev]};
231 amrex::FillPatchTwoLevels(mf_cc_vel[lev], mf_cc_vel[lev].nGrowVect(), IntVect(0,0,0),
232 t_new[lev], cmf, ctime, fmf, ftime,
233 0, 0, mf_cc_vel[lev].nComp(), geom[lev-1], geom[lev],
239 for (
int lev = 0; lev <= finest_level; ++lev)
245 for (
int i = 0; i <
NCONS; ++i) {
248 amrex::Abort(
"Found while writing output: Cons (salt, temp, or scalar, etc) contains nan or inf");
250 MultiFab::Copy(mf[lev],*
cons_new[lev],i,mf_comp,1,ngrow_vars);
257 if (mf_cc_vel[lev].contains_nan(0,1) || mf_cc_vel[lev].contains_inf(0,1)) {
258 amrex::Abort(
"Found while writing output: u velocity contains nan or inf");
260 MultiFab::Copy(mf[lev], mf_cc_vel[lev], 0, mf_comp, 1, 0);
264 if (mf_cc_vel[lev].contains_nan(1,1) || mf_cc_vel[lev].contains_inf(1,1)) {
265 amrex::Abort(
"Found while writing output: v velocity contains nan or inf");
267 MultiFab::Copy(mf[lev], mf_cc_vel[lev], 1, mf_comp, 1, 0);
271 if (mf_cc_vel[lev].contains_nan(2,1) || mf_cc_vel[lev].contains_inf(2,1)) {
272 amrex::Abort(
"Found while writing output: z velocity contains nan or inf");
274 MultiFab::Copy(mf[lev], mf_cc_vel[lev], 2, mf_comp, 1, 0);
279 auto calculate_derived = [&](
const std::string& der_name,
283 MultiFab dmf(mf[lev], make_alias, mf_comp, 1);
285#pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
287 for (MFIter mfi(dmf, TilingIfNotGPU()); mfi.isValid(); ++mfi)
289 const Box& bx = mfi.tilebox();
290 auto& dfab = dmf[mfi];
292 if (der_name ==
"vorticity") {
293 auto const& sfab = mf_cc_vel[lev][mfi];
294 der_function(bx, dfab, 0, 1, sfab,
vec_pm[lev]->const_array(mfi),
vec_pn[lev]->const_array(mfi), Geom(lev),
t_new[0],
nullptr, lev);
296 auto const& sfab = (*
cons_new[lev])[mfi];
297 der_function(bx, dfab, 0, 1, sfab,
vec_pm[lev]->const_array(mfi),
vec_pn[lev]->const_array(mfi), Geom(lev),
t_new[0],
nullptr, lev);
309 Real dx = Geom()[lev].CellSizeArray()[0];
310 Real dy = Geom()[lev].CellSizeArray()[1];
317 MultiFab dmf(mf[lev], make_alias, mf_comp, AMREX_SPACEDIM);
319#pragma omp parallel if (Gpu::notInLaunchRegion())
321 for (MFIter mfi(dmf, TilingIfNotGPU()); mfi.isValid(); ++mfi) {
322 const Box& bx = mfi.tilebox();
323 const Array4<Real> loc_arr = dmf.array(mfi);
324 const Array4<Real const> zp_arr =
vec_z_phys_nd[lev]->const_array(mfi);
326 ParallelFor(bx, [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) {
327 loc_arr(i,j,k,0) = (i+0.5_rt) * dx;
328 loc_arr(i,j,k,1) = (j+0.5_rt) * dy;
329 loc_arr(i,j,k,2) = 0.125_rt * (zp_arr(i,j ,k ) + zp_arr(i+1,j ,k ) +
330 zp_arr(i,j+1,k ) + zp_arr(i+1,j+1,k ) +
331 zp_arr(i,j ,k+1) + zp_arr(i+1,j ,k+1) +
332 zp_arr(i,j+1,k+1) + zp_arr(i+1,j+1,k+1) );
335 mf_comp += AMREX_SPACEDIM;
338#ifdef REMORA_USE_PARTICLES
339 const auto& particles_namelist( particleData.getNames() );
340 for (ParticlesNamesVector::size_type i = 0; i < particles_namelist.size(); i++) {
342 MultiFab temp_dat(mf[lev].boxArray(), mf[lev].DistributionMap(), 1, 0);
344 particleData[particles_namelist[i]]->Increment(temp_dat, lev);
345 MultiFab::Copy(mf[lev], temp_dat, 0, mf_comp, 1, 0);
350 Vector<std::string> particle_mesh_plot_names(0);
351 particleData.GetMeshPlotVarNames( particle_mesh_plot_names );
352 for (
int i = 0; i < particle_mesh_plot_names.size(); i++) {
353 std::string plot_var_name(particle_mesh_plot_names[i]);
355 MultiFab temp_dat(mf[lev].boxArray(), mf[lev].DistributionMap(), 1, 1);
357 particleData.GetMeshPlotVar(plot_var_name, temp_dat, lev);
358 MultiFab::Copy(mf[lev], temp_dat, 0, mf_comp, 1, 0);
365 Real dz = Geom()[lev].CellSizeArray()[2];
366 int N = Geom()[lev].Domain().size()[2];
369#pragma omp parallel if (Gpu::notInLaunchRegion())
371 for (MFIter mfi(mf_nd[lev], TilingIfNotGPU()); mfi.isValid(); ++mfi)
373 const Box& bx = mfi.tilebox();
374 Array4<Real> mf_arr = mf_nd[lev].array(mfi);
375 ParallelFor(bx, [=] AMREX_GPU_DEVICE (
int i,
int j,
int k) {
376 mf_arr(i,j,k,2) = mf_arr(i,j,k,2) + (N-k) * dz;
384 if (finest_level == 0)
387 amrex::Print() <<
"Writing plotfile " << plotfilename <<
"\n";
389 GetVecOfConstPtrs(mf),
390 GetVecOfConstPtrs(mf_nd),
391 GetVecOfConstPtrs(mf_u),
392 GetVecOfConstPtrs(mf_v),
393 GetVecOfConstPtrs(mf_w),
399#ifdef REMORA_USE_PARTICLES
400 particleData.Checkpoint(plotfilename);
403#ifdef REMORA_USE_HDF5
405 amrex::Print() <<
"Writing plotfile " << plotfilename+
"d01.h5" <<
"\n";
406 WriteMultiLevelPlotfileHDF5(plotfilename, finest_level+1,
407 GetVecOfConstPtrs(mf),
412 amrex::Abort(
"User specified unknown plot_filetype");
417 amrex::Print() <<
"Writing plotfile " << plotfilename <<
"\n";
419 [[maybe_unused]]
int desired_ratio = std::max(std::max(ref_ratio[lev0][0],ref_ratio[lev0][1]),ref_ratio[lev0][2]);
420 bool any_ratio_one = ( ( (ref_ratio[lev0][0] == 1) || (ref_ratio[lev0][1] == 1) ) ||
421 (ref_ratio[lev0][2] == 1) );
422 for (
int lev = 1; lev < finest_level; lev++) {
423 any_ratio_one = any_ratio_one ||
424 ( ( (ref_ratio[lev][0] == 1) || (ref_ratio[lev][1] == 1) ) ||
425 (ref_ratio[lev][2] == 1) );
428 Vector<IntVect> r2(finest_level);
429 Vector<Geometry> g2(finest_level+1);
430 Vector<MultiFab> mf2(finest_level+1);
432 mf2[0].define(grids[0], dmap[0], ncomp_mf, 0);
435 MultiFab::Copy(mf2[0],mf[0],0,0,mf[0].nComp(),0);
438 Array<int,AMREX_SPACEDIM> periodicity =
439 {Geom()[0].isPeriodic(0),Geom()[0].isPeriodic(1),Geom()[0].isPeriodic(2)};
440 g2[0].define(Geom()[0].Domain(),&(Geom()[0].ProbDomain()),0,periodicity.data());
442 r2[0] = IntVect(1,1,ref_ratio[0][0]);
443 for (
int lev = 1; lev <= finest_level; ++lev) {
447 r2[lev-1][2] = r2[lev-2][2] * ref_ratio[lev-1][0];
450 mf2[lev].define(refine(grids[lev],r2[lev-1]), dmap[lev], ncomp_mf, 0);
453 Box d2(Geom()[lev].Domain());
454 d2.refine(r2[lev-1]);
456 g2[lev].define(d2,&(Geom()[lev].ProbDomain()),0,periodicity.data());
461 amrex::Vector<amrex::BCRec> null_dom_bcs;
462 null_dom_bcs.resize(mf2[0].nComp());
463 for (
int n = 0; n < mf2[0].nComp(); n++) {
464 for (
int dir = 0; dir < AMREX_SPACEDIM; dir++) {
471 for (
int lev = 1; lev <= finest_level; ++lev) {
472 Interpolater* mapper_c = &pc_interp;
473 InterpFromCoarseLevel(mf2[lev],
t_new[lev], mf[lev],
474 0, 0, mf2[lev].nComp(),
477 r2[lev-1], mapper_c, null_dom_bcs, 0);
481 Vector<IntVect> rr(finest_level);
482 for (
int lev = 0; lev < finest_level; ++lev) {
483 rr[lev] = IntVect(ref_ratio[lev][0],ref_ratio[lev][1],ref_ratio[lev][0]);
487 GetVecOfConstPtrs(mf2),
488 GetVecOfConstPtrs(mf_nd),
489 GetVecOfConstPtrs(mf_u),
490 GetVecOfConstPtrs(mf_v),
491 GetVecOfConstPtrs(mf_w),
497#ifdef REMORA_USE_PARTICLES
498 particleData.Checkpoint(plotfilename);
502 GetVecOfConstPtrs(mf),
503 GetVecOfConstPtrs(mf_nd),
504 GetVecOfConstPtrs(mf_u),
505 GetVecOfConstPtrs(mf_v),
506 GetVecOfConstPtrs(mf_w),
511#ifdef REMORA_USE_PARTICLES
512 particleData.Checkpoint(plotfilename);
517 for (
int lev = 0; lev <= finest_level; ++lev) {
539 const Vector<const MultiFab*>& mf,
540 const Vector<const MultiFab*>& mf_nd,
541 const Vector<const MultiFab*>& mf_u,
542 const Vector<const MultiFab*>& mf_v,
543 const Vector<const MultiFab*>& mf_w,
544 const Vector<std::string>& varnames,
545 const Vector<Geometry>& my_geom,
547 const Vector<int>& level_steps,
548 const Vector<IntVect>& rr,
549 const std::string &versionName,
550 const std::string &levelPrefix,
551 const std::string &mfPrefix,
552 const Vector<std::string>& extra_dirs)
const
554 BL_PROFILE(
"WriteMultiLevelPlotfileWithBathymetry()");
556 BL_ASSERT(nlevels <= mf.size());
557 BL_ASSERT(nlevels <= ref_ratio.size()+1);
558 BL_ASSERT(nlevels <= level_steps.size());
559 BL_ASSERT(mf[0]->nComp() == varnames.size());
561 bool callBarrier(
false);
562 PreBuildDirectorHierarchy(plotfilename, levelPrefix, nlevels, callBarrier);
563 if (!extra_dirs.empty()) {
564 for (
const auto& d : extra_dirs) {
565 const std::string ed = plotfilename+
"/"+d;
566 PreBuildDirectorHierarchy(ed, levelPrefix, nlevels, callBarrier);
569 ParallelDescriptor::Barrier();
571 if (ParallelDescriptor::MyProc() == ParallelDescriptor::NProcs()-1) {
572 Vector<BoxArray> boxArrays(nlevels);
573 for(
int level(0); level < boxArrays.size(); ++level) {
574 boxArrays[level] = mf[level]->boxArray();
578 VisMF::IO_Buffer io_buffer(VisMF::IO_Buffer_Size);
579 std::string HeaderFileName(plotfilename +
"/Header");
580 std::ofstream HeaderFile;
581 HeaderFile.rdbuf()->pubsetbuf(io_buffer.dataPtr(), io_buffer.size());
582 HeaderFile.open(HeaderFileName.c_str(), std::ofstream::out |
583 std::ofstream::trunc |
584 std::ofstream::binary);
585 if( ! HeaderFile.good()) FileOpenFailed(HeaderFileName);
587 my_geom, time, level_steps, rr, versionName,
588 levelPrefix, mfPrefix);
591 if (AsyncOut::UseAsyncOut()) {
592 AsyncOut::Submit(std::move(f));
598 std::string mf_nodal_prefix =
"Nu_nd";
599 std::string mf_uface_prefix =
"UFace";
600 std::string mf_vface_prefix =
"VFace";
601 std::string mf_wface_prefix =
"WFace";
603 for (
int level = 0; level <= finest_level; ++level)
605 if (AsyncOut::UseAsyncOut()) {
606 VisMF::AsyncWrite(*mf[level],
607 MultiFabFileFullPrefix(level, plotfilename, levelPrefix, mfPrefix),
609 VisMF::AsyncWrite(*mf_nd[level],
610 MultiFabFileFullPrefix(level, plotfilename, levelPrefix, mf_nodal_prefix),
613 VisMF::AsyncWrite(*mf_u[level],
614 MultiFabFileFullPrefix(level, plotfilename, levelPrefix, mf_uface_prefix),
616 VisMF::AsyncWrite(*mf_v[level],
617 MultiFabFileFullPrefix(level, plotfilename, levelPrefix, mf_vface_prefix),
619 VisMF::AsyncWrite(*mf_w[level],
620 MultiFabFileFullPrefix(level, plotfilename, levelPrefix, mf_wface_prefix),
623 }
else {
const MultiFab* data;
624 std::unique_ptr<MultiFab> mf_tmp;
625 if (mf[level]->nGrowVect() != 0) {
626 mf_tmp = std::make_unique<MultiFab>(mf[level]->boxArray(),
627 mf[level]->DistributionMap(),
628 mf[level]->nComp(), 0, MFInfo(),
629 mf[level]->Factory());
630 MultiFab::Copy(*mf_tmp, *mf[level], 0, 0, mf[level]->nComp(), 0);
635 VisMF::Write(*data , MultiFabFileFullPrefix(level, plotfilename, levelPrefix, mfPrefix));
636 VisMF::Write(*mf_nd[level], MultiFabFileFullPrefix(level, plotfilename, levelPrefix, mf_nodal_prefix));
638 VisMF::Write(*mf_u[level], MultiFabFileFullPrefix(level, plotfilename, levelPrefix, mf_uface_prefix));
639 VisMF::Write(*mf_v[level], MultiFabFileFullPrefix(level, plotfilename, levelPrefix, mf_vface_prefix));
640 VisMF::Write(*mf_w[level], MultiFabFileFullPrefix(level, plotfilename, levelPrefix, mf_wface_prefix));
661 [[maybe_unused]]
int nlevels,
662 const Vector<BoxArray> &bArray,
663 const Vector<std::string> &varnames,
664 const Vector<Geometry>& my_geom,
666 const Vector<int> &level_steps,
667 const Vector<IntVect>& my_ref_ratio,
668 const std::string &versionName,
669 const std::string &levelPrefix,
670 const std::string &mfPrefix)
const
672 BL_ASSERT(nlevels <= bArray.size());
673 BL_ASSERT(nlevels <= ref_ratio.size()+1);
674 BL_ASSERT(nlevels <= level_steps.size());
676 int num_extra_mfs = 1;
681 HeaderFile.precision(17);
684 HeaderFile << versionName <<
'\n';
686 HeaderFile << varnames.size() <<
'\n';
688 for (
int ivar = 0; ivar < varnames.size(); ++ivar) {
689 HeaderFile << varnames[ivar] <<
"\n";
691 HeaderFile << AMREX_SPACEDIM <<
'\n';
692 HeaderFile << time <<
'\n';
693 HeaderFile << finest_level <<
'\n';
694 for (
int i = 0; i < AMREX_SPACEDIM; ++i) {
695 HeaderFile << my_geom[0].ProbLo(i) <<
' ';
698 for (
int i = 0; i < AMREX_SPACEDIM; ++i) {
699 HeaderFile << my_geom[0].ProbHi(i) <<
' ';
702 for (
int i = 0; i < finest_level; ++i) {
703 HeaderFile << my_ref_ratio[i][0] <<
' ';
706 for (
int i = 0; i <= finest_level; ++i) {
707 HeaderFile << my_geom[i].Domain() <<
' ';
710 for (
int i = 0; i <= finest_level; ++i) {
711 HeaderFile << level_steps[i] <<
' ';
714 for (
int i = 0; i <= finest_level; ++i) {
715 for (
int k = 0; k < AMREX_SPACEDIM; ++k) {
716 HeaderFile << my_geom[i].CellSize()[k] <<
' ';
720 HeaderFile << (int) my_geom[0].
Coord() <<
'\n';
723 for (
int level = 0; level <= finest_level; ++level) {
724 HeaderFile << level <<
' ' << bArray[level].size() <<
' ' << time <<
'\n';
725 HeaderFile << level_steps[level] <<
'\n';
727 const IntVect& domain_lo = my_geom[level].Domain().smallEnd();
728 for (
int i = 0; i < bArray[level].size(); ++i)
733 const Box& b = shift(bArray[level][i], -domain_lo);
734 RealBox loc = RealBox(b, my_geom[level].CellSize(), my_geom[level].ProbLo());
735 for (
int n = 0; n < AMREX_SPACEDIM; ++n) {
736 HeaderFile << loc.lo(n) <<
' ' << loc.hi(n) <<
'\n';
740 HeaderFile << MultiFabHeaderPath(level, levelPrefix, mfPrefix) <<
'\n';
742 HeaderFile << num_extra_mfs <<
"\n";
743 HeaderFile <<
"3" <<
"\n";
744 HeaderFile <<
"amrexvec_nu_x" <<
"\n";
745 HeaderFile <<
"amrexvec_nu_y" <<
"\n";
746 HeaderFile <<
"amrexvec_nu_z" <<
"\n";
747 std::string mf_nodal_prefix =
"Nu_nd";
748 for (
int level = 0; level <= finest_level; ++level) {
749 HeaderFile << MultiFabHeaderPath(level, levelPrefix, mf_nodal_prefix) <<
'\n';
752 HeaderFile <<
"1" <<
"\n";
753 HeaderFile <<
"u_vel" <<
"\n";
754 std::string mf_uface_prefix =
"UFace";
755 for (
int level = 0; level <= finest_level; ++level) {
756 HeaderFile << MultiFabHeaderPath(level, levelPrefix, mf_uface_prefix) <<
'\n';
758 HeaderFile <<
"1" <<
"\n";
759 HeaderFile <<
"v_vel" <<
"\n";
760 std::string mf_vface_prefix =
"VFace";
761 for (
int level = 0; level <= finest_level; ++level) {
762 HeaderFile << MultiFabHeaderPath(level, levelPrefix, mf_vface_prefix) <<
'\n';
764 HeaderFile <<
"1" <<
"\n";
765 HeaderFile <<
"w_vel" <<
"\n";
766 std::string mf_wface_prefix =
"WFace";
767 for (
int level = 0; level <= finest_level; ++level) {
768 HeaderFile << MultiFabHeaderPath(level, levelPrefix, mf_wface_prefix) <<
'\n';
780 for (MFIter mfi(*
cons_new[lev],
false); mfi.isValid(); ++mfi) {
781 Box gbx1 = mfi.growntilebox(IntVect(
NGROW+1,
NGROW+1,0));
782 Box ubx = mfi.grownnodaltilebox(0,IntVect(
NGROW,
NGROW,0));
783 Box vbx = mfi.grownnodaltilebox(1,IntVect(
NGROW,
NGROW,0));
785 Array4<Real>
const& Zt_avg1 =
vec_Zt_avg1[lev]->array(mfi);
786 Array4<Real>
const& ubar =
vec_ubar[lev]->array(mfi);
787 Array4<Real>
const& vbar =
vec_vbar[lev]->array(mfi);
788 Array4<Real>
const& xvel =
xvel_new[lev]->array(mfi);
789 Array4<Real>
const& yvel =
yvel_new[lev]->array(mfi);
793 Array4<Real const>
const& mskr =
vec_mskr[lev]->array(mfi);
794 Array4<Real const>
const& msku =
vec_msku[lev]->array(mfi);
795 Array4<Real const>
const& mskv =
vec_mskv[lev]->array(mfi);
797 ParallelFor(makeSlab(gbx1,2,0), [=] AMREX_GPU_DEVICE (
int i,
int j,
int )
799 if (mskr(i,j,0) == 0.0) {
800 Zt_avg1(i,j,0) = fill_value;
803 ParallelFor(gbx1, [=] AMREX_GPU_DEVICE (
int i,
int j,
int k)
805 if (mskr(i,j,0) == 0.0) {
806 temp(i,j,k) = fill_value;
807 salt(i,j,k) = fill_value;
810 ParallelFor(makeSlab(ubx,2,0), 3, [=] AMREX_GPU_DEVICE (
int i,
int j,
int ,
int n)
812 if (msku(i,j,0) == 0.0 && ubar(i,j,0)==fill_where) {
813 ubar(i,j,0,n) = fill_value;
816 ParallelFor(makeSlab(vbx,2,0), 3, [=] AMREX_GPU_DEVICE (
int i,
int j,
int ,
int n)
818 if (mskv(i,j,0) == 0.0 && vbar(i,j,0)==fill_where) {
819 vbar(i,j,0,n) = fill_value;
822 ParallelFor(ubx, [=] AMREX_GPU_DEVICE (
int i,
int j,
int k)
824 if (msku(i,j,0) == 0.0 && xvel(i,j,k)==fill_where) {
825 xvel(i,j,k) = fill_value;
828 ParallelFor(vbx, [=] AMREX_GPU_DEVICE (
int i,
int j,
int k)
830 if (mskv(i,j,0) == 0.0 && yvel(i,j,k)==fill_where) {
831 yvel(i,j,k) = fill_value;
835 Gpu::streamSynchronize();