Merge add absolute function, complex/real converter on class

This PR adds the absolute (inplace/outplace) of the matrix and complex<-> real on class

Summary
- add `to_complex` and to_real as remove_complex alias
- make `to_complex/real` and `remove_complex` work for class or scalar/complex
Note. Can not use `remove_complex<class>` as friend class. Need to use `class<remove_complex<V>, I>`. this type helper should work in the other situation.
- add `inplace_absolute_array` and `outplace_absolute_array`
- `format->compute_absolute_inplace()` is inplace absolute function
- `auto result = format->compute_absolute()` is outplace absolute function: format -> remove_complex<format>
- Add `AbsoluteComputable` interface and `compute_absolute_linop` can be used with LinOp.
- After this PR, need /bigobj for all MSVC.

Related PR: #636

.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

common/base/math.hpp.inc

@@ -54,10 +54,14 @@ struct is_complex_impl<thrust::complex<T>>
     : public std::integral_constant<bool, true> {};
 
 
+template <typename T>
+struct is_complex_or_scalar_impl<thrust::complex<T>> : std::is_scalar<T> {};
+
+
 template <typename T>
 struct truncate_type_impl<thrust::complex<T>> {
     using type = thrust::complex<typename truncate_type_impl<T>::type>;
 };
 
 
-}  // namespace detail
+}  // namespace detail

common/components/absolute_array.hpp.inc

@@ -0,0 +1,61 @@
+/*******************************<GINKGO LICENSE>******************************
+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.
+******************************<GINKGO LICENSE>*******************************/
+
+namespace kernel {
+
+
+template <typename ValueType>
+__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 <typename ValueType>
+__global__
+    __launch_bounds__(default_block_size) void outplace_absolute_array_kernel(
+        const size_type n, const ValueType *__restrict__ in,
+        remove_complex<ValueType> *__restrict__ out)
+{
+    const auto tidx = thread::get_thread_id_flat();
+    if (tidx < n) {
+        out[tidx] = abs(in[tidx]);
+    }
+}
+
+
+}  // namespace kernel

common/matrix/dense_kernels.hpp.inc

@@ -510,4 +510,33 @@ __global__ __launch_bounds__(default_block_size) void extract_diagonal(
 }
 
 
+template <typename ValueType>
+__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 <typename ValueType>
+__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<ValueType> *__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

core/base/extended_float.hpp

@@ -70,18 +70,17 @@ template <std::size_t, typename = void>
 struct uint_of_impl {};
 
 template <std::size_t Bits>
-struct uint_of_impl<Bits, xstd::void_t<std::enable_if_t<(Bits <= 16)>>> {
+struct uint_of_impl<Bits, std::enable_if_t<(Bits <= 16)>> {
     using type = uint16;
 };
 
 template <std::size_t Bits>
-struct uint_of_impl<Bits,
-                    xstd::void_t<std::enable_if_t<(16 < Bits && Bits <= 32)>>> {
+struct uint_of_impl<Bits, std::enable_if_t<(16 < Bits && Bits <= 32)>> {
     using type = uint32;
 };
 
 template <std::size_t Bits>
-struct uint_of_impl<Bits, xstd::void_t<std::enable_if_t<(32 < Bits)>>> {
+struct uint_of_impl<Bits, std::enable_if_t<(32 < Bits)>> {
     using type = uint64;
 };
 

core/components/absolute_array.hpp

@@ -0,0 +1,109 @@
+/*******************************<GINKGO LICENSE>******************************
+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.
+******************************<GINKGO LICENSE>*******************************/
+
+#ifndef GKO_CORE_COMPONENTS_ABSOLUTE_ARRAY_HPP_
+#define GKO_CORE_COMPONENTS_ABSOLUTE_ARRAY_HPP_
+
+
+#include <memory>
+
+
+#include <ginkgo/core/base/executor.hpp>
+#include <ginkgo/core/base/math.hpp>
+#include <ginkgo/core/base/types.hpp>
+
+
+namespace gko {
+namespace kernels {
+
+
+#define GKO_DECLARE_INPLACE_ABSOLUTE_ARRAY_KERNEL(ValueType)                 \
+    void inplace_absolute_array(std::shared_ptr<const DefaultExecutor> exec, \
+                                ValueType *data, size_type num_entries)
+
+#define GKO_DECLARE_OUTPLACE_ABSOLUTE_ARRAY_KERNEL(ValueType)                 \
+    void outplace_absolute_array(std::shared_ptr<const DefaultExecutor> exec, \
+                                 const ValueType *in, size_type num_entries,  \
+                                 remove_complex<ValueType> *out)
+
+
+#define GKO_DECLARE_ALL_AS_TEMPLATES                      \
+    template <typename ValueType>                         \
+    GKO_DECLARE_INPLACE_ABSOLUTE_ARRAY_KERNEL(ValueType); \
+    template <typename ValueType>                         \
+    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_

core/device_hooks/common_kernels.inc.cpp

@@ -33,6 +33,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include <ginkgo/core/base/exception_helpers.hpp>
 
 
+#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 <typename ValueType>
+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 <typename ValueType>
+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 <typename ValueType>
+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 <typename ValueType>
+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
 

core/matrix/coo.cpp

@@ -45,6 +45,7 @@ OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 #include <ginkgo/core/matrix/dense.hpp>
 
 
+#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<ValueType, IndexType>::extract_diagonal() const
 }
 
 
+template <typename ValueType, typename IndexType>
+void Coo<ValueType, IndexType>::compute_absolute_inplace()
+{
+    auto exec = this->get_executor();
+
+    exec->run(coo::make_inplace_absolute_array(
+        this->get_values(), this->get_num_stored_elements()));
+}
+
+
+template <typename ValueType, typename IndexType>
+std::unique_ptr<typename Coo<ValueType, IndexType>::absolute_type>
+Coo<ValueType, IndexType>::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<ValueType, IndexType>
 GKO_INSTANTIATE_FOR_EACH_VALUE_AND_INDEX_TYPE(GKO_DECLARE_COO_MATRIX);