Make a CPU code work on GPU in AMRex #3329

lw5366 · 2023-05-24T12:26:28Z

lw5366
May 24, 2023

I have successfully developed a CFD solver using AMRex, it works well on cpu, and I am now trying to make it work on GPU as well. The code can be compiled successfully by using CUDA, but it encounters some issues at the running stage. I am trying to get help from here to fix it, some major potential problems are as follows:

In this solver, "particle" is used for fluid-structure interaction. From my understanding, if we want the code work on GPU, we do not need to manage the memory on GPU, as AMRex will handle it automactically when the calculation is wokring on GPU. If this is not the case, please point it out.
After the particle is initialized, then "Redistribute()" is called to distribute the particles to the right place, errors happened in calling Redistribute() using GPU. Note, it works well on CPU, does it mean we need to do something else when run it on GPU, as I can see Redistribute() will call RedistributeGPU() when GPU is used.

When I was trying to transverse the particles using the following code, errors also happened. From the GPU example "ElectromagneticPIC" provided, it also traverses the particle like this, it seems that we do not need to manage the particle memory, and AMRex will handle the data transfer between CPU and GPU automatically when required.

#ifdef AMREX_USE_OMP
#pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
#endif
 for (MyParIter pti(*this, lev); pti.isValid(); ++pti) 
 {
 	// Particle structs
 	auto& particles = pti.GetArrayOfStructs();
 	const amrex::Long np  = pti.numParticles();
 	auto p_ptr = particles().data();
 	auto& attribs = pti.GetAttribs();
 	
 	auto fxa = attribs[PIdx::fxa].data();
 	auto fya = attribs[PIdx::fya].data();
 	auto fza = attribs[PIdx::fza].data();
                  ....
         }

8: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x4265d6]
amrex::ParIter<0, 0, 28, 0, amrex::ArenaAllocator>::~ParIter()
/Exec/../../amrex/Src/Particle/AMReX_ParIter.H:99:7

This is the defninition of MyParIter:

class MyParIter
: public amrex::ParIter<0, 0, PIdx::nattribs>

{
public:
using amrex::ParIter<0, 0, PIdx::nattribs>::ParIter;
using RealVector = ParIter::ContainerType::RealVector;

     const std::array<RealVector, PIdx::nattribs>& GetAttribs () const {
     return GetStructOfArrays().GetRealData();
 }

 std::array<RealVector, PIdx::nattribs>& GetAttribs () {
     return GetStructOfArrays().GetRealData();
 }

 const RealVector& GetAttribs (int comp) const {
     return GetStructOfArrays().GetRealData(comp);
 }

 RealVector& GetAttribs (int comp) {
     return GetStructOfArrays().GetRealData(comp);
 }

};

If you have any thoughts in terms of the adaptation of the CPU code to GPU, please let me know. Thanks very much!

atmyers · 2023-05-24T17:04:23Z

atmyers
May 24, 2023
Maintainer

What kind of GPU are you trying this on? An NVIDIA one using the CUDA backend?

3 replies

lw5366 May 25, 2023
Author

Yes, it is an NVIDIA one using CUDA. Thanks!

atmyers May 25, 2023
Maintainer

For CUDA, it should place all particles in managed memory by default, so you should not need to handle any memory traffic yourself.

The reason why the error is reported in the destructor for ParIter is that's where a stream synchronize is called, I think the actual problem is in the code inside the ParIter loop. Could you try running your code using the CUDA_LAUNCH_BLOCKING=1, which will sync after every kernel launch? Also, you could try running under cuda-gdb or cuda-memcheck to see where the problem is first occuring.

lw5366 May 25, 2023
Author

Thanks for your quick answer!
CUDA_LAUNCH_BLOCKING=1 gives the same error.

Error in ParIter loop

Yes, the error happened at the destructor for ParIter. I checked the error log, and it is always "amrex::Abort::0::GPU last error detected in file ../../amrex/Src/Base/AMReX_MFIter.cpp line 229: invalid device symbol !!!". It seems that only "Gpu::streamSynchronizeAll();" related to GPU is called as I did not call anything using GPU in my ParIter loop.

MFIter::~MFIter ()
{
#ifdef AMREX_USE_OMP
#pragma omp master
#endif
{
depth = 0;
}

#ifdef BL_USE_TEAM
if ( ! (flags & NoTeamBarrier) )
ParallelDescriptor::MyTeam().MemoryBarrier();
#endif

#ifdef AMREX_USE_GPU
if (device_sync) Gpu::streamSynchronizeAll();
#endif

#ifdef AMREX_USE_GPU
AMREX_GPU_ERROR_CHECK(); // This is line 229
Gpu::Device::resetStreamIndex();
#endif
#endif

if (m_fa) {

#ifdef AMREX_USE_OMP
#pragma omp barrier
#pragma omp single
#endif
m_fa->clearThisBD();
}
}

