Tpetra::MultiVector pack & unpack uses atomic ops even for Kokkos::Serial
Created by: ibaned
I'm separating this issue out from kokkos/kokkos#549. @mhoemmen wrote:
If Kokkos::Serial is Tpetra's execution space, Tpetra does not use atomic_assign by default. If Kokkos::OpenMP is Tpetra's execution space, but OpenMP is only using 1 thread, then Tpetra does use atomic_assign by default.
I'm pretty sure Kokkos::Serial
is the only Kokkos execution space I have enabled.
I then get this Valgrind report from a parallel Albany run:
==26619== Conditional jump or move depends on uninitialised value(s)
==26619== at 0x1F8B76DE: void Kokkos::parallel_for<Tpetra::KokkosRefactor::Details::UnpackArrayMultiColumn<Kokkos::View<double**, Kokkos::LayoutLeft, Kokkos::Serial, void>, Kokkos::View<double const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void>, Kokkos::View<int const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void>, Tpetra::KokkosRefactor::Details::InsertOp> >(unsigned long, Tpetra::KokkosRefactor::Details::UnpackArrayMultiColumn<Kokkos::View<double**, Kokkos::LayoutLeft, Kokkos::Serial, void>, Kokkos::View<double const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void>, Kokkos::View<int const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void>, Tpetra::KokkosRefactor::Details::InsertOp> const&, std::string const&) (Kokkos_Atomic_Exchange.hpp:282)
==26619== by 0x1F8B78BE: Tpetra::KokkosRefactor::Details::UnpackArrayMultiColumn<Kokkos::View<double**, Kokkos::LayoutLeft, Kokkos::Serial, void>, Kokkos::View<double const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void>, Kokkos::View<int const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void>, Tpetra::KokkosRefactor::Details::InsertOp>::unpack(Kokkos::View<double**, Kokkos::LayoutLeft, Kokkos::Serial, void> const&, Kokkos::View<double const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void> const&, Kokkos::View<int const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void> const&, Tpetra::KokkosRefactor::Details::InsertOp const&, unsigned long) (Tpetra_KokkosRefactor_Details_MultiVectorDistObjectKernels.hpp:746)
==26619== by 0x1F980D22: Tpetra::MultiVector<double, int, long long, Kokkos::Compat::KokkosDeviceWrapperNode<Kokkos::Serial, Kokkos::HostSpace>, false>::unpackAndCombineNew(Kokkos::DualView<int const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void> const&, Kokkos::DualView<double const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void> const&, Kokkos::DualView<unsigned long const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void> const&, unsigned long, Tpetra::Distributor&, Tpetra::CombineMode) (Tpetra_KokkosRefactor_Details_MultiVectorDistObjectKernels.hpp:897)
==26619== by 0x1FBB7B28: Tpetra::DistObject<double, int, long long, Kokkos::Compat::KokkosDeviceWrapperNode<Kokkos::Serial, Kokkos::HostSpace>, false>::doTransferNew(Tpetra::SrcDistObject const&, Tpetra::CombineMode, unsigned long, Kokkos::DualView<int const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void> const&, Kokkos::DualView<int const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void> const&, Kokkos::DualView<int const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void> const&, Kokkos::DualView<int const*, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, void, void> const&, Tpetra::Distributor&, Tpetra::DistObject<double, int, long long, Kokkos::Compat::KokkosDeviceWrapperNode<Kokkos::Serial, Kokkos::HostSpace>, false>::ReverseOption, bool) (Tpetra_DistObject_def.hpp:1210)
==26619== by 0x1FBB4B5B: Tpetra::DistObject<double, int, long long, Kokkos::Compat::KokkosDeviceWrapperNode<Kokkos::Serial, Kokkos::HostSpace>, false>::doTransfer(Tpetra::SrcDistObject const&, Tpetra::CombineMode, unsigned long, Teuchos::ArrayView<int const> const&, Teuchos::ArrayView<int const> const&, Teuchos::ArrayView<int const> const&, Teuchos::ArrayView<int const> const&, Tpetra::Distributor&, Tpetra::DistObject<double, int, long long, Kokkos::Compat::KokkosDeviceWrapperNode<Kokkos::Serial, Kokkos::HostSpace>, false>::ReverseOption) (Tpetra_DistObject_def.hpp:432)
==26619== by 0x1FBB3759: Tpetra::DistObject<double, int, long long, Kokkos::Compat::KokkosDeviceWrapperNode<Kokkos::Serial, Kokkos::HostSpace>, false>::doImport(Tpetra::SrcDistObject const&, Tpetra::Import<int, long long, Kokkos::Compat::KokkosDeviceWrapperNode<Kokkos::Serial, Kokkos::HostSpace> > const&, Tpetra::CombineMode) (Tpetra_DistObject_def.hpp:271)
==26619== by 0x6310C96: AAdapt::AdaptiveSolutionManagerT::AdaptiveSolutionManagerT(Teuchos::RCP<Teuchos::ParameterList> const&, Teuchos::RCP<Tpetra::Vector<double, int, long long, Kokkos::Compat::KokkosDeviceWrapperNode<Kokkos::Serial, Kokkos::HostSpace>, false> const> const&, Teuchos::RCP<Sacado::ScalarParameterLibrary<SPL_Traits> > const&, Albany::StateManager const&, Teuchos::RCP<AAdapt::rc::Manager> const&, Teuchos::RCP<Teuchos::Comm<int> const> const&) (AAdapt_AdaptiveSolutionManagerT.cpp:127)
The actual uninitialized value issue is from this comparison in Kokkos::atomic_assign
, which is one of the few times comparing to an uninitialized value makes sense, so technically a false positive, nothing wrong with Kokkos.
The second line points to this code in Tpetra, which is calling Kokkos::atomic_assign
.
I'm just creating this issue in case Tpetra should not be doing this for performance reasons, but besides the Valgrind output it doesn't affect me much.