Access pointer to MultiFAB data for ALL boxes/fabs at once

[tony3601000]

Hi there,

I need to perform operations for all cells in parallel using cuBLAS. Since the cells are grouped into boxes/fabs, I currently have the following code (simplified):

// Loop through each box/fab
for (amrex::MFIter mfi(multifab); mfi.isValid(); ++mfi)
{   
  // 1. Get device pointer to multifab data
  amrex::Real* AMREX_RESTRICT multifab_ptr = multifab[mfi].dataPtr();

  // 2. Create cublas cublas
  cublasHandle_t cublas;
  cublasCreate(&cublas);
  
  // 3. Call cublas function 
  cublasDgemv(...);

  // 4. Destroy cublas
  cublasDestroy(cublas);
}

Ideally, I'd like to perform something like this instead:

// 1. Get pointer to all data at all boxes/fabs (contiguous in memory)
amrex::Real* AMREX_RESTRICT multifab_ptrs = multifab.dataPtrs();

// 2. Create cublas cublas
cublasHandle_t cublas;
cublasCreate(&cublas);
   
// 3. Call cublas function ON ALL FABS AT ONCE
cublasDgemv(...);

// 4. Destroy cublas
cublasDestroy(cublas);

I'm aware of the following strategies:

auto multifab_arrays = multifab.arrays();
auto multifab_const_arrays = multifab.const_arrays();

But these do not seem to be stored in a contiguous format. Otherwise, what is the best approach to implementing parallel GPU computing for a cluster of small fabs, each with a few cells?

Thank you for the help!

Best,

Anthony

[WeiqunZhang]

Maybe you can use the batched version of cublasDgemv.

Or maybe it's not too bad to copy everything to a contiguous chunk of memory and perform cublas there.

Or you can allocate the memory and give MultiFab the pointers. In that case, MultiFab just holds non-owning pointers provided by you. So you need to mange the memory allocation and deallocation.

Or we maybe could add an option to allocate everything in a contiguous chunk of memory. It's not hard to do.

[tony3601000]

I'll look into copying everything to a contiguous chunk of memory. Ideally, an AMReX option to allocate all multifab data to a contiguous memory block would be best. Thank you!

[rbielaw]

Is there a significant downside to having multiple or all MultiFabs on a level being in a contiguous chunk of memory? Having them be contiguous would open-up some optimization options for GPUs.

[WeiqunZhang]

I think that would be very error prone and inefficient in memory usage.

[WeiqunZhang]

@tony3601000 Could you give this a try? #3857

[WeiqunZhang]

Note that the new feature is not turned on by default.