Reorder krylov bases of gmres #523

yhmtsai · 2020-04-30T13:43:46Z

This PR changes the order of storing krylov bases.

Originally, we store the Krylov bases as following:
K1 = the first vector of Krylov basis for each rhs => num_rows * num_rhs
K = [K1 K2 ... Kn], which lead the stride of Krylov still be larger than 1 even if the num_rhs = 1.
Thus, we can not get coalesced memory access.
Without changing the implementation of algorithm, change the Krylov order like
K = [K1; K2; ...; Kn] to get coalesced memory access when num_rhs = 1

When num_rhs = 1, use vendor dot not our own multidot_kernel which is parallel on #rhs.

Moreover, the stride of krylov_bases is same as stride of rhs.
Do not face the problem mentioned in #481, so next_krylov_basis is not needed now.

Todo:

change the algorithm comment

common/solver/gmres_kernels.hpp.inc

codecov · 2020-04-30T19:32:26Z

Codecov Report

Merging #523 into develop will increase coverage by 0.05%.
The diff coverage is 74.64%.

@@             Coverage Diff             @@
##           develop     #523      +/-   ##
===========================================
+ Coverage    88.47%   88.53%   +0.05%     
===========================================
  Files          276      276              
  Lines        17454    17400      -54     
===========================================
- Hits         15443    15405      -38     
+ Misses        2011     1995      -16

Impacted Files	Coverage Δ
omp/solver/gmres_kernels.cpp	`0.00% <0.00%> (ø)`
core/solver/gmres.cpp	`94.62% <100.00%> (+0.50%)`	⬆️
omp/test/solver/gmres_kernels.cpp	`100.00% <100.00%> (ø)`
reference/solver/gmres_kernels.cpp	`95.45% <100.00%> (-0.24%)`	⬇️
omp/components/format_conversion.hpp	`40.00% <0.00%> (-60.00%)`	⬇️
include/ginkgo/core/matrix/ell.hpp	`68.75% <0.00%> (-31.25%)`	⬇️
include/ginkgo/core/base/utils.hpp	`88.46% <0.00%> (-1.74%)`	⬇️
include/ginkgo/core/base/polymorphic_object.hpp	`95.23% <0.00%> (-1.59%)`	⬇️
core/solver/ir.cpp	`95.12% <0.00%> (-0.54%)`	⬇️
core/test/base/lin_op.cpp	`92.30% <0.00%> (-0.10%)`	⬇️
... and 20 more

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 bebd252...966568c. Read the comment docs.

common/solver/gmres_kernels.hpp.inc

upsj · 2020-05-04T12:46:17Z

I think the complex Givens rotations are still not 100% correct. What it looks like to me:
$\begin{pmatrix}c&s\-\overline{s}&\overline{c}\end{pmatrix}\begin{pmatrix}a\b\end{pmatrix}=\begin{pmatrix}r\0\end{pmatrix}$
We need this form for the matrix to be unitary. Ansatz:
$r=\sqrt{|a|^2 + |b|^2}, c=\overline{a} / r, s=\overline{b} / r$
Then the update looks like this:
$\begin{pmatrix}c&s\-\overline{s}&\overline{c}\end{pmatrix}\begin{pmatrix}a\b\end{pmatrix}=\begin{pmatrix}ca+sb\-\overline{s}a + \overline{c}b\end{pmatrix}$
So we need to also conjugate sin and cos for the second entry of H we are updating.

Please feel free to check my calculations!

yhmtsai · 2020-05-05T14:12:33Z

@upsj The formula looks correct.
the current given rotation will rotate to the a direction.
That is, in your first formula -> [r (unit_vector of a) 0]'

I prefer your formula because of the same behaviour of cos and sin.

yhmtsai · 2020-05-07T12:29:58Z

common/solver/gmres_kernels.hpp.inc