Error in using "Restribute()"

When I call "Redistribute()" function, it also gives the error "invalid device symbol", but at a different place.
amrex::Abort::0::GPU last error detected in file ../../amrex/Src/Base/AMReX_GpuLaunchFunctsG.H line 754: invalid device symbol !!!
SIGABRT

template
void launch (int nblocks, int nthreads_per_block, std::size_t shared_mem_bytes,
gpuStream_t stream, L&& f) noexcept
{
AMREX_ASSERT(nthreads_per_block <= AMREX_GPU_MAX_THREADS);
AMREX_LAUNCH_KERNEL(AMREX_GPU_MAX_THREADS, nblocks, nthreads_per_block, shared_mem_bytes,
stream, [=] AMREX_GPU_DEVICE () noexcept { f(); });
AMREX_GPU_ERROR_CHECK(); //This is line 754
}

Full error log:

=== If no file names and line numbers are shown below, one can run
addr2line -Cpfie my_exefile my_line_address
to convert my_line_address (e.g., 0x4a6b) into file name and line number.
Or one can use amrex/Tools/Backtrace/parse_bt.py.

=== Please note that the line number reported by addr2line may not be accurate.
One can use
readelf -wl my_exefile | grep my_line_address'
to find out the offset for that line.

0: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x83ff76]
amrex::BLBackTrace::print_backtrace_info(_IO_FILE*)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Base/AMReX_BLBackTrace.cpp:186:25

1: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x83fb53]
amrex::BLBackTrace::handler(int)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Base/AMReX_BLBackTrace.cpp:96:7

2: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x5dbc5f]
amrex::ParallelDescriptor::Abort(int, bool)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Base/AMReX_ParallelDescriptor.cpp:206:21

3: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x56aa42]
amrex::Abort_host(char const*)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Base/AMReX.cpp:242:1

4: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x56a822]
amrex::Abort(std::__cxx11::basic_string<char, std::char_traits, std::allocator > const&)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Base/AMReX.cpp:194:1

5: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x4245b8]
amrex::Gpu::ErrorCheck(char const*, int)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Base/AMReX_GpuError.H:54:240

