From 75551c4afa0997105504e05e1b186aeb8f14fe9c Mon Sep 17 00:00:00 2001 From: Pratik Nayak Date: Thu, 12 Sep 2019 22:57:40 +0200 Subject: [PATCH 1/8] Add reference kernels and tests. + Add conversions from dense to Sparsity pattern. --- core/device_hooks/common_kernels.inc.cpp | 6 + core/matrix/dense.cpp | 52 ++++ core/matrix/dense_kernels.hpp | 75 ++--- core/matrix/sparsity.cpp | 2 + cuda/matrix/dense_kernels.cu | 11 + include/ginkgo/core/matrix/dense.hpp | 16 + include/ginkgo/core/matrix/sparsity.hpp | 1 + omp/matrix/dense_kernels.cpp | 35 +++ reference/matrix/dense_kernels.cpp | 31 ++ reference/matrix/sparsity_kernels.cpp | 128 +++++++- reference/test/matrix/CMakeLists.txt | 1 + reference/test/matrix/sparsity_kernels.cpp | 335 +++++++++++++++++++++ 12 files changed, 647 insertions(+), 46 deletions(-) create mode 100644 reference/test/matrix/sparsity_kernels.cpp diff --git a/core/device_hooks/common_kernels.inc.cpp b/core/device_hooks/common_kernels.inc.cpp index ebe128f6dca..e338c731ed4 100644 --- a/core/device_hooks/common_kernels.inc.cpp +++ b/core/device_hooks/common_kernels.inc.cpp @@ -125,6 +125,12 @@ GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL); +template +GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL(ValueType, IndexType) +GKO_NOT_COMPILED(GKO_HOOK_MODULE); +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL); + template GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL(ValueType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); diff --git a/core/matrix/dense.cpp b/core/matrix/dense.cpp index 703b89c08fa..c24afc1ba12 100644 --- a/core/matrix/dense.cpp +++ b/core/matrix/dense.cpp @@ -43,6 +43,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include "core/matrix/dense_kernels.hpp" @@ -75,6 +76,7 @@ GKO_REGISTER_OPERATION(convert_to_csr, dense::convert_to_csr); GKO_REGISTER_OPERATION(convert_to_ell, dense::convert_to_ell); GKO_REGISTER_OPERATION(convert_to_hybrid, dense::convert_to_hybrid); GKO_REGISTER_OPERATION(convert_to_sellp, dense::convert_to_sellp); +GKO_REGISTER_OPERATION(convert_to_sparsity, dense::convert_to_sparsity); } // namespace dense @@ -181,6 +183,22 @@ inline void conversion_helper(Sellp *result, } +template +inline void conversion_helper(Sparsity *result, + MatrixType *source, const OperationType &op) +{ + auto exec = source->get_executor(); + + size_type num_stored_nonzeros = 0; + exec->run(dense::make_count_nonzeros(source, &num_stored_nonzeros)); + auto tmp = Sparsity::create(exec, source->get_size(), + num_stored_nonzeros); + exec->run(op(tmp.get(), source)); + tmp->move_to(result); +} + + } // namespace @@ -423,6 +441,40 @@ void Dense::move_to(Sellp *result) } +template +void Dense::convert_to(Sparsity *result) const +{ + conversion_helper( + result, this, + dense::template make_convert_to_sparsity *&>); +} + + +template +void Dense::move_to(Sparsity *result) +{ + this->convert_to(result); +} + + +template +void Dense::convert_to(Sparsity *result) const +{ + conversion_helper( + result, this, + dense::template make_convert_to_sparsity *&>); +} + + +template +void Dense::move_to(Sparsity *result) +{ + this->convert_to(result); +} + + namespace { diff --git a/core/matrix/dense_kernels.hpp b/core/matrix/dense_kernels.hpp index f5555a8c32d..66f0ca41111 100644 --- a/core/matrix/dense_kernels.hpp +++ b/core/matrix/dense_kernels.hpp @@ -98,6 +98,11 @@ namespace kernels { matrix::Sellp<_type, _prec> *other, \ const matrix::Dense<_type> *source) +#define GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL(_type, _prec) \ + void convert_to_sparsity(std::shared_ptr exec, \ + matrix::Sparsity<_type, _prec> *other, \ + const matrix::Dense<_type> *source) + #define GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL(_type) \ void count_nonzeros(std::shared_ptr exec, \ const matrix::Dense<_type> *source, size_type *result) @@ -128,40 +133,42 @@ namespace kernels { matrix::Dense<_type> *trans, \ const matrix::Dense<_type> *orig) -#define GKO_DECLARE_ALL_AS_TEMPLATES \ - template \ - GKO_DECLARE_DENSE_SIMPLE_APPLY_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_APPLY_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_SCALE_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_ADD_SCALED_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_COMPUTE_DOT_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_COMPUTE_NORM2_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_CONVERT_TO_COO_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_DENSE_CONVERT_TO_CSR_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_DENSE_CONVERT_TO_ELL_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_DENSE_CONVERT_TO_HYBRID_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_CALCULATE_MAX_NNZ_PER_ROW_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_CALCULATE_NONZEROS_PER_ROW_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_CALCULATE_TOTAL_COLS_KERNEL(ValueType); \ - template \ - GKO_DECLARE_TRANSPOSE_KERNEL(ValueType); \ - template \ +#define GKO_DECLARE_ALL_AS_TEMPLATES \ + template \ + GKO_DECLARE_DENSE_SIMPLE_APPLY_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_APPLY_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_SCALE_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_ADD_SCALED_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_COMPUTE_DOT_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_COMPUTE_NORM2_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_CONVERT_TO_COO_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_DENSE_CONVERT_TO_CSR_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_DENSE_CONVERT_TO_ELL_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_DENSE_CONVERT_TO_HYBRID_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_CALCULATE_MAX_NNZ_PER_ROW_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_CALCULATE_NONZEROS_PER_ROW_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_CALCULATE_TOTAL_COLS_KERNEL(ValueType); \ + template \ + GKO_DECLARE_TRANSPOSE_KERNEL(ValueType); \ + template \ GKO_DECLARE_CONJ_TRANSPOSE_KERNEL(ValueType) diff --git a/core/matrix/sparsity.cpp b/core/matrix/sparsity.cpp index 5228b973df0..d98d6172d2d 100644 --- a/core/matrix/sparsity.cpp +++ b/core/matrix/sparsity.cpp @@ -164,6 +164,8 @@ template std::unique_ptr> Sparsity::to_adjacency_matrix() const { + // Adjacency matrix has to be square. + GKO_ASSERT_IS_SQUARE_MATRIX(this); auto exec = this->get_executor(); GKO_ASSERT_IS_SQUARE_MATRIX(this); size_type num_diagonal_elements = 0; diff --git a/cuda/matrix/dense_kernels.cu b/cuda/matrix/dense_kernels.cu index 4b410a4ae63..c66cf6ae427 100644 --- a/cuda/matrix/dense_kernels.cu +++ b/cuda/matrix/dense_kernels.cu @@ -39,6 +39,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include "cuda/base/cublas_bindings.hpp" @@ -730,6 +731,16 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL); +template +void convert_to_sparsity(std::shared_ptr exec, + matrix::Sparsity *result, + const matrix::Dense *source) + GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL); + + template void count_nonzeros(std::shared_ptr exec, const matrix::Dense *source, size_type *result) diff --git a/include/ginkgo/core/matrix/dense.hpp b/include/ginkgo/core/matrix/dense.hpp index 3e23ebaa23d..ea568a02f3a 100644 --- a/include/ginkgo/core/matrix/dense.hpp +++ b/include/ginkgo/core/matrix/dense.hpp @@ -65,6 +65,9 @@ class Hybrid; template class Sellp; +template +class Sparsity; + /** * Dense is a matrix format which explicitly stores all values of the matrix. @@ -94,6 +97,8 @@ class Dense : public EnableLinOp>, public ConvertibleTo>, public ConvertibleTo>, public ConvertibleTo>, + public ConvertibleTo>, + public ConvertibleTo>, public ReadableFromMatrixData, public ReadableFromMatrixData, public WritableToMatrixData, @@ -111,6 +116,8 @@ class Dense : public EnableLinOp>, friend class Hybrid; friend class Sellp; friend class Sellp; + friend class Sparsity; + friend class Sparsity; public: using EnableLinOp::convert_to; @@ -179,6 +186,15 @@ class Dense : public EnableLinOp>, void move_to(Sellp *result) override; + void convert_to(Sparsity *result) const override; + + void move_to(Sparsity *result) override; + + void convert_to(Sparsity *result) const override; + + void move_to(Sparsity *result) override; + + void read(const mat_data &data) override; void read(const mat_data32 &data) override; diff --git a/include/ginkgo/core/matrix/sparsity.hpp b/include/ginkgo/core/matrix/sparsity.hpp index e8e6a6e4c75..76a3e8b8ee7 100644 --- a/include/ginkgo/core/matrix/sparsity.hpp +++ b/include/ginkgo/core/matrix/sparsity.hpp @@ -73,6 +73,7 @@ class Sparsity : public EnableLinOp>, using EnableLinOp::convert_to; using EnableLinOp::move_to; + using value_type = ValueType; using index_type = IndexType; using value_type = ValueType; diff --git a/omp/matrix/dense_kernels.cpp b/omp/matrix/dense_kernels.cpp index 12f0b8cf17d..07d33cb92b4 100644 --- a/omp/matrix/dense_kernels.cpp +++ b/omp/matrix/dense_kernels.cpp @@ -46,6 +46,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include namespace gko { @@ -455,6 +456,40 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL); +template +void convert_to_sparsity(std::shared_ptr exec, + matrix::Sparsity *result, + const matrix::Dense *source) +{ + auto num_rows = result->get_size()[0]; + auto num_cols = result->get_size()[1]; + auto num_nonzeros = result->get_num_nonzeros(); + + auto row_ptrs = result->get_row_ptrs(); + auto col_idxs = result->get_col_idxs(); + + size_type cur_ptr = 0; + row_ptrs[0] = cur_ptr; + for (size_type row = 0; row < num_rows; ++row) { +#pragma omp parallel for + for (size_type col = 0; col < num_cols; ++col) { + auto val = source->at(row, col); + if (val != zero()) { +#pragma omp critical + { + col_idxs[cur_ptr] = col; + ++cur_ptr; + } + } + } + row_ptrs[row + 1] = cur_ptr; + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL); + + template void count_nonzeros(std::shared_ptr exec, const matrix::Dense *source, size_type *result) diff --git a/reference/matrix/dense_kernels.cpp b/reference/matrix/dense_kernels.cpp index ba975b720c9..ab3d76318b7 100644 --- a/reference/matrix/dense_kernels.cpp +++ b/reference/matrix/dense_kernels.cpp @@ -40,6 +40,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include @@ -414,6 +415,36 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL); +template +void convert_to_sparsity(std::shared_ptr exec, + matrix::Sparsity *result, + const matrix::Dense *source) +{ + auto num_rows = result->get_size()[0]; + auto num_cols = result->get_size()[1]; + auto num_nonzeros = result->get_num_nonzeros(); + + auto row_ptrs = result->get_row_ptrs(); + auto col_idxs = result->get_col_idxs(); + + size_type cur_ptr = 0; + row_ptrs[0] = cur_ptr; + for (size_type row = 0; row < num_rows; ++row) { + for (size_type col = 0; col < num_cols; ++col) { + auto val = source->at(row, col); + if (val != zero()) { + col_idxs[cur_ptr] = col; + ++cur_ptr; + } + } + row_ptrs[row + 1] = cur_ptr; + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL); + + template void count_nonzeros(std::shared_ptr exec, const matrix::Dense *source, size_type *result) diff --git a/reference/matrix/sparsity_kernels.cpp b/reference/matrix/sparsity_kernels.cpp index 87a51d211ca..72dd4282e69 100644 --- a/reference/matrix/sparsity_kernels.cpp +++ b/reference/matrix/sparsity_kernels.cpp @@ -51,7 +51,7 @@ namespace gko { namespace kernels { namespace reference { /** - * @brief The Compressed sparse row matrix format namespace. + * @brief The Sparsity pattern format namespace. * @ref Sparsity * @ingroup sparsity */ @@ -61,8 +61,25 @@ namespace sparsity { template void spmv(std::shared_ptr exec, const matrix::Sparsity *a, - const matrix::Dense *b, - matrix::Dense *c) GKO_NOT_IMPLEMENTED; + const matrix::Dense *b, matrix::Dense *c) +{ + auto row_ptrs = a->get_const_row_ptrs(); + auto col_idxs = a->get_const_col_idxs(); + + for (size_type row = 0; row < a->get_size()[0]; ++row) { + for (size_type j = 0; j < c->get_size()[1]; ++j) { + c->at(row, j) = zero(); + } + for (size_type k = row_ptrs[row]; + k < static_cast(row_ptrs[row + 1]); ++k) { + auto val = one(); + auto col = col_idxs[k]; + for (size_type j = 0; j < c->get_size()[1]; ++j) { + c->at(row, j) += val * b->at(col, j); + } + } + } +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SPARSITY_SPMV_KERNEL); @@ -73,7 +90,27 @@ void advanced_spmv(std::shared_ptr exec, const matrix::Sparsity *a, const matrix::Dense *b, const matrix::Dense *beta, - matrix::Dense *c) GKO_NOT_IMPLEMENTED; + matrix::Dense *c) +{ + auto row_ptrs = a->get_const_row_ptrs(); + auto col_idxs = a->get_const_col_idxs(); + auto valpha = alpha->at(0, 0); + auto vbeta = beta->at(0, 0); + + for (size_type row = 0; row < a->get_size()[0]; ++row) { + for (size_type j = 0; j < c->get_size()[1]; ++j) { + c->at(row, j) *= vbeta; + } + for (size_type k = row_ptrs[row]; + k < static_cast(row_ptrs[row + 1]); ++k) { + auto val = one(); + auto col = col_idxs[k]; + for (size_type j = 0; j < c->get_size()[1]; ++j) { + c->at(row, j) += valpha * val * b->at(col, j); + } + } + } +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_ADVANCED_SPMV_KERNEL); @@ -93,7 +130,26 @@ template void remove_diagonal_elements(std::shared_ptr exec, matrix::Sparsity *matrix, const IndexType *row_ptrs, - const IndexType *col_idxs) GKO_NOT_IMPLEMENTED; + const IndexType *col_idxs) +{ + auto num_rows = matrix->get_size()[0]; + auto adj_ptrs = matrix->get_row_ptrs(); + auto adj_idxs = matrix->get_col_idxs(); + for (auto i = 0; i <= num_rows; ++i) { + adj_ptrs[i] = row_ptrs[i] - i; + } + std::vector temp_idxs; + for (auto i = 0; i < num_rows; ++i) { + for (auto j = row_ptrs[i]; j < row_ptrs[i + 1]; ++j) { + if (col_idxs[j] != i) { + temp_idxs.push_back(col_idxs[j]); + } + } + } + for (auto i = 0; i < temp_idxs.size(); ++i) { + adj_idxs[i] = temp_idxs[i]; + } +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_REMOVE_DIAGONAL_ELEMENTS_KERNEL); @@ -103,22 +159,47 @@ template inline void convert_sparsity_to_csc(size_type num_rows, const IndexType *row_ptrs, const IndexType *col_idxs, - IndexType *row_idxs, - IndexType *col_ptrs) GKO_NOT_IMPLEMENTED; + IndexType *row_idxs, IndexType *col_ptrs) +{ + for (size_type row = 0; row < num_rows; ++row) { + for (auto i = row_ptrs[row]; i < row_ptrs[row + 1]; ++i) { + const auto dest_idx = col_ptrs[col_idxs[i]]++; + row_idxs[dest_idx] = row; + } + } +} template void transpose_and_transform(std::shared_ptr exec, matrix::Sparsity *trans, const matrix::Sparsity *orig) - GKO_NOT_IMPLEMENTED; +{ + auto trans_row_ptrs = trans->get_row_ptrs(); + auto orig_row_ptrs = orig->get_const_row_ptrs(); + auto trans_col_idxs = trans->get_col_idxs(); + auto orig_col_idxs = orig->get_const_col_idxs(); + + auto orig_num_cols = orig->get_size()[1]; + auto orig_num_rows = orig->get_size()[0]; + auto orig_nnz = orig_row_ptrs[orig_num_rows]; + + trans_row_ptrs[0] = 0; + convert_idxs_to_ptrs(orig_col_idxs, orig_nnz, trans_row_ptrs + 1, + orig_num_cols); + + convert_sparsity_to_csc(orig_num_rows, orig_row_ptrs, orig_col_idxs, + trans_col_idxs, trans_row_ptrs + 1); +} template void transpose(std::shared_ptr exec, matrix::Sparsity *trans, const matrix::Sparsity *orig) - GKO_NOT_IMPLEMENTED; +{ + transpose_and_transform(exec, trans, orig); +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_TRANSPOSE_KERNEL); @@ -127,7 +208,16 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( template void sort_by_column_index(std::shared_ptr exec, matrix::Sparsity *to_sort) - GKO_NOT_IMPLEMENTED; +{ + auto row_ptrs = to_sort->get_row_ptrs(); + auto col_idxs = to_sort->get_col_idxs(); + const auto number_rows = to_sort->get_size()[0]; + for (size_type i = 0; i < number_rows; ++i) { + auto start_row_idx = row_ptrs[i]; + auto row_nnz = row_ptrs[i + 1] - start_row_idx; + std::sort(col_idxs + start_row_idx, col_idxs + start_row_idx + row_nnz); + } +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_SORT_BY_COLUMN_INDEX); @@ -136,8 +226,22 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( template void is_sorted_by_column_index( std::shared_ptr exec, - const matrix::Sparsity *to_check, - bool *is_sorted) GKO_NOT_IMPLEMENTED; + const matrix::Sparsity *to_check, bool *is_sorted) +{ + const auto row_ptrs = to_check->get_const_row_ptrs(); + const auto col_idxs = to_check->get_const_col_idxs(); + const auto size = to_check->get_size(); + for (size_type i = 0; i < size[0]; ++i) { + for (auto idx = row_ptrs[i] + 1; idx < row_ptrs[i + 1]; ++idx) { + if (col_idxs[idx - 1] > col_idxs[idx]) { + *is_sorted = false; + return; + } + } + } + *is_sorted = true; + return; +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_IS_SORTED_BY_COLUMN_INDEX); diff --git a/reference/test/matrix/CMakeLists.txt b/reference/test/matrix/CMakeLists.txt index 25b7597c3f3..e105acbc3d6 100644 --- a/reference/test/matrix/CMakeLists.txt +++ b/reference/test/matrix/CMakeLists.txt @@ -5,3 +5,4 @@ ginkgo_create_test(ell_kernels) ginkgo_create_test(hybrid_kernels) ginkgo_create_test(identity) ginkgo_create_test(sellp_kernels) +ginkgo_create_test(sparsity_kernels) diff --git a/reference/test/matrix/sparsity_kernels.cpp b/reference/test/matrix/sparsity_kernels.cpp new file mode 100644 index 00000000000..81bb392950c --- /dev/null +++ b/reference/test/matrix/sparsity_kernels.cpp @@ -0,0 +1,335 @@ +/************************************************************* +Copyright (c) 2017-2019, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/matrix/sparsity_kernels.hpp" + + +#include + + +#include + + +#include +#include +#include +#include +#include + + +#include "core/test/utils/assertions.hpp" + + +namespace { + + +class Sparsity : public ::testing::Test { +protected: + using Mtx = gko::matrix::Sparsity<>; + using Vec = gko::matrix::Dense<>; + + Sparsity() + : exec(gko::ReferenceExecutor::create()), + mtx(Mtx::create(exec, gko::dim<2>{2, 3}, 4)), + mtx2(Mtx::create(exec, gko::dim<2>{2, 3}, 5)), + mtx3_sorted(Mtx::create(exec, gko::dim<2>(3, 3), 7)), + mtx3_unsorted(Mtx::create(exec, gko::dim<2>(3, 3), 7)) + { + this->create_mtx(mtx.get()); + this->create_mtx2(mtx2.get()); + this->create_mtx3(mtx3_sorted.get(), mtx3_unsorted.get()); + } + + void create_mtx(Mtx *m) + { + Mtx::index_type *c = m->get_col_idxs(); + Mtx::index_type *r = m->get_row_ptrs(); + /* + * 1 1 1 + * 0 1 0 + */ + r[0] = 0; + r[1] = 3; + r[2] = 4; + c[0] = 0; + c[1] = 1; + c[2] = 2; + c[3] = 1; + } + + void create_mtx2(Mtx *m) + { + Mtx::index_type *c = m->get_col_idxs(); + Mtx::index_type *r = m->get_row_ptrs(); + // It keeps an explict zero + /* + * 1 1 1 + * {0} 1 0 + */ + r[0] = 0; + r[1] = 3; + r[2] = 5; + c[0] = 0; + c[1] = 1; + c[2] = 2; + c[3] = 0; + c[4] = 1; + } + + void create_mtx3(Mtx *sorted, Mtx *unsorted) + { + auto cols_s = sorted->get_col_idxs(); + auto rows_s = sorted->get_row_ptrs(); + auto cols_u = unsorted->get_col_idxs(); + auto rows_u = unsorted->get_row_ptrs(); + /* For both versions (sorted and unsorted), this matrix is stored: + * 0 1 1 + * 1 1 1 + * 1 0 1 + * The unsorted matrix will have the (value, column) pair per row not + * sorted, which we still consider a valid SPARSITY format. + */ + rows_s[0] = 0; + rows_s[1] = 2; + rows_s[2] = 5; + rows_s[3] = 7; + rows_u[0] = 0; + rows_u[1] = 2; + rows_u[2] = 5; + rows_u[3] = 7; + + cols_s[0] = 1; + cols_s[1] = 2; + cols_s[2] = 0; + cols_s[3] = 1; + cols_s[4] = 2; + cols_s[5] = 0; + cols_s[6] = 2; + // The same applies for the columns + cols_u[0] = 2; + cols_u[1] = 1; + cols_u[2] = 1; + cols_u[3] = 2; + cols_u[4] = 0; + cols_u[5] = 2; + cols_u[6] = 0; + } + + std::shared_ptr exec; + std::unique_ptr mtx; + std::unique_ptr mtx2; + std::unique_ptr mtx3_sorted; + std::unique_ptr mtx3_unsorted; +}; + + +TEST_F(Sparsity, AppliesToDenseVector) +{ + auto x = gko::initialize({2.0, 1.0, 4.0}, exec); + auto y = Vec::create(exec, gko::dim<2>{2, 1}); + + mtx->apply(x.get(), y.get()); + + EXPECT_EQ(y->at(0), 7.0); + EXPECT_EQ(y->at(1), 1.0); +} + + +TEST_F(Sparsity, AppliesToDenseMatrix) +{ + auto x = gko::initialize({{2.0, 3.0}, {1.0, -1.5}, {4.0, 2.5}}, exec); + auto y = Vec::create(exec, gko::dim<2>{2}); + + mtx->apply(x.get(), y.get()); + + EXPECT_EQ(y->at(0, 0), 7.0); + EXPECT_EQ(y->at(1, 0), 1.0); + EXPECT_EQ(y->at(0, 1), 4.0); + EXPECT_EQ(y->at(1, 1), -1.5); +} + + +TEST_F(Sparsity, AppliesLinearCombinationToDenseVector) +{ + auto alpha = gko::initialize({-1.0}, exec); + auto beta = gko::initialize({2.0}, exec); + auto x = gko::initialize({2.0, 1.0, 4.0}, exec); + auto y = gko::initialize({1.0, 2.0}, exec); + + mtx->apply(alpha.get(), x.get(), beta.get(), y.get()); + + EXPECT_EQ(y->at(0), -5.0); + EXPECT_EQ(y->at(1), 3.0); +} + + +TEST_F(Sparsity, AppliesLinearCombinationToDenseMatrix) +{ + auto alpha = gko::initialize({-1.0}, exec); + auto beta = gko::initialize({2.0}, exec); + auto x = gko::initialize({{2.0, 3.0}, {1.0, -1.5}, {4.0, 2.5}}, exec); + auto y = gko::initialize({{1.0, 0.5}, {2.0, -1.5}}, exec); + + mtx->apply(alpha.get(), x.get(), beta.get(), y.get()); + + EXPECT_EQ(y->at(0, 0), -5.0); + EXPECT_EQ(y->at(1, 0), 3.0); + EXPECT_EQ(y->at(0, 1), -3.0); + EXPECT_EQ(y->at(1, 1), -1.5); +} + + +TEST_F(Sparsity, ApplyFailsOnWrongInnerDimension) +{ + auto x = Vec::create(exec, gko::dim<2>{2}); + auto y = Vec::create(exec, gko::dim<2>{2}); + + ASSERT_THROW(mtx->apply(x.get(), y.get()), gko::DimensionMismatch); +} + + +TEST_F(Sparsity, ApplyFailsOnWrongNumberOfRows) +{ + auto x = Vec::create(exec, gko::dim<2>{3, 2}); + auto y = Vec::create(exec, gko::dim<2>{3, 2}); + + ASSERT_THROW(mtx->apply(x.get(), y.get()), gko::DimensionMismatch); +} + + +TEST_F(Sparsity, ApplyFailsOnWrongNumberOfCols) +{ + auto x = Vec::create(exec, gko::dim<2>{3}); + auto y = Vec::create(exec, gko::dim<2>{2}); + + ASSERT_THROW(mtx->apply(x.get(), y.get()), gko::DimensionMismatch); +} + + +TEST_F(Sparsity, SquareMtxIsTransposable) +{ + // clang-format off + auto mtx2 = gko::initialize>( + {{1.0, 1.0, 1.0}, + {0.0, 1.0, 0.0}, + {0.0, 1.0, 1.0}}, exec); + // clang-format on + + auto trans = mtx2->transpose(); + auto trans_as_sparsity = + static_cast *>(trans.get()); + + // clang-format off + GKO_ASSERT_MTX_NEAR(trans_as_sparsity, + l({{1.0, 0.0, 0.0}, + {1.0, 1.0, 1.0}, + {1.0, 0.0, 1.0}}), 0.0); + // clang-format on +} + + +TEST_F(Sparsity, NonSquareMtxIsTransposable) +{ + auto trans = mtx->transpose(); + auto trans_as_sparsity = + static_cast *>(trans.get()); + + // clang-format off + GKO_ASSERT_MTX_NEAR(trans_as_sparsity, + l({{1.0, 0.0}, + {1.0, 1.0}, + {1.0, 0.0}}), 0.0); + // clang-format on +} + + +TEST_F(Sparsity, SquareMtxIsConvertibleToAdjacencyMatrix) +{ + // clang-format off + auto mtx2 = gko::initialize>( + {{1.0, 1.0, 1.0}, + {0.0, 1.0, 0.0}, + {0.0, 1.0, 1.0}}, exec); + // clang-format on + + auto adj_mat = mtx2->to_adjacency_matrix(); + + // clang-format off + GKO_ASSERT_MTX_NEAR(adj_mat.get(), + l({{0.0, 1.0, 1.0}, + {0.0, 0.0, 0.0}, + {0.0, 1.0, 0.0}}), 0.0); + // clang-format on +} + + +TEST_F(Sparsity, NonSquareMtxIsNotConvertibleToAdjacencyMatrix) +{ + ASSERT_THROW(mtx->to_adjacency_matrix(), gko::DimensionMismatch); +} + + +TEST_F(Sparsity, RecognizeSortedMatrix) +{ + ASSERT_TRUE(mtx->is_sorted_by_column_index()); + ASSERT_TRUE(mtx2->is_sorted_by_column_index()); + ASSERT_TRUE(mtx3_sorted->is_sorted_by_column_index()); +} + + +TEST_F(Sparsity, RecognizeUnsortedMatrix) +{ + ASSERT_FALSE(mtx3_unsorted->is_sorted_by_column_index()); +} + + +TEST_F(Sparsity, SortSortedMatrix) +{ + auto matrix = mtx3_sorted->clone(); + + matrix->sort_by_column_index(); + + GKO_ASSERT_MTX_NEAR(matrix, mtx3_sorted, 0.0); +} + + +TEST_F(Sparsity, SortUnsortedMatrix) +{ + auto matrix = mtx3_unsorted->clone(); + + matrix->sort_by_column_index(); + + GKO_ASSERT_MTX_NEAR(matrix, mtx3_sorted, 0.0); +} + + +} // namespace From 9275cd7c22067501f09631b9145aaac19c681331 Mon Sep 17 00:00:00 2001 From: Pratik Nayak Date: Fri, 13 Sep 2019 00:24:28 +0200 Subject: [PATCH 2/8] Add OpenMP kernels and tests. --- omp/matrix/sparsity_kernels.cpp | 136 ++++++++++++-- omp/test/matrix/CMakeLists.txt | 1 + omp/test/matrix/sparsity_kernels.cpp | 270 +++++++++++++++++++++++++++ 3 files changed, 395 insertions(+), 12 deletions(-) create mode 100644 omp/test/matrix/sparsity_kernels.cpp diff --git a/omp/matrix/sparsity_kernels.cpp b/omp/matrix/sparsity_kernels.cpp index 4a3118ebf8f..4c8e8944e1d 100644 --- a/omp/matrix/sparsity_kernels.cpp +++ b/omp/matrix/sparsity_kernels.cpp @@ -54,7 +54,7 @@ namespace gko { namespace kernels { namespace omp { /** - * @brief The Compressed sparse row matrix format namespace. + * @brief The Sparsity pattern format namespace. * * @ingroup sparsity */ @@ -64,8 +64,26 @@ namespace sparsity { template void spmv(std::shared_ptr exec, const matrix::Sparsity *a, - const matrix::Dense *b, - matrix::Dense *c) GKO_NOT_IMPLEMENTED; + const matrix::Dense *b, matrix::Dense *c) +{ + auto row_ptrs = a->get_const_row_ptrs(); + auto col_idxs = a->get_const_col_idxs(); + +#pragma omp parallel for + for (size_type row = 0; row < a->get_size()[0]; ++row) { + for (size_type j = 0; j < c->get_size()[1]; ++j) { + c->at(row, j) = zero(); + } + for (size_type k = row_ptrs[row]; + k < static_cast(row_ptrs[row + 1]); ++k) { + auto val = one(); + auto col = col_idxs[k]; + for (size_type j = 0; j < c->get_size()[1]; ++j) { + c->at(row, j) += val * b->at(col, j); + } + } + } +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SPARSITY_SPMV_KERNEL); @@ -76,7 +94,28 @@ void advanced_spmv(std::shared_ptr exec, const matrix::Sparsity *a, const matrix::Dense *b, const matrix::Dense *beta, - matrix::Dense *c) GKO_NOT_IMPLEMENTED; + matrix::Dense *c) +{ + auto row_ptrs = a->get_const_row_ptrs(); + auto col_idxs = a->get_const_col_idxs(); + auto valpha = alpha->at(0, 0); + auto vbeta = beta->at(0, 0); + +#pragma omp parallel for + for (size_type row = 0; row < a->get_size()[0]; ++row) { + for (size_type j = 0; j < c->get_size()[1]; ++j) { + c->at(row, j) *= vbeta; + } + for (size_type k = row_ptrs[row]; + k < static_cast(row_ptrs[row + 1]); ++k) { + auto val = one(); + auto col = col_idxs[k]; + for (size_type j = 0; j < c->get_size()[1]; ++j) { + c->at(row, j) += valpha * val * b->at(col, j); + } + } + } +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_ADVANCED_SPMV_KERNEL); @@ -96,7 +135,26 @@ template void remove_diagonal_elements(std::shared_ptr exec, matrix::Sparsity *matrix, const IndexType *row_ptrs, - const IndexType *col_idxs) GKO_NOT_IMPLEMENTED; + const IndexType *col_idxs) +{ + auto num_rows = matrix->get_size()[0]; + auto adj_ptrs = matrix->get_row_ptrs(); + auto adj_idxs = matrix->get_col_idxs(); + for (auto i = 0; i <= num_rows; ++i) { + adj_ptrs[i] = row_ptrs[i] - i; + } + std::vector temp_idxs; + for (auto i = 0; i < num_rows; ++i) { + for (auto j = row_ptrs[i]; j < row_ptrs[i + 1]; ++j) { + if (col_idxs[j] != i) { + temp_idxs.push_back(col_idxs[j]); + } + } + } + for (auto i = 0; i < temp_idxs.size(); ++i) { + adj_idxs[i] = temp_idxs[i]; + } +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_REMOVE_DIAGONAL_ELEMENTS_KERNEL); @@ -106,22 +164,47 @@ template inline void convert_sparsity_to_csc(size_type num_rows, const IndexType *row_ptrs, const IndexType *col_idxs, - IndexType *row_idxs, - IndexType *col_ptrs) GKO_NOT_IMPLEMENTED; + IndexType *row_idxs, IndexType *col_ptrs) +{ + for (size_type row = 0; row < num_rows; ++row) { + for (auto i = row_ptrs[row]; i < row_ptrs[row + 1]; ++i) { + const auto dest_idx = col_ptrs[col_idxs[i]]++; + row_idxs[dest_idx] = row; + } + } +} template void transpose_and_transform(std::shared_ptr exec, matrix::Sparsity *trans, const matrix::Sparsity *orig) - GKO_NOT_IMPLEMENTED; +{ + auto trans_row_ptrs = trans->get_row_ptrs(); + auto orig_row_ptrs = orig->get_const_row_ptrs(); + auto trans_col_idxs = trans->get_col_idxs(); + auto orig_col_idxs = orig->get_const_col_idxs(); + + auto orig_num_cols = orig->get_size()[1]; + auto orig_num_rows = orig->get_size()[0]; + auto orig_nnz = orig_row_ptrs[orig_num_rows]; + + trans_row_ptrs[0] = 0; + convert_unsorted_idxs_to_ptrs(orig_col_idxs, orig_nnz, trans_row_ptrs + 1, + orig_num_cols); + + convert_sparsity_to_csc(orig_num_rows, orig_row_ptrs, orig_col_idxs, + trans_col_idxs, trans_row_ptrs + 1); +} template void transpose(std::shared_ptr exec, matrix::Sparsity *trans, const matrix::Sparsity *orig) - GKO_NOT_IMPLEMENTED; +{ + transpose_and_transform(exec, trans, orig); +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_TRANSPOSE_KERNEL); @@ -130,7 +213,17 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( template void sort_by_column_index(std::shared_ptr exec, matrix::Sparsity *to_sort) - GKO_NOT_IMPLEMENTED; +{ + auto row_ptrs = to_sort->get_row_ptrs(); + auto col_idxs = to_sort->get_col_idxs(); + const auto number_rows = to_sort->get_size()[0]; +#pragma omp parallel for + for (size_type i = 0; i < number_rows; ++i) { + auto start_row_idx = row_ptrs[i]; + auto row_nnz = row_ptrs[i + 1] - start_row_idx; + std::sort(col_idxs + start_row_idx, col_idxs + start_row_idx + row_nnz); + } +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_SORT_BY_COLUMN_INDEX); @@ -139,8 +232,27 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( template void is_sorted_by_column_index( std::shared_ptr exec, - const matrix::Sparsity *to_check, - bool *is_sorted) GKO_NOT_IMPLEMENTED; + const matrix::Sparsity *to_check, bool *is_sorted) +{ + const auto row_ptrs = to_check->get_const_row_ptrs(); + const auto col_idxs = to_check->get_const_col_idxs(); + const auto size = to_check->get_size(); + bool local_is_sorted = true; +#pragma omp parallel for shared(local_is_sorted) + for (size_type i = 0; i < size[0]; ++i) { +#pragma omp flush(local_is_sorted) + // Skip comparison if any thread detects that it is not sorted + if (local_is_sorted) { + for (auto idx = row_ptrs[i] + 1; idx < row_ptrs[i + 1]; ++idx) { + if (col_idxs[idx - 1] > col_idxs[idx]) { + local_is_sorted = false; + break; + } + } + } + } + *is_sorted = local_is_sorted; +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_IS_SORTED_BY_COLUMN_INDEX); diff --git a/omp/test/matrix/CMakeLists.txt b/omp/test/matrix/CMakeLists.txt index 7cdf64bc6b5..108cfc1f44d 100644 --- a/omp/test/matrix/CMakeLists.txt +++ b/omp/test/matrix/CMakeLists.txt @@ -4,3 +4,4 @@ ginkgo_create_test(dense_kernels) ginkgo_create_test(ell_kernels) ginkgo_create_test(hybrid_kernels) ginkgo_create_test(sellp_kernels) +ginkgo_create_test(sparsity_kernels) diff --git a/omp/test/matrix/sparsity_kernels.cpp b/omp/test/matrix/sparsity_kernels.cpp new file mode 100644 index 00000000000..10f2c406189 --- /dev/null +++ b/omp/test/matrix/sparsity_kernels.cpp @@ -0,0 +1,270 @@ +/************************************************************* +Copyright (c) 2017-2019, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include "core/matrix/sparsity_kernels.hpp" + + +#include +#include +#include + +#include + + +#include +#include +#include +#include + + +#include "core/test/utils.hpp" + + +namespace { + + +class Sparsity : public ::testing::Test { +protected: + using Mtx = gko::matrix::Sparsity<>; + using Vec = gko::matrix::Dense<>; + using ComplexVec = gko::matrix::Dense>; + using ComplexMtx = gko::matrix::Sparsity>; + + Sparsity() : mtx_size(532, 231), rand_engine(42) {} + + void SetUp() + { + ref = gko::ReferenceExecutor::create(); + omp = gko::OmpExecutor::create(); + } + + void TearDown() + { + if (omp != nullptr) { + ASSERT_NO_THROW(omp->synchronize()); + } + } + + template + std::unique_ptr gen_mtx(int num_rows, int num_cols, + int min_nnz_row) + { + return gko::test::generate_random_sparsity_matrix( + num_rows, num_cols, + std::uniform_int_distribution<>(min_nnz_row, num_cols), rand_engine, + ref); + } + + void set_up_apply_data(int num_vectors = 1) + { + mtx = Mtx::create(ref); + mtx->copy_from(gen_mtx(mtx_size[0], mtx_size[1], 1)); + complex_mtx = ComplexMtx::create(ref); + complex_mtx->copy_from( + gen_mtx(mtx_size[0], mtx_size[1], 1)); + expected = gen_mtx(mtx_size[0], num_vectors, 1); + y = gen_mtx(mtx_size[1], num_vectors, 1); + alpha = gko::initialize({2.0}, ref); + beta = gko::initialize({-1.0}, ref); + dmtx = Mtx::create(omp); + dmtx->copy_from(mtx.get()); + complex_dmtx = ComplexMtx::create(omp); + complex_dmtx->copy_from(complex_mtx.get()); + dresult = Vec::create(omp); + dresult->copy_from(expected.get()); + dy = Vec::create(omp); + dy->copy_from(y.get()); + dalpha = Vec::create(omp); + dalpha->copy_from(alpha.get()); + dbeta = Vec::create(omp); + dbeta->copy_from(beta.get()); + } + + struct matrix_pair { + std::unique_ptr ref; + std::unique_ptr omp; + }; + + matrix_pair gen_unsorted_mtx() + { + constexpr int min_nnz_per_row = 2; // Must be larger/equal than 2 + auto local_mtx_ref = + gen_mtx(mtx_size[0], mtx_size[1], min_nnz_per_row); + for (size_t row = 0; row < mtx_size[0]; ++row) { + const auto row_ptrs = local_mtx_ref->get_const_row_ptrs(); + const auto start_row = row_ptrs[row]; + auto col_idx = local_mtx_ref->get_col_idxs() + start_row; + const auto nnz_in_this_row = row_ptrs[row + 1] - row_ptrs[row]; + auto swap_idx_dist = + std::uniform_int_distribution<>(0, nnz_in_this_row - 1); + // shuffle `nnz_in_this_row / 2` times + for (size_t perm = 0; perm < nnz_in_this_row; perm += 2) { + const auto idx1 = swap_idx_dist(rand_engine); + const auto idx2 = swap_idx_dist(rand_engine); + std::swap(col_idx[idx1], col_idx[idx2]); + } + } + auto local_mtx_omp = Mtx::create(omp); + local_mtx_omp->copy_from(local_mtx_ref.get()); + + return {std::move(local_mtx_ref), std::move(local_mtx_omp)}; + } + + std::shared_ptr ref; + std::shared_ptr omp; + + const gko::dim<2> mtx_size; + std::ranlux48 rand_engine; + + std::unique_ptr mtx; + std::unique_ptr complex_mtx; + std::unique_ptr expected; + std::unique_ptr y; + std::unique_ptr alpha; + std::unique_ptr beta; + + std::unique_ptr dmtx; + std::unique_ptr complex_dmtx; + std::unique_ptr dresult; + std::unique_ptr dy; + std::unique_ptr dalpha; + std::unique_ptr dbeta; +}; + + +TEST_F(Sparsity, SimpleApplyIsEquivalentToRef) +{ + set_up_apply_data(); + + mtx->apply(y.get(), expected.get()); + dmtx->apply(dy.get(), dresult.get()); + + GKO_ASSERT_MTX_NEAR(dresult, expected, 1e-14); +} + + +TEST_F(Sparsity, AdvancedApplyIsEquivalentToRef) +{ + set_up_apply_data(); + + mtx->apply(alpha.get(), y.get(), beta.get(), expected.get()); + dmtx->apply(dalpha.get(), dy.get(), dbeta.get(), dresult.get()); + + GKO_ASSERT_MTX_NEAR(dresult, expected, 1e-14); +} + + +TEST_F(Sparsity, SimpleApplyToDenseMatrixIsEquivalentToRef) +{ + set_up_apply_data(3); + + mtx->apply(y.get(), expected.get()); + dmtx->apply(dy.get(), dresult.get()); + + GKO_ASSERT_MTX_NEAR(dresult, expected, 1e-14); +} + + +TEST_F(Sparsity, AdvancedApplyToDenseMatrixIsEquivalentToRef) +{ + set_up_apply_data(3); + + mtx->apply(alpha.get(), y.get(), beta.get(), expected.get()); + dmtx->apply(dalpha.get(), dy.get(), dbeta.get(), dresult.get()); + + GKO_ASSERT_MTX_NEAR(dresult, expected, 1e-14); +} + + +TEST_F(Sparsity, TransposeIsEquivalentToRef) +{ + set_up_apply_data(); + + auto trans = mtx->transpose(); + auto d_trans = dmtx->transpose(); + + GKO_ASSERT_MTX_NEAR(static_cast(d_trans.get()), + static_cast(trans.get()), 0.0); +} + + +TEST_F(Sparsity, RecognizeSortedMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + bool is_sorted_omp{}; + bool is_sorted_ref{}; + + is_sorted_ref = mtx->is_sorted_by_column_index(); + is_sorted_omp = dmtx->is_sorted_by_column_index(); + + ASSERT_EQ(is_sorted_ref, is_sorted_omp); +} + + +TEST_F(Sparsity, RecognizeUnsortedMatrixIsEquivalentToRef) +{ + auto uns_mtx = gen_unsorted_mtx(); + bool is_sorted_omp{}; + bool is_sorted_ref{}; + + is_sorted_ref = uns_mtx.ref->is_sorted_by_column_index(); + is_sorted_omp = uns_mtx.omp->is_sorted_by_column_index(); + + ASSERT_EQ(is_sorted_ref, is_sorted_omp); +} + + +TEST_F(Sparsity, SortSortedMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + mtx->sort_by_column_index(); + dmtx->sort_by_column_index(); + + // Values must be unchanged, therefore, tolerance is `0` + GKO_ASSERT_MTX_NEAR(mtx, dmtx, 0); +} + + +TEST_F(Sparsity, SortUnsortedMatrixIsEquivalentToRef) +{ + auto uns_mtx = gen_unsorted_mtx(); + + uns_mtx.ref->sort_by_column_index(); + uns_mtx.omp->sort_by_column_index(); + + // Values must be unchanged, therefore, tolerance is `0` + GKO_ASSERT_MTX_NEAR(uns_mtx.ref, uns_mtx.omp, 0); +} + + +} // namespace From 03d28d35c7fd473794d72e82e362fad3356ec805 Mon Sep 17 00:00:00 2001 From: Pratik Nayak Date: Fri, 13 Sep 2019 01:49:50 +0200 Subject: [PATCH 3/8] Add conversions to Sparsity matrix format from CSR for all execs. --- core/matrix/csr.cpp | 21 +++++++++++++ core/matrix/csr_kernels.hpp | 1 + cuda/test/matrix/csr_kernels.cpp | 26 ++++++++++++++++ include/ginkgo/core/matrix/csr.hpp | 9 ++++++ include/ginkgo/core/matrix/sparsity.hpp | 5 ++++ omp/test/matrix/csr_kernels.cpp | 35 +++++++++++++++++++--- reference/test/matrix/csr_kernels.cpp | 40 +++++++++++++++++++++++++ 7 files changed, 133 insertions(+), 4 deletions(-) diff --git a/core/matrix/csr.cpp b/core/matrix/csr.cpp index 3ae2e50e564..b367f784c5b 100644 --- a/core/matrix/csr.cpp +++ b/core/matrix/csr.cpp @@ -41,6 +41,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include "core/matrix/csr_kernels.hpp" @@ -187,6 +188,26 @@ void Csr::move_to(Sellp *result) } +template +void Csr::convert_to( + Sparsity *result) const +{ + auto exec = this->get_executor(); + auto tmp = Sparsity::create( + exec, this->get_size(), this->get_num_stored_elements()); + tmp->col_idxs_ = this->col_idxs_; + tmp->row_ptrs_ = this->row_ptrs_; + tmp->move_to(result); +} + + +template +void Csr::move_to(Sparsity *result) +{ + this->convert_to(result); +} + + template void Csr::convert_to( Ell *result) const diff --git a/core/matrix/csr_kernels.hpp b/core/matrix/csr_kernels.hpp index 0049011935a..db0c5445495 100644 --- a/core/matrix/csr_kernels.hpp +++ b/core/matrix/csr_kernels.hpp @@ -41,6 +41,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include namespace gko { diff --git a/cuda/test/matrix/csr_kernels.cpp b/cuda/test/matrix/csr_kernels.cpp index 77de599c723..54fd5803f57 100644 --- a/cuda/test/matrix/csr_kernels.cpp +++ b/cuda/test/matrix/csr_kernels.cpp @@ -47,6 +47,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include "core/test/utils.hpp" @@ -388,6 +389,31 @@ TEST_F(Csr, MoveToEllIsEquivalentToRef) GKO_ASSERT_MTX_NEAR(ell_mtx.get(), dell_mtx.get(), 1e-14); } +TEST_F(Csr, ConvertToSparsityIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared()); + auto sparsity_mtx = gko::matrix::Sparsity<>::create(ref); + auto d_sparsity_mtx = gko::matrix::Sparsity<>::create(cuda); + + mtx->convert_to(sparsity_mtx.get()); + dmtx->convert_to(d_sparsity_mtx.get()); + + GKO_ASSERT_MTX_NEAR(sparsity_mtx.get(), d_sparsity_mtx.get(), 1e-14); +} + + +TEST_F(Csr, MoveToSparsityIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared()); + auto sparsity_mtx = gko::matrix::Sparsity<>::create(ref); + auto d_sparsity_mtx = gko::matrix::Sparsity<>::create(cuda); + + mtx->move_to(sparsity_mtx.get()); + dmtx->move_to(d_sparsity_mtx.get()); + + GKO_ASSERT_MTX_NEAR(sparsity_mtx.get(), d_sparsity_mtx.get(), 1e-14); +} + TEST_F(Csr, CalculateMaxNnzPerRowIsEquivalentToRef) { diff --git a/include/ginkgo/core/matrix/csr.hpp b/include/ginkgo/core/matrix/csr.hpp index 936c0180cb3..3c09cbe1cdc 100644 --- a/include/ginkgo/core/matrix/csr.hpp +++ b/include/ginkgo/core/matrix/csr.hpp @@ -57,6 +57,9 @@ class Hybrid; template class Sellp; +template +class Sparsity; + /** * CSR is a matrix format which stores only the nonzero coefficients by @@ -82,6 +85,7 @@ class Csr : public EnableLinOp>, public ConvertibleTo>, public ConvertibleTo>, public ConvertibleTo>, + public ConvertibleTo>, public ReadableFromMatrixData, public WritableToMatrixData, public Transposable { @@ -92,6 +96,7 @@ class Csr : public EnableLinOp>, friend class Ell; friend class Hybrid; friend class Sellp; + friend class Sparsity; public: using EnableLinOp::convert_to; @@ -299,6 +304,10 @@ class Csr : public EnableLinOp>, void move_to(Sellp *result) override; + void convert_to(Sparsity *result) const override; + + void move_to(Sparsity *result) override; + void read(const mat_data &data) override; void write(mat_data &data) const override; diff --git a/include/ginkgo/core/matrix/sparsity.hpp b/include/ginkgo/core/matrix/sparsity.hpp index 76a3e8b8ee7..1b6f0c21503 100644 --- a/include/ginkgo/core/matrix/sparsity.hpp +++ b/include/ginkgo/core/matrix/sparsity.hpp @@ -42,6 +42,10 @@ namespace gko { namespace matrix { +template +class Csr; + + /** * Sparsity is a matrix format which stores only the sparsity pattern of a * sparse matrix by compressing each row of the matrix (compressed sparse row @@ -68,6 +72,7 @@ class Sparsity : public EnableLinOp>, public Transposable { friend class EnableCreateMethod; friend class EnablePolymorphicObject; + friend class Csr; public: using EnableLinOp::convert_to; diff --git a/omp/test/matrix/csr_kernels.cpp b/omp/test/matrix/csr_kernels.cpp index e5254f64ed3..3f3491f30e4 100644 --- a/omp/test/matrix/csr_kernels.cpp +++ b/omp/test/matrix/csr_kernels.cpp @@ -47,6 +47,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include "core/test/utils.hpp" @@ -266,10 +267,10 @@ TEST_F(Csr, ConvertToDenseIsEquivalentToRef) auto dense_mtx = gko::matrix::Dense<>::create(ref); auto ddense_mtx = gko::matrix::Dense<>::create(omp); - mtx->convert_to(ddense_mtx.get()); + mtx->convert_to(dense_mtx.get()); dmtx->convert_to(ddense_mtx.get()); - GKO_ASSERT_MTX_NEAR(dense_mtx.get(), dense_mtx.get(), 1e-14); + GKO_ASSERT_MTX_NEAR(ddense_mtx.get(), dense_mtx.get(), 1e-14); } @@ -279,10 +280,36 @@ TEST_F(Csr, MoveToDenseIsEquivalentToRef) auto dense_mtx = gko::matrix::Dense<>::create(ref); auto ddense_mtx = gko::matrix::Dense<>::create(omp); - mtx->move_to(ddense_mtx.get()); + mtx->move_to(dense_mtx.get()); dmtx->move_to(ddense_mtx.get()); - GKO_ASSERT_MTX_NEAR(dense_mtx.get(), dense_mtx.get(), 1e-14); + GKO_ASSERT_MTX_NEAR(ddense_mtx.get(), dense_mtx.get(), 1e-14); +} + + +TEST_F(Csr, ConvertToSparsityIsEquivalentToRef) +{ + set_up_apply_data(); + auto sparsity_mtx = gko::matrix::Sparsity<>::create(ref); + auto d_sparsity_mtx = gko::matrix::Sparsity<>::create(omp); + + mtx->convert_to(sparsity_mtx.get()); + dmtx->convert_to(d_sparsity_mtx.get()); + + GKO_ASSERT_MTX_NEAR(d_sparsity_mtx.get(), sparsity_mtx.get(), 1e-14); +} + + +TEST_F(Csr, MoveToSparsityIsEquivalentToRef) +{ + set_up_apply_data(); + auto sparsity_mtx = gko::matrix::Sparsity<>::create(ref); + auto d_sparsity_mtx = gko::matrix::Sparsity<>::create(omp); + + mtx->move_to(sparsity_mtx.get()); + dmtx->move_to(d_sparsity_mtx.get()); + + GKO_ASSERT_MTX_NEAR(d_sparsity_mtx.get(), sparsity_mtx.get(), 1e-14); } diff --git a/reference/test/matrix/csr_kernels.cpp b/reference/test/matrix/csr_kernels.cpp index b03e8e810d0..264ee48d931 100644 --- a/reference/test/matrix/csr_kernels.cpp +++ b/reference/test/matrix/csr_kernels.cpp @@ -48,6 +48,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include "core/test/utils/assertions.hpp" @@ -61,6 +62,7 @@ class Csr : public ::testing::Test { using Coo = gko::matrix::Coo<>; using Mtx = gko::matrix::Csr<>; using Sellp = gko::matrix::Sellp<>; + using Sparsity = gko::matrix::Sparsity<>; using Ell = gko::matrix::Ell<>; using Hybrid = gko::matrix::Hybrid<>; using ComplexMtx = gko::matrix::Csr>; @@ -238,6 +240,22 @@ class Csr : public ::testing::Test { EXPECT_EQ(v[129], 0.0); } + void assert_equal_to_mtx(const Sparsity *m) + { + auto *c = m->get_const_col_idxs(); + auto *r = m->get_const_row_ptrs(); + + ASSERT_EQ(m->get_size(), gko::dim<2>(2, 3)); + ASSERT_EQ(m->get_num_nonzeros(), 4); + EXPECT_EQ(r[0], 0); + EXPECT_EQ(r[1], 3); + EXPECT_EQ(r[2], 4); + EXPECT_EQ(c[0], 0); + EXPECT_EQ(c[1], 1); + EXPECT_EQ(c[2], 2); + EXPECT_EQ(c[3], 1); + } + void assert_equal_to_mtx(const Ell *m) { auto v = m->get_const_values(); @@ -474,6 +492,28 @@ TEST_F(Csr, MovesToSellp) } +TEST_F(Csr, ConvertsToSparsity) +{ + auto sparsity_mtx = gko::matrix::Sparsity<>::create(mtx->get_executor()); + + mtx->convert_to(sparsity_mtx.get()); + + assert_equal_to_mtx(sparsity_mtx.get()); +} + + +TEST_F(Csr, MovesToSparsity) +{ + auto sparsity_mtx = gko::matrix::Sparsity<>::create(mtx->get_executor()); + auto csr_ref = gko::matrix::Csr<>::create(mtx->get_executor()); + + csr_ref->copy_from(mtx.get()); + csr_ref->move_to(sparsity_mtx.get()); + + assert_equal_to_mtx(sparsity_mtx.get()); +} + + TEST_F(Csr, ConvertsToHybridAutomatically) { auto hybrid_mtx = gko::matrix::Hybrid<>::create(mtx->get_executor()); From 31fcefd5958524a782a5be857f8f7e844658a230 Mon Sep 17 00:00:00 2001 From: Pratik Nayak Date: Fri, 20 Sep 2019 10:53:22 +0200 Subject: [PATCH 4/8] Add a new constructor that automatically converts from the input LinOp. --- include/ginkgo/core/matrix/sparsity.hpp | 18 ++++ reference/test/matrix/CMakeLists.txt | 1 + reference/test/matrix/sparsity.cpp | 106 ++++++++++++++++++++++++ 3 files changed, 125 insertions(+) create mode 100644 reference/test/matrix/sparsity.cpp diff --git a/include/ginkgo/core/matrix/sparsity.hpp b/include/ginkgo/core/matrix/sparsity.hpp index 1b6f0c21503..32913421c0f 100644 --- a/include/ginkgo/core/matrix/sparsity.hpp +++ b/include/ginkgo/core/matrix/sparsity.hpp @@ -34,6 +34,9 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #define GKO_CORE_MATRIX_SPARSITY_HPP_ +#include + + #include #include @@ -242,6 +245,21 @@ class Sparsity : public EnableLinOp>, GKO_ENSURE_IN_BOUNDS(this->get_size()[0], row_ptrs_.get_num_elems()); } + /** + * Creates a SPARSITY matrix from an existing matrix. Uses the + * `copy_and_convert_to` functionality. + * + * @param exec Executor associated to the matrix + * @param matrix The input matrix + */ + Sparsity(std::shared_ptr exec, + std::shared_ptr matrix) + : EnableLinOp(exec, matrix->get_size()) + { + auto tmp_ = copy_and_convert_to(exec, matrix); + this->copy_from(std::move(tmp_.get())); + } + void apply_impl(const LinOp *b, LinOp *x) const override; void apply_impl(const LinOp *alpha, const LinOp *b, const LinOp *beta, diff --git a/reference/test/matrix/CMakeLists.txt b/reference/test/matrix/CMakeLists.txt index e105acbc3d6..d69c75090b1 100644 --- a/reference/test/matrix/CMakeLists.txt +++ b/reference/test/matrix/CMakeLists.txt @@ -5,4 +5,5 @@ ginkgo_create_test(ell_kernels) ginkgo_create_test(hybrid_kernels) ginkgo_create_test(identity) ginkgo_create_test(sellp_kernels) +ginkgo_create_test(sparsity) ginkgo_create_test(sparsity_kernels) diff --git a/reference/test/matrix/sparsity.cpp b/reference/test/matrix/sparsity.cpp new file mode 100644 index 00000000000..3188d63219b --- /dev/null +++ b/reference/test/matrix/sparsity.cpp @@ -0,0 +1,106 @@ +/************************************************************* +Copyright (c) 2017-2019, the Ginkgo authors +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions +are met: + +1. Redistributions of source code must retain the above copyright +notice, this list of conditions and the following disclaimer. + +2. Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +3. Neither the name of the copyright holder nor the names of its +contributors may be used to endorse or promote products derived from +this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS +IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED +TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A +PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT +HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, +SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT +LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, +DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY +THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT +(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +*************************************************************/ + +#include + + +#include + + +#include +#include +#include +#include +#include + + +#include "core/test/utils/assertions.hpp" + + +namespace { + + +class Sparsity : public ::testing::Test { +protected: + using v_type = double; + using i_type = int; + using Mtx = gko::matrix::Sparsity; + using Csr = gko::matrix::Csr; + using DenseMtx = gko::matrix::Dense; + + Sparsity() + : exec(gko::ReferenceExecutor::create()), + mtx(Mtx::create(exec, gko::dim<2>{2, 3}, 4)) + { + Mtx::index_type *c = mtx->get_col_idxs(); + Mtx::index_type *r = mtx->get_row_ptrs(); + r[0] = 0; + r[1] = 3; + r[2] = 4; + c[0] = 0; + c[1] = 1; + c[2] = 2; + c[3] = 1; + } + + std::shared_ptr exec; + std::unique_ptr mtx; +}; + + +TEST_F(Sparsity, CanBeCreatedFromExistingCsrMatrix) +{ + auto csr_mtx = gko::initialize( + {{2.0, 3.0, 0.0}, {0.0, 1.0, 1.0}, {0.0, 0.0, -3.0}}, exec); + auto comp_mtx = gko::initialize( + {{1.0, 1.0, 0.0}, {0.0, 1.0, 1.0}, {0.0, 0.0, 1.0}}, exec); + + auto mtx = Mtx::create(exec, std::move(csr_mtx)); + + GKO_ASSERT_MTX_NEAR(comp_mtx.get(), mtx.get(), 0.0) +} + + +TEST_F(Sparsity, CanBeCreatedFromExistingDenseMatrix) +{ + auto dense_mtx = gko::initialize( + {{2.0, 3.0, 0.0}, {0.0, 1.0, 1.0}, {0.0, 0.0, -3.0}}, exec); + auto comp_mtx = gko::initialize( + {{1.0, 1.0, 0.0}, {0.0, 1.0, 1.0}, {0.0, 0.0, 1.0}}, exec); + + auto mtx = Mtx::create(exec, std::move(dense_mtx)); + + GKO_ASSERT_MTX_NEAR(comp_mtx.get(), mtx.get(), 0.0) +} + + +} // namespace From 01e798983b4edcac48855282461b37021a9b8f4b Mon Sep 17 00:00:00 2001 From: Pratik Nayak Date: Mon, 30 Sep 2019 15:32:36 +0200 Subject: [PATCH 5/8] Add ref and omp count diag kernel, tests for rem diag kernel. --- include/ginkgo/core/matrix/sparsity.hpp | 1 - omp/matrix/sparsity_kernels.cpp | 17 +++++++- omp/test/matrix/sparsity_kernels.cpp | 40 +++++++++++++++++- reference/matrix/sparsity_kernels.cpp | 16 +++++++- reference/test/matrix/sparsity_kernels.cpp | 47 ++++++++++++++++++++++ 5 files changed, 116 insertions(+), 5 deletions(-) diff --git a/include/ginkgo/core/matrix/sparsity.hpp b/include/ginkgo/core/matrix/sparsity.hpp index 32913421c0f..eac37f99052 100644 --- a/include/ginkgo/core/matrix/sparsity.hpp +++ b/include/ginkgo/core/matrix/sparsity.hpp @@ -83,7 +83,6 @@ class Sparsity : public EnableLinOp>, using value_type = ValueType; using index_type = IndexType; - using value_type = ValueType; using mat_data = matrix_data; diff --git a/omp/matrix/sparsity_kernels.cpp b/omp/matrix/sparsity_kernels.cpp index 4c8e8944e1d..acb27783ee7 100644 --- a/omp/matrix/sparsity_kernels.cpp +++ b/omp/matrix/sparsity_kernels.cpp @@ -125,7 +125,21 @@ template void count_num_diagonal_elements( std::shared_ptr exec, const matrix::Sparsity *matrix, - size_type &num_diagonal_elements) GKO_NOT_IMPLEMENTED; + size_type &num_diagonal_elements) +{ + auto num_rows = matrix->get_size()[0]; + auto row_ptrs = matrix->get_const_row_ptrs(); + auto col_idxs = matrix->get_const_col_idxs(); + size_type num_diag = 0; + for (auto i = 0; i < num_rows; ++i) { + for (auto j = row_ptrs[i]; j < row_ptrs[i + 1]; ++j) { + if (col_idxs[j] == i) { + num_diag++; + } + } + } + num_diagonal_elements = num_diag; +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); @@ -151,6 +165,7 @@ void remove_diagonal_elements(std::shared_ptr exec, } } } +#pragma omp parallel for for (auto i = 0; i < temp_idxs.size(); ++i) { adj_idxs[i] = temp_idxs[i]; } diff --git a/omp/test/matrix/sparsity_kernels.cpp b/omp/test/matrix/sparsity_kernels.cpp index 10f2c406189..ae3cf29f660 100644 --- a/omp/test/matrix/sparsity_kernels.cpp +++ b/omp/test/matrix/sparsity_kernels.cpp @@ -37,6 +37,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include + #include @@ -80,8 +81,8 @@ class Sparsity : public ::testing::Test { { return gko::test::generate_random_sparsity_matrix( num_rows, num_cols, - std::uniform_int_distribution<>(min_nnz_row, num_cols), rand_engine, - ref); + std::uniform_int_distribution<>(min_nnz_row, num_cols), 1.0, + rand_engine, ref); } void set_up_apply_data(int num_vectors = 1) @@ -217,6 +218,41 @@ TEST_F(Sparsity, TransposeIsEquivalentToRef) } +TEST_F(Sparsity, CountsNumberOfDiagElementsIsEqualToRef) +{ + set_up_apply_data(); + gko::size_type num_diags = 0; + gko::size_type d_num_diags = 0; + + gko::kernels::reference::sparsity::count_num_diagonal_elements( + ref, mtx.get(), num_diags); + gko::kernels::omp::sparsity::count_num_diagonal_elements(omp, dmtx.get(), + d_num_diags); + + ASSERT_EQ(d_num_diags, num_diags); +} + + +TEST_F(Sparsity, RemovesDiagElementsKernelIsEquivalentToRef) +{ + set_up_apply_data(); + gko::size_type num_diags = 0; + gko::kernels::reference::sparsity::count_num_diagonal_elements( + ref, mtx.get(), num_diags); + auto tmp = + Mtx::create(ref, mtx->get_size(), mtx->get_num_nonzeros() - num_diags); + auto d_tmp = Mtx::create(omp, dmtx->get_size(), + dmtx->get_num_nonzeros() - num_diags); + gko::kernels::reference::sparsity::remove_diagonal_elements( + ref, tmp.get(), mtx->get_const_row_ptrs(), mtx->get_const_col_idxs()); + gko::kernels::omp::sparsity::remove_diagonal_elements( + omp, d_tmp.get(), dmtx->get_const_row_ptrs(), + dmtx->get_const_col_idxs()); + + GKO_ASSERT_MTX_NEAR(tmp.get(), d_tmp.get(), 0.0); +} + + TEST_F(Sparsity, RecognizeSortedMatrixIsEquivalentToRef) { set_up_apply_data(); diff --git a/reference/matrix/sparsity_kernels.cpp b/reference/matrix/sparsity_kernels.cpp index 72dd4282e69..65e52cc5f83 100644 --- a/reference/matrix/sparsity_kernels.cpp +++ b/reference/matrix/sparsity_kernels.cpp @@ -120,7 +120,21 @@ template void count_num_diagonal_elements( std::shared_ptr exec, const matrix::Sparsity *matrix, - size_type &num_diagonal_elements) GKO_NOT_IMPLEMENTED; + size_type &num_diagonal_elements) +{ + auto num_rows = matrix->get_size()[0]; + auto row_ptrs = matrix->get_const_row_ptrs(); + auto col_idxs = matrix->get_const_col_idxs(); + size_type num_diag = 0; + for (auto i = 0; i < num_rows; ++i) { + for (auto j = row_ptrs[i]; j < row_ptrs[i + 1]; ++j) { + if (col_idxs[j] == i) { + num_diag++; + } + } + } + num_diagonal_elements = num_diag; +} GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); diff --git a/reference/test/matrix/sparsity_kernels.cpp b/reference/test/matrix/sparsity_kernels.cpp index 81bb392950c..5b4cdfdb584 100644 --- a/reference/test/matrix/sparsity_kernels.cpp +++ b/reference/test/matrix/sparsity_kernels.cpp @@ -272,6 +272,53 @@ TEST_F(Sparsity, NonSquareMtxIsTransposable) } +TEST_F(Sparsity, CountsCorrectNumberOfDiagonalElements) +{ + // clang-format off + auto mtx2 = gko::initialize>( + {{1.0, 1.0, 1.0}, + {0.0, 1.0, 0.0}, + {0.0, 1.0, 1.0}}, exec); + auto mtx_s = gko::initialize>( + {{1.0, 1.0, 1.0}, + {0.0, 0.0, 0.0}, + {0.0, 1.0, 1.0}}, exec); + // clang-format on + gko::size_type m1_num_diags = 0; + gko::size_type ms_num_diags = 0; + gko::kernels::reference::sparsity::count_num_diagonal_elements( + exec, mtx2.get(), m1_num_diags); + gko::kernels::reference::sparsity::count_num_diagonal_elements( + exec, mtx_s.get(), ms_num_diags); + + ASSERT_EQ(m1_num_diags, 3); + ASSERT_EQ(ms_num_diags, 2); +} + + +TEST_F(Sparsity, RemovesDiagonalElements) +{ + // clang-format off + auto mtx2 = gko::initialize>( + {{1.0, 1.0, 1.0}, + {0.0, 1.0, 0.0}, + {0.0, 1.0, 1.0}}, exec); + auto mtx_s = gko::initialize>( + {{0.0, 1.0, 1.0}, + {0.0, 0.0, 0.0}, + {0.0, 1.0, 0.0}}, exec); + // clang-format on + auto tmp_mtx = gko::matrix::Sparsity<>::create(exec, mtx_s->get_size(), + mtx_s->get_num_nonzeros()); + tmp_mtx->copy_from(mtx2.get()); + gko::kernels::reference::sparsity::remove_diagonal_elements( + exec, tmp_mtx.get(), mtx2->get_const_row_ptrs(), + mtx2->get_const_col_idxs()); + + GKO_ASSERT_MTX_NEAR(tmp_mtx.get(), mtx_s.get(), 0.0); +} + + TEST_F(Sparsity, SquareMtxIsConvertibleToAdjacencyMatrix) { // clang-format off From 8a588b87772241bd5ed95353f332bbc659c26878 Mon Sep 17 00:00:00 2001 From: Pratik Nayak Date: Tue, 1 Oct 2019 10:55:35 +0200 Subject: [PATCH 6/8] Review update + Fix bugs in diag removal, add more tests. + Fix value setting in CSR and dense conversions to Sparsity --- core/matrix/csr.cpp | 1 + core/matrix/sparsity.cpp | 4 +- include/ginkgo/core/matrix/dense.hpp | 1 - omp/matrix/dense_kernels.cpp | 3 +- omp/matrix/sparsity_kernels.cpp | 11 ++- reference/matrix/dense_kernels.cpp | 4 +- reference/matrix/sparsity_kernels.cpp | 11 ++- reference/test/matrix/sparsity_kernels.cpp | 85 ++++++++++++++-------- 8 files changed, 79 insertions(+), 41 deletions(-) diff --git a/core/matrix/csr.cpp b/core/matrix/csr.cpp index b367f784c5b..2aff276bce5 100644 --- a/core/matrix/csr.cpp +++ b/core/matrix/csr.cpp @@ -197,6 +197,7 @@ void Csr::convert_to( exec, this->get_size(), this->get_num_stored_elements()); tmp->col_idxs_ = this->col_idxs_; tmp->row_ptrs_ = this->row_ptrs_; + tmp->value_ = gko::Array(exec, {one()}); tmp->move_to(result); } diff --git a/core/matrix/sparsity.cpp b/core/matrix/sparsity.cpp index d98d6172d2d..a74f7069c7a 100644 --- a/core/matrix/sparsity.cpp +++ b/core/matrix/sparsity.cpp @@ -164,14 +164,12 @@ template std::unique_ptr> Sparsity::to_adjacency_matrix() const { - // Adjacency matrix has to be square. - GKO_ASSERT_IS_SQUARE_MATRIX(this); auto exec = this->get_executor(); + // Adjacency matrix has to be square. GKO_ASSERT_IS_SQUARE_MATRIX(this); size_type num_diagonal_elements = 0; exec->run(sparsity::make_count_num_diagonal_elements( this, num_diagonal_elements)); - ValueType one = 1.0; auto adj_mat = Sparsity::create(exec, this->get_size(), this->get_num_nonzeros() - num_diagonal_elements); diff --git a/include/ginkgo/core/matrix/dense.hpp b/include/ginkgo/core/matrix/dense.hpp index ea568a02f3a..ce0f835ba93 100644 --- a/include/ginkgo/core/matrix/dense.hpp +++ b/include/ginkgo/core/matrix/dense.hpp @@ -194,7 +194,6 @@ class Dense : public EnableLinOp>, void move_to(Sparsity *result) override; - void read(const mat_data &data) override; void read(const mat_data32 &data) override; diff --git a/omp/matrix/dense_kernels.cpp b/omp/matrix/dense_kernels.cpp index 07d33cb92b4..c698f92fb40 100644 --- a/omp/matrix/dense_kernels.cpp +++ b/omp/matrix/dense_kernels.cpp @@ -463,10 +463,11 @@ void convert_to_sparsity(std::shared_ptr exec, { auto num_rows = result->get_size()[0]; auto num_cols = result->get_size()[1]; - auto num_nonzeros = result->get_num_nonzeros(); auto row_ptrs = result->get_row_ptrs(); auto col_idxs = result->get_col_idxs(); + auto value = result->get_value(); + value[0] = one(); size_type cur_ptr = 0; row_ptrs[0] = cur_ptr; diff --git a/omp/matrix/sparsity_kernels.cpp b/omp/matrix/sparsity_kernels.cpp index acb27783ee7..bc781f6b429 100644 --- a/omp/matrix/sparsity_kernels.cpp +++ b/omp/matrix/sparsity_kernels.cpp @@ -154,8 +154,15 @@ void remove_diagonal_elements(std::shared_ptr exec, auto num_rows = matrix->get_size()[0]; auto adj_ptrs = matrix->get_row_ptrs(); auto adj_idxs = matrix->get_col_idxs(); - for (auto i = 0; i <= num_rows; ++i) { - adj_ptrs[i] = row_ptrs[i] - i; + size_type num_diag = 0; + adj_ptrs[0] = row_ptrs[0]; + for (auto i = 0; i < num_rows; ++i) { + for (auto j = row_ptrs[i]; j < row_ptrs[i + 1]; ++j) { + if (col_idxs[j] == i) { + num_diag++; + } + } + adj_ptrs[i + 1] = row_ptrs[i + 1] - num_diag; } std::vector temp_idxs; for (auto i = 0; i < num_rows; ++i) { diff --git a/reference/matrix/dense_kernels.cpp b/reference/matrix/dense_kernels.cpp index ab3d76318b7..57f45e580e8 100644 --- a/reference/matrix/dense_kernels.cpp +++ b/reference/matrix/dense_kernels.cpp @@ -422,11 +422,11 @@ void convert_to_sparsity(std::shared_ptr exec, { auto num_rows = result->get_size()[0]; auto num_cols = result->get_size()[1]; - auto num_nonzeros = result->get_num_nonzeros(); auto row_ptrs = result->get_row_ptrs(); auto col_idxs = result->get_col_idxs(); - + auto value = result->get_value(); + value[0] = one(); size_type cur_ptr = 0; row_ptrs[0] = cur_ptr; for (size_type row = 0; row < num_rows; ++row) { diff --git a/reference/matrix/sparsity_kernels.cpp b/reference/matrix/sparsity_kernels.cpp index 65e52cc5f83..aa8a0ccfdf2 100644 --- a/reference/matrix/sparsity_kernels.cpp +++ b/reference/matrix/sparsity_kernels.cpp @@ -149,8 +149,15 @@ void remove_diagonal_elements(std::shared_ptr exec, auto num_rows = matrix->get_size()[0]; auto adj_ptrs = matrix->get_row_ptrs(); auto adj_idxs = matrix->get_col_idxs(); - for (auto i = 0; i <= num_rows; ++i) { - adj_ptrs[i] = row_ptrs[i] - i; + size_type num_diag = 0; + adj_ptrs[0] = row_ptrs[0]; + for (auto i = 0; i < num_rows; ++i) { + for (auto j = row_ptrs[i]; j < row_ptrs[i + 1]; ++j) { + if (col_idxs[j] == i) { + num_diag++; + } + } + adj_ptrs[i + 1] = row_ptrs[i + 1] - num_diag; } std::vector temp_idxs; for (auto i = 0; i < num_rows; ++i) { diff --git a/reference/test/matrix/sparsity_kernels.cpp b/reference/test/matrix/sparsity_kernels.cpp index 5b4cdfdb584..ffa18d1968e 100644 --- a/reference/test/matrix/sparsity_kernels.cpp +++ b/reference/test/matrix/sparsity_kernels.cpp @@ -275,42 +275,44 @@ TEST_F(Sparsity, NonSquareMtxIsTransposable) TEST_F(Sparsity, CountsCorrectNumberOfDiagonalElements) { // clang-format off - auto mtx2 = gko::initialize>( - {{1.0, 1.0, 1.0}, - {0.0, 1.0, 0.0}, - {0.0, 1.0, 1.0}}, exec); - auto mtx_s = gko::initialize>( - {{1.0, 1.0, 1.0}, - {0.0, 0.0, 0.0}, - {0.0, 1.0, 1.0}}, exec); + auto mtx2 = gko::initialize>( + {{1.0, 1.0, 1.0}, + {0.0, 1.0, 0.0}, + {0.0, 1.0, 1.0}}, exec); + auto mtx_s = gko::initialize>( + {{1.0, 1.0, 1.0}, + {0.0, 0.0, 0.0}, + {0.0, 1.0, 1.0}}, exec); // clang-format on - gko::size_type m1_num_diags = 0; + gko::size_type m2_num_diags = 0; gko::size_type ms_num_diags = 0; + gko::kernels::reference::sparsity::count_num_diagonal_elements( - exec, mtx2.get(), m1_num_diags); + exec, mtx2.get(), m2_num_diags); gko::kernels::reference::sparsity::count_num_diagonal_elements( exec, mtx_s.get(), ms_num_diags); - ASSERT_EQ(m1_num_diags, 3); + ASSERT_EQ(m2_num_diags, 3); ASSERT_EQ(ms_num_diags, 2); } -TEST_F(Sparsity, RemovesDiagonalElements) +TEST_F(Sparsity, RemovesDiagonalElementsForFullRankMatrix) { // clang-format off - auto mtx2 = gko::initialize>( - {{1.0, 1.0, 1.0}, - {0.0, 1.0, 0.0}, - {0.0, 1.0, 1.0}}, exec); - auto mtx_s = gko::initialize>( - {{0.0, 1.0, 1.0}, - {0.0, 0.0, 0.0}, - {0.0, 1.0, 0.0}}, exec); + auto mtx2 = gko::initialize>( + {{1.0, 1.0, 1.0}, + {0.0, 1.0, 0.0}, + {0.0, 1.0, 1.0}}, exec); + auto mtx_s = gko::initialize>( + {{0.0, 1.0, 1.0}, + {0.0, 0.0, 0.0}, + {0.0, 1.0, 0.0}}, exec); // clang-format on auto tmp_mtx = gko::matrix::Sparsity<>::create(exec, mtx_s->get_size(), mtx_s->get_num_nonzeros()); tmp_mtx->copy_from(mtx2.get()); + gko::kernels::reference::sparsity::remove_diagonal_elements( exec, tmp_mtx.get(), mtx2->get_const_row_ptrs(), mtx2->get_const_col_idxs()); @@ -319,23 +321,46 @@ TEST_F(Sparsity, RemovesDiagonalElements) } -TEST_F(Sparsity, SquareMtxIsConvertibleToAdjacencyMatrix) +TEST_F(Sparsity, RemovesDiagonalElementsForIncompleteRankMatrix) { // clang-format off - auto mtx2 = gko::initialize>( - {{1.0, 1.0, 1.0}, - {0.0, 1.0, 0.0}, - {0.0, 1.0, 1.0}}, exec); + auto mtx2 = gko::initialize>( + {{1.0, 1.0, 1.0}, + {0.0, 0.0, 0.0}, + {0.0, 1.0, 1.0}}, exec); + auto mtx_s = gko::initialize>( + {{0.0, 1.0, 1.0}, + {0.0, 0.0, 0.0}, + {0.0, 1.0, 0.0}}, exec); // clang-format on + auto tmp_mtx = gko::matrix::Sparsity<>::create(exec, mtx_s->get_size(), + mtx_s->get_num_nonzeros()); + tmp_mtx->copy_from(mtx2.get()); - auto adj_mat = mtx2->to_adjacency_matrix(); + gko::kernels::reference::sparsity::remove_diagonal_elements( + exec, tmp_mtx.get(), mtx2->get_const_row_ptrs(), + mtx2->get_const_col_idxs()); + GKO_ASSERT_MTX_NEAR(tmp_mtx.get(), mtx_s.get(), 0.0); +} + + +TEST_F(Sparsity, SquareMtxIsConvertibleToAdjacencyMatrix) +{ // clang-format off - GKO_ASSERT_MTX_NEAR(adj_mat.get(), - l({{0.0, 1.0, 1.0}, - {0.0, 0.0, 0.0}, - {0.0, 1.0, 0.0}}), 0.0); + auto mtx2 = gko::initialize>( + {{1.0, 1.0, 1.0}, + {0.0, 1.0, 0.0}, + {0.0, 1.0, 1.0}}, exec); + auto mtx_s = gko::initialize>( + {{0.0, 1.0, 1.0}, + {0.0, 0.0, 0.0}, + {0.0, 1.0, 0.0}}, exec); // clang-format on + + auto adj_mat = mtx2->to_adjacency_matrix(); + + GKO_ASSERT_MTX_NEAR(adj_mat.get(), mtx_s.get(), 0.0); } From b8bd4773d90d4bed7b2ef588f8a8c93085e1e989 Mon Sep 17 00:00:00 2001 From: Pratik Nayak Date: Wed, 2 Oct 2019 12:11:48 +0200 Subject: [PATCH 7/8] Rename Sparsity to SparsityCsr. --- core/CMakeLists.txt | 2 +- core/device_hooks/common_kernels.inc.cpp | 40 +++--- core/matrix/csr.cpp | 9 +- core/matrix/csr_kernels.hpp | 2 +- core/matrix/dense.cpp | 16 +-- core/matrix/dense_kernels.hpp | 76 +++++------ .../matrix/{sparsity.cpp => sparsity_csr.cpp} | 80 ++++++------ ...y_kernels.hpp => sparsity_csr_kernels.hpp} | 120 +++++++++--------- core/test/matrix/CMakeLists.txt | 2 +- .../matrix/{sparsity.cpp => sparsity_csr.cpp} | 38 +++--- cuda/CMakeLists.txt | 2 +- cuda/matrix/dense_kernels.cu | 6 +- ...ity_kernels.cu => sparsity_csr_kernels.cu} | 37 +++--- cuda/test/matrix/csr_kernels.cpp | 14 +- include/ginkgo/core/matrix/csr.hpp | 10 +- include/ginkgo/core/matrix/dense.hpp | 18 +-- .../matrix/{sparsity.hpp => sparsity_csr.hpp} | 67 +++++----- include/ginkgo/ginkgo.hpp | 2 +- omp/CMakeLists.txt | 2 +- omp/matrix/dense_kernels.cpp | 6 +- ...y_kernels.cpp => sparsity_csr_kernels.cpp} | 46 +++---- omp/test/matrix/CMakeLists.txt | 2 +- omp/test/matrix/csr_kernels.cpp | 14 +- ...y_kernels.cpp => sparsity_csr_kernels.cpp} | 46 +++---- reference/CMakeLists.txt | 2 +- reference/matrix/dense_kernels.cpp | 6 +- ...y_kernels.cpp => sparsity_csr_kernels.cpp} | 48 +++---- reference/test/matrix/CMakeLists.txt | 4 +- reference/test/matrix/csr_kernels.cpp | 14 +- .../matrix/{sparsity.cpp => sparsity_csr.cpp} | 12 +- ...y_kernels.cpp => sparsity_csr_kernels.cpp} | 84 ++++++------ test_install/test_install.cpp | 4 +- 32 files changed, 422 insertions(+), 409 deletions(-) rename core/matrix/{sparsity.cpp => sparsity_csr.cpp} (66%) rename core/matrix/{sparsity_kernels.hpp => sparsity_csr_kernels.hpp} (57%) rename core/test/matrix/{sparsity.cpp => sparsity_csr.cpp} (86%) rename cuda/matrix/{sparsity_kernels.cu => sparsity_csr_kernels.cu} (79%) rename include/ginkgo/core/matrix/{sparsity.hpp => sparsity_csr.hpp} (80%) rename omp/matrix/{sparsity_kernels.cpp => sparsity_csr_kernels.cpp} (86%) rename omp/test/matrix/{sparsity_kernels.cpp => sparsity_csr_kernels.cpp} (85%) rename reference/matrix/{sparsity_kernels.cpp => sparsity_csr_kernels.cpp} (85%) rename reference/test/matrix/{sparsity.cpp => sparsity_csr.cpp} (92%) rename reference/test/matrix/{sparsity_kernels.cpp => sparsity_csr_kernels.cpp} (81%) diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt index b147d64291d..a976e362d97 100644 --- a/core/CMakeLists.txt +++ b/core/CMakeLists.txt @@ -22,7 +22,7 @@ target_sources(ginkgo matrix/hybrid.cpp matrix/identity.cpp matrix/sellp.cpp - matrix/sparsity.cpp + matrix/sparsity_csr.cpp preconditioner/jacobi.cpp solver/bicgstab.cpp solver/cg.cpp diff --git a/core/device_hooks/common_kernels.inc.cpp b/core/device_hooks/common_kernels.inc.cpp index e338c731ed4..05a7f8bc136 100644 --- a/core/device_hooks/common_kernels.inc.cpp +++ b/core/device_hooks/common_kernels.inc.cpp @@ -40,7 +40,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "core/matrix/ell_kernels.hpp" #include "core/matrix/hybrid_kernels.hpp" #include "core/matrix/sellp_kernels.hpp" -#include "core/matrix/sparsity_kernels.hpp" +#include "core/matrix/sparsity_csr_kernels.hpp" #include "core/preconditioner/jacobi_kernels.hpp" #include "core/solver/bicgstab_kernels.hpp" #include "core/solver/cg_kernels.hpp" @@ -126,10 +126,10 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL); template -GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL(ValueType, IndexType) +GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL); + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL); template GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL(ValueType) @@ -358,52 +358,54 @@ GKO_NOT_COMPILED(GKO_HOOK_MODULE); } // namespace ir -namespace sparsity { +namespace sparsity_csr { template -GKO_DECLARE_SPARSITY_SPMV_KERNEL(ValueType, IndexType) +GKO_DECLARE_SPARSITY_CSR_SPMV_KERNEL(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SPARSITY_SPMV_KERNEL); +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_SPMV_KERNEL); template -GKO_DECLARE_SPARSITY_ADVANCED_SPMV_KERNEL(ValueType, IndexType) +GKO_DECLARE_SPARSITY_CSR_ADVANCED_SPMV_KERNEL(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_ADVANCED_SPMV_KERNEL); + GKO_DECLARE_SPARSITY_CSR_ADVANCED_SPMV_KERNEL); template -GKO_DECLARE_SPARSITY_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL(ValueType, IndexType) +GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL(ValueType, + IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); + GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); template -GKO_DECLARE_SPARSITY_REMOVE_DIAGONAL_ELEMENTS_KERNEL(ValueType, IndexType) +GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_REMOVE_DIAGONAL_ELEMENTS_KERNEL); + GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL); template -GKO_DECLARE_SPARSITY_TRANSPOSE_KERNEL(ValueType, IndexType) +GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_TRANSPOSE_KERNEL); + GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL); template -GKO_DECLARE_SPARSITY_SORT_BY_COLUMN_INDEX(ValueType, IndexType) +GKO_DECLARE_SPARSITY_CSR_SORT_BY_COLUMN_INDEX(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_SORT_BY_COLUMN_INDEX); + GKO_DECLARE_SPARSITY_CSR_SORT_BY_COLUMN_INDEX); template -GKO_DECLARE_SPARSITY_IS_SORTED_BY_COLUMN_INDEX(ValueType, IndexType) +GKO_DECLARE_SPARSITY_CSR_IS_SORTED_BY_COLUMN_INDEX(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_IS_SORTED_BY_COLUMN_INDEX); + GKO_DECLARE_SPARSITY_CSR_IS_SORTED_BY_COLUMN_INDEX); -} // namespace sparsity +} // namespace sparsity_csr namespace csr { diff --git a/core/matrix/csr.cpp b/core/matrix/csr.cpp index 2aff276bce5..d7e56eb11a5 100644 --- a/core/matrix/csr.cpp +++ b/core/matrix/csr.cpp @@ -41,7 +41,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include +#include #include "core/matrix/csr_kernels.hpp" @@ -190,10 +190,10 @@ void Csr::move_to(Sellp *result) template void Csr::convert_to( - Sparsity *result) const + SparsityCsr *result) const { auto exec = this->get_executor(); - auto tmp = Sparsity::create( + auto tmp = SparsityCsr::create( exec, this->get_size(), this->get_num_stored_elements()); tmp->col_idxs_ = this->col_idxs_; tmp->row_ptrs_ = this->row_ptrs_; @@ -203,7 +203,8 @@ void Csr::convert_to( template -void Csr::move_to(Sparsity *result) +void Csr::move_to( + SparsityCsr *result) { this->convert_to(result); } diff --git a/core/matrix/csr_kernels.hpp b/core/matrix/csr_kernels.hpp index db0c5445495..2e099594e70 100644 --- a/core/matrix/csr_kernels.hpp +++ b/core/matrix/csr_kernels.hpp @@ -41,7 +41,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include +#include namespace gko { diff --git a/core/matrix/dense.cpp b/core/matrix/dense.cpp index c24afc1ba12..32e1c35fa02 100644 --- a/core/matrix/dense.cpp +++ b/core/matrix/dense.cpp @@ -43,7 +43,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include +#include #include "core/matrix/dense_kernels.hpp" @@ -185,15 +185,15 @@ inline void conversion_helper(Sellp *result, template -inline void conversion_helper(Sparsity *result, +inline void conversion_helper(SparsityCsr *result, MatrixType *source, const OperationType &op) { auto exec = source->get_executor(); size_type num_stored_nonzeros = 0; exec->run(dense::make_count_nonzeros(source, &num_stored_nonzeros)); - auto tmp = Sparsity::create(exec, source->get_size(), - num_stored_nonzeros); + auto tmp = SparsityCsr::create( + exec, source->get_size(), num_stored_nonzeros); exec->run(op(tmp.get(), source)); tmp->move_to(result); } @@ -442,7 +442,7 @@ void Dense::move_to(Sellp *result) template -void Dense::convert_to(Sparsity *result) const +void Dense::convert_to(SparsityCsr *result) const { conversion_helper( result, this, @@ -452,14 +452,14 @@ void Dense::convert_to(Sparsity *result) const template -void Dense::move_to(Sparsity *result) +void Dense::move_to(SparsityCsr *result) { this->convert_to(result); } template -void Dense::convert_to(Sparsity *result) const +void Dense::convert_to(SparsityCsr *result) const { conversion_helper( result, this, @@ -469,7 +469,7 @@ void Dense::convert_to(Sparsity *result) const template -void Dense::move_to(Sparsity *result) +void Dense::move_to(SparsityCsr *result) { this->convert_to(result); } diff --git a/core/matrix/dense_kernels.hpp b/core/matrix/dense_kernels.hpp index 66f0ca41111..5e6169d4732 100644 --- a/core/matrix/dense_kernels.hpp +++ b/core/matrix/dense_kernels.hpp @@ -98,9 +98,9 @@ namespace kernels { matrix::Sellp<_type, _prec> *other, \ const matrix::Dense<_type> *source) -#define GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL(_type, _prec) \ +#define GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL(_type, _prec) \ void convert_to_sparsity(std::shared_ptr exec, \ - matrix::Sparsity<_type, _prec> *other, \ + matrix::SparsityCsr<_type, _prec> *other, \ const matrix::Dense<_type> *source) #define GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL(_type) \ @@ -133,42 +133,42 @@ namespace kernels { matrix::Dense<_type> *trans, \ const matrix::Dense<_type> *orig) -#define GKO_DECLARE_ALL_AS_TEMPLATES \ - template \ - GKO_DECLARE_DENSE_SIMPLE_APPLY_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_APPLY_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_SCALE_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_ADD_SCALED_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_COMPUTE_DOT_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_COMPUTE_NORM2_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_CONVERT_TO_COO_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_DENSE_CONVERT_TO_CSR_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_DENSE_CONVERT_TO_ELL_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_DENSE_CONVERT_TO_HYBRID_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_CALCULATE_MAX_NNZ_PER_ROW_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_CALCULATE_NONZEROS_PER_ROW_KERNEL(ValueType); \ - template \ - GKO_DECLARE_DENSE_CALCULATE_TOTAL_COLS_KERNEL(ValueType); \ - template \ - GKO_DECLARE_TRANSPOSE_KERNEL(ValueType); \ - template \ +#define GKO_DECLARE_ALL_AS_TEMPLATES \ + template \ + GKO_DECLARE_DENSE_SIMPLE_APPLY_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_APPLY_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_SCALE_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_ADD_SCALED_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_COMPUTE_DOT_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_COMPUTE_NORM2_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_CONVERT_TO_COO_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_DENSE_CONVERT_TO_CSR_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_DENSE_CONVERT_TO_ELL_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_DENSE_CONVERT_TO_HYBRID_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_CALCULATE_MAX_NNZ_PER_ROW_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_CALCULATE_NONZEROS_PER_ROW_KERNEL(ValueType); \ + template \ + GKO_DECLARE_DENSE_CALCULATE_TOTAL_COLS_KERNEL(ValueType); \ + template \ + GKO_DECLARE_TRANSPOSE_KERNEL(ValueType); \ + template \ GKO_DECLARE_CONJ_TRANSPOSE_KERNEL(ValueType) diff --git a/core/matrix/sparsity.cpp b/core/matrix/sparsity_csr.cpp similarity index 66% rename from core/matrix/sparsity.cpp rename to core/matrix/sparsity_csr.cpp index a74f7069c7a..985fd03bc1e 100644 --- a/core/matrix/sparsity.cpp +++ b/core/matrix/sparsity_csr.cpp @@ -30,7 +30,7 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include +#include #include @@ -43,59 +43,61 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include -#include "core/matrix/sparsity_kernels.hpp" +#include "core/matrix/sparsity_csr_kernels.hpp" namespace gko { namespace matrix { -namespace sparsity { +namespace sparsity_csr { -GKO_REGISTER_OPERATION(spmv, sparsity::spmv); -GKO_REGISTER_OPERATION(advanced_spmv, sparsity::advanced_spmv); -GKO_REGISTER_OPERATION(transpose, sparsity::transpose); +GKO_REGISTER_OPERATION(spmv, sparsity_csr::spmv); +GKO_REGISTER_OPERATION(advanced_spmv, sparsity_csr::advanced_spmv); +GKO_REGISTER_OPERATION(transpose, sparsity_csr::transpose); GKO_REGISTER_OPERATION(count_num_diagonal_elements, - sparsity::count_num_diagonal_elements); + sparsity_csr::count_num_diagonal_elements); GKO_REGISTER_OPERATION(remove_diagonal_elements, - sparsity::remove_diagonal_elements); -GKO_REGISTER_OPERATION(sort_by_column_index, sparsity::sort_by_column_index); + sparsity_csr::remove_diagonal_elements); +GKO_REGISTER_OPERATION(sort_by_column_index, + sparsity_csr::sort_by_column_index); GKO_REGISTER_OPERATION(is_sorted_by_column_index, - sparsity::is_sorted_by_column_index); + sparsity_csr::is_sorted_by_column_index); -} // namespace sparsity +} // namespace sparsity_csr template -void Sparsity::apply_impl(const LinOp *b, LinOp *x) const +void SparsityCsr::apply_impl(const LinOp *b, + LinOp *x) const { using Dense = Dense; this->get_executor()->run( - sparsity::make_spmv(this, as(b), as(x))); + sparsity_csr::make_spmv(this, as(b), as(x))); } template -void Sparsity::apply_impl(const LinOp *alpha, - const LinOp *b, - const LinOp *beta, - LinOp *x) const +void SparsityCsr::apply_impl(const LinOp *alpha, + const LinOp *b, + const LinOp *beta, + LinOp *x) const { using Dense = Dense; - this->get_executor()->run(sparsity::make_advanced_spmv( + this->get_executor()->run(sparsity_csr::make_advanced_spmv( as(alpha), this, as(b), as(beta), as(x))); } template -void Sparsity::read(const mat_data &data) +void SparsityCsr::read(const mat_data &data) { size_type nnz = 0; for (const auto &elem : data.nonzeros) { nnz += (elem.value != zero()); } auto tmp = - Sparsity::create(this->get_executor()->get_master(), data.size, nnz); + SparsityCsr::create(this->get_executor()->get_master(), data.size, nnz); size_type ind = 0; size_type cur_ptr = 0; tmp->get_row_ptrs()[0] = cur_ptr; @@ -118,13 +120,13 @@ void Sparsity::read(const mat_data &data) template -void Sparsity::write(mat_data &data) const +void SparsityCsr::write(mat_data &data) const { std::unique_ptr op{}; - const Sparsity *tmp{}; + const SparsityCsr *tmp{}; if (this->get_executor()->get_master() != this->get_executor()) { op = this->clone(this->get_executor()->get_master()); - tmp = static_cast(op.get()); + tmp = static_cast(op.get()); } else { tmp = this; } @@ -144,62 +146,62 @@ void Sparsity::write(mat_data &data) const template -std::unique_ptr Sparsity::transpose() const +std::unique_ptr SparsityCsr::transpose() const { auto exec = this->get_executor(); - auto trans_cpy = Sparsity::create(exec, gko::transpose(this->get_size()), - this->get_num_nonzeros()); + auto trans_cpy = SparsityCsr::create(exec, gko::transpose(this->get_size()), + this->get_num_nonzeros()); - exec->run(sparsity::make_transpose(trans_cpy.get(), this)); + exec->run(sparsity_csr::make_transpose(trans_cpy.get(), this)); return std::move(trans_cpy); } template -std::unique_ptr Sparsity::conj_transpose() const +std::unique_ptr SparsityCsr::conj_transpose() const GKO_NOT_IMPLEMENTED; template -std::unique_ptr> -Sparsity::to_adjacency_matrix() const +std::unique_ptr> +SparsityCsr::to_adjacency_matrix() const { auto exec = this->get_executor(); // Adjacency matrix has to be square. GKO_ASSERT_IS_SQUARE_MATRIX(this); size_type num_diagonal_elements = 0; - exec->run(sparsity::make_count_num_diagonal_elements( + exec->run(sparsity_csr::make_count_num_diagonal_elements( this, num_diagonal_elements)); auto adj_mat = - Sparsity::create(exec, this->get_size(), - this->get_num_nonzeros() - num_diagonal_elements); + SparsityCsr::create(exec, this->get_size(), + this->get_num_nonzeros() - num_diagonal_elements); - exec->run(sparsity::make_remove_diagonal_elements( + exec->run(sparsity_csr::make_remove_diagonal_elements( adj_mat.get(), this->get_const_row_ptrs(), this->get_const_col_idxs())); return std::move(adj_mat); } template -void Sparsity::sort_by_column_index() +void SparsityCsr::sort_by_column_index() { auto exec = this->get_executor(); - exec->run(sparsity::make_sort_by_column_index(this)); + exec->run(sparsity_csr::make_sort_by_column_index(this)); } template -bool Sparsity::is_sorted_by_column_index() const +bool SparsityCsr::is_sorted_by_column_index() const { auto exec = this->get_executor(); bool is_sorted; - exec->run(sparsity::make_is_sorted_by_column_index(this, &is_sorted)); + exec->run(sparsity_csr::make_is_sorted_by_column_index(this, &is_sorted)); return is_sorted; } #define GKO_DECLARE_SPARSITY_MATRIX(ValueType, IndexType) \ - class Sparsity + class SparsityCsr GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SPARSITY_MATRIX); diff --git a/core/matrix/sparsity_kernels.hpp b/core/matrix/sparsity_csr_kernels.hpp similarity index 57% rename from core/matrix/sparsity_kernels.hpp rename to core/matrix/sparsity_csr_kernels.hpp index 4f50a7ee62f..1fd40ca06a3 100644 --- a/core/matrix/sparsity_kernels.hpp +++ b/core/matrix/sparsity_csr_kernels.hpp @@ -30,104 +30,106 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#ifndef GKO_CORE_MATRIX_SPARSITY_KERNELS_HPP_ -#define GKO_CORE_MATRIX_SPARSITY_KERNELS_HPP_ +#ifndef GKO_CORE_MATRIX_SPARSITY_CSR_KERNELS_HPP_ +#define GKO_CORE_MATRIX_SPARSITY_CSR_KERNELS_HPP_ #include #include -#include +#include namespace gko { namespace kernels { -#define GKO_DECLARE_SPARSITY_SPMV_KERNEL(ValueType, IndexType) \ - void spmv(std::shared_ptr exec, \ - const matrix::Sparsity *a, \ +#define GKO_DECLARE_SPARSITY_CSR_SPMV_KERNEL(ValueType, IndexType) \ + void spmv(std::shared_ptr exec, \ + const matrix::SparsityCsr *a, \ const matrix::Dense *b, matrix::Dense *c) -#define GKO_DECLARE_SPARSITY_ADVANCED_SPMV_KERNEL(ValueType, IndexType) \ - void advanced_spmv(std::shared_ptr exec, \ - const matrix::Dense *alpha, \ - const matrix::Sparsity *a, \ - const matrix::Dense *b, \ - const matrix::Dense *beta, \ +#define GKO_DECLARE_SPARSITY_CSR_ADVANCED_SPMV_KERNEL(ValueType, IndexType) \ + void advanced_spmv(std::shared_ptr exec, \ + const matrix::Dense *alpha, \ + const matrix::SparsityCsr *a, \ + const matrix::Dense *b, \ + const matrix::Dense *beta, \ matrix::Dense *c) -#define GKO_DECLARE_SPARSITY_REMOVE_DIAGONAL_ELEMENTS_KERNEL(ValueType, \ - IndexType) \ - void remove_diagonal_elements( \ - std::shared_ptr exec, \ - matrix::Sparsity *matrix, \ +#define GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL(ValueType, \ + IndexType) \ + void remove_diagonal_elements( \ + std::shared_ptr exec, \ + matrix::SparsityCsr *matrix, \ const IndexType *row_ptrs, const IndexType *col_idxs) -#define GKO_DECLARE_SPARSITY_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL(ValueType, \ - IndexType) \ - void count_num_diagonal_elements( \ - std::shared_ptr exec, \ - const matrix::Sparsity *matrix, \ +#define GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL(ValueType, \ + IndexType) \ + void count_num_diagonal_elements( \ + std::shared_ptr exec, \ + const matrix::SparsityCsr *matrix, \ size_type &num_diagonal_elements) -#define GKO_DECLARE_SPARSITY_TRANSPOSE_KERNEL(ValueType, IndexType) \ - void transpose(std::shared_ptr exec, \ - matrix::Sparsity *trans, \ - const matrix::Sparsity *orig) - -#define GKO_DECLARE_SPARSITY_SORT_BY_COLUMN_INDEX(ValueType, IndexType) \ - void sort_by_column_index(std::shared_ptr exec, \ - matrix::Sparsity *to_sort) - -#define GKO_DECLARE_SPARSITY_IS_SORTED_BY_COLUMN_INDEX(ValueType, IndexType) \ - void is_sorted_by_column_index( \ - std::shared_ptr exec, \ - const matrix::Sparsity *to_check, \ +#define GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL(ValueType, IndexType) \ + void transpose(std::shared_ptr exec, \ + matrix::SparsityCsr *trans, \ + const matrix::SparsityCsr *orig) + +#define GKO_DECLARE_SPARSITY_CSR_SORT_BY_COLUMN_INDEX(ValueType, IndexType) \ + void sort_by_column_index( \ + std::shared_ptr exec, \ + matrix::SparsityCsr *to_sort) + +#define GKO_DECLARE_SPARSITY_CSR_IS_SORTED_BY_COLUMN_INDEX(ValueType, \ + IndexType) \ + void is_sorted_by_column_index( \ + std::shared_ptr exec, \ + const matrix::SparsityCsr *to_check, \ bool *is_sorted) -#define GKO_DECLARE_ALL_AS_TEMPLATES \ - template \ - GKO_DECLARE_SPARSITY_SPMV_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_SPARSITY_ADVANCED_SPMV_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_SPARSITY_REMOVE_DIAGONAL_ELEMENTS_KERNEL(ValueType, \ - IndexType); \ - template \ - GKO_DECLARE_SPARSITY_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL(ValueType, \ - IndexType); \ - template \ - GKO_DECLARE_SPARSITY_TRANSPOSE_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_SPARSITY_SORT_BY_COLUMN_INDEX(ValueType, IndexType); \ - template \ - GKO_DECLARE_SPARSITY_IS_SORTED_BY_COLUMN_INDEX(ValueType, IndexType) +#define GKO_DECLARE_ALL_AS_TEMPLATES \ + template \ + GKO_DECLARE_SPARSITY_CSR_SPMV_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_SPARSITY_CSR_ADVANCED_SPMV_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL(ValueType, \ + IndexType); \ + template \ + GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL(ValueType, \ + IndexType); \ + template \ + GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_SPARSITY_CSR_SORT_BY_COLUMN_INDEX(ValueType, IndexType); \ + template \ + GKO_DECLARE_SPARSITY_CSR_IS_SORTED_BY_COLUMN_INDEX(ValueType, IndexType) namespace omp { -namespace sparsity { +namespace sparsity_csr { GKO_DECLARE_ALL_AS_TEMPLATES; -} // namespace sparsity +} // namespace sparsity_csr } // namespace omp namespace cuda { -namespace sparsity { +namespace sparsity_csr { GKO_DECLARE_ALL_AS_TEMPLATES; -} // namespace sparsity +} // namespace sparsity_csr } // namespace cuda namespace reference { -namespace sparsity { +namespace sparsity_csr { GKO_DECLARE_ALL_AS_TEMPLATES; -} // namespace sparsity +} // namespace sparsity_csr } // namespace reference @@ -138,4 +140,4 @@ GKO_DECLARE_ALL_AS_TEMPLATES; } // namespace gko -#endif // GKO_CORE_MATRIX_SPARSITY_KERNELS_HPP_ +#endif // GKO_CORE_MATRIX_SPARSITY_CSR_KERNELS_HPP_ diff --git a/core/test/matrix/CMakeLists.txt b/core/test/matrix/CMakeLists.txt index d62b5c05eba..ba370edb60f 100644 --- a/core/test/matrix/CMakeLists.txt +++ b/core/test/matrix/CMakeLists.txt @@ -5,4 +5,4 @@ ginkgo_create_test(ell) ginkgo_create_test(hybrid) ginkgo_create_test(identity) ginkgo_create_test(sellp) -ginkgo_create_test(sparsity) +ginkgo_create_test(sparsity_csr) diff --git a/core/test/matrix/sparsity.cpp b/core/test/matrix/sparsity_csr.cpp similarity index 86% rename from core/test/matrix/sparsity.cpp rename to core/test/matrix/sparsity_csr.cpp index 641110793b1..6d11271e2d3 100644 --- a/core/test/matrix/sparsity.cpp +++ b/core/test/matrix/sparsity_csr.cpp @@ -30,7 +30,7 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include +#include #include @@ -46,13 +46,13 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. namespace { -class Sparsity : public ::testing::Test { +class SparsityCsr : public ::testing::Test { protected: - using Mtx = gko::matrix::Sparsity<>; + using Mtx = gko::matrix::SparsityCsr<>; - Sparsity() + SparsityCsr() : exec(gko::ReferenceExecutor::create()), - mtx(gko::matrix::Sparsity<>::create(exec, gko::dim<2>{2, 3}, 4)) + mtx(gko::matrix::SparsityCsr<>::create(exec, gko::dim<2>{2, 3}, 4)) { Mtx::index_type *c = mtx->get_col_idxs(); Mtx::index_type *r = mtx->get_row_ptrs(); @@ -99,20 +99,20 @@ class Sparsity : public ::testing::Test { }; -TEST_F(Sparsity, KnowsItsSize) +TEST_F(SparsityCsr, KnowsItsSize) { ASSERT_EQ(mtx->get_size(), gko::dim<2>(2, 3)); ASSERT_EQ(mtx->get_num_nonzeros(), 4); } -TEST_F(Sparsity, ContainsCorrectData) +TEST_F(SparsityCsr, ContainsCorrectData) { assert_equal_to_original_mtx(mtx.get()); } -TEST_F(Sparsity, CanBeEmpty) +TEST_F(SparsityCsr, CanBeEmpty) { auto mtx = Mtx::create(exec); @@ -120,22 +120,22 @@ TEST_F(Sparsity, CanBeEmpty) } -TEST_F(Sparsity, SetsCorrectDefaultValue) +TEST_F(SparsityCsr, SetsCorrectDefaultValue) { - auto mtx = gko::matrix::Sparsity<>::create(exec, gko::dim<2>{3, 2}, - static_cast(0)); + auto mtx = gko::matrix::SparsityCsr<>::create( + exec, gko::dim<2>{3, 2}, static_cast(0)); ASSERT_EQ(mtx->get_const_value()[0], 1.0); ASSERT_EQ(mtx->get_value()[0], 1.0); } -TEST_F(Sparsity, CanBeCreatedFromExistingData) +TEST_F(SparsityCsr, CanBeCreatedFromExistingData) { gko::int32 col_idxs[] = {0, 1, 1, 0}; gko::int32 row_ptrs[] = {0, 2, 3, 4}; - auto mtx = gko::matrix::Sparsity<>::create( + auto mtx = gko::matrix::SparsityCsr<>::create( exec, gko::dim<2>{3, 2}, gko::Array::view(exec, 4, col_idxs), gko::Array::view(exec, 4, row_ptrs), 2.0); @@ -149,7 +149,7 @@ TEST_F(Sparsity, CanBeCreatedFromExistingData) } -TEST_F(Sparsity, CanBeCopied) +TEST_F(SparsityCsr, CanBeCopied) { auto copy = Mtx::create(exec); @@ -160,7 +160,7 @@ TEST_F(Sparsity, CanBeCopied) } -TEST_F(Sparsity, CanBeMoved) +TEST_F(SparsityCsr, CanBeMoved) { auto copy = Mtx::create(exec); @@ -170,7 +170,7 @@ TEST_F(Sparsity, CanBeMoved) } -TEST_F(Sparsity, CanBeCloned) +TEST_F(SparsityCsr, CanBeCloned) { auto clone = mtx->clone(); @@ -179,7 +179,7 @@ TEST_F(Sparsity, CanBeCloned) } -TEST_F(Sparsity, CanBeCleared) +TEST_F(SparsityCsr, CanBeCleared) { mtx->clear(); @@ -187,7 +187,7 @@ TEST_F(Sparsity, CanBeCleared) } -TEST_F(Sparsity, CanBeReadFromMatrixData) +TEST_F(SparsityCsr, CanBeReadFromMatrixData) { auto m = Mtx::create(exec); @@ -203,7 +203,7 @@ TEST_F(Sparsity, CanBeReadFromMatrixData) } -TEST_F(Sparsity, GeneratesCorrectMatrixData) +TEST_F(SparsityCsr, GeneratesCorrectMatrixData) { using tpl = gko::matrix_data<>::nonzero_type; gko::matrix_data<> data; diff --git a/cuda/CMakeLists.txt b/cuda/CMakeLists.txt index f44099c2294..d44728b81d4 100644 --- a/cuda/CMakeLists.txt +++ b/cuda/CMakeLists.txt @@ -37,7 +37,7 @@ target_sources(ginkgo_cuda matrix/ell_kernels.cu matrix/hybrid_kernels.cu matrix/sellp_kernels.cu - matrix/sparsity_kernels.cu + matrix/sparsity_csr_kernels.cu preconditioner/jacobi_advanced_apply_kernel.cu preconditioner/jacobi_generate_kernel.cu preconditioner/jacobi_kernels.cu diff --git a/cuda/matrix/dense_kernels.cu b/cuda/matrix/dense_kernels.cu index c66cf6ae427..0b0ef1e80b3 100644 --- a/cuda/matrix/dense_kernels.cu +++ b/cuda/matrix/dense_kernels.cu @@ -39,7 +39,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include +#include #include "cuda/base/cublas_bindings.hpp" @@ -733,12 +733,12 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( template void convert_to_sparsity(std::shared_ptr exec, - matrix::Sparsity *result, + matrix::SparsityCsr *result, const matrix::Dense *source) GKO_NOT_IMPLEMENTED; GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL); + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL); template diff --git a/cuda/matrix/sparsity_kernels.cu b/cuda/matrix/sparsity_csr_kernels.cu similarity index 79% rename from cuda/matrix/sparsity_kernels.cu rename to cuda/matrix/sparsity_csr_kernels.cu index 26792ea967d..4c5ded1af43 100644 --- a/cuda/matrix/sparsity_kernels.cu +++ b/cuda/matrix/sparsity_csr_kernels.cu @@ -30,7 +30,7 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include "core/matrix/sparsity_kernels.hpp" +#include "core/matrix/sparsity_csr_kernels.hpp" #include @@ -63,81 +63,82 @@ namespace cuda { * * @ingroup sparsity */ -namespace sparsity { +namespace sparsity_csr { template void spmv(std::shared_ptr exec, - const matrix::Sparsity *a, + const matrix::SparsityCsr *a, const matrix::Dense *b, matrix::Dense *c) GKO_NOT_IMPLEMENTED; -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SPARSITY_SPMV_KERNEL); +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_SPMV_KERNEL); template void advanced_spmv(std::shared_ptr exec, const matrix::Dense *alpha, - const matrix::Sparsity *a, + const matrix::SparsityCsr *a, const matrix::Dense *b, const matrix::Dense *beta, matrix::Dense *c) GKO_NOT_IMPLEMENTED; GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_ADVANCED_SPMV_KERNEL); + GKO_DECLARE_SPARSITY_CSR_ADVANCED_SPMV_KERNEL); template void count_num_diagonal_elements( std::shared_ptr exec, - const matrix::Sparsity *matrix, + const matrix::SparsityCsr *matrix, size_type &num_diagonal_elements) GKO_NOT_IMPLEMENTED; GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); + GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); template void remove_diagonal_elements(std::shared_ptr exec, - matrix::Sparsity *matrix, + matrix::SparsityCsr *matrix, const IndexType *row_ptrs, const IndexType *col_idxs) GKO_NOT_IMPLEMENTED; GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_REMOVE_DIAGONAL_ELEMENTS_KERNEL); + GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL); template void transpose(std::shared_ptr exec, - matrix::Sparsity *trans, - const matrix::Sparsity *orig) + matrix::SparsityCsr *trans, + const matrix::SparsityCsr *orig) GKO_NOT_IMPLEMENTED; GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_TRANSPOSE_KERNEL); + GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL); template void sort_by_column_index(std::shared_ptr exec, - matrix::Sparsity *to_sort) + matrix::SparsityCsr *to_sort) GKO_NOT_IMPLEMENTED; GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_SORT_BY_COLUMN_INDEX); + GKO_DECLARE_SPARSITY_CSR_SORT_BY_COLUMN_INDEX); template void is_sorted_by_column_index( std::shared_ptr exec, - const matrix::Sparsity *to_check, + const matrix::SparsityCsr *to_check, bool *is_sorted) GKO_NOT_IMPLEMENTED; GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_IS_SORTED_BY_COLUMN_INDEX); + GKO_DECLARE_SPARSITY_CSR_IS_SORTED_BY_COLUMN_INDEX); -} // namespace sparsity +} // namespace sparsity_csr } // namespace cuda } // namespace kernels } // namespace gko diff --git a/cuda/test/matrix/csr_kernels.cpp b/cuda/test/matrix/csr_kernels.cpp index 54fd5803f57..fd08a070cdb 100644 --- a/cuda/test/matrix/csr_kernels.cpp +++ b/cuda/test/matrix/csr_kernels.cpp @@ -47,7 +47,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include +#include #include "core/test/utils.hpp" @@ -389,11 +389,11 @@ TEST_F(Csr, MoveToEllIsEquivalentToRef) GKO_ASSERT_MTX_NEAR(ell_mtx.get(), dell_mtx.get(), 1e-14); } -TEST_F(Csr, ConvertToSparsityIsEquivalentToRef) +TEST_F(Csr, ConvertToSparsityCsrIsEquivalentToRef) { set_up_apply_data(std::make_shared()); - auto sparsity_mtx = gko::matrix::Sparsity<>::create(ref); - auto d_sparsity_mtx = gko::matrix::Sparsity<>::create(cuda); + auto sparsity_mtx = gko::matrix::SparsityCsr<>::create(ref); + auto d_sparsity_mtx = gko::matrix::SparsityCsr<>::create(cuda); mtx->convert_to(sparsity_mtx.get()); dmtx->convert_to(d_sparsity_mtx.get()); @@ -402,11 +402,11 @@ TEST_F(Csr, ConvertToSparsityIsEquivalentToRef) } -TEST_F(Csr, MoveToSparsityIsEquivalentToRef) +TEST_F(Csr, MoveToSparsityCsrIsEquivalentToRef) { set_up_apply_data(std::make_shared()); - auto sparsity_mtx = gko::matrix::Sparsity<>::create(ref); - auto d_sparsity_mtx = gko::matrix::Sparsity<>::create(cuda); + auto sparsity_mtx = gko::matrix::SparsityCsr<>::create(ref); + auto d_sparsity_mtx = gko::matrix::SparsityCsr<>::create(cuda); mtx->move_to(sparsity_mtx.get()); dmtx->move_to(d_sparsity_mtx.get()); diff --git a/include/ginkgo/core/matrix/csr.hpp b/include/ginkgo/core/matrix/csr.hpp index 3c09cbe1cdc..391308ffd1b 100644 --- a/include/ginkgo/core/matrix/csr.hpp +++ b/include/ginkgo/core/matrix/csr.hpp @@ -58,7 +58,7 @@ template class Sellp; template -class Sparsity; +class SparsityCsr; /** @@ -85,7 +85,7 @@ class Csr : public EnableLinOp>, public ConvertibleTo>, public ConvertibleTo>, public ConvertibleTo>, - public ConvertibleTo>, + public ConvertibleTo>, public ReadableFromMatrixData, public WritableToMatrixData, public Transposable { @@ -96,7 +96,7 @@ class Csr : public EnableLinOp>, friend class Ell; friend class Hybrid; friend class Sellp; - friend class Sparsity; + friend class SparsityCsr; public: using EnableLinOp::convert_to; @@ -304,9 +304,9 @@ class Csr : public EnableLinOp>, void move_to(Sellp *result) override; - void convert_to(Sparsity *result) const override; + void convert_to(SparsityCsr *result) const override; - void move_to(Sparsity *result) override; + void move_to(SparsityCsr *result) override; void read(const mat_data &data) override; diff --git a/include/ginkgo/core/matrix/dense.hpp b/include/ginkgo/core/matrix/dense.hpp index ce0f835ba93..ab0ae4fc9dd 100644 --- a/include/ginkgo/core/matrix/dense.hpp +++ b/include/ginkgo/core/matrix/dense.hpp @@ -66,7 +66,7 @@ template class Sellp; template -class Sparsity; +class SparsityCsr; /** @@ -97,8 +97,8 @@ class Dense : public EnableLinOp>, public ConvertibleTo>, public ConvertibleTo>, public ConvertibleTo>, - public ConvertibleTo>, - public ConvertibleTo>, + public ConvertibleTo>, + public ConvertibleTo>, public ReadableFromMatrixData, public ReadableFromMatrixData, public WritableToMatrixData, @@ -116,8 +116,8 @@ class Dense : public EnableLinOp>, friend class Hybrid; friend class Sellp; friend class Sellp; - friend class Sparsity; - friend class Sparsity; + friend class SparsityCsr; + friend class SparsityCsr; public: using EnableLinOp::convert_to; @@ -186,13 +186,13 @@ class Dense : public EnableLinOp>, void move_to(Sellp *result) override; - void convert_to(Sparsity *result) const override; + void convert_to(SparsityCsr *result) const override; - void move_to(Sparsity *result) override; + void move_to(SparsityCsr *result) override; - void convert_to(Sparsity *result) const override; + void convert_to(SparsityCsr *result) const override; - void move_to(Sparsity *result) override; + void move_to(SparsityCsr *result) override; void read(const mat_data &data) override; diff --git a/include/ginkgo/core/matrix/sparsity.hpp b/include/ginkgo/core/matrix/sparsity_csr.hpp similarity index 80% rename from include/ginkgo/core/matrix/sparsity.hpp rename to include/ginkgo/core/matrix/sparsity_csr.hpp index eac37f99052..1e43ec0930f 100644 --- a/include/ginkgo/core/matrix/sparsity.hpp +++ b/include/ginkgo/core/matrix/sparsity_csr.hpp @@ -30,8 +30,8 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#ifndef GKO_CORE_MATRIX_SPARSITY_HPP_ -#define GKO_CORE_MATRIX_SPARSITY_HPP_ +#ifndef GKO_CORE_MATRIX_SPARSITY_CSR_HPP_ +#define GKO_CORE_MATRIX_SPARSITY_CSR_HPP_ #include @@ -50,7 +50,7 @@ class Csr; /** - * Sparsity is a matrix format which stores only the sparsity pattern of a + * SparsityCsr is a matrix format which stores only the sparsity pattern of a * sparse matrix by compressing each row of the matrix (compressed sparse row * format). * @@ -68,18 +68,19 @@ class Csr; * @ingroup LinOp */ template -class Sparsity : public EnableLinOp>, - public EnableCreateMethod>, - public ReadableFromMatrixData, - public WritableToMatrixData, - public Transposable { - friend class EnableCreateMethod; - friend class EnablePolymorphicObject; +class SparsityCsr + : public EnableLinOp>, + public EnableCreateMethod>, + public ReadableFromMatrixData, + public WritableToMatrixData, + public Transposable { + friend class EnableCreateMethod; + friend class EnablePolymorphicObject; friend class Csr; public: - using EnableLinOp::convert_to; - using EnableLinOp::move_to; + using EnableLinOp::convert_to; + using EnableLinOp::move_to; using value_type = ValueType; using index_type = IndexType; @@ -96,14 +97,14 @@ class Sparsity : public EnableLinOp>, /** * Transforms the sparsity matrix to an adjacency matrix. As the adjacency - * matrix has to be square, the input Sparsity matrix for this function to - * work has to be square. + * matrix has to be square, the input SparsityCsr matrix for this function + * to work has to be square. * * @note The adjacency matrix in this case is the sparsity pattern but with * the diagonal ones removed. This is mainly used for the * reordering/partitioning as taken in by graph libraries such as METIS. */ - std::unique_ptr to_adjacency_matrix() const; + std::unique_ptr to_adjacency_matrix() const; /** * Sorts each row by column index @@ -125,7 +126,7 @@ class Sparsity : public EnableLinOp>, index_type *get_col_idxs() noexcept { return col_idxs_.get_data(); } /** - * @copydoc Sparsity::get_col_idxs() + * @copydoc SparsityCsr::get_col_idxs() * * @note This is the constant version of the function, which can be * significantly more memory efficient than the non-constant version, @@ -144,7 +145,7 @@ class Sparsity : public EnableLinOp>, index_type *get_row_ptrs() noexcept { return row_ptrs_.get_data(); } /** - * @copydoc Sparsity::get_row_ptrs() + * @copydoc SparsityCsr::get_row_ptrs() * * @note This is the constant version of the function, which can be * significantly more memory efficient than the non-constant version, @@ -163,7 +164,7 @@ class Sparsity : public EnableLinOp>, value_type *get_value() noexcept { return value_.get_data(); } /** - * @copydoc Sparsity::get_value() + * @copydoc SparsityCsr::get_value() * * @note This is the constant version of the function, which can be * significantly more memory efficient than the non-constant version, @@ -187,15 +188,15 @@ class Sparsity : public EnableLinOp>, protected: /** - * Creates an uninitialized Sparsity matrix of the specified size. + * Creates an uninitialized SparsityCsr matrix of the specified size. * * @param exec Executor associated to the matrix * @param size size of the matrix * @param num_nonzeros number of nonzeros */ - Sparsity(std::shared_ptr exec, - const dim<2> &size = dim<2>{}, size_type num_nonzeros = {}) - : EnableLinOp(exec, size), + SparsityCsr(std::shared_ptr exec, + const dim<2> &size = dim<2>{}, size_type num_nonzeros = {}) + : EnableLinOp(exec, size), col_idxs_(exec, num_nonzeros), // avoid allocation for empty matrix row_ptrs_(exec, size[0] + (size[0] > 0)) @@ -210,7 +211,7 @@ class Sparsity : public EnableLinOp>, } /** - * Creates a Sparsity matrix from already allocated (and initialized) row + * Creates a SparsityCsr matrix from already allocated (and initialized) row * pointer and column index arrays. * * @tparam ColIdxsArray type of `col_idxs` array @@ -229,10 +230,10 @@ class Sparsity : public EnableLinOp>, * matrix. */ template - Sparsity(std::shared_ptr exec, const dim<2> &size, - ColIdxsArray &&col_idxs, RowPtrsArray &&row_ptrs, - value_type value = one()) - : EnableLinOp(exec, size), + SparsityCsr(std::shared_ptr exec, const dim<2> &size, + ColIdxsArray &&col_idxs, RowPtrsArray &&row_ptrs, + value_type value = one()) + : EnableLinOp(exec, size), col_idxs_{exec, std::forward(col_idxs)}, row_ptrs_{exec, std::forward(row_ptrs)} { @@ -245,17 +246,17 @@ class Sparsity : public EnableLinOp>, } /** - * Creates a SPARSITY matrix from an existing matrix. Uses the + * Creates a Sparsity matrix from an existing matrix. Uses the * `copy_and_convert_to` functionality. * * @param exec Executor associated to the matrix * @param matrix The input matrix */ - Sparsity(std::shared_ptr exec, - std::shared_ptr matrix) - : EnableLinOp(exec, matrix->get_size()) + SparsityCsr(std::shared_ptr exec, + std::shared_ptr matrix) + : EnableLinOp(exec, matrix->get_size()) { - auto tmp_ = copy_and_convert_to(exec, matrix); + auto tmp_ = copy_and_convert_to(exec, matrix); this->copy_from(std::move(tmp_.get())); } @@ -275,4 +276,4 @@ class Sparsity : public EnableLinOp>, } // namespace gko -#endif // GKO_CORE_MATRIX_SPARSITY_HPP_ +#endif // GKO_CORE_MATRIX_SPARSITY_CSR_HPP_ diff --git a/include/ginkgo/ginkgo.hpp b/include/ginkgo/ginkgo.hpp index 8ded1aba34e..95a7dc3c734 100644 --- a/include/ginkgo/ginkgo.hpp +++ b/include/ginkgo/ginkgo.hpp @@ -74,7 +74,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include +#include #include #include diff --git a/omp/CMakeLists.txt b/omp/CMakeLists.txt index b771b6d7465..3f8705b2d2b 100644 --- a/omp/CMakeLists.txt +++ b/omp/CMakeLists.txt @@ -11,7 +11,7 @@ target_sources(ginkgo_omp matrix/ell_kernels.cpp matrix/hybrid_kernels.cpp matrix/sellp_kernels.cpp - matrix/sparsity_kernels.cpp + matrix/sparsity_csr_kernels.cpp preconditioner/jacobi_kernels.cpp solver/bicgstab_kernels.cpp solver/cg_kernels.cpp diff --git a/omp/matrix/dense_kernels.cpp b/omp/matrix/dense_kernels.cpp index c698f92fb40..2e2047604bf 100644 --- a/omp/matrix/dense_kernels.cpp +++ b/omp/matrix/dense_kernels.cpp @@ -46,7 +46,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include +#include namespace gko { @@ -458,7 +458,7 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( template void convert_to_sparsity(std::shared_ptr exec, - matrix::Sparsity *result, + matrix::SparsityCsr *result, const matrix::Dense *source) { auto num_rows = result->get_size()[0]; @@ -488,7 +488,7 @@ void convert_to_sparsity(std::shared_ptr exec, } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL); + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL); template diff --git a/omp/matrix/sparsity_kernels.cpp b/omp/matrix/sparsity_csr_kernels.cpp similarity index 86% rename from omp/matrix/sparsity_kernels.cpp rename to omp/matrix/sparsity_csr_kernels.cpp index bc781f6b429..3800faf5167 100644 --- a/omp/matrix/sparsity_kernels.cpp +++ b/omp/matrix/sparsity_csr_kernels.cpp @@ -30,7 +30,7 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include "core/matrix/sparsity_kernels.hpp" +#include "core/matrix/sparsity_csr_kernels.hpp" #include @@ -54,16 +54,16 @@ namespace gko { namespace kernels { namespace omp { /** - * @brief The Sparsity pattern format namespace. + * @brief The SparsityCsr pattern format namespace. * * @ingroup sparsity */ -namespace sparsity { +namespace sparsity_csr { template void spmv(std::shared_ptr exec, - const matrix::Sparsity *a, + const matrix::SparsityCsr *a, const matrix::Dense *b, matrix::Dense *c) { auto row_ptrs = a->get_const_row_ptrs(); @@ -85,13 +85,14 @@ void spmv(std::shared_ptr exec, } } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SPARSITY_SPMV_KERNEL); +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_SPMV_KERNEL); template void advanced_spmv(std::shared_ptr exec, const matrix::Dense *alpha, - const matrix::Sparsity *a, + const matrix::SparsityCsr *a, const matrix::Dense *b, const matrix::Dense *beta, matrix::Dense *c) @@ -118,13 +119,13 @@ void advanced_spmv(std::shared_ptr exec, } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_ADVANCED_SPMV_KERNEL); + GKO_DECLARE_SPARSITY_CSR_ADVANCED_SPMV_KERNEL); template void count_num_diagonal_elements( std::shared_ptr exec, - const matrix::Sparsity *matrix, + const matrix::SparsityCsr *matrix, size_type &num_diagonal_elements) { auto num_rows = matrix->get_size()[0]; @@ -142,12 +143,12 @@ void count_num_diagonal_elements( } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); + GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); template void remove_diagonal_elements(std::shared_ptr exec, - matrix::Sparsity *matrix, + matrix::SparsityCsr *matrix, const IndexType *row_ptrs, const IndexType *col_idxs) { @@ -179,7 +180,7 @@ void remove_diagonal_elements(std::shared_ptr exec, } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_REMOVE_DIAGONAL_ELEMENTS_KERNEL); + GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL); template @@ -198,9 +199,10 @@ inline void convert_sparsity_to_csc(size_type num_rows, template -void transpose_and_transform(std::shared_ptr exec, - matrix::Sparsity *trans, - const matrix::Sparsity *orig) +void transpose_and_transform( + std::shared_ptr exec, + matrix::SparsityCsr *trans, + const matrix::SparsityCsr *orig) { auto trans_row_ptrs = trans->get_row_ptrs(); auto orig_row_ptrs = orig->get_const_row_ptrs(); @@ -222,19 +224,19 @@ void transpose_and_transform(std::shared_ptr exec, template void transpose(std::shared_ptr exec, - matrix::Sparsity *trans, - const matrix::Sparsity *orig) + matrix::SparsityCsr *trans, + const matrix::SparsityCsr *orig) { transpose_and_transform(exec, trans, orig); } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_TRANSPOSE_KERNEL); + GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL); template void sort_by_column_index(std::shared_ptr exec, - matrix::Sparsity *to_sort) + matrix::SparsityCsr *to_sort) { auto row_ptrs = to_sort->get_row_ptrs(); auto col_idxs = to_sort->get_col_idxs(); @@ -248,13 +250,13 @@ void sort_by_column_index(std::shared_ptr exec, } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_SORT_BY_COLUMN_INDEX); + GKO_DECLARE_SPARSITY_CSR_SORT_BY_COLUMN_INDEX); template void is_sorted_by_column_index( std::shared_ptr exec, - const matrix::Sparsity *to_check, bool *is_sorted) + const matrix::SparsityCsr *to_check, bool *is_sorted) { const auto row_ptrs = to_check->get_const_row_ptrs(); const auto col_idxs = to_check->get_const_col_idxs(); @@ -277,10 +279,10 @@ void is_sorted_by_column_index( } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_IS_SORTED_BY_COLUMN_INDEX); + GKO_DECLARE_SPARSITY_CSR_IS_SORTED_BY_COLUMN_INDEX); -} // namespace sparsity +} // namespace sparsity_csr } // namespace omp } // namespace kernels } // namespace gko diff --git a/omp/test/matrix/CMakeLists.txt b/omp/test/matrix/CMakeLists.txt index 108cfc1f44d..b160723ab94 100644 --- a/omp/test/matrix/CMakeLists.txt +++ b/omp/test/matrix/CMakeLists.txt @@ -4,4 +4,4 @@ ginkgo_create_test(dense_kernels) ginkgo_create_test(ell_kernels) ginkgo_create_test(hybrid_kernels) ginkgo_create_test(sellp_kernels) -ginkgo_create_test(sparsity_kernels) +ginkgo_create_test(sparsity_csr_kernels) diff --git a/omp/test/matrix/csr_kernels.cpp b/omp/test/matrix/csr_kernels.cpp index 3f3491f30e4..72fa988ff5e 100644 --- a/omp/test/matrix/csr_kernels.cpp +++ b/omp/test/matrix/csr_kernels.cpp @@ -47,7 +47,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include +#include #include "core/test/utils.hpp" @@ -287,11 +287,11 @@ TEST_F(Csr, MoveToDenseIsEquivalentToRef) } -TEST_F(Csr, ConvertToSparsityIsEquivalentToRef) +TEST_F(Csr, ConvertToSparsityCsrIsEquivalentToRef) { set_up_apply_data(); - auto sparsity_mtx = gko::matrix::Sparsity<>::create(ref); - auto d_sparsity_mtx = gko::matrix::Sparsity<>::create(omp); + auto sparsity_mtx = gko::matrix::SparsityCsr<>::create(ref); + auto d_sparsity_mtx = gko::matrix::SparsityCsr<>::create(omp); mtx->convert_to(sparsity_mtx.get()); dmtx->convert_to(d_sparsity_mtx.get()); @@ -300,11 +300,11 @@ TEST_F(Csr, ConvertToSparsityIsEquivalentToRef) } -TEST_F(Csr, MoveToSparsityIsEquivalentToRef) +TEST_F(Csr, MoveToSparsityCsrIsEquivalentToRef) { set_up_apply_data(); - auto sparsity_mtx = gko::matrix::Sparsity<>::create(ref); - auto d_sparsity_mtx = gko::matrix::Sparsity<>::create(omp); + auto sparsity_mtx = gko::matrix::SparsityCsr<>::create(ref); + auto d_sparsity_mtx = gko::matrix::SparsityCsr<>::create(omp); mtx->move_to(sparsity_mtx.get()); dmtx->move_to(d_sparsity_mtx.get()); diff --git a/omp/test/matrix/sparsity_kernels.cpp b/omp/test/matrix/sparsity_csr_kernels.cpp similarity index 85% rename from omp/test/matrix/sparsity_kernels.cpp rename to omp/test/matrix/sparsity_csr_kernels.cpp index ae3cf29f660..4e2a189dbf7 100644 --- a/omp/test/matrix/sparsity_kernels.cpp +++ b/omp/test/matrix/sparsity_csr_kernels.cpp @@ -30,7 +30,7 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include "core/matrix/sparsity_kernels.hpp" +#include "core/matrix/sparsity_csr_kernels.hpp" #include @@ -44,7 +44,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include +#include #include "core/test/utils.hpp" @@ -53,14 +53,14 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. namespace { -class Sparsity : public ::testing::Test { +class SparsityCsr : public ::testing::Test { protected: - using Mtx = gko::matrix::Sparsity<>; + using Mtx = gko::matrix::SparsityCsr<>; using Vec = gko::matrix::Dense<>; using ComplexVec = gko::matrix::Dense>; - using ComplexMtx = gko::matrix::Sparsity>; + using ComplexMtx = gko::matrix::SparsityCsr>; - Sparsity() : mtx_size(532, 231), rand_engine(42) {} + SparsityCsr() : mtx_size(532, 231), rand_engine(42) {} void SetUp() { @@ -162,7 +162,7 @@ class Sparsity : public ::testing::Test { }; -TEST_F(Sparsity, SimpleApplyIsEquivalentToRef) +TEST_F(SparsityCsr, SimpleApplyIsEquivalentToRef) { set_up_apply_data(); @@ -173,7 +173,7 @@ TEST_F(Sparsity, SimpleApplyIsEquivalentToRef) } -TEST_F(Sparsity, AdvancedApplyIsEquivalentToRef) +TEST_F(SparsityCsr, AdvancedApplyIsEquivalentToRef) { set_up_apply_data(); @@ -184,7 +184,7 @@ TEST_F(Sparsity, AdvancedApplyIsEquivalentToRef) } -TEST_F(Sparsity, SimpleApplyToDenseMatrixIsEquivalentToRef) +TEST_F(SparsityCsr, SimpleApplyToDenseMatrixIsEquivalentToRef) { set_up_apply_data(3); @@ -195,7 +195,7 @@ TEST_F(Sparsity, SimpleApplyToDenseMatrixIsEquivalentToRef) } -TEST_F(Sparsity, AdvancedApplyToDenseMatrixIsEquivalentToRef) +TEST_F(SparsityCsr, AdvancedApplyToDenseMatrixIsEquivalentToRef) { set_up_apply_data(3); @@ -206,7 +206,7 @@ TEST_F(Sparsity, AdvancedApplyToDenseMatrixIsEquivalentToRef) } -TEST_F(Sparsity, TransposeIsEquivalentToRef) +TEST_F(SparsityCsr, TransposeIsEquivalentToRef) { set_up_apply_data(); @@ -218,34 +218,34 @@ TEST_F(Sparsity, TransposeIsEquivalentToRef) } -TEST_F(Sparsity, CountsNumberOfDiagElementsIsEqualToRef) +TEST_F(SparsityCsr, CountsNumberOfDiagElementsIsEqualToRef) { set_up_apply_data(); gko::size_type num_diags = 0; gko::size_type d_num_diags = 0; - gko::kernels::reference::sparsity::count_num_diagonal_elements( + gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( ref, mtx.get(), num_diags); - gko::kernels::omp::sparsity::count_num_diagonal_elements(omp, dmtx.get(), - d_num_diags); + gko::kernels::omp::sparsity_csr::count_num_diagonal_elements( + omp, dmtx.get(), d_num_diags); ASSERT_EQ(d_num_diags, num_diags); } -TEST_F(Sparsity, RemovesDiagElementsKernelIsEquivalentToRef) +TEST_F(SparsityCsr, RemovesDiagElementsKernelIsEquivalentToRef) { set_up_apply_data(); gko::size_type num_diags = 0; - gko::kernels::reference::sparsity::count_num_diagonal_elements( + gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( ref, mtx.get(), num_diags); auto tmp = Mtx::create(ref, mtx->get_size(), mtx->get_num_nonzeros() - num_diags); auto d_tmp = Mtx::create(omp, dmtx->get_size(), dmtx->get_num_nonzeros() - num_diags); - gko::kernels::reference::sparsity::remove_diagonal_elements( + gko::kernels::reference::sparsity_csr::remove_diagonal_elements( ref, tmp.get(), mtx->get_const_row_ptrs(), mtx->get_const_col_idxs()); - gko::kernels::omp::sparsity::remove_diagonal_elements( + gko::kernels::omp::sparsity_csr::remove_diagonal_elements( omp, d_tmp.get(), dmtx->get_const_row_ptrs(), dmtx->get_const_col_idxs()); @@ -253,7 +253,7 @@ TEST_F(Sparsity, RemovesDiagElementsKernelIsEquivalentToRef) } -TEST_F(Sparsity, RecognizeSortedMatrixIsEquivalentToRef) +TEST_F(SparsityCsr, RecognizeSortedMatrixIsEquivalentToRef) { set_up_apply_data(); bool is_sorted_omp{}; @@ -266,7 +266,7 @@ TEST_F(Sparsity, RecognizeSortedMatrixIsEquivalentToRef) } -TEST_F(Sparsity, RecognizeUnsortedMatrixIsEquivalentToRef) +TEST_F(SparsityCsr, RecognizeUnsortedMatrixIsEquivalentToRef) { auto uns_mtx = gen_unsorted_mtx(); bool is_sorted_omp{}; @@ -279,7 +279,7 @@ TEST_F(Sparsity, RecognizeUnsortedMatrixIsEquivalentToRef) } -TEST_F(Sparsity, SortSortedMatrixIsEquivalentToRef) +TEST_F(SparsityCsr, SortSortedMatrixIsEquivalentToRef) { set_up_apply_data(); @@ -291,7 +291,7 @@ TEST_F(Sparsity, SortSortedMatrixIsEquivalentToRef) } -TEST_F(Sparsity, SortUnsortedMatrixIsEquivalentToRef) +TEST_F(SparsityCsr, SortUnsortedMatrixIsEquivalentToRef) { auto uns_mtx = gen_unsorted_mtx(); diff --git a/reference/CMakeLists.txt b/reference/CMakeLists.txt index fdd3a1e9948..7516fc8641f 100644 --- a/reference/CMakeLists.txt +++ b/reference/CMakeLists.txt @@ -9,7 +9,7 @@ target_sources(ginkgo_reference matrix/ell_kernels.cpp matrix/hybrid_kernels.cpp matrix/sellp_kernels.cpp - matrix/sparsity_kernels.cpp + matrix/sparsity_csr_kernels.cpp preconditioner/jacobi_kernels.cpp solver/bicgstab_kernels.cpp solver/cg_kernels.cpp diff --git a/reference/matrix/dense_kernels.cpp b/reference/matrix/dense_kernels.cpp index 57f45e580e8..ca1c0921ad1 100644 --- a/reference/matrix/dense_kernels.cpp +++ b/reference/matrix/dense_kernels.cpp @@ -40,7 +40,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include +#include #include @@ -417,7 +417,7 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( template void convert_to_sparsity(std::shared_ptr exec, - matrix::Sparsity *result, + matrix::SparsityCsr *result, const matrix::Dense *source) { auto num_rows = result->get_size()[0]; @@ -442,7 +442,7 @@ void convert_to_sparsity(std::shared_ptr exec, } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_KERNEL); + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL); template diff --git a/reference/matrix/sparsity_kernels.cpp b/reference/matrix/sparsity_csr_kernels.cpp similarity index 85% rename from reference/matrix/sparsity_kernels.cpp rename to reference/matrix/sparsity_csr_kernels.cpp index aa8a0ccfdf2..99698d488ff 100644 --- a/reference/matrix/sparsity_kernels.cpp +++ b/reference/matrix/sparsity_csr_kernels.cpp @@ -30,7 +30,7 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include "core/matrix/sparsity_kernels.hpp" +#include "core/matrix/sparsity_csr_kernels.hpp" #include @@ -51,16 +51,16 @@ namespace gko { namespace kernels { namespace reference { /** - * @brief The Sparsity pattern format namespace. - * @ref Sparsity + * @brief The SparsityCsr pattern format namespace. + * @ref SparsityCsr * @ingroup sparsity */ -namespace sparsity { +namespace sparsity_csr { template void spmv(std::shared_ptr exec, - const matrix::Sparsity *a, + const matrix::SparsityCsr *a, const matrix::Dense *b, matrix::Dense *c) { auto row_ptrs = a->get_const_row_ptrs(); @@ -81,13 +81,14 @@ void spmv(std::shared_ptr exec, } } -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SPARSITY_SPMV_KERNEL); +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_SPMV_KERNEL); template void advanced_spmv(std::shared_ptr exec, const matrix::Dense *alpha, - const matrix::Sparsity *a, + const matrix::SparsityCsr *a, const matrix::Dense *b, const matrix::Dense *beta, matrix::Dense *c) @@ -113,13 +114,13 @@ void advanced_spmv(std::shared_ptr exec, } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_ADVANCED_SPMV_KERNEL); + GKO_DECLARE_SPARSITY_CSR_ADVANCED_SPMV_KERNEL); template void count_num_diagonal_elements( std::shared_ptr exec, - const matrix::Sparsity *matrix, + const matrix::SparsityCsr *matrix, size_type &num_diagonal_elements) { auto num_rows = matrix->get_size()[0]; @@ -137,12 +138,12 @@ void count_num_diagonal_elements( } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); + GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); template void remove_diagonal_elements(std::shared_ptr exec, - matrix::Sparsity *matrix, + matrix::SparsityCsr *matrix, const IndexType *row_ptrs, const IndexType *col_idxs) { @@ -173,7 +174,7 @@ void remove_diagonal_elements(std::shared_ptr exec, } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_REMOVE_DIAGONAL_ELEMENTS_KERNEL); + GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL); template @@ -192,9 +193,10 @@ inline void convert_sparsity_to_csc(size_type num_rows, template -void transpose_and_transform(std::shared_ptr exec, - matrix::Sparsity *trans, - const matrix::Sparsity *orig) +void transpose_and_transform( + std::shared_ptr exec, + matrix::SparsityCsr *trans, + const matrix::SparsityCsr *orig) { auto trans_row_ptrs = trans->get_row_ptrs(); auto orig_row_ptrs = orig->get_const_row_ptrs(); @@ -216,19 +218,19 @@ void transpose_and_transform(std::shared_ptr exec, template void transpose(std::shared_ptr exec, - matrix::Sparsity *trans, - const matrix::Sparsity *orig) + matrix::SparsityCsr *trans, + const matrix::SparsityCsr *orig) { transpose_and_transform(exec, trans, orig); } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_TRANSPOSE_KERNEL); + GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL); template void sort_by_column_index(std::shared_ptr exec, - matrix::Sparsity *to_sort) + matrix::SparsityCsr *to_sort) { auto row_ptrs = to_sort->get_row_ptrs(); auto col_idxs = to_sort->get_col_idxs(); @@ -241,13 +243,13 @@ void sort_by_column_index(std::shared_ptr exec, } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_SORT_BY_COLUMN_INDEX); + GKO_DECLARE_SPARSITY_CSR_SORT_BY_COLUMN_INDEX); template void is_sorted_by_column_index( std::shared_ptr exec, - const matrix::Sparsity *to_check, bool *is_sorted) + const matrix::SparsityCsr *to_check, bool *is_sorted) { const auto row_ptrs = to_check->get_const_row_ptrs(); const auto col_idxs = to_check->get_const_col_idxs(); @@ -265,10 +267,10 @@ void is_sorted_by_column_index( } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_IS_SORTED_BY_COLUMN_INDEX); + GKO_DECLARE_SPARSITY_CSR_IS_SORTED_BY_COLUMN_INDEX); -} // namespace sparsity +} // namespace sparsity_csr } // namespace reference } // namespace kernels } // namespace gko diff --git a/reference/test/matrix/CMakeLists.txt b/reference/test/matrix/CMakeLists.txt index d69c75090b1..7c0b5742eed 100644 --- a/reference/test/matrix/CMakeLists.txt +++ b/reference/test/matrix/CMakeLists.txt @@ -5,5 +5,5 @@ ginkgo_create_test(ell_kernels) ginkgo_create_test(hybrid_kernels) ginkgo_create_test(identity) ginkgo_create_test(sellp_kernels) -ginkgo_create_test(sparsity) -ginkgo_create_test(sparsity_kernels) +ginkgo_create_test(sparsity_csr) +ginkgo_create_test(sparsity_csr_kernels) diff --git a/reference/test/matrix/csr_kernels.cpp b/reference/test/matrix/csr_kernels.cpp index 264ee48d931..f775ea018ee 100644 --- a/reference/test/matrix/csr_kernels.cpp +++ b/reference/test/matrix/csr_kernels.cpp @@ -48,7 +48,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include +#include #include "core/test/utils/assertions.hpp" @@ -62,7 +62,7 @@ class Csr : public ::testing::Test { using Coo = gko::matrix::Coo<>; using Mtx = gko::matrix::Csr<>; using Sellp = gko::matrix::Sellp<>; - using Sparsity = gko::matrix::Sparsity<>; + using SparsityCsr = gko::matrix::SparsityCsr<>; using Ell = gko::matrix::Ell<>; using Hybrid = gko::matrix::Hybrid<>; using ComplexMtx = gko::matrix::Csr>; @@ -240,7 +240,7 @@ class Csr : public ::testing::Test { EXPECT_EQ(v[129], 0.0); } - void assert_equal_to_mtx(const Sparsity *m) + void assert_equal_to_mtx(const SparsityCsr *m) { auto *c = m->get_const_col_idxs(); auto *r = m->get_const_row_ptrs(); @@ -492,9 +492,9 @@ TEST_F(Csr, MovesToSellp) } -TEST_F(Csr, ConvertsToSparsity) +TEST_F(Csr, ConvertsToSparsityCsr) { - auto sparsity_mtx = gko::matrix::Sparsity<>::create(mtx->get_executor()); + auto sparsity_mtx = gko::matrix::SparsityCsr<>::create(mtx->get_executor()); mtx->convert_to(sparsity_mtx.get()); @@ -502,9 +502,9 @@ TEST_F(Csr, ConvertsToSparsity) } -TEST_F(Csr, MovesToSparsity) +TEST_F(Csr, MovesToSparsityCsr) { - auto sparsity_mtx = gko::matrix::Sparsity<>::create(mtx->get_executor()); + auto sparsity_mtx = gko::matrix::SparsityCsr<>::create(mtx->get_executor()); auto csr_ref = gko::matrix::Csr<>::create(mtx->get_executor()); csr_ref->copy_from(mtx.get()); diff --git a/reference/test/matrix/sparsity.cpp b/reference/test/matrix/sparsity_csr.cpp similarity index 92% rename from reference/test/matrix/sparsity.cpp rename to reference/test/matrix/sparsity_csr.cpp index 3188d63219b..c58d15d5337 100644 --- a/reference/test/matrix/sparsity.cpp +++ b/reference/test/matrix/sparsity_csr.cpp @@ -30,7 +30,7 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include +#include #include @@ -49,15 +49,15 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. namespace { -class Sparsity : public ::testing::Test { +class SparsityCsr : public ::testing::Test { protected: using v_type = double; using i_type = int; - using Mtx = gko::matrix::Sparsity; + using Mtx = gko::matrix::SparsityCsr; using Csr = gko::matrix::Csr; using DenseMtx = gko::matrix::Dense; - Sparsity() + SparsityCsr() : exec(gko::ReferenceExecutor::create()), mtx(Mtx::create(exec, gko::dim<2>{2, 3}, 4)) { @@ -77,7 +77,7 @@ class Sparsity : public ::testing::Test { }; -TEST_F(Sparsity, CanBeCreatedFromExistingCsrMatrix) +TEST_F(SparsityCsr, CanBeCreatedFromExistingCsrMatrix) { auto csr_mtx = gko::initialize( {{2.0, 3.0, 0.0}, {0.0, 1.0, 1.0}, {0.0, 0.0, -3.0}}, exec); @@ -90,7 +90,7 @@ TEST_F(Sparsity, CanBeCreatedFromExistingCsrMatrix) } -TEST_F(Sparsity, CanBeCreatedFromExistingDenseMatrix) +TEST_F(SparsityCsr, CanBeCreatedFromExistingDenseMatrix) { auto dense_mtx = gko::initialize( {{2.0, 3.0, 0.0}, {0.0, 1.0, 1.0}, {0.0, 0.0, -3.0}}, exec); diff --git a/reference/test/matrix/sparsity_kernels.cpp b/reference/test/matrix/sparsity_csr_kernels.cpp similarity index 81% rename from reference/test/matrix/sparsity_kernels.cpp rename to reference/test/matrix/sparsity_csr_kernels.cpp index ffa18d1968e..4af0f932880 100644 --- a/reference/test/matrix/sparsity_kernels.cpp +++ b/reference/test/matrix/sparsity_csr_kernels.cpp @@ -30,7 +30,7 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include "core/matrix/sparsity_kernels.hpp" +#include "core/matrix/sparsity_csr_kernels.hpp" #include @@ -43,7 +43,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include +#include #include "core/test/utils/assertions.hpp" @@ -52,12 +52,12 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. namespace { -class Sparsity : public ::testing::Test { +class SparsityCsr : public ::testing::Test { protected: - using Mtx = gko::matrix::Sparsity<>; + using Mtx = gko::matrix::SparsityCsr<>; using Vec = gko::matrix::Dense<>; - Sparsity() + SparsityCsr() : exec(gko::ReferenceExecutor::create()), mtx(Mtx::create(exec, gko::dim<2>{2, 3}, 4)), mtx2(Mtx::create(exec, gko::dim<2>{2, 3}, 5)), @@ -152,7 +152,7 @@ class Sparsity : public ::testing::Test { }; -TEST_F(Sparsity, AppliesToDenseVector) +TEST_F(SparsityCsr, AppliesToDenseVector) { auto x = gko::initialize({2.0, 1.0, 4.0}, exec); auto y = Vec::create(exec, gko::dim<2>{2, 1}); @@ -164,7 +164,7 @@ TEST_F(Sparsity, AppliesToDenseVector) } -TEST_F(Sparsity, AppliesToDenseMatrix) +TEST_F(SparsityCsr, AppliesToDenseMatrix) { auto x = gko::initialize({{2.0, 3.0}, {1.0, -1.5}, {4.0, 2.5}}, exec); auto y = Vec::create(exec, gko::dim<2>{2}); @@ -178,7 +178,7 @@ TEST_F(Sparsity, AppliesToDenseMatrix) } -TEST_F(Sparsity, AppliesLinearCombinationToDenseVector) +TEST_F(SparsityCsr, AppliesLinearCombinationToDenseVector) { auto alpha = gko::initialize({-1.0}, exec); auto beta = gko::initialize({2.0}, exec); @@ -192,7 +192,7 @@ TEST_F(Sparsity, AppliesLinearCombinationToDenseVector) } -TEST_F(Sparsity, AppliesLinearCombinationToDenseMatrix) +TEST_F(SparsityCsr, AppliesLinearCombinationToDenseMatrix) { auto alpha = gko::initialize({-1.0}, exec); auto beta = gko::initialize({2.0}, exec); @@ -208,7 +208,7 @@ TEST_F(Sparsity, AppliesLinearCombinationToDenseMatrix) } -TEST_F(Sparsity, ApplyFailsOnWrongInnerDimension) +TEST_F(SparsityCsr, ApplyFailsOnWrongInnerDimension) { auto x = Vec::create(exec, gko::dim<2>{2}); auto y = Vec::create(exec, gko::dim<2>{2}); @@ -217,7 +217,7 @@ TEST_F(Sparsity, ApplyFailsOnWrongInnerDimension) } -TEST_F(Sparsity, ApplyFailsOnWrongNumberOfRows) +TEST_F(SparsityCsr, ApplyFailsOnWrongNumberOfRows) { auto x = Vec::create(exec, gko::dim<2>{3, 2}); auto y = Vec::create(exec, gko::dim<2>{3, 2}); @@ -226,7 +226,7 @@ TEST_F(Sparsity, ApplyFailsOnWrongNumberOfRows) } -TEST_F(Sparsity, ApplyFailsOnWrongNumberOfCols) +TEST_F(SparsityCsr, ApplyFailsOnWrongNumberOfCols) { auto x = Vec::create(exec, gko::dim<2>{3}); auto y = Vec::create(exec, gko::dim<2>{2}); @@ -235,10 +235,10 @@ TEST_F(Sparsity, ApplyFailsOnWrongNumberOfCols) } -TEST_F(Sparsity, SquareMtxIsTransposable) +TEST_F(SparsityCsr, SquareMtxIsTransposable) { // clang-format off - auto mtx2 = gko::initialize>( + auto mtx2 = gko::initialize>( {{1.0, 1.0, 1.0}, {0.0, 1.0, 0.0}, {0.0, 1.0, 1.0}}, exec); @@ -246,7 +246,7 @@ TEST_F(Sparsity, SquareMtxIsTransposable) auto trans = mtx2->transpose(); auto trans_as_sparsity = - static_cast *>(trans.get()); + static_cast *>(trans.get()); // clang-format off GKO_ASSERT_MTX_NEAR(trans_as_sparsity, @@ -257,11 +257,11 @@ TEST_F(Sparsity, SquareMtxIsTransposable) } -TEST_F(Sparsity, NonSquareMtxIsTransposable) +TEST_F(SparsityCsr, NonSquareMtxIsTransposable) { auto trans = mtx->transpose(); auto trans_as_sparsity = - static_cast *>(trans.get()); + static_cast *>(trans.get()); // clang-format off GKO_ASSERT_MTX_NEAR(trans_as_sparsity, @@ -272,14 +272,14 @@ TEST_F(Sparsity, NonSquareMtxIsTransposable) } -TEST_F(Sparsity, CountsCorrectNumberOfDiagonalElements) +TEST_F(SparsityCsr, CountsCorrectNumberOfDiagonalElements) { // clang-format off - auto mtx2 = gko::initialize>( + auto mtx2 = gko::initialize>( {{1.0, 1.0, 1.0}, {0.0, 1.0, 0.0}, {0.0, 1.0, 1.0}}, exec); - auto mtx_s = gko::initialize>( + auto mtx_s = gko::initialize>( {{1.0, 1.0, 1.0}, {0.0, 0.0, 0.0}, {0.0, 1.0, 1.0}}, exec); @@ -287,9 +287,9 @@ TEST_F(Sparsity, CountsCorrectNumberOfDiagonalElements) gko::size_type m2_num_diags = 0; gko::size_type ms_num_diags = 0; - gko::kernels::reference::sparsity::count_num_diagonal_elements( + gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( exec, mtx2.get(), m2_num_diags); - gko::kernels::reference::sparsity::count_num_diagonal_elements( + gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( exec, mtx_s.get(), ms_num_diags); ASSERT_EQ(m2_num_diags, 3); @@ -297,23 +297,23 @@ TEST_F(Sparsity, CountsCorrectNumberOfDiagonalElements) } -TEST_F(Sparsity, RemovesDiagonalElementsForFullRankMatrix) +TEST_F(SparsityCsr, RemovesDiagonalElementsForFullRankMatrix) { // clang-format off - auto mtx2 = gko::initialize>( + auto mtx2 = gko::initialize>( {{1.0, 1.0, 1.0}, {0.0, 1.0, 0.0}, {0.0, 1.0, 1.0}}, exec); - auto mtx_s = gko::initialize>( + auto mtx_s = gko::initialize>( {{0.0, 1.0, 1.0}, {0.0, 0.0, 0.0}, {0.0, 1.0, 0.0}}, exec); // clang-format on - auto tmp_mtx = gko::matrix::Sparsity<>::create(exec, mtx_s->get_size(), - mtx_s->get_num_nonzeros()); + auto tmp_mtx = gko::matrix::SparsityCsr<>::create( + exec, mtx_s->get_size(), mtx_s->get_num_nonzeros()); tmp_mtx->copy_from(mtx2.get()); - gko::kernels::reference::sparsity::remove_diagonal_elements( + gko::kernels::reference::sparsity_csr::remove_diagonal_elements( exec, tmp_mtx.get(), mtx2->get_const_row_ptrs(), mtx2->get_const_col_idxs()); @@ -321,23 +321,23 @@ TEST_F(Sparsity, RemovesDiagonalElementsForFullRankMatrix) } -TEST_F(Sparsity, RemovesDiagonalElementsForIncompleteRankMatrix) +TEST_F(SparsityCsr, RemovesDiagonalElementsForIncompleteRankMatrix) { // clang-format off - auto mtx2 = gko::initialize>( + auto mtx2 = gko::initialize>( {{1.0, 1.0, 1.0}, {0.0, 0.0, 0.0}, {0.0, 1.0, 1.0}}, exec); - auto mtx_s = gko::initialize>( + auto mtx_s = gko::initialize>( {{0.0, 1.0, 1.0}, {0.0, 0.0, 0.0}, {0.0, 1.0, 0.0}}, exec); // clang-format on - auto tmp_mtx = gko::matrix::Sparsity<>::create(exec, mtx_s->get_size(), - mtx_s->get_num_nonzeros()); + auto tmp_mtx = gko::matrix::SparsityCsr<>::create( + exec, mtx_s->get_size(), mtx_s->get_num_nonzeros()); tmp_mtx->copy_from(mtx2.get()); - gko::kernels::reference::sparsity::remove_diagonal_elements( + gko::kernels::reference::sparsity_csr::remove_diagonal_elements( exec, tmp_mtx.get(), mtx2->get_const_row_ptrs(), mtx2->get_const_col_idxs()); @@ -345,14 +345,14 @@ TEST_F(Sparsity, RemovesDiagonalElementsForIncompleteRankMatrix) } -TEST_F(Sparsity, SquareMtxIsConvertibleToAdjacencyMatrix) +TEST_F(SparsityCsr, SquareMtxIsConvertibleToAdjacencyMatrix) { // clang-format off - auto mtx2 = gko::initialize>( + auto mtx2 = gko::initialize>( {{1.0, 1.0, 1.0}, {0.0, 1.0, 0.0}, {0.0, 1.0, 1.0}}, exec); - auto mtx_s = gko::initialize>( + auto mtx_s = gko::initialize>( {{0.0, 1.0, 1.0}, {0.0, 0.0, 0.0}, {0.0, 1.0, 0.0}}, exec); @@ -364,13 +364,13 @@ TEST_F(Sparsity, SquareMtxIsConvertibleToAdjacencyMatrix) } -TEST_F(Sparsity, NonSquareMtxIsNotConvertibleToAdjacencyMatrix) +TEST_F(SparsityCsr, NonSquareMtxIsNotConvertibleToAdjacencyMatrix) { ASSERT_THROW(mtx->to_adjacency_matrix(), gko::DimensionMismatch); } -TEST_F(Sparsity, RecognizeSortedMatrix) +TEST_F(SparsityCsr, RecognizeSortedMatrix) { ASSERT_TRUE(mtx->is_sorted_by_column_index()); ASSERT_TRUE(mtx2->is_sorted_by_column_index()); @@ -378,13 +378,13 @@ TEST_F(Sparsity, RecognizeSortedMatrix) } -TEST_F(Sparsity, RecognizeUnsortedMatrix) +TEST_F(SparsityCsr, RecognizeUnsortedMatrix) { ASSERT_FALSE(mtx3_unsorted->is_sorted_by_column_index()); } -TEST_F(Sparsity, SortSortedMatrix) +TEST_F(SparsityCsr, SortSortedMatrix) { auto matrix = mtx3_sorted->clone(); @@ -394,7 +394,7 @@ TEST_F(Sparsity, SortSortedMatrix) } -TEST_F(Sparsity, SortUnsortedMatrix) +TEST_F(SparsityCsr, SortUnsortedMatrix) { auto matrix = mtx3_unsorted->clone(); diff --git a/test_install/test_install.cpp b/test_install/test_install.cpp index 5115a7405c6..5205743cd7c 100644 --- a/test_install/test_install.cpp +++ b/test_install/test_install.cpp @@ -244,9 +244,9 @@ int main(int, char **) auto test = Mtx::create(refExec, gko::dim<2>{2, 2}, 2); } - // core/matrix/sparsity.hpp + // core/matrix/sparsity_csr.hpp { - using Mtx = gko::matrix::Sparsity<>; + using Mtx = gko::matrix::SparsityCsr<>; auto test = Mtx::create(refExec, gko::dim<2>{2, 2}); // core/preconditioner/ilu.hpp From 745eb45615e547e84b5b91b7954479cc267a312f Mon Sep 17 00:00:00 2001 From: Pratik Nayak Date: Tue, 8 Oct 2019 14:58:03 +0200 Subject: [PATCH 8/8] Review update + Update some forgotten names. + Improve some kernels. + Add some dense to sparsity csr conversion tests. --- core/matrix/csr.cpp | 6 ++- core/matrix/dense.cpp | 16 +++---- core/matrix/dense_kernels.hpp | 8 ++-- core/matrix/sparsity_csr.cpp | 2 +- core/matrix/sparsity_csr_kernels.hpp | 2 +- cuda/matrix/dense_kernels.cu | 6 +-- cuda/matrix/sparsity_csr_kernels.cu | 2 +- omp/matrix/dense_kernels.cpp | 6 +-- omp/matrix/sparsity_csr_kernels.cpp | 17 +++---- omp/test/matrix/dense_kernels.cpp | 31 ++++++++++++ omp/test/matrix/sparsity_csr_kernels.cpp | 7 +-- reference/matrix/dense_kernels.cpp | 6 +-- reference/matrix/sparsity_csr_kernels.cpp | 16 +++---- reference/test/matrix/dense_kernels.cpp | 47 +++++++++++++++++++ reference/test/matrix/sparsity_csr.cpp | 4 +- .../test/matrix/sparsity_csr_kernels.cpp | 4 +- 16 files changed, 128 insertions(+), 52 deletions(-) diff --git a/core/matrix/csr.cpp b/core/matrix/csr.cpp index d7e56eb11a5..53dd2f02b12 100644 --- a/core/matrix/csr.cpp +++ b/core/matrix/csr.cpp @@ -197,7 +197,11 @@ void Csr::convert_to( exec, this->get_size(), this->get_num_stored_elements()); tmp->col_idxs_ = this->col_idxs_; tmp->row_ptrs_ = this->row_ptrs_; - tmp->value_ = gko::Array(exec, {one()}); + if (result->value_.get_data()) { + tmp->value_ = result->value_; + } else { + tmp->value_ = gko::Array(exec, {one()}); + } tmp->move_to(result); } diff --git a/core/matrix/dense.cpp b/core/matrix/dense.cpp index 32e1c35fa02..72c23343857 100644 --- a/core/matrix/dense.cpp +++ b/core/matrix/dense.cpp @@ -76,7 +76,7 @@ GKO_REGISTER_OPERATION(convert_to_csr, dense::convert_to_csr); GKO_REGISTER_OPERATION(convert_to_ell, dense::convert_to_ell); GKO_REGISTER_OPERATION(convert_to_hybrid, dense::convert_to_hybrid); GKO_REGISTER_OPERATION(convert_to_sellp, dense::convert_to_sellp); -GKO_REGISTER_OPERATION(convert_to_sparsity, dense::convert_to_sparsity); +GKO_REGISTER_OPERATION(convert_to_sparsity_csr, dense::convert_to_sparsity_csr); } // namespace dense @@ -444,10 +444,9 @@ void Dense::move_to(Sellp *result) template void Dense::convert_to(SparsityCsr *result) const { - conversion_helper( - result, this, - dense::template make_convert_to_sparsity *&>); + conversion_helper(result, this, + dense::template make_convert_to_sparsity_csr< + decltype(result), const Dense *&>); } @@ -461,10 +460,9 @@ void Dense::move_to(SparsityCsr *result) template void Dense::convert_to(SparsityCsr *result) const { - conversion_helper( - result, this, - dense::template make_convert_to_sparsity *&>); + conversion_helper(result, this, + dense::template make_convert_to_sparsity_csr< + decltype(result), const Dense *&>); } diff --git a/core/matrix/dense_kernels.hpp b/core/matrix/dense_kernels.hpp index 5e6169d4732..4857fb81db9 100644 --- a/core/matrix/dense_kernels.hpp +++ b/core/matrix/dense_kernels.hpp @@ -98,10 +98,10 @@ namespace kernels { matrix::Sellp<_type, _prec> *other, \ const matrix::Dense<_type> *source) -#define GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL(_type, _prec) \ - void convert_to_sparsity(std::shared_ptr exec, \ - matrix::SparsityCsr<_type, _prec> *other, \ - const matrix::Dense<_type> *source) +#define GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL(_type, _prec) \ + void convert_to_sparsity_csr(std::shared_ptr exec, \ + matrix::SparsityCsr<_type, _prec> *other, \ + const matrix::Dense<_type> *source) #define GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL(_type) \ void count_nonzeros(std::shared_ptr exec, \ diff --git a/core/matrix/sparsity_csr.cpp b/core/matrix/sparsity_csr.cpp index 985fd03bc1e..3c7cfa1363f 100644 --- a/core/matrix/sparsity_csr.cpp +++ b/core/matrix/sparsity_csr.cpp @@ -171,7 +171,7 @@ SparsityCsr::to_adjacency_matrix() const GKO_ASSERT_IS_SQUARE_MATRIX(this); size_type num_diagonal_elements = 0; exec->run(sparsity_csr::make_count_num_diagonal_elements( - this, num_diagonal_elements)); + this, &num_diagonal_elements)); auto adj_mat = SparsityCsr::create(exec, this->get_size(), this->get_num_nonzeros() - num_diagonal_elements); diff --git a/core/matrix/sparsity_csr_kernels.hpp b/core/matrix/sparsity_csr_kernels.hpp index 1fd40ca06a3..f9af3dcdffa 100644 --- a/core/matrix/sparsity_csr_kernels.hpp +++ b/core/matrix/sparsity_csr_kernels.hpp @@ -68,7 +68,7 @@ namespace kernels { void count_num_diagonal_elements( \ std::shared_ptr exec, \ const matrix::SparsityCsr *matrix, \ - size_type &num_diagonal_elements) + size_type *num_diagonal_elements) #define GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL(ValueType, IndexType) \ void transpose(std::shared_ptr exec, \ diff --git a/cuda/matrix/dense_kernels.cu b/cuda/matrix/dense_kernels.cu index 0b0ef1e80b3..a74c431599a 100644 --- a/cuda/matrix/dense_kernels.cu +++ b/cuda/matrix/dense_kernels.cu @@ -732,9 +732,9 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( template -void convert_to_sparsity(std::shared_ptr exec, - matrix::SparsityCsr *result, - const matrix::Dense *source) +void convert_to_sparsity_csr(std::shared_ptr exec, + matrix::SparsityCsr *result, + const matrix::Dense *source) GKO_NOT_IMPLEMENTED; GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( diff --git a/cuda/matrix/sparsity_csr_kernels.cu b/cuda/matrix/sparsity_csr_kernels.cu index 4c5ded1af43..2bdb8372030 100644 --- a/cuda/matrix/sparsity_csr_kernels.cu +++ b/cuda/matrix/sparsity_csr_kernels.cu @@ -92,7 +92,7 @@ template void count_num_diagonal_elements( std::shared_ptr exec, const matrix::SparsityCsr *matrix, - size_type &num_diagonal_elements) GKO_NOT_IMPLEMENTED; + size_type *num_diagonal_elements) GKO_NOT_IMPLEMENTED; GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); diff --git a/omp/matrix/dense_kernels.cpp b/omp/matrix/dense_kernels.cpp index 2e2047604bf..5654b40d753 100644 --- a/omp/matrix/dense_kernels.cpp +++ b/omp/matrix/dense_kernels.cpp @@ -457,9 +457,9 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( template -void convert_to_sparsity(std::shared_ptr exec, - matrix::SparsityCsr *result, - const matrix::Dense *source) +void convert_to_sparsity_csr(std::shared_ptr exec, + matrix::SparsityCsr *result, + const matrix::Dense *source) { auto num_rows = result->get_size()[0]; auto num_cols = result->get_size()[1]; diff --git a/omp/matrix/sparsity_csr_kernels.cpp b/omp/matrix/sparsity_csr_kernels.cpp index 3800faf5167..830435fc376 100644 --- a/omp/matrix/sparsity_csr_kernels.cpp +++ b/omp/matrix/sparsity_csr_kernels.cpp @@ -68,6 +68,7 @@ void spmv(std::shared_ptr exec, { auto row_ptrs = a->get_const_row_ptrs(); auto col_idxs = a->get_const_col_idxs(); + auto val = a->get_const_value()[0]; #pragma omp parallel for for (size_type row = 0; row < a->get_size()[0]; ++row) { @@ -76,7 +77,6 @@ void spmv(std::shared_ptr exec, } for (size_type k = row_ptrs[row]; k < static_cast(row_ptrs[row + 1]); ++k) { - auto val = one(); auto col = col_idxs[k]; for (size_type j = 0; j < c->get_size()[1]; ++j) { c->at(row, j) += val * b->at(col, j); @@ -101,6 +101,7 @@ void advanced_spmv(std::shared_ptr exec, auto col_idxs = a->get_const_col_idxs(); auto valpha = alpha->at(0, 0); auto vbeta = beta->at(0, 0); + auto val = a->get_const_value()[0]; #pragma omp parallel for for (size_type row = 0; row < a->get_size()[0]; ++row) { @@ -109,7 +110,6 @@ void advanced_spmv(std::shared_ptr exec, } for (size_type k = row_ptrs[row]; k < static_cast(row_ptrs[row + 1]); ++k) { - auto val = one(); auto col = col_idxs[k]; for (size_type j = 0; j < c->get_size()[1]; ++j) { c->at(row, j) += valpha * val * b->at(col, j); @@ -126,7 +126,7 @@ template void count_num_diagonal_elements( std::shared_ptr exec, const matrix::SparsityCsr *matrix, - size_type &num_diagonal_elements) + size_type *num_diagonal_elements) { auto num_rows = matrix->get_size()[0]; auto row_ptrs = matrix->get_const_row_ptrs(); @@ -139,7 +139,7 @@ void count_num_diagonal_elements( } } } - num_diagonal_elements = num_diag; + *num_diagonal_elements = num_diag; } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( @@ -165,18 +165,15 @@ void remove_diagonal_elements(std::shared_ptr exec, } adj_ptrs[i + 1] = row_ptrs[i + 1] - num_diag; } - std::vector temp_idxs; + auto nnz = 0; for (auto i = 0; i < num_rows; ++i) { for (auto j = row_ptrs[i]; j < row_ptrs[i + 1]; ++j) { if (col_idxs[j] != i) { - temp_idxs.push_back(col_idxs[j]); + adj_idxs[nnz] = col_idxs[j]; + nnz++; } } } -#pragma omp parallel for - for (auto i = 0; i < temp_idxs.size(); ++i) { - adj_idxs[i] = temp_idxs[i]; - } } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( diff --git a/omp/test/matrix/dense_kernels.cpp b/omp/test/matrix/dense_kernels.cpp index b0f0b75170b..b1d601291f1 100644 --- a/omp/test/matrix/dense_kernels.cpp +++ b/omp/test/matrix/dense_kernels.cpp @@ -47,6 +47,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include "core/test/utils.hpp" @@ -363,6 +364,36 @@ TEST_F(Dense, MoveToCsrIsEquivalentToRef) } +TEST_F(Dense, ConvertToSparsityCsrIsEquivalentToRef) +{ + auto mtx = gen_mtx(532, 231); + auto dmtx = Mtx::create(omp); + dmtx->copy_from(mtx.get()); + auto sparsity_mtx = gko::matrix::SparsityCsr<>::create(ref); + auto d_sparsity_mtx = gko::matrix::SparsityCsr<>::create(omp); + + mtx->convert_to(sparsity_mtx.get()); + dmtx->convert_to(d_sparsity_mtx.get()); + + GKO_ASSERT_MTX_NEAR(d_sparsity_mtx.get(), sparsity_mtx.get(), 1e-14); +} + + +TEST_F(Dense, MoveToSparsityCsrIsEquivalentToRef) +{ + auto mtx = gen_mtx(532, 231); + auto dmtx = Mtx::create(omp); + dmtx->copy_from(mtx.get()); + auto sparsity_mtx = gko::matrix::SparsityCsr<>::create(ref); + auto d_sparsity_mtx = gko::matrix::SparsityCsr<>::create(omp); + + mtx->move_to(sparsity_mtx.get()); + dmtx->move_to(d_sparsity_mtx.get()); + + GKO_ASSERT_MTX_NEAR(d_sparsity_mtx.get(), sparsity_mtx.get(), 1e-14); +} + + TEST_F(Dense, ConvertToEllIsEquivalentToRef) { auto rmtx = gen_mtx(532, 231); diff --git a/omp/test/matrix/sparsity_csr_kernels.cpp b/omp/test/matrix/sparsity_csr_kernels.cpp index 4e2a189dbf7..91852dbcb53 100644 --- a/omp/test/matrix/sparsity_csr_kernels.cpp +++ b/omp/test/matrix/sparsity_csr_kernels.cpp @@ -225,9 +225,9 @@ TEST_F(SparsityCsr, CountsNumberOfDiagElementsIsEqualToRef) gko::size_type d_num_diags = 0; gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( - ref, mtx.get(), num_diags); + ref, mtx.get(), &num_diags); gko::kernels::omp::sparsity_csr::count_num_diagonal_elements( - omp, dmtx.get(), d_num_diags); + omp, dmtx.get(), &d_num_diags); ASSERT_EQ(d_num_diags, num_diags); } @@ -238,11 +238,12 @@ TEST_F(SparsityCsr, RemovesDiagElementsKernelIsEquivalentToRef) set_up_apply_data(); gko::size_type num_diags = 0; gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( - ref, mtx.get(), num_diags); + ref, mtx.get(), &num_diags); auto tmp = Mtx::create(ref, mtx->get_size(), mtx->get_num_nonzeros() - num_diags); auto d_tmp = Mtx::create(omp, dmtx->get_size(), dmtx->get_num_nonzeros() - num_diags); + gko::kernels::reference::sparsity_csr::remove_diagonal_elements( ref, tmp.get(), mtx->get_const_row_ptrs(), mtx->get_const_col_idxs()); gko::kernels::omp::sparsity_csr::remove_diagonal_elements( diff --git a/reference/matrix/dense_kernels.cpp b/reference/matrix/dense_kernels.cpp index ca1c0921ad1..82a20a8b1a4 100644 --- a/reference/matrix/dense_kernels.cpp +++ b/reference/matrix/dense_kernels.cpp @@ -416,9 +416,9 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( template -void convert_to_sparsity(std::shared_ptr exec, - matrix::SparsityCsr *result, - const matrix::Dense *source) +void convert_to_sparsity_csr(std::shared_ptr exec, + matrix::SparsityCsr *result, + const matrix::Dense *source) { auto num_rows = result->get_size()[0]; auto num_cols = result->get_size()[1]; diff --git a/reference/matrix/sparsity_csr_kernels.cpp b/reference/matrix/sparsity_csr_kernels.cpp index 99698d488ff..42b4edd88a2 100644 --- a/reference/matrix/sparsity_csr_kernels.cpp +++ b/reference/matrix/sparsity_csr_kernels.cpp @@ -65,6 +65,7 @@ void spmv(std::shared_ptr exec, { auto row_ptrs = a->get_const_row_ptrs(); auto col_idxs = a->get_const_col_idxs(); + auto val = a->get_const_value()[0]; for (size_type row = 0; row < a->get_size()[0]; ++row) { for (size_type j = 0; j < c->get_size()[1]; ++j) { @@ -72,7 +73,6 @@ void spmv(std::shared_ptr exec, } for (size_type k = row_ptrs[row]; k < static_cast(row_ptrs[row + 1]); ++k) { - auto val = one(); auto col = col_idxs[k]; for (size_type j = 0; j < c->get_size()[1]; ++j) { c->at(row, j) += val * b->at(col, j); @@ -97,6 +97,7 @@ void advanced_spmv(std::shared_ptr exec, auto col_idxs = a->get_const_col_idxs(); auto valpha = alpha->at(0, 0); auto vbeta = beta->at(0, 0); + auto val = a->get_const_value()[0]; for (size_type row = 0; row < a->get_size()[0]; ++row) { for (size_type j = 0; j < c->get_size()[1]; ++j) { @@ -104,7 +105,6 @@ void advanced_spmv(std::shared_ptr exec, } for (size_type k = row_ptrs[row]; k < static_cast(row_ptrs[row + 1]); ++k) { - auto val = one(); auto col = col_idxs[k]; for (size_type j = 0; j < c->get_size()[1]; ++j) { c->at(row, j) += valpha * val * b->at(col, j); @@ -121,7 +121,7 @@ template void count_num_diagonal_elements( std::shared_ptr exec, const matrix::SparsityCsr *matrix, - size_type &num_diagonal_elements) + size_type *num_diagonal_elements) { auto num_rows = matrix->get_size()[0]; auto row_ptrs = matrix->get_const_row_ptrs(); @@ -134,7 +134,7 @@ void count_num_diagonal_elements( } } } - num_diagonal_elements = num_diag; + *num_diagonal_elements = num_diag; } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( @@ -160,17 +160,15 @@ void remove_diagonal_elements(std::shared_ptr exec, } adj_ptrs[i + 1] = row_ptrs[i + 1] - num_diag; } - std::vector temp_idxs; + auto nnz = 0; for (auto i = 0; i < num_rows; ++i) { for (auto j = row_ptrs[i]; j < row_ptrs[i + 1]; ++j) { if (col_idxs[j] != i) { - temp_idxs.push_back(col_idxs[j]); + adj_idxs[nnz] = col_idxs[j]; + nnz++; } } } - for (auto i = 0; i < temp_idxs.size(); ++i) { - adj_idxs[i] = temp_idxs[i]; - } } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( diff --git a/reference/test/matrix/dense_kernels.cpp b/reference/test/matrix/dense_kernels.cpp index 892a557a973..fc30c4a89df 100644 --- a/reference/test/matrix/dense_kernels.cpp +++ b/reference/test/matrix/dense_kernels.cpp @@ -48,6 +48,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include "core/test/utils.hpp" @@ -380,6 +381,52 @@ TEST_F(Dense, MovesToCsr) } +TEST_F(Dense, ConvertsToSparsityCsr) +{ + auto sparsity_csr_mtx = + gko::matrix::SparsityCsr<>::create(mtx4->get_executor()); + + mtx4->convert_to(sparsity_csr_mtx.get()); + auto v = sparsity_csr_mtx->get_const_value(); + auto c = sparsity_csr_mtx->get_const_col_idxs(); + auto r = sparsity_csr_mtx->get_const_row_ptrs(); + + ASSERT_EQ(sparsity_csr_mtx->get_size(), gko::dim<2>(2, 3)); + ASSERT_EQ(sparsity_csr_mtx->get_num_nonzeros(), 4); + EXPECT_EQ(r[0], 0); + EXPECT_EQ(r[1], 3); + EXPECT_EQ(r[2], 4); + EXPECT_EQ(c[0], 0); + EXPECT_EQ(c[1], 1); + EXPECT_EQ(c[2], 2); + EXPECT_EQ(c[3], 1); + EXPECT_EQ(v[0], 1.0); +} + + +TEST_F(Dense, MovesToSparsityCsr) +{ + auto sparsity_csr_mtx = + gko::matrix::SparsityCsr<>::create(mtx4->get_executor()); + + mtx4->move_to(sparsity_csr_mtx.get()); + auto v = sparsity_csr_mtx->get_const_value(); + auto c = sparsity_csr_mtx->get_const_col_idxs(); + auto r = sparsity_csr_mtx->get_const_row_ptrs(); + + ASSERT_EQ(sparsity_csr_mtx->get_size(), gko::dim<2>(2, 3)); + ASSERT_EQ(sparsity_csr_mtx->get_num_nonzeros(), 4); + EXPECT_EQ(r[0], 0); + EXPECT_EQ(r[1], 3); + EXPECT_EQ(r[2], 4); + EXPECT_EQ(c[0], 0); + EXPECT_EQ(c[1], 1); + EXPECT_EQ(c[2], 2); + EXPECT_EQ(c[3], 1); + EXPECT_EQ(v[0], 1.0); +} + + TEST_F(Dense, ConvertsToEll) { auto ell_mtx = gko::matrix::Ell<>::create(mtx7->get_executor()); diff --git a/reference/test/matrix/sparsity_csr.cpp b/reference/test/matrix/sparsity_csr.cpp index c58d15d5337..f20e8ba3134 100644 --- a/reference/test/matrix/sparsity_csr.cpp +++ b/reference/test/matrix/sparsity_csr.cpp @@ -86,7 +86,7 @@ TEST_F(SparsityCsr, CanBeCreatedFromExistingCsrMatrix) auto mtx = Mtx::create(exec, std::move(csr_mtx)); - GKO_ASSERT_MTX_NEAR(comp_mtx.get(), mtx.get(), 0.0) + GKO_ASSERT_MTX_NEAR(comp_mtx.get(), mtx.get(), 0.0); } @@ -99,7 +99,7 @@ TEST_F(SparsityCsr, CanBeCreatedFromExistingDenseMatrix) auto mtx = Mtx::create(exec, std::move(dense_mtx)); - GKO_ASSERT_MTX_NEAR(comp_mtx.get(), mtx.get(), 0.0) + GKO_ASSERT_MTX_NEAR(comp_mtx.get(), mtx.get(), 0.0); } diff --git a/reference/test/matrix/sparsity_csr_kernels.cpp b/reference/test/matrix/sparsity_csr_kernels.cpp index 4af0f932880..e63bca5453e 100644 --- a/reference/test/matrix/sparsity_csr_kernels.cpp +++ b/reference/test/matrix/sparsity_csr_kernels.cpp @@ -288,9 +288,9 @@ TEST_F(SparsityCsr, CountsCorrectNumberOfDiagonalElements) gko::size_type ms_num_diags = 0; gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( - exec, mtx2.get(), m2_num_diags); + exec, mtx2.get(), &m2_num_diags); gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( - exec, mtx_s.get(), ms_num_diags); + exec, mtx_s.get(), &ms_num_diags); ASSERT_EQ(m2_num_diags, 3); ASSERT_EQ(ms_num_diags, 2);