Skip to content
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

port cuda coo to amd coo #357

Merged
merged 17 commits into from
Nov 14, 2019
Merged

port cuda coo to amd coo #357

merged 17 commits into from
Nov 14, 2019

Conversation

yhmtsai
Copy link
Member

@yhmtsai yhmtsai commented Oct 8, 2019

This PR ports cuda coo implementation into amd coo.

HIP does not support the warp-level cooperative group.
I add some implementations of warp-level cooperative group, it works fine now but it is not same as cooperative group.
I write the similar interface to wrap the shfl, any, etc functions.
If HIP or we implement the correct cooperative group, it only need to change cooperative_group.hip.hpp and does not hurt any other implementation.
Moreover, the code is messy and the optimization is nothing (I use min(300, nnz/warp_size) for calculating the number of warp, 300 is randomly set).
The optimization is done but need to tune. AMD has 4 SIMD in each multiprocessor. I use it to calculate the number of warps.

There are two failed test (Coo.AdvancedApplyIsEquivalentToRef and Coo.AdvancedApplyToDenseMatrixIsEquivalentToRef) because they need dense hip code.
Edit: comment them until implementing hip dense kernel.
Other tests of Coo are passed.

Summary:

  1. add cooperative_group for hip
  2. split config and type
  3. hip coo codes
  4. Create common folder to handle the duplicated code of hip and cuda.
    the name.hpp.inc in common puts the subnamespace and does not include any headers.
namespace proper {
    namespace parent {
        // some constexpr variable
        #include "common/path/to/name.hpp.inc"
        // other codes
    }
}

TODO:

@yhmtsai yhmtsai added type:matrix-format This is related to the Matrix formats 1:ST:WIP This PR is a work in progress. Not ready for review. mod:hip This is related to the HIP module. labels Oct 8, 2019
@yhmtsai yhmtsai self-assigned this Oct 8, 2019
@codecov
Copy link

codecov bot commented Oct 8, 2019

Codecov Report

❗ No coverage uploaded for pull request base (develop@f4ae401). Click here to learn what that means.
The diff coverage is 32.2%.

Impacted file tree graph

@@            Coverage Diff             @@
##             develop     #357   +/-   ##
==========================================
  Coverage           ?   64.02%           
==========================================
  Files              ?      265           
  Lines              ?    16358           
  Branches           ?        0           
==========================================
  Hits               ?    10474           
  Misses             ?     5884           
  Partials           ?        0
Impacted Files Coverage Δ
include/ginkgo/core/base/lin_op.hpp 78.48% <ø> (ø)
core/test/utils/assertions.hpp 85.03% <ø> (ø)
include/ginkgo/core/base/types.hpp 92.59% <ø> (ø)
reference/test/matrix/hybrid_kernels.cpp 100% <ø> (ø)
include/ginkgo/core/base/exception_helpers.hpp 100% <ø> (ø)
include/ginkgo/core/base/math.hpp 90.32% <ø> (ø)
core/base/extended_float.hpp 79.12% <ø> (ø)
include/ginkgo/core/stop/criterion.hpp 84.61% <ø> (ø)
benchmark/spmv/spmv.cpp 0% <ø> (ø)
benchmark/utils/general.hpp 0% <0%> (ø)
... and 26 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 f4ae401...e0caafb. Read the comment docs.

@yhmtsai yhmtsai force-pushed the amd_coo branch 2 times, most recently from 0f77458 to c653acb Compare October 23, 2019 09:12
@yhmtsai yhmtsai mentioned this pull request Oct 25, 2019
2 tasks
@yhmtsai yhmtsai force-pushed the amd_coo branch 4 times, most recently from c4ae452 to 2a506f9 Compare October 28, 2019 11:46
@yhmtsai yhmtsai mentioned this pull request Oct 29, 2019
2 tasks
@yhmtsai yhmtsai added reg:benchmarking This is related to benchmarking. and removed 1:ST:WIP This PR is a work in progress. Not ready for review. labels Oct 29, 2019
@yhmtsai yhmtsai force-pushed the amd_coo branch 2 times, most recently from 86591d7 to 57b1ce3 Compare October 30, 2019 20:37
@yhmtsai yhmtsai added the 1:ST:ready-for-review This PR is ready for review label Oct 31, 2019
Copy link
Member

@tcojean tcojean 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. I do have some comments though.

CMakeLists.txt Outdated Show resolved Hide resolved
benchmark/spmv/CMakeLists.txt Outdated Show resolved Hide resolved
benchmark/spmv/CMakeLists.txt Outdated Show resolved Hide resolved
hip/CMakeLists.txt Outdated Show resolved Hide resolved
hip/base/config.hip.hpp Show resolved Hide resolved
hip/matrix/coo_kernels.hip.cpp Outdated Show resolved Hide resolved
hip/matrix/coo_kernels.hip.cpp Outdated Show resolved Hide resolved
hip/matrix/coo_kernels.hip.cpp Show resolved Hide resolved
hip/matrix/coo_kernels.hip.cpp Outdated Show resolved Hide resolved
hip/matrix/coo_kernels.hip.cpp Outdated Show resolved Hide resolved
hip/components/cooperative_groups.hip.hpp Outdated Show resolved Hide resolved
hip/components/segment_scan.hip.hpp Outdated Show resolved Hide resolved
include/ginkgo/core/base/executor.hpp Outdated Show resolved Hide resolved
@yhmtsai yhmtsai force-pushed the amd_coo branch 5 times, most recently from f2175d2 to 03f5fe5 Compare November 8, 2019 13:32
@yhmtsai
Copy link
Member Author

yhmtsai commented Nov 8, 2019

This is updated coo vs hipsp_coo
plot

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!

hip/components/atomic.hip.hpp Outdated Show resolved Hide resolved
hip/components/format_conversion.hip.hpp Outdated Show resolved Hide resolved
include/ginkgo/core/base/executor.hpp Show resolved Hide resolved
Copy link
Member

@upsj upsj left a comment

Choose a reason for hiding this comment

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

LGTM!
Most of the changes were code moving between files and tiny refactorings anyways.
A comment and a suggestion:

  1. The tests for HIP COO kernels are missing two test cases that are present in the CUDA tests
  2. Could we not quickly manually implement the any, all and ballot methods for sub-warp tiles?

hip/components/atomic.hip.hpp Outdated Show resolved Hide resolved
hip/components/atomic.hip.hpp Outdated Show resolved Hide resolved
cuda/components/format_conversion.cuh Outdated Show resolved Hide resolved
hip/components/format_conversion.hip.hpp Outdated Show resolved Hide resolved
hip/test/matrix/coo_kernels.hip.cpp Show resolved Hide resolved
hip/test/matrix/coo_kernels.hip.cpp Show resolved Hide resolved
cuda/matrix/dense_kernels.cu Show resolved Hide resolved
hip/components/cooperative_groups.hip.hpp Show resolved Hide resolved
@tcojean tcojean 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 Nov 14, 2019
@yhmtsai yhmtsai merged commit 876234e into develop Nov 14, 2019
@yhmtsai yhmtsai deleted the amd_coo branch November 21, 2019 20:13
@tcojean tcojean mentioned this pull request Jun 23, 2020
tcojean added 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:hip This is related to the HIP module. reg:benchmarking This is related to benchmarking. type:matrix-format This is related to the Matrix formats
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

5 participants