diff --git a/.github/workflows/windows-build.yml b/.github/workflows/windows-build.yml index c1bda6bc2f3..126a029e942 100644 --- a/.github/workflows/windows-build.yml +++ b/.github/workflows/windows-build.yml @@ -22,6 +22,9 @@ jobs: if: matrix.config.version == 'latest' run: | choco install cuda -y + - name: bigobj_env + run: | + echo "::set-env name=CXXFLAGS::/bigobj" - name: configure run: | $env:ChocolateyInstall = Convert-Path "$((Get-Command choco).Path)\..\.." @@ -51,8 +54,7 @@ jobs: run: | echo "::set-env name=origin_path::$env:PATH" echo "::add-path::$pwd\build\windows_shared_library" - - name: debug_env - if: matrix.config.build_type == 'Debug' + - name: bigobj_env run: | echo "::set-env name=CXXFLAGS::/bigobj" - name: configure diff --git a/common/base/math.hpp.inc b/common/base/math.hpp.inc index 3ba49b585c3..576ab32859c 100644 --- a/common/base/math.hpp.inc +++ b/common/base/math.hpp.inc @@ -54,10 +54,14 @@ struct is_complex_impl> : public std::integral_constant {}; +template +struct is_complex_or_scalar_impl> : std::is_scalar {}; + + template struct truncate_type_impl> { using type = thrust::complex::type>; }; -} // namespace detail \ No newline at end of file +} // namespace detail diff --git a/common/components/absolute_array.hpp.inc b/common/components/absolute_array.hpp.inc new file mode 100644 index 00000000000..ee6079fb8fc --- /dev/null +++ b/common/components/absolute_array.hpp.inc @@ -0,0 +1,61 @@ +/************************************************************* +Copyright (c) 2017-2020, 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. +*************************************************************/ + +namespace kernel { + + +template +__global__ + __launch_bounds__(default_block_size) void inplace_absolute_array_kernel( + const size_type n, ValueType *__restrict__ array) +{ + const auto tidx = thread::get_thread_id_flat(); + if (tidx < n) { + array[tidx] = abs(array[tidx]); + } +} + + +template +__global__ + __launch_bounds__(default_block_size) void outplace_absolute_array_kernel( + const size_type n, const ValueType *__restrict__ in, + remove_complex *__restrict__ out) +{ + const auto tidx = thread::get_thread_id_flat(); + if (tidx < n) { + out[tidx] = abs(in[tidx]); + } +} + + +} // namespace kernel diff --git a/common/matrix/dense_kernels.hpp.inc b/common/matrix/dense_kernels.hpp.inc index 9d0285e71f7..c26203115fe 100644 --- a/common/matrix/dense_kernels.hpp.inc +++ b/common/matrix/dense_kernels.hpp.inc @@ -510,4 +510,33 @@ __global__ __launch_bounds__(default_block_size) void extract_diagonal( } +template +__global__ __launch_bounds__(default_block_size) void inplace_absolute_dense( + size_type num_rows, size_type num_cols, ValueType *__restrict__ data, + size_type stride) +{ + const auto tidx = thread::get_thread_id_flat(); + auto row = tidx / num_cols; + auto col = tidx % num_cols; + if (row < num_rows) { + data[row * stride + col] = abs(data[row * stride + col]); + } +} + + +template +__global__ __launch_bounds__(default_block_size) void outplace_absolute_dense( + size_type num_rows, size_type num_cols, const ValueType *__restrict__ in, + size_type stride_in, remove_complex *__restrict__ out, + size_type stride_out) +{ + const auto tidx = thread::get_thread_id_flat(); + auto row = tidx / num_cols; + auto col = tidx % num_cols; + if (row < num_rows) { + out[row * stride_out + col] = abs(in[row * stride_in + col]); + } +} + + } // namespace kernel diff --git a/core/base/extended_float.hpp b/core/base/extended_float.hpp index 8ebf836641b..c83e7b81f20 100644 --- a/core/base/extended_float.hpp +++ b/core/base/extended_float.hpp @@ -70,18 +70,17 @@ template struct uint_of_impl {}; template -struct uint_of_impl>> { +struct uint_of_impl> { using type = uint16; }; template -struct uint_of_impl>> { +struct uint_of_impl> { using type = uint32; }; template -struct uint_of_impl>> { +struct uint_of_impl> { using type = uint64; }; diff --git a/core/components/absolute_array.hpp b/core/components/absolute_array.hpp new file mode 100644 index 00000000000..d48d04335fa --- /dev/null +++ b/core/components/absolute_array.hpp @@ -0,0 +1,109 @@ +/************************************************************* +Copyright (c) 2017-2020, 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_COMPONENTS_ABSOLUTE_ARRAY_HPP_ +#define GKO_CORE_COMPONENTS_ABSOLUTE_ARRAY_HPP_ + + +#include + + +#include +#include +#include + + +namespace gko { +namespace kernels { + + +#define GKO_DECLARE_INPLACE_ABSOLUTE_ARRAY_KERNEL(ValueType) \ + void inplace_absolute_array(std::shared_ptr exec, \ + ValueType *data, size_type num_entries) + +#define GKO_DECLARE_OUTPLACE_ABSOLUTE_ARRAY_KERNEL(ValueType) \ + void outplace_absolute_array(std::shared_ptr exec, \ + const ValueType *in, size_type num_entries, \ + remove_complex *out) + + +#define GKO_DECLARE_ALL_AS_TEMPLATES \ + template \ + GKO_DECLARE_INPLACE_ABSOLUTE_ARRAY_KERNEL(ValueType); \ + template \ + GKO_DECLARE_OUTPLACE_ABSOLUTE_ARRAY_KERNEL(ValueType) + + +namespace omp { +namespace components { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace components +} // namespace omp + + +namespace cuda { +namespace components { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace components +} // namespace cuda + + +namespace reference { +namespace components { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace components +} // namespace reference + + +namespace hip { +namespace components { + +GKO_DECLARE_ALL_AS_TEMPLATES; + +} // namespace components +} // namespace hip + + +#undef GKO_DECLARE_ALL_AS_TEMPLATES + + +} // namespace kernels +} // namespace gko + + +#endif // GKO_CORE_COMPONENTS_ABSOLUTE_ARRAY_HPP_ diff --git a/core/device_hooks/common_kernels.inc.cpp b/core/device_hooks/common_kernels.inc.cpp index 6adb0df0445..2f52d52ce8f 100644 --- a/core/device_hooks/common_kernels.inc.cpp +++ b/core/device_hooks/common_kernels.inc.cpp @@ -33,6 +33,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include "core/components/absolute_array.hpp" #include "core/components/fill_array.hpp" #include "core/components/precision_conversion.hpp" #include "core/components/prefix_sum.hpp" @@ -94,6 +95,16 @@ GKO_INSTANTIATE_FOR_EACH_INDEX_TYPE(GKO_DECLARE_FILL_ARRAY_KERNEL); GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_FILL_ARRAY_KERNEL); template GKO_DECLARE_FILL_ARRAY_KERNEL(size_type); +template +GKO_DECLARE_INPLACE_ABSOLUTE_ARRAY_KERNEL(ValueType) +GKO_NOT_COMPILED(GKO_HOOK_MODULE); +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_INPLACE_ABSOLUTE_ARRAY_KERNEL); + +template +GKO_DECLARE_OUTPLACE_ABSOLUTE_ARRAY_KERNEL(ValueType) +GKO_NOT_COMPILED(GKO_HOOK_MODULE); +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_ARRAY_KERNEL); + } // namespace components @@ -233,6 +244,16 @@ GKO_DECLARE_EXTRACT_DIAGONAL_KERNEL(ValueType) GKO_NOT_COMPILED(GKO_HOOK_MODULE); GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_EXTRACT_DIAGONAL_KERNEL); +template +GKO_DECLARE_INPLACE_ABSOLUTE_DENSE_KERNEL(ValueType) +GKO_NOT_COMPILED(GKO_HOOK_MODULE); +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_INPLACE_ABSOLUTE_DENSE_KERNEL); + +template +GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL(ValueType) +GKO_NOT_COMPILED(GKO_HOOK_MODULE); +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL); + } // namespace dense diff --git a/core/matrix/coo.cpp b/core/matrix/coo.cpp index dd5a4dda553..6ab488d9988 100644 --- a/core/matrix/coo.cpp +++ b/core/matrix/coo.cpp @@ -45,6 +45,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include "core/components/absolute_array.hpp" #include "core/components/fill_array.hpp" #include "core/matrix/coo_kernels.hpp" @@ -64,6 +65,10 @@ GKO_REGISTER_OPERATION(convert_to_csr, coo::convert_to_csr); GKO_REGISTER_OPERATION(convert_to_dense, coo::convert_to_dense); GKO_REGISTER_OPERATION(extract_diagonal, coo::extract_diagonal); GKO_REGISTER_OPERATION(fill_array, components::fill_array); +GKO_REGISTER_OPERATION(inplace_absolute_array, + components::inplace_absolute_array); +GKO_REGISTER_OPERATION(outplace_absolute_array, + components::outplace_absolute_array); } // namespace coo @@ -233,6 +238,35 @@ Coo::extract_diagonal() const } +template +void Coo::compute_absolute_inplace() +{ + auto exec = this->get_executor(); + + exec->run(coo::make_inplace_absolute_array( + this->get_values(), this->get_num_stored_elements())); +} + + +template +std::unique_ptr::absolute_type> +Coo::compute_absolute() const +{ + auto exec = this->get_executor(); + + auto abs_coo = absolute_type::create(exec, this->get_size(), + this->get_num_stored_elements()); + + abs_coo->col_idxs_ = col_idxs_; + abs_coo->row_idxs_ = row_idxs_; + exec->run(coo::make_outplace_absolute_array(this->get_const_values(), + this->get_num_stored_elements(), + abs_coo->get_values())); + + return abs_coo; +} + + #define GKO_DECLARE_COO_MATRIX(ValueType, IndexType) \ class Coo GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_COO_MATRIX); diff --git a/core/matrix/csr.cpp b/core/matrix/csr.cpp index 1863eea2959..32431850132 100644 --- a/core/matrix/csr.cpp +++ b/core/matrix/csr.cpp @@ -46,6 +46,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include "core/components/absolute_array.hpp" #include "core/components/fill_array.hpp" #include "core/matrix/csr_kernels.hpp" @@ -81,6 +82,10 @@ GKO_REGISTER_OPERATION(is_sorted_by_column_index, csr::is_sorted_by_column_index); GKO_REGISTER_OPERATION(extract_diagonal, csr::extract_diagonal); GKO_REGISTER_OPERATION(fill_array, components::fill_array); +GKO_REGISTER_OPERATION(inplace_absolute_array, + components::inplace_absolute_array); +GKO_REGISTER_OPERATION(outplace_absolute_array, + components::outplace_absolute_array); } // namespace csr @@ -482,6 +487,36 @@ Csr::extract_diagonal() const } +template +void Csr::compute_absolute_inplace() +{ + auto exec = this->get_executor(); + + exec->run(csr::make_inplace_absolute_array( + this->get_values(), this->get_num_stored_elements())); +} + + +template +std::unique_ptr::absolute_type> +Csr::compute_absolute() const +{ + auto exec = this->get_executor(); + + auto abs_csr = absolute_type::create(exec, this->get_size(), + this->get_num_stored_elements()); + + abs_csr->col_idxs_ = col_idxs_; + abs_csr->row_ptrs_ = row_ptrs_; + exec->run(csr::make_outplace_absolute_array(this->get_const_values(), + this->get_num_stored_elements(), + abs_csr->get_values())); + + convert_strategy_helper(abs_csr.get()); + return abs_csr; +} + + #define GKO_DECLARE_CSR_MATRIX(ValueType, IndexType) \ class Csr GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_CSR_MATRIX); diff --git a/core/matrix/dense.cpp b/core/matrix/dense.cpp index aab1d2e598a..aaaaf2d01a5 100644 --- a/core/matrix/dense.cpp +++ b/core/matrix/dense.cpp @@ -85,6 +85,8 @@ GKO_REGISTER_OPERATION(convert_to_hybrid, dense::convert_to_hybrid); GKO_REGISTER_OPERATION(convert_to_sellp, dense::convert_to_sellp); GKO_REGISTER_OPERATION(convert_to_sparsity_csr, dense::convert_to_sparsity_csr); GKO_REGISTER_OPERATION(extract_diagonal, dense::extract_diagonal); +GKO_REGISTER_OPERATION(inplace_absolute_dense, dense::inplace_absolute_dense); +GKO_REGISTER_OPERATION(outplace_absolute_dense, dense::outplace_absolute_dense); } // namespace dense @@ -752,6 +754,30 @@ std::unique_ptr> Dense::extract_diagonal() const } +template +void Dense::compute_absolute_inplace() +{ + auto exec = this->get_executor(); + + exec->run(dense::make_inplace_absolute_dense(this)); +} + + +template +std::unique_ptr::absolute_type> +Dense::compute_absolute() const +{ + auto exec = this->get_executor(); + + // do not inherit the stride + auto abs_dense = absolute_type::create(exec, this->get_size()); + + exec->run(dense::make_outplace_absolute_dense(this, abs_dense.get())); + + return abs_dense; +} + + #define GKO_DECLARE_DENSE_MATRIX(_type) class Dense<_type> GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DENSE_MATRIX); diff --git a/core/matrix/dense_kernels.hpp b/core/matrix/dense_kernels.hpp index 8edeec6878e..30ad6ca4729 100644 --- a/core/matrix/dense_kernels.hpp +++ b/core/matrix/dense_kernels.hpp @@ -172,6 +172,17 @@ namespace kernels { const matrix::Dense<_vtype> *orig, \ matrix::Diagonal<_vtype> *diag) +#define GKO_DECLARE_INPLACE_ABSOLUTE_DENSE_KERNEL(_vtype) \ + void inplace_absolute_dense(std::shared_ptr exec, \ + matrix::Dense<_vtype> *source) + +#define GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL(_vtype) \ + void outplace_absolute_dense( \ + std::shared_ptr exec, \ + const matrix::Dense<_vtype> *source, \ + matrix::Dense> *result) + + #define GKO_DECLARE_ALL_AS_TEMPLATES \ template \ GKO_DECLARE_DENSE_SIMPLE_APPLY_KERNEL(ValueType); \ @@ -220,7 +231,11 @@ namespace kernels { template \ GKO_DECLARE_INVERSE_COLUMN_PERMUTE_KERNEL(ValueType, IndexType); \ template \ - GKO_DECLARE_EXTRACT_DIAGONAL_KERNEL(ValueType) + GKO_DECLARE_EXTRACT_DIAGONAL_KERNEL(ValueType); \ + template \ + GKO_DECLARE_INPLACE_ABSOLUTE_DENSE_KERNEL(ValueType); \ + template \ + GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL(ValueType) namespace omp { diff --git a/core/matrix/diagonal.cpp b/core/matrix/diagonal.cpp index 04130cd4604..8df8af6a801 100644 --- a/core/matrix/diagonal.cpp +++ b/core/matrix/diagonal.cpp @@ -38,6 +38,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include "core/components/absolute_array.hpp" #include "core/matrix/diagonal_kernels.hpp" @@ -52,6 +53,10 @@ GKO_REGISTER_OPERATION(apply_to_csr, diagonal::apply_to_csr); GKO_REGISTER_OPERATION(right_apply_to_csr, diagonal::right_apply_to_csr); GKO_REGISTER_OPERATION(convert_to_csr, diagonal::convert_to_csr); GKO_REGISTER_OPERATION(conj_transpose, diagonal::conj_transpose); +GKO_REGISTER_OPERATION(inplace_absolute_array, + components::inplace_absolute_array); +GKO_REGISTER_OPERATION(outplace_absolute_array, + components::outplace_absolute_array); } // namespace diagonal @@ -261,6 +266,32 @@ void Diagonal::write(mat_data32 &data) const } +template +void Diagonal::compute_absolute_inplace() +{ + auto exec = this->get_executor(); + + exec->run(diagonal::make_inplace_absolute_array(this->get_values(), + this->get_size()[0])); +} + + +template +std::unique_ptr::absolute_type> +Diagonal::compute_absolute() const +{ + auto exec = this->get_executor(); + + auto abs_diagonal = absolute_type::create(exec, this->get_size()[0]); + + exec->run(diagonal::make_outplace_absolute_array( + this->get_const_values(), this->get_size()[0], + abs_diagonal->get_values())); + + return abs_diagonal; +} + + #define GKO_DECLARE_DIAGONAL_MATRIX(value_type) class Diagonal GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DIAGONAL_MATRIX); diff --git a/core/matrix/ell.cpp b/core/matrix/ell.cpp index 5f7591421e1..c802fe917e5 100644 --- a/core/matrix/ell.cpp +++ b/core/matrix/ell.cpp @@ -44,6 +44,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include "core/components/absolute_array.hpp" #include "core/components/fill_array.hpp" #include "core/matrix/ell_kernels.hpp" @@ -62,6 +63,10 @@ GKO_REGISTER_OPERATION(calculate_nonzeros_per_row, ell::calculate_nonzeros_per_row); GKO_REGISTER_OPERATION(extract_diagonal, ell::extract_diagonal); GKO_REGISTER_OPERATION(fill_array, components::fill_array); +GKO_REGISTER_OPERATION(inplace_absolute_array, + components::inplace_absolute_array); +GKO_REGISTER_OPERATION(outplace_absolute_array, + components::outplace_absolute_array); } // namespace ell @@ -251,6 +256,35 @@ Ell::extract_diagonal() const } +template +void Ell::compute_absolute_inplace() +{ + auto exec = this->get_executor(); + + exec->run(ell::make_inplace_absolute_array( + this->get_values(), this->get_num_stored_elements())); +} + + +template +std::unique_ptr::absolute_type> +Ell::compute_absolute() const +{ + auto exec = this->get_executor(); + + auto abs_ell = absolute_type::create( + exec, this->get_size(), this->get_num_stored_elements_per_row(), + this->get_stride()); + + abs_ell->col_idxs_ = col_idxs_; + exec->run(ell::make_outplace_absolute_array(this->get_const_values(), + this->get_num_stored_elements(), + abs_ell->get_values())); + + return abs_ell; +} + + #define GKO_DECLARE_ELL_MATRIX(ValueType, IndexType) \ class Ell GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_ELL_MATRIX); diff --git a/core/matrix/hybrid.cpp b/core/matrix/hybrid.cpp index d5f012b79c7..92768a09f89 100644 --- a/core/matrix/hybrid.cpp +++ b/core/matrix/hybrid.cpp @@ -43,6 +43,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include "core/components/absolute_array.hpp" #include "core/components/fill_array.hpp" #include "core/matrix/coo_kernels.hpp" #include "core/matrix/ell_kernels.hpp" @@ -60,6 +61,10 @@ GKO_REGISTER_OPERATION(count_nonzeros, hybrid::count_nonzeros); GKO_REGISTER_OPERATION(extract_coo_diagonal, coo::extract_diagonal); GKO_REGISTER_OPERATION(extract_ell_diagonal, ell::extract_diagonal); GKO_REGISTER_OPERATION(fill_array, components::fill_array); +GKO_REGISTER_OPERATION(inplace_absolute_array, + components::inplace_absolute_array); +GKO_REGISTER_OPERATION(outplace_absolute_array, + components::outplace_absolute_array); } // namespace hybrid @@ -285,6 +290,34 @@ Hybrid::extract_diagonal() const } +template +void Hybrid::compute_absolute_inplace() +{ + auto exec = this->get_executor(); + + exec->run(hybrid::make_inplace_absolute_array( + this->get_ell_values(), this->get_ell_num_stored_elements())); + exec->run(hybrid::make_inplace_absolute_array( + this->get_coo_values(), this->get_coo_num_stored_elements())); +} + + +template +std::unique_ptr::absolute_type> +Hybrid::compute_absolute() const +{ + auto exec = this->get_executor(); + + auto abs_hybrid = absolute_type::create( + exec, this->get_size(), this->get_strategy()); + + abs_hybrid->ell_->copy_from(ell_->compute_absolute()); + abs_hybrid->coo_->copy_from(coo_->compute_absolute()); + + return abs_hybrid; +} + + #define GKO_DECLARE_HYBRID_MATRIX(ValueType, IndexType) \ class Hybrid GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_HYBRID_MATRIX); diff --git a/core/matrix/sellp.cpp b/core/matrix/sellp.cpp index 7164af154f4..94c30544832 100644 --- a/core/matrix/sellp.cpp +++ b/core/matrix/sellp.cpp @@ -42,6 +42,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include "core/base/allocator.hpp" +#include "core/components/absolute_array.hpp" #include "core/components/fill_array.hpp" #include "core/matrix/sellp_kernels.hpp" @@ -58,6 +59,10 @@ GKO_REGISTER_OPERATION(convert_to_csr, sellp::convert_to_csr); GKO_REGISTER_OPERATION(count_nonzeros, sellp::count_nonzeros); GKO_REGISTER_OPERATION(extract_diagonal, sellp::extract_diagonal); GKO_REGISTER_OPERATION(fill_array, components::fill_array); +GKO_REGISTER_OPERATION(inplace_absolute_array, + components::inplace_absolute_array); +GKO_REGISTER_OPERATION(outplace_absolute_array, + components::outplace_absolute_array); } // namespace sellp @@ -304,6 +309,37 @@ Sellp::extract_diagonal() const } +template +void Sellp::compute_absolute_inplace() +{ + auto exec = this->get_executor(); + + exec->run(sellp::make_inplace_absolute_array( + this->get_values(), this->get_num_stored_elements())); +} + + +template +std::unique_ptr::absolute_type> +Sellp::compute_absolute() const +{ + auto exec = this->get_executor(); + + auto abs_sellp = absolute_type::create( + exec, this->get_size(), this->get_slice_size(), + this->get_stride_factor(), this->get_total_cols()); + + abs_sellp->col_idxs_ = col_idxs_; + abs_sellp->slice_lengths_ = slice_lengths_; + abs_sellp->slice_sets_ = slice_sets_; + exec->run(sellp::make_outplace_absolute_array( + this->get_const_values(), this->get_num_stored_elements(), + abs_sellp->get_values())); + + return abs_sellp; +} + + #define GKO_DECLARE_SELLP_MATRIX(ValueType, IndexType) \ class Sellp GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_SELLP_MATRIX); diff --git a/core/test/base/lin_op.cpp b/core/test/base/lin_op.cpp index 622158d9b42..303e8bd6e89 100644 --- a/core/test/base/lin_op.cpp +++ b/core/test/base/lin_op.cpp @@ -33,9 +33,17 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include +#include +#include +#include + + #include +#include + + namespace { @@ -314,4 +322,108 @@ TEST_F(EnableLinOpFactory, PassesParametersToLinOp) } +template +class DummyLinOpWithType + : public gko::EnableLinOp>, + public gko::EnableCreateMethod>, + public gko::EnableAbsoluteComputation< + gko::remove_complex>> { +public: + using absolute_type = gko::remove_complex; + DummyLinOpWithType(std::shared_ptr exec) + : gko::EnableLinOp(exec) + {} + + DummyLinOpWithType(std::shared_ptr exec, + gko::dim<2> size, Type value) + : gko::EnableLinOp(exec, size), value_(value) + {} + + void compute_absolute_inplace() override { value_ = gko::abs(value_); } + + std::unique_ptr compute_absolute() const override + { + return std::make_unique( + this->get_executor(), this->get_size(), gko::abs(value_)); + } + + Type get_value() const { return value_; } + +protected: + void apply_impl(const gko::LinOp *b, gko::LinOp *x) const override {} + + void apply_impl(const gko::LinOp *alpha, const gko::LinOp *b, + const gko::LinOp *beta, gko::LinOp *x) const override + {} + +private: + Type value_; +}; + + +class EnableAbsoluteComputation : public ::testing::Test { +protected: + using dummy_type = DummyLinOpWithType>; + EnableAbsoluteComputation() + : ref{gko::ReferenceExecutor::create()}, + op{dummy_type::create(ref, gko::dim<2>{1, 1}, + std::complex{-3.0, 4.0})} + {} + + std::shared_ptr ref; + std::shared_ptr op; +}; + + +TEST_F(EnableAbsoluteComputation, InplaceAbsoluteOnConcreteType) +{ + op->compute_absolute_inplace(); + + ASSERT_EQ(op->get_value(), std::complex{5.0}); +} + + +TEST_F(EnableAbsoluteComputation, OutplaceAbsoluteOnConcreteType) +{ + auto abs_op = op->compute_absolute(); + + static_assert( + std::is_same>>::value, + "Types must match."); + ASSERT_EQ(abs_op->get_value(), 5.0); +} + + +TEST_F(EnableAbsoluteComputation, InplaceAbsoluteOnAbsoluteComputable) +{ + auto linop = gko::as(op); + + gko::as(linop)->compute_absolute_inplace(); + + ASSERT_EQ(gko::as(linop)->get_value(), + std::complex{5.0}); +} + + +TEST_F(EnableAbsoluteComputation, OutplaceAbsoluteOnAbsoluteComputable) +{ + auto abs_op = op->compute_absolute(); + + static_assert( + std::is_same>>::value, + "Types must match."); + ASSERT_EQ(abs_op->get_value(), 5.0); +} + + +TEST_F(EnableAbsoluteComputation, ThrowWithoutAbsoluteComputableInterface) +{ + std::shared_ptr linop = DummyLinOp::create(ref); + + ASSERT_THROW(gko::as(linop), gko::NotSupported); +} + + } // namespace diff --git a/core/test/base/math.cpp b/core/test/base/math.cpp index c63cd4ae8e9..288bea41576 100644 --- a/core/test/base/math.cpp +++ b/core/test/base/math.cpp @@ -53,6 +53,10 @@ static_assert( "imag must return a real type"); +template +class DummyClass {}; + + template void test_real_is_finite() { @@ -135,4 +139,48 @@ TEST(Conjugate, DoubleComplex) } +TEST(RemoveComplexClass, Float) +{ + using origin = DummyClass; + using expect = DummyClass; + + bool check = std::is_same>::value; + + ASSERT_TRUE(check); +} + + +TEST(RemoveComplexClass, Double) +{ + using origin = DummyClass; + using expect = DummyClass; + + bool check = std::is_same>::value; + + ASSERT_TRUE(check); +} + + +TEST(RemoveComplexClass, FloatComplex) +{ + using origin = DummyClass, int>; + using expect = DummyClass; + + bool check = std::is_same>::value; + + ASSERT_TRUE(check); +} + + +TEST(RemoveComplexClass, DoubleComplex) +{ + using origin = DummyClass, int>; + using expect = DummyClass; + + bool check = std::is_same>::value; + + ASSERT_TRUE(check); +} + + } // namespace diff --git a/core/test/matrix/hybrid.cpp b/core/test/matrix/hybrid.cpp index dac9da86167..9a5ad45f0c7 100644 --- a/core/test/matrix/hybrid.cpp +++ b/core/test/matrix/hybrid.cpp @@ -42,6 +42,21 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. namespace { +template +struct change_index_s { + using type = gko::int32; +}; + +template <> +struct change_index_s { + using type = gko::int64; +}; + + +template +using change_index = typename change_index_s::type; + + template class Hybrid : public ::testing::Test { protected: @@ -297,4 +312,102 @@ TYPED_TEST(Hybrid, GeneratesCorrectMatrixData) } +TYPED_TEST(Hybrid, GetCorrectColumnLimit) +{ + using Mtx = typename TestFixture::Mtx; + using Mtx2 = gko::remove_complex; + using strategy = typename Mtx::column_limit; + using strategy2 = typename Mtx2::column_limit; + + auto mtx = Mtx::create(this->exec, std::make_shared(2)); + auto mtx_stra = gko::as(mtx->get_strategy()); + auto mtx2_stra = gko::as(mtx->template get_strategy()); + + EXPECT_EQ(mtx_stra->get_num_columns(), 2); + EXPECT_EQ(mtx2_stra->get_num_columns(), 2); +} + + +TYPED_TEST(Hybrid, GetCorrectImbalanceLimit) +{ + using Mtx = typename TestFixture::Mtx; + using Mtx2 = gko::remove_complex; + using strategy = typename Mtx::imbalance_limit; + using strategy2 = typename Mtx2::imbalance_limit; + + auto mtx = Mtx::create(this->exec, std::make_shared(0.4)); + auto mtx_stra = gko::as(mtx->get_strategy()); + auto mtx2_stra = gko::as(mtx->template get_strategy()); + + EXPECT_EQ(mtx_stra->get_percentage(), 0.4f); + EXPECT_EQ(mtx2_stra->get_percentage(), 0.4f); +} + + +TYPED_TEST(Hybrid, GetCorrectImbalanceBoundedLimit) +{ + using Mtx = typename TestFixture::Mtx; + using Mtx2 = gko::remove_complex; + using strategy = typename Mtx::imbalance_bounded_limit; + using strategy2 = typename Mtx2::imbalance_bounded_limit; + + auto mtx = Mtx::create(this->exec, std::make_shared(0.4, 0.1)); + auto mtx_stra = gko::as(mtx->get_strategy()); + auto mtx2_stra = gko::as(mtx->template get_strategy()); + + EXPECT_EQ(mtx_stra->get_percentage(), 0.4f); + EXPECT_EQ(mtx_stra->get_ratio(), 0.1f); + EXPECT_EQ(mtx2_stra->get_percentage(), 0.4f); + EXPECT_EQ(mtx2_stra->get_ratio(), 0.1f); +} + + +TYPED_TEST(Hybrid, GetCorrectMinimalStorageLimitWithDifferentHybType) +{ + using Mtx = typename TestFixture::Mtx; + using value_type = typename TestFixture::value_type; + using index_type = typename TestFixture::index_type; + using Mtx2 = gko::matrix::Hybrid>; + using strategy = typename Mtx::minimal_storage_limit; + using strategy2 = typename Mtx2::imbalance_limit; + + auto mtx = Mtx::create(this->exec, std::make_shared()); + auto mtx_stra = gko::as(mtx->get_strategy()); + auto mtx2_stra = gko::as(mtx->template get_strategy()); + + EXPECT_EQ(mtx2_stra->get_percentage(), mtx_stra->get_percentage()); +} + + +TYPED_TEST(Hybrid, GetCorrectMinimalStorageLimitWithSameHybType) +{ + using Mtx = typename TestFixture::Mtx; + using Mtx2 = Mtx; + using strategy = typename Mtx::minimal_storage_limit; + using strategy2 = typename Mtx2::minimal_storage_limit; + + auto mtx = Mtx::create(this->exec, std::make_shared()); + auto mtx_stra = gko::as(mtx->get_strategy()); + auto mtx2_stra = gko::as(mtx->template get_strategy()); + + EXPECT_EQ(mtx2_stra->get_percentage(), mtx_stra->get_percentage()); +} + + +TYPED_TEST(Hybrid, GetCorrectAutomatic) +{ + using Mtx = typename TestFixture::Mtx; + using value_type = typename TestFixture::value_type; + using index_type = typename TestFixture::index_type; + using Mtx2 = Mtx; + using strategy = typename Mtx::automatic; + using strategy2 = typename Mtx2::automatic; + + auto mtx = Mtx::create(this->exec, std::make_shared()); + auto mtx_stra = gko::as(mtx->get_strategy()); + + ASSERT_NO_THROW(gko::as(mtx->template get_strategy())); +} + + } // namespace diff --git a/cuda/CMakeLists.txt b/cuda/CMakeLists.txt index 52560f7be07..d3a73ab536c 100644 --- a/cuda/CMakeLists.txt +++ b/cuda/CMakeLists.txt @@ -70,6 +70,7 @@ target_sources(ginkgo_cuda base/exception.cpp base/executor.cpp base/version.cpp + components/absolute_array.cu components/fill_array.cu components/precision_conversion.cu components/prefix_sum.cu diff --git a/cuda/components/absolute_array.cu b/cuda/components/absolute_array.cu new file mode 100644 index 00000000000..8ef4901bf2c --- /dev/null +++ b/cuda/components/absolute_array.cu @@ -0,0 +1,82 @@ +/************************************************************* +Copyright (c) 2017-2020, 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/components/absolute_array.hpp" + + +#include "cuda/base/types.hpp" +#include "cuda/components/thread_ids.cuh" + + +namespace gko { +namespace kernels { +namespace cuda { +namespace components { + + +constexpr int default_block_size = 512; + + +#include "common/components/absolute_array.hpp.inc" + + +template +void inplace_absolute_array(std::shared_ptr exec, + ValueType *data, size_type n) +{ + const dim3 block_size(default_block_size, 1, 1); + const dim3 grid_size(ceildiv(n, block_size.x), 1, 1); + kernel::inplace_absolute_array_kernel<<>>( + n, as_cuda_type(data)); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_INPLACE_ABSOLUTE_ARRAY_KERNEL); + + +template +void outplace_absolute_array(std::shared_ptr exec, + const ValueType *in, size_type n, + remove_complex *out) +{ + const dim3 block_size(default_block_size, 1, 1); + const dim3 grid_size(ceildiv(n, block_size.x), 1, 1); + kernel::outplace_absolute_array_kernel<<>>( + n, as_cuda_type(in), as_cuda_type(out)); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_ARRAY_KERNEL); + + +} // namespace components +} // namespace cuda +} // namespace kernels +} // namespace gko diff --git a/cuda/matrix/dense_kernels.cu b/cuda/matrix/dense_kernels.cu index f0fe2cfae7e..b1c4ae747f8 100644 --- a/cuda/matrix/dense_kernels.cu +++ b/cuda/matrix/dense_kernels.cu @@ -701,6 +701,38 @@ void extract_diagonal(std::shared_ptr exec, GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_EXTRACT_DIAGONAL_KERNEL); +template +void inplace_absolute_dense(std::shared_ptr exec, + matrix::Dense *source) +{ + auto dim = source->get_size(); + const dim3 grid_dim = ceildiv(dim[0] * dim[1], default_block_size); + + kernel::inplace_absolute_dense<<>>( + dim[0], dim[1], as_cuda_type(source->get_values()), + source->get_stride()); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_INPLACE_ABSOLUTE_DENSE_KERNEL); + + +template +void outplace_absolute_dense(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Dense> *result) +{ + auto dim = source->get_size(); + const dim3 grid_dim = ceildiv(dim[0] * dim[1], default_block_size); + + kernel::outplace_absolute_dense<<>>( + dim[0], dim[1], as_cuda_type(source->get_const_values()), + source->get_stride(), as_cuda_type(result->get_values()), + result->get_stride()); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL); + + } // namespace dense } // namespace cuda } // namespace kernels diff --git a/cuda/test/components/CMakeLists.txt b/cuda/test/components/CMakeLists.txt index 154a39e963e..1d76a251948 100644 --- a/cuda/test/components/CMakeLists.txt +++ b/cuda/test/components/CMakeLists.txt @@ -2,6 +2,7 @@ ginkgo_create_cuda_test(cooperative_groups_kernels) ginkgo_create_cuda_test(merging_kernels) ginkgo_create_cuda_test(searching_kernels) ginkgo_create_cuda_test(sorting_kernels) +ginkgo_create_test(absolute_array) ginkgo_create_test(fill_array) ginkgo_create_test(precision_conversion) ginkgo_create_test(prefix_sum) diff --git a/cuda/test/components/absolute_array.cpp b/cuda/test/components/absolute_array.cpp new file mode 100644 index 00000000000..919c042cff0 --- /dev/null +++ b/cuda/test/components/absolute_array.cpp @@ -0,0 +1,132 @@ +/************************************************************* +Copyright (c) 2017-2020, 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/components/absolute_array.hpp" + + +#include +#include +#include + + +#include + + +#include + + +#include "core/test/utils/assertions.hpp" + + +namespace { + + +class AbsoluteArray : public ::testing::Test { +protected: + using value_type = double; + using complex_type = std::complex; + AbsoluteArray() + : ref(gko::ReferenceExecutor::create()), + exec(gko::CudaExecutor::create(0, ref)), + total_size(6344), + vals(ref, total_size), + dvals(exec, total_size), + complex_vals(ref, total_size), + dcomplex_vals(exec, total_size) + { + std::fill_n(vals.get_data(), total_size, -1234.0); + dvals = vals; + std::fill_n(complex_vals.get_data(), total_size, complex_type{3, 4}); + dcomplex_vals = complex_vals; + } + + std::shared_ptr ref; + std::shared_ptr exec; + gko::size_type total_size; + gko::Array vals; + gko::Array dvals; + gko::Array complex_vals; + gko::Array dcomplex_vals; +}; + + +TEST_F(AbsoluteArray, InplaceEqualsReference) +{ + gko::kernels::cuda::components::inplace_absolute_array( + exec, dvals.get_data(), total_size); + gko::kernels::reference::components::inplace_absolute_array( + ref, vals.get_data(), total_size); + + GKO_ASSERT_ARRAY_EQ(vals, dvals); +} + + +TEST_F(AbsoluteArray, InplaceComplexEqualsReference) +{ + gko::kernels::cuda::components::inplace_absolute_array( + exec, dcomplex_vals.get_data(), total_size); + gko::kernels::reference::components::inplace_absolute_array( + ref, complex_vals.get_data(), total_size); + + GKO_ASSERT_ARRAY_EQ(complex_vals, dcomplex_vals); +} + + +TEST_F(AbsoluteArray, OutplaceEqualsReference) +{ + gko::Array abs_vals(ref, total_size); + gko::Array dabs_vals(exec, total_size); + + gko::kernels::cuda::components::outplace_absolute_array( + exec, dvals.get_const_data(), total_size, dabs_vals.get_data()); + gko::kernels::reference::components::outplace_absolute_array( + ref, vals.get_const_data(), total_size, abs_vals.get_data()); + + GKO_ASSERT_ARRAY_EQ(abs_vals, dabs_vals); +} + + +TEST_F(AbsoluteArray, OutplaceComplexEqualsReference) +{ + gko::Array abs_vals(ref, total_size); + gko::Array dabs_vals(exec, total_size); + + gko::kernels::cuda::components::outplace_absolute_array( + exec, dcomplex_vals.get_const_data(), total_size, dabs_vals.get_data()); + gko::kernels::reference::components::outplace_absolute_array( + ref, complex_vals.get_const_data(), total_size, abs_vals.get_data()); + + GKO_ASSERT_ARRAY_EQ(abs_vals, dabs_vals); +} + + +} // namespace diff --git a/cuda/test/matrix/coo_kernels.cpp b/cuda/test/matrix/coo_kernels.cpp index 3522e918cc6..6815a9d3df1 100644 --- a/cuda/test/matrix/coo_kernels.cpp +++ b/cuda/test/matrix/coo_kernels.cpp @@ -271,4 +271,26 @@ TEST_F(Coo, ExtractDiagonalIsEquivalentToRef) } +TEST_F(Coo, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + mtx->compute_absolute_inplace(); + dmtx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(mtx, dmtx, 1e-14); +} + + +TEST_F(Coo, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + auto abs_mtx = mtx->compute_absolute(); + auto dabs_mtx = dmtx->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, dabs_mtx, 1e-14); +} + + } // namespace diff --git a/cuda/test/matrix/csr_kernels.cpp b/cuda/test/matrix/csr_kernels.cpp index 2f1bb544ac3..4205f5c5d97 100644 --- a/cuda/test/matrix/csr_kernels.cpp +++ b/cuda/test/matrix/csr_kernels.cpp @@ -60,8 +60,8 @@ namespace { class Csr : public ::testing::Test { protected: - using Mtx = gko::matrix::Csr<>; using Vec = gko::matrix::Dense<>; + using Mtx = gko::matrix::Csr<>; using ComplexVec = gko::matrix::Dense>; using ComplexMtx = gko::matrix::Csr>; @@ -740,4 +740,48 @@ TEST_F(Csr, ExtractDiagonalIsEquivalentToRef) } +TEST_F(Csr, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared(cuda)); + + mtx->compute_absolute_inplace(); + dmtx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(mtx, dmtx, 1e-14); +} + + +TEST_F(Csr, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared(cuda)); + + auto abs_mtx = mtx->compute_absolute(); + auto dabs_mtx = dmtx->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, dabs_mtx, 1e-14); +} + + +TEST_F(Csr, InplaceAbsoluteComplexMatrixIsEquivalentToRef) +{ + set_up_apply_complex_data(std::make_shared(cuda)); + + complex_mtx->compute_absolute_inplace(); + complex_dmtx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(complex_mtx, complex_dmtx, 1e-14); +} + + +TEST_F(Csr, OutplaceAbsoluteComplexMatrixIsEquivalentToRef) +{ + set_up_apply_complex_data(std::make_shared(cuda)); + + auto abs_mtx = complex_mtx->compute_absolute(); + auto dabs_mtx = complex_dmtx->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, dabs_mtx, 1e-14); +} + + } // namespace diff --git a/cuda/test/matrix/dense_kernels.cpp b/cuda/test/matrix/dense_kernels.cpp index c520dfcf2a3..e8732c13136 100644 --- a/cuda/test/matrix/dense_kernels.cpp +++ b/cuda/test/matrix/dense_kernels.cpp @@ -587,4 +587,26 @@ TEST_F(Dense, ExtractDiagonalIsEquivalentToRef) } +TEST_F(Dense, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + x->compute_absolute_inplace(); + dx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(x, dx, 1e-14); +} + + +TEST_F(Dense, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + auto abs_x = x->compute_absolute(); + auto dabs_x = dx->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_x, dabs_x, 1e-14); +} + + } // namespace diff --git a/cuda/test/matrix/diagonal_kernels.cpp b/cuda/test/matrix/diagonal_kernels.cpp index 8c8bbe207e2..695eee1a880 100644 --- a/cuda/test/matrix/diagonal_kernels.cpp +++ b/cuda/test/matrix/diagonal_kernels.cpp @@ -249,4 +249,26 @@ TEST_F(Diagonal, ConjTransposeIsEquivalentToRef) } +TEST_F(Diagonal, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + diag->compute_absolute_inplace(); + ddiag->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(diag, ddiag, 1e-14); +} + + +TEST_F(Diagonal, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + auto abs_diag = diag->compute_absolute(); + auto dabs_diag = ddiag->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_diag, dabs_diag, 1e-14); +} + + } // namespace diff --git a/cuda/test/matrix/ell_kernels.cpp b/cuda/test/matrix/ell_kernels.cpp index 960faaed20e..ce48a80e634 100644 --- a/cuda/test/matrix/ell_kernels.cpp +++ b/cuda/test/matrix/ell_kernels.cpp @@ -356,4 +356,26 @@ TEST_F(Ell, ExtractDiagonalIsEquivalentToRef) } +TEST_F(Ell, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + mtx->compute_absolute_inplace(); + dmtx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(mtx, dmtx, 1e-14); +} + + +TEST_F(Ell, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + auto abs_mtx = mtx->compute_absolute(); + auto dabs_mtx = dmtx->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, dabs_mtx, 1e-14); +} + + } // namespace diff --git a/cuda/test/matrix/hybrid_kernels.cpp b/cuda/test/matrix/hybrid_kernels.cpp index 5e7048632bc..366686f27a0 100644 --- a/cuda/test/matrix/hybrid_kernels.cpp +++ b/cuda/test/matrix/hybrid_kernels.cpp @@ -231,4 +231,33 @@ TEST_F(Hybrid, ExtractDiagonalIsEquivalentToRef) } +TEST_F(Hybrid, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + mtx->compute_absolute_inplace(); + dmtx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(mtx, dmtx, 1e-14); +} + + +TEST_F(Hybrid, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(1, std::make_shared(2)); + using AbsMtx = gko::remove_complex; + + auto abs_mtx = mtx->compute_absolute(); + auto dabs_mtx = dmtx->compute_absolute(); + auto abs_strategy = gko::as(abs_mtx->get_strategy()); + auto dabs_strategy = + gko::as(dabs_mtx->get_strategy()); + + GKO_ASSERT_MTX_NEAR(abs_mtx, dabs_mtx, 1e-14); + GKO_ASSERT_EQ(abs_strategy->get_num_columns(), + dabs_strategy->get_num_columns()); + GKO_ASSERT_EQ(abs_strategy->get_num_columns(), 2); +} + + } // namespace diff --git a/cuda/test/matrix/sellp_kernels.cpp b/cuda/test/matrix/sellp_kernels.cpp index a3cdeed95ad..7a00b63d0c1 100644 --- a/cuda/test/matrix/sellp_kernels.cpp +++ b/cuda/test/matrix/sellp_kernels.cpp @@ -350,4 +350,26 @@ TEST_F(Sellp, ExtractDiagonalWithSliceSizeAndStrideFactorIsEquivalentToRef) } +TEST_F(Sellp, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_matrix(32, 2); + + mtx->compute_absolute_inplace(); + dmtx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(mtx, dmtx, 1e-14); +} + + +TEST_F(Sellp, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_matrix(32, 2); + + auto abs_mtx = mtx->compute_absolute(); + auto dabs_mtx = dmtx->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, dabs_mtx, 1e-14); +} + + } // namespace diff --git a/hip/CMakeLists.txt b/hip/CMakeLists.txt index 57c6503ce36..20fcb0d79a0 100644 --- a/hip/CMakeLists.txt +++ b/hip/CMakeLists.txt @@ -143,6 +143,7 @@ set(GINKGO_HIP_SOURCES base/exception.hip.cpp base/executor.hip.cpp base/version.hip.cpp + components/absolute_array.hip.cpp components/fill_array.hip.cpp components/precision_conversion.hip.cpp components/prefix_sum.hip.cpp diff --git a/hip/components/absolute_array.hip.cpp b/hip/components/absolute_array.hip.cpp new file mode 100644 index 00000000000..3f7ee4c012c --- /dev/null +++ b/hip/components/absolute_array.hip.cpp @@ -0,0 +1,86 @@ +/************************************************************* +Copyright (c) 2017-2020, 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/components/absolute_array.hpp" + + +#include + + +#include "hip/base/types.hip.hpp" +#include "hip/components/thread_ids.hip.hpp" + + +namespace gko { +namespace kernels { +namespace hip { +namespace components { + + +constexpr int default_block_size = 512; + + +#include "common/components/absolute_array.hpp.inc" + + +template +void inplace_absolute_array(std::shared_ptr exec, + ValueType *data, size_type n) +{ + const dim3 block_size(default_block_size, 1, 1); + const dim3 grid_size(ceildiv(n, block_size.x), 1, 1); + hipLaunchKernelGGL(kernel::inplace_absolute_array_kernel, dim3(grid_size), + dim3(block_size), 0, 0, n, as_hip_type(data)); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_INPLACE_ABSOLUTE_ARRAY_KERNEL); + + +template +void outplace_absolute_array(std::shared_ptr exec, + const ValueType *in, size_type n, + remove_complex *out) +{ + const dim3 block_size(default_block_size, 1, 1); + const dim3 grid_size(ceildiv(n, block_size.x), 1, 1); + hipLaunchKernelGGL(kernel::outplace_absolute_array_kernel, dim3(grid_size), + dim3(block_size), 0, 0, n, as_hip_type(in), + as_hip_type(out)); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_ARRAY_KERNEL); + + +} // namespace components +} // namespace hip +} // namespace kernels +} // namespace gko diff --git a/hip/matrix/dense_kernels.hip.cpp b/hip/matrix/dense_kernels.hip.cpp index 824d2712b7f..2793d09832f 100644 --- a/hip/matrix/dense_kernels.hip.cpp +++ b/hip/matrix/dense_kernels.hip.cpp @@ -729,6 +729,39 @@ void extract_diagonal(std::shared_ptr exec, GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_EXTRACT_DIAGONAL_KERNEL); +template +void inplace_absolute_dense(std::shared_ptr exec, + matrix::Dense *source) +{ + auto dim = source->get_size(); + const dim3 grid_dim = ceildiv(dim[0] * dim[1], default_block_size); + + hipLaunchKernelGGL(kernel::inplace_absolute_dense, dim3(grid_dim), + dim3(default_block_size), 0, 0, dim[0], dim[1], + as_hip_type(source->get_values()), source->get_stride()); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_INPLACE_ABSOLUTE_DENSE_KERNEL); + + +template +void outplace_absolute_dense(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Dense> *result) +{ + auto dim = source->get_size(); + const dim3 grid_dim = ceildiv(dim[0] * dim[1], default_block_size); + + hipLaunchKernelGGL(kernel::outplace_absolute_dense, dim3(grid_dim), + dim3(default_block_size), 0, 0, dim[0], dim[1], + as_hip_type(source->get_const_values()), + source->get_stride(), as_hip_type(result->get_values()), + result->get_stride()); +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL); + + } // namespace dense } // namespace hip } // namespace kernels diff --git a/hip/test/components/CMakeLists.txt b/hip/test/components/CMakeLists.txt index b3bec2595f9..f0bae60fe60 100644 --- a/hip/test/components/CMakeLists.txt +++ b/hip/test/components/CMakeLists.txt @@ -1,3 +1,4 @@ +ginkgo_create_hip_test(absolute_array) ginkgo_create_hip_test(cooperative_groups_kernels) ginkgo_create_hip_test(fill_array) ginkgo_create_hip_test(merging_kernels) diff --git a/hip/test/components/absolute_array.hip.cpp b/hip/test/components/absolute_array.hip.cpp new file mode 100644 index 00000000000..f718805d5eb --- /dev/null +++ b/hip/test/components/absolute_array.hip.cpp @@ -0,0 +1,132 @@ +/************************************************************* +Copyright (c) 2017-2020, 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/components/absolute_array.hpp" + + +#include +#include +#include + + +#include + + +#include + + +#include "core/test/utils/assertions.hpp" + + +namespace { + + +class AbsoluteArray : public ::testing::Test { +protected: + using value_type = double; + using complex_type = std::complex; + AbsoluteArray() + : ref(gko::ReferenceExecutor::create()), + exec(gko::HipExecutor::create(0, ref)), + total_size(6344), + vals(ref, total_size), + dvals(exec, total_size), + complex_vals(ref, total_size), + dcomplex_vals(exec, total_size) + { + std::fill_n(vals.get_data(), total_size, -1234.0); + dvals = vals; + std::fill_n(complex_vals.get_data(), total_size, complex_type{3, 4}); + dcomplex_vals = complex_vals; + } + + std::shared_ptr ref; + std::shared_ptr exec; + gko::size_type total_size; + gko::Array vals; + gko::Array dvals; + gko::Array complex_vals; + gko::Array dcomplex_vals; +}; + + +TEST_F(AbsoluteArray, InplaceEqualsReference) +{ + gko::kernels::hip::components::inplace_absolute_array( + exec, dvals.get_data(), total_size); + gko::kernels::reference::components::inplace_absolute_array( + ref, vals.get_data(), total_size); + + GKO_ASSERT_ARRAY_EQ(vals, dvals); +} + + +TEST_F(AbsoluteArray, InplaceComplexEqualsReference) +{ + gko::kernels::hip::components::inplace_absolute_array( + exec, dcomplex_vals.get_data(), total_size); + gko::kernels::reference::components::inplace_absolute_array( + ref, complex_vals.get_data(), total_size); + + GKO_ASSERT_ARRAY_EQ(complex_vals, dcomplex_vals); +} + + +TEST_F(AbsoluteArray, OutplaceEqualsReference) +{ + gko::Array abs_vals(ref, total_size); + gko::Array dabs_vals(exec, total_size); + + gko::kernels::hip::components::outplace_absolute_array( + exec, dvals.get_const_data(), total_size, dabs_vals.get_data()); + gko::kernels::reference::components::outplace_absolute_array( + ref, vals.get_const_data(), total_size, abs_vals.get_data()); + + GKO_ASSERT_ARRAY_EQ(abs_vals, dabs_vals); +} + + +TEST_F(AbsoluteArray, OutplaceComplexEqualsReference) +{ + gko::Array abs_vals(ref, total_size); + gko::Array dabs_vals(exec, total_size); + + gko::kernels::hip::components::outplace_absolute_array( + exec, dcomplex_vals.get_const_data(), total_size, dabs_vals.get_data()); + gko::kernels::reference::components::outplace_absolute_array( + ref, complex_vals.get_const_data(), total_size, abs_vals.get_data()); + + GKO_ASSERT_ARRAY_EQ(abs_vals, dabs_vals); +} + + +} // namespace diff --git a/hip/test/matrix/coo_kernels.hip.cpp b/hip/test/matrix/coo_kernels.hip.cpp index 3e23bb6f4a4..6d0a0935cd6 100644 --- a/hip/test/matrix/coo_kernels.hip.cpp +++ b/hip/test/matrix/coo_kernels.hip.cpp @@ -271,4 +271,26 @@ TEST_F(Coo, ExtractDiagonalIsEquivalentToRef) } +TEST_F(Coo, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + mtx->compute_absolute_inplace(); + dmtx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(mtx, dmtx, 1e-14); +} + + +TEST_F(Coo, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + auto abs_mtx = mtx->compute_absolute(); + auto dabs_mtx = dmtx->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, dabs_mtx, 1e-14); +} + + } // namespace diff --git a/hip/test/matrix/csr_kernels.hip.cpp b/hip/test/matrix/csr_kernels.hip.cpp index 14ec43d6a4a..54b0f4cda65 100644 --- a/hip/test/matrix/csr_kernels.hip.cpp +++ b/hip/test/matrix/csr_kernels.hip.cpp @@ -60,8 +60,10 @@ namespace { class Csr : public ::testing::Test { protected: - using Mtx = gko::matrix::Csr<>; using Vec = gko::matrix::Dense<>; + using Mtx = gko::matrix::Csr<>; + using ComplexVec = gko::matrix::Dense>; + using ComplexMtx = gko::matrix::Csr>; Csr() : mtx_size(532, 231), rand_engine(42) {} @@ -114,6 +116,16 @@ class Csr : public ::testing::Test { dbeta->copy_from(beta.get()); } + void set_up_apply_complex_data( + std::shared_ptr strategy) + { + complex_mtx = ComplexMtx::create(ref, strategy); + complex_mtx->copy_from( + gen_mtx(mtx_size[0], mtx_size[1], 1)); + complex_dmtx = ComplexMtx::create(hip, strategy); + complex_dmtx->copy_from(complex_mtx.get()); + } + struct matrix_pair { std::unique_ptr ref; std::unique_ptr hip; @@ -153,6 +165,7 @@ class Csr : public ::testing::Test { std::ranlux48 rand_engine; std::unique_ptr mtx; + std::unique_ptr complex_mtx; std::unique_ptr square_mtx; std::unique_ptr expected; std::unique_ptr y; @@ -160,6 +173,7 @@ class Csr : public ::testing::Test { std::unique_ptr beta; std::unique_ptr dmtx; + std::unique_ptr complex_dmtx; std::unique_ptr square_dmtx; std::unique_ptr dresult; std::unique_ptr dy; @@ -711,4 +725,48 @@ TEST_F(Csr, ExtractDiagonalIsEquivalentToRef) } +TEST_F(Csr, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared(hip)); + + mtx->compute_absolute_inplace(); + dmtx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(mtx, dmtx, 1e-14); +} + + +TEST_F(Csr, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(std::make_shared(hip)); + + auto abs_mtx = mtx->compute_absolute(); + auto dabs_mtx = dmtx->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, dabs_mtx, 1e-14); +} + + +TEST_F(Csr, InplaceAbsoluteComplexMatrixIsEquivalentToRef) +{ + set_up_apply_complex_data(std::make_shared(hip)); + + complex_mtx->compute_absolute_inplace(); + complex_dmtx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(complex_mtx, complex_dmtx, 1e-14); +} + + +TEST_F(Csr, OutplaceAbsoluteComplexMatrixIsEquivalentToRef) +{ + set_up_apply_complex_data(std::make_shared(hip)); + + auto abs_mtx = complex_mtx->compute_absolute(); + auto dabs_mtx = complex_dmtx->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, dabs_mtx, 1e-14); +} + + } // namespace diff --git a/hip/test/matrix/dense_kernels.hip.cpp b/hip/test/matrix/dense_kernels.hip.cpp index 40de01b14e8..c14809ed79e 100644 --- a/hip/test/matrix/dense_kernels.hip.cpp +++ b/hip/test/matrix/dense_kernels.hip.cpp @@ -569,4 +569,26 @@ TEST_F(Dense, ExtractDiagonalIsEquivalentToRef) } +TEST_F(Dense, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + x->compute_absolute_inplace(); + dx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(x, dx, 1e-14); +} + + +TEST_F(Dense, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + auto abs_x = x->compute_absolute(); + auto dabs_x = dx->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_x, dabs_x, 1e-14); +} + + } // namespace diff --git a/hip/test/matrix/diagonal_kernels.hip.cpp b/hip/test/matrix/diagonal_kernels.hip.cpp index 6b87034006e..1bd1b3f0b8a 100644 --- a/hip/test/matrix/diagonal_kernels.hip.cpp +++ b/hip/test/matrix/diagonal_kernels.hip.cpp @@ -249,4 +249,26 @@ TEST_F(Diagonal, ConjTransposeIsEquivalentToRef) } +TEST_F(Diagonal, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + diag->compute_absolute_inplace(); + ddiag->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(diag, ddiag, 1e-14); +} + + +TEST_F(Diagonal, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + auto abs_diag = diag->compute_absolute(); + auto dabs_diag = ddiag->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_diag, dabs_diag, 1e-14); +} + + } // namespace diff --git a/hip/test/matrix/ell_kernels.hip.cpp b/hip/test/matrix/ell_kernels.hip.cpp index fa9a2df717a..74d22dbb738 100644 --- a/hip/test/matrix/ell_kernels.hip.cpp +++ b/hip/test/matrix/ell_kernels.hip.cpp @@ -356,4 +356,26 @@ TEST_F(Ell, ExtractDiagonalIsEquivalentToRef) } +TEST_F(Ell, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + mtx->compute_absolute_inplace(); + dmtx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(mtx, dmtx, 1e-14); +} + + +TEST_F(Ell, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + auto abs_mtx = mtx->compute_absolute(); + auto dabs_mtx = dmtx->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, dabs_mtx, 1e-14); +} + + } // namespace diff --git a/hip/test/matrix/hybrid_kernels.hip.cpp b/hip/test/matrix/hybrid_kernels.hip.cpp index a765cbc2717..f7c60b59b3f 100644 --- a/hip/test/matrix/hybrid_kernels.hip.cpp +++ b/hip/test/matrix/hybrid_kernels.hip.cpp @@ -231,4 +231,33 @@ TEST_F(Hybrid, ExtractDiagonalIsEquivalentToRef) } +TEST_F(Hybrid, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(); + + mtx->compute_absolute_inplace(); + dmtx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(mtx, dmtx, 1e-14); +} + + +TEST_F(Hybrid, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_data(1, std::make_shared(2)); + using AbsMtx = gko::remove_complex; + + auto abs_mtx = mtx->compute_absolute(); + auto dabs_mtx = dmtx->compute_absolute(); + auto abs_strategy = gko::as(abs_mtx->get_strategy()); + auto dabs_strategy = + gko::as(dabs_mtx->get_strategy()); + + GKO_ASSERT_MTX_NEAR(abs_mtx, dabs_mtx, 1e-14); + GKO_ASSERT_EQ(abs_strategy->get_num_columns(), + dabs_strategy->get_num_columns()); + GKO_ASSERT_EQ(abs_strategy->get_num_columns(), 2); +} + + } // namespace diff --git a/hip/test/matrix/sellp_kernels.hip.cpp b/hip/test/matrix/sellp_kernels.hip.cpp index 80c3cf7c113..6c99519b71d 100644 --- a/hip/test/matrix/sellp_kernels.hip.cpp +++ b/hip/test/matrix/sellp_kernels.hip.cpp @@ -349,4 +349,26 @@ TEST_F(Sellp, ExtractDiagonalWithSliceSizeAndStrideFactorIsEquivalentToRef) } +TEST_F(Sellp, InplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_matrix(32, 2); + + mtx->compute_absolute_inplace(); + dmtx->compute_absolute_inplace(); + + GKO_ASSERT_MTX_NEAR(mtx, dmtx, 1e-14); +} + + +TEST_F(Sellp, OutplaceAbsoluteMatrixIsEquivalentToRef) +{ + set_up_apply_matrix(32, 2); + + auto abs_mtx = mtx->compute_absolute(); + auto dabs_mtx = dmtx->compute_absolute(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, dabs_mtx, 1e-14); +} + + } // namespace diff --git a/include/ginkgo/core/base/lin_op.hpp b/include/ginkgo/core/base/lin_op.hpp index dcf5a1a88cf..6c873198b6e 100644 --- a/include/ginkgo/core/base/lin_op.hpp +++ b/include/ginkgo/core/base/lin_op.hpp @@ -42,6 +42,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. #include #include #include +#include #include #include #include @@ -622,6 +623,60 @@ class DiagonalExtractable { const = 0; }; + +/** + * The AbsoluteComputable is an interface that allows to get the component wise + * absolute of a LinOp. Use EnableAbsoluteComputation to + * implement this interface. + */ +class AbsoluteComputable { +public: + /** + * Gets the absolute LinOp + * + * @return a pointer to the new absolute LinOp + */ + virtual std::unique_ptr compute_absolute_linop() const = 0; + + /** + * Compute absolute inplace on each element. + */ + virtual void compute_absolute_inplace() = 0; +}; + + +/** + * The EnableAbsoluteComputation mixin provides the default implementations of + * `compute_absolute_linop` and the absolute interface. `compute_absolute` gets + * a new AbsoluteLinOp. `compute_absolute_inplace` applies absolute + * inplace, so it still keeps the value_type of the class. + * + * @tparam AbsoluteLinOp the absolute LinOp which is being returned + * [CRTP parameter] + * + * @ingroup LinOp + */ +template +class EnableAbsoluteComputation : public AbsoluteComputable { +public: + using absolute_type = AbsoluteLinOp; + + virtual ~EnableAbsoluteComputation() = default; + + std::unique_ptr compute_absolute_linop() const override + { + return this->compute_absolute(); + } + + /** + * Gets the AbsoluteLinOp + * + * @return a pointer to the new absolute object + */ + virtual std::unique_ptr compute_absolute() const = 0; +}; + + /** * The EnableLinOp mixin can be used to provide sensible default implementations * of the majority of the LinOp and PolymorphicObject interface. diff --git a/include/ginkgo/core/base/math.hpp b/include/ginkgo/core/base/math.hpp index 9f9dabb8f75..0bb6ac9b5a9 100644 --- a/include/ginkgo/core/base/math.hpp +++ b/include/ginkgo/core/base/math.hpp @@ -133,6 +133,27 @@ struct remove_complex_impl> { }; +/** + * Use the complex type if it is not complex. + * + * @tparam T the type being made complex + */ +template +struct to_complex_impl { + using type = std::complex; +}; + +/** + * Use the same type if it is complex type. + * + * @tparam T the type being made complex + */ +template +struct to_complex_impl> { + using type = std::complex; +}; + + template struct is_complex_impl : public std::integral_constant {}; @@ -141,6 +162,96 @@ struct is_complex_impl> : public std::integral_constant {}; +template +struct is_complex_or_scalar_impl : std::is_scalar {}; + +template +struct is_complex_or_scalar_impl> : std::is_scalar {}; + + +/** + * template_converter is converting the template parameters of a class by + * converter. + * + * @tparam converter which convert one type to another type + * @tparam T type + */ +template