CUDA API error in ParticleContainer<...>::Redistribute() + workaround

[mirenradia]

Disclaimer

This is more of an issue than a discussion but I don't think this is really a problem with AMReX itself and probably something related to the software stack of the system I was running on. Nevertheless, I thought it might be helpful to document the issue and my workaround in case anybody else encounters it or something similar hence why I have opened this discussion.

I spent quite a long time trying to debug this problem and didn't get anywhere so unless someone is confident they know what is going wrong, I don't want to invest much more time in trying to figure out the problem. I'm mainly writing this here just in case anybody else encounters a similar issue.

Summary

I was trying to use the PunctureTracker capabilities of SpacetimeX/CarpetX. When the code called Redistribute on an amrex::ParticleContainer, I got the following CUDA API error:

warning: Cuda API error detected: cudaLaunchKernel returned (0x1)

warning: Cuda API error detected: cudaPeekAtLastError returned (0x1)

warning: Cuda API error detected: cudaPeekAtLastError returned (0x1)

warning: Cuda API error detected: cudaGetLastError returned (0x1)

terminate called after throwing an instance of 'std::runtime_error'
  what():  GPU error in file /path/to/AMReX/include/AMReX_Scan.H line 1242 invalid argument

For more details, see this issue I opened in the Einstein Toolkit issue tracker.

Investigation

Running through cuda-gdb, I got the following backtrace (only AMReX part included). See above issue for more complete backtrace.

Open backtrace

#0  0x000015554c8de52f in raise () from /lib64/libc.so.6
#1  0x000015554c8b1e65 in abort () from /lib64/libc.so.6
#2  0x000015554d176bd9 in __gnu_cxx::__verbose_terminate_handler ()
    at ../../../../libstdc++-v3/libsupc++/vterminate.cc:95
#3  0x000015554d18225a in __cxxabiv1::__terminate (handler=<optimized out>)
    at ../../../../libstdc++-v3/libsupc++/eh_terminate.cc:48
#4  0x000015554d1822c5 in std::terminate ()
    at ../../../../libstdc++-v3/libsupc++/eh_terminate.cc:58
#5  0x000015554d182517 in __cxa_throw (obj=<optimized out>, tinfo=0x337e8dc0 <typeinfo for std::runtime_error@@GLIBCXX_3.4>, dest=0x41da90 <std::runtime_error::~runtime_error()@plt>)
    at ../../../../libstdc++-v3/libsupc++/eh_throw.cc:98
#6  0x00000000004b290b in amrex::Error_host (
    msg=0x4256d8e0 "GPU error in file /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/external/AMReX/include/AMReX_Scan.H line 1242 invalid argument", type=0x2d4fba9 "Abort")
    at /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/build/AMReX/amrex-24.10/Src/Base/AMReX.cpp:242
#7  amrex::Abort (
    msg=0x4256d8e0 "GPU error in file /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/external/AMReX/include/AMReX_Scan.H line 1242 invalid argument")
    at /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/build/AMReX/amrex-24.10/Src/Base/AMReX.H:159
#8  amrex::Abort (msg=...)
    at /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/build/AMReX/amrex-24.10/Src/Base/AMReX.cpp:223
#9  0x00000000010de45f in amrex::Scan::ExclusiveSum<long, int, void> (a_ret_sum=..., out=<optimized out>, in=0x15461cd98000, n=<optimized out>)
    at /home/y07/shared/utils/core/gcc/12.2.0/include/c++/12.2.0/bits/new_allocator.h:90
#10 amrex::Gpu::exclusive_scan<int*, int*> (result=<optimized out>, end=<optimized out>, begin=0x15461cd98000)
    at /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/external/AMReX/include/AMReX_Scan.H:1381
#11 amrex::DenseBins<amrex::BoxND<3> >::build<int, __nv_dl_wrapper_t<__nv_dl_tag<void (amrex::DenseBins<amrex::BoxND<3> >::*)(amrex::BinPolicy::GPUBinPolicy, int, amrex::BoxND<3> const*, amrex::BoxND<3> const&, __nv_dl_wrapper_t<__nv_dl_trailing_return_tag<void (amrex::ParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::*)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::build, amrex::IntVectND<3>, 2u>, amrex::IntVectND<3>, amrex::IntVectND<3> > const&), &amrex::DenseBins<amrex::BoxND<3> >::build<int, __nv_dl_wrapper_t<__nv_dl_trailing_return_tag<void (amrex::ParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::*)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::build, amrex::IntVectND<3>, 2u>, amrex::IntVectND<3>, amrex::IntVectND<3> > >, 1u>, __nv_dl_wrapper_t<__nv_dl_trailing_return_tag<void (amrex::ParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::*)(amrex::BoxArray const&, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::build, amrex::IntVectND<3>, 2u>, amrex::IntVectND<3>, amrex::IntVectND<3> > const, amrex::Dim3 const, amrex::Dim3 const> > (f=..., nbins=<optimized out>, v=<optimized out>, nitems=1, this=0x4337d478)
    at /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/external/AMReX/include/AMReX_DenseBins.H:252
