ReduceOps.eval with BoxND

[AlexanderSinn]

Summary

This PR enables the use of ReduceOps.eval() with a BoxND of any dimension. This is achieved using BoxIndexerND, just like ParallelForRNG, which also provides support for 64-bit indexing to support very large boxes.

There was issue where the extra size added by BoxIndexerND pushes an EB kernel over the parameter size limit for SYCL. 
This was fixed by using MultiArray4 for some of the arrays instead.

error: Total size of kernel arguments exceeds limit! Total arguments size: 2064, limit: 2048
in kernel: 'typeinfo name for void amrex::launch<256, void amrex::ReduceOps<amrex::ReduceOpMax>::eval<amrex::ReduceData<int>, amrex::EB2::Level::coarsenFromFine(amrex::EB2::Level&, bool)::'lambda4'(int, int, int), 3>(amrex::BoxND<3> const&, amrex::ReduceData<int>&, amrex::EB2::Level::coarsenFromFine(amrex::EB2::Level&, bool)::'lambda4'(int, int, int) const&)::'lambda'(amrex::Gpu::Handler const&)>(int, unsigned long, amrex::gpuStream_t, amrex::EB2::Level::coarsenFromFine(amrex::EB2::Level&, bool)::'lambda4'(int, int, int) const&)::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::'lambda'(sycl::_V1::nd_item<1>)'
error: backend compiler failed build.

amrex/Src/EB/AMReX_EB2_Level.cpp

Lines 364 to 379 in 1fe7883

	reduce_op.eval(ndgbx, reduce_data,
	[=] AMREX_GPU_DEVICE (int i, int j, int k) -> ReduceTuple
	{
	amrex::ignore_unused(j,k);
	int ierr = coarsen_from_fine(AMREX_D_DECL(i,j,k), bx, 2,
	cvol,ccent,cba,cbc,cbn,
	AMREX_D_DECL(capx,capy,capz),
	AMREX_D_DECL(cfcx,cfcy,cfcz),
	AMREX_D_DECL(cecx,cecy,cecz),
	cflag,fvol,fcent,fba,fbc,fbn,
	AMREX_D_DECL(fapx,fapy,fapz),
	AMREX_D_DECL(ffcx,ffcy,ffcz),
	AMREX_D_DECL(fecx,fecy,fecz),
	fflag);
	return {ierr};
	});

Additional background

Checklist

The proposed changes:

• fix a bug or incorrect behavior in AMReX
• add new capabilities to AMReX
• changes answers in the test suite to more than roundoff level
• are likely to significantly affect the results of downstream AMReX users
• include documentation in the code and/or rst files, if appropriate