From 974b176384282a8a2d4b11de6f45932028ed390d Mon Sep 17 00:00:00 2001 From: Fritz Goebel Date: Thu, 22 Oct 2020 15:27:37 +0200 Subject: [PATCH 1/7] Add conversion to complex and extraction of real and imaginary part to dense --- common/matrix/dense_kernels.hpp.inc | 45 ++++++++++ core/device_hooks/common_kernels.inc.cpp | 15 ++++ core/matrix/dense.cpp | 45 ++++++++++ core/matrix/dense_kernels.hpp | 23 ++++- cuda/matrix/dense_kernels.cu | 51 +++++++++++ cuda/test/matrix/dense_kernels.cpp | 33 ++++++++ hip/matrix/dense_kernels.hip.cpp | 54 ++++++++++++ hip/test/matrix/dense_kernels.hip.cpp | 33 ++++++++ include/ginkgo/core/matrix/dense.hpp | 7 ++ omp/matrix/dense_kernels.cpp | 54 ++++++++++++ omp/test/matrix/dense_kernels.cpp | 33 ++++++++ reference/matrix/dense_kernels.cpp | 48 +++++++++++ reference/test/matrix/dense_kernels.cpp | 103 +++++++++++++++++++++++ 13 files changed, 543 insertions(+), 1 deletion(-) diff --git a/common/matrix/dense_kernels.hpp.inc b/common/matrix/dense_kernels.hpp.inc index c26203115fe..b4649c40fdd 100644 --- a/common/matrix/dense_kernels.hpp.inc +++ b/common/matrix/dense_kernels.hpp.inc @@ -539,4 +539,49 @@ __global__ __launch_bounds__(default_block_size) void outplace_absolute_dense( } +template +__global__ __launch_bounds__(default_block_size) void make_complex( + size_type num_rows, size_type num_cols, const ValueType *__restrict__ in, + size_type stride_in, to_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] = to_complex{in[row * stride_in + col]}; + } +} + + +template +__global__ __launch_bounds__(default_block_size) void get_real( + size_type num_rows, size_type num_cols, const ValueType *__restrict__ in, + size_type stride_in, to_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] = real(in[row * stride_in + col]); + } +} + + +template +__global__ __launch_bounds__(default_block_size) void get_imag( + size_type num_rows, size_type num_cols, const ValueType *__restrict__ in, + size_type stride_in, to_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] = imag(in[row * stride_in + col]); + } +} + + } // namespace kernel diff --git a/core/device_hooks/common_kernels.inc.cpp b/core/device_hooks/common_kernels.inc.cpp index a981372c630..34992284498 100644 --- a/core/device_hooks/common_kernels.inc.cpp +++ b/core/device_hooks/common_kernels.inc.cpp @@ -256,6 +256,21 @@ 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); +template +GKO_DECLARE_MAKE_COMPLEX_KERNEL(ValueType) +GKO_NOT_COMPILED(GKO_HOOK_MODULE); +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_MAKE_COMPLEX_KERNEL); + +template +GKO_DECLARE_GET_REAL_KERNEL(ValueType) +GKO_NOT_COMPILED(GKO_HOOK_MODULE); +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_REAL_KERNEL); + +template +GKO_DECLARE_GET_IMAG_KERNEL(ValueType) +GKO_NOT_COMPILED(GKO_HOOK_MODULE); +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_IMAG_KERNEL); + } // namespace dense diff --git a/core/matrix/dense.cpp b/core/matrix/dense.cpp index 2e0b663966f..398b5192035 100644 --- a/core/matrix/dense.cpp +++ b/core/matrix/dense.cpp @@ -88,6 +88,9 @@ 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); +GKO_REGISTER_OPERATION(make_complex, dense::make_complex); +GKO_REGISTER_OPERATION(get_real, dense::get_real); +GKO_REGISTER_OPERATION(get_imag, dense::get_imag); } // namespace dense @@ -795,6 +798,48 @@ Dense::compute_absolute() const } +template +std::unique_ptr::complex_type> +Dense::make_complex() const +{ + auto exec = this->get_executor(); + + auto complex_dense = complex_type::create(exec, this->get_size()); + + exec->run(dense::make_make_complex(this, complex_dense.get())); + + return complex_dense; +} + + +template +std::unique_ptr::absolute_type> +Dense::get_real() const +{ + auto exec = this->get_executor(); + + auto real_dense = absolute_type::create(exec, this->get_size()); + + exec->run(dense::make_get_real(this, real_dense.get())); + + return real_dense; +} + + +template +std::unique_ptr::absolute_type> +Dense::get_imag() const +{ + auto exec = this->get_executor(); + + auto imag_dense = absolute_type::create(exec, this->get_size()); + + exec->run(dense::make_get_imag(this, imag_dense.get())); + + return imag_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 1f9bdae1cc2..e99d455aa77 100644 --- a/core/matrix/dense_kernels.hpp +++ b/core/matrix/dense_kernels.hpp @@ -182,6 +182,21 @@ namespace kernels { const matrix::Dense<_vtype> *source, \ matrix::Dense> *result) +#define GKO_DECLARE_MAKE_COMPLEX_KERNEL(_vtype) \ + void make_complex(std::shared_ptr exec, \ + const matrix::Dense<_vtype> *source, \ + matrix::Dense> *result) + +#define GKO_DECLARE_GET_REAL_KERNEL(_vtype) \ + void get_real(std::shared_ptr exec, \ + const matrix::Dense<_vtype> *source, \ + matrix::Dense> *result) + +#define GKO_DECLARE_GET_IMAG_KERNEL(_vtype) \ + void get_imag(std::shared_ptr exec, \ + const matrix::Dense<_vtype> *source, \ + matrix::Dense> *result) + #define GKO_DECLARE_ALL_AS_TEMPLATES \ template \ @@ -235,7 +250,13 @@ namespace kernels { template \ GKO_DECLARE_INPLACE_ABSOLUTE_DENSE_KERNEL(ValueType); \ template \ - GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL(ValueType) + GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL(ValueType); \ + template \ + GKO_DECLARE_MAKE_COMPLEX_KERNEL(ValueType); \ + template \ + GKO_DECLARE_GET_REAL_KERNEL(ValueType); \ + template \ + GKO_DECLARE_GET_IMAG_KERNEL(ValueType) namespace omp { diff --git a/cuda/matrix/dense_kernels.cu b/cuda/matrix/dense_kernels.cu index 72d557e7dd1..83804b43083 100644 --- a/cuda/matrix/dense_kernels.cu +++ b/cuda/matrix/dense_kernels.cu @@ -734,6 +734,57 @@ void outplace_absolute_dense(std::shared_ptr exec, GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL); +template +void make_complex(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::make_complex<<>>( + 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_MAKE_COMPLEX_KERNEL); + + +template +void get_real(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::get_real<<>>( + 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_GET_REAL_KERNEL); + + +template +void get_imag(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::get_imag<<>>( + 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_GET_IMAG_KERNEL); + + } // namespace dense } // namespace cuda } // namespace kernels diff --git a/cuda/test/matrix/dense_kernels.cpp b/cuda/test/matrix/dense_kernels.cpp index 1dee85ba561..ac8f0262978 100644 --- a/cuda/test/matrix/dense_kernels.cpp +++ b/cuda/test/matrix/dense_kernels.cpp @@ -643,4 +643,37 @@ TEST_F(Dense, OutplaceAbsoluteMatrixIsEquivalentToRef) } +TEST_F(Dense, MakeComplexIsEquivalentToRef) +{ + set_up_apply_data(); + + auto complex_x = x->make_complex(); + auto dcomplex_x = dx->make_complex(); + + GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 1e-14); +} + + +TEST_F(Dense, GetRealIsEquivalentToRef) +{ + set_up_apply_data(); + + auto real_x = x->get_real(); + auto dreal_x = dx->get_real(); + + GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 1e-14); +} + + +TEST_F(Dense, GetImagIsEquivalentToRef) +{ + set_up_apply_data(); + + auto imag_x = x->get_imag(); + auto dimag_x = dx->get_imag(); + + GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 1e-14); +} + + } // namespace diff --git a/hip/matrix/dense_kernels.hip.cpp b/hip/matrix/dense_kernels.hip.cpp index fa732e1e86a..e5de9f7df83 100644 --- a/hip/matrix/dense_kernels.hip.cpp +++ b/hip/matrix/dense_kernels.hip.cpp @@ -763,6 +763,60 @@ void outplace_absolute_dense(std::shared_ptr exec, GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL); +template +void make_complex(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::make_complex, 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_MAKE_COMPLEX_KERNEL); + + +template +void get_real(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::get_real, 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_GET_REAL_KERNEL); + + +template +void get_imag(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::get_imag, 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_GET_IMAG_KERNEL); + + } // namespace dense } // namespace hip } // namespace kernels diff --git a/hip/test/matrix/dense_kernels.hip.cpp b/hip/test/matrix/dense_kernels.hip.cpp index fb40a3819fa..8385fd21205 100644 --- a/hip/test/matrix/dense_kernels.hip.cpp +++ b/hip/test/matrix/dense_kernels.hip.cpp @@ -626,4 +626,37 @@ TEST_F(Dense, OutplaceAbsoluteMatrixIsEquivalentToRef) } +TEST_F(Dense, MakeComplexIsEquivalentToRef) +{ + set_up_apply_data(); + + auto complex_x = x->make_complex(); + auto dcomplex_x = dx->make_complex(); + + GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 1e-14); +} + + +TEST_F(Dense, GetRealIsEquivalentToRef) +{ + set_up_apply_data(); + + auto real_x = x->get_real(); + auto dreal_x = dx->get_real(); + + GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 1e-14); +} + + +TEST_F(Dense, GetImagIsEquivalentToRef) +{ + set_up_apply_data(); + + auto imag_x = x->get_imag(); + auto dimag_x = dx->get_imag(); + + GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 1e-14); +} + + } // namespace diff --git a/include/ginkgo/core/matrix/dense.hpp b/include/ginkgo/core/matrix/dense.hpp index 65113ec204b..ce9e85d5d16 100644 --- a/include/ginkgo/core/matrix/dense.hpp +++ b/include/ginkgo/core/matrix/dense.hpp @@ -142,6 +142,7 @@ class Dense using mat_data = gko::matrix_data; using mat_data32 = gko::matrix_data; using absolute_type = remove_complex; + using complex_type = to_complex; using row_major_range = gko::range>; @@ -255,6 +256,12 @@ class Dense void compute_absolute_inplace() override; + std::unique_ptr make_complex() const; + + std::unique_ptr get_real() const; + + std::unique_ptr get_imag() const; + /** * Returns a pointer to the array of values of the matrix. * diff --git a/omp/matrix/dense_kernels.cpp b/omp/matrix/dense_kernels.cpp index afe74e3c5e4..15a9d6e9a99 100644 --- a/omp/matrix/dense_kernels.cpp +++ b/omp/matrix/dense_kernels.cpp @@ -805,6 +805,60 @@ void outplace_absolute_dense(std::shared_ptr exec, GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL); +template +void make_complex(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Dense> *result) +{ + auto dim = source->get_size(); + +#pragma omp parallel for + for (size_type row = 0; row < dim[0]; row++) { + for (size_type col = 0; col < dim[1]; col++) { + result->at(row, col) = to_complex{source->at(row, col)}; + } + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_MAKE_COMPLEX_KERNEL); + + +template +void get_real(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Dense> *result) +{ + auto dim = source->get_size(); + +#pragma omp parallel for + for (size_type row = 0; row < dim[0]; row++) { + for (size_type col = 0; col < dim[1]; col++) { + result->at(row, col) = real(source->at(row, col)); + } + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_REAL_KERNEL); + + +template +void get_imag(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Dense> *result) +{ + auto dim = source->get_size(); + +#pragma omp parallel for + for (size_type row = 0; row < dim[0]; row++) { + for (size_type col = 0; col < dim[1]; col++) { + result->at(row, col) = imag(source->at(row, col)); + } + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_IMAG_KERNEL); + + } // namespace dense } // namespace omp } // namespace kernels diff --git a/omp/test/matrix/dense_kernels.cpp b/omp/test/matrix/dense_kernels.cpp index 06999748b9f..90c7880bd85 100644 --- a/omp/test/matrix/dense_kernels.cpp +++ b/omp/test/matrix/dense_kernels.cpp @@ -758,4 +758,37 @@ TEST_F(Dense, OutplaceAbsoluteMatrixIsEquivalentToRef) } +TEST_F(Dense, MakeComplexIsEquivalentToRef) +{ + set_up_apply_data(); + + auto complex_x = x->make_complex(); + auto dcomplex_x = dx->make_complex(); + + GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 1e-14); +} + + +TEST_F(Dense, GetRealIsEquivalentToRef) +{ + set_up_apply_data(); + + auto real_x = x->get_real(); + auto dreal_x = dx->get_real(); + + GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 1e-14); +} + + +TEST_F(Dense, GetImagIsEquivalentToRef) +{ + set_up_apply_data(); + + auto imag_x = x->get_imag(); + auto dimag_x = dx->get_imag(); + + GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 1e-14); +} + + } // namespace diff --git a/reference/matrix/dense_kernels.cpp b/reference/matrix/dense_kernels.cpp index f0c8f938d05..147204157e7 100644 --- a/reference/matrix/dense_kernels.cpp +++ b/reference/matrix/dense_kernels.cpp @@ -716,6 +716,54 @@ void outplace_absolute_dense(std::shared_ptr exec, GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL); +template +void make_complex(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Dense> *result) +{ + auto dim = source->get_size(); + for (size_type row = 0; row < dim[0]; row++) { + for (size_type col = 0; col < dim[1]; col++) { + result->at(row, col) = to_complex{source->at(row, col)}; + } + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_MAKE_COMPLEX_KERNEL); + + +template +void get_real(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Dense> *result) +{ + auto dim = source->get_size(); + for (size_type row = 0; row < dim[0]; row++) { + for (size_type col = 0; col < dim[1]; col++) { + result->at(row, col) = real(source->at(row, col)); + } + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_REAL_KERNEL); + + +template +void get_imag(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Dense> *result) +{ + auto dim = source->get_size(); + for (size_type row = 0; row < dim[0]; row++) { + for (size_type col = 0; col < dim[1]; col++) { + result->at(row, col) = imag(source->at(row, col)); + } + } +} + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_IMAG_KERNEL); + + } // namespace dense } // namespace reference } // namespace kernels diff --git a/reference/test/matrix/dense_kernels.cpp b/reference/test/matrix/dense_kernels.cpp index b6195ffa8b0..b31fd51e52a 100644 --- a/reference/test/matrix/dense_kernels.cpp +++ b/reference/test/matrix/dense_kernels.cpp @@ -2151,6 +2151,53 @@ TYPED_TEST(Dense, AdvancedAppliesToComplex) } +TYPED_TEST(Dense, MakeComplex) +{ + using T = typename TestFixture::value_type; + // clang-format off + // {1.0, -1.0, -0.5}, + // {-2.0, 2.0, 4.5}, + // {2.1, 3.4, 1.2} + // clang-format on + + auto abs_mtx = this->mtx5->make_complex(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, this->mtx5, r::value); +} + + +TYPED_TEST(Dense, GetReal) +{ + using T = typename TestFixture::value_type; + // clang-format off + // {1.0, -1.0, -0.5}, + // {-2.0, 2.0, 4.5}, + // {2.1, 3.4, 1.2} + // clang-format on + + auto abs_mtx = this->mtx5->get_real(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, this->mtx5, r::value); +} + + +TYPED_TEST(Dense, GetImag) +{ + using T = typename TestFixture::value_type; + // clang-format off + // {1.0, -1.0, -0.5}, + // {-2.0, 2.0, 4.5}, + // {2.1, 3.4, 1.2} + // clang-format on + + auto abs_mtx = this->mtx5->get_imag(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, + l({{0.0, 0.0, 0.0}, {0.0, 0.0, 0.0}, {0.0, 0.0, 0.0}}), + r::value); +} + + template class DenseComplex : public ::testing::Test { protected: @@ -2220,4 +2267,60 @@ TYPED_TEST(DenseComplex, OutplaceAbsolute) } +TYPED_TEST(DenseComplex, MakeComplex) +{ + using Mtx = typename TestFixture::Mtx; + using T = typename TestFixture::value_type; + auto exec = gko::ReferenceExecutor::create(); + // clang-format off + auto mtx = gko::initialize( + {{T{1.0, 0.0}, T{3.0, 4.0}, T{0.0, 2.0}}, + {T{-4.0, -3.0}, T{-1.0, 0}, T{0.0, 0.0}}, + {T{0.0, 0.0}, T{0.0, -1.5}, T{2.0, 0.0}}}, exec); + // clang-format on + + auto abs_mtx = mtx->make_complex(); + + GKO_ASSERT_MTX_NEAR(abs_mtx, mtx, 0.0); +} + + +TYPED_TEST(DenseComplex, GetReal) +{ + using Mtx = typename TestFixture::Mtx; + using T = typename TestFixture::value_type; + auto exec = gko::ReferenceExecutor::create(); + // clang-format off + auto mtx = gko::initialize( + {{T{1.0, 0.0}, T{3.0, 4.0}, T{0.0, 2.0}}, + {T{-4.0, -3.0}, T{-1.0, 0}, T{0.0, 0.0}}, + {T{0.0, 0.0}, T{0.0, -1.5}, T{2.0, 0.0}}}, exec); + // clang-format on + + auto abs_mtx = mtx->get_real(); + + GKO_ASSERT_MTX_NEAR( + abs_mtx, l({{1.0, 3.0, 0.0}, {-4.0, -1.0, 0.0}, {0.0, 0.0, 2.0}}), 0.0); +} + + +TYPED_TEST(DenseComplex, GetImag) +{ + using Mtx = typename TestFixture::Mtx; + using T = typename TestFixture::value_type; + auto exec = gko::ReferenceExecutor::create(); + // clang-format off + auto mtx = gko::initialize( + {{T{1.0, 0.0}, T{3.0, 4.0}, T{0.0, 2.0}}, + {T{-4.0, -3.0}, T{-1.0, 0}, T{0.0, 0.0}}, + {T{0.0, 0.0}, T{0.0, -1.5}, T{2.0, 0.0}}}, exec); + // clang-format on + + auto abs_mtx = mtx->get_imag(); + + GKO_ASSERT_MTX_NEAR( + abs_mtx, l({{0.0, 4.0, 2.0}, {-3.0, 0.0, 0.0}, {0.0, -1.5, 0.0}}), 0.0); +} + + } // namespace From 39b39e3a43c7750c78ae83d225a0e612d6da7c0e Mon Sep 17 00:00:00 2001 From: Fritz Goebel Date: Mon, 26 Oct 2020 11:55:49 +0100 Subject: [PATCH 2/7] Fix wrong parameter types in gpu kernels --- common/matrix/dense_kernels.hpp.inc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/common/matrix/dense_kernels.hpp.inc b/common/matrix/dense_kernels.hpp.inc index b4649c40fdd..8adeaed9521 100644 --- a/common/matrix/dense_kernels.hpp.inc +++ b/common/matrix/dense_kernels.hpp.inc @@ -557,7 +557,7 @@ __global__ __launch_bounds__(default_block_size) void make_complex( template __global__ __launch_bounds__(default_block_size) void get_real( size_type num_rows, size_type num_cols, const ValueType *__restrict__ in, - size_type stride_in, to_complex *__restrict__ out, + size_type stride_in, remove_complex *__restrict__ out, size_type stride_out) { const auto tidx = thread::get_thread_id_flat(); @@ -572,7 +572,7 @@ __global__ __launch_bounds__(default_block_size) void get_real( template __global__ __launch_bounds__(default_block_size) void get_imag( size_type num_rows, size_type num_cols, const ValueType *__restrict__ in, - size_type stride_in, to_complex *__restrict__ out, + size_type stride_in, remove_complex *__restrict__ out, size_type stride_out) { const auto tidx = thread::get_thread_id_flat(); From a07a28d5315039da18de3d304f8818596cde544b Mon Sep 17 00:00:00 2001 From: Fritz Goebel Date: Tue, 27 Oct 2020 15:34:48 +0100 Subject: [PATCH 3/7] Apply suggestions from code review --- common/matrix/dense_kernels.hpp.inc | 7 +++---- cuda/test/matrix/dense_kernels.cpp | 6 +++--- hip/test/matrix/dense_kernels.hip.cpp | 6 +++--- include/ginkgo/core/matrix/dense.hpp | 12 ++++++++++++ omp/test/matrix/dense_kernels.cpp | 6 +++--- 5 files changed, 24 insertions(+), 13 deletions(-) diff --git a/common/matrix/dense_kernels.hpp.inc b/common/matrix/dense_kernels.hpp.inc index 8adeaed9521..a9947ad891f 100644 --- a/common/matrix/dense_kernels.hpp.inc +++ b/common/matrix/dense_kernels.hpp.inc @@ -539,17 +539,16 @@ __global__ __launch_bounds__(default_block_size) void outplace_absolute_dense( } -template +template __global__ __launch_bounds__(default_block_size) void make_complex( size_type num_rows, size_type num_cols, const ValueType *__restrict__ in, - size_type stride_in, to_complex *__restrict__ out, - size_type stride_out) + size_type stride_in, ComplexType *__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] = to_complex{in[row * stride_in + col]}; + out[row * stride_out + col] = in[row * stride_in + col]; } } diff --git a/cuda/test/matrix/dense_kernels.cpp b/cuda/test/matrix/dense_kernels.cpp index ac8f0262978..ea067231ece 100644 --- a/cuda/test/matrix/dense_kernels.cpp +++ b/cuda/test/matrix/dense_kernels.cpp @@ -650,7 +650,7 @@ TEST_F(Dense, MakeComplexIsEquivalentToRef) auto complex_x = x->make_complex(); auto dcomplex_x = dx->make_complex(); - GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 1e-14); + GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 0); } @@ -661,7 +661,7 @@ TEST_F(Dense, GetRealIsEquivalentToRef) auto real_x = x->get_real(); auto dreal_x = dx->get_real(); - GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 1e-14); + GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 0); } @@ -672,7 +672,7 @@ TEST_F(Dense, GetImagIsEquivalentToRef) auto imag_x = x->get_imag(); auto dimag_x = dx->get_imag(); - GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 1e-14); + GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 0); } diff --git a/hip/test/matrix/dense_kernels.hip.cpp b/hip/test/matrix/dense_kernels.hip.cpp index 8385fd21205..5268e6d1878 100644 --- a/hip/test/matrix/dense_kernels.hip.cpp +++ b/hip/test/matrix/dense_kernels.hip.cpp @@ -633,7 +633,7 @@ TEST_F(Dense, MakeComplexIsEquivalentToRef) auto complex_x = x->make_complex(); auto dcomplex_x = dx->make_complex(); - GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 1e-14); + GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 0); } @@ -644,7 +644,7 @@ TEST_F(Dense, GetRealIsEquivalentToRef) auto real_x = x->get_real(); auto dreal_x = dx->get_real(); - GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 1e-14); + GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 0); } @@ -655,7 +655,7 @@ TEST_F(Dense, GetImagIsEquivalentToRef) auto imag_x = x->get_imag(); auto dimag_x = dx->get_imag(); - GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 1e-14); + GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 0); } diff --git a/include/ginkgo/core/matrix/dense.hpp b/include/ginkgo/core/matrix/dense.hpp index ce9e85d5d16..b7359599104 100644 --- a/include/ginkgo/core/matrix/dense.hpp +++ b/include/ginkgo/core/matrix/dense.hpp @@ -256,10 +256,22 @@ class Dense void compute_absolute_inplace() override; + /** + * Creates a complex copy of the original matrix. If the original matrix + * was real, the imaginary part of the result will be zero. + */ std::unique_ptr make_complex() const; + /** + * Creates a new real matrix and extracts the real part of the original + * matrix into that. + */ std::unique_ptr get_real() const; + /** + * Creates a new real matrix and extracts the imaginary part of the + * original matrix into that. + */ std::unique_ptr get_imag() const; /** diff --git a/omp/test/matrix/dense_kernels.cpp b/omp/test/matrix/dense_kernels.cpp index 90c7880bd85..fc5089fcf8a 100644 --- a/omp/test/matrix/dense_kernels.cpp +++ b/omp/test/matrix/dense_kernels.cpp @@ -765,7 +765,7 @@ TEST_F(Dense, MakeComplexIsEquivalentToRef) auto complex_x = x->make_complex(); auto dcomplex_x = dx->make_complex(); - GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 1e-14); + GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 0); } @@ -776,7 +776,7 @@ TEST_F(Dense, GetRealIsEquivalentToRef) auto real_x = x->get_real(); auto dreal_x = dx->get_real(); - GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 1e-14); + GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 0); } @@ -787,7 +787,7 @@ TEST_F(Dense, GetImagIsEquivalentToRef) auto imag_x = x->get_imag(); auto dimag_x = dx->get_imag(); - GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 1e-14); + GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 0); } From db511c595f78f35be35397a2a0fb767b2fcc272c Mon Sep 17 00:00:00 2001 From: Fritz Goebel Date: Wed, 28 Oct 2020 11:13:28 +0100 Subject: [PATCH 4/7] Add conversions into given matrix --- core/matrix/dense.cpp | 33 ++++ cuda/test/matrix/dense_kernels.cpp | 39 +++++ hip/test/matrix/dense_kernels.hip.cpp | 39 +++++ include/ginkgo/core/matrix/dense.hpp | 18 +++ omp/test/matrix/dense_kernels.cpp | 39 +++++ reference/test/matrix/dense_kernels.cpp | 200 ++++++++++++++++++++++-- 6 files changed, 354 insertions(+), 14 deletions(-) diff --git a/core/matrix/dense.cpp b/core/matrix/dense.cpp index 398b5192035..2cb02fb4361 100644 --- a/core/matrix/dense.cpp +++ b/core/matrix/dense.cpp @@ -812,6 +812,17 @@ Dense::make_complex() const } +template +void Dense::make_complex(Dense> *result) const +{ + auto exec = this->get_executor(); + + GKO_ASSERT_EQUAL_DIMENSIONS(this, result); + + exec->run(dense::make_make_complex(this, result)); +} + + template std::unique_ptr::absolute_type> Dense::get_real() const @@ -826,6 +837,17 @@ Dense::get_real() const } +template +void Dense::get_real(Dense> *result) const +{ + auto exec = this->get_executor(); + + GKO_ASSERT_EQUAL_DIMENSIONS(this, result); + + exec->run(dense::make_get_real(this, result)); +} + + template std::unique_ptr::absolute_type> Dense::get_imag() const @@ -840,6 +862,17 @@ Dense::get_imag() const } +template +void Dense::get_imag(Dense> *result) const +{ + auto exec = this->get_executor(); + + GKO_ASSERT_EQUAL_DIMENSIONS(this, result); + + exec->run(dense::make_get_imag(this, result)); +} + + #define GKO_DECLARE_DENSE_MATRIX(_type) class Dense<_type> GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_DENSE_MATRIX); diff --git a/cuda/test/matrix/dense_kernels.cpp b/cuda/test/matrix/dense_kernels.cpp index ea067231ece..823aabe9b2b 100644 --- a/cuda/test/matrix/dense_kernels.cpp +++ b/cuda/test/matrix/dense_kernels.cpp @@ -654,6 +654,19 @@ TEST_F(Dense, MakeComplexIsEquivalentToRef) } +TEST_F(Dense, MakeComplexWithGivenResultIsEquivalentToRef) +{ + set_up_apply_data(); + + auto complex_x = ComplexMtx::create(ref, x->get_size()); + x->make_complex(complex_x.get()); + auto dcomplex_x = ComplexMtx::create(cuda, x->get_size()); + dx->make_complex(dcomplex_x.get()); + + GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 0); +} + + TEST_F(Dense, GetRealIsEquivalentToRef) { set_up_apply_data(); @@ -665,6 +678,19 @@ TEST_F(Dense, GetRealIsEquivalentToRef) } +TEST_F(Dense, GetRealWithGivenResultIsEquivalentToRef) +{ + set_up_apply_data(); + + auto real_x = Mtx::create(ref, x->get_size()); + x->get_real(real_x.get()); + auto dreal_x = Mtx::create(cuda, dx->get_size()); + dx->get_real(dreal_x.get()); + + GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 0); +} + + TEST_F(Dense, GetImagIsEquivalentToRef) { set_up_apply_data(); @@ -676,4 +702,17 @@ TEST_F(Dense, GetImagIsEquivalentToRef) } +TEST_F(Dense, GetImagWithGivenResultIsEquivalentToRef) +{ + set_up_apply_data(); + + auto imag_x = Mtx::create(ref, x->get_size()); + x->get_imag(imag_x.get()); + auto dimag_x = Mtx::create(cuda, dx->get_size()); + dx->get_imag(dimag_x.get()); + + GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 0); +} + + } // namespace diff --git a/hip/test/matrix/dense_kernels.hip.cpp b/hip/test/matrix/dense_kernels.hip.cpp index 5268e6d1878..a4f4b4e58f1 100644 --- a/hip/test/matrix/dense_kernels.hip.cpp +++ b/hip/test/matrix/dense_kernels.hip.cpp @@ -637,6 +637,19 @@ TEST_F(Dense, MakeComplexIsEquivalentToRef) } +TEST_F(Dense, MakeComplexWithGivenResultIsEquivalentToRef) +{ + set_up_apply_data(); + + auto complex_x = ComplexMtx::create(ref, x->get_size()); + x->make_complex(complex_x.get()); + auto dcomplex_x = ComplexMtx::create(hip, x->get_size()); + dx->make_complex(dcomplex_x.get()); + + GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 0); +} + + TEST_F(Dense, GetRealIsEquivalentToRef) { set_up_apply_data(); @@ -648,6 +661,19 @@ TEST_F(Dense, GetRealIsEquivalentToRef) } +TEST_F(Dense, GetRealWithGivenResultIsEquivalentToRef) +{ + set_up_apply_data(); + + auto real_x = Mtx::create(ref, x->get_size()); + x->get_real(real_x.get()); + auto dreal_x = Mtx::create(hip, dx->get_size()); + dx->get_real(dreal_x.get()); + + GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 0); +} + + TEST_F(Dense, GetImagIsEquivalentToRef) { set_up_apply_data(); @@ -659,4 +685,17 @@ TEST_F(Dense, GetImagIsEquivalentToRef) } +TEST_F(Dense, GetImagWithGivenResultIsEquivalentToRef) +{ + set_up_apply_data(); + + auto imag_x = Mtx::create(ref, x->get_size()); + x->get_imag(imag_x.get()); + auto dimag_x = Mtx::create(hip, dx->get_size()); + dx->get_imag(dimag_x.get()); + + GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 0); +} + + } // namespace diff --git a/include/ginkgo/core/matrix/dense.hpp b/include/ginkgo/core/matrix/dense.hpp index b7359599104..7bb439afc34 100644 --- a/include/ginkgo/core/matrix/dense.hpp +++ b/include/ginkgo/core/matrix/dense.hpp @@ -262,18 +262,36 @@ class Dense */ std::unique_ptr make_complex() const; + /** + * Writes a complex copy of the original matrix to a given complex matrix. + * If the original matrix was real, the imaginary part of the result will + * be zero. + */ + void make_complex(Dense> *result) const; + /** * Creates a new real matrix and extracts the real part of the original * matrix into that. */ std::unique_ptr get_real() const; + /** + * Extracts the real part of the original matrix into a given real matrix. + */ + void get_real(Dense> *result) const; + /** * Creates a new real matrix and extracts the imaginary part of the * original matrix into that. */ std::unique_ptr get_imag() const; + /** + * Extracts the imaginary part of the original matrix into a given real + * matrix. + */ + void get_imag(Dense> *result) const; + /** * Returns a pointer to the array of values of the matrix. * diff --git a/omp/test/matrix/dense_kernels.cpp b/omp/test/matrix/dense_kernels.cpp index fc5089fcf8a..cf2d975c2ef 100644 --- a/omp/test/matrix/dense_kernels.cpp +++ b/omp/test/matrix/dense_kernels.cpp @@ -769,6 +769,19 @@ TEST_F(Dense, MakeComplexIsEquivalentToRef) } +TEST_F(Dense, MakeComplexWithGivenResultIsEquivalentToRef) +{ + set_up_apply_data(); + + auto complex_x = ComplexMtx::create(ref, x->get_size()); + x->make_complex(complex_x.get()); + auto dcomplex_x = ComplexMtx::create(omp, x->get_size()); + dx->make_complex(dcomplex_x.get()); + + GKO_ASSERT_MTX_NEAR(complex_x, dcomplex_x, 0); +} + + TEST_F(Dense, GetRealIsEquivalentToRef) { set_up_apply_data(); @@ -780,6 +793,19 @@ TEST_F(Dense, GetRealIsEquivalentToRef) } +TEST_F(Dense, GetRealWithGivenResultIsEquivalentToRef) +{ + set_up_apply_data(); + + auto real_x = Mtx::create(ref, x->get_size()); + x->get_real(real_x.get()); + auto dreal_x = Mtx::create(omp, dx->get_size()); + dx->get_real(dreal_x.get()); + + GKO_ASSERT_MTX_NEAR(real_x, dreal_x, 0); +} + + TEST_F(Dense, GetImagIsEquivalentToRef) { set_up_apply_data(); @@ -791,4 +817,17 @@ TEST_F(Dense, GetImagIsEquivalentToRef) } +TEST_F(Dense, GetImagWithGivenResultIsEquivalentToRef) +{ + set_up_apply_data(); + + auto imag_x = Mtx::create(ref, x->get_size()); + x->get_imag(imag_x.get()); + auto dimag_x = Mtx::create(omp, dx->get_size()); + dx->get_imag(dimag_x.get()); + + GKO_ASSERT_MTX_NEAR(imag_x, dimag_x, 0); +} + + } // namespace diff --git a/reference/test/matrix/dense_kernels.cpp b/reference/test/matrix/dense_kernels.cpp index b31fd51e52a..562176aefbe 100644 --- a/reference/test/matrix/dense_kernels.cpp +++ b/reference/test/matrix/dense_kernels.cpp @@ -65,6 +65,8 @@ class Dense : public ::testing::Test { protected: using value_type = T; using Mtx = gko::matrix::Dense; + using ComplexMtx = gko::matrix::Dense>; + using RealMtx = gko::matrix::Dense>; Dense() : exec(gko::ReferenceExecutor::create()), mtx1(gko::initialize(4, {{1.0, 2.0, 3.0}, {1.5, 2.5, 3.5}}, @@ -2160,9 +2162,45 @@ TYPED_TEST(Dense, MakeComplex) // {2.1, 3.4, 1.2} // clang-format on - auto abs_mtx = this->mtx5->make_complex(); + auto complex_mtx = this->mtx5->make_complex(); - GKO_ASSERT_MTX_NEAR(abs_mtx, this->mtx5, r::value); + GKO_ASSERT_MTX_NEAR(complex_mtx, this->mtx5, 0); +} + + +TYPED_TEST(Dense, MakeComplexWithGivenResult) +{ + using T = typename TestFixture::value_type; + using ComplexMtx = typename TestFixture::ComplexMtx; + auto exec = this->mtx5->get_executor(); + // clang-format off + // {1.0, -1.0, -0.5}, + // {-2.0, 2.0, 4.5}, + // {2.1, 3.4, 1.2} + // clang-format on + + auto complex_mtx = ComplexMtx::create(exec, this->mtx5->get_size()); + this->mtx5->make_complex(complex_mtx.get()); + + GKO_ASSERT_MTX_NEAR(complex_mtx, this->mtx5, 0); +} + + +TYPED_TEST(Dense, MakeComplexWithGivenResultFailsForWrongDimensions) +{ + using T = typename TestFixture::value_type; + using ComplexMtx = typename TestFixture::ComplexMtx; + auto exec = this->mtx5->get_executor(); + // clang-format off + // {1.0, -1.0, -0.5}, + // {-2.0, 2.0, 4.5}, + // {2.1, 3.4, 1.2} + // clang-format on + + auto complex_mtx = ComplexMtx::create(exec); + + ASSERT_THROW(this->mtx5->make_complex(complex_mtx.get()), + gko::DimensionMismatch); } @@ -2175,9 +2213,43 @@ TYPED_TEST(Dense, GetReal) // {2.1, 3.4, 1.2} // clang-format on - auto abs_mtx = this->mtx5->get_real(); + auto real_mtx = this->mtx5->get_real(); + + GKO_ASSERT_MTX_NEAR(real_mtx, this->mtx5, 0); +} + + +TYPED_TEST(Dense, GetRealWithGivenResult) +{ + using T = typename TestFixture::value_type; + using RealMtx = typename TestFixture::RealMtx; + auto exec = this->mtx5->get_executor(); + // clang-format off + // {1.0, -1.0, -0.5}, + // {-2.0, 2.0, 4.5}, + // {2.1, 3.4, 1.2} + // clang-format on + + auto real_mtx = RealMtx::create(exec, this->mtx5->get_size()); + this->mtx5->get_real(real_mtx.get()); - GKO_ASSERT_MTX_NEAR(abs_mtx, this->mtx5, r::value); + GKO_ASSERT_MTX_NEAR(real_mtx, this->mtx5, 0); +} + + +TYPED_TEST(Dense, GetRealWithGivenResultFailsForWrongDimensions) +{ + using T = typename TestFixture::value_type; + using RealMtx = typename TestFixture::RealMtx; + auto exec = this->mtx5->get_executor(); + // clang-format off + // {1.0, -1.0, -0.5}, + // {-2.0, 2.0, 4.5}, + // {2.1, 3.4, 1.2} + // clang-format on + + auto real_mtx = RealMtx::create(exec); + ASSERT_THROW(this->mtx5->get_real(real_mtx.get()), gko::DimensionMismatch); } @@ -2190,11 +2262,45 @@ TYPED_TEST(Dense, GetImag) // {2.1, 3.4, 1.2} // clang-format on - auto abs_mtx = this->mtx5->get_imag(); + auto imag_mtx = this->mtx5->get_imag(); - GKO_ASSERT_MTX_NEAR(abs_mtx, - l({{0.0, 0.0, 0.0}, {0.0, 0.0, 0.0}, {0.0, 0.0, 0.0}}), - r::value); + GKO_ASSERT_MTX_NEAR( + imag_mtx, l({{0.0, 0.0, 0.0}, {0.0, 0.0, 0.0}, {0.0, 0.0, 0.0}}), 0); +} + + +TYPED_TEST(Dense, GetImagWithGivenResult) +{ + using T = typename TestFixture::value_type; + using RealMtx = typename TestFixture::RealMtx; + auto exec = this->mtx5->get_executor(); + // clang-format off + // {1.0, -1.0, -0.5}, + // {-2.0, 2.0, 4.5}, + // {2.1, 3.4, 1.2} + // clang-format on + + auto imag_mtx = RealMtx::create(exec, this->mtx5->get_size()); + this->mtx5->get_imag(imag_mtx.get()); + + GKO_ASSERT_MTX_NEAR( + imag_mtx, l({{0.0, 0.0, 0.0}, {0.0, 0.0, 0.0}, {0.0, 0.0, 0.0}}), 0); +} + + +TYPED_TEST(Dense, GetImagWithGivenResultFailsForWrongDimensions) +{ + using T = typename TestFixture::value_type; + using RealMtx = typename TestFixture::RealMtx; + auto exec = this->mtx5->get_executor(); + // clang-format off + // {1.0, -1.0, -0.5}, + // {-2.0, 2.0, 4.5}, + // {2.1, 3.4, 1.2} + // clang-format on + + auto imag_mtx = RealMtx::create(exec); + ASSERT_THROW(this->mtx5->get_imag(imag_mtx.get()), gko::DimensionMismatch); } @@ -2203,6 +2309,7 @@ class DenseComplex : public ::testing::Test { protected: using value_type = T; using Mtx = gko::matrix::Dense; + using RealMtx = gko::matrix::Dense>; }; @@ -2279,9 +2386,28 @@ TYPED_TEST(DenseComplex, MakeComplex) {T{0.0, 0.0}, T{0.0, -1.5}, T{2.0, 0.0}}}, exec); // clang-format on - auto abs_mtx = mtx->make_complex(); + auto complex_mtx = mtx->make_complex(); - GKO_ASSERT_MTX_NEAR(abs_mtx, mtx, 0.0); + GKO_ASSERT_MTX_NEAR(complex_mtx, mtx, 0.0); +} + + +TYPED_TEST(DenseComplex, MakeComplexWithGivenResult) +{ + using Mtx = typename TestFixture::Mtx; + using T = typename TestFixture::value_type; + auto exec = gko::ReferenceExecutor::create(); + // clang-format off + auto mtx = gko::initialize( + {{T{1.0, 0.0}, T{3.0, 4.0}, T{0.0, 2.0}}, + {T{-4.0, -3.0}, T{-1.0, 0}, T{0.0, 0.0}}, + {T{0.0, 0.0}, T{0.0, -1.5}, T{2.0, 0.0}}}, exec); + // clang-format on + + auto complex_mtx = Mtx::create(exec, mtx->get_size()); + mtx->make_complex(complex_mtx.get()); + + GKO_ASSERT_MTX_NEAR(complex_mtx, mtx, 0.0); } @@ -2297,10 +2423,33 @@ TYPED_TEST(DenseComplex, GetReal) {T{0.0, 0.0}, T{0.0, -1.5}, T{2.0, 0.0}}}, exec); // clang-format on - auto abs_mtx = mtx->get_real(); + auto real_mtx = mtx->get_real(); + + GKO_ASSERT_MTX_NEAR( + real_mtx, l({{1.0, 3.0, 0.0}, {-4.0, -1.0, 0.0}, {0.0, 0.0, 2.0}}), + 0.0); +} + + +TYPED_TEST(DenseComplex, GetRealWithGivenResult) +{ + using Mtx = typename TestFixture::Mtx; + using RealMtx = typename TestFixture::RealMtx; + using T = typename TestFixture::value_type; + auto exec = gko::ReferenceExecutor::create(); + // clang-format off + auto mtx = gko::initialize( + {{T{1.0, 0.0}, T{3.0, 4.0}, T{0.0, 2.0}}, + {T{-4.0, -3.0}, T{-1.0, 0}, T{0.0, 0.0}}, + {T{0.0, 0.0}, T{0.0, -1.5}, T{2.0, 0.0}}}, exec); + // clang-format on + + auto real_mtx = RealMtx::create(exec, mtx->get_size()); + mtx->get_real(real_mtx.get()); GKO_ASSERT_MTX_NEAR( - abs_mtx, l({{1.0, 3.0, 0.0}, {-4.0, -1.0, 0.0}, {0.0, 0.0, 2.0}}), 0.0); + real_mtx, l({{1.0, 3.0, 0.0}, {-4.0, -1.0, 0.0}, {0.0, 0.0, 2.0}}), + 0.0); } @@ -2316,10 +2465,33 @@ TYPED_TEST(DenseComplex, GetImag) {T{0.0, 0.0}, T{0.0, -1.5}, T{2.0, 0.0}}}, exec); // clang-format on - auto abs_mtx = mtx->get_imag(); + auto imag_mtx = mtx->get_imag(); GKO_ASSERT_MTX_NEAR( - abs_mtx, l({{0.0, 4.0, 2.0}, {-3.0, 0.0, 0.0}, {0.0, -1.5, 0.0}}), 0.0); + imag_mtx, l({{0.0, 4.0, 2.0}, {-3.0, 0.0, 0.0}, {0.0, -1.5, 0.0}}), + 0.0); +} + + +TYPED_TEST(DenseComplex, GetImagWithGivenResult) +{ + using Mtx = typename TestFixture::Mtx; + using RealMtx = typename TestFixture::RealMtx; + using T = typename TestFixture::value_type; + auto exec = gko::ReferenceExecutor::create(); + // clang-format off + auto mtx = gko::initialize( + {{T{1.0, 0.0}, T{3.0, 4.0}, T{0.0, 2.0}}, + {T{-4.0, -3.0}, T{-1.0, 0}, T{0.0, 0.0}}, + {T{0.0, 0.0}, T{0.0, -1.5}, T{2.0, 0.0}}}, exec); + // clang-format on + + auto imag_mtx = RealMtx::create(exec, mtx->get_size()); + mtx->get_imag(imag_mtx.get()); + + GKO_ASSERT_MTX_NEAR( + imag_mtx, l({{0.0, 4.0, 2.0}, {-3.0, 0.0, 0.0}, {0.0, -1.5, 0.0}}), + 0.0); } From a63cb7bd06f3acc638ecdf83da5d0d13c7790c66 Mon Sep 17 00:00:00 2001 From: Fritz Goebel Date: Thu, 29 Oct 2020 13:22:54 +0100 Subject: [PATCH 5/7] Use temporary clone in conversions with given result --- core/matrix/dense.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/core/matrix/dense.cpp b/core/matrix/dense.cpp index 2cb02fb4361..cb350c31628 100644 --- a/core/matrix/dense.cpp +++ b/core/matrix/dense.cpp @@ -819,7 +819,8 @@ void Dense::make_complex(Dense> *result) const GKO_ASSERT_EQUAL_DIMENSIONS(this, result); - exec->run(dense::make_make_complex(this, result)); + exec->run(dense::make_make_complex( + this, make_temporary_clone(exec, result).get())); } @@ -844,7 +845,8 @@ void Dense::get_real(Dense> *result) const GKO_ASSERT_EQUAL_DIMENSIONS(this, result); - exec->run(dense::make_get_real(this, result)); + exec->run( + dense::make_get_real(this, make_temporary_clone(exec, result).get())); } @@ -869,7 +871,8 @@ void Dense::get_imag(Dense> *result) const GKO_ASSERT_EQUAL_DIMENSIONS(this, result); - exec->run(dense::make_get_imag(this, result)); + exec->run( + dense::make_get_imag(this, make_temporary_clone(exec, result).get())); } From 339a065325e8a8b1b09205dad2c31b1f50374a49 Mon Sep 17 00:00:00 2001 From: Fritz Goebel Date: Thu, 5 Nov 2020 14:49:57 +0100 Subject: [PATCH 6/7] Add dpcpp GKO_NOT_IMPLEMENTED kernels, code review --- dpcpp/matrix/dense_kernels.dp.cpp | 27 +++++++++++++++++++++++++ reference/test/matrix/dense_kernels.cpp | 4 ++-- 2 files changed, 29 insertions(+), 2 deletions(-) diff --git a/dpcpp/matrix/dense_kernels.dp.cpp b/dpcpp/matrix/dense_kernels.dp.cpp index e84519e8192..3f2f823d799 100644 --- a/dpcpp/matrix/dense_kernels.dp.cpp +++ b/dpcpp/matrix/dense_kernels.dp.cpp @@ -308,6 +308,33 @@ void outplace_absolute_dense(std::shared_ptr exec, GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_OUTPLACE_ABSOLUTE_DENSE_KERNEL); +template +void make_complex(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Dense> *result) + GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_MAKE_COMPLEX_KERNEL); + + +template +void get_real(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Dense> *result) + GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_REAL_KERNEL); + + +template +void get_imag(std::shared_ptr exec, + const matrix::Dense *source, + matrix::Dense> *result) + GKO_NOT_IMPLEMENTED; + +GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_IMAG_KERNEL); + + } // namespace dense } // namespace dpcpp } // namespace kernels diff --git a/reference/test/matrix/dense_kernels.cpp b/reference/test/matrix/dense_kernels.cpp index 562176aefbe..d6618866de0 100644 --- a/reference/test/matrix/dense_kernels.cpp +++ b/reference/test/matrix/dense_kernels.cpp @@ -65,8 +65,8 @@ class Dense : public ::testing::Test { protected: using value_type = T; using Mtx = gko::matrix::Dense; - using ComplexMtx = gko::matrix::Dense>; - using RealMtx = gko::matrix::Dense>; + using ComplexMtx = gko::to_complex; + using RealMtx = gko::remove_complex; Dense() : exec(gko::ReferenceExecutor::create()), mtx1(gko::initialize(4, {{1.0, 2.0, 3.0}, {1.5, 2.5, 3.5}}, From c393e0b7a25eeeac3b578c69624ddb55ec349e56 Mon Sep 17 00:00:00 2001 From: Fritz Goebel Date: Thu, 5 Nov 2020 19:43:31 +0100 Subject: [PATCH 7/7] Fix missed dpcpp kernel --- dpcpp/matrix/dense_kernels.dp.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/dpcpp/matrix/dense_kernels.dp.cpp b/dpcpp/matrix/dense_kernels.dp.cpp index 3f2f823d799..fac85ba3a14 100644 --- a/dpcpp/matrix/dense_kernels.dp.cpp +++ b/dpcpp/matrix/dense_kernels.dp.cpp @@ -327,7 +327,7 @@ GKO_INSTANTIATE_FOR_EACH_VALUE_TYPE(GKO_DECLARE_GET_REAL_KERNEL); template -void get_imag(std::shared_ptr exec, +void get_imag(std::shared_ptr exec, const matrix::Dense *source, matrix::Dense> *result) GKO_NOT_IMPLEMENTED;