Skip to content

Add Compressed Basis GMRES (CB-GMRES) #693

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
thoasm merged 87 commits into from
Feb 21, 2021
Merged

Add Compressed Basis GMRES (CB-GMRES) #693

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
87 commits
Select commit Hold shift + click to select a range
cd1ffc9
First commit for the GmresMixed class:
Apr 27, 2020
d7ba04c
Inclusion of the omp executor in the repository:
Apr 27, 2020
969358f
Inclusion of the cuda executor in the repository:
Apr 29, 2020
2b7122f
The test files are finally included, but:
Apr 29, 2020
6dc6470
The first unoptimized version is done. Next tasks:
Apr 30, 2020
4d4fab1
The default value for ValueTypeKrylovBasis has been changed to avoid …
josealiaga Apr 30, 2020
e0d39a6
Definition of the CG2 variant of the finish_arnoldi routine for omp a…
josealiaga May 3, 2020
27019da
Add GMRES_mixed to benchmark
May 7, 2020
f2b5a3a
Definition of the CGS2 version of finish_arnoldi method, for omp and …
josealiaga May 7, 2020
bfd4196
Finally a good implementation of the multidot_kernels_num_iters_1 is …
josealiaga May 12, 2020
4ad491e
Add Accessor support and extend reference test
May 13, 2020
fbcc4f2
Made GmresMixed compile with complex types
May 13, 2020
911fd5c
The update routines have been improved. Now the computational time is…
josealiaga May 19, 2020
db081ac
Add specialization for integer types for Accessor
May 19, 2020
7db4796
Make the scale work with integer types
May 20, 2020
9bbcf87
Add helper to determine if we need a scale or not
May 20, 2020
97689ab
Add a helper structure to manage the scale writing in common
May 20, 2020
c0aa2ed
Testing the push command
josealiaga May 20, 2020
da8a56e
Definition of norm2 and norminf routines in CUDA. Only the first one …
josealiaga May 22, 2020
23efe7a
remove_complex has been added to the norms variables, and multinorm2_…
josealiaga May 24, 2020
f3cdbd7
Fixed cuda step2 to take a view into account
May 25, 2020
d4865f2
Change const accessor to non-const in check_arnoldi_norms_new
May 25, 2020
42a779b
The set_scale method finally works!!
josealiaga May 27, 2020
6b25d2e
Change storage layout of krylov_bases
Jun 1, 2020
e3ca7cc
Make memory access to krylov_bases coalesced again
Jun 2, 2020
911fccb
Transpose grid when launching singledot kernel
Jun 4, 2020
5dd8bf2
Reversed the transpose of the grid dim
Jun 4, 2020
b214a5f
Add half precision support to GmresMixed
Jun 15, 2020
f1f178e
Hopefully improve singledot performance
Jun 15, 2020
fb64f5c
Infinity norm only computed when scale is present
Jun 17, 2020
155c552
Add another RHS generation in the benchmark
Jun 23, 2020
d660c47
Fix residual_norm calculation in GmresMixed
Jun 23, 2020
8ede6d6
Make sure GmresMixed does not exit early
Jun 25, 2020
af6fb48
Add benchmark parameter for GMRES krylov_dim
Jun 26, 2020
49f1763
Add forced iterations when convergence is detected
Jun 26, 2020
188ff18
Add debug output to forced iterations
Jun 26, 2020
297004b
Fix reference bug in GmresMixed
Jul 8, 2020
b7765c3
DEBUG: Add write output for integral accessor
Jul 30, 2020
f31e87d
DEBUG: Move towards `at` with accessor
Aug 14, 2020
6f82f47
Remove Accessor3dConst
Aug 17, 2020
c385a2b
Adopt OpenMP support to new Accessor
Aug 17, 2020
9a0a0f0
Remove unused GMRES_mixed code from Ref & OMP
Aug 17, 2020
3d44a21
Adopt CUDA to the new accessor format (NOT `at`)
Aug 17, 2020
e1654b4
Make HIP and CUDA work with new accessor (NOT at)
Aug 17, 2020
5d09f2a
Remove unused code from CUDA
Aug 17, 2020
48e0899
CUDA implementation is now using `at`
Aug 18, 2020
b542646
Re-add ConstAccessor
Aug 18, 2020
b0a2ac3
Fix accessor by adding additional __restrict__
Aug 19, 2020
5d1f173
GmresMixed storage prec is now a factory parameter
Aug 25, 2020
267bbf1
Improve reference test and include the enum there
Aug 25, 2020
6b92a4f
Fix the reference test to pass
Aug 26, 2020
1d4a773
Adopt to new parameter macros
Sep 1, 2020
30381e2
Update the helper to throw when complex
Sep 7, 2020
6d30e36
Make GmresMixed work properly with multiple RHS
Sep 9, 2020
7f86d9c
Fix benchmark to work with new GmresMixed layout
Sep 10, 2020
b2b9ebc
Use new reduced_row_major Accessor in GmresMixed
Sep 11, 2020
db3046f
Remove unnecessary code from CUDA GmresMixed
Sep 14, 2020
da87d06
Add HIP kernels
Sep 14, 2020
65a8a8b
Fix GmresMixed core problem
Sep 17, 2020
a9646e8
Improve force-reset behavior
Dec 12, 2020
c6020db
Rename GmresMixed to CbGmres
Jan 22, 2021
36a1a19
Format files
ginkgo-bot Jan 23, 2021
c4a7335
Add DPCPP stubs to allow compilation
Jan 24, 2021
fa1fc39
Make cb-gmres benchmarks dependent on etype
Jan 28, 2021
c509263
Fix implementation and reference test for CB-GMRES
Feb 1, 2021
6bb09da
Update tolerance for one reference CB-GMRES test
Feb 1, 2021
3655021
Update atomic_max
Feb 2, 2021
5270478
Remove unnecessary kernels and properly name them
Feb 2, 2021
2f0a32e
Review update
Feb 2, 2021
58e1104
Add Helper INSTANTIATE macro for CB-GMRES
Feb 3, 2021
c47a4ca
Remove CB-GMRES and GMRES example
Feb 3, 2021
2b587ba
Review update
Feb 3, 2021
da8a6b4
Remove unnecessary includes of iostream and time.h
Feb 3, 2021
c4c1270
Remove circular dependency of compute_norm2 in (CB)-GMRES
Feb 3, 2021
00d33b5
Update solver generation in benchmark
Feb 3, 2021
b163893
Update eta and arnoldi_norms in CB-GMRES
Feb 4, 2021
342957a
Remove CUDA 9.0 exception for constexpr parameter
Feb 4, 2021
370d208
Review Update
Feb 5, 2021
1b2071d
Sonarcloud update
Feb 6, 2021
b4a6fc9
Review update; Improve run_all_benchmarks.sh
Feb 8, 2021
599e261
Put storage_precision enum into cb_gmres namespace
Feb 9, 2021
5126858
Add CB-GMRES example
Feb 9, 2021
169040d
Remove unnecessary included files for CB-GMRES
Feb 9, 2021
05374fd
Review update
Feb 15, 2021
389d038
Review update
Feb 16, 2021
6d6bbab
Update contributors.txt
josealiaga Feb 19, 2021
4d722f1
Update contributors.txt
josealiaga Feb 19, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
Inclusion of the cuda executor in the repository: 
    * For CGS, a loop of kernels is used instead of a kernel with a loop.
    * The test files are not repaired yet.
  • Loading branch information
