Review update

+ Update some forgotten names. + Improve some kernels. + Add some dense to sparsity csr conversion tests.
ginkgo-project · Oct 8, 2019 · 745eb45 · 745eb45
1 parent b8bd477
commit 745eb45
Show file tree

Hide file tree

Showing 16 changed files with 128 additions and 52 deletions.
diff --git a/core/matrix/csr.cpp b/core/matrix/csr.cpp
@@ -197,7 +197,11 @@ void Csr<ValueType, IndexType>::convert_to(
  exec, this->get_size(), this->get_num_stored_elements());
  tmp->col_idxs_ = this->col_idxs_;
  tmp->row_ptrs_ = this->row_ptrs_;
- tmp->value_ = gko::Array<ValueType>(exec, {one<ValueType>()});
+ if (result->value_.get_data()) {
+ tmp->value_ = result->value_;
+ } else {
+ tmp->value_ = gko::Array<ValueType>(exec, {one<ValueType>()});
+ }
  tmp->move_to(result);
 }
 

diff --git a/core/matrix/dense.cpp b/core/matrix/dense.cpp
@@ -76,7 +76,7 @@ GKO_REGISTER_OPERATION(convert_to_csr, dense::convert_to_csr);
 GKO_REGISTER_OPERATION(convert_to_ell, dense::convert_to_ell);
 GKO_REGISTER_OPERATION(convert_to_hybrid, dense::convert_to_hybrid);
 GKO_REGISTER_OPERATION(convert_to_sellp, dense::convert_to_sellp);
-GKO_REGISTER_OPERATION(convert_to_sparsity, dense::convert_to_sparsity);
+GKO_REGISTER_OPERATION(convert_to_sparsity_csr, dense::convert_to_sparsity_csr);
 
 
 } // namespace dense
@@ -444,10 +444,9 @@ void Dense<ValueType>::move_to(Sellp<ValueType, int64> *result)
 template <typename ValueType>
 void Dense<ValueType>::convert_to(SparsityCsr<ValueType, int32> *result) const
 {
- conversion_helper(
- result, this,
- dense::template make_convert_to_sparsity<decltype(result),
- const Dense<ValueType> *&>);
+ conversion_helper(result, this,
+ dense::template make_convert_to_sparsity_csr<
+ decltype(result), const Dense<ValueType> *&>);
 }
 
 
@@ -461,10 +460,9 @@ void Dense<ValueType>::move_to(SparsityCsr<ValueType, int32> *result)
 template <typename ValueType>
 void Dense<ValueType>::convert_to(SparsityCsr<ValueType, int64> *result) const
 {
- conversion_helper(
- result, this,
- dense::template make_convert_to_sparsity<decltype(result),
- const Dense<ValueType> *&>);
+ conversion_helper(result, this,
+ dense::template make_convert_to_sparsity_csr<
+ decltype(result), const Dense<ValueType> *&>);
 }
 
 

diff --git a/core/matrix/dense_kernels.hpp b/core/matrix/dense_kernels.hpp
@@ -98,10 +98,10 @@ namespace kernels {
  matrix::Sellp<_type, _prec> *other, \
  const matrix::Dense<_type> *source)
 
-#define GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL(_type, _prec) \
- void convert_to_sparsity(std::shared_ptr<const DefaultExecutor> exec, \
- matrix::SparsityCsr<_type, _prec> *other, \
- const matrix::Dense<_type> *source)
+#define GKO_DECLARE_DENSE_CONVERT_TO_SPARSITY_CSR_KERNEL(_type, _prec)  \
+ void convert_to_sparsity_csr(std::shared_ptr<const DefaultExecutor> 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<const DefaultExecutor> exec, \

diff --git a/core/matrix/sparsity_csr.cpp b/core/matrix/sparsity_csr.cpp
@@ -171,7 +171,7 @@ SparsityCsr<ValueType, IndexType>::to_adjacency_matrix() const
  GKO_ASSERT_IS_SQUARE_MATRIX(this);
  size_type num_diagonal_elements = 0;
  exec->run(sparsity_csr::make_count_num_diagonal_elements(
- this, num_diagonal_elements));
+ this, &num_diagonal_elements));
  auto adj_mat =
  SparsityCsr::create(exec, this->get_size(),
  this->get_num_nonzeros() - num_diagonal_elements);

