Using AMReX for GPUs - is_device_ptr equivalent?

[ElloiseFangelLloyd]

Edit: removed confusing leftovers in code.

Hello

I have a Monte Carlo particle code which is well optimised for GPU usage. I'd like to use the MLPoisson solver to update the fields at given intervals. I'm using OpenMP for parallelism and offloading.

My implementation stores the particles - I'm simulating electrons - exclusively on the device - there is no copy on the host. When I do my simulations, I can specify is_device_ptr() in the #pragma omp target statement, and then pass a pointer to the location of the particles in device memory. Then I can do operations on them. This all works well.

I have a scheme to deposit these particles onto a MultiFab using the MFIter, which is then fed to the solver. The scheme works well on the CPU, but I am encountering issues when working on the GPU. Passing the particle pointer gives a Segmentation fault, presumably because AMReX looks in host memory and finds nothing. Is there a way to specify to AMReX that the information exists only on the device - and pass it an address? Similar to the is_device_ptr() clause from OpenMP?

I have looked at the very nice Electrostatic PIC in the AMReX example codes for inspiration - however, particle collisions are a very large part of my project here, and require careful handling on the GPU, hence building something from scratch.

My particle deposition scheme:

void MapElectrons(MultiFab &grid, ElectronSystem* electrons) {

        for (MFIter mfi(grid, TilingIfNotGPU()); mfi.isValid(); ++mfi)
        {
            
            const Box& bx = mfi.growntilebox(1);
            Array4<Real> const &domain = grid.array(mfi);

            size_t current_size = electrons->num_particles;
            amrex::ParallelFor(current_size, [=] AMREX_GPU_DEVICE (int index)
            {
                Real x = electrons->x[index];
                Real y = electrons->y[index];
                Real z = electrons->z[index];

                Real half_cell = Real(InitConditions::grid_delta/2);

                int i = (x >= half_cell) ? static_cast<int>(amrex::Math::floor((x - half_cell) / InitConditions::grid_delta)) : 0;
                int j = (y >= half_cell) ? static_cast<int>(amrex::Math::floor((y - half_cell) / InitConditions::grid_delta)) : 0;
                int k = (z >= half_cell) ? static_cast<int>(amrex::Math::floor((z - half_cell) / InitConditions::grid_delta)) : 0;

                CloudInCell(domain, x, y, z, i, j, k);
                
            });
        }

    }

I get the segmentation fault in the line size_t current_size = electrons->num_particles, so when it tries to access electrons and finds nothing.

I have also tried enclosing the whole thing in a #pragma omp target region, but this gave compiler linker errors to the MFIter file, and does not seem like the right approach.

Does anyone know of any similar/equivalent approaches to the is_device_ptr() implemented in OpenMP - or other ideas?

Thank you very much in advance!!

[atmyers]

Hi Elloise,

The ElectronSystem object is your own particle implementation, correct?

The domain object should already be accessible on the device. But you do need to know the size of the kernel to launch on the host. Could you use amrex::Gpu::copy to copy the size of the electron system to a host variable, and use that to launch the kernel?

-Andrew

[WeiqunZhang]

Most likely you already know num_particles on the host, otherwise how did you allocate memory for particles?