aliaga@uji.es authored and Thomas Grützmacher committed Feb 19, 2021
commit 969358f726042e83302d7a94495e786c134e404b
2 changes: 2 additions & 0 deletions common/components/atomic.hpp.inc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -81,6 +81,7 @@ __forceinline__ __device__ ResultType reinterpret(ValueType val)

// Support 64-bit ATOMIC_ADD
Comment thread
thoasm marked this conversation as resolved.
Outdated
GKO_BIND_ATOMIC_HELPER_STRUCTURE(unsigned long long int);
// GKO_BIND_ATOMIC_HELPER_STRUCTURE(size_type);
Comment thread
thoasm marked this conversation as resolved.
Outdated
// Support 32-bit ATOMIC_ADD
Comment thread
thoasm marked this conversation as resolved.
Outdated
GKO_BIND_ATOMIC_HELPER_STRUCTURE(unsigned int);

Expand Down Expand Up @@ -115,6 +116,7 @@ GKO_BIND_ATOMIC_ADD(unsigned int);
GKO_BIND_ATOMIC_ADD(unsigned long long int);
GKO_BIND_ATOMIC_ADD(float);

// GKO_BIND_ATOMIC_ADD(size_type);
Comment thread
thoasm marked this conversation as resolved.
Outdated

#if !defined(__HIPCC__) || \
(defined(__HIP_DEVICE_COMPILE__) && GINKGO_HIP_PLATFORM_NVCC)
Expand Down
597 changes: 597 additions & 0 deletions common/solver/gmres_mixed_kernels.hpp.inc
View file
Open in desktop

