Dense: conversion to complex and extraction of real and imaginary part #658

fritzgoebel · 2020-10-26T10:49:00Z

This PR enables dense to:

convert real to complex: add all zero imaginary part
complex to complex: return a copy
extract the real part: if real, return a copy, if complex, return the real part
extract the imaginary part: if real, return an all zero matrix, if complex return the imaginary part.

upsj

LGTM!

cuda/test/matrix/dense_kernels.cpp

yhmtsai

I forget the downside point if Convertible<to_complex> and Convertible<to_real>.
Does it destroy the default convert_to?
and maybe possible another operation is make_imag_complex for real->imag and complex->imag(complex type)

cuda/test/matrix/dense_kernels.cpp

include/ginkgo/core/matrix/dense.hpp

upsj · 2020-10-27T12:07:10Z

@yhmtsai The issue here is that the conversion is not symmetric: double -> complex<double> is unambiguous, but complex<double> -> double has two candidates (real and imag), so we can't just say complex->convert_to(double), which is the only interface ConvertibleTo provides. We also can't easily have just a one-sided conversion double -> complex, since this would require us to conditionally inherit from ConvertibleTo and conditionally override convert_to based on whether ValueType is complex.

codecov · 2020-10-27T20:16:18Z

Codecov Report

Merging #658 into develop will increase coverage by 0.01%.
The diff coverage is 95.89%.

@@             Coverage Diff             @@
##           develop     #658      +/-   ##
===========================================
+ Coverage    92.81%   92.82%   +0.01%     
===========================================
  Files          322      322              
  Lines        22904    23050     +146     
===========================================
+ Hits         21259    21397     +138     
- Misses        1645     1653       +8

Impacted Files	Coverage Δ
core/device_hooks/common_kernels.inc.cpp	`0.00% <0.00%> (ø)`
include/ginkgo/core/matrix/dense.hpp	`98.13% <ø> (ø)`
core/matrix/dense.cpp	`99.34% <100.00%> (+0.07%)`	⬆️
omp/matrix/dense_kernels.cpp	`97.56% <100.00%> (+0.14%)`	⬆️
omp/test/matrix/dense_kernels.cpp	`99.72% <100.00%> (+0.02%)`	⬆️
reference/matrix/dense_kernels.cpp	`100.00% <100.00%> (ø)`
reference/test/matrix/dense_kernels.cpp	`100.00% <100.00%> (ø)`
core/base/extended_float.hpp	`91.26% <0.00%> (-0.98%)`	⬇️
omp/reorder/rcm_kernels.cpp	`94.44% <0.00%> (-0.62%)`	⬇️

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update 99c2004...c393e0b. Read the comment docs.

core/matrix/dense.cpp

pratikvn

LGTM!

yhmtsai

LGTM in general, some minor things are in the comments.
After rebasing, you will need to add these functionalities GKO_NOT_IMPLEMENTED in dpcpp kernel.

yhmtsai · 2020-11-05T13:01:36Z

common/matrix/dense_kernels.hpp.inc

+template <typename ValueType, typename ComplexType>
+__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, ComplexType *__restrict__ out, size_type stride_out)


If this ComplexType is only for complex, I prefer use one template and use to_complex to represent complex.
Before Ginkgo 2.0.0, I think we need to discuss the precision(ValueType) conversion...

This doesn't work, since to_complex returns the CPU complex type, not thrust::complex. An alternative would be thrust::complex<remove_complex<ValueType>>

you are right. maybe does another way cuda_type<to_complex<ValueType>> work?

reference/test/matrix/dense_kernels.cpp

…o dense

upsj · 2020-11-05T17:00:59Z

dpcpp/matrix/dense_kernels.dp.cpp

