diff --git a/common/CMakeLists.txt b/common/CMakeLists.txt index 0630f91e849..b4f0437eb1a 100644 --- a/common/CMakeLists.txt +++ b/common/CMakeLists.txt @@ -13,6 +13,7 @@ set(UNIFIED_SOURCES matrix/ell_kernels.cpp matrix/hybrid_kernels.cpp matrix/sellp_kernels.cpp + matrix/sparsity_csr_kernels.cpp matrix/diagonal_kernels.cpp multigrid/pgm_kernels.cpp preconditioner/jacobi_kernels.cpp diff --git a/common/cuda_hip/matrix/sparsity_csr_kernels.hpp.inc b/common/cuda_hip/matrix/sparsity_csr_kernels.hpp.inc index 7ac83a5694e..dddd7946a04 100644 --- a/common/cuda_hip/matrix/sparsity_csr_kernels.hpp.inc +++ b/common/cuda_hip/matrix/sparsity_csr_kernels.hpp.inc @@ -110,35 +110,6 @@ __global__ __launch_bounds__(spmv_block_size) void abstract_classical_spmv( } // namespace kernel -template -void fill_in_dense(std::shared_ptr exec, - const matrix::SparsityCsr* input, - matrix::Dense* output) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_CSR_FILL_IN_DENSE_KERNEL); - - -template -void count_num_diagonal_elements( - std::shared_ptr exec, - const matrix::SparsityCsr* matrix, - 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); - - -template -void remove_diagonal_elements( - std::shared_ptr exec, const IndexType* row_ptrs, - const IndexType* col_idxs, - matrix::SparsityCsr* matrix) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL); - - template void transpose(std::shared_ptr exec, const matrix::SparsityCsr* orig, diff --git a/common/unified/matrix/sparsity_csr_kernels.cpp b/common/unified/matrix/sparsity_csr_kernels.cpp new file mode 100644 index 00000000000..fca2c904383 --- /dev/null +++ b/common/unified/matrix/sparsity_csr_kernels.cpp @@ -0,0 +1,146 @@ +/************************************************************* +Copyright (c) 2017-2023, 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 "common/unified/base/kernel_launch.hpp" +#include "core/components/prefix_sum_kernels.hpp" + + +namespace gko { +namespace kernels { +namespace GKO_DEVICE_NAMESPACE { +/** + * @brief The SparsityCsr matrix format namespace. + * + * @ingroup sparsity_csr + */ +namespace sparsity_csr { + + +template +void fill_in_dense(std::shared_ptr exec, + const matrix::SparsityCsr* input, + matrix::Dense* output) +{ + run_kernel( + exec, + [] GKO_KERNEL(auto row, auto row_ptrs, auto col_idxs, auto value, + auto output) { + const auto begin = row_ptrs[row]; + const auto end = row_ptrs[row + 1]; + const auto val = *value; + for (auto nz = begin; nz < end; nz++) { + output(row, col_idxs[nz]) = val; + } + }, + input->get_size()[0], input->get_const_row_ptrs(), + input->get_const_col_idxs(), input->get_const_value(), output); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_FILL_IN_DENSE_KERNEL); + + +template +void diagonal_element_prefix_sum( + std::shared_ptr exec, + const matrix::SparsityCsr* matrix, + IndexType* prefix_sum) +{ + const auto num_rows = matrix->get_size()[0]; + run_kernel( + exec, + [] GKO_KERNEL(auto row, auto row_ptrs, auto col_idxs, auto prefix_sum) { + const auto begin = row_ptrs[row]; + const auto end = row_ptrs[row + 1]; + IndexType count = 0; + for (auto nz = begin; nz < end; nz++) { + if (col_idxs[nz] == row) { + count++; + } + } + prefix_sum[row] = count; + }, + num_rows, matrix->get_const_row_ptrs(), matrix->get_const_col_idxs(), + prefix_sum); + components::prefix_sum(exec, prefix_sum, num_rows + 1); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_DIAGONAL_ELEMENT_PREFIX_SUM_KERNEL); + + +template +void remove_diagonal_elements(std::shared_ptr exec, + const IndexType* row_ptrs, + const IndexType* col_idxs, + const IndexType* diag_prefix_sum, + matrix::SparsityCsr* matrix) +{ + const auto num_rows = matrix->get_size()[0]; + run_kernel( + exec, + [] GKO_KERNEL(auto row, auto in_row_ptrs, auto in_col_idxs, + auto diag_prefix_sum, auto out_row_ptrs, + auto out_col_idxs) { + const auto in_begin = in_row_ptrs[row]; + const auto in_end = in_row_ptrs[row + 1]; + auto out_idx = in_begin - diag_prefix_sum[row]; + for (auto nz = in_begin; nz < in_end; nz++) { + const auto col = in_col_idxs[nz]; + if (col != row) { + out_col_idxs[out_idx] = col; + out_idx++; + } + } + if (row == 0) { + out_row_ptrs[0] = 0; + } + out_row_ptrs[row + 1] = out_idx; + }, + num_rows, row_ptrs, col_idxs, diag_prefix_sum, matrix->get_row_ptrs(), + matrix->get_col_idxs()); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( + GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL); + + +} // namespace sparsity_csr +} // namespace GKO_DEVICE_NAMESPACE +} // namespace kernels +} // namespace gko diff --git a/core/device_hooks/common_kernels.inc.cpp b/core/device_hooks/common_kernels.inc.cpp index b64c5e22b9e..708546d2243 100644 --- a/core/device_hooks/common_kernels.inc.cpp +++ b/core/device_hooks/common_kernels.inc.cpp @@ -499,7 +499,7 @@ GKO_STUB_MIXED_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_CSR_ADVANCED_SPMV_KERNEL); GKO_STUB_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SPARSITY_CSR_FILL_IN_DENSE_KERNEL); GKO_STUB_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); + GKO_DECLARE_SPARSITY_CSR_DIAGONAL_ELEMENT_PREFIX_SUM_KERNEL); GKO_STUB_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL); GKO_STUB_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL); diff --git a/core/matrix/sparsity_csr.cpp b/core/matrix/sparsity_csr.cpp index 0db304e07a5..939d0201415 100644 --- a/core/matrix/sparsity_csr.cpp +++ b/core/matrix/sparsity_csr.cpp @@ -56,8 +56,8 @@ namespace { 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_csr::count_num_diagonal_elements); +GKO_REGISTER_OPERATION(diagonal_element_prefix_sum, + sparsity_csr::diagonal_element_prefix_sum); GKO_REGISTER_OPERATION(convert_idxs_to_ptrs, components::convert_idxs_to_ptrs); GKO_REGISTER_OPERATION(fill_in_dense, sparsity_csr::fill_in_dense); GKO_REGISTER_OPERATION(remove_diagonal_elements, @@ -270,15 +270,19 @@ 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_csr::make_count_num_diagonal_elements( - this, &num_diagonal_elements)); + const auto num_rows = this->get_size()[0]; + array diag_prefix_sum{exec, num_rows + 1}; + exec->run(sparsity_csr::make_diagonal_element_prefix_sum( + this, diag_prefix_sum.get_data())); + const auto num_diagonal_elements = static_cast( + exec->copy_val_to_host(diag_prefix_sum.get_const_data() + num_rows)); auto adj_mat = SparsityCsr::create(exec, this->get_size(), this->get_num_nonzeros() - num_diagonal_elements); exec->run(sparsity_csr::make_remove_diagonal_elements( - this->get_const_row_ptrs(), this->get_const_col_idxs(), adj_mat.get())); + this->get_const_row_ptrs(), this->get_const_col_idxs(), + diag_prefix_sum.get_const_data(), adj_mat.get())); return std::move(adj_mat); } diff --git a/core/matrix/sparsity_csr_kernels.hpp b/core/matrix/sparsity_csr_kernels.hpp index 6cac29d5d0e..8f80e738b91 100644 --- a/core/matrix/sparsity_csr_kernels.hpp +++ b/core/matrix/sparsity_csr_kernels.hpp @@ -75,14 +75,15 @@ namespace kernels { void remove_diagonal_elements( \ std::shared_ptr exec, \ const IndexType* row_ptrs, const IndexType* col_idxs, \ + const IndexType* diag_prefix_sum, \ matrix::SparsityCsr* matrix) -#define GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL(ValueType, \ +#define GKO_DECLARE_SPARSITY_CSR_DIAGONAL_ELEMENT_PREFIX_SUM_KERNEL(ValueType, \ IndexType) \ - void count_num_diagonal_elements( \ + void diagonal_element_prefix_sum( \ std::shared_ptr exec, \ const matrix::SparsityCsr* matrix, \ - size_type* num_diagonal_elements) + IndexType* prefix_sum) #define GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL(ValueType, IndexType) \ void transpose(std::shared_ptr exec, \ @@ -116,7 +117,7 @@ namespace kernels { GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL(ValueType, \ IndexType); \ template \ - GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL(ValueType, \ + GKO_DECLARE_SPARSITY_CSR_DIAGONAL_ELEMENT_PREFIX_SUM_KERNEL(ValueType, \ IndexType); \ template \ GKO_DECLARE_SPARSITY_CSR_TRANSPOSE_KERNEL(ValueType, IndexType); \ diff --git a/dpcpp/matrix/sparsity_csr_kernels.dp.cpp b/dpcpp/matrix/sparsity_csr_kernels.dp.cpp index ca48553990f..769bab5108d 100644 --- a/dpcpp/matrix/sparsity_csr_kernels.dp.cpp +++ b/dpcpp/matrix/sparsity_csr_kernels.dp.cpp @@ -290,35 +290,6 @@ GKO_INSTANTIATE_FOR_EACH_MIXED_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_CSR_ADVANCED_SPMV_KERNEL); -template -void fill_in_dense(std::shared_ptr exec, - const matrix::SparsityCsr* input, - matrix::Dense* output) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_CSR_FILL_IN_DENSE_KERNEL); - - -template -void count_num_diagonal_elements( - std::shared_ptr exec, - const matrix::SparsityCsr* matrix, - 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); - - -template -void remove_diagonal_elements( - std::shared_ptr exec, const IndexType* row_ptrs, - const IndexType* col_idxs, - matrix::SparsityCsr* matrix) GKO_NOT_IMPLEMENTED; - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_CSR_REMOVE_DIAGONAL_ELEMENTS_KERNEL); - - template void transpose(std::shared_ptr exec, const matrix::SparsityCsr* orig, diff --git a/include/ginkgo/ginkgo.hpp b/include/ginkgo/ginkgo.hpp index d9b4c348d11..5b6df4088ec 100644 --- a/include/ginkgo/ginkgo.hpp +++ b/include/ginkgo/ginkgo.hpp @@ -123,6 +123,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include +#include #include #include #include diff --git a/omp/matrix/sparsity_csr_kernels.cpp b/omp/matrix/sparsity_csr_kernels.cpp index 8b50a5e7413..9c8307e233c 100644 --- a/omp/matrix/sparsity_csr_kernels.cpp +++ b/omp/matrix/sparsity_csr_kernels.cpp @@ -131,87 +131,6 @@ GKO_INSTANTIATE_FOR_EACH_MIXED_VALUE_AND_INDEX_TYPE( GKO_DECLARE_SPARSITY_CSR_ADVANCED_SPMV_KERNEL); -template -void fill_in_dense(std::shared_ptr exec, - const matrix::SparsityCsr* input, - matrix::Dense* output) -{ - auto row_ptrs = input->get_const_row_ptrs(); - auto col_idxs = input->get_const_col_idxs(); - auto val = input->get_const_value()[0]; - const auto num_rows = input->get_size()[0]; - -#pragma omp parallel for - for (size_type row = 0; row < num_rows; ++row) { - for (auto k = row_ptrs[row]; k < row_ptrs[row + 1]; ++k) { - auto col = col_idxs[k]; - output->at(row, col) = val; - } - } -} - -GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_CSR_FILL_IN_DENSE_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, - const IndexType* row_ptrs, - const IndexType* col_idxs, - matrix::SparsityCsr* matrix) -{ - 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, diff --git a/reference/matrix/sparsity_csr_kernels.cpp b/reference/matrix/sparsity_csr_kernels.cpp index 01bafb65023..7aacc9d079c 100644 --- a/reference/matrix/sparsity_csr_kernels.cpp +++ b/reference/matrix/sparsity_csr_kernels.cpp @@ -149,33 +149,35 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( template -void count_num_diagonal_elements( +void diagonal_element_prefix_sum( std::shared_ptr exec, const matrix::SparsityCsr* matrix, - size_type* num_diagonal_elements) + IndexType* prefix_sum) { 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) { + IndexType num_diag = 0; + for (size_type i = 0; i < num_rows; ++i) { + prefix_sum[i] = num_diag; for (auto j = row_ptrs[i]; j < row_ptrs[i + 1]; ++j) { if (col_idxs[j] == i) { num_diag++; } } } - *num_diagonal_elements = num_diag; + prefix_sum[num_rows] = num_diag; } GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE( - GKO_DECLARE_SPARSITY_CSR_COUNT_NUM_DIAGONAL_ELEMENTS_KERNEL); + GKO_DECLARE_SPARSITY_CSR_DIAGONAL_ELEMENT_PREFIX_SUM_KERNEL); template void remove_diagonal_elements(std::shared_ptr exec, const IndexType* row_ptrs, const IndexType* col_idxs, + const IndexType* diag_prefix_sum, matrix::SparsityCsr* matrix) { auto num_rows = matrix->get_size()[0]; @@ -183,7 +185,7 @@ void remove_diagonal_elements(std::shared_ptr exec, 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 (size_type 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++; @@ -192,7 +194,7 @@ void remove_diagonal_elements(std::shared_ptr exec, adj_ptrs[i + 1] = row_ptrs[i + 1] - num_diag; } auto nnz = 0; - for (auto i = 0; i < num_rows; ++i) { + for (size_type 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]; diff --git a/reference/test/matrix/sparsity_csr_kernels.cpp b/reference/test/matrix/sparsity_csr_kernels.cpp index 0388d31910d..4d356ffd828 100644 --- a/reference/test/matrix/sparsity_csr_kernels.cpp +++ b/reference/test/matrix/sparsity_csr_kernels.cpp @@ -395,9 +395,10 @@ TYPED_TEST(SparsityCsr, NonSquareMtxIsTransposable) } -TYPED_TEST(SparsityCsr, CountsCorrectNumberOfDiagonalElements) +TYPED_TEST(SparsityCsr, ComputesCorrectDiagonalElementPrefixSum) { using Mtx = typename TestFixture::Mtx; + using index_type = typename TestFixture::index_type; // clang-format off auto mtx2 = gko::initialize({{1.0, 1.0, 1.0}, {0.0, 1.0, 0.0}, @@ -406,22 +407,23 @@ TYPED_TEST(SparsityCsr, CountsCorrectNumberOfDiagonalElements) {0.0, 0.0, 0.0}, {0.0, 1.0, 1.0}}, this->exec); // clang-format on - gko::size_type m2_num_diags = 0; - gko::size_type ms_num_diags = 0; + gko::array prefix_sum1{this->exec, 4}; + gko::array prefix_sum2{this->exec, 4}; - gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( - this->exec, mtx2.get(), &m2_num_diags); - gko::kernels::reference::sparsity_csr::count_num_diagonal_elements( - this->exec, mtx_s.get(), &ms_num_diags); + gko::kernels::reference::sparsity_csr::diagonal_element_prefix_sum( + this->exec, mtx2.get(), prefix_sum1.get_data()); + gko::kernels::reference::sparsity_csr::diagonal_element_prefix_sum( + this->exec, mtx_s.get(), prefix_sum2.get_data()); - ASSERT_EQ(m2_num_diags, 3); - ASSERT_EQ(ms_num_diags, 2); + GKO_ASSERT_ARRAY_EQ(prefix_sum1, I({0, 1, 2, 3})); + GKO_ASSERT_ARRAY_EQ(prefix_sum2, I({0, 1, 1, 2})); } TYPED_TEST(SparsityCsr, RemovesDiagonalElementsForFullRankMatrix) { using Mtx = typename TestFixture::Mtx; + using index_type = typename TestFixture::index_type; // clang-format off auto mtx2 = gko::initialize({{1.0, 1.0, 1.0}, {0.0, 1.0, 0.0}, @@ -433,10 +435,13 @@ TYPED_TEST(SparsityCsr, RemovesDiagonalElementsForFullRankMatrix) auto tmp_mtx = Mtx::create(this->exec, mtx_s->get_size(), mtx_s->get_num_nonzeros()); tmp_mtx->copy_from(mtx2); + gko::array prefix_sum{this->exec, 4}; + gko::kernels::reference::sparsity_csr::diagonal_element_prefix_sum( + this->exec, mtx2.get(), prefix_sum.get_data()); gko::kernels::reference::sparsity_csr::remove_diagonal_elements( this->exec, mtx2->get_const_row_ptrs(), mtx2->get_const_col_idxs(), - tmp_mtx.get()); + prefix_sum.get_const_data(), tmp_mtx.get()); GKO_ASSERT_MTX_NEAR(tmp_mtx, mtx_s, 0.0); } @@ -445,6 +450,7 @@ TYPED_TEST(SparsityCsr, RemovesDiagonalElementsForFullRankMatrix) TYPED_TEST(SparsityCsr, RemovesDiagonalElementsForIncompleteRankMatrix) { using Mtx = typename TestFixture::Mtx; + using index_type = typename TestFixture::index_type; // clang-format off auto mtx2 = gko::initialize({{1.0, 1.0, 1.0}, {0.0, 0.0, 0.0}, @@ -456,10 +462,13 @@ TYPED_TEST(SparsityCsr, RemovesDiagonalElementsForIncompleteRankMatrix) auto tmp_mtx = Mtx::create(this->exec, mtx_s->get_size(), mtx_s->get_num_nonzeros()); tmp_mtx->copy_from(mtx2); + gko::array prefix_sum{this->exec, 4}; + gko::kernels::reference::sparsity_csr::diagonal_element_prefix_sum( + this->exec, mtx2.get(), prefix_sum.get_data()); gko::kernels::reference::sparsity_csr::remove_diagonal_elements( this->exec, mtx2->get_const_row_ptrs(), mtx2->get_const_col_idxs(), - tmp_mtx.get()); + prefix_sum.get_const_data(), tmp_mtx.get()); GKO_ASSERT_MTX_NEAR(tmp_mtx, mtx_s, 0.0); } diff --git a/test/matrix/CMakeLists.txt b/test/matrix/CMakeLists.txt index e906a1b4242..a9cf267a3c8 100644 --- a/test/matrix/CMakeLists.txt +++ b/test/matrix/CMakeLists.txt @@ -13,3 +13,4 @@ endif() ginkgo_create_common_test(hybrid_kernels) ginkgo_create_common_test(matrix) ginkgo_create_common_test(sellp_kernels) +ginkgo_create_common_test(sparsity_csr_kernels) diff --git a/test/matrix/sparsity_csr_kernels.cpp b/test/matrix/sparsity_csr_kernels.cpp new file mode 100644 index 00000000000..34f01f8a02b --- /dev/null +++ b/test/matrix/sparsity_csr_kernels.cpp @@ -0,0 +1,150 @@ +/************************************************************* +Copyright (c) 2017-2023, 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 + + +#include "core/matrix/sparsity_csr_kernels.hpp" +#include "core/test/utils.hpp" +#include "core/test/utils/assertions.hpp" +#include "core/test/utils/matrix_generator.hpp" +#include "test/utils/executor.hpp" + + +namespace { + + +class SparsityCsr : public CommonTestFixture { +protected: + using value_type = double; + using index_type = int; + using Mtx = gko::matrix::SparsityCsr; + + SparsityCsr() : rng{9312} + { + auto data = + gko::test::generate_random_matrix_data( + 100, 100, std::uniform_int_distribution(1, 10), + std::uniform_real_distribution(0.0, 1.0), rng); + // make sure the matrix contains a few diagonal entries + for (int i = 0; i < 10; i++) { + data.nonzeros.emplace_back(i * 3, i * 3, 0.0); + } + data.sum_duplicates(); + mtx = Mtx::create(ref); + mtx->read(data); + dmtx = gko::clone(exec, mtx); + } + + std::default_random_engine rng; + std::unique_ptr mtx; + std::unique_ptr dmtx; +}; + + +TEST_F(SparsityCsr, KernelDiagonalElementPrefixSumIsEquivalentToRef) +{ + gko::array prefix_sum{this->ref, this->mtx->get_size()[0] + 1}; + gko::array dprefix_sum{this->exec, + this->mtx->get_size()[0] + 1}; + + gko::kernels::reference::sparsity_csr::diagonal_element_prefix_sum( + ref, mtx.get(), prefix_sum.get_data()); + gko::kernels::EXEC_NAMESPACE::sparsity_csr::diagonal_element_prefix_sum( + exec, dmtx.get(), dprefix_sum.get_data()); + + GKO_ASSERT_ARRAY_EQ(prefix_sum, dprefix_sum); +} + + +TEST_F(SparsityCsr, KernelRemoveDiagonalElementsIsEquivalentToRef) +{ + const auto num_rows = this->mtx->get_size()[0]; + gko::array prefix_sum{this->ref, num_rows + 1}; + gko::kernels::reference::sparsity_csr::diagonal_element_prefix_sum( + ref, mtx.get(), prefix_sum.get_data()); + gko::array dprefix_sum{this->exec, prefix_sum}; + const auto out_mtx = Mtx::create( + ref, mtx->get_size(), + mtx->get_num_nonzeros() - prefix_sum.get_const_data()[num_rows]); + const auto dout_mtx = Mtx::create( + exec, mtx->get_size(), + mtx->get_num_nonzeros() - prefix_sum.get_const_data()[num_rows]); + + gko::kernels::reference::sparsity_csr::remove_diagonal_elements( + ref, mtx->get_const_row_ptrs(), mtx->get_const_col_idxs(), + prefix_sum.get_const_data(), out_mtx.get()); + gko::kernels::EXEC_NAMESPACE::sparsity_csr::remove_diagonal_elements( + exec, dmtx->get_const_row_ptrs(), dmtx->get_const_col_idxs(), + dprefix_sum.get_const_data(), dout_mtx.get()); + + GKO_ASSERT_MTX_NEAR(out_mtx, dout_mtx, 0.0); +} + + +TEST_F(SparsityCsr, ToAdjacencyMatrixIsEquivalentToRef) +{ + const auto out_mtx = mtx->to_adjacency_matrix(); + const auto dout_mtx = dmtx->to_adjacency_matrix(); + + GKO_ASSERT_MTX_NEAR(out_mtx, dout_mtx, 0.0); +} + + +TEST_F(SparsityCsr, ConvertToDenseIsEquivalentToRef) +{ + const auto out_dense = gko::matrix::Dense::create( + exec, mtx->get_size(), mtx->get_size()[1] + 2); + const auto dout_dense = gko::matrix::Dense::create( + exec, mtx->get_size(), mtx->get_size()[1] + 2); + + mtx->convert_to(out_dense); + dmtx->convert_to(dout_dense); + + GKO_ASSERT_MTX_NEAR(out_dense, dout_dense, 0.0); +} + + +} // namespace diff --git a/test/reorder/CMakeLists.txt b/test/reorder/CMakeLists.txt index 2df2d6e544c..4e1b90b997b 100644 --- a/test/reorder/CMakeLists.txt +++ b/test/reorder/CMakeLists.txt @@ -1,3 +1,3 @@ if (GINKGO_HAVE_METIS) - ginkgo_create_common_test(nested_dissection DISABLE_EXECUTORS cuda hip dpcpp) + ginkgo_create_common_test(nested_dissection) endif() \ No newline at end of file