review updates

Co-authored-by: Fritz Goebel <fritz.goebel@kit.edu>

core/components/disjoint_sets.hpp

@@ -91,7 +91,7 @@ class disjoint_sets {
 
     /**
      * Returns the representative of the set containing element x.
-     * Also performs path-compresson on the corresponding path.
+     * Also performs path-compression on the corresponding path.
      *
      * @param x  the element
      * @return the element representing the set containing x

core/factorization/elimination_forest.cpp

@@ -60,7 +60,7 @@ void compute_elim_forest_parent_impl(std::shared_ptr<const Executor> host_exec,
         auto row_rep = row;
         for (auto nz = row_ptrs[row]; nz < row_ptrs[row + 1]; nz++) {
             const auto col = cols[nz];
-            // for each lower diagonal entry
+            // for each lower triangular entry
             if (col < row) {
                 // find the subtree it is contained in
                 const auto col_rep = subtrees.find(col);

core/test/components/disjoint_sets.cpp

@@ -34,6 +34,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 
 #include <algorithm>
+#include <bitset>
 #include <type_traits>
 
 
@@ -125,12 +126,7 @@ TYPED_TEST(DisjointSets, SequentialMergesIntoSingleSet)
 
 int popcount(int i)
 {
-    int count{};
-    while (i) {
-        count += i & 1;
-        i >>= 1;
-    }
-    return count;
+    return std::bitset<32>(static_cast<unsigned>(i)).count();
 }
 
 

core/test/factorization/elimination_forest.cpp

@@ -139,6 +139,40 @@ TYPED_TEST(EliminationForest, WorksForSeparable)
 }
 
 
+TYPED_TEST(EliminationForest, WorksForPostOrderNotSelfInverse)
+{
+    using matrix_type = typename TestFixture::matrix_type;
+    using index_type = typename TestFixture::index_type;
+    auto mtx = gko::initialize<typename TestFixture::matrix_type>(
+        {
+            {1, 0, 1, 0, 0, 0, 0, 0, 0, 0},
+            {0, 1, 0, 0, 1, 0, 0, 0, 0, 0},
+            {1, 0, 1, 0, 0, 0, 1, 0, 0, 0},
+            {0, 0, 0, 1, 0, 0, 0, 0, 1, 0},
+            {0, 1, 0, 0, 1, 1, 0, 0, 0, 0},
+            {0, 0, 0, 0, 1, 1, 1, 0, 0, 0},
+            {0, 0, 1, 0, 0, 1, 1, 1, 0, 0},
+            {0, 0, 0, 0, 0, 0, 1, 1, 1, 0},
+            {0, 0, 0, 1, 0, 0, 0, 1, 1, 1},
+            {0, 0, 0, 0, 0, 0, 1, 0, 1, 1},
+        },
+        this->ref);
+    auto forest = gko::factorization::compute_elim_forest(mtx.get());
+    GKO_ASSERT_ARRAY_EQ(forest.parents,
+                        I<index_type>({2, 4, 6, 8, 5, 6, 7, 8, 9, 10}));
+    GKO_ASSERT_ARRAY_EQ(forest.child_ptrs,
+                        I<index_type>({0, 0, 0, 1, 1, 2, 3, 5, 6, 8, 9, 10}));
+    GKO_ASSERT_ARRAY_EQ(forest.children,
+                        I<index_type>({0, 1, 4, 2, 5, 6, 3, 7, 8, 9}));
+    GKO_ASSERT_ARRAY_EQ(forest.postorder,
+                        I<index_type>({3, 0, 2, 1, 4, 5, 6, 7, 8, 9}));
+    GKO_ASSERT_ARRAY_EQ(forest.inv_postorder,
+                        I<index_type>({1, 3, 2, 0, 4, 5, 6, 7, 8, 9}));
+    GKO_ASSERT_ARRAY_EQ(forest.postorder_parents,
+                        I<index_type>({8, 2, 6, 4, 5, 6, 7, 8, 9, 10}));
+}
+
+
 TYPED_TEST(EliminationForest, WorksForAni1)
 {
     using matrix_type = typename TestFixture::matrix_type;

cuda/factorization/cholesky_kernels.cu

@@ -109,7 +109,7 @@ void cholesky_symbolic_count(
                           postorder_cols, permutation, buffer);
         cusparse::destroy(descr);
     }
-    // count nonzeros in L
+    // count nonzeros per row of L
     {
         const auto num_blocks =
             ceildiv(num_rows, default_block_size / config::warp_size);

dev_tools/scripts/config

@@ -30,6 +30,10 @@
     - FixInclude: "common/unified/base/kernel_launch.hpp"
 - "(cuda|hip|dpcpp|omp)/base/kernel_launch_solver\."
     - FixInclude: "common/unified/base/kernel_launch_solver.hpp"
+- "(cuda|hip|dpcpp|omp)/base/kernel_launch_solver\."
+    - FixInclude: "common/unified/base/kernel_launch_solver.hpp"
+- "elimination_forest\.cpp"
+    - FixInclude: "core/factorization/elimination_forest.hpp"
 - "common/unified/.*.cpp"
     - PathIgnore: "2"
     - PathPrefix: "core"

dpcpp/factorization/cholesky_kernels.dp.cpp

@@ -96,7 +96,7 @@ void cholesky_symbolic_count(
             lower_ends[row] = lower_end;
         });
     });
-    // count nonzeros per row
+    // count nonzeros per row of L
     queue->submit([&](sycl::handler& cgh) {
         cgh.parallel_for(sycl::range<1>{num_rows}, [=](sycl::id<1> idx_id) {
             const auto row = idx_id[0];

hip/factorization/cholesky_kernels.hip.cpp

@@ -109,7 +109,7 @@ void cholesky_symbolic_count(
                            postorder_cols, permutation, buffer);
         hipsparse::destroy(descr);
     }
-    // count nonzeros in L
+    // count nonzeros per row of L
     {
         const auto num_blocks =
             ceildiv(num_rows, default_block_size / config::warp_size);

reference/test/factorization/cholesky_kernels.cpp

@@ -30,9 +30,6 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
 OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 ******************************<GINKGO LICENSE>*******************************/
 
-#include "core/factorization/cholesky_kernels.hpp"
-
-
 #include <algorithm>
 #include <memory>
 
@@ -44,6 +41,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 
 #include "core/components/prefix_sum_kernels.hpp"
+#include "core/factorization/cholesky_kernels.hpp"
 #include "core/factorization/elimination_forest.hpp"
 #include "core/test/utils.hpp"
 #include "core/test/utils/assertions.hpp"