QR on a single matrix: valgrind reports invalid reads and writes #2328

cwsmith · 2024-09-05T19:47:07Z

Hello,

Calling SerialQR on a single matrix defined as Kokkos::View<double[16][10]> and running with the Kokkos Serial backend results in valgrind invalid read and write errors (pasted below). The reproducer is pasted below.

Interestingly, when using the CUDA backend in an expanded version of the reproducer (which includes a result comparison after applying the QR factorization via ApplyQ and Trsv) there are no obvious issues.

Note, I'm still figuring out how the QR interface works, hence the single matrix input to QR.

Am I doing anything obviously wrong here? Any help is appreciated.

reproducer

#include <KokkosBatched_QR_Decl.hpp>     //KokkosBlas::QR
#include <KokkosBatched_Util.hpp>        //KokkosBlas::Algo
#include <Kokkos_Core.hpp>

void testQR() {
  typedef Kokkos::View<double[16][10]> MatrixViewType;
  typedef Kokkos::View<double[10]> ColVectorViewType;
  typedef Kokkos::View<double[10]> ColWorkViewType;

  MatrixViewType A("A");
  ColVectorViewType t("t");
  ColWorkViewType w("w");

  // roughly following
  // kokkos-kernels/batched/dense/unit_test/Test_Batched_TeamVectorQR.hpp
  typedef KokkosBlas::Algo::QR::Unblocked AlgoTagType;
  Kokkos::parallel_for("serialQR", 1, KOKKOS_LAMBDA(int) {
        // compute the QR factorization of A and store the results in A and t
        // (tau) - see the lapack dgeqp3(...) documentation:
        // www.netlib.org/lapack/explore-html-3.6.1/dd/d9a/group__double_g_ecomputational_ga1b0500f49e03d2771b797c6e88adabbb.html
        KokkosBatched::SerialQR<AlgoTagType>::invoke(A, t, w);
      });
}

int main(int argc, char **argv) {
  Kokkos::ScopeGuard scope_gaurd(argc, argv);
  testQR();
}

kokkos and kokkos-kernels build

I'm building kokkos (develop @ c2a342b26) and kokkos-kernels (develop @ f26fbca) with the following cmake commands using GCC 12.3.0 on a RHEL9 system.

bdir=buildKokkosSerial
cmake -S kokkos -B $bdir \
  -DBUILD_SHARED_LIBS=on \
  -DCMAKE_CXX_COMPILER=g++ \
  -DKokkos_ENABLE_SERIAL=ON \
  -DKokkos_ENABLE_OPENMP=off \
  -DKokkos_ENABLE_DEBUG=off \
  -DCMAKE_INSTALL_PREFIX=$PWD/$bdir/install
cmake --build $bdir -j 24 --target install

bdir=buildKokkosKernelsSerial
cmake -S kokkos-kernels -B $bdir \
  -DCMAKE_CXX_COMPILER=g++ \
  -DKokkos_ROOT=buildKokkosSerial/install \
  -DCMAKE_INSTALL_PREFIX=$bdir/install
cmake --build $bdir -j 24 --target install

valgrind errors

