port cuda matrix into hip #373

yhmtsai · 2019-10-29T08:29:38Z

This PR ports cuda csr/dense/ell/hybrid/sellp/sparsity_csr into hip

add cuda2hip script to convert cuda code to hip code.
It uses hipify-perl first, and then replace all name corresponding to cuda to hip's (cuda->hip, CUBLAS -> HIPBLAS etc)
It can not handle the optimization, and the format. cuda2hip will add #include "hip/hip_runtime.hpp" in the first line, so need to modify it manually.
add sparselib strategy which is same as cusparse
(hip and cuda csr kernels need to handle cusparse and sparselib for converting to each other)
Update strategy parameter according to executor parameter.
Note: If setting up nwarps manually, strategy will be replaced when converting.
split __host__ __device__ of zero/one. The original __host__ __device__ zero will make hipcc fail but __host__ zero, __device__ zero are successful. I do not know why, but it works.
make ell selection kernel handle 32/64 warp_size
Merge Initial code for the HIP executor #344 (HIP Executor)
Merge port cuda coo to amd coo #357 (Coo part, benchmark, and cooperative group)

upsj

LGTM! Mainly small questions, only the inconsistent call to the HIP API in transpose and the overflow in the merge_path binary search might have to be addressed.

cuda/matrix/csr_kernels.cu

include/ginkgo/core/matrix/csr.hpp

hip/base/hipblas_bindings.hip.hpp

hip/matrix/dense_kernels.hip.cpp

hip/matrix/ell_kernels.hip.cpp

common/matrix/csr_kernels.hpp.inc

hip/matrix/dense_kernels.hip.cpp

common/components/reduction.hpp.inc

dev_tools/scripts/cuda2hip.sh

hip/components/uninitialized_array.hip.hpp

hip/test/matrix/csr_kernels.hip.cpp

hip/test/matrix/hybrid_kernels.hip.cpp

include/ginkgo/core/matrix/csr.hpp

tcojean

LGTM overall. Some minor changes.

common/components/warp_blas.hpp.inc

common/components/prefix_sum.hpp.inc

hip/matrix/ell_kernels.hip.cpp

include/ginkgo/core/base/math.hpp

include/ginkgo/core/matrix/csr.hpp

pratikvn

Thanks for adding the docs on the strategies for CSR. Can you please add them for Hybrid as well ? That would be great.

hip/matrix/ell_kernels.hip.cpp

include/ginkgo/core/matrix/csr.hpp

cuda/base/cublas_bindings.hpp

tcojean

LGTM.

pratikvn

LGTM!

tcojean · 2019-11-15T13:03:45Z

As sonarqube says, it's probably better to have a namespace for the helper in csr:
https://sonarcloud.io/project/issues?id=ginkgo-project_ginkgo&pullRequest=373&resolved=false&types=BUG

Co-Authored-By: Tobias Ribizel <[email protected]>

Co-Authored-By: Pratik Nayak Co-Authored-By: tcojean

upsj

LGTM! Just a small question that I'm unsure about, as I don't know the code-base that well.

include/ginkgo/core/matrix/csr.hpp

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 added type:matrix-format This is related to the Matrix formats mod:hip This is related to the HIP module. labels Oct 29, 2019

yhmtsai self-assigned this Oct 29, 2019

yhmtsai force-pushed the cuda2hip_matrix branch 5 times, most recently from 4126b06 to eace3f2 Compare October 31, 2019 14:02

yhmtsai force-pushed the cuda2hip_matrix branch 7 times, most recently from ac5129b to 867e8ff Compare November 12, 2019 20:25

thoasm mentioned this pull request Nov 12, 2019

Improve CSR strategies #379

Merged

yhmtsai force-pushed the cuda2hip_matrix branch from c3075ac to 377cd33 Compare November 13, 2019 10:33

upsj mentioned this pull request Nov 14, 2019

port cuda coo to amd coo #357

Merged

2 tasks

upsj requested changes Nov 14, 2019

View reviewed changes

upsj reviewed Nov 14, 2019

View reviewed changes