diff --git a/core/matrix/sparsity_csr_kernels.hpp b/core/matrix/sparsity_csr_kernels.hpp
@@ -68,7 +68,7 @@ namespace kernels {
  void count_num_diagonal_elements( \
  std::shared_ptr<const DefaultExecutor> exec, \
  const matrix::SparsityCsr<ValueType, IndexType> *matrix, \
- size_type &num_diagonal_elements)
+ size_type *num_diagonal_elements)
 
 #define GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL(ValueType, IndexType) \
  void transpose(std::shared_ptr<const DefaultExecutor> exec, \

diff --git a/cuda/matrix/dense_kernels.cu b/cuda/matrix/dense_kernels.cu
@@ -732,9 +732,9 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
 
 
 template <typename ValueType, typename IndexType>
-void convert_to_sparsity(std::shared_ptr<const CudaExecutor> exec,
- matrix::SparsityCsr<ValueType, IndexType> *result,
- const matrix::Dense<ValueType> *source)
+void convert_to_sparsity_csr(std::shared_ptr<const CudaExecutor> exec,
+  matrix::SparsityCsr<ValueType, IndexType> *result,
+  const matrix::Dense<ValueType> *source)
  GKO_NOT_IMPLEMENTED;
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(

diff --git a/cuda/matrix/sparsity_csr_kernels.cu b/cuda/matrix/sparsity_csr_kernels.cu
@@ -92,7 +92,7 @@ template <typename ValueType, typename IndexType>
 void count_num_diagonal_elements(
  std::shared_ptr<const CudaExecutor> exec,
  const matrix::SparsityCsr<ValueType, IndexType> *matrix,
- size_type &num_diagonal_elements) GKO_NOT_IMPLEMENTED;
+ size_type *num_diagonal_elements) GKO_NOT_IMPLEMENTED;
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
  GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL);

diff --git a/omp/matrix/dense_kernels.cpp b/omp/matrix/dense_kernels.cpp
@@ -457,9 +457,9 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
 
 
 template <typename ValueType, typename IndexType>
