diff --git a/common/cuda_hip/matrix/batch_ell_kernel_launcher.hpp.inc b/common/cuda_hip/matrix/batch_ell_kernel_launcher.hpp.inc new file mode 100644 index 00000000000..f8da432aa4d --- /dev/null +++ b/common/cuda_hip/matrix/batch_ell_kernel_launcher.hpp.inc @@ -0,0 +1,78 @@ +/************************************************************* +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. +*************************************************************/ + +template +void simple_apply(std::shared_ptr exec, + const batch::matrix::Ell* mat, + const batch::MultiVector* b, + batch::MultiVector* x) +{ + const auto num_blocks = mat->get_num_batch_items(); + const auto b_ub = get_batch_struct(b); + const auto x_ub = get_batch_struct(x); + const auto mat_ub = get_batch_struct(mat); + if (b->get_common_size()[1] > 1) { + GKO_NOT_IMPLEMENTED; + } + simple_apply_kernel<<get_stream()>>>(mat_ub, b_ub, x_ub); +} + + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INT32_TYPE( + GKO_DECLARE_BATCH_ELL_SIMPLE_APPLY_KERNEL); + + +template +void advanced_apply(std::shared_ptr exec, + const batch::MultiVector* alpha, + const batch::matrix::Ell* mat, + const batch::MultiVector* b, + const batch::MultiVector* beta, + batch::MultiVector* x) +{ + const auto num_blocks = mat->get_num_batch_items(); + const auto b_ub = get_batch_struct(b); + const auto x_ub = get_batch_struct(x); + const auto mat_ub = get_batch_struct(mat); + const auto alpha_ub = get_batch_struct(alpha); + const auto beta_ub = get_batch_struct(beta); + if (b->get_common_size()[1] > 1) { + GKO_NOT_IMPLEMENTED; + } + advanced_apply_kernel<<get_stream()>>>(alpha_ub, mat_ub, b_ub, + beta_ub, x_ub); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INT32_TYPE( + GKO_DECLARE_BATCH_ELL_ADVANCED_APPLY_KERNEL); diff --git a/common/cuda_hip/matrix/batch_ell_kernels.hpp.inc b/common/cuda_hip/matrix/batch_ell_kernels.hpp.inc new file mode 100644 index 00000000000..de6ca879890 --- /dev/null +++ b/common/cuda_hip/matrix/batch_ell_kernels.hpp.inc @@ -0,0 +1,156 @@ +/************************************************************* +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. +*************************************************************/ + + +template +__device__ __forceinline__ void simple_apply( + const gko::batch::matrix::ell::batch_item& mat, + const ValueType* const __restrict__ b, ValueType* const __restrict__ x) +{ + const auto num_rows = mat.num_rows; + const auto num_stored_elements_per_row = mat.num_stored_elems_per_row; + const auto stride = mat.stride; + const auto val = mat.values; + const auto col = mat.col_idxs; + for (int tidx = threadIdx.x; tidx < num_rows; tidx += blockDim.x) { + auto temp = zero(); + for (size_type idx = 0; idx < num_stored_elements_per_row; idx++) { + const auto ind = tidx + idx * stride; + const auto col_idx = col[ind]; + if (col_idx == invalid_index()) { + break; + } else { + temp += val[ind] * b[col_idx]; + } + } + x[tidx] = temp; + } +} + +template +__global__ __launch_bounds__( + default_block_size, + sm_oversubscription) void simple_apply_kernel(const gko::batch::matrix:: + ell::uniform_batch< + const ValueType, + IndexType> + mat, + const gko::batch:: + multi_vector:: + uniform_batch< + const ValueType> + b, + const gko::batch:: + multi_vector:: + uniform_batch< + ValueType> + x) +{ + for (size_type batch_id = blockIdx.x; batch_id < mat.num_batch_items; + batch_id += gridDim.x) { + const auto mat_b = + gko::batch::matrix::extract_batch_item(mat, batch_id); + const auto b_b = gko::batch::extract_batch_item(b, batch_id); + const auto x_b = gko::batch::extract_batch_item(x, batch_id); + simple_apply(mat_b, b_b.values, x_b.values); + } +} + + +template +__device__ __forceinline__ void advanced_apply( + const ValueType alpha, + const gko::batch::matrix::ell::batch_item& mat, + const ValueType* const __restrict__ b, const ValueType beta, + ValueType* const __restrict__ x) +{ + const auto num_rows = mat.num_rows; + const auto num_stored_elements_per_row = mat.num_stored_elems_per_row; + const auto stride = mat.stride; + const auto val = mat.values; + const auto col = mat.col_idxs; + for (int tidx = threadIdx.x; tidx < num_rows; tidx += blockDim.x) { + auto temp = zero(); + for (size_type idx = 0; idx < num_stored_elements_per_row; idx++) { + const auto ind = tidx + idx * stride; + const auto col_idx = col[ind]; + if (col_idx == invalid_index()) { + break; + } else { + temp += alpha * val[ind] * b[col_idx]; + } + } + x[tidx] = temp + beta * x[tidx]; + } +} + +template +__global__ __launch_bounds__( + default_block_size, + sm_oversubscription) void advanced_apply_kernel(const gko::batch:: + multi_vector:: + uniform_batch< + const ValueType> + alpha, + const gko::batch::matrix:: + ell::uniform_batch< + const ValueType, + IndexType> + mat, + const gko::batch:: + multi_vector:: + uniform_batch< + const ValueType> + b, + const gko::batch:: + multi_vector:: + uniform_batch< + const ValueType> + beta, + const gko::batch:: + multi_vector:: + uniform_batch< + ValueType> + x) +{ + for (size_type batch_id = blockIdx.x; batch_id < mat.num_batch_items; + batch_id += gridDim.x) { + const auto mat_b = + gko::batch::matrix::extract_batch_item(mat, batch_id); + const auto b_b = gko::batch::extract_batch_item(b, batch_id); + const auto x_b = gko::batch::extract_batch_item(x, batch_id); + const auto alpha_b = gko::batch::extract_batch_item(alpha, batch_id); + const auto beta_b = gko::batch::extract_batch_item(beta, batch_id); + advanced_apply(alpha_b.values[0], mat_b, b_b.values, beta_b.values[0], + x_b.values); + } +} diff --git a/core/CMakeLists.txt b/core/CMakeLists.txt index 46ea67abc65..ae8035bcbf9 100644 --- a/core/CMakeLists.txt +++ b/core/CMakeLists.txt @@ -40,6 +40,7 @@ target_sources(ginkgo log/record.cpp log/stream.cpp matrix/batch_dense.cpp + matrix/batch_ell.cpp matrix/coo.cpp matrix/csr.cpp matrix/dense.cpp diff --git a/core/base/batch_multi_vector.cpp b/core/base/batch_multi_vector.cpp index 6a14919bf2f..6dcf8dd90b5 100644 --- a/core/base/batch_multi_vector.cpp +++ b/core/base/batch_multi_vector.cpp @@ -291,27 +291,6 @@ void MultiVector::move_to( } -template -void MultiVector::convert_to(matrix::Dense* result) const -{ - auto exec = result->get_executor() == nullptr ? this->get_executor() - : result->get_executor(); - auto tmp = gko::batch::matrix::Dense::create_const( - exec, this->get_size(), - make_const_array_view(this->get_executor(), - this->get_num_stored_elements(), - this->get_const_values())); - result->copy_from(tmp); -} - - -template -void MultiVector::move_to(matrix::Dense* result) -{ - this->convert_to(result); -} - - #define GKO_DECLARE_BATCH_MULTI_VECTOR(_type) class MultiVector<_type> GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_BATCH_MULTI_VECTOR); diff --git a/core/base/batch_utilities.hpp b/core/base/batch_utilities.hpp index 834e89c8358..b4e380a4162 100644 --- a/core/base/batch_utilities.hpp +++ b/core/base/batch_utilities.hpp @@ -47,21 +47,28 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include namespace gko { namespace batch { -template +/** + * Duplicate a given input batch object. + */ +template std::unique_ptr duplicate(std::shared_ptr exec, size_type num_duplications, - const OutputType* input) + const OutputType* input, + TArgs&&... create_args) { auto num_batch_items = input->get_num_batch_items(); - auto tmp = OutputType::create( - exec, batch_dim<2>(num_batch_items * num_duplications, - input->get_common_size())); + auto tmp = + OutputType::create(exec, + batch_dim<2>(num_batch_items * num_duplications, + input->get_common_size()), + std::forward(create_args)...); for (size_type i = 0; i < num_duplications; ++i) { for (size_type b = 0; b < num_batch_items; ++b) { @@ -74,14 +81,18 @@ std::unique_ptr duplicate(std::shared_ptr exec, } -template +/** + * Duplicate a monolithic matrix and create a batch object. + */ +template std::unique_ptr create_from_item( std::shared_ptr exec, const size_type num_duplications, - const typename OutputType::unbatch_type* input) + const typename OutputType::unbatch_type* input, TArgs&&... create_args) { auto num_batch_items = num_duplications; auto tmp = OutputType::create( - exec, batch_dim<2>(num_batch_items, input->get_size())); + exec, batch_dim<2>(num_batch_items, input->get_size()), + std::forward(create_args)...); for (size_type b = 0; b < num_batch_items; ++b) { tmp->create_view_for_item(b)->copy_from(input); @@ -91,14 +102,23 @@ std::unique_ptr create_from_item( } -template +/** + * Create a batch object from a vector of monolithic object that share the same + * sparsity pattern. + * + * @note The sparsity of the elements in the input vector of matrices needs to + * be the same. TODO: Check for same sparsity among the different input items + */ +template std::unique_ptr create_from_item( std::shared_ptr exec, - const std::vector& input) + const std::vector& input, + TArgs&&... create_args) { auto num_batch_items = input.size(); auto tmp = OutputType::create( - exec, batch_dim<2>(num_batch_items, input[0]->get_size())); + exec, batch_dim<2>(num_batch_items, input[0]->get_size()), + std::forward(create_args)...); for (size_type b = 0; b < num_batch_items; ++b) { tmp->create_view_for_item(b)->copy_from(input[b]); @@ -108,6 +128,9 @@ std::unique_ptr create_from_item( } +/** + * Unbatch a batched object into a vector of items of its unbatch_type. + */ template auto unbatch(const InputType* batch_object) { @@ -121,14 +144,61 @@ auto unbatch(const InputType* batch_object) } -template +namespace detail { + + +template +void assert_same_sparsity_in_batched_data( + const std::vector>& data) +{ + if (data.empty()) { + return; + } + auto num_nnz = data.at(0).nonzeros.size(); + auto base_data = data.at(0); + base_data.ensure_row_major_order(); + for (int b = 1; b < data.size(); ++b) { + if (data[b].nonzeros.size() != num_nnz) { + GKO_NOT_IMPLEMENTED; + } + auto temp_data = data.at(b); + temp_data.ensure_row_major_order(); + for (int nnz = 0; nnz < num_nnz; ++nnz) { + if (temp_data.nonzeros.at(nnz).row != + base_data.nonzeros.at(nnz).row || + temp_data.nonzeros.at(nnz).column != + base_data.nonzeros.at(nnz).column) { + GKO_NOT_IMPLEMENTED; + } + } + } +} + + +} // namespace detail + + +/** + * Create a batch object from a vector of gko::matrix_data objects. Each item of + * the vector needs to store the same sparsity pattern. + */ +template std::unique_ptr read( std::shared_ptr exec, - const std::vector>& data) + const std::vector>& data, + TArgs&&... create_args) { auto num_batch_items = data.size(); + // Throw if all the items in the batch dont have same sparsity. + if (!std::is_same>::value && + !std::is_same>::value) { + detail::assert_same_sparsity_in_batched_data(data); + } auto tmp = - OutputType::create(exec, batch_dim<2>(num_batch_items, data[0].size)); + OutputType::create(exec, batch_dim<2>(num_batch_items, data.at(0).size), + std::forward(create_args)...); for (size_type b = 0; b < num_batch_items; ++b) { tmp->create_view_for_item(b)->read(data[b]); @@ -138,6 +208,9 @@ std::unique_ptr read( } +/** + * Write a vector of matrix data objects from an input batch object. + */ template std::vector> write( const OutputType* mvec) @@ -154,6 +227,220 @@ std::vector> write( } +/** + * Creates and initializes a batch of the specified Matrix type from a series of + * single column-vectors. + * + * @tparam Matrix matrix type to initialize (It has to implement the + * read function) + * @tparam TArgs argument types for Matrix::create method + * (not including the implied Executor as the first argument) + * + * @param vals values used to initialize the batch vector + * @param exec Executor associated to the vector + * @param create_args additional arguments passed to Matrix::create, not + * including the Executor, which is passed as the first + * argument + * + * @ingroup mat_formats + */ +template +std::unique_ptr initialize( + std::initializer_list> + vals, + std::shared_ptr exec, TArgs&&... create_args) +{ + using value_type = typename Matrix::value_type; + using index_type = typename Matrix::index_type; + using mat_data = gko::matrix_data; + size_type num_batch_items = vals.size(); + GKO_THROW_IF_INVALID(num_batch_items > 0, "Input data is empty"); + auto vals_begin = begin(vals); + size_type common_num_rows = vals_begin ? vals_begin->size() : 0; + auto common_size = dim<2>(common_num_rows, 1); + for (auto& val : vals) { + GKO_ASSERT_EQ(common_num_rows, val.size()); + } + auto b_size = batch_dim<2>(num_batch_items, common_size); + size_type batch = 0; + std::vector input_mat_data(num_batch_items, common_size); + for (const auto& b : vals) { + input_mat_data[batch].nonzeros.reserve(b.size()); + size_type idx = 0; + for (const auto& elem : b) { + if (elem != zero()) { + input_mat_data[batch].nonzeros.emplace_back(idx, 0, elem); + } + ++idx; + } + ++batch; + } + return read( + exec, input_mat_data, std::forward(create_args)...); +} + + +/** + * Creates and initializes a batch of matrices. + * + * @tparam Matrix matrix type to initialize (It has to implement the + * read function) + * @tparam TArgs argument types for Matrix::create method + * (not including the implied Executor as the first argument) + * + * @param vals values used to initialize the matrix + * @param exec Executor associated with the matrix + * @param create_args additional arguments passed to Matrix::create, not + * including the Executor, which is passed as the first + * argument + * + * @ingroup mat_formats + */ +template +std::unique_ptr initialize( + std::initializer_list>> + vals, + std::shared_ptr exec, TArgs&&... create_args) +{ + using value_type = typename Matrix::value_type; + using index_type = typename Matrix::index_type; + using mat_data = gko::matrix_data; + size_type num_batch_items = vals.size(); + GKO_THROW_IF_INVALID(num_batch_items > 0, "Input data is empty"); + auto vals_begin = begin(vals); + size_type common_num_rows = vals_begin ? vals_begin->size() : 0; + size_type common_num_cols = + vals_begin->begin() ? vals_begin->begin()->size() : 0; + auto common_size = dim<2>(common_num_rows, common_num_cols); + for (const auto& b : vals) { + auto num_rows = b.size(); + auto num_cols = begin(b)->size(); + auto b_size = dim<2>(num_rows, num_cols); + GKO_ASSERT_EQUAL_DIMENSIONS(b_size, common_size); + } + + auto b_size = batch_dim<2>(num_batch_items, common_size); + size_type batch = 0; + std::vector input_mat_data(num_batch_items, common_size); + for (const auto& b : vals) { + size_type ridx = 0; + for (const auto& row : b) { + size_type cidx = 0; + for (const auto& elem : row) { + if (elem != zero()) { + input_mat_data[batch].nonzeros.emplace_back(ridx, cidx, + elem); + } + ++cidx; + } + ++ridx; + } + ++batch; + } + return read( + exec, input_mat_data, std::forward(create_args)...); +} + + +/** + * Creates and initializes a batch of specified Matrix type with a single + * column-vector by making copies of the single input column vector. + * + * @tparam Matrix matrix type to initialize (It has to implement the + * read function) + * @tparam TArgs argument types for Matrix::create method + * (not including the implied Executor as the first argument) + * + * @param num_batch_items The number of times the input vector is to be + * duplicated + * @param vals values used to initialize each vector in the temp. batch + * @param exec Executor associated with the matrix + * @param create_args additional arguments passed to Matrix::create, not + * including the Executor, which is passed as the first + * argument + * + * @ingroup mat_formats + */ +template +std::unique_ptr initialize( + const size_type num_batch_items, + std::initializer_list vals, + std::shared_ptr exec, TArgs&&... create_args) +{ + using value_type = typename Matrix::value_type; + using index_type = typename Matrix::index_type; + using mat_data = gko::matrix_data; + GKO_THROW_IF_INVALID(num_batch_items > 0 && vals.size() > 0, + "Input data is empty"); + auto num_rows = begin(vals) ? vals.size() : 0; + auto common_size = dim<2>(num_rows, 1); + auto b_size = batch_dim<2>(num_batch_items, common_size); + mat_data single_mat_data(common_size); + single_mat_data.nonzeros.reserve(num_rows); + size_type idx = 0; + for (const auto& elem : vals) { + if (elem != zero()) { + single_mat_data.nonzeros.emplace_back(idx, 0, elem); + } + ++idx; + } + std::vector input_mat_data(num_batch_items, single_mat_data); + return read( + exec, input_mat_data, std::forward(create_args)...); +} + + +/** + * Creates and initializes a matrix from copies of a given matrix. + * + * @tparam Matrix matrix type to initialize (It has to implement the + * read function) + * @tparam TArgs argument types for Matrix::create method + * (not including the implied Executor as the first argument) + * + * @param num_batch_items The number of times the input matrix is duplicated + * @param vals values used to initialize each matrix in the temp. batch + * @param exec Executor associated to the matrix + * @param create_args additional arguments passed to Matrix::create, not + * including the Executor, which is passed as the first + * argument + * + * @ingroup mat_formats + */ +template +std::unique_ptr initialize( + const size_type num_batch_items, + std::initializer_list> + vals, + std::shared_ptr exec, TArgs&&... create_args) +{ + using value_type = typename Matrix::value_type; + using index_type = typename Matrix::index_type; + using mat_data = gko::matrix_data; + GKO_THROW_IF_INVALID(num_batch_items > 0 && vals.size() > 0, + "Input data is empty"); + auto common_size = dim<2>(begin(vals) ? vals.size() : 0, + begin(vals) ? begin(vals)->size() : 0); + batch_dim<2> b_size(num_batch_items, common_size); + mat_data single_mat_data(common_size); + size_type ridx = 0; + for (const auto& row : vals) { + size_type cidx = 0; + for (const auto& elem : row) { + if (elem != zero()) { + single_mat_data.nonzeros.emplace_back(ridx, cidx, elem); + } + ++cidx; + } + ++ridx; + } + std::vector input_mat_data(num_batch_items, single_mat_data); + return read( + exec, input_mat_data, std::forward(create_args)...); +} + + } // namespace batch } // namespace gko diff --git a/core/device_hooks/common_kernels.inc.cpp b/core/device_hooks/common_kernels.inc.cpp index 87cab3dcf0b..462675c15db 100644 --- a/core/device_hooks/common_kernels.inc.cpp +++ b/core/device_hooks/common_kernels.inc.cpp @@ -58,6 +58,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "core/factorization/par_ilu_kernels.hpp" #include "core/factorization/par_ilut_kernels.hpp" #include "core/matrix/batch_dense_kernels.hpp" +#include "core/matrix/batch_ell_kernels.hpp" #include "core/matrix/coo_kernels.hpp" #include "core/matrix/csr_kernels.hpp" #include "core/matrix/dense_kernels.hpp" @@ -137,6 +138,11 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. _macro(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); \ GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(_macro) +#define GKO_STUB_VALUE_AND_INT32_TYPE(_macro) \ + template \ + _macro(ValueType, IndexType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); \ + GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INT32_TYPE(_macro) + #define GKO_STUB_MIXED_VALUE_AND_INDEX_TYPE(_macro) \ template \ @@ -310,6 +316,16 @@ GKO_STUB_VALUE_TYPE(GKO_DECLARE_BATCH_DENSE_ADVANCED_APPLY_KERNEL); } // namespace batch_dense +namespace batch_ell { + + +GKO_STUB_VALUE_AND_INT32_TYPE(GKO_DECLARE_BATCH_ELL_SIMPLE_APPLY_KERNEL); +GKO_STUB_VALUE_AND_INT32_TYPE(GKO_DECLARE_BATCH_ELL_ADVANCED_APPLY_KERNEL); + + +} // namespace batch_ell + + namespace dense { diff --git a/core/matrix/batch_dense.cpp b/core/matrix/batch_dense.cpp index 758635cea7f..58c7fa25cea 100644 --- a/core/matrix/batch_dense.cpp +++ b/core/matrix/batch_dense.cpp @@ -96,14 +96,6 @@ Dense::create_const_view_for_item(size_type item_id) const } -template -std::unique_ptr> Dense::create_with_config_of( - ptr_param> other) -{ - return Dense::create(other->get_executor(), other->get_size()); -} - - template std::unique_ptr> Dense::create_const( std::shared_ptr exec, const batch_dim<2>& sizes, @@ -124,11 +116,72 @@ Dense::Dense(std::shared_ptr exec, {} +template +Dense* Dense::apply( + ptr_param> b, + ptr_param> x) +{ + this->validate_application_parameters(b.get(), x.get()); + auto exec = this->get_executor(); + this->apply_impl(make_temporary_clone(exec, b).get(), + make_temporary_clone(exec, x).get()); + return this; +} + + +template +const Dense* Dense::apply( + ptr_param> b, + ptr_param> x) const +{ + this->validate_application_parameters(b.get(), x.get()); + auto exec = this->get_executor(); + this->apply_impl(make_temporary_clone(exec, b).get(), + make_temporary_clone(exec, x).get()); + return this; +} + + +template +Dense* Dense::apply( + ptr_param> alpha, + ptr_param> b, + ptr_param> beta, + ptr_param> x) +{ + this->validate_application_parameters(alpha.get(), b.get(), beta.get(), + x.get()); + auto exec = this->get_executor(); + this->apply_impl(make_temporary_clone(exec, alpha).get(), + make_temporary_clone(exec, b).get(), + make_temporary_clone(exec, beta).get(), + make_temporary_clone(exec, x).get()); + return this; +} + + +template +const Dense* Dense::apply( + ptr_param> alpha, + ptr_param> b, + ptr_param> beta, + ptr_param> x) const +{ + this->validate_application_parameters(alpha.get(), b.get(), beta.get(), + x.get()); + auto exec = this->get_executor(); + this->apply_impl(make_temporary_clone(exec, alpha).get(), + make_temporary_clone(exec, b).get(), + make_temporary_clone(exec, beta).get(), + make_temporary_clone(exec, x).get()); + return this; +} + + template void Dense::apply_impl(const MultiVector* b, MultiVector* x) const { - this->validate_application_parameters(b, x); this->get_executor()->run(dense::make_simple_apply(this, b, x)); } @@ -139,7 +192,6 @@ void Dense::apply_impl(const MultiVector* alpha, const MultiVector* beta, MultiVector* x) const { - this->validate_application_parameters(alpha, b, beta, x); this->get_executor()->run( dense::make_advanced_apply(alpha, this, b, beta, x)); } diff --git a/core/matrix/batch_ell.cpp b/core/matrix/batch_ell.cpp new file mode 100644 index 00000000000..b2987e741d9 --- /dev/null +++ b/core/matrix/batch_ell.cpp @@ -0,0 +1,229 @@ +/************************************************************* +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 +#include + + +#include "core/matrix/batch_ell_kernels.hpp" + + +namespace gko { +namespace batch { +namespace matrix { +namespace ell { +namespace { + + +GKO_REGISTER_OPERATION(simple_apply, batch_ell::simple_apply); +GKO_REGISTER_OPERATION(advanced_apply, batch_ell::advanced_apply); + + +} // namespace +} // namespace ell + + +template +std::unique_ptr> +Ell::create_view_for_item(size_type item_id) +{ + auto exec = this->get_executor(); + auto num_rows = this->get_common_size()[0]; + auto stride = this->get_common_size()[0]; + auto mat = unbatch_type::create( + exec, this->get_common_size(), + make_array_view(exec, this->get_num_elements_per_item(), + this->get_values_for_item(item_id)), + make_array_view(exec, this->get_num_elements_per_item(), + this->get_col_idxs()), + this->get_num_stored_elements_per_row(), stride); + return mat; +} + + +template +std::unique_ptr> +Ell::create_const_view_for_item(size_type item_id) const +{ + auto exec = this->get_executor(); + auto num_rows = this->get_common_size()[0]; + auto stride = this->get_common_size()[0]; + auto mat = unbatch_type::create_const( + exec, this->get_common_size(), + make_const_array_view(exec, this->get_num_elements_per_item(), + this->get_const_values_for_item(item_id)), + make_const_array_view(exec, this->get_num_elements_per_item(), + this->get_const_col_idxs()), + this->get_num_stored_elements_per_row(), stride); + return mat; +} + + +template +std::unique_ptr> +Ell::create_const( + std::shared_ptr exec, const batch_dim<2>& sizes, + const IndexType num_elems_per_row, + gko::detail::const_array_view&& values, + gko::detail::const_array_view&& col_idxs) +{ + // cast const-ness away, but return a const object afterwards, + // so we can ensure that no modifications take place. + return std::unique_ptr( + new Ell{exec, sizes, num_elems_per_row, + gko::detail::array_const_cast(std::move(values)), + gko::detail::array_const_cast(std::move(col_idxs))}); +} + + +template +Ell::Ell(std::shared_ptr exec, + const batch_dim<2>& size, + IndexType num_elems_per_row) + : EnableBatchLinOp>(exec, size), + num_elems_per_row_(num_elems_per_row == 0 ? size.get_common_size()[1] + : num_elems_per_row), + values_(exec, compute_num_elems(size, num_elems_per_row_)), + col_idxs_(exec, this->get_common_size()[0] * num_elems_per_row_) +{} + + +template +Ell* Ell::apply( + ptr_param> b, + ptr_param> x) +{ + this->validate_application_parameters(b.get(), x.get()); + auto exec = this->get_executor(); + this->apply_impl(make_temporary_clone(exec, b).get(), + make_temporary_clone(exec, x).get()); + return this; +} + + +template +const Ell* Ell::apply( + ptr_param> b, + ptr_param> x) const +{ + this->apply(b, x); + return this; +} + + +template +Ell* Ell::apply( + ptr_param> alpha, + ptr_param> b, + ptr_param> beta, + ptr_param> x) +{ + this->validate_application_parameters(alpha.get(), b.get(), beta.get(), + x.get()); + auto exec = this->get_executor(); + this->apply_impl(make_temporary_clone(exec, alpha).get(), + make_temporary_clone(exec, b).get(), + make_temporary_clone(exec, beta).get(), + make_temporary_clone(exec, x).get()); + return this; +} + + +template +const Ell* Ell::apply( + ptr_param> alpha, + ptr_param> b, + ptr_param> beta, + ptr_param> x) const +{ + this->apply(alpha, b, beta, x); + return this; +} + + +template +void Ell::apply_impl(const MultiVector* b, + MultiVector* x) const +{ + this->get_executor()->run(ell::make_simple_apply(this, b, x)); +} + + +template +void Ell::apply_impl(const MultiVector* alpha, + const MultiVector* b, + const MultiVector* beta, + MultiVector* x) const +{ + this->get_executor()->run( + ell::make_advanced_apply(alpha, this, b, beta, x)); +} + + +template +void Ell::convert_to( + Ell, IndexType>* result) const +{ + result->values_ = this->values_; + result->col_idxs_ = this->col_idxs_; + result->num_elems_per_row_ = this->num_elems_per_row_; + result->set_size(this->get_size()); +} + + +template +void Ell::move_to( + Ell, IndexType>* result) +{ + this->convert_to(result); +} + + +#define GKO_DECLARE_BATCH_ELL_MATRIX(ValueType) class Ell +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_BATCH_ELL_MATRIX); + + +} // namespace matrix +} // namespace batch +} // namespace gko diff --git a/core/matrix/batch_ell_kernels.hpp b/core/matrix/batch_ell_kernels.hpp new file mode 100644 index 00000000000..d3acc582f9b --- /dev/null +++ b/core/matrix/batch_ell_kernels.hpp @@ -0,0 +1,84 @@ +/************************************************************* +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. +*************************************************************/ + +#ifndef GKO_CORE_MATRIX_BATCH_ELL_KERNELS_HPP_ +#define GKO_CORE_MATRIX_BATCH_ELL_KERNELS_HPP_ + + +#include + + +#include +#include +#include + + +#include "core/base/kernel_declaration.hpp" + + +namespace gko { +namespace kernels { + + +#define GKO_DECLARE_BATCH_ELL_SIMPLE_APPLY_KERNEL(_vtype, _itype) \ + void simple_apply(std::shared_ptr exec, \ + const batch::matrix::Ell<_vtype, _itype>* a, \ + const batch::MultiVector<_vtype>* b, \ + batch::MultiVector<_vtype>* c) + +#define GKO_DECLARE_BATCH_ELL_ADVANCED_APPLY_KERNEL(_vtype, _itype) \ + void advanced_apply(std::shared_ptr exec, \ + const batch::MultiVector<_vtype>* alpha, \ + const batch::matrix::Ell<_vtype, _itype>* a, \ + const batch::MultiVector<_vtype>* b, \ + const batch::MultiVector<_vtype>* beta, \ + batch::MultiVector<_vtype>* c) + +#define GKO_DECLARE_ALL_AS_TEMPLATES \ + template \ + GKO_DECLARE_BATCH_ELL_SIMPLE_APPLY_KERNEL(ValueType, IndexType); \ + template \ + GKO_DECLARE_BATCH_ELL_ADVANCED_APPLY_KERNEL(ValueType, IndexType) + + +GKO_DECLARE_FOR_ALL_EXECUTOR_NAMESPACES(batch_ell, + GKO_DECLARE_ALL_AS_TEMPLATES); + + +#undef GKO_DECLARE_ALL_AS_TEMPLATES + + +} // namespace kernels +} // namespace gko + + +#endif // GKO_CORE_MATRIX_BATCH_ELL_KERNELS_HPP_ diff --git a/core/matrix/batch_struct.hpp b/core/matrix/batch_struct.hpp index 0bbfde40cc9..f208f5ff078 100644 --- a/core/matrix/batch_struct.hpp +++ b/core/matrix/batch_struct.hpp @@ -37,6 +37,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include namespace gko { @@ -82,6 +83,53 @@ struct uniform_batch { } // namespace dense +namespace ell { + + +/** + * Encapsulates one matrix from a batch of ell matrices. + */ +template +struct batch_item { + using value_type = ValueType; + using index_type = IndexType; + + ValueType* values; + const index_type* col_idxs; + index_type stride; + index_type num_rows; + index_type num_cols; + index_type num_stored_elems_per_row; +}; + + +/** + * A 'simple' structure to store a global uniform batch of ell matrices. + */ +template +struct uniform_batch { + using value_type = ValueType; + using index_type = IndexType; + using entry_type = batch_item; + + ValueType* values; + const index_type* col_idxs; + size_type num_batch_items; + index_type stride; + index_type num_rows; + index_type num_cols; + index_type num_stored_elems_per_row; + + size_type get_entry_storage() const + { + return num_rows * num_stored_elems_per_row * sizeof(value_type); + } +}; + + +} // namespace ell + + template GKO_ATTRIBUTES GKO_INLINE dense::batch_item to_const( const dense::batch_item& b) @@ -116,6 +164,54 @@ GKO_ATTRIBUTES GKO_INLINE dense::batch_item extract_batch_item( } +template +GKO_ATTRIBUTES GKO_INLINE ell::batch_item to_const( + const ell::batch_item& b) +{ + return {b.values, b.col_idxs, b.stride, + b.num_rows, b.num_cols, b.num_stored_elems_per_row}; +} + + +template +GKO_ATTRIBUTES GKO_INLINE ell::uniform_batch +to_const(const ell::uniform_batch& ub) +{ + return {ub.values, ub.col_idxs, ub.num_batch_items, ub.stride, + ub.num_rows, ub.num_cols, ub.num_stored_elems_per_row}; +} + + +template +GKO_ATTRIBUTES GKO_INLINE ell::batch_item +extract_batch_item(const ell::uniform_batch& batch, + const size_type batch_idx) +{ + return {batch.values + + batch_idx * batch.num_stored_elems_per_row * batch.num_rows, + batch.col_idxs, + batch.stride, + batch.num_rows, + batch.num_cols, + batch.num_stored_elems_per_row}; +} + +template +GKO_ATTRIBUTES GKO_INLINE ell::batch_item +extract_batch_item(ValueType* const batch_values, + IndexType* const batch_col_idxs, const int stride, + const int num_rows, const int num_cols, + int num_elems_per_row, const size_type batch_idx) +{ + return {batch_values + batch_idx * num_elems_per_row * num_rows, + batch_col_idxs, + stride, + num_rows, + num_cols, + num_elems_per_row}; +} + + } // namespace matrix } // namespace batch } // namespace gko diff --git a/core/test/matrix/CMakeLists.txt b/core/test/matrix/CMakeLists.txt index cca4b8da1c0..ec7ef93e517 100644 --- a/core/test/matrix/CMakeLists.txt +++ b/core/test/matrix/CMakeLists.txt @@ -1,4 +1,5 @@ ginkgo_create_test(batch_dense) +ginkgo_create_test(batch_ell) ginkgo_create_test(coo) ginkgo_create_test(coo_builder) ginkgo_create_test(csr) diff --git a/core/test/matrix/batch_ell.cpp b/core/test/matrix/batch_ell.cpp new file mode 100644 index 00000000000..2c8166aa023 --- /dev/null +++ b/core/test/matrix/batch_ell.cpp @@ -0,0 +1,525 @@ +/************************************************************* +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 "core/base/batch_utilities.hpp" +#include "core/test/utils.hpp" +#include "core/test/utils/batch_helpers.hpp" + + +template +class Ell : public ::testing::Test { +protected: + using value_type = T; + using index_type = gko::int32; + using BatchEllMtx = gko::batch::matrix::Ell; + using EllMtx = gko::matrix::Ell; + using size_type = gko::size_type; + Ell() + : exec(gko::ReferenceExecutor::create()), + mtx(gko::batch::initialize( + {{{-1.0, 2.0, 3.0}, {-1.5, 2.5, 3.5}}, + {{1.0, 2.5, 3.0}, {1.0, 2.0, 3.0}}}, + exec, 3)), + sp_mtx(gko::batch::initialize( + {{{-1.0, 0.0, 0.0}, {0.0, 2.5, 3.5}}, + {{1.0, 0.0, 0.0}, {0.0, 2.0, 3.0}}}, + exec, 2)), + ell_mtx(gko::initialize({{1.0, 2.5, 3.0}, {1.0, 2.0, 3.0}}, + exec, gko::dim<2>(2, 3), 3)), + sp_ell_mtx(gko::initialize({{1.0, 0.0, 0.0}, {0.0, 2.0, 3.0}}, + exec, gko::dim<2>(2, 3), 2)) + {} + + static void assert_equal_to_original_sparse_mtx(const BatchEllMtx* m) + { + ASSERT_EQ(m->get_num_batch_items(), 2); + ASSERT_EQ(m->get_common_size(), gko::dim<2>(2, 3)); + ASSERT_EQ(m->get_num_stored_elements(), 2 * (2 * 2)); + ASSERT_EQ(m->get_num_stored_elements_per_row(), 2); + EXPECT_EQ(m->get_const_values()[0], value_type{-1.0}); + EXPECT_EQ(m->get_const_values()[1], value_type{2.5}); + EXPECT_EQ(m->get_const_values()[2], value_type{0.0}); + EXPECT_EQ(m->get_const_values()[3], value_type{3.5}); + EXPECT_EQ(m->get_const_values()[4], value_type{1.0}); + EXPECT_EQ(m->get_const_values()[5], value_type{2.0}); + EXPECT_EQ(m->get_const_values()[6], value_type{0.0}); + EXPECT_EQ(m->get_const_values()[7], value_type{3.0}); + EXPECT_EQ(m->get_const_col_idxs()[0], index_type{0}); + EXPECT_EQ(m->get_const_col_idxs()[1], index_type{1}); + EXPECT_EQ(m->get_const_col_idxs()[2], index_type{-1}); + ASSERT_EQ(m->get_const_col_idxs()[3], index_type{2}); + } + + static void assert_equal_to_original_mtx(const BatchEllMtx* m) + { + ASSERT_EQ(m->get_num_batch_items(), 2); + ASSERT_EQ(m->get_common_size(), gko::dim<2>(2, 3)); + ASSERT_EQ(m->get_num_stored_elements(), 2 * (2 * 3)); + ASSERT_EQ(m->get_num_stored_elements_per_row(), 3); + EXPECT_EQ(m->get_const_values()[0], value_type{-1.0}); + EXPECT_EQ(m->get_const_values()[1], value_type{-1.5}); + EXPECT_EQ(m->get_const_values()[2], value_type{2.0}); + EXPECT_EQ(m->get_const_values()[3], value_type{2.5}); + EXPECT_EQ(m->get_const_values()[4], value_type{3.0}); + EXPECT_EQ(m->get_const_values()[5], value_type{3.5}); + EXPECT_EQ(m->get_const_values()[6], value_type{1.0}); + EXPECT_EQ(m->get_const_values()[7], value_type{1.0}); + EXPECT_EQ(m->get_const_values()[8], value_type{2.5}); + EXPECT_EQ(m->get_const_values()[9], value_type{2.0}); + EXPECT_EQ(m->get_const_values()[10], value_type{3.0}); + ASSERT_EQ(m->get_const_values()[11], value_type{3.0}); + } + + static void assert_empty(BatchEllMtx* m) + { + ASSERT_EQ(m->get_num_batch_items(), 0); + ASSERT_EQ(m->get_num_stored_elements(), 0); + ASSERT_EQ(m->get_num_stored_elements_per_row(), 0); + } + + std::shared_ptr exec; + std::unique_ptr mtx; + std::unique_ptr sp_mtx; + std::unique_ptr ell_mtx; + std::unique_ptr sp_ell_mtx; +}; + +TYPED_TEST_SUITE(Ell, gko::test::ValueTypes); + + +TYPED_TEST(Ell, KnowsItsSizeAndValues) +{ + this->assert_equal_to_original_mtx(this->mtx.get()); +} + + +TYPED_TEST(Ell, SparseMtxKnowsItsSizeAndValues) +{ + this->assert_equal_to_original_sparse_mtx(this->sp_mtx.get()); +} + + +TYPED_TEST(Ell, CanBeEmpty) +{ + using BatchEllMtx = typename TestFixture::BatchEllMtx; + + auto empty = BatchEllMtx::create(this->exec); + + this->assert_empty(empty.get()); + ASSERT_EQ(empty->get_const_values(), nullptr); +} + + +TYPED_TEST(Ell, CanGetValuesForEntry) +{ + using value_type = typename TestFixture::value_type; + + ASSERT_EQ(this->mtx->get_values_for_item(1)[0], value_type{1.0}); +} + + +TYPED_TEST(Ell, CanCreateEllItemView) +{ + GKO_ASSERT_MTX_NEAR(this->mtx->create_view_for_item(1), this->ell_mtx, 0.0); +} + + +TYPED_TEST(Ell, CanCreateSpEllItemView) +{ + GKO_ASSERT_MTX_NEAR(this->sp_mtx->create_view_for_item(1), this->sp_ell_mtx, + 0.0); +} + + +TYPED_TEST(Ell, CanBeCopied) +{ + using BatchEllMtx = typename TestFixture::BatchEllMtx; + + auto mtx_copy = BatchEllMtx::create(this->exec); + + mtx_copy->copy_from(this->mtx.get()); + + this->assert_equal_to_original_mtx(this->mtx.get()); + this->mtx->get_values()[0] = 7; + this->assert_equal_to_original_mtx(mtx_copy.get()); +} + + +TYPED_TEST(Ell, CanBeMoved) +{ + using BatchEllMtx = typename TestFixture::BatchEllMtx; + + auto mtx_copy = BatchEllMtx::create(this->exec); + + this->mtx->move_to(mtx_copy); + + this->assert_equal_to_original_mtx(mtx_copy.get()); +} + + +TYPED_TEST(Ell, CanBeCloned) +{ + auto mtx_clone = this->mtx->clone(); + + this->assert_equal_to_original_mtx( + dynamic_castmtx.get())>(mtx_clone.get())); +} + + +TYPED_TEST(Ell, CanBeCleared) +{ + this->mtx->clear(); + + this->assert_empty(this->mtx.get()); +} + + +TYPED_TEST(Ell, CanBeConstructedWithSize) +{ + using BatchEllMtx = typename TestFixture::BatchEllMtx; + + auto m = BatchEllMtx::create(this->exec, + gko::batch_dim<2>(2, gko::dim<2>{5, 3}), 2); + + ASSERT_EQ(m->get_num_batch_items(), 2); + ASSERT_EQ(m->get_common_size(), gko::dim<2>(5, 3)); + ASSERT_EQ(m->get_num_stored_elements_per_row(), 2); + ASSERT_EQ(m->get_num_stored_elements(), 20); +} + + +TYPED_TEST(Ell, CanBeConstructedFromExistingData) +{ + using value_type = typename TestFixture::value_type; + using index_type = typename TestFixture::index_type; + using BatchEllMtx = typename TestFixture::BatchEllMtx; + // clang-format off + value_type values[] = { + -1.0, 2.5, + 0.0, 3.5, + 1.0, 2.0, + 0.0, 3.0}; + index_type col_idxs[] = { + 0, 1, + -1, 2}; + // clang-format on + + auto m = BatchEllMtx::create( + this->exec, gko::batch_dim<2>(2, gko::dim<2>(2, 3)), 2, + gko::array::view(this->exec, 8, values), + gko::array::view(this->exec, 4, col_idxs)); + + this->assert_equal_to_original_sparse_mtx(m.get()); +} + + +TYPED_TEST(Ell, CanBeConstructedFromExistingConstData) +{ + using value_type = typename TestFixture::value_type; + using index_type = typename TestFixture::index_type; + using BatchEllMtx = typename TestFixture::BatchEllMtx; + // clang-format off + value_type values[] = { + -1.0, 2.5, + 0.0, 3.5, + 1.0, 2.0, + 0.0, 3.0}; + index_type col_idxs[] = { + 0, 1, + -1, 2}; + // clang-format on + + auto m = BatchEllMtx::create_const( + this->exec, gko::batch_dim<2>(2, gko::dim<2>(2, 3)), 2, + gko::array::const_view(this->exec, 8, values), + gko::array::const_view(this->exec, 4, col_idxs)); + + this->assert_equal_to_original_sparse_mtx(m.get()); +} + + +TYPED_TEST(Ell, CanBeConstructedFromEllMatrices) +{ + using BatchEllMtx = typename TestFixture::BatchEllMtx; + using EllMtx = typename TestFixture::EllMtx; + auto mat1 = gko::initialize({{-1.0, 0.0, 0.0}, {0.0, 2.5, 3.5}}, + this->exec); + auto mat2 = + gko::initialize({{1.0, 0.0, 0.0}, {0.0, 2.0, 3.0}}, this->exec); + + auto m = gko::batch::create_from_item( + this->exec, std::vector{mat1.get(), mat2.get()}, + mat1->get_num_stored_elements_per_row()); + + this->assert_equal_to_original_sparse_mtx(m.get()); +} + + +TYPED_TEST(Ell, CanBeConstructedFromEllMatricesByDuplication) +{ + using BatchEllMtx = typename TestFixture::BatchEllMtx; + using EllMtx = typename TestFixture::EllMtx; + auto mat1 = + gko::initialize({{1.0, 0.0, 0.0}, {0.0, 2.0, 0.0}}, this->exec); + auto bat_m = gko::batch::create_from_item( + this->exec, std::vector{mat1.get(), mat1.get(), mat1.get()}, + mat1->get_num_stored_elements_per_row()); + + auto m = gko::batch::create_from_item( + this->exec, 3, mat1.get(), mat1->get_num_stored_elements_per_row()); + + GKO_ASSERT_BATCH_MTX_NEAR(bat_m.get(), m.get(), 1e-14); +} + + +TYPED_TEST(Ell, CanBeConstructedByDuplicatingEllMatrices) +{ + using BatchEllMtx = typename TestFixture::BatchEllMtx; + using EllMtx = typename TestFixture::EllMtx; + auto mat1 = gko::initialize({{-1.0, 0.0, 0.0}, {0.0, 2.5, 0.0}}, + this->exec); + auto mat2 = + gko::initialize({{1.0, 0.0, 0.0}, {0.0, 2.0, 0.0}}, this->exec); + + auto m = gko::batch::create_from_item( + this->exec, std::vector{mat1.get(), mat2.get()}, + mat1->get_num_stored_elements_per_row()); + auto m_ref = gko::batch::create_from_item( + this->exec, + std::vector{mat1.get(), mat2.get(), mat1.get(), mat2.get(), + mat1.get(), mat2.get()}, + mat1->get_num_stored_elements_per_row()); + + auto m2 = gko::batch::duplicate( + this->exec, 3, m.get(), mat1->get_num_stored_elements_per_row()); + + GKO_ASSERT_BATCH_MTX_NEAR(m2.get(), m_ref.get(), 1e-14); +} + + +TYPED_TEST(Ell, CanBeUnbatchedIntoEllMatrices) +{ + using BatchEllMtx = typename TestFixture::BatchEllMtx; + using EllMtx = typename TestFixture::EllMtx; + auto mat1 = gko::initialize({{-1.0, 0.0, 0.0}, {0.0, 2.5, 3.5}}, + this->exec); + auto mat2 = + gko::initialize({{1.0, 0.0, 0.0}, {0.0, 2.0, 3.0}}, this->exec); + + auto ell_mats = gko::batch::unbatch(this->sp_mtx.get()); + + GKO_ASSERT_MTX_NEAR(ell_mats[0].get(), mat1.get(), 0.); + GKO_ASSERT_MTX_NEAR(ell_mats[1].get(), mat2.get(), 0.); +} + + +TYPED_TEST(Ell, CanBeListConstructed) +{ + using value_type = typename TestFixture::value_type; + using index_type = typename TestFixture::index_type; + using BatchEllMtx = typename TestFixture::BatchEllMtx; + using EllMtx = typename TestFixture::EllMtx; + + auto m = gko::batch::initialize({{0.0, -1.0}, {0.0, -5.0}}, + this->exec); + + ASSERT_EQ(m->get_num_batch_items(), 2); + ASSERT_EQ(m->get_common_size(), gko::dim<2>(2, 1)); + ASSERT_EQ(m->get_num_stored_elements(), 4); + ASSERT_EQ(m->get_num_stored_elements_per_row(), 1); + EXPECT_EQ(m->get_values()[0], value_type{0.0}); + EXPECT_EQ(m->get_values()[1], value_type{-1.0}); + EXPECT_EQ(m->get_values()[2], value_type{0.0}); + EXPECT_EQ(m->get_values()[3], value_type{-5.0}); + EXPECT_EQ(m->get_col_idxs()[0], index_type{-1}); + EXPECT_EQ(m->get_col_idxs()[1], index_type{0}); +} + + +TYPED_TEST(Ell, CanBeListConstructedByCopies) +{ + using value_type = typename TestFixture::value_type; + using index_type = typename TestFixture::index_type; + using BatchEllMtx = typename TestFixture::BatchEllMtx; + + auto m = gko::batch::initialize(2, I({0.0, -1.0}), + this->exec, 1); + + ASSERT_EQ(m->get_num_batch_items(), 2); + ASSERT_EQ(m->get_common_size(), gko::dim<2>(2, 1)); + ASSERT_EQ(m->get_num_stored_elements(), 4); + ASSERT_EQ(m->get_num_stored_elements_per_row(), 1); + EXPECT_EQ(m->get_values()[0], value_type{0.0}); + EXPECT_EQ(m->get_values()[1], value_type{-1.0}); + EXPECT_EQ(m->get_values()[2], value_type{0.0}); + EXPECT_EQ(m->get_values()[3], value_type{-1.0}); + EXPECT_EQ(m->get_col_idxs()[0], index_type{-1}); + EXPECT_EQ(m->get_col_idxs()[1], index_type{0}); +} + + +TYPED_TEST(Ell, CanBeDoubleListConstructed) +{ + using value_type = typename TestFixture::value_type; + using index_type = typename TestFixture::index_type; + using BatchEllMtx = typename TestFixture::BatchEllMtx; + using T = value_type; + + auto m = gko::batch::initialize( + // clang-format off + {{I{1.0, 0.0, 0.0}, + I{2.0, 0.0, 3.0}, + I{3.0, 6.0, 0.0}}, + {I{1.0, 0.0, 0.0}, + I{3.0, 0.0, -2.0}, + I{5.0, 8.0, 0.0}}}, + // clang-format on + this->exec, 2); + + ASSERT_EQ(m->get_num_batch_items(), 2); + ASSERT_EQ(m->get_common_size(), gko::dim<2>(3, 3)); + ASSERT_EQ(m->get_num_stored_elements(), 2 * (2 * 3)); + ASSERT_EQ(m->get_num_stored_elements_per_row(), 2); + EXPECT_EQ(m->get_values()[0], value_type{1.0}); + EXPECT_EQ(m->get_values()[1], value_type{2.0}); + EXPECT_EQ(m->get_values()[2], value_type{3.0}); + EXPECT_EQ(m->get_values()[3], value_type{0.0}); + EXPECT_EQ(m->get_values()[4], value_type{3.0}); + EXPECT_EQ(m->get_values()[5], value_type{6.0}); + EXPECT_EQ(m->get_values()[6], value_type{1.0}); + EXPECT_EQ(m->get_values()[7], value_type{3.0}); + EXPECT_EQ(m->get_values()[8], value_type{5.0}); + EXPECT_EQ(m->get_values()[9], value_type{0.0}); + EXPECT_EQ(m->get_values()[10], value_type{-2.0}); + EXPECT_EQ(m->get_values()[11], value_type{8.0}); + EXPECT_EQ(m->get_col_idxs()[0], index_type{0}); + EXPECT_EQ(m->get_col_idxs()[1], index_type{0}); + EXPECT_EQ(m->get_col_idxs()[2], index_type{0}); + EXPECT_EQ(m->get_col_idxs()[3], index_type{-1}); + EXPECT_EQ(m->get_col_idxs()[4], index_type{2}); + EXPECT_EQ(m->get_col_idxs()[5], index_type{1}); +} + + +TYPED_TEST(Ell, CanBeReadFromMatrixData) +{ + using value_type = typename TestFixture::value_type; + using index_type = typename TestFixture::index_type; + using BatchEllMtx = typename TestFixture::BatchEllMtx; + auto vec_data = std::vector>{}; + vec_data.emplace_back(gko::matrix_data( + {2, 3}, {{0, 0, -1.0}, {1, 1, 2.5}, {1, 2, 3.5}})); + vec_data.emplace_back(gko::matrix_data( + {2, 3}, {{0, 0, 1.0}, {1, 1, 2.0}, {1, 2, 3.0}})); + + auto m = gko::batch::read(this->exec, + vec_data, 2); + + this->assert_equal_to_original_sparse_mtx(m.get()); +} + + +TYPED_TEST(Ell, ThrowsForDataWithDifferentNnz) +{ + using value_type = typename TestFixture::value_type; + using index_type = typename TestFixture::index_type; + using BatchEllMtx = typename TestFixture::BatchEllMtx; + auto vec_data = std::vector>{}; + vec_data.emplace_back( + gko::matrix_data({2, 3}, { + {0, 0, -1.0}, + {1, 1, 2.5}, + {1, 2, 0.5}, + {2, 2, -3.0}, + })); + vec_data.emplace_back(gko::matrix_data( + {2, 3}, {{0, 0, 1.0}, {1, 1, 2.0}, {1, 2, 3.0}})); + + EXPECT_THROW( + gko::batch::detail::assert_same_sparsity_in_batched_data(vec_data), + gko::NotImplemented); +} + + +TYPED_TEST(Ell, ThrowsForDataWithDifferentSparsity) +{ + using value_type = typename TestFixture::value_type; + using index_type = typename TestFixture::index_type; + using BatchEllMtx = typename TestFixture::BatchEllMtx; + auto vec_data = std::vector>{}; + vec_data.emplace_back( + gko::matrix_data({2, 3}, { + {0, 0, -1.0}, + {1, 1, 2.5}, + {2, 2, -3.0}, + })); + vec_data.emplace_back(gko::matrix_data( + {2, 3}, {{0, 0, 1.0}, {1, 1, 2.0}, {1, 2, 3.0}})); + + EXPECT_THROW( + gko::batch::detail::assert_same_sparsity_in_batched_data(vec_data), + gko::NotImplemented); +} + + +TYPED_TEST(Ell, GeneratesCorrectMatrixData) +{ + using value_type = typename TestFixture::value_type; + using index_type = typename TestFixture::index_type; + using BatchEllMtx = typename TestFixture::BatchEllMtx; + using tpl = typename gko::matrix_data::nonzero_type; + + auto data = gko::batch::write( + this->sp_mtx.get()); + + ASSERT_EQ(data[0].size, gko::dim<2>(2, 3)); + ASSERT_EQ(data[0].nonzeros.size(), 3); + EXPECT_EQ(data[0].nonzeros[0], tpl(0, 0, value_type{-1.0})); + EXPECT_EQ(data[0].nonzeros[1], tpl(1, 1, value_type{2.5})); + EXPECT_EQ(data[0].nonzeros[2], tpl(1, 2, value_type{3.5})); + ASSERT_EQ(data[1].size, gko::dim<2>(2, 3)); + ASSERT_EQ(data[1].nonzeros.size(), 3); + EXPECT_EQ(data[1].nonzeros[0], tpl(0, 0, value_type{1.0})); + EXPECT_EQ(data[1].nonzeros[1], tpl(1, 1, value_type{2.0})); + EXPECT_EQ(data[1].nonzeros[2], tpl(1, 2, value_type{3.0})); +} diff --git a/core/test/utils/batch_helpers.hpp b/core/test/utils/batch_helpers.hpp index 4cf9d4973e2..5b1fa60ed36 100644 --- a/core/test/utils/batch_helpers.hpp +++ b/core/test/utils/batch_helpers.hpp @@ -82,13 +82,23 @@ std::unique_ptr generate_random_batch_matrix( auto result = MatrixType::create( exec, batch_dim<2>(num_batch_items, dim<2>(num_rows, num_cols)), std::forward(args)...); + auto sp_mat = generate_random_device_matrix_data( + num_rows, num_cols, nonzero_dist, value_dist, engine, + exec->get_master()); + auto row_idxs = gko::array::const_view( + exec->get_master(), sp_mat.get_num_elems(), + sp_mat.get_const_row_idxs()) + .copy_to_array(); + auto col_idxs = gko::array::const_view( + exec->get_master(), sp_mat.get_num_elems(), + sp_mat.get_const_col_idxs()) + .copy_to_array(); - // TODO: Need to preserve sparsity pattern across batch items for batched - // sparse matrix formats for (size_type b = 0; b < num_batch_items; b++) { auto rand_mat = - generate_random_matrix( - num_rows, num_cols, nonzero_dist, value_dist, engine, exec); + fill_random_matrix( + num_rows, num_cols, row_idxs, col_idxs, value_dist, engine, + exec); result->create_view_for_item(b)->copy_from(rand_mat.get()); } diff --git a/core/test/utils/matrix_generator.hpp b/core/test/utils/matrix_generator.hpp index 6928c5424a5..d5370c6ef6a 100644 --- a/core/test/utils/matrix_generator.hpp +++ b/core/test/utils/matrix_generator.hpp @@ -42,6 +42,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include #include #include #include @@ -54,6 +55,49 @@ namespace gko { namespace test { +/** + * Fills matrix data for a random matrix given a sparsity pattern + * + * @tparam ValueType the type for matrix values + * @tparam IndexType the type for row and column indices + * @tparam ValueDistribution type of value distribution + * @tparam Engine type of random engine + * + * @param num_rows number of rows + * @param num_cols number of columns + * @param row_idxs the row indices of the matrix + * @param col_idxs the column indices of the matrix + * @param value_dist distribution of matrix values + * @param engine a random engine + * + * @return the generated matrix_data with entries according to the given + * dimensions and nonzero count and value distributions. + */ +template +matrix_data fill_random_matrix_data( + size_type num_rows, size_type num_cols, + const gko::array& row_indices, + const gko::array& col_indices, ValueDistribution&& value_dist, + Engine&& engine) +{ + matrix_data data{gko::dim<2>{num_rows, num_cols}, {}}; + auto host_exec = row_indices.get_executor()->get_master(); + auto host_row_indices = make_temporary_clone(host_exec, &row_indices); + auto host_col_indices = make_temporary_clone(host_exec, &col_indices); + + for (int nnz = 0; nnz < row_indices.get_num_elems(); ++nnz) { + data.nonzeros.emplace_back( + host_row_indices->get_const_data()[nnz], + host_col_indices->get_const_data()[nnz], + detail::get_rand_value(value_dist, engine)); + } + + data.ensure_row_major_order(); + return data; +} + + /** * Generates matrix data for a random matrix. * @@ -156,6 +200,49 @@ generate_random_device_matrix_data(gko::size_type num_rows, } +/** + * Fills a random matrix with given sparsity pattern. + * + * @tparam MatrixType type of matrix to generate (must implement + * the interface `ReadableFromMatrixData<>` and provide + * matching `value_type` and `index_type` type aliases) + * @tparam IndexType the type for row and column indices + * @tparam ValueDistribution type of value distribution + * @tparam Engine type of random engine + * + * @param num_rows number of rows + * @param num_cols number of columns + * @param row_idxs the row indices of the matrix + * @param col_idxs the column indices of the matrix + * @param value_dist distribution of matrix values + * @param exec executor where the matrix should be allocated + * @param args additional arguments for the matrix constructor + * + * @return the unique pointer of MatrixType + */ +template , + typename IndexType = typename MatrixType::index_type, + typename ValueDistribution, typename Engine, typename... MatrixArgs> +std::unique_ptr fill_random_matrix( + size_type num_rows, size_type num_cols, + const gko::array& row_idxs, + const gko::array& col_idxs, ValueDistribution&& value_dist, + Engine&& engine, std::shared_ptr exec, MatrixArgs&&... args) +{ + using value_type = typename MatrixType::value_type; + using index_type = IndexType; + + GKO_ASSERT(row_idxs.get_num_elems() == col_idxs.get_num_elems()); + GKO_ASSERT(row_idxs.get_num_elems() <= (num_rows * num_cols)); + auto result = MatrixType::create(exec, std::forward(args)...); + result->read(fill_random_matrix_data( + num_rows, num_cols, row_idxs, col_idxs, + std::forward(value_dist), + std::forward(engine))); + return result; +} + + /** * Generates a random matrix. * @@ -163,6 +250,10 @@ generate_random_device_matrix_data(gko::size_type num_rows, * the interface `ReadableFromMatrixData<>` and provide * matching `value_type` and `index_type` type aliases) * + * @param num_rows number of rows + * @param num_cols number of columns + * @param nonzero_dist distribution of nonzeros per row + * @param value_dist distribution of matrix values * @param exec executor where the matrix should be allocated * @param args additional arguments for the matrix constructor * diff --git a/cuda/CMakeLists.txt b/cuda/CMakeLists.txt index dfa1b2177ee..f5b7932ed39 100644 --- a/cuda/CMakeLists.txt +++ b/cuda/CMakeLists.txt @@ -39,6 +39,7 @@ target_sources(ginkgo_cuda factorization/par_ilut_spgeam_kernel.cu factorization/par_ilut_sweep_kernel.cu matrix/batch_dense_kernels.cu + matrix/batch_ell_kernels.cu matrix/coo_kernels.cu ${CSR_INSTANTIATE} matrix/dense_kernels.cu diff --git a/cuda/matrix/batch_dense_kernels.cu b/cuda/matrix/batch_dense_kernels.cu index dd82e15b8cc..c693a3ae861 100644 --- a/cuda/matrix/batch_dense_kernels.cu +++ b/cuda/matrix/batch_dense_kernels.cu @@ -36,7 +36,8 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include -#include +#include +#include #include "core/base/batch_struct.hpp" diff --git a/cuda/matrix/batch_ell_kernels.cu b/cuda/matrix/batch_ell_kernels.cu new file mode 100644 index 00000000000..5cadd7755a2 --- /dev/null +++ b/cuda/matrix/batch_ell_kernels.cu @@ -0,0 +1,85 @@ +/************************************************************* +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/batch_ell_kernels.hpp" + + +#include + + +#include +#include +#include + + +#include "core/base/batch_struct.hpp" +#include "core/matrix/batch_struct.hpp" +#include "cuda/base/batch_struct.hpp" +#include "cuda/base/config.hpp" +#include "cuda/base/thrust.cuh" +#include "cuda/components/cooperative_groups.cuh" +#include "cuda/components/reduction.cuh" +#include "cuda/components/thread_ids.cuh" +#include "cuda/components/uninitialized_array.hpp" +#include "cuda/matrix/batch_struct.hpp" + + +namespace gko { +namespace kernels { +namespace cuda { +/** + * @brief The Ell matrix format namespace. + * @ref Ell + * @ingroup batch_ell + */ +namespace batch_ell { + + +constexpr auto default_block_size = 256; +constexpr int sm_oversubscription = 4; + +// clang-format off + +// NOTE: DO NOT CHANGE THE ORDERING OF THE INCLUDES + +#include "common/cuda_hip/matrix/batch_ell_kernels.hpp.inc" + + +#include "common/cuda_hip/matrix/batch_ell_kernel_launcher.hpp.inc" + +// clang-format on + + +} // namespace batch_ell +} // namespace cuda +} // namespace kernels +} // namespace gko diff --git a/cuda/matrix/batch_struct.hpp b/cuda/matrix/batch_struct.hpp index 73712a7b81b..4a2a1835961 100644 --- a/cuda/matrix/batch_struct.hpp +++ b/cuda/matrix/batch_struct.hpp @@ -38,6 +38,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include #include "core/base/batch_struct.hpp" @@ -87,6 +88,40 @@ get_batch_struct(batch::matrix::Dense* const op) } +/** + * Generates an immutable uniform batch struct from a batch of ell matrices. + */ +template +inline batch::matrix::ell::uniform_batch, IndexType> +get_batch_struct(const batch::matrix::Ell* const op) +{ + return {as_cuda_type(op->get_const_values()), + op->get_const_col_idxs(), + op->get_num_batch_items(), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[1]), + static_cast(op->get_num_stored_elements_per_row())}; +} + + +/** + * Generates a uniform batch struct from a batch of ell matrices. + */ +template +inline batch::matrix::ell::uniform_batch, IndexType> +get_batch_struct(batch::matrix::Ell* const op) +{ + return {as_cuda_type(op->get_values()), + op->get_col_idxs(), + op->get_num_batch_items(), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[1]), + static_cast(op->get_num_stored_elements_per_row())}; +} + + } // namespace cuda } // namespace kernels } // namespace gko diff --git a/dpcpp/CMakeLists.txt b/dpcpp/CMakeLists.txt index 9990496c98f..9c2e799ede9 100644 --- a/dpcpp/CMakeLists.txt +++ b/dpcpp/CMakeLists.txt @@ -37,6 +37,7 @@ target_sources(ginkgo_dpcpp factorization/par_ilut_spgeam_kernel.dp.cpp factorization/par_ilut_sweep_kernel.dp.cpp matrix/batch_dense_kernels.dp.cpp + matrix/batch_ell_kernels.dp.cpp matrix/coo_kernels.dp.cpp matrix/csr_kernels.dp.cpp matrix/fbcsr_kernels.dp.cpp diff --git a/dpcpp/matrix/batch_ell_kernels.dp.cpp b/dpcpp/matrix/batch_ell_kernels.dp.cpp new file mode 100644 index 00000000000..5a69bbd3d5d --- /dev/null +++ b/dpcpp/matrix/batch_ell_kernels.dp.cpp @@ -0,0 +1,174 @@ +/************************************************************* +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/batch_ell_kernels.hpp" + + +#include + + +#include + + +#include +#include + + +#include "core/base/batch_struct.hpp" +#include "core/matrix/batch_struct.hpp" +#include "dpcpp/base/batch_struct.hpp" +#include "dpcpp/base/dim3.dp.hpp" +#include "dpcpp/base/dpct.hpp" +#include "dpcpp/base/helper.hpp" +#include "dpcpp/components/cooperative_groups.dp.hpp" +#include "dpcpp/components/intrinsics.dp.hpp" +#include "dpcpp/components/reduction.dp.hpp" +#include "dpcpp/components/thread_ids.dp.hpp" +#include "dpcpp/matrix/batch_struct.hpp" + + +namespace gko { +namespace kernels { +namespace dpcpp { +/** + * @brief The Ell matrix format namespace. + * @ref Ell + * @ingroup batch_ell + */ +namespace batch_ell { + + +#include "dpcpp/matrix/batch_ell_kernels.hpp.inc" + + +template +void simple_apply(std::shared_ptr exec, + const batch::matrix::Ell* mat, + const batch::MultiVector* b, + batch::MultiVector* x) +{ + const size_type num_rows = mat->get_common_size()[0]; + const size_type num_cols = mat->get_common_size()[1]; + + const auto num_batch_items = mat->get_num_batch_items(); + auto device = exec->get_queue()->get_device(); + // TODO: use runtime selection of group size based on num_rows. + auto group_size = + device.get_info(); + + const dim3 block(group_size); + const dim3 grid(num_batch_items); + const auto x_ub = get_batch_struct(x); + const auto b_ub = get_batch_struct(b); + const auto mat_ub = get_batch_struct(mat); + if (b_ub.num_rhs > 1) { + GKO_NOT_IMPLEMENTED; + } + + // Launch a kernel that has nbatches blocks, each block has max group size + exec->get_queue()->submit([&](sycl::handler& cgh) { + cgh.parallel_for( + sycl_nd_range(grid, block), + [=](sycl::nd_item<3> item_ct1) + [[sycl::reqd_sub_group_size(config::warp_size)]] { + auto group = item_ct1.get_group(); + auto group_id = group.get_group_linear_id(); + const auto mat_b = + batch::matrix::extract_batch_item(mat_ub, group_id); + const auto b_b = batch::extract_batch_item(b_ub, group_id); + const auto x_b = batch::extract_batch_item(x_ub, group_id); + simple_apply_kernel(mat_b, b_b, x_b, item_ct1); + }); + }); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INT32_TYPE( + GKO_DECLARE_BATCH_ELL_SIMPLE_APPLY_KERNEL); + + +template +void advanced_apply(std::shared_ptr exec, + const batch::MultiVector* alpha, + const batch::matrix::Ell* mat, + const batch::MultiVector* b, + const batch::MultiVector* beta, + batch::MultiVector* x) +{ + const auto mat_ub = get_batch_struct(mat); + const auto b_ub = get_batch_struct(b); + const auto x_ub = get_batch_struct(x); + const auto alpha_ub = get_batch_struct(alpha); + const auto beta_ub = get_batch_struct(beta); + + if (b_ub.num_rhs > 1) { + GKO_NOT_IMPLEMENTED; + } + + const auto num_batch_items = mat_ub.num_batch_items; + auto device = exec->get_queue()->get_device(); + // TODO: use runtime selection of group size based on num_rows. + auto group_size = + device.get_info(); + + const dim3 block(group_size); + const dim3 grid(num_batch_items); + + // Launch a kernel that has nbatches blocks, each block has max group size + exec->get_queue()->submit([&](sycl::handler& cgh) { + cgh.parallel_for( + sycl_nd_range(grid, block), + [=](sycl::nd_item<3> item_ct1) + [[sycl::reqd_sub_group_size(config::warp_size)]] { + auto group = item_ct1.get_group(); + auto group_id = group.get_group_linear_id(); + const auto mat_b = + batch::matrix::extract_batch_item(mat_ub, group_id); + const auto b_b = batch::extract_batch_item(b_ub, group_id); + const auto x_b = batch::extract_batch_item(x_ub, group_id); + const auto alpha_b = + batch::extract_batch_item(alpha_ub, group_id); + const auto beta_b = + batch::extract_batch_item(beta_ub, group_id); + advanced_apply_kernel(alpha_b, mat_b, b_b, beta_b, x_b, + item_ct1); + }); + }); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INT32_TYPE( + GKO_DECLARE_BATCH_ELL_ADVANCED_APPLY_KERNEL); + + +} // namespace batch_ell +} // namespace dpcpp +} // namespace kernels +} // namespace gko diff --git a/dpcpp/matrix/batch_ell_kernels.hpp.inc b/dpcpp/matrix/batch_ell_kernels.hpp.inc new file mode 100644 index 00000000000..64d71710dbb --- /dev/null +++ b/dpcpp/matrix/batch_ell_kernels.hpp.inc @@ -0,0 +1,81 @@ +/************************************************************* +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. +*************************************************************/ + +template +__dpct_inline__ void simple_apply_kernel( + const gko::batch::matrix::ell::batch_item& mat, + const gko::batch::multi_vector::batch_item& b, + const gko::batch::multi_vector::batch_item& x, + sycl::nd_item<3>& item_ct1) +{ + for (int tidx = item_ct1.get_local_linear_id(); tidx < mat.num_rows; + tidx += item_ct1.get_local_range().size()) { + auto temp = zero(); + for (size_type idx = 0; idx < mat.num_stored_elems_per_row; idx++) { + const auto col_idx = mat.col_idxs[tidx + idx * mat.stride]; + if (col_idx == invalid_index()) { + break; + } else { + temp += mat.values[tidx + idx * mat.stride] * + b.values[col_idx * b.stride]; + } + } + x.values[tidx * x.stride] = temp; + } +} + + +template +__dpct_inline__ void advanced_apply_kernel( + const gko::batch::multi_vector::batch_item& alpha, + const gko::batch::matrix::ell::batch_item& mat, + const gko::batch::multi_vector::batch_item& b, + const gko::batch::multi_vector::batch_item& beta, + const gko::batch::multi_vector::batch_item& x, + sycl::nd_item<3>& item_ct1) +{ + for (int tidx = item_ct1.get_local_linear_id(); tidx < mat.num_rows; + tidx += item_ct1.get_local_range().size()) { + auto temp = zero(); + for (size_type idx = 0; idx < mat.num_stored_elems_per_row; idx++) { + const auto col_idx = mat.col_idxs[tidx + idx * mat.stride]; + if (col_idx == invalid_index()) { + break; + } else { + temp += mat.values[tidx + idx * mat.stride] * + b.values[col_idx * b.stride]; + } + } + x.values[tidx * x.stride] = + alpha.values[0] * temp + beta.values[0] * x.values[tidx * x.stride]; + } +} diff --git a/dpcpp/matrix/batch_struct.hpp b/dpcpp/matrix/batch_struct.hpp index b0393daf55d..fe04407d82d 100644 --- a/dpcpp/matrix/batch_struct.hpp +++ b/dpcpp/matrix/batch_struct.hpp @@ -38,6 +38,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include #include "core/base/batch_struct.hpp" @@ -86,6 +87,40 @@ inline batch::matrix::dense::uniform_batch get_batch_struct( } +/** + * Generates an immutable uniform batch struct from a batch of ell matrices. + */ +template +inline batch::matrix::ell::uniform_batch +get_batch_struct(const batch::matrix::Ell* const op) +{ + return {op->get_const_values(), + op->get_const_col_idxs(), + op->get_num_batch_items(), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[1]), + static_cast(op->get_num_stored_elements_per_row())}; +} + + +/** + * Generates a uniform batch struct from a batch of ell matrices. + */ +template +inline batch::matrix::ell::uniform_batch get_batch_struct( + batch::matrix::Ell* const op) +{ + return {op->get_values(), + op->get_col_idxs(), + op->get_num_batch_items(), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[1]), + static_cast(op->get_num_stored_elements_per_row())}; +} + + } // namespace dpcpp } // namespace kernels } // namespace gko diff --git a/hip/CMakeLists.txt b/hip/CMakeLists.txt index 21b573b6cd0..ccc88769a4e 100644 --- a/hip/CMakeLists.txt +++ b/hip/CMakeLists.txt @@ -36,6 +36,7 @@ set(GINKGO_HIP_SOURCES factorization/par_ilut_spgeam_kernel.hip.cpp factorization/par_ilut_sweep_kernel.hip.cpp matrix/batch_dense_kernels.hip.cpp + matrix/batch_ell_kernels.hip.cpp matrix/coo_kernels.hip.cpp ${CSR_INSTANTIATE} matrix/dense_kernels.hip.cpp diff --git a/hip/matrix/batch_dense_kernels.hip.cpp b/hip/matrix/batch_dense_kernels.hip.cpp index eb3da83760a..3361feeb8b8 100644 --- a/hip/matrix/batch_dense_kernels.hip.cpp +++ b/hip/matrix/batch_dense_kernels.hip.cpp @@ -38,7 +38,6 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include -#include #include "core/base/batch_struct.hpp" diff --git a/hip/matrix/batch_ell_kernels.hip.cpp b/hip/matrix/batch_ell_kernels.hip.cpp new file mode 100644 index 00000000000..96e7cdb298e --- /dev/null +++ b/hip/matrix/batch_ell_kernels.hip.cpp @@ -0,0 +1,86 @@ +/************************************************************* +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/batch_ell_kernels.hpp" + + +#include +#include + + +#include +#include +#include + + +#include "core/base/batch_struct.hpp" +#include "core/matrix/batch_struct.hpp" +#include "hip/base/batch_struct.hip.hpp" +#include "hip/base/config.hip.hpp" +#include "hip/base/thrust.hip.hpp" +#include "hip/components/cooperative_groups.hip.hpp" +#include "hip/components/reduction.hip.hpp" +#include "hip/components/thread_ids.hip.hpp" +#include "hip/components/uninitialized_array.hip.hpp" +#include "hip/matrix/batch_struct.hip.hpp" + + +namespace gko { +namespace kernels { +namespace hip { +/** + * @brief The Ell matrix format namespace. + * @ref Ell + * @ingroup batch_ell + */ +namespace batch_ell { + + +constexpr auto default_block_size = 256; +constexpr int sm_oversubscription = 4; + +// clang-format off + +// NOTE: DO NOT CHANGE THE ORDERING OF THE INCLUDES + +#include "common/cuda_hip/matrix/batch_ell_kernels.hpp.inc" + + +#include "common/cuda_hip/matrix/batch_ell_kernel_launcher.hpp.inc" + +// clang-format on + + +} // namespace batch_ell +} // namespace hip +} // namespace kernels +} // namespace gko diff --git a/hip/matrix/batch_struct.hip.hpp b/hip/matrix/batch_struct.hip.hpp index 4670cf0988b..e35f13f1249 100644 --- a/hip/matrix/batch_struct.hip.hpp +++ b/hip/matrix/batch_struct.hip.hpp @@ -38,6 +38,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include #include "core/base/batch_struct.hpp" @@ -87,6 +88,40 @@ get_batch_struct(batch::matrix::Dense* const op) } +/** + * Generates an immutable uniform batch struct from a batch of ell matrices. + */ +template +inline batch::matrix::ell::uniform_batch, IndexType> +get_batch_struct(const batch::matrix::Ell* const op) +{ + return {as_hip_type(op->get_const_values()), + op->get_const_col_idxs(), + op->get_num_batch_items(), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[1]), + static_cast(op->get_num_stored_elements_per_row())}; +} + + +/** + * Generates a uniform batch struct from a batch of ell matrices. + */ +template +inline batch::matrix::ell::uniform_batch, IndexType> +get_batch_struct(batch::matrix::Ell* const op) +{ + return {as_hip_type(op->get_values()), + op->get_col_idxs(), + op->get_num_batch_items(), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[1]), + static_cast(op->get_num_stored_elements_per_row())}; +} + + } // namespace hip } // namespace kernels } // namespace gko diff --git a/include/ginkgo/core/base/batch_multi_vector.hpp b/include/ginkgo/core/base/batch_multi_vector.hpp index 61dffba3193..405603269ff 100644 --- a/include/ginkgo/core/base/batch_multi_vector.hpp +++ b/include/ginkgo/core/base/batch_multi_vector.hpp @@ -52,14 +52,6 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. namespace gko { namespace batch { -namespace matrix { - - -template -class Dense; - - -} /** @@ -90,21 +82,17 @@ class MultiVector : public EnablePolymorphicObject>, public EnablePolymorphicAssignment>, public EnableCreateMethod>, - public ConvertibleTo>>, - public ConvertibleTo> { + public ConvertibleTo>> { friend class EnableCreateMethod; friend class EnablePolymorphicObject; friend class MultiVector>; friend class MultiVector>; - friend class matrix::Dense; public: using EnablePolymorphicAssignment::convert_to; using EnablePolymorphicAssignment::move_to; using ConvertibleTo>>::convert_to; using ConvertibleTo>>::move_to; - using ConvertibleTo>::convert_to; - using ConvertibleTo>::move_to; using value_type = ValueType; using index_type = int32; @@ -126,10 +114,6 @@ class MultiVector void move_to(MultiVector>* result) override; - void convert_to(matrix::Dense* result) const override; - - void move_to(matrix::Dense* result) override; - /** * Creates a mutable view (of matrix::Dense type) of one item of the Batch * MultiVector object. Does not perform any deep copies, but only returns a @@ -462,224 +446,6 @@ class MultiVector }; -/** - * Creates and initializes a batch of single column-vectors. - * - * This function first creates a temporary MultiVector, fills it with - * passed in values, and then converts the vector to the requested type. - * - * @tparam Matrix matrix type to initialize - * (MultiVector has to implement the ConvertibleTo - * interface) - * @tparam TArgs argument types for Matrix::create method - * (not including the implied Executor as the first argument) - * - * @param vals values used to initialize the batch vector - * @param exec Executor associated to the vector - * @param create_args additional arguments passed to Matrix::create, not - * including the Executor, which is passed as the first - * argument - * - * @ingroup MultiVector - * @ingroup mat_formats - */ -template -std::unique_ptr initialize( - std::initializer_list> - vals, - std::shared_ptr exec, TArgs&&... create_args) -{ - using batch_multi_vector = MultiVector; - size_type num_batch_items = vals.size(); - GKO_THROW_IF_INVALID(num_batch_items > 0, "Input data is empty"); - auto vals_begin = begin(vals); - size_type common_num_rows = vals_begin ? vals_begin->size() : 0; - auto common_size = dim<2>(common_num_rows, 1); - for (auto& val : vals) { - GKO_ASSERT_EQ(common_num_rows, val.size()); - } - auto b_size = batch_dim<2>(num_batch_items, common_size); - auto tmp = batch_multi_vector::create(exec->get_master(), b_size); - size_type batch = 0; - for (const auto& b : vals) { - size_type idx = 0; - for (const auto& elem : b) { - tmp->at(batch, idx) = elem; - ++idx; - } - ++batch; - } - auto mtx = Matrix::create(exec, std::forward(create_args)...); - tmp->move_to(mtx); - return mtx; -} - - -/** - * Creates and initializes a batch of multi-vectors. - * - * This function first creates a temporary MultiVector, fills it with - * passed in values, and then converts the vector to the requested type. - * - * @tparam Matrix matrix type to initialize - * (Dense has to implement the ConvertibleTo interface) - * @tparam TArgs argument types for Matrix::create method - * (not including the implied Executor as the first argument) - * - * @param vals values used to initialize the vector - * @param exec Executor associated to the vector - * @param create_args additional arguments passed to Matrix::create, not - * including the Executor, which is passed as the first - * argument - * - * @ingroup MultiVector - * @ingroup mat_formats - */ -template -std::unique_ptr initialize( - std::initializer_list>> - vals, - std::shared_ptr exec, TArgs&&... create_args) -{ - using batch_multi_vector = MultiVector; - size_type num_batch_items = vals.size(); - GKO_THROW_IF_INVALID(num_batch_items > 0, "Input data is empty"); - auto vals_begin = begin(vals); - size_type common_num_rows = vals_begin ? vals_begin->size() : 0; - size_type common_num_cols = - vals_begin->begin() ? vals_begin->begin()->size() : 0; - auto common_size = dim<2>(common_num_rows, common_num_cols); - for (const auto& b : vals) { - auto num_rows = b.size(); - auto num_cols = begin(b)->size(); - auto b_size = dim<2>(num_rows, num_cols); - GKO_ASSERT_EQUAL_DIMENSIONS(b_size, common_size); - } - - auto b_size = batch_dim<2>(num_batch_items, common_size); - auto tmp = batch_multi_vector::create(exec->get_master(), b_size); - size_type batch = 0; - for (const auto& b : vals) { - size_type ridx = 0; - for (const auto& row : b) { - size_type cidx = 0; - for (const auto& elem : row) { - tmp->at(batch, ridx, cidx) = elem; - ++cidx; - } - ++ridx; - } - ++batch; - } - auto mtx = Matrix::create(exec, std::forward(create_args)...); - tmp->move_to(mtx); - return mtx; -} - - -/** - * Creates and initializes a batch single column-vector by making copies of the - * single input column vector. - * - * This function first creates a temporary batch multi-vector, fills it with - * passed in values, and then converts the vector to the requested type. - * - * @tparam Matrix matrix type to initialize - * (MultiVector has to implement the ConvertibleTo - * interface) - * @tparam TArgs argument types for Matrix::create method - * (not including the implied Executor as the first argument) - * - * @param num_vectors The number of times the input vector is to be duplicated - * @param vals values used to initialize each vector in the temp. batch - * @param exec Executor associated to the vector - * @param create_args additional arguments passed to Matrix::create, not - * including the Executor, which is passed as the first - * argument - * - * @ingroup MultiVector - * @ingroup mat_formats - */ -template -std::unique_ptr initialize( - const size_type num_vectors, - std::initializer_list vals, - std::shared_ptr exec, TArgs&&... create_args) -{ - using batch_multi_vector = MultiVector; - size_type num_batch_items = num_vectors; - GKO_THROW_IF_INVALID(num_batch_items > 0 && vals.size() > 0, - "Input data is empty"); - auto b_size = - batch_dim<2>(num_batch_items, dim<2>(begin(vals) ? vals.size() : 0, 1)); - auto tmp = batch_multi_vector::create(exec->get_master(), b_size); - for (size_type batch = 0; batch < num_vectors; batch++) { - size_type idx = 0; - for (const auto& elem : vals) { - tmp->at(batch, idx) = elem; - ++idx; - } - } - auto mtx = Matrix::create(exec, std::forward(create_args)...); - tmp->move_to(mtx); - return mtx; -} - - -/** - * Creates and initializes a matrix from copies of a given matrix. - * - * This function first creates a temporary batch multi-vector, fills it with - * passed in values, and then converts the vector to the requested type. - * - * @tparam Matrix matrix type to initialize - * (MultiVector has to implement the ConvertibleTo - * interface) - * @tparam TArgs argument types for Matrix::create method - * (not including the implied Executor as the first argument) - * - * @param num_batch_items The number of times the input matrix is duplicated - * @param vals values used to initialize each vector in the temp. batch - * @param exec Executor associated to the vector - * @param create_args additional arguments passed to Matrix::create, not - * including the Executor, which is passed as the first - * argument - * - * @ingroup LinOp - * @ingroup mat_formats - */ -template -std::unique_ptr initialize( - const size_type num_batch_items, - std::initializer_list> - vals, - std::shared_ptr exec, TArgs&&... create_args) -{ - using batch_multi_vector = MultiVector; - GKO_THROW_IF_INVALID(num_batch_items > 0 && vals.size() > 0, - "Input data is empty"); - auto common_size = dim<2>(begin(vals) ? vals.size() : 0, - begin(vals) ? begin(vals)->size() : 0); - batch_dim<2> b_size(num_batch_items, common_size); - auto tmp = batch_multi_vector::create(exec->get_master(), b_size); - for (size_type batch = 0; batch < num_batch_items; batch++) { - size_type ridx = 0; - for (const auto& row : vals) { - size_type cidx = 0; - for (const auto& elem : row) { - tmp->at(batch, ridx, cidx) = elem; - ++cidx; - } - ++ridx; - } - } - auto mtx = Matrix::create(exec, std::forward(create_args)...); - tmp->move_to(mtx); - return mtx; -} - - } // namespace batch } // namespace gko diff --git a/include/ginkgo/core/base/types.hpp b/include/ginkgo/core/base/types.hpp index 68b5da6e3eb..f5a75c7448e 100644 --- a/include/ginkgo/core/base/types.hpp +++ b/include/ginkgo/core/base/types.hpp @@ -531,6 +531,22 @@ GKO_ATTRIBUTES constexpr bool operator!=(precision_reduction x, template _macro(double, int64) #endif +#if GINKGO_DPCPP_SINGLE_MODE +#define GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INT32_TYPE(_macro) \ + template _macro(float, int32); \ + template <> \ + _macro(double, int32) GKO_NOT_IMPLEMENTED; \ + template _macro(std::complex, int32); \ + template <> \ + _macro(std::complex, int32) GKO_NOT_IMPLEMENTED +#else +#define GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INT32_TYPE(_macro) \ + template _macro(float, int32); \ + template _macro(double, int32); \ + template _macro(std::complex, int32); \ + template _macro(std::complex, int32) +#endif + /** * Instantiates a template for each value and index type compiled by Ginkgo. diff --git a/include/ginkgo/core/matrix/batch_dense.hpp b/include/ginkgo/core/matrix/batch_dense.hpp index 7f3ce5890e4..47230c24e32 100644 --- a/include/ginkgo/core/matrix/batch_dense.hpp +++ b/include/ginkgo/core/matrix/batch_dense.hpp @@ -93,15 +93,6 @@ class Dense final : public EnableBatchLinOp>, using absolute_type = remove_complex; using complex_type = to_complex; - /** - * Creates a Dense matrix with the configuration of another Dense - * matrix. - * - * @param other The other matrix whose configuration needs to copied. - */ - static std::unique_ptr create_with_config_of( - ptr_param other); - void convert_to(Dense>* result) const override; void move_to(Dense>* result) override; @@ -275,11 +266,8 @@ class Dense final : public EnableBatchLinOp>, * @param b the multi-vector to be applied to * @param x the output multi-vector */ - void apply(const MultiVector* b, - MultiVector* x) const - { - this->apply_impl(b, x); - } + Dense* apply(ptr_param> b, + ptr_param> x); /** * Apply the matrix to a multi-vector with a linear combination of the given @@ -291,13 +279,26 @@ class Dense final : public EnableBatchLinOp>, * @param beta the scalar to scale the x vector with * @param x the output multi-vector */ - void apply(const MultiVector* alpha, - const MultiVector* b, - const MultiVector* beta, - MultiVector* x) const - { - this->apply_impl(alpha, b, beta, x); - } + Dense* apply(ptr_param> alpha, + ptr_param> b, + ptr_param> beta, + ptr_param> x); + + /** + * @copydoc apply(const MultiVector*, MultiVector*) + */ + const Dense* apply(ptr_param> b, + ptr_param> x) const; + + /** + * @copydoc apply(const MultiVector*, const + * MultiVector*, const MultiVector*, + * MultiVector*) + */ + const Dense* apply(ptr_param> alpha, + ptr_param> b, + ptr_param> beta, + ptr_param> x) const; private: inline size_type compute_num_elems(const batch_dim<2>& size) @@ -306,7 +307,6 @@ class Dense final : public EnableBatchLinOp>, size.get_common_size()[1]; } -protected: /** * Creates an uninitialized Dense matrix of the specified size. * @@ -362,7 +362,6 @@ class Dense final : public EnableBatchLinOp>, idx % this->get_common_size()[1]); } -private: array values_; }; diff --git a/include/ginkgo/core/matrix/batch_ell.hpp b/include/ginkgo/core/matrix/batch_ell.hpp new file mode 100644 index 00000000000..fa00a0631fd --- /dev/null +++ b/include/ginkgo/core/matrix/batch_ell.hpp @@ -0,0 +1,385 @@ +/************************************************************* +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. +*************************************************************/ + +#ifndef GKO_PUBLIC_CORE_MATRIX_BATCH_ELL_HPP_ +#define GKO_PUBLIC_CORE_MATRIX_BATCH_ELL_HPP_ + + +#include +#include + + +#include +#include +#include +#include +#include +#include +#include +#include +#include + + +namespace gko { +namespace batch { +namespace matrix { + + +/** + * Ell is a sparse matrix format that stores the same number of nonzeros in each + * row, enabling coalesced accesses. It is suitable for sparsity patterns that + * have a similar number of nonzeros in every row. The values are stored in a + * column-major fashion similar to the monolithic gko::matrix::Ell class. + * + * Similar to the monolithic gko::matrix::Ell class, invalid_index is + * used as the column index for padded zero entries. + * + * @note It is also assumed that the sparsity pattern of all the items in the + * batch is the same and therefore only a single copy of the sparsity pattern is + * stored. + * + * @note Currently only IndexType of int32 is supported. + * + * @tparam ValueType value precision of matrix elements + * @tparam IndexType index precision of matrix elements + * + * @ingroup batch_ell + * @ingroup mat_formats + * @ingroup BatchLinOp + */ +template +class Ell final + : public EnableBatchLinOp>, + public EnableCreateMethod>, + public ConvertibleTo, IndexType>> { + friend class EnableCreateMethod; + friend class EnablePolymorphicObject; + friend class Ell, IndexType>; + friend class Ell, IndexType>; + static_assert(std::is_same::value, + "IndexType must be a 32 bit integer"); + +public: + using EnableBatchLinOp::convert_to; + using EnableBatchLinOp::move_to; + + using value_type = ValueType; + using index_type = IndexType; + using unbatch_type = gko::matrix::Ell; + using absolute_type = remove_complex; + using complex_type = to_complex; + + void convert_to( + Ell, IndexType>* result) const override; + + void move_to(Ell, IndexType>* result) override; + + /** + * Creates a mutable view (of matrix::Ell type) of one item of the + * batch::matrix::Ell object. Does not perform any deep + * copies, but only returns a view of the data. + * + * @param item_id The index of the batch item + * + * @return a batch::matrix::Ell object with the data from the batch item + * at the given index. + */ + std::unique_ptr create_view_for_item(size_type item_id); + + /** + * @copydoc create_view_for_item(size_type) + */ + std::unique_ptr create_const_view_for_item( + size_type item_id) const; + + /** + * Returns a pointer to the array of values of the matrix + * + * @return the pointer to the array of values + */ + value_type* get_values() noexcept { return values_.get_data(); } + + /** + * @copydoc get_values() + * + * @note This is the constant version of the function, which can be + * significantly more memory efficient than the non-constant version, + * so always prefer this version. + */ + const value_type* get_const_values() const noexcept + { + return values_.get_const_data(); + } + + /** + * Returns a pointer to the array of column indices of the matrix + * + * @return the pointer to the array of column indices + */ + index_type* get_col_idxs() noexcept { return col_idxs_.get_data(); } + + /** + * @copydoc get_col_idxs() + * + * @note This is the constant version of the function, which can be + * significantly more memory efficient than the non-constant version, + * so always prefer this version. + */ + const index_type* get_const_col_idxs() const noexcept + { + return col_idxs_.get_const_data(); + } + + /** + * Returns the number of elements per row explicitly stored. + * + * @return the number of elements stored in each row of the ELL matrix. Same + * for each batch item + */ + index_type get_num_stored_elements_per_row() const noexcept + { + return num_elems_per_row_; + } + + /** + * Returns the number of elements explicitly stored in the batch matrix, + * cumulative across all the batch items. + * + * @return the number of elements explicitly stored in the vector, + * cumulative across all the batch items + */ + size_type get_num_stored_elements() const noexcept + { + return values_.get_num_elems(); + } + + /** + * Returns the number of stored elements in each batch item. + * + * @return the number of stored elements per batch item. + */ + size_type get_num_elements_per_item() const noexcept + { + return this->get_num_stored_elements() / this->get_num_batch_items(); + } + + /** + * Returns a pointer to the array of col_idxs of the matrix. This is shared + * across all batch items. + * + * @param batch_id the id of the batch item. + * + * @return the pointer to the array of col_idxs + */ + index_type* get_col_idxs_for_item(size_type batch_id) noexcept + { + GKO_ASSERT(batch_id < this->get_num_batch_items()); + return col_idxs_.get_data(); + } + + /** + * @copydoc get_col_idxs_for_item(size_type) + * + * @note This is the constant version of the function, which can be + * significantly more memory efficient than the non-constant version, + * so always prefer this version. + */ + const index_type* get_const_col_idxs_for_item( + size_type batch_id) const noexcept + { + GKO_ASSERT(batch_id < this->get_num_batch_items()); + return col_idxs_.get_const_data(); + } + + /** + * Returns a pointer to the array of values of the matrix for a + * specific batch item. + * + * @param batch_id the id of the batch item. + * + * @return the pointer to the array of values + */ + value_type* get_values_for_item(size_type batch_id) noexcept + { + GKO_ASSERT(batch_id < this->get_num_batch_items()); + return values_.get_data() + + batch_id * this->get_num_elements_per_item(); + } + + /** + * @copydoc get_values_for_item(size_type) + * + * @note This is the constant version of the function, which can be + * significantly more memory efficient than the non-constant version, + * so always prefer this version. + */ + const value_type* get_const_values_for_item( + size_type batch_id) const noexcept + { + GKO_ASSERT(batch_id < this->get_num_batch_items()); + return values_.get_const_data() + + batch_id * this->get_num_elements_per_item(); + } + + /** + * Creates a constant (immutable) batch ell matrix from a constant + * array. The column indices array needs to be the same for all batch items. + * + * @param exec the executor to create the matrix on + * @param size the dimensions of the matrix + * @param num_elems_per_row the number of elements to be stored in each row + * @param values the value array of the matrix + * @param col_idxs the col_idxs array of a single batch item of the matrix. + * + * @return A smart pointer to the constant matrix wrapping the input + * array (if it resides on the same executor as the matrix) or a copy of the + * array on the correct executor. + */ + static std::unique_ptr create_const( + std::shared_ptr exec, const batch_dim<2>& sizes, + const index_type num_elems_per_row, + gko::detail::const_array_view&& values, + gko::detail::const_array_view&& col_idxs); + + /** + * Apply the matrix to a multi-vector. Represents the matrix vector + * multiplication, x = A * b, where x and b are both multi-vectors. + * + * @param b the multi-vector to be applied to + * @param x the output multi-vector + */ + Ell* apply(ptr_param> b, + ptr_param> x); + + /** + * Apply the matrix to a multi-vector with a linear combination of the given + * input vector. Represents the matrix vector multiplication, x = alpha * A + * * b + beta * x, where x and b are both multi-vectors. + * + * @param alpha the scalar to scale the matrix-vector product with + * @param b the multi-vector to be applied to + * @param beta the scalar to scale the x vector with + * @param x the output multi-vector + */ + Ell* apply(ptr_param> alpha, + ptr_param> b, + ptr_param> beta, + ptr_param> x); + + /** + * @copydoc apply(const MultiVector*, MultiVector*) + */ + const Ell* apply(ptr_param> b, + ptr_param> x) const; + + /** + * @copydoc apply(const MultiVector*, const + * MultiVector*, const MultiVector*, + * MultiVector*) + */ + const Ell* apply(ptr_param> alpha, + ptr_param> b, + ptr_param> beta, + ptr_param> x) const; + +private: + size_type compute_num_elems(const batch_dim<2>& size, + IndexType num_elems_per_row) + { + return size.get_num_batch_items() * size.get_common_size()[0] * + num_elems_per_row; + } + + /** + * Creates an uninitialized Ell matrix of the specified size. + * + * @param exec Executor associated to the matrix + * @param size size of the matrix + * @param num_elems_per_row the number of elements to be stored in each row + */ + Ell(std::shared_ptr exec, + const batch_dim<2>& size = batch_dim<2>{}, + const IndexType num_elems_per_row = 0); + + /** + * Creates a Ell matrix from an already allocated (and initialized) + * array. The column indices array needs to be the same for all batch items. + * + * @tparam ValuesArray type of array of values + * + * @param exec Executor associated to the matrix + * @param size size of the matrix + * @param num_elems_per_row the number of elements to be stored in each row + * @param values array of matrix values + * @param col_idxs the col_idxs array of a single batch item of the matrix. + * + * @note If `values` is not an rvalue, not an array of ValueType, or is on + * the wrong executor, an internal copy will be created, and the + * original array data will not be used in the matrix. + */ + template + Ell(std::shared_ptr exec, const batch_dim<2>& size, + const IndexType num_elems_per_row, ValuesArray&& values, + IndicesArray&& col_idxs) + : EnableBatchLinOp(exec, size), + num_elems_per_row_{num_elems_per_row}, + values_{exec, std::forward(values)}, + col_idxs_{exec, std::forward(col_idxs)} + { + // Ensure that the value and col_idxs arrays have the correct size + auto num_elems = this->get_common_size()[0] * num_elems_per_row * + this->get_num_batch_items(); + GKO_ASSERT_EQ(num_elems, values_.get_num_elems()); + GKO_ASSERT_EQ(this->get_num_elements_per_item(), + col_idxs_.get_num_elems()); + } + + void apply_impl(const MultiVector* b, + MultiVector* x) const; + + void apply_impl(const MultiVector* alpha, + const MultiVector* b, + const MultiVector* beta, + MultiVector* x) const; + + index_type num_elems_per_row_; + array values_; + array col_idxs_; +}; + + +} // namespace matrix +} // namespace batch +} // namespace gko + + +#endif // GKO_PUBLIC_CORE_MATRIX_BATCH_ELL_HPP_ diff --git a/include/ginkgo/ginkgo.hpp b/include/ginkgo/ginkgo.hpp index 8bb29242e88..ad90e264189 100644 --- a/include/ginkgo/ginkgo.hpp +++ b/include/ginkgo/ginkgo.hpp @@ -109,6 +109,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include +#include #include #include #include diff --git a/omp/CMakeLists.txt b/omp/CMakeLists.txt index d87399492f5..aa8e30cd590 100644 --- a/omp/CMakeLists.txt +++ b/omp/CMakeLists.txt @@ -24,6 +24,7 @@ target_sources(ginkgo_omp factorization/par_ilu_kernels.cpp factorization/par_ilut_kernels.cpp matrix/batch_dense_kernels.cpp + matrix/batch_ell_kernels.cpp matrix/coo_kernels.cpp matrix/csr_kernels.cpp matrix/dense_kernels.cpp diff --git a/omp/matrix/batch_dense_kernels.cpp b/omp/matrix/batch_dense_kernels.cpp index 2d0b7ed4d40..b91a4133dba 100644 --- a/omp/matrix/batch_dense_kernels.cpp +++ b/omp/matrix/batch_dense_kernels.cpp @@ -36,8 +36,8 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include -#include -#include +#include +#include #include "core/base/batch_struct.hpp" diff --git a/omp/matrix/batch_ell_kernels.cpp b/omp/matrix/batch_ell_kernels.cpp new file mode 100644 index 00000000000..17710a97366 --- /dev/null +++ b/omp/matrix/batch_ell_kernels.cpp @@ -0,0 +1,117 @@ +/************************************************************* +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/batch_ell_kernels.hpp" + + +#include + + +#include +#include + + +#include "core/base/batch_struct.hpp" +#include "core/matrix/batch_struct.hpp" +#include "reference/base/batch_struct.hpp" +#include "reference/matrix/batch_struct.hpp" + + +namespace gko { +namespace kernels { +namespace omp { +/** + * @brief The Ell matrix format namespace. + * @ref Ell + * @ingroup batch_ell + */ +namespace batch_ell { + + +#include "reference/matrix/batch_ell_kernels.hpp.inc" + + +template +void simple_apply(std::shared_ptr exec, + const batch::matrix::Ell* mat, + const batch::MultiVector* b, + batch::MultiVector* x) +{ + const auto b_ub = host::get_batch_struct(b); + const auto x_ub = host::get_batch_struct(x); + const auto mat_ub = host::get_batch_struct(mat); +#pragma omp parallel for + for (size_type batch = 0; batch < x->get_num_batch_items(); ++batch) { + const auto mat_item = batch::matrix::extract_batch_item(mat_ub, batch); + const auto b_item = batch::extract_batch_item(b_ub, batch); + const auto x_item = batch::extract_batch_item(x_ub, batch); + simple_apply_kernel(mat_item, b_item, x_item); + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INT32_TYPE( + GKO_DECLARE_BATCH_ELL_SIMPLE_APPLY_KERNEL); + + +template +void advanced_apply(std::shared_ptr exec, + const batch::MultiVector* alpha, + const batch::matrix::Ell* mat, + const batch::MultiVector* b, + const batch::MultiVector* beta, + batch::MultiVector* x) +{ + const auto b_ub = host::get_batch_struct(b); + const auto x_ub = host::get_batch_struct(x); + const auto mat_ub = host::get_batch_struct(mat); + const auto alpha_ub = host::get_batch_struct(alpha); + const auto beta_ub = host::get_batch_struct(beta); +#pragma omp parallel for + for (size_type batch = 0; batch < x->get_num_batch_items(); ++batch) { + const auto mat_item = batch::matrix::extract_batch_item(mat_ub, batch); + const auto b_item = batch::extract_batch_item(b_ub, batch); + const auto x_item = batch::extract_batch_item(x_ub, batch); + const auto alpha_item = batch::extract_batch_item(alpha_ub, batch); + const auto beta_item = batch::extract_batch_item(beta_ub, batch); + advanced_apply_kernel(alpha_item.values[0], mat_item, b_item, + beta_item.values[0], x_item); + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INT32_TYPE( + GKO_DECLARE_BATCH_ELL_ADVANCED_APPLY_KERNEL); + + +} // namespace batch_ell +} // namespace omp +} // namespace kernels +} // namespace gko diff --git a/reference/CMakeLists.txt b/reference/CMakeLists.txt index 37498588ca7..21dfc0dfb5a 100644 --- a/reference/CMakeLists.txt +++ b/reference/CMakeLists.txt @@ -26,6 +26,7 @@ target_sources(ginkgo_reference factorization/par_ilu_kernels.cpp factorization/par_ilut_kernels.cpp matrix/batch_dense_kernels.cpp + matrix/batch_ell_kernels.cpp matrix/coo_kernels.cpp matrix/csr_kernels.cpp matrix/dense_kernels.cpp diff --git a/reference/matrix/batch_dense_kernels.cpp b/reference/matrix/batch_dense_kernels.cpp index 3d7ef03a3bd..87d73bb8e34 100644 --- a/reference/matrix/batch_dense_kernels.cpp +++ b/reference/matrix/batch_dense_kernels.cpp @@ -36,9 +36,8 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include -#include -#include -#include +#include +#include #include "core/base/batch_struct.hpp" diff --git a/reference/matrix/batch_ell_kernels.cpp b/reference/matrix/batch_ell_kernels.cpp new file mode 100644 index 00000000000..1d3a0e1ef94 --- /dev/null +++ b/reference/matrix/batch_ell_kernels.cpp @@ -0,0 +1,115 @@ +/************************************************************* +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/batch_ell_kernels.hpp" + + +#include + + +#include +#include + + +#include "core/base/batch_struct.hpp" +#include "core/matrix/batch_struct.hpp" +#include "reference/base/batch_struct.hpp" +#include "reference/matrix/batch_struct.hpp" + + +namespace gko { +namespace kernels { +namespace reference { +/** + * @brief The Ell matrix format namespace. + * @ref Ell + * @ingroup batch_ell + */ +namespace batch_ell { + + +#include "reference/matrix/batch_ell_kernels.hpp.inc" + + +template +void simple_apply(std::shared_ptr exec, + const batch::matrix::Ell* mat, + const batch::MultiVector* b, + batch::MultiVector* x) +{ + const auto b_ub = host::get_batch_struct(b); + const auto x_ub = host::get_batch_struct(x); + const auto mat_ub = host::get_batch_struct(mat); + for (size_type batch = 0; batch < x->get_num_batch_items(); ++batch) { + const auto mat_item = batch::matrix::extract_batch_item(mat_ub, batch); + const auto b_item = batch::extract_batch_item(b_ub, batch); + const auto x_item = batch::extract_batch_item(x_ub, batch); + simple_apply_kernel(mat_item, b_item, x_item); + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INT32_TYPE( + GKO_DECLARE_BATCH_ELL_SIMPLE_APPLY_KERNEL); + + +template +void advanced_apply(std::shared_ptr exec, + const batch::MultiVector* alpha, + const batch::matrix::Ell* mat, + const batch::MultiVector* b, + const batch::MultiVector* beta, + batch::MultiVector* x) +{ + const auto b_ub = host::get_batch_struct(b); + const auto x_ub = host::get_batch_struct(x); + const auto mat_ub = host::get_batch_struct(mat); + const auto alpha_ub = host::get_batch_struct(alpha); + const auto beta_ub = host::get_batch_struct(beta); + for (size_type batch = 0; batch < x->get_num_batch_items(); ++batch) { + const auto mat_item = batch::matrix::extract_batch_item(mat_ub, batch); + const auto b_item = batch::extract_batch_item(b_ub, batch); + const auto x_item = batch::extract_batch_item(x_ub, batch); + const auto alpha_item = batch::extract_batch_item(alpha_ub, batch); + const auto beta_item = batch::extract_batch_item(beta_ub, batch); + advanced_apply_kernel(alpha_item.values[0], mat_item, b_item, + beta_item.values[0], x_item); + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INT32_TYPE( + GKO_DECLARE_BATCH_ELL_ADVANCED_APPLY_KERNEL); + + +} // namespace batch_ell +} // namespace reference +} // namespace kernels +} // namespace gko diff --git a/reference/matrix/batch_ell_kernels.hpp.inc b/reference/matrix/batch_ell_kernels.hpp.inc new file mode 100644 index 00000000000..979df1a19bd --- /dev/null +++ b/reference/matrix/batch_ell_kernels.hpp.inc @@ -0,0 +1,80 @@ +/************************************************************* +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. +*************************************************************/ + +template +inline void simple_apply_kernel( + const gko::batch::matrix::ell::batch_item& a, + const gko::batch::multi_vector::batch_item& b, + const gko::batch::multi_vector::batch_item& c) +{ + for (int row = 0; row < c.num_rows; ++row) { + for (int j = 0; j < c.num_rhs; ++j) { + c.values[row * c.stride + j] = zero(); + } + for (auto k = 0; k < a.num_stored_elems_per_row; ++k) { + auto val = a.values[row + k * a.stride]; + auto col = a.col_idxs[row + k * a.stride]; + if (col != invalid_index()) { + for (int j = 0; j < c.num_rhs; ++j) { + c.values[row * c.stride + j] += + val * b.values[col * b.stride + j]; + } + } + } + } +} + + +template +inline void advanced_apply_kernel( + const ValueType alpha, + const gko::batch::matrix::ell::batch_item& a, + const gko::batch::multi_vector::batch_item& b, + const ValueType beta, + const gko::batch::multi_vector::batch_item& c) +{ + for (int row = 0; row < a.num_rows; ++row) { + for (int j = 0; j < c.num_rhs; ++j) { + c.values[row * c.stride + j] *= beta; + } + for (auto k = 0; k < a.num_stored_elems_per_row; ++k) { + auto val = a.values[row + k * a.stride]; + auto col = a.col_idxs[row + k * a.stride]; + if (col != invalid_index()) { + for (int j = 0; j < b.num_rhs; ++j) { + c.values[row * c.stride + j] += + alpha * val * b.values[col * b.stride + j]; + } + } + } + } +} diff --git a/reference/matrix/batch_struct.hpp b/reference/matrix/batch_struct.hpp index 483d7717718..bb7680d1493 100644 --- a/reference/matrix/batch_struct.hpp +++ b/reference/matrix/batch_struct.hpp @@ -39,6 +39,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include +#include #include "core/base/batch_struct.hpp" @@ -90,6 +91,40 @@ inline batch::matrix::dense::uniform_batch get_batch_struct( } +/** + * Generates an immutable uniform batch struct from a batch of ell matrices. + */ +template +inline batch::matrix::ell::uniform_batch +get_batch_struct(const batch::matrix::Ell* const op) +{ + return {op->get_const_values(), + op->get_const_col_idxs(), + op->get_num_batch_items(), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[1]), + static_cast(op->get_num_stored_elements_per_row())}; +} + + +/** + * Generates a uniform batch struct from a batch of ell matrices. + */ +template +inline batch::matrix::ell::uniform_batch get_batch_struct( + batch::matrix::Ell* const op) +{ + return {op->get_values(), + op->get_col_idxs(), + op->get_num_batch_items(), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[0]), + static_cast(op->get_common_size()[1]), + static_cast(op->get_num_stored_elements_per_row())}; +} + + } // namespace host } // namespace kernels } // namespace gko diff --git a/reference/test/matrix/CMakeLists.txt b/reference/test/matrix/CMakeLists.txt index 18634de662d..05498cbadc4 100644 --- a/reference/test/matrix/CMakeLists.txt +++ b/reference/test/matrix/CMakeLists.txt @@ -1,4 +1,5 @@ ginkgo_create_test(batch_dense_kernels) +ginkgo_create_test(batch_ell_kernels) ginkgo_create_test(coo_kernels) ginkgo_create_test(csr_kernels) ginkgo_create_test(dense_kernels) diff --git a/reference/test/matrix/batch_ell_kernels.cpp b/reference/test/matrix/batch_ell_kernels.cpp new file mode 100644 index 00000000000..81f189c3e02 --- /dev/null +++ b/reference/test/matrix/batch_ell_kernels.cpp @@ -0,0 +1,217 @@ +/************************************************************* +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 +#include +#include + + +#include "core/matrix/batch_ell_kernels.hpp" +#include "core/test/utils.hpp" + + +template +class Ell : public ::testing::Test { +protected: + using value_type = T; + using size_type = gko::size_type; + using BMtx = gko::batch::matrix::Ell; + using BMVec = gko::batch::MultiVector; + using EllMtx = gko::matrix::Ell; + using DenseMtx = gko::matrix::Dense; + Ell() + : exec(gko::ReferenceExecutor::create()), + mtx_0(gko::batch::initialize( + {{I({1.0, -1.0, 1.5}), I({-2.0, 2.0, 3.0})}, + {{1.0, -2.0, -0.5}, {1.0, -2.5, 4.0}}}, + exec)), + mtx_00(gko::initialize( + {I({1.0, -1.0, 1.5}), I({-2.0, 2.0, 3.0})}, exec)), + mtx_01(gko::initialize( + {I({1.0, -2.0, -0.5}), I({1.0, -2.5, 4.0})}, exec)), + b_0(gko::batch::initialize( + {{I({1.0, 0.0, 1.0}), I({2.0, 0.0, 1.0}), + I({1.0, 0.0, 2.0})}, + {I({-1.0, 1.0, 1.0}), I({1.0, -1.0, 1.0}), + I({1.0, 0.0, 2.0})}}, + exec)), + b_00(gko::initialize( + {I({1.0, 0.0, 1.0}), I({2.0, 0.0, 1.0}), + I({1.0, 0.0, 2.0})}, + exec)), + b_01(gko::initialize( + {I({-1.0, 1.0, 1.0}), I({1.0, -1.0, 1.0}), + I({1.0, 0.0, 2.0})}, + exec)), + x_0(gko::batch::initialize( + {{I({2.0, 0.0, 1.0}), I({2.0, 0.0, 2.0})}, + {I({-2.0, 1.0, 1.0}), I({1.0, -1.0, -1.0})}}, + exec)), + x_00(gko::initialize( + {I({2.0, 0.0, 1.0}), I({2.0, 0.0, 2.0})}, exec)), + x_01(gko::initialize( + {I({-2.0, 1.0, 1.0}), I({1.0, -1.0, -1.0})}, exec)) + {} + + std::shared_ptr exec; + std::unique_ptr mtx_0; + std::unique_ptr mtx_00; + std::unique_ptr mtx_01; + std::unique_ptr b_0; + std::unique_ptr b_00; + std::unique_ptr b_01; + std::unique_ptr x_0; + std::unique_ptr x_00; + std::unique_ptr x_01; + + std::ranlux48 rand_engine; +}; + + +TYPED_TEST_SUITE(Ell, gko::test::ValueTypes); + + +TYPED_TEST(Ell, AppliesToBatchMultiVector) +{ + using T = typename TestFixture::value_type; + + this->mtx_0->apply(this->b_0.get(), this->x_0.get()); + + this->mtx_00->apply(this->b_00.get(), this->x_00.get()); + this->mtx_01->apply(this->b_01.get(), this->x_01.get()); + auto res = gko::batch::unbatch>(this->x_0.get()); + GKO_ASSERT_MTX_NEAR(res[0].get(), this->x_00.get(), r::value); + GKO_ASSERT_MTX_NEAR(res[1].get(), this->x_01.get(), r::value); +} + + +TYPED_TEST(Ell, AppliesLinearCombinationToBatchMultiVector) +{ + using BMtx = typename TestFixture::BMtx; + using BMVec = typename TestFixture::BMVec; + using DenseMtx = typename TestFixture::DenseMtx; + using T = typename TestFixture::value_type; + auto alpha = gko::batch::initialize({{1.5}, {-1.0}}, this->exec); + auto beta = gko::batch::initialize({{2.5}, {-4.0}}, this->exec); + auto alpha0 = gko::initialize({1.5}, this->exec); + auto alpha1 = gko::initialize({-1.0}, this->exec); + auto beta0 = gko::initialize({2.5}, this->exec); + auto beta1 = gko::initialize({-4.0}, this->exec); + + this->mtx_0->apply(alpha.get(), this->b_0.get(), beta.get(), + this->x_0.get()); + + this->mtx_00->apply(alpha0.get(), this->b_00.get(), beta0.get(), + this->x_00.get()); + this->mtx_01->apply(alpha1.get(), this->b_01.get(), beta1.get(), + this->x_01.get()); + auto res = gko::batch::unbatch>(this->x_0.get()); + GKO_ASSERT_MTX_NEAR(res[0].get(), this->x_00.get(), r::value); + GKO_ASSERT_MTX_NEAR(res[1].get(), this->x_01.get(), r::value); +} + + +TYPED_TEST(Ell, ApplyFailsOnWrongNumberOfResultCols) +{ + using BMVec = typename TestFixture::BMVec; + auto res = BMVec::create(this->exec, gko::batch_dim<2>{2, gko::dim<2>{2}}); + + ASSERT_THROW(this->mtx_0->apply(this->b_0.get(), res.get()), + gko::DimensionMismatch); +} + + +TYPED_TEST(Ell, ApplyFailsOnWrongNumberOfResultRows) +{ + using BMVec = typename TestFixture::BMVec; + auto res = BMVec::create(this->exec, gko::batch_dim<2>{2, gko::dim<2>{3}}); + + ASSERT_THROW(this->mtx_0->apply(this->b_0.get(), res.get()), + gko::DimensionMismatch); +} + + +TYPED_TEST(Ell, ApplyFailsOnWrongInnerDimension) +{ + using BMVec = typename TestFixture::BMVec; + auto res = + BMVec::create(this->exec, gko::batch_dim<2>{2, gko::dim<2>{2, 3}}); + + ASSERT_THROW(this->mtx_0->apply(res.get(), this->x_0.get()), + gko::DimensionMismatch); +} + + +TYPED_TEST(Ell, AdvancedApplyFailsOnWrongInnerDimension) +{ + using BMVec = typename TestFixture::BMVec; + auto res = + BMVec::create(this->exec, gko::batch_dim<2>{2, gko::dim<2>{2, 3}}); + auto alpha = + BMVec::create(this->exec, gko::batch_dim<2>{2, gko::dim<2>{1, 1}}); + auto beta = + BMVec::create(this->exec, gko::batch_dim<2>{2, gko::dim<2>{1, 1}}); + + ASSERT_THROW( + this->mtx_0->apply(alpha.get(), res.get(), beta.get(), this->x_0.get()), + gko::DimensionMismatch); +} + + +TYPED_TEST(Ell, AdvancedApplyFailsOnWrongAlphaDimension) +{ + using BMVec = typename TestFixture::BMVec; + auto res = + BMVec::create(this->exec, gko::batch_dim<2>{2, gko::dim<2>{3, 3}}); + auto alpha = + BMVec::create(this->exec, gko::batch_dim<2>{2, gko::dim<2>{2, 1}}); + auto beta = + BMVec::create(this->exec, gko::batch_dim<2>{2, gko::dim<2>{1, 1}}); + + ASSERT_THROW( + this->mtx_0->apply(alpha.get(), res.get(), beta.get(), this->x_0.get()), + gko::DimensionMismatch); +} diff --git a/test/matrix/CMakeLists.txt b/test/matrix/CMakeLists.txt index 9f3b17cd858..a03a0a0bb4e 100644 --- a/test/matrix/CMakeLists.txt +++ b/test/matrix/CMakeLists.txt @@ -1,4 +1,5 @@ ginkgo_create_common_test(batch_dense_kernels) +ginkgo_create_common_test(batch_ell_kernels) ginkgo_create_common_device_test(csr_kernels) ginkgo_create_common_test(csr_kernels2) ginkgo_create_common_test(coo_kernels) diff --git a/test/matrix/batch_ell_kernels.cpp b/test/matrix/batch_ell_kernels.cpp new file mode 100644 index 00000000000..572f47ba47d --- /dev/null +++ b/test/matrix/batch_ell_kernels.cpp @@ -0,0 +1,143 @@ +/************************************************************* +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/batch_ell_kernels.hpp" + + +#include +#include + + +#include + + +#include +#include +#include +#include + + +#include "core/base/batch_utilities.hpp" +#include "core/test/utils.hpp" +#include "core/test/utils/assertions.hpp" +#include "core/test/utils/batch_helpers.hpp" +#include "test/utils/executor.hpp" + + +class Ell : public CommonTestFixture { +protected: + using BMtx = gko::batch::matrix::Ell; + using BMVec = gko::batch::MultiVector; + + Ell() : rand_engine(15) {} + + template + std::unique_ptr gen_mtx(const gko::size_type num_batch_items, + gko::size_type num_rows, + gko::size_type num_cols, + int num_elems_per_row) + { + return gko::test::generate_random_batch_matrix( + num_batch_items, num_rows, num_cols, + std::uniform_int_distribution<>(num_elems_per_row, + num_elems_per_row), + std::normal_distribution<>(-1.0, 1.0), rand_engine, ref, + num_elems_per_row); + } + + std::unique_ptr gen_mvec(const gko::size_type num_batch_items, + gko::size_type num_rows, + gko::size_type num_cols) + { + return gko::test::generate_random_batch_matrix( + num_batch_items, num_rows, num_cols, + std::uniform_int_distribution<>(num_cols, num_cols), + std::normal_distribution<>(-1.0, 1.0), rand_engine, ref); + } + + void set_up_apply_data(gko::size_type num_vecs = 1, + int num_elems_per_row = 5) + { + const int num_rows = 252; + const int num_cols = 32; + GKO_ASSERT(num_elems_per_row <= num_cols); + mat = gen_mtx(batch_size, num_rows, num_cols, num_elems_per_row); + y = gen_mvec(batch_size, num_cols, num_vecs); + alpha = gen_mvec(batch_size, 1, 1); + beta = gen_mvec(batch_size, 1, 1); + dmat = gko::clone(exec, mat); + dy = gko::clone(exec, y); + dalpha = gko::clone(exec, alpha); + dbeta = gko::clone(exec, beta); + expected = BMVec::create( + ref, + gko::batch_dim<2>(batch_size, gko::dim<2>{num_rows, num_vecs})); + expected->fill(gko::one()); + dresult = gko::clone(exec, expected); + } + + std::ranlux48 rand_engine; + + const size_t batch_size = 11; + std::unique_ptr mat; + std::unique_ptr y; + std::unique_ptr alpha; + std::unique_ptr beta; + std::unique_ptr expected; + std::unique_ptr dresult; + std::unique_ptr dmat; + std::unique_ptr dy; + std::unique_ptr dalpha; + std::unique_ptr dbeta; +}; + + +TEST_F(Ell, SingleVectorApplyIsEquivalentToRef) +{ + set_up_apply_data(1); + + mat->apply(y.get(), expected.get()); + dmat->apply(dy.get(), dresult.get()); + + GKO_ASSERT_BATCH_MTX_NEAR(dresult, expected, r::value); +} + + +TEST_F(Ell, SingleVectorAdvancedApplyIsEquivalentToRef) +{ + set_up_apply_data(1); + + mat->apply(alpha.get(), y.get(), beta.get(), expected.get()); + dmat->apply(dalpha.get(), dy.get(), dbeta.get(), dresult.get()); + + GKO_ASSERT_BATCH_MTX_NEAR(dresult, expected, r::value); +} diff --git a/test/test_install/test_install.cpp b/test/test_install/test_install.cpp index 7e53ea8f165..c00bb594ecd 100644 --- a/test/test_install/test_install.cpp +++ b/test/test_install/test_install.cpp @@ -219,13 +219,20 @@ int main() auto test = batch_multi_vector_type::create(exec); } - // core/base/batch_dense.hpp + // core/matrix/batch_dense.hpp { using type1 = float; using batch_dense_type = gko::batch::matrix::Dense; auto test = batch_dense_type::create(exec); } + // core/matrix/batch_ell.hpp + { + using type1 = float; + using batch_ell_type = gko::batch::matrix::Ell; + auto test = batch_ell_type::create(exec); + } + // core/base/combination.hpp { using type1 = int;