From c40d147e046879e1b960c2769e274269f5266e37 Mon Sep 17 00:00:00 2001 From: Tobias Ribizel Date: Wed, 9 Feb 2022 18:52:10 +0100 Subject: [PATCH] review updates Co-authored-by: Fritz Goebel --- core/components/disjoint_sets.hpp | 2 +- core/factorization/elimination_forest.cpp | 2 +- core/test/components/disjoint_sets.cpp | 8 ++--- .../test/factorization/elimination_forest.cpp | 34 +++++++++++++++++++ cuda/factorization/cholesky_kernels.cu | 2 +- dev_tools/scripts/config | 4 +++ dpcpp/factorization/cholesky_kernels.dp.cpp | 2 +- hip/factorization/cholesky_kernels.hip.cpp | 2 +- .../test/factorization/cholesky_kernels.cpp | 4 +-- 9 files changed, 46 insertions(+), 14 deletions(-) diff --git a/core/components/disjoint_sets.hpp b/core/components/disjoint_sets.hpp index 6dcaf673b6e..0640c2f2e10 100644 --- a/core/components/disjoint_sets.hpp +++ b/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 diff --git a/core/factorization/elimination_forest.cpp b/core/factorization/elimination_forest.cpp index 71752b112c7..c26a0eaf498 100644 --- a/core/factorization/elimination_forest.cpp +++ b/core/factorization/elimination_forest.cpp @@ -60,7 +60,7 @@ void compute_elim_forest_parent_impl(std::shared_ptr 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); diff --git a/core/test/components/disjoint_sets.cpp b/core/test/components/disjoint_sets.cpp index 4d2c59463a0..0de97e6b290 100644 --- a/core/test/components/disjoint_sets.cpp +++ b/core/test/components/disjoint_sets.cpp @@ -34,6 +34,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include #include @@ -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(i)).count(); } diff --git a/core/test/factorization/elimination_forest.cpp b/core/test/factorization/elimination_forest.cpp index 5ac5c7413c1..3e4a9fffa97 100644 --- a/core/test/factorization/elimination_forest.cpp +++ b/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( + { + {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({2, 4, 6, 8, 5, 6, 7, 8, 9, 10})); + GKO_ASSERT_ARRAY_EQ(forest.child_ptrs, + I({0, 0, 0, 1, 1, 2, 3, 5, 6, 8, 9, 10})); + GKO_ASSERT_ARRAY_EQ(forest.children, + I({0, 1, 4, 2, 5, 6, 3, 7, 8, 9})); + GKO_ASSERT_ARRAY_EQ(forest.postorder, + I({3, 0, 2, 1, 4, 5, 6, 7, 8, 9})); + GKO_ASSERT_ARRAY_EQ(forest.inv_postorder, + I({1, 3, 2, 0, 4, 5, 6, 7, 8, 9})); + GKO_ASSERT_ARRAY_EQ(forest.postorder_parents, + I({8, 2, 6, 4, 5, 6, 7, 8, 9, 10})); +} + + TYPED_TEST(EliminationForest, WorksForAni1) { using matrix_type = typename TestFixture::matrix_type; diff --git a/cuda/factorization/cholesky_kernels.cu b/cuda/factorization/cholesky_kernels.cu index 99abdce4bf6..2097fa68335 100644 --- a/cuda/factorization/cholesky_kernels.cu +++ b/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); diff --git a/dev_tools/scripts/config b/dev_tools/scripts/config index 36ff9ca13d8..7f43fb78089 100644 --- a/dev_tools/scripts/config +++ b/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" diff --git a/dpcpp/factorization/cholesky_kernels.dp.cpp b/dpcpp/factorization/cholesky_kernels.dp.cpp index 53acbd9727e..61660b98422 100644 --- a/dpcpp/factorization/cholesky_kernels.dp.cpp +++ b/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]; diff --git a/hip/factorization/cholesky_kernels.hip.cpp b/hip/factorization/cholesky_kernels.hip.cpp index f069dd64b90..182b05f0032 100644 --- a/hip/factorization/cholesky_kernels.hip.cpp +++ b/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); diff --git a/reference/test/factorization/cholesky_kernels.cpp b/reference/test/factorization/cholesky_kernels.cpp index 98b4c5a1959..59efb0ea4b2 100644 --- a/reference/test/factorization/cholesky_kernels.cpp +++ b/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. *************************************************************/ -#include "core/factorization/cholesky_kernels.hpp" - - #include #include @@ -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"