-void convert_to_sparsity(std::shared_ptr<const OmpExecutor> exec,
- matrix::SparsityCsr<ValueType, IndexType> *result,
- const matrix::Dense<ValueType> *source)
+void convert_to_sparsity_csr(std::shared_ptr<const OmpExecutor> exec,
+  matrix::SparsityCsr<ValueType, IndexType> *result,
+  const matrix::Dense<ValueType> *source)
 {
  auto num_rows = result->get_size()[0];
  auto num_cols = result->get_size()[1];

diff --git a/omp/matrix/sparsity_csr_kernels.cpp b/omp/matrix/sparsity_csr_kernels.cpp
@@ -68,6 +68,7 @@ void spmv(std::shared_ptr<const OmpExecutor> exec,
 {
  auto row_ptrs = a->get_const_row_ptrs();
  auto col_idxs = a->get_const_col_idxs();
+ auto val = a->get_const_value()[0];
 
 #pragma omp parallel for
  for (size_type row = 0; row < a->get_size()[0]; ++row) {
@@ -76,7 +77,6 @@ void spmv(std::shared_ptr<const OmpExecutor> exec,
  }
  for (size_type k = row_ptrs[row];
  k < static_cast<size_type>(row_ptrs[row + 1]); ++k) {
- auto val = one<ValueType>();
  auto col = col_idxs[k];
  for (size_type j = 0; j < c->get_size()[1]; ++j) {
  c->at(row, j) += val * b->at(col, j);
@@ -101,6 +101,7 @@ void advanced_spmv(std::shared_ptr<const OmpExecutor> exec,
  auto col_idxs = a->get_const_col_idxs();
  auto valpha = alpha->at(0, 0);
  auto vbeta = beta->at(0, 0);
+ auto val = a->get_const_value()[0];
 
 #pragma omp parallel for
  for (size_type row = 0; row < a->get_size()[0]; ++row) {
@@ -109,7 +110,6 @@ void advanced_spmv(std::shared_ptr<const OmpExecutor> exec,
  }
  for (size_type k = row_ptrs[row];
  k < static_cast<size_type>(row_ptrs[row + 1]); ++k) {
- auto val = one<ValueType>();
  auto col = col_idxs[k];
  for (size_type j = 0; j < c->get_size()[1]; ++j) {
  c->at(row, j) += valpha * val * b->at(col, j);
@@ -126,7 +126,7 @@ template <typename ValueType, typename IndexType>
 void count_num_diagonal_elements(
  std::shared_ptr<const OmpExecutor> exec,
  const matrix::SparsityCsr<ValueType, IndexType> *matrix,
- size_type &num_diagonal_elements)
+ size_type *num_diagonal_elements)
 {
  auto num_rows = matrix->get_size()[0];
  auto row_ptrs = matrix->get_const_row_ptrs();
@@ -139,7 +139,7 @@ void count_num_diagonal_elements(
  }
  }
  }
- num_diagonal_elements = num_diag;
+ *num_diagonal_elements = num_diag;
 }
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
@@ -165,18 +165,15 @@ void remove_diagonal_elements(std::shared_ptr<const OmpExecutor> exec,
  }
  adj_ptrs[i + 1] = row_ptrs[i + 1] - num_diag;
  }
- std::vector<IndexType> temp_idxs;
+ auto nnz = 0;
  for (auto i = 0; i < num_rows; ++i) {
  for (auto j = row_ptrs[i]; j < row_ptrs[i + 1]; ++j) {
  if (col_idxs[j] != i) {
- temp_idxs.push_back(col_idxs[j]);
+ adj_idxs[nnz] = col_idxs[j];
+ nnz++;
  }
  }
  }
-#pragma omp parallel for
- for (auto i = 0; i < temp_idxs.size(); ++i) {
- adj_idxs[i] = temp_idxs[i];
- }
 }
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(

diff --git a/omp/test/matrix/dense_kernels.cpp b/omp/test/matrix/dense_kernels.cpp
@@ -47,6 +47,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include <ginkgo/core/matrix/ell.hpp>
 #include <ginkgo/core/matrix/hybrid.hpp>
 #include <ginkgo/core/matrix/sellp.hpp>
+#include <ginkgo/core/matrix/sparsity_csr.hpp>
 
 
 #include "core/test/utils.hpp"
@@ -363,6 +364,36 @@ TEST_F(Dense, MoveToCsrIsEquivalentToRef)
 }
 
 
+TEST_F(Dense, ConvertToSparsityCsrIsEquivalentToRef)
+{
+ auto mtx = gen_mtx<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<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<Mtx>(532, 231);

diff --git a/omp/test/matrix/sparsity_csr_kernels.cpp b/omp/test/matrix/sparsity_csr_kernels.cpp
@@ -225,9 +225,9 @@ TEST_F(SparsityCsr, CountsNumberOfDiagElementsIsEqualToRef)
  gko::size_type d_num_diags = 0;
 
  gko::kernels::reference::sparsity_csr::count_num_diagonal_elements(
- ref, mtx.get(), num_diags);
+ ref, mtx.get(), &num_diags);
  gko::kernels::omp::sparsity_csr::count_num_diagonal_elements(
- omp, dmtx.get(), d_num_diags);
+ omp, dmtx.get(), &d_num_diags);
 
  ASSERT_EQ(d_num_diags, num_diags);
 }
@@ -238,11 +238,12 @@ TEST_F(SparsityCsr, RemovesDiagElementsKernelIsEquivalentToRef)
  set_up_apply_data();
  gko::size_type num_diags = 0;
  gko::kernels::reference::sparsity_csr::count_num_diagonal_elements(
- ref, mtx.get(), num_diags);
+ ref, mtx.get(), &num_diags);
  auto tmp =
  Mtx::create(ref, mtx->get_size(), mtx->get_num_nonzeros() - num_diags);
  auto d_tmp = Mtx::create(omp, dmtx->get_size(),
  dmtx->get_num_nonzeros() - num_diags);
+
  gko::kernels::reference::sparsity_csr::remove_diagonal_elements(
  ref, tmp.get(), mtx->get_const_row_ptrs(), mtx->get_const_col_idxs());
  gko::kernels::omp::sparsity_csr::remove_diagonal_elements(

diff --git a/reference/matrix/dense_kernels.cpp b/reference/matrix/dense_kernels.cpp
@@ -416,9 +416,9 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
 
 
 template <typename ValueType, typename IndexType>
-void convert_to_sparsity(std::shared_ptr<const ReferenceExecutor> exec,
- matrix::SparsityCsr<ValueType, IndexType> *result,
- const matrix::Dense<ValueType> *source)
+void convert_to_sparsity_csr(std::shared_ptr<const ReferenceExecutor> exec,
+  matrix::SparsityCsr<ValueType, IndexType> *result,
+  const matrix::Dense<ValueType> *source)
 {
  auto num_rows = result->get_size()[0];
  auto num_cols = result->get_size()[1];

diff --git a/reference/matrix/sparsity_csr_kernels.cpp b/reference/matrix/sparsity_csr_kernels.cpp
@@ -65,14 +65,14 @@ void spmv(std::shared_ptr<const ReferenceExecutor> exec,
 {
  auto row_ptrs = a->get_const_row_ptrs();
  auto col_idxs = a->get_const_col_idxs();
+ auto val = a->get_const_value()[0];
 
  for (size_type row = 0; row < a->get_size()[0]; ++row) {
  for (size_type j = 0; j < c->get_size()[1]; ++j) {
  c->at(row, j) = zero<ValueType>();
  }
  for (size_type k = row_ptrs[row];
  k < static_cast<size_type>(row_ptrs[row + 1]); ++k) {
- auto val = one<ValueType>();
  auto col = col_idxs[k];
  for (size_type j = 0; j < c->get_size()[1]; ++j) {
  c->at(row, j) += val * b->at(col, j);
@@ -97,14 +97,14 @@ void advanced_spmv(std::shared_ptr<const ReferenceExecutor> exec,
  auto col_idxs = a->get_const_col_idxs();
  auto valpha = alpha->at(0, 0);
  auto vbeta = beta->at(0, 0);
+ auto val = a->get_const_value()[0];
 
  for (size_type row = 0; row < a->get_size()[0]; ++row) {
  for (size_type j = 0; j < c->get_size()[1]; ++j) {
  c->at(row, j) *= vbeta;
  }
  for (size_type k = row_ptrs[row];
  k < static_cast<size_type>(row_ptrs[row + 1]); ++k) {
- auto val = one<ValueType>();
  auto col = col_idxs[k];
  for (size_type j = 0; j < c->get_size()[1]; ++j) {
  c->at(row, j) += valpha * val * b->at(col, j);
@@ -121,7 +121,7 @@ template <typename ValueType, typename IndexType>
 void count_num_diagonal_elements(
  std::shared_ptr<const ReferenceExecutor> exec,
  const matrix::SparsityCsr<ValueType, IndexType> *matrix,
- size_type &num_diagonal_elements)
+ size_type *num_diagonal_elements)
 {
  auto num_rows = matrix->get_size()[0];
  auto row_ptrs = matrix->get_const_row_ptrs();
@@ -134,7 +134,7 @@ void count_num_diagonal_elements(
  }
  }
  }
- num_diagonal_elements = num_diag;
+ *num_diagonal_elements = num_diag;
 }
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(
@@ -160,17 +160,15 @@ void remove_diagonal_elements(std::shared_ptr<const ReferenceExecutor> exec,
  }
  adj_ptrs[i + 1] = row_ptrs[i + 1] - num_diag;
  }
- std::vector<IndexType> temp_idxs;
+ auto nnz = 0;
  for (auto i = 0; i < num_rows; ++i) {
  for (auto j = row_ptrs[i]; j < row_ptrs[i + 1]; ++j) {
  if (col_idxs[j] != i) {
- temp_idxs.push_back(col_idxs[j]);
+ adj_idxs[nnz] = col_idxs[j];
+ nnz++;
  }
  }
  }
- for (auto i = 0; i < temp_idxs.size(); ++i) {
- adj_idxs[i] = temp_idxs[i];
- }
 }
 
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(