Large diffs are not rendered by default.

6 changes: 5 additions & 1 deletion core/solver/gmres_mixed.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -117,6 +117,10 @@ void GmresMixed<ValueType, ValueTypeKrylovBases>::apply_impl(const LinOp *b,
bool one_changed{};
Array<stopping_status> stop_status(this->get_executor(),
dense_b->get_size()[1]);
Array<stopping_status> reorth_status(this->get_executor(),
dense_b->get_size()[1]);
// krylov_dim_mixed_ + 1);
Array<size_type> num_reorth(this->get_executor(), 1);

// Initialization
exec->run(gmres_mixed::make_initialize_1(
Expand Down Expand Up @@ -217,7 +221,7 @@ void GmresMixed<ValueType, ValueTypeKrylovBases>::apply_impl(const LinOp *b,
residual_norm.get(), residual_norm_collection.get(),
krylov_bases.get(), hessenberg_iter.get(), buffer_iter.get(),
b_norm.get(), arnoldi_norm.get(), restart_iter, &final_iter_nums,
&stop_status));
&stop_status, &reorth_status, &num_reorth));
// for i in 0:restart_iter
// hessenberg(restart_iter, i) = next_krylov_basis' *
// krylov_bases(:, i) next_krylov_basis -= hessenberg(restart_iter,
Expand Down
29 changes: 15 additions & 14 deletions core/solver/gmres_mixed_kernels.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -65,20 +65,21 @@ namespace gmres_mixed {


// #define GKO_DECLARE_GMRES_MIXED_STEP_1_KERNEL(_type1, _type2, _bool1, _bool2)
#define GKO_DECLARE_GMRES_MIXED_STEP_1_KERNEL(_type1, _type2) \
void step_1(std::shared_ptr<const DefaultExecutor> exec, \
matrix::Dense<_type1> *next_krylov_basis, \
matrix::Dense<_type1> *givens_sin, \
matrix::Dense<_type1> *givens_cos, \
matrix::Dense<_type1> *residual_norm, \
matrix::Dense<_type1> *residual_norm_collection, \
matrix::Dense<_type2> *krylov_bases, \
matrix::Dense<_type1> *hessenberg_iter, \
matrix::Dense<_type1> *buffer_iter, \
const matrix::Dense<_type1> *b_norm, \
matrix::Dense<_type1> *arnoldi_norm, size_type iter, \
Array<size_type> *final_iter_nums, \
const Array<stopping_status> *stop_status)
#define GKO_DECLARE_GMRES_MIXED_STEP_1_KERNEL(_type1, _type2) \
void step_1( \
std::shared_ptr<const DefaultExecutor> exec, \
matrix::Dense<_type1> *next_krylov_basis, \
matrix::Dense<_type1> *givens_sin, matrix::Dense<_type1> *givens_cos, \
matrix::Dense<_type1> *residual_norm, \
matrix::Dense<_type1> *residual_norm_collection, \
matrix::Dense<_type2> *krylov_bases, \
matrix::Dense<_type1> *hessenberg_iter, \
matrix::Dense<_type1> *buffer_iter, \
const matrix::Dense<_type1> *b_norm, \
matrix::Dense<_type1> *arnoldi_norm, size_type iter, \
Array<size_type> *final_iter_nums, \
const Array<stopping_status> *stop_status, \
Array<stopping_status> *reorth_status, Array<size_type> *num_reorth)


#define GKO_DECLARE_GMRES_MIXED_STEP_2_KERNEL(_type1, _type2) \
Expand Down
Loading