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.

Currently there is an issue where the extra size added by BoxIndexerND pushes an EB kernel over the parameter size limit for SYCL. 

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.

I think it is this one, which indeed uses a lot of Array4 objects:

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