hip/matrix/dense_kernels.hip.cpp Outdated Show resolved Hide resolved

yhmtsai force-pushed the cuda2hip_matrix branch from ac3b637 to 81eddb7 Compare November 14, 2019 10:23

pratikvn reviewed Nov 14, 2019

View reviewed changes

common/components/reduction.hpp.inc Outdated Show resolved Hide resolved

dev_tools/scripts/cuda2hip.sh Outdated Show resolved Hide resolved

hip/components/uninitialized_array.hip.hpp Outdated Show resolved Hide resolved

pratikvn requested changes Nov 14, 2019

View reviewed changes

hip/test/matrix/csr_kernels.hip.cpp Outdated Show resolved Hide resolved

hip/test/matrix/hybrid_kernels.hip.cpp Show resolved Hide resolved

include/ginkgo/core/matrix/csr.hpp Show resolved Hide resolved

tcojean requested changes Nov 14, 2019

View reviewed changes

yhmtsai force-pushed the cuda2hip_matrix branch from 5611f0a to 2e8dd58 Compare November 14, 2019 19:07

pratikvn reviewed Nov 14, 2019

View reviewed changes

tcojean reviewed Nov 15, 2019

View reviewed changes

cuda/base/cublas_bindings.hpp Outdated Show resolved Hide resolved

tcojean approved these changes Nov 15, 2019

View reviewed changes

tcojean added the 1:ST:ready-for-review This PR is ready for review label Nov 15, 2019

pratikvn approved these changes Nov 15, 2019

View reviewed changes

yhmtsai force-pushed the cuda2hip_matrix branch from d490704 to 0bf2a44 Compare November 15, 2019 15:30

yhmtsai and others added 16 commits November 15, 2019 23:46

add cuda2hip script

0c11fcd

add the matrix related file from cuda2hip.sh

fb3a28f

fix the error of hip and add CMakeList

5187aa4

fix math.hip.hpp

f6a8c53

fix format

70809ce

split __host__ __device__, include thrust in test

8ca796f

move the duplicated code into common

95fd12a

add csr strategy in hip

083d039

move cuda2hip.sh into dev_tools/scripts

9d2ec11

fix mindpoint overflow problem

53cc65b

Co-Authored-By: Tobias Ribizel <[email protected]>

review update, fix reduce in complex norm2

83ecf7d

Apply suggestions from code review

3ba176a

Co-Authored-By: Pratik Nayak Co-Authored-By: tcojean

cuda2hip.sh handle namespace update, review update

792cd22

review update

e95424e

move duplicated math into common

8622838

review update and fix sonar issue

2342c02

yhmtsai force-pushed the cuda2hip_matrix branch from 0bf2a44 to 2342c02 Compare November 15, 2019 22:51

upsj approved these changes Nov 16, 2019

View reviewed changes

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

fix csr strategy

28f9fa6

yhmtsai force-pushed the cuda2hip_matrix branch from 3ca2519 to 28f9fa6 Compare November 16, 2019 18:42

yhmtsai merged commit d38e846 into develop Nov 16, 2019

yhmtsai deleted the cuda2hip_matrix branch November 21, 2019 20:13

daineAMD mentioned this pull request Jan 8, 2020

Hipblasversionbug ROCm/hipBLAS#143

Merged

tcojean mentioned this pull request Jun 23, 2020

Release/1.2.0 #576

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

port cuda matrix into hip #373

port cuda matrix into hip #373

yhmtsai commented Oct 29, 2019 •

edited

Loading

upsj left a comment

tcojean left a comment

pratikvn left a comment

tcojean left a comment

pratikvn left a comment

tcojean commented Nov 15, 2019

upsj left a comment

port cuda matrix into hip #373

port cuda matrix into hip #373

Conversation

yhmtsai commented Oct 29, 2019 • edited Loading

upsj left a comment

Choose a reason for hiding this comment

tcojean left a comment

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

pratikvn left a comment

Choose a reason for hiding this comment

tcojean commented Nov 15, 2019

upsj left a comment

Choose a reason for hiding this comment

yhmtsai commented Oct 29, 2019 •

edited

Loading