==3560139== Memcheck, a memory error detector
==3560139== Copyright (C) 2002-2022, and GNU GPL'd, by Julian Seward et al.
==3560139== Using Valgrind-3.22.0 and LibVEX; rerun with -h for copyright info
==3560139== Command: ./QRTests
==3560139==
==3560139== Invalid read of size 8
==3560139==    at 0x413BF5: int KokkosBatched::SerialLeftHouseholderInternal::invoke<double>(int, double*, double*, int, double*) (KokkosBatched_Householder_Serial_Internal.hpp:48)
==3560139==    by 0x4129B9: int KokkosBatched::SerialQR_Internal::invoke<double>(int, int, double*, int, int, double*, int, double*) (KokkosBatched_QR_Serial_Internal.hpp:68)
==3560139==    by 0x4118B0: int KokkosBatched::SerialQR<KokkosBlas::Algo::Level3::Unblocked>::invoke<Kokkos::View<double [16][10]>, Kokkos::View<double [10]>, Kokkos::View<double [10]> >(Kokkos::View<double [16][10]> const&, Kokkos::View<double [10]> const&, Kokkos::View<double [10]> const&) (KokkosBatched_QR_Serial_Impl.hpp:34)
==3560139==    by 0x40F5DF: testQR()::{lambda(int)#1}::operator()(int) const (testQR.cpp:21)
==3560139==    by 0x410015: std::enable_if<is_void_v<void>, void>::type Kokkos::Impl::ParallelFor<testQR()::{lambda(int)#1}, Kokkos::RangePolicy<Kokkos::Serial>, Kokkos::Serial>::exec<void>() const (Kokkos_Serial_Parallel_Range.hpp:37)
==3560139==    by 0x40FEBF: Kokkos::Impl::ParallelFor<testQR()::{lambda(int)#1}, Kokkos::RangePolicy<Kokkos::Serial>, Kokkos::Serial>::execute() const (Kokkos_Serial_Parallel_Range.hpp:56)
==3560139==    by 0x40FC85: void Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::Serial>, testQR()::{lambda(int)#1}, void>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, Kokkos::RangePolicy<Kokkos::Serial> const&, testQR()::{lambda(int)#1} const&) (Kokkos_Parallel.hpp:146)
==3560139==    by 0x40FA98: void Kokkos::parallel_for<testQR()::{lambda(int)#1}>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, unsigned long, testQR()::{lambda(int)#1} const&) (Kokkos_Parallel.hpp:167)
==3560139==    by 0x40F786: testQR() (testQR.cpp:17)
==3560139==    by 0x40F8C2: main (testQR.cpp:27)
==3560139==  Address 0x8991400 is 0 bytes after a block of size 1,408 alloc'd
==3560139==    at 0x484615B: operator new(unsigned long, std::align_val_t, std::nothrow_t const&) (vg_replace_malloc.c:663)
==3560139==    by 0x5C8C0A9: Kokkos::HostSpace::impl_allocate(char const*, unsigned long, unsigned long, Kokkos_Profiling_SpaceHandle) const (Kokkos_HostSpace.cpp:79)
==3560139==    by 0x5C8C284: Kokkos::HostSpace::allocate(char const*, unsigned long, unsigned long) const (Kokkos_HostSpace.cpp:58)
==3560139==    by 0x5C8C767: checked_allocation_with_header<Kokkos::HostSpace> (Kokkos_SharedAlloc.hpp:203)
==3560139==    by 0x5C8C767: Kokkos::Impl::SharedAllocationRecordCommon<Kokkos::HostSpace>::SharedAllocationRecordCommon(Kokkos::HostSpace const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, unsigned long, void (*)(Kokkos::Impl::SharedAllocationRecord<void, void>*)) (Kokkos_SharedAlloc_timpl.hpp:62)
==3560139==    by 0x4150F8: Kokkos::Impl::SharedAllocationRecord<Kokkos::HostSpace, void>::SharedAllocationRecordCommon(Kokkos::HostSpace const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, unsigned long, void (*)(Kokkos::Impl::SharedAllocationRecord<void, void>*)) (Kokkos_HostSpace.hpp:178)
==3560139==    by 0x41513E: Kokkos::Impl::SharedAllocationRecord<Kokkos::HostSpace, Kokkos::Impl::ViewValueFunctor<Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, double> >::SharedAllocationRecord(Kokkos::HostSpace const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, unsigned long) (Kokkos_SharedAlloc.hpp:419)
==3560139==    by 0x4144C9: Kokkos::Impl::SharedAllocationRecord<Kokkos::HostSpace, Kokkos::Impl::ViewValueFunctor<Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, double> >::allocate(Kokkos::HostSpace const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, unsigned long) (Kokkos_SharedAlloc.hpp:434)
==3560139==    by 0x4132B8: Kokkos::Impl::SharedAllocationRecord<void, void>* Kokkos::Impl::ViewMapping<Kokkos::ViewTraits<double [16][10]>, void>::allocate_shared<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, Kokkos::HostSpace, Kokkos::Serial>(Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<cha
==3560139==    by 0x4124F2: Kokkos::View<double [16][10]>::View<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > >(Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > const&, std::enable_if<!Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char
==3560139==    by 0x4116EC: Kokkos::View<double [16][10]>::View<char [2]>(char const (&) [2], std::enable_if<Kokkos::Impl::is_view_label<char [2]>::value, unsigned long const>::type, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long) (Kokkos_ViewLegacy.hpp:1158)
==3560139==    by 0x40F666: testQR() (testQR.cpp:10)
==3560139==    by 0x40F8C2: main (testQR.cpp:27)
==3560139==
==3560139== Invalid write of size 8
==3560139==    at 0x413C62: int KokkosBatched::SerialLeftHouseholderInternal::invoke<double>(int, double*, double*, int, double*) (KokkosBatched_Householder_Serial_Internal.hpp:55)
==3560139==    by 0x4129B9: int KokkosBatched::SerialQR_Internal::invoke<double>(int, int, double*, int, int, double*, int, double*) (KokkosBatched_QR_Serial_Internal.hpp:68)
==3560139==    by 0x4118B0: int KokkosBatched::SerialQR<KokkosBlas::Algo::Level3::Unblocked>::invoke<Kokkos::View<double [16][10]>, Kokkos::View<double [10]>, Kokkos::View<double [10]> >(Kokkos::View<double [16][10]> const&, Kokkos::View<double [10]> const&, Kokkos::View<double [10]> const&) (KokkosBatched_QR_Serial_Impl.hpp:34)
==3560139==    by 0x40F5DF: testQR()::{lambda(int)#1}::operator()(int) const (testQR.cpp:21)
==3560139==    by 0x410015: std::enable_if<is_void_v<void>, void>::type Kokkos::Impl::ParallelFor<testQR()::{lambda(int)#1}, Kokkos::RangePolicy<Kokkos::Serial>, Kokkos::Serial>::exec<void>() const (Kokkos_Serial_Parallel_Range.hpp:37)
==3560139==    by 0x40FEBF: Kokkos::Impl::ParallelFor<testQR()::{lambda(int)#1}, Kokkos::RangePolicy<Kokkos::Serial>, Kokkos::Serial>::execute() const (Kokkos_Serial_Parallel_Range.hpp:56)
==3560139==    by 0x40FC85: void Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::Serial>, testQR()::{lambda(int)#1}, void>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, Kokkos::RangePolicy<Kokkos::Serial> const&, testQR()::{lambda(int)#1} const&) (Kokkos_Parallel.hpp:146)
==3560139==    by 0x40FA98: void Kokkos::parallel_for<testQR()::{lambda(int)#1}>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, unsigned long, testQR()::{lambda(int)#1} const&) (Kokkos_Parallel.hpp:167)
==3560139==    by 0x40F786: testQR() (testQR.cpp:17)
==3560139==    by 0x40F8C2: main (testQR.cpp:27)
==3560139==  Address 0x8991890 is 0 bytes after a block of size 208 alloc'd
==3560139==    at 0x484615B: operator new(unsigned long, std::align_val_t, std::nothrow_t const&) (vg_replace_malloc.c:663)
==3560139==    by 0x5C8C0A9: Kokkos::HostSpace::impl_allocate(char const*, unsigned long, unsigned long, Kokkos_Profiling_SpaceHandle) const (Kokkos_HostSpace.cpp:79)
==3560139==    by 0x5C8C284: Kokkos::HostSpace::allocate(char const*, unsigned long, unsigned long) const (Kokkos_HostSpace.cpp:58)
==3560139==    by 0x5C8C767: checked_allocation_with_header<Kokkos::HostSpace> (Kokkos_SharedAlloc.hpp:203)
==3560139==    by 0x5C8C767: Kokkos::Impl::SharedAllocationRecordCommon<Kokkos::HostSpace>::SharedAllocationRecordCommon(Kokkos::HostSpace const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, unsigned long, void (*)(Kokkos::Impl::SharedAllocationRecord<void, void>*)) (Kokkos_SharedAlloc_timpl.hpp:62)
==3560139==    by 0x4150F8: Kokkos::Impl::SharedAllocationRecord<Kokkos::HostSpace, void>::SharedAllocationRecordCommon(Kokkos::HostSpace const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, unsigned long, void (*)(Kokkos::Impl::SharedAllocationRecord<void, void>*)) (Kokkos_HostSpace.hpp:178)
==3560139==    by 0x41513E: Kokkos::Impl::SharedAllocationRecord<Kokkos::HostSpace, Kokkos::Impl::ViewValueFunctor<Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, double> >::SharedAllocationRecord(Kokkos::HostSpace const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, unsigned long) (Kokkos_SharedAlloc.hpp:419)
==3560139==    by 0x4144C9: Kokkos::Impl::SharedAllocationRecord<Kokkos::HostSpace, Kokkos::Impl::ViewValueFunctor<Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, double> >::allocate(Kokkos::HostSpace const&, std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, unsigned long) (Kokkos_SharedAlloc.hpp:434)
==3560139==    by 0x413584: Kokkos::Impl::SharedAllocationRecord<void, void>* Kokkos::Impl::ViewMapping<Kokkos::ViewTraits<double [10]>, void>::allocate_shared<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, Kokkos::HostSpace, Kokkos::Serial>(Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >
==3560139==    by 0x4126C6: Kokkos::View<double [10]>::View<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > >(Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > const&, std::enable_if<!Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >
==3560139==    by 0x4117B8: Kokkos::View<double [10]>::View<char [2]>(char const (&) [2], std::enable_if<Kokkos::Impl::is_view_label<char [2]>::value, unsigned long const>::type, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long) (Kokkos_ViewLegacy.hpp:1158)
==3560139==    by 0x40F6A2: testQR() (testQR.cpp:11)
==3560139==    by 0x40F8C2: main (testQR.cpp:27)

... snip

The text was updated successfully, but these errors were encountered:

lucbv · 2024-09-06T17:05:39Z

Thanks for reporting, we will have a look at this

cwsmith · 2024-09-06T18:00:22Z

Hi @lucbv . Thanks for looking into this.

It seems like the problem is related to handling rectangular matrices. If the MatrixViewType is defined to be 10x10:

typedef Kokkos::View<double[16][10]> MatrixViewType;

there are no errors under valgrind.

Digging into the code a bit, but without a full understanding of it, I see that this loop over matrix rows:

kokkos-kernels/batched/dense/impl/KokkosBatched_QR_Serial_Internal.hpp

Lines 56 to 60 in 2c4dd7e

    
           for (int m_atl = 0; m_atl < m; ++m_atl) { 
        
             // part 2x2 into 3x3 
        
             A_part3x3.partWithABR(A_part2x2, 1, 1); 
        
             const int m_A22 = m - m_atl - 1; 
        
             const int n_A22 = n - m_atl - 1;

that successively removes one row and one column to form the 3x3 partitioned matrix A_part3x3 via the call A_part3x3.partWithABR(A_part2x2, 1, 1);.

In the original case of the 16x10 matrix, running valgrind with the gdbserver I see that the first invalid read occurs in SerialLeftHouseholderInternal::invoke(...) when m_atl=10 in SerialQR_Internal::invoke(...). This seems to make sense as all 10 columns have been removed.

The test case in the repo for QR appears to only run with square matrices:

test without column pivoting:

kokkos-kernels/batched/dense/unit_test/Test_Batched_TeamVectorQR.hpp

Line 111 in 2c4dd7e

MatrixViewType a("a", N, BlkSize, BlkSize);

test 'WithColumnPivoting':

kokkos-kernels/batched/dense/unit_test/Test_Batched_TeamVectorQR_WithColumnPivoting.hpp

Line 121 in 2c4dd7e

MatrixViewType a("a", N, BlkSize, BlkSize);

lucbv · 2024-09-06T19:16:07Z

Okay, thanks for digging a bit into this, I will run the code in valgrind / gdb as well and hopefully can reproduce and report my observation. The algorithm indeed uses a partitioning in the matrix to perform some operations but it should still work for rectangular matrices. Once I find something promising I will let you know about it : )

The serial QR algorithms does not have unit-tests and is failing for non square matrices. See issue kokkos#2328. This first commit fixes the issue with rectangular matrices and adds a basic test for that use case. Next will work on adding a test that exercises the interfaces on multiple matrices of different sizes within a parallel_for. Finally equivalent tests will be added for the square case as well.

lucbv · 2024-09-19T20:13:25Z

The PR above, #2342, has a fix for the rectangular matrices and introduces more tests for the Serial QR feature. The tests are not fully implemented yet but the fix seems to be okay if you want to give it a try.

cwsmith · 2024-09-20T15:09:44Z

This is great. Thank you @lucbv.

Using the PR branch (9121f0a) I ran the reproducer under valgrind again and the SerialQR errors are gone.

Running the expanded version of the reproducer:
https://github.com/SCOREC/meshFields/blob/20a68919b338003ff8792ce7d2cc6c5df3f13613/test/testQR.cpp
under valgrind reports an invalid read in the call to SerialTrsv. Some additional details on the first invalid read and the valgrind log are below.

stack at first Trsv invalid read

Running under gdb reports the following values of variables at the point of the first reported invalid read.

Reading symbols from /opt/scorec/spack/rhel9/v0201_4/install/linux-rhel9-x86_64/gcc-12.3.0/libiconv-1.17-oylnknwv5m57zkfjde7op6ne3pqdkhxg/lib/libiconv.so.2...
0x0000000000427b8f in KokkosBatched::SerialTrsvInternalUpper<KokkosBlas::Algo::Level2::Unblocked>::invoke<double, double> (use_unit_diag=false, m=16, alpha=1, A=0x8991f00, as0=1, as1=16, b=0x8993940, bs0=1)
    at /space/cwsmith/meshFields/buildKokkosKernelsSerial/install/include/KokkosBatched_Trsv_Serial_Internal.hpp:161
161           if (!use_unit_diag) *beta1 = *beta1 / A[p * as0 + p * as1];
(ins)(gdb) where
#0  0x0000000000427b8f in KokkosBatched::SerialTrsvInternalUpper<KokkosBlas::Algo::Level2::Unblocked>::invoke<double, double> (use_unit_diag=false, m=16, alpha=1, A=0x8991f00, as0=1, as1=16, b=0x8993940, bs0=1)
    at /space/cwsmith/meshFields/buildKokkosKernelsSerial/install/include/KokkosBatched_Trsv_Serial_Internal.hpp:161
#1  0x0000000000423408 in KokkosBatched::SerialTrsv<KokkosBatched::Uplo::Upper, KokkosBlas::Trans::NoTranspose, KokkosBatched::Diag::NonUnit, KokkosBlas::Algo::Level2::Unblocked>::invoke<double, Kokkos::View<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> >, Kokkos::View<double [16], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> > > (alpha=1, A=..., b=...)
    at /space/cwsmith/meshFields/buildKokkosKernelsSerial/install/include/KokkosBatched_Trsv_Serial_Impl.hpp:191
#2  0x000000000041a7c8 in operator() (__closure=0x1ffefef0d0)
    at /space/cwsmith/meshFields/meshFields/test/testQR.cpp:127
#3  0x000000000041bf36 in Kokkos::Impl::ParallelFor<testSolveQR()::<lambda(int)>, Kokkos::RangePolicy<Kokkos::Serial>, Kokkos::Serial>::exec<void>(void) const (this=0x1ffefef0d0)
    at /space/cwsmith/meshFields/buildKokkosSerial/install/include/Serial/Kokkos_Serial_Parallel_Range.hpp:37
#4  0x000000000041bb86 in Kokkos::Impl::ParallelFor<testSolveQR()::<lambda(int)>, Kokkos::RangePolicy<Kokkos::Serial>, Kokkos::Serial>::execute(void) const (this=0x1ffefef0d0)
    at /space/cwsmith/meshFields/buildKokkosSerial/install/include/Serial/Kokkos_Serial_Parallel_Range.hpp:56
#5  0x000000000041b71e in Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::Serial>, testSolveQR()::<lambda(int)> >(const std::string &, const Kokkos::RangePolicy<Kokkos::Serial> &, const struct {...} &) (str=..., policy=..., functor=...)
    at /space/cwsmith/meshFields/buildKokkosSerial/install/include/Kokkos_Parallel.hpp:146
#6  0x000000000041b4c3 in Kokkos::parallel_for<testSolveQR()::<lambda(int)> >(const std::string &, size_t, const struct {...} &) (str=..., work_count=1, functor=...)
    at /space/cwsmith/meshFields/buildKokkosSerial/install/include/Kokkos_Parallel.hpp:167
#7  0x000000000041afbc in testSolveQR () at /space/cwsmith/meshFields/meshFields/test/testQR.cpp:113
#8  0x000000000041b2ed in main (argc=1, argv=0x1ffefef608) at /space/cwsmith/meshFields/meshFields/test/testQR.cpp:140
(ins)(gdb) p p
$1 = 15
(ins)(gdb) p m
$2 = 16

Given this loop from p=m-1:0

kokkos-kernels/batched/dense/impl/KokkosBatched_Trsv_Serial_Internal.hpp

Line 152 in b221005

for (int p = (m - 1); p >= 0; --p) {

and the use of p twice to compute the index into A here (where the invalid read occurs):

kokkos-kernels/batched/dense/impl/KokkosBatched_Trsv_Serial_Internal.hpp

Line 161 in b221005

if (!use_unit_diag) *beta1 = *beta1 / A[p * as0 + p * as1];

, it looks that the assumption that A is square is made here as well.

valgrind log

==265289== Memcheck, a memory error detector
==265289== Copyright (C) 2002-2022, and GNU GPL'd, by Julian Seward et al. 
==265289== Using Valgrind-3.22.0 and LibVEX; rerun with -h for copyright info
==265289== Command: ./QRTests
==265289== 
==265289== Invalid read of size 8
==265289==    at 0x427B8F: int KokkosBatched::SerialTrsvInternalUpper<KokkosBlas::Algo::Level2::Unblocked>::invoke<double, double>(bool, int, double, double const*, int, int, double*, int) (KokkosBatched_Trsv_Serial_Internal.hpp:161)
==265289==    by 0x423407: int KokkosBatched::SerialTrsv<KokkosBatched::Uplo::Upper, KokkosBlas::Trans::NoTranspose, KokkosBatched::Diag::NonUnit, KokkosBlas::Algo::Level2::Unblocked>::invoke<double, Kokkos::View<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> >, Kokkos::View<double [16], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> > >(double, Kokkos::View<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> > const&, Kokkos::View<double [16], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> > const&) (KokkosBatched_Trsv_Serial_Impl.hpp:191)
==265289==    by 0x41A7C7: testSolveQR()::{lambda(int)#1}::operator()(int) const (testQR.cpp:127)
==265289==    by 0x41BF35: std::enable_if<is_void_v<void>, void>::type Kokkos::Impl::ParallelFor<testSolveQR()::{lambda(int)#1}, Kokkos::RangePolicy<Kokkos::Serial>, Kokkos::Serial>::exec<void>() const (Kokkos_Serial_Parallel_Range.hpp:37)
==265289==    by 0x41BB85: Kokkos::Impl::ParallelFor<testSolveQR()::{lambda(int)#1}, Kokkos::RangePolicy<Kokkos::Serial>, Kokkos::Serial>::execute() const (Kokkos_Serial_Parallel_Range.hpp:56)
==265289==    by 0x41B71D: void Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::Serial>, testSolveQR()::{lambda(int)#1}, void>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, Kokkos::RangePolicy<Kokkos::Serial> const&, testSolveQR()::{lambda(int)#1} const&) (Kokkos_Parallel.hpp:146)
==265289==    by 0x41B4C2: void Kokkos::parallel_for<testSolveQR()::{lambda(int)#1}>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, unsigned long, testSolveQR()::{lambda(int)#1} const&) (Kokkos_Parallel.hpp:167)
==265289==    by 0x41AFBB: testSolveQR() (testQR.cpp:113)
==265289==    by 0x41B2EC: main (testQR.cpp:140)
==265289==  Address 0x89926f8 is 16 bytes after a block of size 40 free'd
==265289==    at 0x484893D: operator delete(void*, unsigned long) (vg_replace_malloc.c:1101)
==265289==    by 0x423D0C: Kokkos::Impl::HostSharedPtr<Kokkos::Impl::SerialInternal>::cleanup() (Kokkos_HostSharedPtr.hpp:120)
==265289==    by 0x41DDA1: Kokkos::Impl::HostSharedPtr<Kokkos::Impl::SerialInternal>::~HostSharedPtr() (Kokkos_HostSharedPtr.hpp:92)
==265289==    by 0x41CC4F: Kokkos::Serial::~Serial() (Kokkos_Serial.hpp:95)
==265289==    by 0x423E87: Kokkos::Impl::ViewCtorProp<void, Kokkos::Serial>::~ViewCtorProp() (Kokkos_ViewCtor.hpp:122)
==265289==    by 0x423EA7: Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >, Kokkos::HostSpace, Kokkos::Serial>::~ViewCtorProp() (Kokkos_ViewCtor.hpp:182)
==265289==    by 0x425342: Kokkos::View<double [10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> >::View<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > >(Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > const&, std::enable_if<!Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > >::has_pointer, Kokkos::LayoutLeft>::type const&) (Kokkos_ViewLegacy.hpp:1048)
==265289==    by 0x41F9BC: Kokkos::View<double [10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> >::View<char [3]>(char const (&) [3], std::enable_if<Kokkos::Impl::is_view_label<char [3]>::value, unsigned long const>::type, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long) (Kokkos_ViewLegacy.hpp:1158)
==265289==    by 0x41AC8F: testSolveQR() (testQR.cpp:89)
==265289==    by 0x41B2EC: main (testQR.cpp:140)
==265289==  Block was alloc'd at
==265289==    at 0x4844F95: operator new(unsigned long) (vg_replace_malloc.c:483)
==265289==    by 0x5C94902: HostSharedPtr<Kokkos::Serial::Serial()::<lambda(Kokkos::Impl::SerialInternal*)> > (Kokkos_HostSharedPtr.hpp:47)
==265289==    by 0x5C94902: Kokkos::Serial::Serial() (Kokkos_Serial.cpp:160)
==265289==    by 0x4251A9: Kokkos::View<double [10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> >::View<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > >(Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > > const&, std::enable_if<!Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > >::has_pointer, Kokkos::LayoutLeft>::type const&) (Kokkos_ViewLegacy.hpp:1002)
==265289==    by 0x41F9BC: Kokkos::View<double [10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> >::View<char [3]>(char const (&) [3], std::enable_if<Kokkos::Impl::is_view_label<char [3]>::value, unsigned long const>::type, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long) (Kokkos_ViewLegacy.hpp:1158)
==265289==    by 0x41AC8F: testSolveQR() (testQR.cpp:89)
==265289==    by 0x41B2EC: main (testQR.cpp:140)
==265289== 
==265289== Invalid read of size 8
==265289==    at 0x427BDC: int KokkosBatched::SerialTrsvInternalUpper<KokkosBlas::Algo::Level2::Unblocked>::invoke<double, double>(bool, int, double, double const*, int, int, double*, int) (KokkosBatched_Trsv_Serial_Internal.hpp:163)
==265289==    by 0x423407: int KokkosBatched::SerialTrsv<KokkosBatched::Uplo::Upper, KokkosBlas::Trans::NoTranspose, KokkosBatched::Diag::NonUnit, KokkosBlas::Algo::Level2::Unblocked>::invoke<double, Kokkos::View<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> >, Kokkos::View<double [16], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> > >(double, Kokkos::View<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> > const&, Kokkos::View<double [16], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> > const&) (KokkosBatched_Trsv_Serial_Impl.hpp:191)
==265289==    by 0x41A7C7: testSolveQR()::{lambda(int)#1}::operator()(int) const (testQR.cpp:127)
==265289==    by 0x41BF35: std::enable_if<is_void_v<void>, void>::type Kokkos::Impl::ParallelFor<testSolveQR()::{lambda(int)#1}, Kokkos::RangePolicy<Kokkos::Serial>, Kokkos::Serial>::exec<void>() const (Kokkos_Serial_Parallel_Range.hpp:37)
==265289==    by 0x41BB85: Kokkos::Impl::ParallelFor<testSolveQR()::{lambda(int)#1}, Kokkos::RangePolicy<Kokkos::Serial>, Kokkos::Serial>::execute() const (Kokkos_Serial_Parallel_Range.hpp:56)
==265289==    by 0x41B71D: void Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::Serial>, testSolveQR()::{lambda(int)#1}, void>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, Kokkos::RangePolicy<Kokkos::Serial> const&, testSolveQR()::{lambda(int)#1} const&) (Kokkos_Parallel.hpp:146)
==265289==    by 0x41B4C2: void Kokkos::parallel_for<testSolveQR()::{lambda(int)#1}>(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&, unsigned long, testSolveQR()::{lambda(int)#1} const&) (Kokkos_Parallel.hpp:167)
==265289==    by 0x41AFBB: testSolveQR() (testQR.cpp:113)
==265289==    by 0x41B2EC: main (testQR.cpp:140)
==265289==  Address 0x8992680 is 6 bytes after a block of size 74 free'd
==265289==    at 0x484893D: operator delete(void*, unsigned long) (vg_replace_malloc.c:1101)
==265289==    by 0x5C94C5D: deallocate (new_allocator.h:158)
==265289==    by 0x5C94C5D: deallocate (alloc_traits.h:496)
==265289==    by 0x5C94C5D: _M_destroy (basic_string.h:300)
==265289==    by 0x5C94C5D: _M_dispose (basic_string.h:294)
==265289==    by 0x5C94C5D: ~basic_string (basic_string.h:803)
==265289==    by 0x5C94C5D: profile_fence_event<Kokkos::Serial, Kokkos::Serial::impl_static_fence(const std::string&)::<lambda()> > (Kokkos_Profiling.hpp:219)
==265289==    by 0x5C94C5D: impl_static_fence (Kokkos_Serial.hpp:147)
==265289==    by 0x5C94C5D: Kokkos::Impl::ExecSpaceDerived<Kokkos::Serial>::static_fence(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) (Kokkos_ExecSpaceManager.hpp:131)
==265289==    by 0x5C85684: Kokkos::Impl::ExecSpaceManager::static_fence(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) (Kokkos_Core.cpp:243)
==265289==    by 0x41E8BB: void Kokkos::deep_copy<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, Kokkos::Experimental::EmptyViewHooks>(Kokkos::View<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> > const&, Kokkos::View<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, Kokkos::Experimental::EmptyViewHooks> const&, std::enable_if<((is_void_v<Kokkos::ViewTraits<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> >::specialize>)&&(is_void_v<Kokkos::ViewTraits<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, Kokkos::Experimental::EmptyViewHooks>::specialize>))&&((((unsigned int)Kokkos::ViewTraits<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> >::rank)!=(0))||(((unsigned int)Kokkos::ViewTraits<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, Kokkos::Experimental::EmptyViewHooks>::rank)!=(0))), void>::type*) (Kokkos_CopyViews.hpp:1709)
==265289==    by 0x41AC57: testSolveQR() (testQR.cpp:87)
==265289==    by 0x41B2EC: main (testQR.cpp:140)
==265289==  Block was alloc'd at
==265289==    at 0x4844F95: operator new(unsigned long) (vg_replace_malloc.c:483)
==265289==    by 0x5C943FB: void std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> >::_M_construct<char*>(char*, char*, std::forward_iterator_tag) [clone .isra.0] (basic_string.tcc:225)
==265289==    by 0x5C94C36: basic_string (basic_string.h:552)
==265289==    by 0x5C94C36: profile_fence_event<Kokkos::Serial, Kokkos::Serial::impl_static_fence(const std::string&)::<lambda()> > (Kokkos_Profiling.hpp:219)
==265289==    by 0x5C94C36: impl_static_fence (Kokkos_Serial.hpp:147)
==265289==    by 0x5C94C36: Kokkos::Impl::ExecSpaceDerived<Kokkos::Serial>::static_fence(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) (Kokkos_ExecSpaceManager.hpp:131)
==265289==    by 0x5C85684: Kokkos::Impl::ExecSpaceManager::static_fence(std::__cxx11::basic_string<char, std::char_traits<char>, std::allocator<char> > const&) (Kokkos_Core.cpp:243)
==265289==    by 0x41E8BB: void Kokkos::deep_copy<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, Kokkos::Experimental::EmptyViewHooks>(Kokkos::View<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> > const&, Kokkos::View<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, Kokkos::Experimental::EmptyViewHooks> const&, std::enable_if<((is_void_v<Kokkos::ViewTraits<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> >::specialize>)&&(is_void_v<Kokkos::ViewTraits<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, Kokkos::Experimental::EmptyViewHooks>::specialize>))&&((((unsigned int)Kokkos::ViewTraits<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace> >::rank)!=(0))||(((unsigned int)Kokkos::ViewTraits<double [16][10], Kokkos::LayoutLeft, Kokkos::Device<Kokkos::Serial, Kokkos::HostSpace>, Kokkos::Experimental::EmptyViewHooks>::rank)!=(0))), void>::type*) (Kokkos_CopyViews.hpp:1709)
==265289==    by 0x41AC57: testSolveQR() (testQR.cpp:87)
==265289==    by 0x41B2EC: main (testQR.cpp:140)
==265289== 
==265289== 
==265289== HEAP SUMMARY:
==265289==     in use at exit: 176 bytes in 1 blocks
==265289==   total heap usage: 230 allocs, 229 frees, 115,919 bytes allocated
==265289== 
==265289== LEAK SUMMARY:
==265289==    definitely lost: 0 bytes in 0 blocks
==265289==    indirectly lost: 0 bytes in 0 blocks
==265289==      possibly lost: 0 bytes in 0 blocks
==265289==    still reachable: 176 bytes in 1 blocks
==265289==         suppressed: 0 bytes in 0 blocks
==265289== Rerun with --leak-check=full to see details of leaked memory
==265289== 
==265289== For lists of detected and suppressed errors, rerun with: -s
==265289== ERROR SUMMARY: 80 errors from 2 contexts (suppressed: 0 from 0)

lucbv · 2024-09-20T20:09:24Z

Okay, I will try to wrap up the PR and get that tested and merged, then I can move on to trsv, hopefully it's not more complicated than the QR fix but writing proper tests is what takes time!

lucbv · 2024-09-22T18:47:55Z

So I have not looked at it in detail but my guess is that we are assuming the triangular matrix to be stored in a square matrix, size mxm. Since yours is coming for the QR factorization of a rectangular matrix we need to fix the code so that it works for a mxn input matrix, basically we will ignore the non-square part of the input. Should not be too bad hopefully. I will create a PR once I have confirmed that this is the issue and we have a fix for you...

The serial QR algorithms does not have unit-tests and is failing for non square matrices. See issue kokkos#2328. This first commit fixes the issue with rectangular matrices and adds a basic test for that use case. Next will work on adding a test that exercises the interfaces on multiple matrices of different sizes within a parallel_for. Finally equivalent tests will be added for the square case as well.

The serial QR algorithms does not have unit-tests and is failing for non square matrices. See issue kokkos#2328. This first commit fixes the issue with rectangular matrices and adds a basic test for that use case. Next will work on adding a test that exercises the interfaces on multiple matrices of different sizes within a parallel_for. Finally equivalent tests will be added for the square case as well. Signed-off-by: Luc <lberge@sandia.gov>

The serial QR algorithms does not have unit-tests and is failing for non square matrices. See issue kokkos#2328. This first commit fixes the issue with rectangular matrices and adds a basic test for that use case. Next will work on adding a test that exercises the interfaces on multiple matrices of different sizes within a parallel_for. Finally equivalent tests will be added for the square case as well. Fixing unused variable error It looks like the Left NoTranspose ApplyQ is not doing the correct thing. Will have a look at that next. Signed-off-by: Luc <lberge@sandia.gov>

The serial QR algorithms does not have unit-tests and is failing for non square matrices. See issue kokkos#2328. This first commit fixes the issue with rectangular matrices and adds a basic test for that use case. Next will work on adding a test that exercises the interfaces on multiple matrices of different sizes within a parallel_for. Finally equivalent tests will be added for the square case as well. Fixing unused variable error It looks like the Left NoTranspose ApplyQ is not doing the correct thing. Will have a look at that next. Signed-off-by: Luc <lberge@sandia.gov> Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov>

jacobmerson · 2025-03-03T16:48:44Z

@lucbv what's the status of this? We would like to make use of this.

The serial QR algorithms does not have unit-tests and is failing for non square matrices. See issue kokkos#2328. This first commit fixes the issue with rectangular matrices and adds a basic test for that use case. Next will work on adding a test that exercises the interfaces on multiple matrices of different sizes within a parallel_for. Finally equivalent tests will be added for the square case as well. Fixing unused variable error It looks like the Left NoTranspose ApplyQ is not doing the correct thing. Will have a look at that next. Spliting the tests a bit better, looking at Q and Qt together to detect issues with consistency between the two. Also eventually allows to use GEMM to figure out which one is wrong. Final fixes and adjusting the tolerance Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov>

lucbv · 2025-03-07T22:25:57Z

@jacobmerson I am cleaning it up now, it should merge on Monday

jacobmerson · 2025-03-07T22:46:12Z

Great thank you!

* batched - dense: Testing and fixing Serial QR The serial QR algorithms does not have unit-tests and is failing for non square matrices. See issue #2328. This first commit fixes the issue with rectangular matrices and adds a basic test for that use case. Next will work on adding a test that exercises the interfaces on multiple matrices of different sizes within a parallel_for. Finally equivalent tests will be added for the square case as well. Fixing unused variable error It looks like the Left NoTranspose ApplyQ is not doing the correct thing. Will have a look at that next. Spliting the tests a bit better, looking at Q and Qt together to detect issues with consistency between the two. Also eventually allows to use GEMM to figure out which one is wrong. Final fixes and adjusting the tolerance Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov> * Removing the stride for the workspace Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov> * Applying clang-format Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov> * Adding checks on inputs Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov> * More clean-ups Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov> * clang-format Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov> * Batched - QR: adding static assertion on struct templates Checking that the side, mode and algo are all valid. Cleaning up the test by using create_mirror_view_and_copy when appropriate. Somehow the semantic of create_mirror_view_and_copy is different from the regular deep_copy so need to make a few changes after testing on GPUs... Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov> * Batched - QR: adding check for apply Q on the right side. Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov> * clang-format Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov> * clang-format Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov> --------- Signed-off-by: Luc Berger-Vergiat <lberge@sandia.gov>

lucbv · 2025-03-11T20:45:50Z

The PR has been merged so you should be able to test again, please let us know if you see any issues?

lucbv self-assigned this Sep 6, 2024

lucbv added the bug label Sep 6, 2024

lucbv mentioned this issue Sep 19, 2024

batched - dense: Testing and fixing Serial QR #2342

Merged

16 tasks

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

QR on a single matrix: valgrind reports invalid reads and writes #2328

QR on a single matrix: valgrind reports invalid reads and writes #2328

cwsmith commented Sep 5, 2024 •

edited

Loading

lucbv commented Sep 6, 2024

cwsmith commented Sep 6, 2024

lucbv commented Sep 6, 2024

lucbv commented Sep 19, 2024

cwsmith commented Sep 20, 2024 •

edited

Loading

lucbv commented Sep 20, 2024

lucbv commented Sep 22, 2024 •

edited

Loading

jacobmerson commented Mar 3, 2025

lucbv commented Mar 7, 2025

jacobmerson commented Mar 7, 2025

lucbv commented Mar 11, 2025

QR on a single matrix: valgrind reports invalid reads and writes #2328

QR on a single matrix: valgrind reports invalid reads and writes #2328

Comments

cwsmith commented Sep 5, 2024 • edited Loading

reproducer

kokkos and kokkos-kernels build

valgrind errors

lucbv commented Sep 6, 2024

cwsmith commented Sep 6, 2024

lucbv commented Sep 6, 2024

lucbv commented Sep 19, 2024

cwsmith commented Sep 20, 2024 • edited Loading

stack at first Trsv invalid read

valgrind log

lucbv commented Sep 20, 2024

lucbv commented Sep 22, 2024 • edited Loading

jacobmerson commented Mar 3, 2025

lucbv commented Mar 7, 2025

jacobmerson commented Mar 7, 2025

lucbv commented Mar 11, 2025

cwsmith commented Sep 5, 2024 •

edited

Loading

cwsmith commented Sep 20, 2024 •

edited

Loading

lucbv commented Sep 22, 2024 •

edited

Loading