From 5545d209619412507d1aed5234346a05b5fad64c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Thomas=20Gr=C3=BCtzmacher?= Date: Wed, 3 Apr 2019 17:07:28 +0200 Subject: [PATCH] Removed `move_to_XXX` kernels and `COO::transpose` Removed unused `move_to_XXX` kernels from Dense and Csr. All calls to `move_to(Mtx *)`, which previously used the mentioned kernels, now use `convert_to_XXX` kernels. Since we did not implement any transpose functionality in Coo, the inheritance from `Transposable` and all related kernels were removed. Additionally, the tests for Coo, Csr & Dense were modified as follows: - Made include order compliant to guidelines - Newlines in the tests are also according to guidelines - Added MoveTo tests where they were not present - Added tests for Dense conversions (used already generated variables) The conversion from Dense to Hybrid with OpenMP was also slightly improved. --- core/device_hooks/common_kernels.inc.cpp | 42 ---------- core/matrix/coo.cpp | 26 ------- core/matrix/coo_kernels.hpp | 40 ++++------ core/matrix/csr.cpp | 16 +--- core/matrix/csr_kernels.hpp | 7 -- core/matrix/dense.cpp | 46 +++-------- core/matrix/dense_kernels.hpp | 28 ------- cuda/matrix/coo_kernels.cu | 19 ----- cuda/matrix/csr_kernels.cu | 9 --- cuda/matrix/dense_kernels.cu | 36 --------- cuda/test/matrix/coo_kernels.cpp | 13 ++-- cuda/test/matrix/csr_kernels.cpp | 73 ++++++++++++++---- cuda/test/matrix/dense_kernels.cpp | 91 ++++++++++++++++------ include/ginkgo/core/matrix/coo.hpp | 7 +- omp/matrix/coo_kernels.cpp | 19 ----- omp/matrix/csr_kernels.cpp | 12 --- omp/matrix/dense_kernels.cpp | 62 +++------------ omp/test/matrix/coo_kernels.cpp | 9 ++- omp/test/matrix/csr_kernels.cpp | 47 +++++++++++- omp/test/matrix/dense_kernels.cpp | 97 ++++++++++++------------ reference/matrix/coo_kernels.cpp | 19 ----- reference/matrix/csr_kernels.cpp | 24 ++---- reference/matrix/dense_kernels.cpp | 48 ------------ reference/test/matrix/coo_kernels.cpp | 12 ++- reference/test/matrix/csr_kernels.cpp | 39 +++++++--- reference/test/matrix/dense_kernels.cpp | 57 +++++--------- 26 files changed, 333 insertions(+), 565 deletions(-) diff --git a/core/device_hooks/common_kernels.inc.cpp b/core/device_hooks/common_kernels.inc.cpp index 0c89b78b9e3..c222e4fbc9e 100644 --- a/core/device_hooks/common_kernels.inc.cpp +++ b/core/device_hooks/common_kernels.inc.cpp @@ -103,49 +103,24 @@ GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_CSR_KERNEL); -template -GKO_DECLARE_DENSE_MOVE_TO_CSR_KERNEL(ValueType, IndexType) -GKO_NOT_COMPILED(GKO_HOOK_MODULE); -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_CSR_KERNEL); - - template GKO_DECLARE_DENSE_CONVERT_TO_ELL_KERNEL(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_ELL_KERNEL); -template -GKO_DECLARE_DENSE_MOVE_TO_ELL_KERNEL(ValueType, IndexType) -GKO_NOT_COMPILED(GKO_HOOK_MODULE); -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_ELL_KERNEL); - template GKO_DECLARE_DENSE_CONVERT_TO_HYBRID_KERNEL(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_HYBRID_KERNEL); -template -GKO_DECLARE_DENSE_MOVE_TO_HYBRID_KERNEL(ValueType, IndexType) -GKO_NOT_COMPILED(GKO_HOOK_MODULE); -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_HYBRID_KERNEL); - template GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL(ValueType, IndexType) 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_MOVE_TO_SELLP_KERNEL(ValueType, IndexType) -GKO_NOT_COMPILED(GKO_HOOK_MODULE); -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_SELLP_KERNEL); - template GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL(ValueType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); @@ -349,12 +324,6 @@ GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_CSR_CONVERT_TO_COO_KERNEL); -template -GKO_DECLARE_CSR_MOVE_TO_DENSE_KERNEL(ValueType, IndexType) -GKO_NOT_COMPILED(GKO_HOOK_MODULE); -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_CSR_MOVE_TO_DENSE_KERNEL); - template GKO_DECLARE_CSR_CONVERT_TO_SELLP_KERNEL(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); @@ -431,17 +400,6 @@ GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_COO_CONVERT_TO_DENSE_KERNEL); -template -GKO_DECLARE_COO_TRANSPOSE_KERNEL(ValueType, IndexType) -GKO_NOT_COMPILED(GKO_HOOK_MODULE); -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_COO_TRANSPOSE_KERNEL); - -template -GKO_DECLARE_COO_CONJ_TRANSPOSE_KERNEL(ValueType, IndexType) -GKO_NOT_COMPILED(GKO_HOOK_MODULE); -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_COO_CONJ_TRANSPOSE_KERNEL); - } // namespace coo diff --git a/core/matrix/coo.cpp b/core/matrix/coo.cpp index d8d83e8cb18..a9dddadf5ab 100644 --- a/core/matrix/coo.cpp +++ b/core/matrix/coo.cpp @@ -61,8 +61,6 @@ GKO_REGISTER_OPERATION(spmv2, coo::spmv2); GKO_REGISTER_OPERATION(advanced_spmv2, coo::advanced_spmv2); GKO_REGISTER_OPERATION(convert_to_csr, coo::convert_to_csr); GKO_REGISTER_OPERATION(convert_to_dense, coo::convert_to_dense); -GKO_REGISTER_OPERATION(transpose, coo::transpose); -GKO_REGISTER_OPERATION(conj_transpose, coo::conj_transpose); } // namespace coo @@ -192,30 +190,6 @@ void Coo::write(mat_data &data) const } -template -std::unique_ptr Coo::transpose() const -{ - auto exec = this->get_executor(); - auto trans_cpy = Coo::create(exec, gko::transpose(this->get_size()), - this->get_num_stored_elements()); - - exec->run(coo::make_transpose(trans_cpy.get(), this)); - return std::move(trans_cpy); -} - - -template -std::unique_ptr Coo::conj_transpose() const -{ - auto exec = this->get_executor(); - auto trans_cpy = Coo::create(exec, gko::transpose(this->get_size()), - this->get_num_stored_elements()); - - exec->run(coo::make_conj_transpose(trans_cpy.get(), this)); - return std::move(trans_cpy); -} - - #define GKO_DECLARE_COO_MATRIX(ValueType, IndexType) \ class Coo GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_COO_MATRIX); diff --git a/core/matrix/coo_kernels.hpp b/core/matrix/coo_kernels.hpp index 1b1ab376ed3..e8508ac1c8c 100644 --- a/core/matrix/coo_kernels.hpp +++ b/core/matrix/coo_kernels.hpp @@ -79,33 +79,19 @@ namespace kernels { matrix::Csr *result, \ const matrix::Coo *source) -#define GKO_DECLARE_COO_TRANSPOSE_KERNEL(ValueType, IndexType) \ - void transpose(std::shared_ptr exec, \ - matrix::Coo *trans, \ - const matrix::Coo *orig) - -#define GKO_DECLARE_COO_CONJ_TRANSPOSE_KERNEL(ValueType, IndexType) \ - void conj_transpose(std::shared_ptr exec, \ - matrix::Coo *trans, \ - const matrix::Coo *orig) - -#define GKO_DECLARE_ALL_AS_TEMPLATES \ - template \ - GKO_DECLARE_COO_SPMV_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_COO_ADVANCED_SPMV_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_COO_SPMV2_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_COO_ADVANCED_SPMV2_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_COO_CONVERT_TO_CSR_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_COO_CONVERT_TO_DENSE_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_COO_TRANSPOSE_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_COO_CONJ_TRANSPOSE_KERNEL(ValueType, IndexType) +#define GKO_DECLARE_ALL_AS_TEMPLATES \ + template \ + GKO_DECLARE_COO_SPMV_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_COO_ADVANCED_SPMV_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_COO_SPMV2_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_COO_ADVANCED_SPMV2_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_COO_CONVERT_TO_CSR_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_COO_CONVERT_TO_DENSE_KERNEL(ValueType, IndexType) namespace omp { diff --git a/core/matrix/csr.cpp b/core/matrix/csr.cpp index 1fae28e7c6c..aa70f311b98 100644 --- a/core/matrix/csr.cpp +++ b/core/matrix/csr.cpp @@ -39,8 +39,8 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include -#include #include +#include #include "core/matrix/csr_kernels.hpp" @@ -55,7 +55,6 @@ GKO_REGISTER_OPERATION(spmv, csr::spmv); GKO_REGISTER_OPERATION(advanced_spmv, csr::advanced_spmv); GKO_REGISTER_OPERATION(convert_to_coo, csr::convert_to_coo); GKO_REGISTER_OPERATION(convert_to_dense, csr::convert_to_dense); -GKO_REGISTER_OPERATION(move_to_dense, csr::move_to_dense); GKO_REGISTER_OPERATION(convert_to_sellp, csr::convert_to_sellp); GKO_REGISTER_OPERATION(calculate_total_cols, csr::calculate_total_cols); GKO_REGISTER_OPERATION(convert_to_ell, csr::convert_to_ell); @@ -103,13 +102,7 @@ void Csr::convert_to( template void Csr::move_to(Coo *result) { - auto exec = this->get_executor(); - auto tmp = Coo::create( - exec, this->get_size(), this->get_num_stored_elements()); - tmp->values_ = std::move(this->values_); - tmp->col_idxs_ = std::move(this->col_idxs_); - exec->run(csr::make_convert_to_coo(tmp.get(), this)); - tmp->move_to(result); + this->convert_to(result); } @@ -126,10 +119,7 @@ void Csr::convert_to(Dense *result) const template void Csr::move_to(Dense *result) { - auto exec = this->get_executor(); - auto tmp = Dense::create(exec, this->get_size()); - exec->run(csr::make_move_to_dense(tmp.get(), this)); - tmp->move_to(result); + this->convert_to(result); } diff --git a/core/matrix/csr_kernels.hpp b/core/matrix/csr_kernels.hpp index 189f0160185..b8799a5696f 100644 --- a/core/matrix/csr_kernels.hpp +++ b/core/matrix/csr_kernels.hpp @@ -64,11 +64,6 @@ namespace kernels { matrix::Dense *result, \ const matrix::Csr *source) -#define GKO_DECLARE_CSR_MOVE_TO_DENSE_KERNEL(ValueType, IndexType) \ - void move_to_dense(std::shared_ptr exec, \ - matrix::Dense *result, \ - matrix::Csr *source) - #define GKO_DECLARE_CSR_CONVERT_TO_COO_KERNEL(ValueType, IndexType) \ void convert_to_coo(std::shared_ptr exec, \ matrix::Coo *result, \ @@ -113,8 +108,6 @@ namespace kernels { template \ GKO_DECLARE_CSR_CONVERT_TO_DENSE_KERNEL(ValueType, IndexType); \ template \ - GKO_DECLARE_CSR_MOVE_TO_DENSE_KERNEL(ValueType, IndexType); \ - template \ GKO_DECLARE_CSR_CONVERT_TO_COO_KERNEL(ValueType, IndexType); \ template \ GKO_DECLARE_CSR_CONVERT_TO_SELLP_KERNEL(ValueType, IndexType); \ diff --git a/core/matrix/dense.cpp b/core/matrix/dense.cpp index 5f9deb23f61..703b89c08fa 100644 --- a/core/matrix/dense.cpp +++ b/core/matrix/dense.cpp @@ -72,13 +72,9 @@ GKO_REGISTER_OPERATION(transpose, dense::transpose); GKO_REGISTER_OPERATION(conj_transpose, dense::conj_transpose); GKO_REGISTER_OPERATION(convert_to_coo, dense::convert_to_coo); GKO_REGISTER_OPERATION(convert_to_csr, dense::convert_to_csr); -GKO_REGISTER_OPERATION(move_to_csr, dense::move_to_csr); GKO_REGISTER_OPERATION(convert_to_ell, dense::convert_to_ell); -GKO_REGISTER_OPERATION(move_to_ell, dense::move_to_ell); GKO_REGISTER_OPERATION(convert_to_hybrid, dense::convert_to_hybrid); -GKO_REGISTER_OPERATION(move_to_hybrid, dense::move_to_hybrid); GKO_REGISTER_OPERATION(convert_to_sellp, dense::convert_to_sellp); -GKO_REGISTER_OPERATION(move_to_sellp, dense::move_to_sellp); } // namespace dense @@ -268,9 +264,7 @@ void Dense::convert_to(Coo *result) const template void Dense::move_to(Coo *result) { - conversion_helper(result, this, - dense::template make_convert_to_coo *&>); + this->convert_to(result); } @@ -287,9 +281,7 @@ void Dense::convert_to(Coo *result) const template void Dense::move_to(Coo *result) { - conversion_helper(result, this, - dense::template make_convert_to_coo *&>); + this->convert_to(result); } @@ -307,10 +299,7 @@ void Dense::convert_to(Csr *result) const template void Dense::move_to(Csr *result) { - conversion_helper(result, this, - dense::template make_move_to_csr *&>); - result->make_srow(); + this->convert_to(result); } @@ -328,10 +317,7 @@ void Dense::convert_to(Csr *result) const template void Dense::move_to(Csr *result) { - conversion_helper(result, this, - dense::template make_move_to_csr *&>); - result->make_srow(); + this->convert_to(result); } @@ -348,9 +334,7 @@ void Dense::convert_to(Ell *result) const template void Dense::move_to(Ell *result) { - conversion_helper(result, this, - dense::template make_move_to_ell *&>); + this->convert_to(result); } @@ -367,9 +351,7 @@ void Dense::convert_to(Ell *result) const template void Dense::move_to(Ell *result) { - conversion_helper(result, this, - dense::template make_move_to_ell *&>); + this->convert_to(result); } @@ -386,9 +368,7 @@ void Dense::convert_to(Hybrid *result) const template void Dense::move_to(Hybrid *result) { - conversion_helper(result, this, - dense::template make_move_to_hybrid *&>); + this->convert_to(result); } @@ -405,9 +385,7 @@ void Dense::convert_to(Hybrid *result) const template void Dense::move_to(Hybrid *result) { - conversion_helper(result, this, - dense::template make_move_to_hybrid *&>); + this->convert_to(result); } @@ -424,9 +402,7 @@ void Dense::convert_to(Sellp *result) const template void Dense::move_to(Sellp *result) { - conversion_helper(result, this, - dense::template make_move_to_sellp *&>); + this->convert_to(result); } @@ -443,9 +419,7 @@ void Dense::convert_to(Sellp *result) const template void Dense::move_to(Sellp *result) { - conversion_helper(result, this, - dense::template make_move_to_sellp *&>); + this->convert_to(result); } diff --git a/core/matrix/dense_kernels.hpp b/core/matrix/dense_kernels.hpp index 9420b7be3c1..f5555a8c32d 100644 --- a/core/matrix/dense_kernels.hpp +++ b/core/matrix/dense_kernels.hpp @@ -83,41 +83,21 @@ namespace kernels { matrix::Csr<_type, _prec> *other, \ const matrix::Dense<_type> *source) -#define GKO_DECLARE_DENSE_MOVE_TO_CSR_KERNEL(_type, _prec) \ - void move_to_csr(std::shared_ptr exec, \ - matrix::Csr<_type, _prec> *other, \ - const matrix::Dense<_type> *source) - #define GKO_DECLARE_DENSE_CONVERT_TO_ELL_KERNEL(_type, _prec) \ void convert_to_ell(std::shared_ptr exec, \ matrix::Ell<_type, _prec> *other, \ const matrix::Dense<_type> *source) -#define GKO_DECLARE_DENSE_MOVE_TO_ELL_KERNEL(_type, _prec) \ - void move_to_ell(std::shared_ptr exec, \ - matrix::Ell<_type, _prec> *other, \ - const matrix::Dense<_type> *source) - #define GKO_DECLARE_DENSE_CONVERT_TO_HYBRID_KERNEL(_type, _prec) \ void convert_to_hybrid(std::shared_ptr exec, \ matrix::Hybrid<_type, _prec> *other, \ const matrix::Dense<_type> *source) -#define GKO_DECLARE_DENSE_MOVE_TO_HYBRID_KERNEL(_type, _prec) \ - void move_to_hybrid(std::shared_ptr exec, \ - matrix::Hybrid<_type, _prec> *other, \ - const matrix::Dense<_type> *source) - #define GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL(_type, _prec) \ void convert_to_sellp(std::shared_ptr exec, \ matrix::Sellp<_type, _prec> *other, \ const matrix::Dense<_type> *source) -#define GKO_DECLARE_DENSE_MOVE_TO_SELLP_KERNEL(_type, _prec) \ - void move_to_sellp(std::shared_ptr exec, \ - matrix::Sellp<_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) @@ -166,19 +146,11 @@ namespace kernels { template \ GKO_DECLARE_DENSE_CONVERT_TO_CSR_KERNEL(ValueType, IndexType); \ template \ - GKO_DECLARE_DENSE_MOVE_TO_CSR_KERNEL(ValueType, IndexType); \ - template \ GKO_DECLARE_DENSE_CONVERT_TO_ELL_KERNEL(ValueType, IndexType); \ template \ - GKO_DECLARE_DENSE_MOVE_TO_ELL_KERNEL(ValueType, IndexType); \ - template \ GKO_DECLARE_DENSE_CONVERT_TO_HYBRID_KERNEL(ValueType, IndexType); \ template \ - GKO_DECLARE_DENSE_MOVE_TO_HYBRID_KERNEL(ValueType, IndexType); \ - template \ GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL(ValueType, IndexType); \ - template \ - GKO_DECLARE_DENSE_MOVE_TO_SELLP_KERNEL(ValueType, IndexType); \ template \ GKO_DECLARE_DENSE_COUNT_NONZEROS_KERNEL(ValueType); \ template \ diff --git a/cuda/matrix/coo_kernels.cu b/cuda/matrix/coo_kernels.cu index b7dda054f77..a4a1acc6727 100644 --- a/cuda/matrix/coo_kernels.cu +++ b/cuda/matrix/coo_kernels.cu @@ -346,25 +346,6 @@ void convert_row_idxs_to_ptrs(std::shared_ptr exec, } -template -void transpose(std::shared_ptr exec, - matrix::Coo *trans, - const matrix::Coo *orig) - GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_COO_TRANSPOSE_KERNEL); - - -template -void conj_transpose(std::shared_ptr exec, - matrix::Coo *trans, - const matrix::Coo *orig) - GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_COO_CONJ_TRANSPOSE_KERNEL); - - template void convert_to_csr(std::shared_ptr exec, matrix::Csr *result, diff --git a/cuda/matrix/csr_kernels.cu b/cuda/matrix/csr_kernels.cu index 81fe1bac2fd..d4107f980e9 100644 --- a/cuda/matrix/csr_kernels.cu +++ b/cuda/matrix/csr_kernels.cu @@ -940,15 +940,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_CSR_CONVERT_TO_DENSE_KERNEL); -template -void move_to_dense( - std::shared_ptr exec, matrix::Dense *result, - matrix::Csr *source) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_CSR_MOVE_TO_DENSE_KERNEL); - - namespace kernel { diff --git a/cuda/matrix/dense_kernels.cu b/cuda/matrix/dense_kernels.cu index 11039b23037..aea60a903e4 100644 --- a/cuda/matrix/dense_kernels.cu +++ b/cuda/matrix/dense_kernels.cu @@ -524,15 +524,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_CSR_KERNEL); -template -void move_to_csr(std::shared_ptr exec, - matrix::Csr *result, - const matrix::Dense *source) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_CSR_KERNEL); - - namespace kernel { @@ -596,15 +587,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_ELL_KERNEL); -template -void move_to_ell(std::shared_ptr exec, - matrix::Ell *result, - const matrix::Dense *source) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_ELL_KERNEL); - - template void convert_to_hybrid(std::shared_ptr exec, matrix::Hybrid *result, @@ -615,15 +597,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_HYBRID_KERNEL); -template -void move_to_hybrid(std::shared_ptr exec, - matrix::Hybrid *result, - const matrix::Dense *source) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_HYBRID_KERNEL); - - namespace kernel { @@ -757,15 +730,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL); -template -void move_to_sellp(std::shared_ptr exec, - matrix::Sellp *result, - const matrix::Dense *source) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_SELLP_KERNEL); - - namespace kernel { diff --git a/cuda/test/matrix/coo_kernels.cpp b/cuda/test/matrix/coo_kernels.cpp index fb4d493b981..9958d623156 100644 --- a/cuda/test/matrix/coo_kernels.cpp +++ b/cuda/test/matrix/coo_kernels.cpp @@ -30,9 +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 +#include "core/matrix/coo_kernels.hpp" #include @@ -41,13 +39,17 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include -#include #include #include #include +#include +#include #include +#include "core/test/utils.hpp" + + namespace { @@ -177,6 +179,7 @@ TEST_F(Coo, SimpleApplyToDenseMatrixIsEquivalentToRef) TEST_F(Coo, 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()); @@ -209,7 +212,6 @@ TEST_F(Coo, AdvancedApplyAddToDenseMatrixIsEquivalentToRef) TEST_F(Coo, ConvertToDenseIsEquivalentToRef) { set_up_apply_data(); - auto dense_mtx = gko::matrix::Dense<>::create(ref); auto ddense_mtx = gko::matrix::Dense<>::create(cuda); @@ -223,7 +225,6 @@ TEST_F(Coo, ConvertToDenseIsEquivalentToRef) TEST_F(Coo, ConvertToCsrIsEquivalentToRef) { set_up_apply_data(); - auto dense_mtx = gko::matrix::Dense<>::create(ref); auto csr_mtx = gko::matrix::Csr<>::create(ref); auto dcsr_mtx = gko::matrix::Csr<>::create(cuda); diff --git a/cuda/test/matrix/csr_kernels.cpp b/cuda/test/matrix/csr_kernels.cpp index cf064d310bf..74eac036439 100644 --- a/cuda/test/matrix/csr_kernels.cpp +++ b/cuda/test/matrix/csr_kernels.cpp @@ -30,8 +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 "core/matrix/csr_kernels.hpp" #include @@ -42,13 +41,13 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include +#include +#include #include -#include #include +#include -#include "core/matrix/dense_kernels.hpp" -#include "core/matrix/csr_kernels.hpp" #include "core/test/utils.hpp" @@ -142,6 +141,7 @@ class Csr : public ::testing::Test { TEST_F(Csr, StrategyAfterCopyIsEquivalentToRef) { set_up_apply_data(std::make_shared(32)); + ASSERT_EQ(mtx->get_strategy()->get_name(), dmtx->get_strategy()->get_name()); } @@ -339,7 +339,6 @@ TEST_F(Csr, ConjugateTransposeIsEquivalentToRef) TEST_F(Csr, ConvertToDenseIsEquivalentToRef) { set_up_apply_data(std::make_shared()); - auto dense_mtx = gko::matrix::Dense<>::create(ref); auto ddense_mtx = gko::matrix::Dense<>::create(cuda); @@ -350,10 +349,22 @@ TEST_F(Csr, ConvertToDenseIsEquivalentToRef) } -TEST_F(Csr, ConvertToEllIsEquivalentToRef) +TEST_F(Csr, MoveToDenseIsEquivalentToRef) { set_up_apply_data(std::make_shared()); + auto dense_mtx = gko::matrix::Dense<>::create(ref); + auto ddense_mtx = gko::matrix::Dense<>::create(cuda); + + mtx->move_to(dense_mtx.get()); + dmtx->move_to(ddense_mtx.get()); + GKO_ASSERT_MTX_NEAR(dense_mtx.get(), ddense_mtx.get(), 1e-14); +} + + +TEST_F(Csr, ConvertToEllIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared()); auto ell_mtx = gko::matrix::Ell<>::create(ref); auto dell_mtx = gko::matrix::Ell<>::create(cuda); @@ -364,16 +375,27 @@ TEST_F(Csr, ConvertToEllIsEquivalentToRef) } -TEST_F(Csr, CalculateMaxNnzPerRowIsEquivalentToRef) +TEST_F(Csr, MoveToEllIsEquivalentToRef) { set_up_apply_data(std::make_shared()); + auto ell_mtx = gko::matrix::Ell<>::create(ref); + auto dell_mtx = gko::matrix::Ell<>::create(cuda); + + mtx->move_to(ell_mtx.get()); + dmtx->move_to(dell_mtx.get()); + + GKO_ASSERT_MTX_NEAR(ell_mtx.get(), dell_mtx.get(), 1e-14); +} + +TEST_F(Csr, CalculateMaxNnzPerRowIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared()); gko::size_type max_nnz_per_row; gko::size_type dmax_nnz_per_row; gko::kernels::reference::csr::calculate_max_nnz_per_row(ref, mtx.get(), &max_nnz_per_row); - gko::kernels::cuda::csr::calculate_max_nnz_per_row(cuda, dmtx.get(), &dmax_nnz_per_row); @@ -384,23 +406,32 @@ TEST_F(Csr, CalculateMaxNnzPerRowIsEquivalentToRef) TEST_F(Csr, ConvertToCooIsEquivalentToRef) { set_up_apply_data(std::make_shared()); - - auto dense_mtx = gko::matrix::Dense<>::create(ref); auto coo_mtx = gko::matrix::Coo<>::create(ref); auto dcoo_mtx = gko::matrix::Coo<>::create(cuda); - mtx->convert_to(dense_mtx.get()); - dense_mtx->convert_to(coo_mtx.get()); + mtx->convert_to(coo_mtx.get()); dmtx->convert_to(dcoo_mtx.get()); GKO_ASSERT_MTX_NEAR(coo_mtx.get(), dcoo_mtx.get(), 1e-14); } -TEST_F(Csr, ConvertToSellpIsEquivalentToRef) +TEST_F(Csr, MoveToCooIsEquivalentToRef) { set_up_apply_data(std::make_shared()); + auto coo_mtx = gko::matrix::Coo<>::create(ref); + auto dcoo_mtx = gko::matrix::Coo<>::create(cuda); + + mtx->move_to(coo_mtx.get()); + dmtx->move_to(dcoo_mtx.get()); + + GKO_ASSERT_MTX_NEAR(coo_mtx.get(), dcoo_mtx.get(), 1e-14); +} + +TEST_F(Csr, ConvertToSellpIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared()); auto sellp_mtx = gko::matrix::Sellp<>::create(ref); auto dsellp_mtx = gko::matrix::Sellp<>::create(cuda); @@ -411,10 +442,22 @@ TEST_F(Csr, ConvertToSellpIsEquivalentToRef) } -TEST_F(Csr, CalculateTotalColsIsEquivalentToRef) +TEST_F(Csr, MoveToSellpIsEquivalentToRef) { set_up_apply_data(std::make_shared()); + auto sellp_mtx = gko::matrix::Sellp<>::create(ref); + auto dsellp_mtx = gko::matrix::Sellp<>::create(cuda); + + mtx->move_to(sellp_mtx.get()); + dmtx->move_to(dsellp_mtx.get()); + GKO_ASSERT_MTX_NEAR(sellp_mtx.get(), dsellp_mtx.get(), 1e-14); +} + + +TEST_F(Csr, CalculateTotalColsIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared()); gko::size_type total_cols; gko::size_type dtotal_cols; diff --git a/cuda/test/matrix/dense_kernels.cpp b/cuda/test/matrix/dense_kernels.cpp index 1cd0a276650..c153e1a5c74 100644 --- a/cuda/test/matrix/dense_kernels.cpp +++ b/cuda/test/matrix/dense_kernels.cpp @@ -30,24 +30,24 @@ 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 -#include -#include -#include +#include "core/matrix/dense_kernels.hpp" -#include "core/matrix/dense_kernels.hpp" +#include #include -#include +#include +#include +#include +#include +#include +#include -#include +#include "core/test/utils.hpp" namespace { @@ -148,12 +148,12 @@ class Dense : public ::testing::Test { TEST_F(Dense, SingleVectorCudaScaleIsEquivalentToRef) { set_up_vector_data(1); + auto result = Mtx::create(ref); x->scale(alpha.get()); dx->scale(dalpha.get()); - - auto result = Mtx::create(ref); result->copy_from(dx.get()); + GKO_ASSERT_MTX_NEAR(result, x, 1e-14); } @@ -295,7 +295,6 @@ TEST_F(Dense, IsConjugateTransposable) TEST_F(Dense, ConvertToCooIsEquivalentToRef) { set_up_apply_data(); - auto coo_mtx = gko::matrix::Coo<>::create(ref); auto dcoo_mtx = gko::matrix::Coo<>::create(cuda); @@ -308,10 +307,24 @@ TEST_F(Dense, ConvertToCooIsEquivalentToRef) } -TEST_F(Dense, ConvertToCsrIsEquivalentToRef) +TEST_F(Dense, MoveToCooIsEquivalentToRef) { set_up_apply_data(); + auto coo_mtx = gko::matrix::Coo<>::create(ref); + auto dcoo_mtx = gko::matrix::Coo<>::create(cuda); + + x->move_to(coo_mtx.get()); + dx->move_to(dcoo_mtx.get()); + + ASSERT_EQ(dcoo_mtx->get_num_stored_elements(), + coo_mtx->get_num_stored_elements()); + GKO_ASSERT_MTX_NEAR(dcoo_mtx.get(), coo_mtx.get(), 1e-14); +} + +TEST_F(Dense, ConvertToCsrIsEquivalentToRef) +{ + set_up_apply_data(); auto csr_mtx = gko::matrix::Csr<>::create(ref); auto dcsr_mtx = gko::matrix::Csr<>::create(cuda); @@ -322,10 +335,22 @@ TEST_F(Dense, ConvertToCsrIsEquivalentToRef) } -TEST_F(Dense, ConvertToEllIsEquivalentToRef) +TEST_F(Dense, MoveToCsrIsEquivalentToRef) { set_up_apply_data(); + auto csr_mtx = gko::matrix::Csr<>::create(ref); + auto dcsr_mtx = gko::matrix::Csr<>::create(cuda); + + x->move_to(csr_mtx.get()); + dx->move_to(dcsr_mtx.get()); + + GKO_ASSERT_MTX_NEAR(dcsr_mtx.get(), csr_mtx.get(), 1e-14); +} + +TEST_F(Dense, ConvertToEllIsEquivalentToRef) +{ + set_up_apply_data(); auto ell_mtx = gko::matrix::Ell<>::create(ref); auto dell_mtx = gko::matrix::Ell<>::create(cuda); @@ -336,10 +361,22 @@ TEST_F(Dense, ConvertToEllIsEquivalentToRef) } -TEST_F(Dense, ConvertToSellpIsEquivalentToRef) +TEST_F(Dense, MoveToEllIsEquivalentToRef) { set_up_apply_data(); + auto ell_mtx = gko::matrix::Ell<>::create(ref); + auto dell_mtx = gko::matrix::Ell<>::create(cuda); + x->move_to(ell_mtx.get()); + dx->move_to(dell_mtx.get()); + + GKO_ASSERT_MTX_NEAR(dell_mtx.get(), ell_mtx.get(), 1e-14); +} + + +TEST_F(Dense, ConvertToSellpIsEquivalentToRef) +{ + set_up_apply_data(); auto sellp_mtx = gko::matrix::Sellp<>::create(ref); auto dsellp_mtx = gko::matrix::Sellp<>::create(cuda); @@ -350,10 +387,22 @@ TEST_F(Dense, ConvertToSellpIsEquivalentToRef) } -TEST_F(Dense, CountNNZIsEquivalentToRef) +TEST_F(Dense, MoveToSellpIsEquivalentToRef) { set_up_apply_data(); + auto sellp_mtx = gko::matrix::Sellp<>::create(ref); + auto dsellp_mtx = gko::matrix::Sellp<>::create(cuda); + + x->move_to(sellp_mtx.get()); + dx->move_to(dsellp_mtx.get()); + + GKO_ASSERT_MTX_NEAR(sellp_mtx, dsellp_mtx, 1e-14); +} + +TEST_F(Dense, CountNNZIsEquivalentToRef) +{ + set_up_apply_data(); gko::size_type nnz; gko::size_type dnnz; @@ -367,13 +416,9 @@ TEST_F(Dense, CountNNZIsEquivalentToRef) TEST_F(Dense, CalculateNNZPerRowIsEquivalentToRef) { set_up_apply_data(); - - gko::Array nnz_per_row; - nnz_per_row.set_executor(ref); + gko::Array nnz_per_row(ref); nnz_per_row.resize_and_reset(x->get_size()[0]); - - gko::Array dnnz_per_row; - dnnz_per_row.set_executor(cuda); + gko::Array dnnz_per_row(cuda); dnnz_per_row.resize_and_reset(dx->get_size()[0]); gko::kernels::reference::dense::calculate_nonzeros_per_row(ref, x.get(), @@ -391,7 +436,6 @@ TEST_F(Dense, CalculateNNZPerRowIsEquivalentToRef) TEST_F(Dense, CalculateMaxNNZPerRowIsEquivalentToRef) { set_up_apply_data(); - gko::size_type max_nnz; gko::size_type dmax_nnz; @@ -407,7 +451,6 @@ TEST_F(Dense, CalculateMaxNNZPerRowIsEquivalentToRef) TEST_F(Dense, CalculateTotalColsIsEquivalentToRef) { set_up_apply_data(); - gko::size_type total_cols; gko::size_type dtotal_cols; diff --git a/include/ginkgo/core/matrix/coo.hpp b/include/ginkgo/core/matrix/coo.hpp index 9605a6121d1..ac6f271a153 100644 --- a/include/ginkgo/core/matrix/coo.hpp +++ b/include/ginkgo/core/matrix/coo.hpp @@ -75,8 +75,7 @@ class Coo : public EnableLinOp>, public ConvertibleTo>, public ConvertibleTo>, public ReadableFromMatrixData, - public WritableToMatrixData, - public Transposable { + public WritableToMatrixData { friend class EnableCreateMethod; friend class EnablePolymorphicObject; friend class Csr; @@ -102,10 +101,6 @@ class Coo : public EnableLinOp>, void write(mat_data &data) const override; - std::unique_ptr transpose() const override; - - std::unique_ptr conj_transpose() const override; - /** * Returns the values of the matrix. * diff --git a/omp/matrix/coo_kernels.cpp b/omp/matrix/coo_kernels.cpp index 32b38303de3..8f1df3cfbc2 100644 --- a/omp/matrix/coo_kernels.cpp +++ b/omp/matrix/coo_kernels.cpp @@ -154,25 +154,6 @@ void convert_row_idxs_to_ptrs(std::shared_ptr exec, } -template -void transpose(std::shared_ptr exec, - matrix::Coo *trans, - const matrix::Coo *orig) - GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_COO_TRANSPOSE_KERNEL); - - -template -void conj_transpose(std::shared_ptr exec, - matrix::Coo *trans, - const matrix::Coo *orig) - GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_COO_CONJ_TRANSPOSE_KERNEL); - - template void convert_to_csr(std::shared_ptr exec, matrix::Csr *result, diff --git a/omp/matrix/csr_kernels.cpp b/omp/matrix/csr_kernels.cpp index a79889b07a1..f923ae2d093 100644 --- a/omp/matrix/csr_kernels.cpp +++ b/omp/matrix/csr_kernels.cpp @@ -177,18 +177,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_CSR_CONVERT_TO_DENSE_KERNEL); -template -void move_to_dense(std::shared_ptr exec, - matrix::Dense *result, - matrix::Csr *source) -{ - omp::csr::convert_to_dense(exec, result, source); -} - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_CSR_MOVE_TO_DENSE_KERNEL); - - template void convert_to_sellp(std::shared_ptr exec, matrix::Sellp *result, diff --git a/omp/matrix/dense_kernels.cpp b/omp/matrix/dense_kernels.cpp index 438a93238d9..12f0b8cf17d 100644 --- a/omp/matrix/dense_kernels.cpp +++ b/omp/matrix/dense_kernels.cpp @@ -281,18 +281,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_CSR_KERNEL); -template -void move_to_csr(std::shared_ptr exec, - matrix::Csr *result, - const matrix::Dense *source) -{ - omp::dense::convert_to_csr(exec, result, source); -} - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_CSR_KERNEL); - - template void convert_to_ell(std::shared_ptr exec, matrix::Ell *result, @@ -326,18 +314,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_ELL_KERNEL); -template -void move_to_ell(std::shared_ptr exec, - matrix::Ell *result, - const matrix::Dense *source) -{ - omp::dense::convert_to_ell(exec, result, source); -} - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_ELL_KERNEL); - - template void convert_to_hybrid(std::shared_ptr exec, matrix::Hybrid *result, @@ -368,6 +344,8 @@ void convert_to_hybrid(std::shared_ptr exec, } size_type coo_idx = 0; + // FIXME: This parallelization may cause the COO part to not being sorted by + // row idx #pragma omp parallel for for (size_type row = 0; row < num_rows; row++) { size_type col_idx = 0; @@ -384,13 +362,17 @@ void convert_to_hybrid(std::shared_ptr exec, while (col < num_cols) { auto val = source->at(row, col); if (val != zero()) { + size_type current_coo_idx; + // Use the critical section for accessing the coo_idx only, the + // rest can be performed in parallel since the index is unique #pragma omp critical { - coo_val[coo_idx] = val; - coo_col[coo_idx] = col; - coo_row[coo_idx] = row; - coo_idx++; + current_coo_idx = coo_idx; + ++coo_idx; } + coo_val[current_coo_idx] = val; + coo_col[current_coo_idx] = col; + coo_row[current_coo_idx] = row; } col++; } @@ -401,18 +383,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_HYBRID_KERNEL); -template -void move_to_hybrid(std::shared_ptr exec, - matrix::Hybrid *result, - const matrix::Dense *source) -{ - omp::dense::convert_to_hybrid(exec, result, source); -} - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_HYBRID_KERNEL); - - template void convert_to_sellp(std::shared_ptr exec, matrix::Sellp *result, @@ -485,18 +455,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL); -template -void move_to_sellp(std::shared_ptr exec, - matrix::Sellp *result, - const matrix::Dense *source) -{ - omp::dense::convert_to_sellp(exec, result, source); -} - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_SELLP_KERNEL); - - template void count_nonzeros(std::shared_ptr exec, const matrix::Dense *source, size_type *result) diff --git a/omp/test/matrix/coo_kernels.cpp b/omp/test/matrix/coo_kernels.cpp index d46d04a1657..355f31fcfa4 100644 --- a/omp/test/matrix/coo_kernels.cpp +++ b/omp/test/matrix/coo_kernels.cpp @@ -30,7 +30,7 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include +#include "core/matrix/coo_kernels.hpp" #include @@ -39,14 +39,17 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include -#include #include #include #include +#include #include #include +#include "core/test/utils.hpp" + + namespace { @@ -176,6 +179,7 @@ TEST_F(Coo, SimpleApplyToDenseMatrixIsEquivalentToRef) TEST_F(Coo, 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()); @@ -208,7 +212,6 @@ TEST_F(Coo, AdvancedApplyAddToDenseMatrixIsEquivalentToRef) TEST_F(Coo, ConvertToCsrIsEquivalentToRef) { set_up_apply_data(); - auto csr_mtx = gko::matrix::Csr<>::create(ref); auto dcsr_mtx = gko::matrix::Csr<>::create(omp); diff --git a/omp/test/matrix/csr_kernels.cpp b/omp/test/matrix/csr_kernels.cpp index 9939949dbd2..02182e111ea 100644 --- a/omp/test/matrix/csr_kernels.cpp +++ b/omp/test/matrix/csr_kernels.cpp @@ -30,7 +30,7 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include +#include "core/matrix/csr_kernels.hpp" #include @@ -39,13 +39,16 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include -#include #include #include #include +#include #include +#include "core/test/utils.hpp" + + namespace { @@ -197,7 +200,6 @@ TEST_F(Csr, ConjugateTransposeIsEquivalentToRef) TEST_F(Csr, ConvertToCooIsEquivalentToRef) { set_up_apply_data(); - auto coo_mtx = gko::matrix::Coo<>::create(ref); auto dcoo_mtx = gko::matrix::Coo<>::create(omp); @@ -208,4 +210,43 @@ TEST_F(Csr, ConvertToCooIsEquivalentToRef) } +TEST_F(Csr, MoveToCooIsEquivalentToRef) +{ + set_up_apply_data(); + auto coo_mtx = gko::matrix::Coo<>::create(ref); + auto dcoo_mtx = gko::matrix::Coo<>::create(omp); + + mtx->move_to(coo_mtx.get()); + dmtx->move_to(dcoo_mtx.get()); + + GKO_ASSERT_MTX_NEAR(coo_mtx.get(), dcoo_mtx.get(), 1e-14); +} + + +TEST_F(Csr, ConvertToDenseIsEquivalentToRef) +{ + set_up_apply_data(); + auto dense_mtx = gko::matrix::Dense<>::create(ref); + auto ddense_mtx = gko::matrix::Dense<>::create(omp); + + mtx->convert_to(ddense_mtx.get()); + dmtx->convert_to(ddense_mtx.get()); + + GKO_ASSERT_MTX_NEAR(dense_mtx.get(), dense_mtx.get(), 1e-14); +} + + +TEST_F(Csr, MoveToDenseIsEquivalentToRef) +{ + set_up_apply_data(); + auto dense_mtx = gko::matrix::Dense<>::create(ref); + auto ddense_mtx = gko::matrix::Dense<>::create(omp); + + mtx->move_to(ddense_mtx.get()); + dmtx->move_to(ddense_mtx.get()); + + GKO_ASSERT_MTX_NEAR(dense_mtx.get(), dense_mtx.get(), 1e-14); +} + + } // namespace diff --git a/omp/test/matrix/dense_kernels.cpp b/omp/test/matrix/dense_kernels.cpp index b88d9139cff..b0f0b75170b 100644 --- a/omp/test/matrix/dense_kernels.cpp +++ b/omp/test/matrix/dense_kernels.cpp @@ -30,17 +30,16 @@ 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 "core/matrix/dense_kernels.hpp" #include #include -#include +#include + + #include #include #include @@ -50,7 +49,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include -#include "core/matrix/dense_kernels.hpp" +#include "core/test/utils.hpp" namespace { @@ -285,19 +284,19 @@ TEST_F(Dense, ConvertToCooIsEquivalentToRef) auto rmtx = gen_mtx(532, 231); auto omtx = Mtx::create(omp); omtx->copy_from(rmtx.get()); - auto srmtx = gko::matrix::Coo<>::create(ref); auto somtx = gko::matrix::Coo<>::create(omp); + auto drmtx = Mtx::create(ref); + auto domtx = Mtx::create(omp); rmtx->convert_to(srmtx.get()); omtx->convert_to(somtx.get()); - - auto drmtx = Mtx::create(ref); - auto domtx = Mtx::create(omp); srmtx->convert_to(drmtx.get()); somtx->convert_to(domtx.get()); GKO_ASSERT_MTX_NEAR(drmtx, domtx, 1e-14); + GKO_ASSERT_MTX_NEAR(srmtx, somtx, 1e-14); + GKO_ASSERT_MTX_NEAR(domtx, omtx, 1e-14); } @@ -306,19 +305,19 @@ TEST_F(Dense, MoveToCooIsEquivalentToRef) auto rmtx = gen_mtx(532, 231); auto omtx = Mtx::create(omp); omtx->copy_from(rmtx.get()); - auto srmtx = gko::matrix::Coo<>::create(ref); auto somtx = gko::matrix::Coo<>::create(omp); + auto drmtx = Mtx::create(ref); + auto domtx = Mtx::create(omp); rmtx->move_to(srmtx.get()); omtx->move_to(somtx.get()); - - auto drmtx = Mtx::create(ref); - auto domtx = Mtx::create(omp); srmtx->move_to(drmtx.get()); somtx->move_to(domtx.get()); GKO_ASSERT_MTX_NEAR(drmtx, domtx, 1e-14); + GKO_ASSERT_MTX_NEAR(srmtx, somtx, 1e-14); + GKO_ASSERT_MTX_NEAR(domtx, omtx, 1e-14); } @@ -327,19 +326,19 @@ TEST_F(Dense, ConvertToCsrIsEquivalentToRef) auto rmtx = gen_mtx(532, 231); auto omtx = Mtx::create(omp); omtx->copy_from(rmtx.get()); - auto srmtx = gko::matrix::Csr<>::create(ref); auto somtx = gko::matrix::Csr<>::create(omp); + auto drmtx = Mtx::create(ref); + auto domtx = Mtx::create(omp); rmtx->convert_to(srmtx.get()); omtx->convert_to(somtx.get()); - - auto drmtx = Mtx::create(ref); - auto domtx = Mtx::create(omp); srmtx->convert_to(drmtx.get()); somtx->convert_to(domtx.get()); GKO_ASSERT_MTX_NEAR(drmtx, domtx, 1e-14); + GKO_ASSERT_MTX_NEAR(srmtx, somtx, 1e-14); + GKO_ASSERT_MTX_NEAR(domtx, omtx, 1e-14); } @@ -348,19 +347,19 @@ TEST_F(Dense, MoveToCsrIsEquivalentToRef) auto rmtx = gen_mtx(532, 231); auto omtx = Mtx::create(omp); omtx->copy_from(rmtx.get()); - auto srmtx = gko::matrix::Csr<>::create(ref); auto somtx = gko::matrix::Csr<>::create(omp); + auto drmtx = Mtx::create(ref); + auto domtx = Mtx::create(omp); rmtx->move_to(srmtx.get()); omtx->move_to(somtx.get()); - - auto drmtx = Mtx::create(ref); - auto domtx = Mtx::create(omp); srmtx->move_to(drmtx.get()); somtx->move_to(domtx.get()); GKO_ASSERT_MTX_NEAR(drmtx, domtx, 1e-14); + GKO_ASSERT_MTX_NEAR(srmtx, somtx, 1e-14); + GKO_ASSERT_MTX_NEAR(domtx, omtx, 1e-14); } @@ -369,19 +368,19 @@ TEST_F(Dense, ConvertToEllIsEquivalentToRef) auto rmtx = gen_mtx(532, 231); auto omtx = Mtx::create(omp); omtx->copy_from(rmtx.get()); - auto srmtx = gko::matrix::Ell<>::create(ref); auto somtx = gko::matrix::Ell<>::create(omp); + auto drmtx = Mtx::create(ref); + auto domtx = Mtx::create(omp); rmtx->convert_to(srmtx.get()); omtx->convert_to(somtx.get()); - - auto drmtx = Mtx::create(ref); - auto domtx = Mtx::create(omp); srmtx->convert_to(drmtx.get()); somtx->convert_to(domtx.get()); GKO_ASSERT_MTX_NEAR(drmtx, domtx, 1e-14); + GKO_ASSERT_MTX_NEAR(srmtx, somtx, 1e-14); + GKO_ASSERT_MTX_NEAR(domtx, omtx, 1e-14); } @@ -390,19 +389,19 @@ TEST_F(Dense, MoveToEllIsEquivalentToRef) auto rmtx = gen_mtx(532, 231); auto omtx = Mtx::create(omp); omtx->copy_from(rmtx.get()); - auto srmtx = gko::matrix::Ell<>::create(ref); auto somtx = gko::matrix::Ell<>::create(omp); + auto drmtx = Mtx::create(ref); + auto domtx = Mtx::create(omp); rmtx->move_to(srmtx.get()); omtx->move_to(somtx.get()); - - auto drmtx = Mtx::create(ref); - auto domtx = Mtx::create(omp); srmtx->move_to(drmtx.get()); somtx->move_to(domtx.get()); GKO_ASSERT_MTX_NEAR(drmtx, domtx, 1e-14); + GKO_ASSERT_MTX_NEAR(srmtx, somtx, 1e-14); + GKO_ASSERT_MTX_NEAR(domtx, omtx, 1e-14); } @@ -411,19 +410,21 @@ TEST_F(Dense, ConvertToHybridIsEquivalentToRef) auto rmtx = gen_mtx(532, 231); auto omtx = Mtx::create(omp); omtx->copy_from(rmtx.get()); - auto srmtx = gko::matrix::Hybrid<>::create(ref); auto somtx = gko::matrix::Hybrid<>::create(omp); + auto drmtx = Mtx::create(ref); + auto domtx = Mtx::create(omp); rmtx->convert_to(srmtx.get()); omtx->convert_to(somtx.get()); - - auto drmtx = Mtx::create(ref); - auto domtx = Mtx::create(omp); srmtx->convert_to(drmtx.get()); somtx->convert_to(domtx.get()); GKO_ASSERT_MTX_NEAR(drmtx, domtx, 1e-14); + // Test between `srmtx` and `somtx` may fail due to the OpenMP + // implementation not sorting the Coo matrix part. + // Therefore, it is not performed. + GKO_ASSERT_MTX_NEAR(domtx, omtx, 1e-14); } @@ -432,19 +433,21 @@ TEST_F(Dense, MoveToHybridIsEquivalentToRef) auto rmtx = gen_mtx(532, 231); auto omtx = Mtx::create(omp); omtx->copy_from(rmtx.get()); - auto srmtx = gko::matrix::Hybrid<>::create(ref); auto somtx = gko::matrix::Hybrid<>::create(omp); + auto drmtx = Mtx::create(ref); + auto domtx = Mtx::create(omp); rmtx->move_to(srmtx.get()); omtx->move_to(somtx.get()); - - auto drmtx = Mtx::create(ref); - auto domtx = Mtx::create(omp); srmtx->move_to(drmtx.get()); somtx->move_to(domtx.get()); GKO_ASSERT_MTX_NEAR(drmtx, domtx, 1e-14); + // Test between `srmtx` and `somtx` may fail due to the OpenMP + // implementation not sorting the Coo matrix part. + // Therefore, it is not performed. + GKO_ASSERT_MTX_NEAR(domtx, omtx, 1e-14); } @@ -453,19 +456,19 @@ TEST_F(Dense, ConvertToSellpIsEquivalentToRef) auto rmtx = gen_mtx(532, 231); auto omtx = Mtx::create(omp); omtx->copy_from(rmtx.get()); - auto srmtx = gko::matrix::Sellp<>::create(ref); auto somtx = gko::matrix::Sellp<>::create(omp); + auto drmtx = Mtx::create(ref); + auto domtx = Mtx::create(omp); rmtx->convert_to(srmtx.get()); omtx->convert_to(somtx.get()); - - auto drmtx = Mtx::create(ref); - auto domtx = Mtx::create(omp); srmtx->convert_to(drmtx.get()); somtx->convert_to(domtx.get()); GKO_ASSERT_MTX_NEAR(drmtx, domtx, 1e-14); + GKO_ASSERT_MTX_NEAR(srmtx, somtx, 1e-14); + GKO_ASSERT_MTX_NEAR(domtx, omtx, 1e-14); } @@ -474,19 +477,19 @@ TEST_F(Dense, MoveToSellpIsEquivalentToRef) auto rmtx = gen_mtx(532, 231); auto omtx = Mtx::create(omp); omtx->copy_from(rmtx.get()); - auto srmtx = gko::matrix::Sellp<>::create(ref); auto somtx = gko::matrix::Sellp<>::create(omp); + auto drmtx = Mtx::create(ref); + auto domtx = Mtx::create(omp); rmtx->move_to(srmtx.get()); omtx->move_to(somtx.get()); - - auto drmtx = Mtx::create(ref); - auto domtx = Mtx::create(omp); srmtx->move_to(drmtx.get()); somtx->move_to(domtx.get()); GKO_ASSERT_MTX_NEAR(drmtx, domtx, 1e-14); + GKO_ASSERT_MTX_NEAR(srmtx, somtx, 1e-14); + GKO_ASSERT_MTX_NEAR(domtx, omtx, 1e-14); } @@ -494,7 +497,6 @@ TEST_F(Dense, CalculateMaxNNZPerRowIsEquivalentToRef) { std::size_t ref_max_nnz_per_row = 0; std::size_t omp_max_nnz_per_row = 0; - auto rmtx = gen_mtx(100, 100, 1); auto omtx = Mtx::create(omp); omtx->copy_from(rmtx.get()); @@ -512,7 +514,6 @@ TEST_F(Dense, CalculateTotalColsIsEquivalentToRef) { std::size_t ref_total_cols = 0; std::size_t omp_total_cols = 0; - auto rmtx = gen_mtx(100, 100, 1); auto omtx = Mtx::create(omp); omtx->copy_from(rmtx.get()); diff --git a/reference/matrix/coo_kernels.cpp b/reference/matrix/coo_kernels.cpp index b347bc7fe54..76f15e45766 100644 --- a/reference/matrix/coo_kernels.cpp +++ b/reference/matrix/coo_kernels.cpp @@ -143,25 +143,6 @@ void convert_row_idxs_to_ptrs(std::shared_ptr exec, } -template -void transpose(std::shared_ptr exec, - matrix::Coo *trans, - const matrix::Coo *orig) - GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_COO_TRANSPOSE_KERNEL); - - -template -void conj_transpose(std::shared_ptr exec, - matrix::Coo *trans, - const matrix::Coo *orig) - GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_COO_CONJ_TRANSPOSE_KERNEL); - - template void convert_to_csr(std::shared_ptr exec, matrix::Csr *result, diff --git a/reference/matrix/csr_kernels.cpp b/reference/matrix/csr_kernels.cpp index e525d824c15..e7107f4a3a6 100644 --- a/reference/matrix/csr_kernels.cpp +++ b/reference/matrix/csr_kernels.cpp @@ -33,6 +33,12 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "core/matrix/csr_kernels.hpp" +#include +#include +#include +#include + + #include #include #include @@ -44,12 +50,6 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "reference/components/format_conversion.hpp" -#include -#include -#include -#include - - namespace gko { namespace kernels { namespace reference { @@ -172,18 +172,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_CSR_CONVERT_TO_DENSE_KERNEL); -template -void move_to_dense(std::shared_ptr exec, - matrix::Dense *result, - matrix::Csr *source) -{ - reference::csr::convert_to_dense(exec, result, source); -} - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_CSR_MOVE_TO_DENSE_KERNEL); - - template void convert_to_sellp(std::shared_ptr exec, matrix::Sellp *result, diff --git a/reference/matrix/dense_kernels.cpp b/reference/matrix/dense_kernels.cpp index 3c6ad1269c3..ba975b720c9 100644 --- a/reference/matrix/dense_kernels.cpp +++ b/reference/matrix/dense_kernels.cpp @@ -256,18 +256,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_CSR_KERNEL); -template -void move_to_csr(std::shared_ptr exec, - matrix::Csr *result, - const matrix::Dense *source) -{ - reference::dense::convert_to_csr(exec, result, source); -} - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_CSR_KERNEL); - - template void convert_to_ell(std::shared_ptr exec, matrix::Ell *result, @@ -300,18 +288,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_ELL_KERNEL); -template -void move_to_ell(std::shared_ptr exec, - matrix::Ell *result, - const matrix::Dense *source) -{ - reference::dense::convert_to_ell(exec, result, source); -} - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_ELL_KERNEL); - - template void convert_to_hybrid(std::shared_ptr exec, matrix::Hybrid *result, @@ -367,18 +343,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_HYBRID_KERNEL); -template -void move_to_hybrid(std::shared_ptr exec, - matrix::Hybrid *result, - const matrix::Dense *source) -{ - reference::dense::convert_to_hybrid(exec, result, source); -} - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_HYBRID_KERNEL); - - template void convert_to_sellp(std::shared_ptr exec, matrix::Sellp *result, @@ -450,18 +414,6 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( GKO_DECLARE_DENSE_CONVERT_TO_SELLP_KERNEL); -template -void move_to_sellp(std::shared_ptr exec, - matrix::Sellp *result, - const matrix::Dense *source) -{ - reference::dense::convert_to_sellp(exec, result, source); -} - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_DENSE_MOVE_TO_SELLP_KERNEL); - - template void count_nonzeros(std::shared_ptr exec, const matrix::Dense *source, size_type *result) diff --git a/reference/test/matrix/coo_kernels.cpp b/reference/test/matrix/coo_kernels.cpp index 4802a0288c3..a98ce5adf78 100644 --- a/reference/test/matrix/coo_kernels.cpp +++ b/reference/test/matrix/coo_kernels.cpp @@ -30,19 +30,22 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include +#include "core/matrix/coo_kernels.hpp" #include -#include #include #include +#include #include #include +#include "core/test/utils/assertions.hpp" + + namespace { @@ -88,7 +91,9 @@ class Coo : public ::testing::Test { TEST_F(Coo, ConvertsToCsr) { auto csr_mtx = gko::matrix::Csr<>::create(mtx->get_executor()); + mtx->convert_to(csr_mtx.get()); + assert_equal_to_mtx_in_csr_format(csr_mtx.get()); } @@ -96,7 +101,9 @@ TEST_F(Coo, ConvertsToCsr) TEST_F(Coo, MovesToCsr) { auto csr_mtx = gko::matrix::Csr<>::create(mtx->get_executor()); + mtx->move_to(csr_mtx.get()); + assert_equal_to_mtx_in_csr_format(csr_mtx.get()); } @@ -228,6 +235,7 @@ TEST_F(Coo, AppliesAddToDenseVector) { auto x = gko::initialize({2.0, 1.0, 4.0}, exec); auto y = gko::initialize({2.0, 1.0}, exec); + mtx->apply2(x.get(), y.get()); GKO_ASSERT_MTX_NEAR(y, l({15.0, 6.0}), 0.0); diff --git a/reference/test/matrix/csr_kernels.cpp b/reference/test/matrix/csr_kernels.cpp index cf99b584651..c96c005dd1e 100644 --- a/reference/test/matrix/csr_kernels.cpp +++ b/reference/test/matrix/csr_kernels.cpp @@ -30,22 +30,22 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include +#include "core/matrix/csr_kernels.hpp" #include -#include #include #include #include +#include #include -#include #include +#include -#include "core/matrix/csr_kernels.hpp" +#include "core/test/utils/assertions.hpp" namespace { @@ -273,7 +273,19 @@ TEST_F(Csr, MovesToDense) TEST_F(Csr, ConvertsToCoo) { auto coo_mtx = gko::matrix::Coo<>::create(mtx->get_executor()); + mtx->convert_to(coo_mtx.get()); + + assert_equal_to_mtx(coo_mtx.get()); +} + + +TEST_F(Csr, MovesToCoo) +{ + auto coo_mtx = gko::matrix::Coo<>::create(mtx->get_executor()); + + mtx->move_to(coo_mtx.get()); + assert_equal_to_mtx(coo_mtx.get()); } @@ -281,7 +293,9 @@ TEST_F(Csr, ConvertsToCoo) TEST_F(Csr, ConvertsToSellp) { auto sellp_mtx = gko::matrix::Sellp<>::create(mtx->get_executor()); + mtx->convert_to(sellp_mtx.get()); + assert_equal_to_mtx(sellp_mtx.get()); } @@ -311,21 +325,26 @@ TEST_F(Csr, CalculatesTotalCols) } -TEST_F(Csr, MovesToCoo) +TEST_F(Csr, ConvertsToEll) { - auto coo_mtx = gko::matrix::Coo<>::create(mtx->get_executor()); - mtx->move_to(coo_mtx.get()); - assert_equal_to_mtx(coo_mtx.get()); + auto ell_mtx = gko::matrix::Ell<>::create(mtx->get_executor()); + auto dense_mtx = gko::matrix::Dense<>::create(mtx->get_executor()); + auto ref_dense_mtx = gko::matrix::Dense<>::create(mtx->get_executor()); + + mtx->convert_to(ell_mtx.get()); + + assert_equal_to_mtx(ell_mtx.get()); } -TEST_F(Csr, ConvertsToEll) +TEST_F(Csr, MovesToEll) { auto ell_mtx = gko::matrix::Ell<>::create(mtx->get_executor()); auto dense_mtx = gko::matrix::Dense<>::create(mtx->get_executor()); auto ref_dense_mtx = gko::matrix::Dense<>::create(mtx->get_executor()); - mtx->convert_to(ell_mtx.get()); + mtx->move_to(ell_mtx.get()); + assert_equal_to_mtx(ell_mtx.get()); } diff --git a/reference/test/matrix/dense_kernels.cpp b/reference/test/matrix/dense_kernels.cpp index c65f2d1cbbb..892a557a973 100644 --- a/reference/test/matrix/dense_kernels.cpp +++ b/reference/test/matrix/dense_kernels.cpp @@ -30,31 +30,29 @@ THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. *************************************************************/ -#include +#include "core/matrix/dense_kernels.hpp" #include - - -#include - - #include -#include +#include -#include #include #include #include #include +#include #include #include #include +#include "core/test/utils.hpp" + + namespace { @@ -285,7 +283,6 @@ TEST_F(Dense, ConvertsToCoo) auto coo_mtx = gko::matrix::Coo<>::create(mtx4->get_executor()); mtx4->convert_to(coo_mtx.get()); - auto v = coo_mtx->get_const_values(); auto c = coo_mtx->get_const_col_idxs(); auto r = coo_mtx->get_const_row_idxs(); @@ -312,7 +309,6 @@ TEST_F(Dense, MovesToCoo) auto coo_mtx = gko::matrix::Coo<>::create(mtx4->get_executor()); mtx4->move_to(coo_mtx.get()); - auto v = coo_mtx->get_const_values(); auto c = coo_mtx->get_const_col_idxs(); auto r = coo_mtx->get_const_row_idxs(); @@ -339,7 +335,6 @@ TEST_F(Dense, ConvertsToCsr) auto csr_mtx = gko::matrix::Csr<>::create(mtx4->get_executor()); mtx4->convert_to(csr_mtx.get()); - auto v = csr_mtx->get_const_values(); auto c = csr_mtx->get_const_col_idxs(); auto r = csr_mtx->get_const_row_ptrs(); @@ -365,7 +360,6 @@ TEST_F(Dense, MovesToCsr) auto csr_mtx = gko::matrix::Csr<>::create(mtx4->get_executor()); mtx4->move_to(csr_mtx.get()); - auto v = csr_mtx->get_const_values(); auto c = csr_mtx->get_const_col_idxs(); auto r = csr_mtx->get_const_row_ptrs(); @@ -391,7 +385,6 @@ TEST_F(Dense, ConvertsToEll) auto ell_mtx = gko::matrix::Ell<>::create(mtx7->get_executor()); mtx7->convert_to(ell_mtx.get()); - auto v = ell_mtx->get_const_values(); auto c = ell_mtx->get_const_col_idxs(); @@ -415,7 +408,6 @@ TEST_F(Dense, MovesToEll) auto ell_mtx = gko::matrix::Ell<>::create(mtx7->get_executor()); mtx7->move_to(ell_mtx.get()); - auto v = ell_mtx->get_const_values(); auto c = ell_mtx->get_const_col_idxs(); @@ -440,7 +432,6 @@ TEST_F(Dense, ConvertsToEllWithStride) gko::matrix::Ell<>::create(mtx7->get_executor(), gko::dim<2>{}, 0, 3); mtx7->convert_to(ell_mtx.get()); - auto v = ell_mtx->get_const_values(); auto c = ell_mtx->get_const_col_idxs(); @@ -469,7 +460,6 @@ TEST_F(Dense, MovesToEllWithStride) gko::matrix::Ell<>::create(mtx7->get_executor(), gko::dim<2>{}, 0, 3); mtx7->move_to(ell_mtx.get()); - auto v = ell_mtx->get_const_values(); auto c = ell_mtx->get_const_col_idxs(); @@ -497,12 +487,12 @@ TEST_F(Dense, MovesToHybridAutomatically) auto hybrid_mtx = gko::matrix::Hybrid<>::create(mtx4->get_executor()); mtx4->move_to(hybrid_mtx.get()); - auto v = hybrid_mtx->get_const_coo_values(); auto c = hybrid_mtx->get_const_coo_col_idxs(); auto r = hybrid_mtx->get_const_coo_row_idxs(); auto n = hybrid_mtx->get_ell_num_stored_elements_per_row(); auto p = hybrid_mtx->get_ell_stride(); + ASSERT_EQ(hybrid_mtx->get_size(), gko::dim<2>(2, 3)); ASSERT_EQ(hybrid_mtx->get_ell_num_stored_elements(), 0); ASSERT_EQ(hybrid_mtx->get_coo_num_stored_elements(), 4); @@ -528,12 +518,12 @@ TEST_F(Dense, ConvertsToHybridAutomatically) auto hybrid_mtx = gko::matrix::Hybrid<>::create(mtx4->get_executor()); mtx4->convert_to(hybrid_mtx.get()); - auto v = hybrid_mtx->get_const_coo_values(); auto c = hybrid_mtx->get_const_coo_col_idxs(); auto r = hybrid_mtx->get_const_coo_row_idxs(); auto n = hybrid_mtx->get_ell_num_stored_elements_per_row(); auto p = hybrid_mtx->get_ell_stride(); + ASSERT_EQ(hybrid_mtx->get_size(), gko::dim<2>(2, 3)); ASSERT_EQ(hybrid_mtx->get_ell_num_stored_elements(), 0); ASSERT_EQ(hybrid_mtx->get_coo_num_stored_elements(), 4); @@ -560,12 +550,12 @@ TEST_F(Dense, MovesToHybridWithStrideAutomatically) gko::dim<2>{}, 0, 3); mtx4->move_to(hybrid_mtx.get()); - auto v = hybrid_mtx->get_const_coo_values(); auto c = hybrid_mtx->get_const_coo_col_idxs(); auto r = hybrid_mtx->get_const_coo_row_idxs(); auto n = hybrid_mtx->get_ell_num_stored_elements_per_row(); auto p = hybrid_mtx->get_ell_stride(); + ASSERT_EQ(hybrid_mtx->get_size(), gko::dim<2>(2, 3)); ASSERT_EQ(hybrid_mtx->get_ell_num_stored_elements(), 0); ASSERT_EQ(hybrid_mtx->get_coo_num_stored_elements(), 4); @@ -592,12 +582,12 @@ TEST_F(Dense, ConvertsToHybridWithStrideAutomatically) gko::dim<2>{}, 0, 3); mtx4->convert_to(hybrid_mtx.get()); - auto v = hybrid_mtx->get_const_coo_values(); auto c = hybrid_mtx->get_const_coo_col_idxs(); auto r = hybrid_mtx->get_const_coo_row_idxs(); auto n = hybrid_mtx->get_ell_num_stored_elements_per_row(); auto p = hybrid_mtx->get_ell_stride(); + ASSERT_EQ(hybrid_mtx->get_size(), gko::dim<2>(2, 3)); ASSERT_EQ(hybrid_mtx->get_ell_num_stored_elements(), 0); ASSERT_EQ(hybrid_mtx->get_coo_num_stored_elements(), 4); @@ -625,11 +615,11 @@ TEST_F(Dense, MovesToHybridWithStrideAndCooLengthByColumns2) std::make_shared::column_limit>(2)); mtx4->move_to(hybrid_mtx.get()); - auto v = hybrid_mtx->get_const_ell_values(); auto c = hybrid_mtx->get_const_ell_col_idxs(); auto n = hybrid_mtx->get_ell_num_stored_elements_per_row(); auto p = hybrid_mtx->get_ell_stride(); + ASSERT_EQ(hybrid_mtx->get_size(), gko::dim<2>(2, 3)); ASSERT_EQ(hybrid_mtx->get_ell_num_stored_elements(), 6); ASSERT_EQ(hybrid_mtx->get_coo_num_stored_elements(), 3); @@ -666,11 +656,11 @@ TEST_F(Dense, ConvertsToHybridWithStrideAndCooLengthByColumns2) std::make_shared::column_limit>(2)); mtx4->convert_to(hybrid_mtx.get()); - auto v = hybrid_mtx->get_const_ell_values(); auto c = hybrid_mtx->get_const_ell_col_idxs(); auto n = hybrid_mtx->get_ell_num_stored_elements_per_row(); auto p = hybrid_mtx->get_ell_stride(); + ASSERT_EQ(hybrid_mtx->get_size(), gko::dim<2>(2, 3)); ASSERT_EQ(hybrid_mtx->get_ell_num_stored_elements(), 6); ASSERT_EQ(hybrid_mtx->get_coo_num_stored_elements(), 3); @@ -707,11 +697,14 @@ TEST_F(Dense, MovesToHybridWithStrideByPercent40) std::make_shared::imbalance_limit>(0.4)); mtx4->move_to(hybrid_mtx.get()); - auto v = hybrid_mtx->get_const_ell_values(); auto c = hybrid_mtx->get_const_ell_col_idxs(); auto n = hybrid_mtx->get_ell_num_stored_elements_per_row(); auto p = hybrid_mtx->get_ell_stride(); + auto coo_v = hybrid_mtx->get_const_coo_values(); + auto coo_c = hybrid_mtx->get_const_coo_col_idxs(); + auto coo_r = hybrid_mtx->get_const_coo_row_idxs(); + ASSERT_EQ(hybrid_mtx->get_size(), gko::dim<2>(2, 3)); ASSERT_EQ(hybrid_mtx->get_ell_num_stored_elements(), 3); EXPECT_EQ(n, 1); @@ -722,10 +715,6 @@ TEST_F(Dense, MovesToHybridWithStrideByPercent40) EXPECT_EQ(v[0], 1.0); EXPECT_EQ(v[1], 5.0); EXPECT_EQ(v[2], 0.0); - - auto coo_v = hybrid_mtx->get_const_coo_values(); - auto coo_c = hybrid_mtx->get_const_coo_col_idxs(); - auto coo_r = hybrid_mtx->get_const_coo_row_idxs(); ASSERT_EQ(hybrid_mtx->get_coo_num_stored_elements(), 2); EXPECT_EQ(coo_v[0], 3.0); EXPECT_EQ(coo_v[1], 2.0); @@ -743,11 +732,14 @@ TEST_F(Dense, ConvertsToHybridWithStrideByPercent40) std::make_shared::imbalance_limit>(0.4)); mtx4->convert_to(hybrid_mtx.get()); - auto v = hybrid_mtx->get_const_ell_values(); auto c = hybrid_mtx->get_const_ell_col_idxs(); auto n = hybrid_mtx->get_ell_num_stored_elements_per_row(); auto p = hybrid_mtx->get_ell_stride(); + auto coo_v = hybrid_mtx->get_const_coo_values(); + auto coo_c = hybrid_mtx->get_const_coo_col_idxs(); + auto coo_r = hybrid_mtx->get_const_coo_row_idxs(); + ASSERT_EQ(hybrid_mtx->get_size(), gko::dim<2>(2, 3)); ASSERT_EQ(hybrid_mtx->get_ell_num_stored_elements(), 3); EXPECT_EQ(n, 1); @@ -758,10 +750,6 @@ TEST_F(Dense, ConvertsToHybridWithStrideByPercent40) EXPECT_EQ(v[0], 1.0); EXPECT_EQ(v[1], 5.0); EXPECT_EQ(v[2], 0.0); - - auto coo_v = hybrid_mtx->get_const_coo_values(); - auto coo_c = hybrid_mtx->get_const_coo_col_idxs(); - auto coo_r = hybrid_mtx->get_const_coo_row_idxs(); ASSERT_EQ(hybrid_mtx->get_coo_num_stored_elements(), 2); EXPECT_EQ(coo_v[0], 3.0); EXPECT_EQ(coo_v[1], 2.0); @@ -777,7 +765,6 @@ TEST_F(Dense, ConvertsToSellp) auto sellp_mtx = gko::matrix::Sellp<>::create(mtx8->get_executor()); mtx8->convert_to(sellp_mtx.get()); - auto v = sellp_mtx->get_const_values(); auto c = sellp_mtx->get_const_col_idxs(); auto s = sellp_mtx->get_const_slice_sets(); @@ -813,7 +800,6 @@ TEST_F(Dense, MovesToSellp) auto sellp_mtx = gko::matrix::Sellp<>::create(mtx8->get_executor()); mtx8->move_to(sellp_mtx.get()); - auto v = sellp_mtx->get_const_values(); auto c = sellp_mtx->get_const_col_idxs(); auto s = sellp_mtx->get_const_slice_sets(); @@ -850,7 +836,6 @@ TEST_F(Dense, ConvertsToSellpWithSliceSizeAndStrideFactor) gko::dim<2>{}, 2, 2, 0); mtx8->convert_to(sellp_mtx.get()); - auto v = sellp_mtx->get_const_values(); auto c = sellp_mtx->get_const_col_idxs(); auto s = sellp_mtx->get_const_slice_sets(); @@ -889,7 +874,6 @@ TEST_F(Dense, MovesToSellpWithSliceSizeAndStrideFactor) gko::dim<2>{}, 2, 2, 0); mtx8->move_to(sellp_mtx.get()); - auto v = sellp_mtx->get_const_values(); auto c = sellp_mtx->get_const_col_idxs(); auto s = sellp_mtx->get_const_slice_sets(); @@ -961,7 +945,6 @@ TEST_F(Dense, ConvertsToAndFromSellpWithMoreThanOneSlice) auto sellp_mtx = gko::matrix::Sellp<>::create(exec); auto dense_mtx = gko::matrix::Dense<>::create(exec); - x->convert_to(sellp_mtx.get()); sellp_mtx->convert_to(dense_mtx.get());