8000 Add support for arbitrary row-lengths in GPU ISAI by upsj · Pull Request #520 · ginkgo-project/ginkgo · GitHub
[go: up one dir, main page]
More Web Proxy on the site http://driver.im/
Skip to content

Add support for arbitrary row-lengths in GPU ISAI #520

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.

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
merged 14 commits into from
May 20, 2020
Merged

Conversation

upsj
Copy link
Member
@upsj upsj commented Apr 29, 2020

This PR optimizes the GPU ISAI in multiple ways:

  • remove unnecessary identity_triangle kernel, since we only need the sparsity pattern of L^k, the values will be overwritten anyways
  • rewrite GPU ISAI kernels so they don't need to transpose the local sparse system
  • build a separate sparse triangular systems for rows longer than warp_size
  • uses a square-and-multiply algorithm instead of the naive approach for ISAI(k)

Also contains a few additional changes:

  • Fixes an off-by-one error in the gdb pretty-printer
  • Ensures that CSR row_ptrs are never empty
  • Renames isfinite to is_finite (I would not consider this interface-breaking, since it is neither documented nor actually our code, but just a wrapper over the equally-named std::isfinite)

Closes #511

@upsj upsj added mod:cuda This is related to the CUDA module. 1:ST:WIP This PR is a work in progress. Not ready for review. mod:hip This is related to the HIP module. labels Apr 29, 2020
@upsj upsj self-assigned this Apr 29, 2020
@upsj upsj added this to the Ginkgo 1.2.0 milestone Apr 30, 2020
@upsj upsj force-pushed the arbitrary_gpu_isai branch from 602516b to aeb7b02 Compare May 7, 2020 13:04
@upsj upsj added 1:ST:ready-for-review This PR is ready for review and removed 1:ST:WIP This PR is a work in progress. Not ready for review. labels May 7, 2020
@upsj upsj added the type:preconditioner This is related to the preconditioners label May 7, 2020
@upsj upsj force-pushed the arbitrary_gpu_isai branch from 19aaa79 to 2afe0a1 Compare May 7, 2020 15:19
@codecov
Copy link
codecov bot commented May 7, 2020

Codecov Report

Merging #520 into develop will increase coverage by 0.21%.
The diff coverage is 89.62%.

Impacted file tree graph

@@             Coverage Diff             @@
##           develop     #520      +/-   ##
===========================================
+ Coverage    88.44%   88.66%   +0.21%     
===========================================
  Files          276      276              
  Lines        17449    17729     +2
8000
80     
===========================================
+ Hits         15433    15719     +286     
+ Misses        2016     2010       -6     
Impacted Files Coverage Δ
core/device_hooks/common_kernels.inc.cpp 0.00% <0.00%> (ø)
core/test/matrix/csr.cpp 98.52% <ø> (ø)
core/test/matrix/sparsity_csr.cpp 98.21% <ø> (ø)
include/ginkgo/core/matrix/csr.hpp 68.62% <0.00%> (-2.08%) ⬇️
omp/factorization/par_ilu_kernels.cpp 100.00% <ø> (ø)
omp/preconditioner/isai_kernels.cpp 0.00% <0.00%> (ø)
core/preconditioner/isai.cpp 82.25% <83.67%> (+19.75%) ⬆️
reference/preconditioner/isai_kernels.cpp 97.58% <97.41%> (+2.34%) ⬆️
reference/test/preconditioner/isai_kernels.cpp 94.59% <99.48%> (+4.42%) ⬆️
core/test/base/math.cpp 100.00% <100.00%> (ø)
... and 18 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 269bf95...3fded2a. Read the comment docs.

Copy link
Member
@yhmtsai yhmtsai left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do not go into detail of merge and excess part yet

Copy link
Member
@yhmtsai yhmtsai left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM in general.
miss __launch_bounds__ and use gko::lend not lend

Copy link
Member
@thoasm thoasm left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM, I mostly have minor comments.

@upsj upsj force-pushed the arbitrary_gpu_isai branch from 960eaba to 1f84596 Compare May 15, 2020 13:29
@upsj upsj changed the title Add support for arbitrary row-lenghts in GPU ISAI Add support for arbitrary row-lengths in GPU ISAI May 17, 2020
Copy link
Member
@pratikvn pratikvn left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM!

@upsj upsj requested a review from yhmtsai May 18, 2020 11:10
Copy link
Member
@yhmtsai yhmtsai left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@upsj upsj force-pushed the arbitrary_gpu_isai branch from 1f84596 to 5150b7b Compare May 18, 2020 13:29
@upsj upsj 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 18, 2020
@upsj upsj force-pushed the arbitrary_gpu_isai branch from 5150b7b to bfce5b8 Compare May 19, 2020 11:42
upsj and others added 13 commits May 19, 2020 20:09
we no longer need it, as we only use
the sparsity pattern of the matrix,
but not its values as input for ISAI
Sonarqube is complaining again
* remove unused variables
* add __launchbounds__
* add __forceinline__

Co-authored-by: Yuhsiang M. Tsai <yhmtsai@gmail.com>
* replace `.get()` by `lend`
* test empty Csr for non-empty row ptrs
* rename ISAI test matrix location variable

Co-authored-by: Thomas Grützmacher <thomas.gruetzmacher@kit.edu>
@upsj upsj force-pushed the arbitrary_gpu_isai branch from bfce5b8 to 3fded2a Compare May 19, 2020 18:09
@sonarqubecloud
Copy link

SonarCloud Quality Gate failed.

Bug A 0 Bugs
Vulnerability A 0 Vulnerabilities (and Security Hotspot 0 Security Hotspots to review)
Code Smell A 13 Code Smells

73.2% 73.2% Coverage
12.7% 12.7% Duplication

@upsj upsj merged commit 6f0a247 into develop May 20, 2020
@upsj upsj deleted the arbitrary_gpu_isai branch May 20, 2020 06:54
@tcojean tcojean mentioned this pull request Jun 23, 2020
tcojean pushed a commit that referenced this pull request Jul 7, 2020
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
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
1:ST:ready-to-merge This PR is ready to merge. mod:cuda This is related to the CUDA module. mod:hip This is related to the HIP module. type:preconditioner This is related to the preconditioners
Projects
None yet
Development

Successfully merging this pull request may close these issues.

Improve ISAI to handle an arbitrary number of elements per row
4 participants
0