GEOS-DEV / GEOS-DEV/GEOS

ALM crash on GPUs

Open
#3,619 2 comments 0 reactions 4 assignees View on GitHub

@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

Open the contributing guide

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.