#12 amrex::DenseBins<amrex::BoxND<3> >::build<int, __nv_dl_wrapper_t<__nv_dl_trailing_return_tag<void (amrex::ParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::*)(amrex::BoxArray const &, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::build, amrex::IntVectND<3>, 2u>, amrex::IntVectND<3>, amrex::IntVectND<3> > > (f=..., bx=<synthetic pointer>..., v=<optimized out>, nitems=1, this=0x4337d478)
    at /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/external/AMReX/include/AMReX_DenseBins.H:189
#13 amrex::DenseBins<amrex::BoxND<3> >::build<int, __nv_dl_wrapper_t<__nv_dl_trailing_return_tag<void (amrex::ParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::*)(amrex::BoxArray const &, amrex::Geometry const&), &amrex::ParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::build, amrex::IntVectND<3>, 2u>, amrex::IntVectND<3>, amrex::IntVectND<3> > > (f=..., bx=<synthetic pointer>..., v=<optimized out>, nitems=1, this=0x4337d478)
    at /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/external/AMReX/include/AMReX_Dens eBins.H:132
#14 amrex::ParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::build (this=0x4337d300, ba=..., geom=...)
    at /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/external/AMReX/include/AMReX_ParticleLocator.H:170
#15 0x00000000010e9982 in amrex::AmrParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::buil d (this=this@entry=0x4337cc58, a_ba=..., a_geom=...)
    at /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/external/AMReX/include/AMReX_ParticleLocator.H:285
#16 0x00000000010f5953 in amrex::AmrParticleLocator<amrex::DenseBins<amrex::BoxND<3> > >::build (a_gdb=<optimized out>, this=0x4337cc58)
    at /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/external/AMReX/include/AMReX_ParticleLocator.H:310
#17 amrex::ParticleContainer_impl<amrex::Particle<3, 2>, 0, 0, amrex::ArenaAllocator, amrex::DefaultAssignor>::RedistributeGPU (this=0x4337cc50, lev_min=0, lev_max=1, nGrow=0, local=0, remove_negative=true)
    at /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/Cactus-CarpetX-main/configs/sim/scratch/external/AMReX/include/AMReX_ParticleContainerI.H:1287
#18 0x00000000010c13f5 in amrex::ParticleContainer_impl<amrex::Particle<3, 2>, 0, 0, amrex::ArenaAllocator, amrex::DefaultAssignor>::Redistribute (remove_negative=true, local=0, nGrow=0, lev_max=-1, lev_min=0, this=0x4337cc50)
    at /mnt/lustre/tursafs1/home/dp415/dp415/dc-radi1/ETK/C actus-CarpetX-main/configs/sim/scratch/external/AMReX/include/AMReX_ParticleContainerI.H:1073

The problem seems to be coming from the call to cub::DeviceScan::ExclusiveSum() here.

I tried several things but none of them worked. These include:

• Running with synchronous kernel execution: CUDA_LAUNCH_BLOCKING=1.
• Running through the CUDA compute sanitizer. This gave me a bit more of the stack where the API failed but nothing really helpful.
• Changing CUDA versions. I was originally using CUDA 12.3.0 and tried 12.6.2.
• Upgrading AMReX. I was originally using 24.10 as that is the default version for this code but I tried 25.10 which was the latest tagged release at the time.

Workaround

In the end, the only thing that worked was removing/commenting out the codepath that used the cub functions and instead relying on this codepath at the end of the function.

System/software information

• System: tursa
• OS: Red Hat Enterprise Linux 8.10
• GPU: Nvidia A100
• Nvidia Driver version: 560.35.03 (CUDA Version 12.6)
• CUDA: 12.3.0 (also 12.6.2)
• AMReX: 24.10 (also 25.10)

[WeiqunZhang]

@mirenradia Thanks for writing this up! I wonder if you want to make it easier to work around the issue by changing the macro to

#if defined(AMREX_USE_CUDA) && defined(__CUDACC__) && (__CUDACC_VER_MAJOR__ >= 11) && !defined(AMREX_NO_CUB_DEVICESCAN)

If you want to do that, please submit a PR. We probably want it for both ExclusiveSum and InclusiveSum.

[mirenradia]

For now, given I've only observed the issue in one code on a single system, I don't think it's worth introducing a new build config option. If anyone else encounters the issue, then maybe I will reconsider.