@@ -300,7 +299,7 @@ __device__ void calculate_residual_norm_kernel(
 const auto this_rnc =
 residual_norm_collection[iter * stride_residual_norm_collection +
 col_idx];
- const auto next_rnc = -register_sin * this_rnc;
+ const auto next_rnc = -conj(register_sin) * this_rnc;
 residual_norm_collection[iter * stride_residual_norm_collection + col_idx] =
 register_cos * this_rnc;


norm_collection(iter) = cos * this_rnc + sin * next_rnc = cos * this_rnc;
norm_collection(iter+1) = -conj(sin) * this_rnc + cos * next_rnc = -conj(sin) * this_rnc;
Is that correct?

pratikvn

LGTM!

tcojean

LGTM. Some minor comments.

cuda/solver/gmres_kernels.cu

hip/solver/gmres_kernels.hip.cpp

sonarcloud · 2020-05-16T00:57:55Z

Kudos, SonarCloud Quality Gate passed!

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

No Coverage information
No Duplication information

The Ginkgo team is proud to announce the new minor release of Ginkgo version 1.2.0. This release brings full HIP support to Ginkgo, new preconditioners (ParILUT, ISAI), conversion between double and float for all LinOps, and many more features and fixes. Supported systems and requirements: + For all platforms, cmake 3.9+ + Linux and MacOS + gcc: 5.3+, 6.3+, 7.3+, all versions after 8.1+ + clang: 3.9+ + Intel compiler: 2017+ + Apple LLVM: 8.0+ + CUDA module: CUDA 9.0+ + HIP module: ROCm 2.8+ + Windows + MinGW and CygWin: gcc 5.3+, 6.3+, 7.3+, all versions after 8.1+ + Microsoft Visual Studio: VS 2017 15.7+ + CUDA module: CUDA 9.0+, Microsoft Visual Studio + OpenMP module: MinGW or CygWin. The current known issues can be found in the [known issues page](https://github.com/ginkgo-project/ginkgo/wiki/Known-Issues). # Additions Here are the main additions to the Ginkgo library. Other thematic additions are listed below. + Add full HIP support to Ginkgo [#344](#344), [#357](#357), [#384](#384), [#373](#373), [#391](#391), [#396](#396), [#395](#395), [#393](#393), [#404](#404), [#439](#439), [#443](#443), [#567](#567) + Add a new ISAI preconditioner [#489](#489), [#502](#502), [#512](#512), [#508](#508), [#520](#520) + Add support for ParILUT and ParICT factorization with ILU preconditioners [#400](#400) + Add a new BiCG solver [#438](#438) + Add a new permutation matrix format [#352](#352), [#469](#469) + Add CSR SpGEMM support [#386](#386), [#398](#398), [#418](#418), [#457](#457) + Add CSR SpGEAM support [#556](#556) + Make all solvers and preconditioners transposable [#535](#535) + Add CsrBuilder and CooBuilder for intrusive access to matrix arrays [#437](#437) + Add a standard-compliant allocator based on the Executors [#504](#504) + Support conversions for all LinOp between double and float [#521](#521) + Add a new boolean to the CUDA and HIP executors to control DeviceReset (default off) [#557](#557) + Add a relaxation factor to IR to represent Richardson Relaxation [#574](#574) + Add two new stopping criteria, for relative (to `norm(b)`) and absolute residual norm [#577](#577) ### Example additions + Templatize all examples to simplify changing the precision [#513](#513) + Add a new adaptive precision block-Jacobi example [#507](#507) + Add a new IR example [#522](#522) + Add a new Mixed Precision Iterative Refinement example [#525](#525) + Add a new example on iterative trisolves in ILU preconditioning [#526](#526), [#536](#536), [#550](#550) ### Compilation and library changes + Auto-detect compilation settings based on environment [#435](#435), [#537](#537) + Add SONAME to shared libraries [#524](#524) + Add clang-cuda support [#543](#543) ### Other additions + Add sorting, searching and merging kernels for GPUs [#403](#403), [#428](#428), [#417](#417), [#455](#455) + Add `gko::as` support for smart pointers [#493](#493) + Add setters and getters for criterion factories [#527](#527) + Add a new method to check whether a solver uses `x` as an initial guess [#531](#531) + Add contribution guidelines [#549](#549) # Fixes ### Algorithms + Improve the classical CSR strategy's performance [#401](#401) + Improve the CSR automatical strategy [#407](#407), [#559](#559) + Memory, speed improvements to the ELL kernel [#411](#411) + Multiple improvements and fixes to ParILU [#419](#419), [#427](#427), [#429](#429), [#456](#456), [#544](#544) + Fix multiple issues with GMRES [#481](#481), [#523](#523), [#575](#575) + Optimize OpenMP matrix conversions [#505](#505) + Ensure the linearity of the ILU preconditioner [#506](#506) + Fix IR's use of the advanced apply [#522](#522) + Fix empty matrices conversions and add tests [#560](#560) ### Other core functionalities + Fix complex number support in our math header [#410](#410) + Fix CUDA compatibility of the main ginkgo header [#450](#450) + Fix isfinite issues [#465](#465) + Fix the Array::view memory leak and the array/view copy/move [#485](#485) + Fix typos preventing use of some interface functions [#496](#496) + Fix the `gko::dim` to abide to the C++ standard [#498](#498) + Simplify the executor copy interface [#516](#516) + Optimize intermediate storage for Composition [#540](#540) + Provide an initial guess for relevant Compositions [#561](#561) + Better management of nullptr as criterion [#562](#562) + Fix the norm calculations for complex support [#564](#564) ### CUDA and HIP specific + Use the return value of the atomic operations in our wrappers [#405](#405) + Improve the portability of warp lane masks [#422](#422) + Extract thread ID computation into a separate function [#464](#464) + Reorder kernel parameters for consistency [#474](#474) + Fix the use of `pragma unroll` in HIP [#492](#492) ### Other + Fix the Ginkgo CMake installation files [#414](#414), [#553](#553) + Fix the Windows compilation [#415](#415) + Always use demangled types in error messages [#434](#434), [#486](#486) + Add CUDA header dependency to appropriate tests [#452](#452) + Fix several sonarqube or compilation warnings [#453](#453), [#463](#463), [#532](#532), [#569](#569) + Add shuffle tests [#460](#460) + Fix MSVC C2398 error [#490](#490) + Fix missing interface tests in test install [#558](#558) # Tools and ecosystem ### Benchmarks + Add better norm support in the benchmarks [#377](#377) + Add CUDA 10.1 generic SpMV support in benchmarks [#468](#468), [#473](#473) + Add sparse library ILU in benchmarks [#487](#487) + Add overhead benchmarking capacities [#501](#501) + Allow benchmarking from a matrix list file [#503](#503) + Fix benchmarking issue with JSON and non-finite numbers [#514](#514) + Fix benchmark logger crashers with OpenMP [#565](#565) ### CI related + Improvements to the CI setup with HIP compilation [#421](#421), [#466](#466) + Add MacOSX CI support [#470](#470), [#488](#488) + Add Windows CI support [#471](#471), [#488](#488), [#510](#510), [#566](#566) + Use sanitizers instead of valgrind [#476](#476) + Add automatic container generation and update facilities [#499](#499) + Fix the CI parallelism settings [#517](#517), [#538](#538), [#539](#539) + Make the codecov patch check informational [#519](#519) + Add support for LLVM sanitizers with improved thread sanitizer support [#578](#578) ### Test suite + Add an assertion for sparsity pattern equality [#416](#416) + Add core and reference multiprecision tests support [#448](#448) + Speed up GPU tests by avoiding device reset [#467](#467) + Change test matrix location string [#494](#494) ### Other + Add Ginkgo badges from our tools [#413](#413) + Update the `create_new_algorithm.sh` script [#420](#420) + Bump copyright and improve license management [#436](#436), [#433](#433) + Set clang-format minimum requirement [#441](#441), [#484](#484) + Update git-cmake-format [#446](#446), [#484](#484) + Disable the development tools by default [#442](#442) + Add a script for automatic header formatting [#447](#447) + Add GDB pretty printer for `gko::Array` [#509](#509) + Improve compilation speed [#533](#533) + Add editorconfig support [#546](#546) + Add a compile-time check for header self-sufficiency [#552](#552) # Related PR: #583

yhmtsai requested review from upsj, pratikvn, thoasm, tcojean and fritzgoebel April 30, 2020 13:43

yhmtsai self-assigned this Apr 30, 2020

pratikvn reviewed Apr 30, 2020

View reviewed changes

common/solver/gmres_kernels.hpp.inc Outdated Show resolved Hide resolved

upsj approved these changes May 1, 2020

View reviewed changes

common/solver/gmres_kernels.hpp.inc Outdated Show resolved Hide resolved

common/solver/gmres_kernels.hpp.inc Outdated Show resolved Hide resolved

hartwiganzt approved these changes May 5, 2020

View reviewed changes

yhmtsai commented May 7, 2020

View reviewed changes

yhmtsai force-pushed the gmres_reorder_krylov branch from c6ad309 to a19aa91 Compare May 7, 2020 13:01

thoasm added the 1:ST:ready-for-review This PR is ready for review label May 8, 2020

pratikvn approved these changes May 11, 2020

View reviewed changes

tcojean approved these changes May 12, 2020

View reviewed changes

cuda/solver/gmres_kernels.cu Show resolved Hide resolved

cuda/solver/gmres_kernels.cu Outdated Show resolved Hide resolved

hip/solver/gmres_kernels.hip.cpp Show resolved Hide resolved

hip/solver/gmres_kernels.hip.cpp Outdated Show resolved Hide resolved

tcojean added this to the Ginkgo 1.2.0 milestone May 14, 2020

yhmtsai force-pushed the gmres_reorder_krylov branch from a19aa91 to 0b78b4d Compare May 14, 2020 08:52

yhmtsai added 5 commits May 15, 2020 10:36

reorder krylov_bases storage

3263eaa

try dot on one rhs

60b9125

make sure using register

2cdf457

remove the usage of next_krylov

c11fb19

move zero_array and use register

b96ffb9

yhmtsai added 5 commits May 15, 2020 10:36

use register

d1de00b

improve hypot stability, fix complex gmres issue

587bdee

remove one unnecessary hessenberg write

6cae85a

fix complex and related algorithm comment

33abb09

review update & add single rhs test

966568c

yhmtsai force-pushed the gmres_reorder_krylov branch from 0b78b4d to 966568c Compare May 15, 2020 08:37

yhmtsai added 1:ST:ready-to-merge This PR is ready to merge. and removed 1:ST:ready-for-review This PR is ready for review labels May 15, 2020

yhmtsai merged commit 2fb46e1 into develop May 15, 2020

yhmtsai deleted the gmres_reorder_krylov branch May 15, 2020 19:13

tcojean mentioned this pull request Jun 23, 2020

Release/1.2.0 #576

Merged

yhmtsai mentioned this pull request May 17, 2022

Simplify GMRES kernels #861

Merged

3 tasks

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Reorder krylov bases of gmres #523

Reorder krylov bases of gmres #523

yhmtsai commented Apr 30, 2020 •

edited

Loading

codecov bot commented Apr 30, 2020 •

edited

Loading

upsj commented May 4, 2020

yhmtsai commented May 5, 2020

yhmtsai May 7, 2020

pratikvn left a comment

tcojean left a comment

sonarcloud bot commented May 16, 2020

Reorder krylov bases of gmres #523

Reorder krylov bases of gmres #523

Conversation

yhmtsai commented Apr 30, 2020 • edited Loading

codecov bot commented Apr 30, 2020 • edited Loading

Codecov Report

upsj commented May 4, 2020

yhmtsai commented May 5, 2020

yhmtsai May 7, 2020

Choose a reason for hiding this comment

pratikvn left a comment

Choose a reason for hiding this comment

tcojean left a comment

Choose a reason for hiding this comment

sonarcloud bot commented May 16, 2020

yhmtsai commented Apr 30, 2020 •

edited

Loading

codecov bot commented Apr 30, 2020 •

edited

Loading