6: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x43cf7a]
void amrex::launch<__nv_dl_wrapper_t<__nv_dl_tag<void (amrex::ReduceOps<amrex::ReduceOpMin, amrex::ReduceOpMin, amrex::ReduceOpMin, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax>::)(int, amrex::ReduceData<int, int, int, int, int, int, int, int, int>&, __nv_dl_wrapper_t<__nv_dl_tag<void (amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::build, 1u>, amrex::Box* const>&&), &(void amrex::ReduceOps<amrex::ReduceOpMin, amrex::ReduceOpMin, amrex::ReduceOpMin, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax>::eval<int, amrex::ReduceData<int, int, int, int, int, int, int, int, int>, __nv_dl_wrapper_t<__nv_dl_tag<void (amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::build, 1u>, amrex::Box const>, void>(int, amrex::ReduceData<int, int, int, int, int, int, int, int, int>&, __nv_dl_wrapper_t<__nv_dl_tag<void (amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::build, 1u>, amrex::Box const>&&)), 1u>, amrex::GpuTuple<int, int, int, int, int, int, int, int, int>, int, int, __nv_dl_wrapper_t<__nv_dl_tag<void (amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::build, 1u>, amrex::Box* const> > >(int, int, unsigned long, CUstream_st*, __nv_dl_wrapper_t<__nv_dl_tag<void (amrex::ReduceOps<amrex::ReduceOpMin, amrex::ReduceOpMin, amrex::ReduceOpMin, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax>::)(int, amrex::ReduceData<int, int, int, int, int, int, int, int, int>&, __nv_dl_wrapper_t<__nv_dl_tag<void (amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::build, 1u>, amrex::Box* const>&&), &(void amrex::ReduceOps<amrex::ReduceOpMin, amrex::ReduceOpMin, amrex::ReduceOpMin, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax>::eval<int, amrex::ReduceData<int, int, int, int, int, int, int, int, int>, __nv_dl_wrapper_t<__nv_dl_tag<void (amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::build, 1u>, amrex::Box const>, void>(int, amrex::ReduceData<int, int, int, int, int, int, int, int, int>&, __nv_dl_wrapper_t<__nv_dl_tag<void (amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::build, 1u>, amrex::Box const>&&)), 1u>, amrex::GpuTuple<int, int, int, int, int, int, int, int, int>, int, int, __nv_dl_wrapper_t<__nv_dl_tag<void (amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::build, 1u>, amrex::Box* const> >&&)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Base/AMReX_GpuLaunchFunctsG.H:755:1

7: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x43ce54]
void amrex::ReduceOps<amrex::ReduceOpMin, amrex::ReduceOpMin, amrex::ReduceOpMin, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax, amrex::ReduceOpMax>::eval<int, amrex::ReduceData<int, int, int, int, int, int, int, int, int>, __nv_dl_wrapper_t<__nv_dl_tag<void (amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::build, 1u>, amrex::Box const>, void>(int, amrex::ReduceData<int, int, int, int, int, int, int, int, int>&, __nv_dl_wrapper_t<__nv_dl_tag<void (amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::build, 1u>, amrex::Box const>&&)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Base/AMReX_Reduce.H:656:14

8: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x43c0a5]
amrex::ParticleLocator<amrex::DenseBinsamrex::Box >::build(amrex::BoxArray const&, amrex::Geometry const&)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Particle/AMReX_ParticleLocator.H:142:35

9: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x475b15]
amrex::AmrParticleLocator<amrex::DenseBinsamrex::Box >::build(amrex::Vector<amrex::BoxArray, std::allocatoramrex::BoxArray > const&, amrex::Vector<amrex::Geometry, std::allocatoramrex::Geometry > const&)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Particle/AMReX_ParticleLocator.H:274:40

10: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x45ce7d]
amrex::AmrParticleLocator<amrex::DenseBinsamrex::Box >::build(amrex::ParGDBBase const*)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Particle/AMReX_ParticleLocator.H:291:20

11: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x43f0d2]
amrex::ParticleContainer<0, 0, 28, 0, amrex::ArenaAllocator>::RedistributeGPU(int, int, int, int, bool)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Particle/AMReX_ParticleContainerI.H:1244:33

12: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x42c6b5]
amrex::ParticleContainer<0, 0, 28, 0, amrex::ArenaAllocator>::Redistribute(int, int, int, int, bool)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/../../amrex/Src/Particle/AMReX_ParticleContainerI.H:1101:1

13: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x55ece6]
StructureParticleContainer::InitParticles(bool, int, double)
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/..//Source/StructureParticleContainer.H:155:62

14: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x556ceb]
main
/scratch/ov50/lw5366/AMReXLBM/LBMGPU/Exec/..//Source/main.cpp:34:9

15: /lib64/libc.so.6(__libc_start_main+0xe5) [0x14bbd7205d85]

16: main3d.gnu.DEBUG.MPI.CUDA.ex() [0x40cebe]
_start at ??:?

I am using CUDA/12.0.0, and only 1 device is used. The device is initialized successfully.
As I am running the code on a remote computer and the cuda-memcheck is not available, will try to negotiate to see if possible.

I appreciate if you have any clues in fixing the issue.

atmyers · 2023-05-25T18:12:26Z

atmyers
May 25, 2023
Maintainer

IME "invalid device symbol" usually means the code has been compiled for the wrong GPU architecture. What model GPU are you using? Do you know what architecture flags are being passed to the compiler, e.g. arch=compute_70 -code=sm_70? Are you able to compile and run some other CUDA code, like a test program?

1 reply

lw5366 May 26, 2023
Author

Thanks for your reponse!
I am using Tesla V100-SXM2-32GB, and have complied and run other codes with -arch=sm_70. From the compiling log, I can see it is using "-arch=compute_70 -code=sm_70", so I believe it is correct?
I am just wondering if you are available recently, and if possible to schdule a meeting for a qucik discussion? Thanks!

WeiqunZhang · 2023-07-13T18:28:29Z

WeiqunZhang
Jul 13, 2023
Maintainer

The issue has been resolved.

0 replies

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Make a CPU code work on GPU in AMRex #3329

{{title}}

{{editor}}'s edit

{{editor}}'s edit

Replies: 3 comments 4 replies

{{title}}

{{title}}

{{title}}

{{title}}

{{editor}}'s edit

{{editor}}'s edit

{{title}}

{{title}}

{{editor}}'s edit

{{editor}}'s edit

{{title}}

Select a reply

Make a CPU code work on GPU in AMRex #3329

lw5366 May 24, 2023

Replies: 3 comments · 4 replies

atmyers May 24, 2023 Maintainer

lw5366 May 25, 2023 Author

atmyers May 25, 2023 Maintainer

lw5366 May 25, 2023 Author

atmyers May 25, 2023 Maintainer

lw5366 May 26, 2023 Author

WeiqunZhang Jul 13, 2023 Maintainer

lw5366
May 24, 2023

Replies: 3 comments 4 replies

atmyers
May 24, 2023
Maintainer

lw5366 May 25, 2023
Author

atmyers May 25, 2023
Maintainer

lw5366 May 25, 2023
Author

atmyers
May 25, 2023
Maintainer

lw5366 May 26, 2023
Author

WeiqunZhang
Jul 13, 2023
Maintainer