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 ebe128f6dca..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" @@ -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_CSR_KERNEL(ValueType, IndexType) +GKO_NOT_COMPILED(GKO_HOOK_MODULE); +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL); + template GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL(ValueType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); @@ -352,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 3ae2e50e564..53dd2f02b12 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,32 @@ void Csr::move_to(Sellp *result) } +template +void Csr::convert_to( + SparsityCsr *result) const +{ + auto exec = this->get_executor(); + 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_; + if (result->value_.get_data()) { + tmp->value_ = result->value_; + } else { + tmp->value_ = gko::Array(exec, {one()}); + } + tmp->move_to(result); +} + + +template +void Csr::move_to( + SparsityCsr *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..2e099594e70 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/core/matrix/dense.cpp b/core/matrix/dense.cpp index 703b89c08fa..72c23343857 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_csr, dense::convert_to_sparsity_csr); } // namespace dense @@ -181,6 +183,22 @@ inline void conversion_helper(Sellp *result, } +template +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 = SparsityCsr::create( + exec, source->get_size(), num_stored_nonzeros); + exec->run(op(tmp.get(), source)); + tmp->move_to(result); +} + + } // namespace @@ -423,6 +441,38 @@ void Dense::move_to(Sellp *result) } +template +void Dense::convert_to(SparsityCsr *result) const +{ + conversion_helper(result, this, + dense::template make_convert_to_sparsity_csr< + decltype(result), const Dense *&>); +} + + +template +void Dense::move_to(SparsityCsr *result) +{ + this->convert_to(result); +} + + +template +void Dense::convert_to(SparsityCsr *result) const +{ + conversion_helper(result, this, + dense::template make_convert_to_sparsity_csr< + decltype(result), const Dense *&>); +} + + +template +void Dense::move_to(SparsityCsr *result) +{ + this->convert_to(result); +} + + namespace { diff --git a/core/matrix/dense_kernels.hpp b/core/matrix/dense_kernels.hpp index f5555a8c32d..4857fb81db9 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_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, \ 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_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 64% rename from core/matrix/sparsity.cpp rename to core/matrix/sparsity_csr.cpp index 5228b973df0..3c7cfa1363f 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( - this, num_diagonal_elements)); - ValueType one = 1.0; + 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 56% rename from core/matrix/sparsity_kernels.hpp rename to core/matrix/sparsity_csr_kernels.hpp index 4f50a7ee62f..f9af3dcdffa 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, \ - 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_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_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 4b410a4ae63..a74c431599a 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_csr(std::shared_ptr exec, + matrix::SparsityCsr *result, + const matrix::Dense *source) + GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL); + + template void count_nonzeros(std::shared_ptr exec, const matrix::Dense *source, size_type *result) diff --git a/cuda/matrix/sparsity_kernels.cu b/cuda/matrix/sparsity_csr_kernels.cu similarity index 78% rename from cuda/matrix/sparsity_kernels.cu rename to cuda/matrix/sparsity_csr_kernels.cu index 26792ea967d..2bdb8372030 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, - size_type &num_diagonal_elements) GKO_NOT_IMPLEMENTED; + 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 77de599c723..fd08a070cdb 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, ConvertToSparsityCsrIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared()); + 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()); + + GKO_ASSERT_MTX_NEAR(sparsity_mtx.get(), d_sparsity_mtx.get(), 1e-14); +} + + +TEST_F(Csr, MoveToSparsityCsrIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared()); + 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()); + + 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..391308ffd1b 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 SparsityCsr; + /** * 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 SparsityCsr; public: using EnableLinOp::convert_to; @@ -299,6 +304,10 @@ class Csr : public EnableLinOp>, void move_to(Sellp *result) override; + void convert_to(SparsityCsr *result) const override; + + void move_to(SparsityCsr *result) override; + void read(const mat_data &data) override; void write(mat_data &data) const override; diff --git a/include/ginkgo/core/matrix/dense.hpp b/include/ginkgo/core/matrix/dense.hpp index 3e23ebaa23d..ab0ae4fc9dd 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 SparsityCsr; + /** * 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 SparsityCsr; + friend class SparsityCsr; public: using EnableLinOp::convert_to; @@ -179,6 +186,14 @@ class Dense : public EnableLinOp>, void move_to(Sellp *result) override; + void convert_to(SparsityCsr *result) const override; + + void move_to(SparsityCsr *result) override; + + void convert_to(SparsityCsr *result) const override; + + void move_to(SparsityCsr *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_csr.hpp similarity index 77% rename from include/ginkgo/core/matrix/sparsity.hpp rename to include/ginkgo/core/matrix/sparsity_csr.hpp index e8e6a6e4c75..1e43ec0930f 100644 --- a/include/ginkgo/core/matrix/sparsity.hpp +++ b/include/ginkgo/core/matrix/sparsity_csr.hpp @@ -30,8 +30,11 @@ 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 #include @@ -42,8 +45,12 @@ namespace gko { namespace matrix { +template +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). * @@ -61,20 +68,22 @@ namespace matrix { * @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 index_type = IndexType; using value_type = ValueType; + using index_type = IndexType; using mat_data = matrix_data; @@ -88,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 @@ -117,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, @@ -136,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, @@ -155,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, @@ -179,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)) @@ -202,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 @@ -221,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)} { @@ -236,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 + */ + SparsityCsr(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, @@ -252,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 12f0b8cf17d..5654b40d753 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,41 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL); +template +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]; + + 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) { +#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_CSR_KERNEL); + + template void count_nonzeros(std::shared_ptr exec, const matrix::Dense *source, size_type *result) diff --git a/omp/matrix/sparsity_csr_kernels.cpp b/omp/matrix/sparsity_csr_kernels.cpp new file mode 100644 index 00000000000..830435fc376 --- /dev/null +++ b/omp/matrix/sparsity_csr_kernels.cpp @@ -0,0 +1,285 @@ +/************************************************************* +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_csr_kernels.hpp" + + +#include +#include +#include + + +#include + + +#include +#include +#include + + +#include "core/base/iterator_factory.hpp" +#include "omp/components/format_conversion.hpp" + + +namespace gko { +namespace kernels { +namespace omp { +/** + * @brief The SparsityCsr pattern format namespace. + * + * @ingroup sparsity + */ +namespace sparsity_csr { + + +template +void spmv(std::shared_ptr exec, + const matrix::SparsityCsr *a, + const matrix::Dense *b, matrix::Dense *c) +{ + 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) { + 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 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_CSR_SPMV_KERNEL); + + +template +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) +{ + 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); + auto val = a->get_const_value()[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 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_CSR_ADVANCED_SPMV_KERNEL); + + +template +void count_num_diagonal_elements( + std::shared_ptr exec, + const matrix::SparsityCsr *matrix, + 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_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); + + +template +void remove_diagonal_elements(std::shared_ptr exec, + matrix::SparsityCsr *matrix, + const IndexType *row_ptrs, + 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(); + 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; + } + 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) { + adj_idxs[nnz] = col_idxs[j]; + nnz++; + } + } + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL); + + +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) +{ + 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::SparsityCsr *trans, + const matrix::SparsityCsr *orig) +{ + 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::SparsityCsr *trans, + const matrix::SparsityCsr *orig) +{ + transpose_and_transform(exec, trans, orig); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL); + + +template +void sort_by_column_index(std::shared_ptr exec, + matrix::SparsityCsr *to_sort) +{ + 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_CSR_SORT_BY_COLUMN_INDEX); + + +template +void is_sorted_by_column_index( + std::shared_ptr exec, + 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(); + 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_CSR_IS_SORTED_BY_COLUMN_INDEX); + + +} // namespace sparsity_csr +} // namespace omp +} // namespace kernels +} // namespace gko diff --git a/omp/matrix/sparsity_kernels.cpp b/omp/matrix/sparsity_kernels.cpp deleted file mode 100644 index 4a3118ebf8f..00000000000 --- a/omp/matrix/sparsity_kernels.cpp +++ /dev/null @@ -1,152 +0,0 @@ -/************************************************************* -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/base/iterator_factory.hpp" -#include "omp/components/format_conversion.hpp" - - -namespace gko { -namespace kernels { -namespace omp { -/** - * @brief The Compressed sparse row matrix format namespace. - * - * @ingroup sparsity - */ -namespace sparsity { - - -template -void spmv(std::shared_ptr exec, - const matrix::Sparsity *a, - const matrix::Dense *b, - matrix::Dense *c) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SPARSITY_SPMV_KERNEL); - - -template -void advanced_spmv(std::shared_ptr exec, - const matrix::Dense *alpha, - const matrix::Sparsity *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); - - -template -void count_num_diagonal_elements( - std::shared_ptr exec, - const matrix::Sparsity *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); - - -template -void remove_diagonal_elements(std::shared_ptr exec, - matrix::Sparsity *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); - - -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; - - -template -void transpose_and_transform(std::shared_ptr exec, - matrix::Sparsity *trans, - const matrix::Sparsity *orig) - GKO_NOT_IMPLEMENTED; - - -template -void transpose(std::shared_ptr exec, - matrix::Sparsity *trans, - const matrix::Sparsity *orig) - GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_TRANSPOSE_KERNEL); - - -template -void sort_by_column_index(std::shared_ptr exec, - matrix::Sparsity *to_sort) - GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_SORT_BY_COLUMN_INDEX); - - -template -void is_sorted_by_column_index( - std::shared_ptr exec, - const matrix::Sparsity *to_check, - bool *is_sorted) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_IS_SORTED_BY_COLUMN_INDEX); - - -} // namespace sparsity -} // namespace omp -} // namespace kernels -} // namespace gko diff --git a/omp/test/matrix/CMakeLists.txt b/omp/test/matrix/CMakeLists.txt index 7cdf64bc6b5..b160723ab94 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_csr_kernels) diff --git a/omp/test/matrix/csr_kernels.cpp b/omp/test/matrix/csr_kernels.cpp index e5254f64ed3..72fa988ff5e 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, ConvertToSparsityCsrIsEquivalentToRef) +{ + set_up_apply_data(); + 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(Csr, MoveToSparsityCsrIsEquivalentToRef) +{ + set_up_apply_data(); + 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); } 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 new file mode 100644 index 00000000000..91852dbcb53 --- /dev/null +++ b/omp/test/matrix/sparsity_csr_kernels.cpp @@ -0,0 +1,307 @@ +/************************************************************* +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_csr_kernels.hpp" + + +#include +#include +#include + + +#include + + +#include +#include +#include +#include + + +#include "core/test/utils.hpp" + + +namespace { + + +class SparsityCsr : public ::testing::Test { +protected: + using Mtx = gko::matrix::SparsityCsr<>; + using Vec = gko::matrix::Dense<>; + using ComplexVec = gko::matrix::Dense>; + using ComplexMtx = gko::matrix::SparsityCsr>; + + SparsityCsr() : 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), 1.0, + 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(SparsityCsr, 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(SparsityCsr, 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(SparsityCsr, 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(SparsityCsr, 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(SparsityCsr, 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(SparsityCsr, CountsNumberOfDiagElementsIsEqualToRef) +{ + set_up_apply_data(); + gko::size_type num_diags = 0; + gko::size_type d_num_diags = 0; + + gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( + ref, mtx.get(), &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(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); + 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( + 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(SparsityCsr, 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(SparsityCsr, 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(SparsityCsr, 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(SparsityCsr, 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 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 ba975b720c9..82a20a8b1a4 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_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]; + + 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) { + 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_CSR_KERNEL); + + template void count_nonzeros(std::shared_ptr exec, const matrix::Dense *source, size_type *result) diff --git a/reference/matrix/sparsity_csr_kernels.cpp b/reference/matrix/sparsity_csr_kernels.cpp new file mode 100644 index 00000000000..42b4edd88a2 --- /dev/null +++ b/reference/matrix/sparsity_csr_kernels.cpp @@ -0,0 +1,274 @@ +/************************************************************* +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_csr_kernels.hpp" + + +#include +#include +#include + + +#include +#include +#include + + +#include "core/base/iterator_factory.hpp" +#include "reference/components/format_conversion.hpp" + + +namespace gko { +namespace kernels { +namespace reference { +/** + * @brief The SparsityCsr pattern format namespace. + * @ref SparsityCsr + * @ingroup sparsity + */ +namespace sparsity_csr { + + +template +void spmv(std::shared_ptr exec, + const matrix::SparsityCsr *a, + const matrix::Dense *b, matrix::Dense *c) +{ + 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) { + c->at(row, j) = zero(); + } + for (size_type k = row_ptrs[row]; + k < static_cast(row_ptrs[row + 1]); ++k) { + 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_CSR_SPMV_KERNEL); + + +template +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) +{ + 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); + 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) { + c->at(row, j) *= vbeta; + } + for (size_type k = row_ptrs[row]; + k < static_cast(row_ptrs[row + 1]); ++k) { + 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_CSR_ADVANCED_SPMV_KERNEL); + + +template +void count_num_diagonal_elements( + std::shared_ptr exec, + const matrix::SparsityCsr *matrix, + 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_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); + + +template +void remove_diagonal_elements(std::shared_ptr exec, + matrix::SparsityCsr *matrix, + const IndexType *row_ptrs, + 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(); + 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; + } + 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) { + adj_idxs[nnz] = col_idxs[j]; + nnz++; + } + } + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL); + + +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) +{ + 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::SparsityCsr *trans, + const matrix::SparsityCsr *orig) +{ + 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::SparsityCsr *trans, + const matrix::SparsityCsr *orig) +{ + transpose_and_transform(exec, trans, orig); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL); + + +template +void sort_by_column_index(std::shared_ptr exec, + matrix::SparsityCsr *to_sort) +{ + 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_CSR_SORT_BY_COLUMN_INDEX); + + +template +void is_sorted_by_column_index( + std::shared_ptr exec, + 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(); + 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_CSR_IS_SORTED_BY_COLUMN_INDEX); + + +} // namespace sparsity_csr +} // namespace reference +} // namespace kernels +} // namespace gko diff --git a/reference/matrix/sparsity_kernels.cpp b/reference/matrix/sparsity_kernels.cpp deleted file mode 100644 index 87a51d211ca..00000000000 --- a/reference/matrix/sparsity_kernels.cpp +++ /dev/null @@ -1,149 +0,0 @@ -/************************************************************* -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 "core/base/iterator_factory.hpp" -#include "reference/components/format_conversion.hpp" - - -namespace gko { -namespace kernels { -namespace reference { -/** - * @brief The Compressed sparse row matrix format namespace. - * @ref Sparsity - * @ingroup sparsity - */ -namespace sparsity { - - -template -void spmv(std::shared_ptr exec, - const matrix::Sparsity *a, - const matrix::Dense *b, - matrix::Dense *c) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SPARSITY_SPMV_KERNEL); - - -template -void advanced_spmv(std::shared_ptr exec, - const matrix::Dense *alpha, - const matrix::Sparsity *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); - - -template -void count_num_diagonal_elements( - std::shared_ptr exec, - const matrix::Sparsity *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); - - -template -void remove_diagonal_elements(std::shared_ptr exec, - matrix::Sparsity *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); - - -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; - - -template -void transpose_and_transform(std::shared_ptr exec, - matrix::Sparsity *trans, - const matrix::Sparsity *orig) - GKO_NOT_IMPLEMENTED; - - -template -void transpose(std::shared_ptr exec, - matrix::Sparsity *trans, - const matrix::Sparsity *orig) - GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_TRANSPOSE_KERNEL); - - -template -void sort_by_column_index(std::shared_ptr exec, - matrix::Sparsity *to_sort) - GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_SORT_BY_COLUMN_INDEX); - - -template -void is_sorted_by_column_index( - std::shared_ptr exec, - const matrix::Sparsity *to_check, - bool *is_sorted) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_IS_SORTED_BY_COLUMN_INDEX); - - -} // namespace sparsity -} // namespace reference -} // namespace kernels -} // namespace gko diff --git a/reference/test/matrix/CMakeLists.txt b/reference/test/matrix/CMakeLists.txt index 25b7597c3f3..7c0b5742eed 100644 --- a/reference/test/matrix/CMakeLists.txt +++ b/reference/test/matrix/CMakeLists.txt @@ -5,3 +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_csr) +ginkgo_create_test(sparsity_csr_kernels) diff --git a/reference/test/matrix/csr_kernels.cpp b/reference/test/matrix/csr_kernels.cpp index b03e8e810d0..f775ea018ee 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 SparsityCsr = gko::matrix::SparsityCsr<>; 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 SparsityCsr *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, ConvertsToSparsityCsr) +{ + auto sparsity_mtx = gko::matrix::SparsityCsr<>::create(mtx->get_executor()); + + mtx->convert_to(sparsity_mtx.get()); + + assert_equal_to_mtx(sparsity_mtx.get()); +} + + +TEST_F(Csr, MovesToSparsityCsr) +{ + 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()); + 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()); 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 new file mode 100644 index 00000000000..f20e8ba3134 --- /dev/null +++ b/reference/test/matrix/sparsity_csr.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 SparsityCsr : public ::testing::Test { +protected: + using v_type = double; + using i_type = int; + using Mtx = gko::matrix::SparsityCsr; + using Csr = gko::matrix::Csr; + using DenseMtx = gko::matrix::Dense; + + SparsityCsr() + : 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(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); + 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(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); + 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 diff --git a/reference/test/matrix/sparsity_csr_kernels.cpp b/reference/test/matrix/sparsity_csr_kernels.cpp new file mode 100644 index 00000000000..e63bca5453e --- /dev/null +++ b/reference/test/matrix/sparsity_csr_kernels.cpp @@ -0,0 +1,407 @@ +/************************************************************* +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_csr_kernels.hpp" + + +#include + + +#include + + +#include +#include +#include +#include +#include + + +#include "core/test/utils/assertions.hpp" + + +namespace { + + +class SparsityCsr : public ::testing::Test { +protected: + using Mtx = gko::matrix::SparsityCsr<>; + using Vec = gko::matrix::Dense<>; + + 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)), + 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(SparsityCsr, 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(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}); + + 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(SparsityCsr, 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(SparsityCsr, 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(SparsityCsr, 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(SparsityCsr, 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(SparsityCsr, 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(SparsityCsr, 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(SparsityCsr, 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(SparsityCsr, 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 m2_num_diags = 0; + gko::size_type ms_num_diags = 0; + + gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( + exec, mtx2.get(), &m2_num_diags); + gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( + exec, mtx_s.get(), &ms_num_diags); + + ASSERT_EQ(m2_num_diags, 3); + ASSERT_EQ(ms_num_diags, 2); +} + + +TEST_F(SparsityCsr, 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); + // clang-format on + 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_csr::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(SparsityCsr, RemovesDiagonalElementsForIncompleteRankMatrix) +{ + // clang-format off + 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::SparsityCsr<>::create( + exec, mtx_s->get_size(), mtx_s->get_num_nonzeros()); + tmp_mtx->copy_from(mtx2.get()); + + gko::kernels::reference::sparsity_csr::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(SparsityCsr, 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); + 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); +} + + +TEST_F(SparsityCsr, NonSquareMtxIsNotConvertibleToAdjacencyMatrix) +{ + ASSERT_THROW(mtx->to_adjacency_matrix(), gko::DimensionMismatch); +} + + +TEST_F(SparsityCsr, 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(SparsityCsr, RecognizeUnsortedMatrix) +{ + ASSERT_FALSE(mtx3_unsorted->is_sorted_by_column_index()); +} + + +TEST_F(SparsityCsr, SortSortedMatrix) +{ + auto matrix = mtx3_sorted->clone(); + + matrix->sort_by_column_index(); + + GKO_ASSERT_MTX_NEAR(matrix, mtx3_sorted, 0.0); +} + + +TEST_F(SparsityCsr, SortUnsortedMatrix) +{ + auto matrix = mtx3_unsorted->clone(); + + matrix->sort_by_column_index(); + + GKO_ASSERT_MTX_NEAR(matrix, mtx3_sorted, 0.0); +} + + +} // namespace 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