ALM crash on GPUs
Open
@jhuang2601 is already working on this.
Since Apr 2, 2025.
type: bug
type: new
- Dominant language
- C++
- Stars
- 287
- Forks
- 109
- Avg merge
- 4d 41m
- Merged PRs (30d)
- 5
Description
Describe the bug
Crash involving ALM on GPUs
To Reproduce
mpirun -np 1 ./geosx -i ALM_SingleFracCompression_benchmark.xml
reproducer.tar.gz
Platform (please complete the following information):
- Machine: lassen
- Compiler: clang-13-cuda-12
- GEOS Version: develop
Additional context
========= Invalid __global__ read of size 4 bytes
========= at void cub::DeviceRadixSortSingleTileKernel<cub::DeviceRadixSortPolicy<int, int, unsigned int>::Policy800, (bool)0, int, int, unsigned int>(const T3 *, T3 *, const T4 *, T4 *, T5, int, int)+0x2b0
========= by thread (0,0,0) in block (0,0,0)
========= Address 0x5d3b6ea59590 is out of bounds
========= and is 24,746,818,955,888 bytes before the nearest allocation at 0x73bd40000000 of size 33,554,432 bytes
========= Saved host backtrace up to driver entry point at kernel launch time
========= Host Frame: [0x2f26ef] in libcuda.so.1
========= Host Frame: __cudart798 [0x1e1690a] in libsolidMechanicsSolvers.so
========= Host Frame: cudaLaunchKernel [0x1e7280a] in libsolidMechanicsSolvers.so
========= Host Frame: void cub::DeviceRadixSortSingleTileKernel<cub::DeviceRadixSortPolicy<int, int, unsigned int>::Policy800, false, int, int, unsigned int>(int const*, int*, int const*, int*, unsigned int, int, int) in dispatch_radix_sort.cuh:259 [0x10bb790] in libsolidMechanicsSolvers.so
========= Host Frame: cudaError thrust::cuda_cub::launcher::triple_chevron::doit_host<void (*)(int const*, int*, int const*, int*, unsigned int, int, int), int*, int*, int*, int*, unsigned int, int, int>(void (*)(int const*, int*, int const*, int*, unsigned int, int, int), int* const&, int* const&, int* const&, int* const&, unsigned int const&, int const&, int const&) const [clone .isra.0] in triple_chevron_launch.h:62 [0x10baadb] in libsolidMechanicsSolvers.so
========= Host Frame: cudaError cub::DeviceRadixSort::SortPairs<int, int, int>(void*, unsigned long&, cub::DoubleBuffer<int>&, cub::DoubleBuffer<int>&, int, int, int, CUstream_st*) in device_radix_sort.cuh:431 [0x10e3480] in libsolidMechanicsSolvers.so
========= Host Frame: std::enable_if<camp::concepts::all_of<RAJA::type_traits::is_execution_policy<RAJA::policy::cuda::cuda_exec_explicit<RAJA::iteration_mapping::Direct, RAJA::cuda::IndexGlobal<(RAJA::named_dim)0, 32, 0>, RAJA::cuda::MaxOccupancyConcretizer, 1ul, false> > >::value, camp::resources::v1::EventProxy<camp::resources::v1::Cuda> >::type RAJA::sort_pairs<RAJA::policy::cuda::cuda_exec_explicit<RAJA::iteration_mapping::Direct, RAJA::cuda::IndexGlobal<(RAJA::named_dim)0, 32, 0>, RAJA::cuda::MaxOccupancyConcretizer, 1ul, false>, LvArray::ArrayView<int, 1, 0, int, LvArray::ChaiBuffer> const&, LvArray::ArrayView<int, 1, 0, int, LvArray::ChaiBuffer> const&, camp::resources::v1::Cuda>(LvArray::ArrayView<int, 1, 0, int, LvArray::ChaiBuffer> const&, LvArray::ArrayView<int, 1, 0, int, LvArray::ChaiBuffer> const&) in sort.hpp:396 [0x10e6660] in libsolidMechanicsSolvers.so
========= Host Frame: geos::SolidMechanicsAugmentedLagrangianContact::createFaceTypeList(geos::DomainPartition const&)::{lambda(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, geos::MeshLevel const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > > const&)#1}::operator()(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, geos::MeshLevel const&, std::vector<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, std::allocator<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > > const&) const [clone .constprop.0] [clone .isra.0] in SolidMechanicsAugmentedLagrangianContact.cpp:1218 [0x11150ce] in libsolidMechanicsSolvers.so
========= Host Frame: geos::SolidMechanicsAugmentedLagrangianContact::createFaceTypeList(geos::DomainPartition const&) in SolidMechanicsAugmentedLagrangianContact.cpp:1164 [0x1116351] in libsolidMechanicsSolvers.so
========= Host Frame: geos::SolidMechanicsAugmentedLagrangianContact::setupSystem(geos::DomainPartition&, geos::DofManager&, LvArray::CRSMatrix<double, long long, int, LvArray::ChaiBuffer>&, geos::HypreVector&, geos::HypreVector&, bool) in SolidMechanicsAugmentedLagrangianContact.cpp:213 [0x1116640] in libsolidMechanicsSolvers.so
========= Host Frame: geos::SolidMechanicsLagrangianFEM::solverStep(double const&, double const&, int, geos::DomainPartition&) in SolidMechanicsLagrangianFEM.cpp:501 [0xbc39f4] in libsolidMechanicsSolvers.so
========= Host Frame: geos::PhysicsSolverBase::execute(double, double, int, int, double, geos::DomainPartition&) in PhysicsSolverBase.cpp:286 [0x12760c] in libphysicsSolversBase.so
========= Host Frame: geos::EventBase::execute(double, double, int, int, double, geos::DomainPartition&) in EventBase.cpp:236 [0x2fcf0] in libevents.so
========= Host Frame: geos::EventManager::run(geos::DomainPartition&) in EventManager.cpp:198 [0x5f1dc] in libevents.so
========= Host Frame: geos::GeosxState::run() in GeosxState.cpp:180 [0x5ad23] in libmainInterface.so
========= Host Frame: main in main.cpp:51 [0x278bc] in geosx
Contributor guide
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Assessment
This issue has not been assessed yet.