+
+
+template <typename ValueType>
+void get_imag(std::shared_ptr<const CudaExecutor> exec,


Suggested change

void get_imag(std::shared_ptr<const CudaExecutor> exec,

void get_imag(std::shared_ptr<const DpcppExecutor> exec,

sonarqubecloud · 2020-11-06T03:51:35Z

Kudos, SonarCloud Quality Gate passed!

0 Bugs
0 Vulnerabilities (and 0 Security Hotspots to review)
6 Code Smells

70.4% Coverage
0.0% Duplication

The version of Java (1.8.0_121) you have used to run this analysis is deprecated and we will stop accepting it accepting it soon.Please update to at least Java 11.
Read more here

Ginkgo release 1.4.0 The Ginkgo team is proud to announce the new Ginkgo minor release 1.4.0. This release brings most of the Ginkgo functionality to the Intel DPC++ ecosystem which enables Intel-GPU and CPU execution. The only Ginkgo features which have not been ported yet are some preconditioners. Ginkgo's mixed-precision support is greatly enhanced thanks to: 1. The new Accessor concept, which allows writing kernels featuring on-the-fly memory compression, among other features. The accessor can be used as header-only, see the [accessor BLAS benchmarks repository](https://github.com/ginkgo-project/accessor-BLAS/tree/develop) as a usage example. 2. All LinOps now transparently support mixed-precision execution. By default, this is done through a temporary copy which may have a performance impact but already allows mixed-precision research. Native mixed-precision ELL kernels are implemented which do not see this cost. The accessor is also leveraged in a new CB-GMRES solver which allows for performance improvements by compressing the Krylov basis vectors. Many other features have been added to Ginkgo, such as reordering support, a new IDR solver, Incomplete Cholesky preconditioner, matrix assembly support (only CPU for now), machine topology information, and more! Supported systems and requirements: + For all platforms, cmake 3.13+ + C++14 compliant compiler + Linux and MacOS + gcc: 5.3+, 6.3+, 7.3+, all versions after 8.1+ + clang: 3.9+ + Intel compiler: 2018+ + Apple LLVM: 8.0+ + CUDA module: CUDA 9.0+ + HIP module: ROCm 3.5+ + DPC++ module: Intel OneAPI 2021.3. Set the CXX compiler to `dpcpp`. + Windows + MinGW and Cygwin: gcc 5.3+, 6.3+, 7.3+, all versions after 8.1+ + Microsoft Visual Studio: VS 2019 + CUDA module: CUDA 9.0+, Microsoft Visual Studio + OpenMP module: MinGW or Cygwin. Algorithm and important feature additions: + Add a new DPC++ Executor for SYCL execution and other base utilities [#648](#648), [#661](#661), [#757](#757), [#832](#832) + Port matrix formats, solvers and related kernels to DPC++. For some kernels, also make use of a shared kernel implementation for all executors (except Reference). [#710](#710), [#799](#799), [#779](#779), [#733](#733), [#844](#844), [#843](#843), [#789](#789), [#845](#845), [#849](#849), [#855](#855), [#856](#856) + Add accessors which allow multi-precision kernels, among other things. [#643](#643), [#708](#708) + Add support for mixed precision operations through apply in all LinOps. [#677](#677) + Add incomplete Cholesky factorizations and preconditioners as well as some improvements to ILU. [#672](#672), [#837](#837), [#846](#846) + Add an AMGX implementation and kernels on all devices but DPC++. [#528](#528), [#695](#695), [#860](#860) + Add a new mixed-precision capability solver, Compressed Basis GMRES (CB-GMRES). [#693](#693), [#763](#763) + Add the IDR(s) solver. [#620](#620) + Add a new fixed-size block CSR matrix format (for the Reference executor). [#671](#671), [#730](#730) + Add native mixed-precision support to the ELL format. [#717](#717), [#780](#780) + Add Reverse Cuthill-McKee reordering [#500](#500), [#649](#649) + Add matrix assembly support on CPUs. [#644](#644) + Extends ISAI from triangular to general and spd matrices. [#690](#690) Other additions: + Add the possibility to apply real matrices to complex vectors. [#655](#655), [#658](#658) + Add functions to compute the absolute of a matrix format. [#636](#636) + Add symmetric permutation and improve existing permutations. [#684](#684), [#657](#657), [#663](#663) + Add a MachineTopology class with HWLOC support [#554](#554), [#697](#697) + Add an implicit residual norm criterion. [#702](#702), [#818](#818), [#850](#850) + Row-major accessor is generalized to more than 2 dimensions and a new "block column-major" accessor has been added. [#707](#707) + Add an heat equation example. [#698](#698), [#706](#706) + Add ccache support in CMake and CI. [#725](#725), [#739](#739) + Allow tuning and benchmarking variables non intrusively. [#692](#692) + Add triangular solver benchmark [#664](#664) + Add benchmarks for BLAS operations [#772](#772), [#829](#829) + Add support for different precisions and consistent index types in benchmarks. [#675](#675), [#828](#828) + Add a Github bot system to facilitate development and PR management. [#667](#667), [#674](#674), [#689](#689), [#853](#853) + Add Intel (DPC++) CI support and enable CI on HPC systems. [#736](#736), [#751](#751), [#781](#781) + Add ssh debugging for Github Actions CI. [#749](#749) + Add pipeline segmentation for better CI speed. [#737](#737) Changes: + Add a Scalar Jacobi specialization and kernels. [#808](#808), [#834](#834), [#854](#854) + Add implicit residual log for solvers and benchmarks. [#714](#714) + Change handling of the conjugate in the dense dot product. [#755](#755) + Improved Dense stride handling. [#774](#774) + Multiple improvements to the OpenMP kernels performance, including COO, an exclusive prefix sum, and more. [#703](#703), [#765](#765), [#740](#740) + Allow specialization of submatrix and other dense creation functions in solvers. [#718](#718) + Improved Identity constructor and treatment of rectangular matrices. [#646](#646) + Allow CUDA/HIP executors to select allocation mode. [#758](#758) + Check if executors share the same memory. [#670](#670) + Improve test install and smoke testing support. [#721](#721) + Update the JOSS paper citation and add publications in the documentation. [#629](#629), [#724](#724) + Improve the version output. [#806](#806) + Add some utilities for dim and span. [#821](#821) + Improved solver and preconditioner benchmarks. [#660](#660) + Improve benchmark timing and output. [#669](#669), [#791](#791), [#801](#801), [#812](#812) Fixes: + Sorting fix for the Jacobi preconditioner. [#659](#659) + Also log the first residual norm in CGS [#735](#735) + Fix BiCG and HIP CSR to work with complex matrices. [#651](#651) + Fix Coo SpMV on strided vectors. [#807](#807) + Fix segfault of extract_diagonal, add short-and-fat test. [#769](#769) + Fix device_reset issue by moving counter/mutex to device. [#810](#810) + Fix `EnableLogging` superclass. [#841](#841) + Support ROCm 4.1.x and breaking HIP_PLATFORM changes. [#726](#726) + Decreased test size for a few device tests. [#742](#742) + Fix multiple issues with our CMake HIP and RPATH setup. [#712](#712), [#745](#745), [#709](#709) + Cleanup our CMake installation step. [#713](#713) + Various simplification and fixes to the Windows CMake setup. [#720](#720), [#785](#785) + Simplify third-party integration. [#786](#786) + Improve Ginkgo device arch flags management. [#696](#696) + Other fixes and improvements to the CMake setup. [#685](#685), [#792](#792), [#705](#705), [#836](#836) + Clarification of dense norm documentation [#784](#784) + Various development tools fixes and improvements [#738](#738), [#830](#830), [#840](#840) + Make multiple operators/constructors explicit. [#650](#650), [#761](#761) + Fix some issues, memory leaks and warnings found by MSVC. [#666](#666), [#731](#731) + Improved solver memory estimates and consistent iteration counts [#691](#691) + Various logger improvements and fixes [#728](#728), [#743](#743), [#754](#754) + Fix for ForwardIterator requirements in iterator_factory. [#665](#665) + Various benchmark fixes. [#647](#647), [#673](#673), [#722](#722) + Various CI fixes and improvements. [#642](#642), [#641](#641), [#795](#795), [#783](#783), [#793](#793), [#852](#852) Related PR: #857

Release 1.4.0 to master The Ginkgo team is proud to announce the new Ginkgo minor release 1.4.0. This release brings most of the Ginkgo functionality to the Intel DPC++ ecosystem which enables Intel-GPU and CPU execution. The only Ginkgo features which have not been ported yet are some preconditioners. Ginkgo's mixed-precision support is greatly enhanced thanks to: 1. The new Accessor concept, which allows writing kernels featuring on-the-fly memory compression, among other features. The accessor can be used as header-only, see the [accessor BLAS benchmarks repository](https://github.com/ginkgo-project/accessor-BLAS/tree/develop) as a usage example. 2. All LinOps now transparently support mixed-precision execution. By default, this is done through a temporary copy which may have a performance impact but already allows mixed-precision research. Native mixed-precision ELL kernels are implemented which do not see this cost. The accessor is also leveraged in a new CB-GMRES solver which allows for performance improvements by compressing the Krylov basis vectors. Many other features have been added to Ginkgo, such as reordering support, a new IDR solver, Incomplete Cholesky preconditioner, matrix assembly support (only CPU for now), machine topology information, and more! Supported systems and requirements: + For all platforms, cmake 3.13+ + C++14 compliant compiler + Linux and MacOS + gcc: 5.3+, 6.3+, 7.3+, all versions after 8.1+ + clang: 3.9+ + Intel compiler: 2018+ + Apple LLVM: 8.0+ + CUDA module: CUDA 9.0+ + HIP module: ROCm 3.5+ + DPC++ module: Intel OneAPI 2021.3. Set the CXX compiler to `dpcpp`. + Windows + MinGW and Cygwin: gcc 5.3+, 6.3+, 7.3+, all versions after 8.1+ + Microsoft Visual Studio: VS 2019 + CUDA module: CUDA 9.0+, Microsoft Visual Studio + OpenMP module: MinGW or Cygwin. Algorithm and important feature additions: + Add a new DPC++ Executor for SYCL execution and other base utilities [#648](#648), [#661](#661), [#757](#757), [#832](#832) + Port matrix formats, solvers and related kernels to DPC++. For some kernels, also make use of a shared kernel implementation for all executors (except Reference). [#710](#710), [#799](#799), [#779](#779), [#733](#733), [#844](#844), [#843](#843), [#789](#789), [#845](#845), [#849](#849), [#855](#855), [#856](#856) + Add accessors which allow multi-precision kernels, among other things. [#643](#643), [#708](#708) + Add support for mixed precision operations through apply in all LinOps. [#677](#677) + Add incomplete Cholesky factorizations and preconditioners as well as some improvements to ILU. [#672](#672), [#837](#837), [#846](#846) + Add an AMGX implementation and kernels on all devices but DPC++. [#528](#528), [#695](#695), [#860](#860) + Add a new mixed-precision capability solver, Compressed Basis GMRES (CB-GMRES). [#693](#693), [#763](#763) + Add the IDR(s) solver. [#620](#620) + Add a new fixed-size block CSR matrix format (for the Reference executor). [#671](#671), [#730](#730) + Add native mixed-precision support to the ELL format. [#717](#717), [#780](#780) + Add Reverse Cuthill-McKee reordering [#500](#500), [#649](#649) + Add matrix assembly support on CPUs. [#644](#644) + Extends ISAI from triangular to general and spd matrices. [#690](#690) Other additions: + Add the possibility to apply real matrices to complex vectors. [#655](#655), [#658](#658) + Add functions to compute the absolute of a matrix format. [#636](#636) + Add symmetric permutation and improve existing permutations. [#684](#684), [#657](#657), [#663](#663) + Add a MachineTopology class with HWLOC support [#554](#554), [#697](#697) + Add an implicit residual norm criterion. [#702](#702), [#818](#818), [#850](#850) + Row-major accessor is generalized to more than 2 dimensions and a new "block column-major" accessor has been added. [#707](#707) + Add an heat equation example. [#698](#698), [#706](#706) + Add ccache support in CMake and CI. [#725](#725), [#739](#739) + Allow tuning and benchmarking variables non intrusively. [#692](#692) + Add triangular solver benchmark [#664](#664) + Add benchmarks for BLAS operations [#772](#772), [#829](#829) + Add support for different precisions and consistent index types in benchmarks. [#675](#675), [#828](#828) + Add a Github bot system to facilitate development and PR management. [#667](#667), [#674](#674), [#689](#689), [#853](#853) + Add Intel (DPC++) CI support and enable CI on HPC systems. [#736](#736), [#751](#751), [#781](#781) + Add ssh debugging for Github Actions CI. [#749](#749) + Add pipeline segmentation for better CI speed. [#737](#737) Changes: + Add a Scalar Jacobi specialization and kernels. [#808](#808), [#834](#834), [#854](#854) + Add implicit residual log for solvers and benchmarks. [#714](#714) + Change handling of the conjugate in the dense dot product. [#755](#755) + Improved Dense stride handling. [#774](#774) + Multiple improvements to the OpenMP kernels performance, including COO, an exclusive prefix sum, and more. [#703](#703), [#765](#765), [#740](#740) + Allow specialization of submatrix and other dense creation functions in solvers. [#718](#718) + Improved Identity constructor and treatment of rectangular matrices. [#646](#646) + Allow CUDA/HIP executors to select allocation mode. [#758](#758) + Check if executors share the same memory. [#670](#670) + Improve test install and smoke testing support. [#721](#721) + Update the JOSS paper citation and add publications in the documentation. [#629](#629), [#724](#724) + Improve the version output. [#806](#806) + Add some utilities for dim and span. [#821](#821) + Improved solver and preconditioner benchmarks. [#660](#660) + Improve benchmark timing and output. [#669](#669), [#791](#791), [#801](#801), [#812](#812) Fixes: + Sorting fix for the Jacobi preconditioner. [#659](#659) + Also log the first residual norm in CGS [#735](#735) + Fix BiCG and HIP CSR to work with complex matrices. [#651](#651) + Fix Coo SpMV on strided vectors. [#807](#807) + Fix segfault of extract_diagonal, add short-and-fat test. [#769](#769) + Fix device_reset issue by moving counter/mutex to device. [#810](#810) + Fix `EnableLogging` superclass. [#841](#841) + Support ROCm 4.1.x and breaking HIP_PLATFORM changes. [#726](#726) + Decreased test size for a few device tests. [#742](#742) + Fix multiple issues with our CMake HIP and RPATH setup. [#712](#712), [#745](#745), [#709](#709) + Cleanup our CMake installation step. [#713](#713) + Various simplification and fixes to the Windows CMake setup. [#720](#720), [#785](#785) + Simplify third-party integration. [#786](#786) + Improve Ginkgo device arch flags management. [#696](#696) + Other fixes and improvements to the CMake setup. [#685](#685), [#792](#792), [#705](#705), [#836](#836) + Clarification of dense norm documentation [#784](#784) + Various development tools fixes and improvements [#738](#738), [#830](#830), [#840](#840) + Make multiple operators/constructors explicit. [#650](#650), [#761](#761) + Fix some issues, memory leaks and warnings found by MSVC. [#666](#666), [#731](#731) + Improved solver memory estimates and consistent iteration counts [#691](#691) + Various logger improvements and fixes [#728](#728), [#743](#743), [#754](#754) + Fix for ForwardIterator requirements in iterator_factory. [#665](#665) + Various benchmark fixes. [#647](#647), [#673](#673), [#722](#722) + Various CI fixes and improvements. [#642](#642), [#641](#641), [#795](#795), [#783](#783), [#793](#793), [#852](#852) Related PR: #866

fritzgoebel added is:new-feature A request or implementation of a feature that does not exist yet. type:matrix-format This is related to the Matrix formats 1:ST:ready-for-review This PR is ready for review mod:all This touches all Ginkgo modules. labels Oct 26, 2020

fritzgoebel requested review from upsj, pratikvn, thoasm, yhmtsai, hartwiganzt and tcojean October 26, 2020 10:49

fritzgoebel self-assigned this Oct 26, 2020

fritzgoebel mentioned this pull request Oct 26, 2020

Idr solver #620

Merged

5 tasks

upsj approved these changes Oct 26, 2020

View reviewed changes

cuda/test/matrix/dense_kernels.cpp Outdated Show resolved Hide resolved

yhmtsai requested changes Oct 27, 2020

View reviewed changes

cuda/test/matrix/dense_kernels.cpp Outdated Show resolved Hide resolved

cuda/test/matrix/dense_kernels.cpp Outdated Show resolved Hide resolved

include/ginkgo/core/matrix/dense.hpp Outdated Show resolved Hide resolved

fritzgoebel force-pushed the dense_complex_real_conversions branch from 0c996eb to ba854c4 Compare October 27, 2020 14:35

upsj requested changes Oct 28, 2020

View reviewed changes

core/matrix/dense.cpp Show resolved Hide resolved

upsj approved these changes Nov 3, 2020

View reviewed changes

pratikvn approved these changes Nov 4, 2020

View reviewed changes

fritzgoebel force-pushed the dense_complex_real_conversions branch from 5ecb5eb to a80b222 Compare November 4, 2020 11:53

yhmtsai approved these changes Nov 5, 2020

View reviewed changes

fritzgoebel added 6 commits November 5, 2020 14:27

Add conversion to complex and extraction of real and imaginary part t…

974b176

…o dense

Fix wrong parameter types in gpu kernels

39b39e3

Apply suggestions from code review

a07a28d

Add conversions into given matrix

db511c5

Use temporary clone in conversions with given result

a63cb7b

Add dpcpp GKO_NOT_IMPLEMENTED kernels, code review

339a065

fritzgoebel force-pushed the dense_complex_real_conversions branch from a80b222 to 339a065 Compare November 5, 2020 13:50

upsj removed the 1:ST:ready-for-review This PR is ready for review label Nov 5, 2020

upsj added the 1:ST:ready-to-merge This PR is ready to merge. label Nov 5, 2020

upsj reviewed Nov 5, 2020

View reviewed changes

Fix missed dpcpp kernel

c393e0b

fritzgoebel merged commit 72902a7 into develop Nov 6, 2020

fritzgoebel deleted the dense_complex_real_conversions branch November 6, 2020 09:09

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Dense: conversion to complex and extraction of real and imaginary part #658

Dense: conversion to complex and extraction of real and imaginary part #658

fritzgoebel commented Oct 26, 2020

upsj left a comment

yhmtsai left a comment •

edited

Loading

upsj commented Oct 27, 2020

codecov bot commented Oct 27, 2020 •

edited

Loading

pratikvn left a comment

yhmtsai left a comment

yhmtsai Nov 5, 2020

upsj Nov 5, 2020

yhmtsai Nov 5, 2020

upsj Nov 5, 2020

sonarqubecloud bot commented Nov 6, 2020



		template <typename ValueType>
		void get_imag(std::shared_ptr<const CudaExecutor> exec,

	void get_imag(std::shared_ptr<const CudaExecutor> exec,
	void get_imag(std::shared_ptr<const DpcppExecutor> exec,

Dense: conversion to complex and extraction of real and imaginary part #658

Dense: conversion to complex and extraction of real and imaginary part #658

Conversation

fritzgoebel commented Oct 26, 2020

upsj left a comment

Choose a reason for hiding this comment

yhmtsai left a comment • edited Loading

Choose a reason for hiding this comment

upsj commented Oct 27, 2020

codecov bot commented Oct 27, 2020 • edited Loading

Codecov Report

pratikvn left a comment

Choose a reason for hiding this comment

yhmtsai left a comment

Choose a reason for hiding this comment

yhmtsai Nov 5, 2020

Choose a reason for hiding this comment

upsj Nov 5, 2020

Choose a reason for hiding this comment

yhmtsai Nov 5, 2020

Choose a reason for hiding this comment

upsj Nov 5, 2020

Choose a reason for hiding this comment

sonarqubecloud bot commented Nov 6, 2020

yhmtsai left a comment •

edited

Loading

codecov bot commented Oct 27, 2020 •

edited

Loading