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 solver without triangular solver to hip #393

Merged
merged 4 commits into from
Nov 27, 2019

Conversation

yhmtsai
Copy link
Member

@yhmtsai yhmtsai commented Nov 19, 2019

This PR ports cuda solver without triangular solver to hip

The test of preconditioner needs to use ginkgo_create_test not ginkgo_create_hip_test.
Otherwise, the stop factory is always empty.

The triangular solver needs some extra works, so I do not put it in this PR.

Note: Some conversion fails from cuda to hip.
The block_size or grid_size is too long to hipify-perl can not recognize it correctly, so need some fixes by hand.
fail1-1 and fail1-2: splitting is failed.
should be dim3(ceildiv(dim_size[0] * stride_next_krylov, default_block_size)), dim3(default_block_size), 0, 0
but get dim3(ceildiv(dim_size[0] * stride_next_krylov)), dim3(default_block_size)), default_block_size, 0

fail2: hipify does not know it is a device call, so hipify does nothing on it.

@yhmtsai yhmtsai added type:solver This is related to the solvers mod:hip This is related to the HIP module. labels Nov 19, 2019
@yhmtsai yhmtsai self-assigned this Nov 19, 2019
@tcojean
Copy link
Member

tcojean commented Nov 19, 2019

@yhmtsai be sure to note down any such problems for a probable future paper

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! But I'm not familiar enough with HIP to understand the test failures.

@tcojean
Copy link
Member

tcojean commented Nov 21, 2019

Here is what I tried to make ginkgo_create_hip_test work with the CG test. Nothing worked so far. I didn't have much time to explore. If possible @yhmtsai can you take over?

  • I can confirm I don't see constructor anywhere in any of the binaries produced. It is most likely inlined (the heavy use of constexpr in the dim class may have something to do with it). The constructor could potentially be everywhere (hip module, ginkgo core module, and the cg test itself should include dim.hpp).
  • in include/ginkgo/core/solver/cg.hpp, change remove system matrix's std:move in case this created a bug, though there should be none: indeed did not change anything.
  • in include/ginkgo/core/base/exception_helpers.hpp, change detail::get_size to return const reference (hope to not use any copy constructor, though compiler should optimize this), did not seem to work.

@yhmtsai
Copy link
Member Author

yhmtsai commented Nov 22, 2019

@tcojean Sure, I will also try to fix it

@yhmtsai yhmtsai force-pushed the cuda2hip_solver_without_trs branch 2 times, most recently from 340cbab to 438bdc7 Compare November 25, 2019 15:35
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! One minor cmake function rename.

cmake/create_test.cmake Outdated Show resolved Hide resolved
@yhmtsai
Copy link
Member Author

yhmtsai commented Nov 26, 2019

I list the output from three type (gcc/hip/gcc compile hcc compile)here.
We can go into it in detail at some points to see why the other two are failed.

link

== compile:gcc link:hip
/opt/rocm/hip/bin/hipcc_cmake_linker_helper /opt/rocm/hcc    CMakeFiles/hip_test_solver_cg_kernels.dir/cg_kernels.cpp.o -o cg_kernels  -L/usr/lib/gcc/x86_64-linux-gnu/7 ../../../core/libginkgo.a ../../../third_party/gtest/build/googlemock/gtest/./libgtest_main.a ../../../third_party/gtest/build/googlemock/gtest/./libgtest.a ../../../core/device_hooks/libginkgo_omp.a ../../../core/device_hooks/libginkgo_cuda.a ../../../reference/libginkgo_reference.a ../../libginkgo_hip.a /opt/rocm/hip/lib/libhip_hcc.so -L /opt/rocm/lib -Wl,-rpath /opt/rocm/lib -Wl,--whole-archive /opt/rocm/lib/libmcwamp.so.2.9 -Wl,--no-whole-archive -ldl -lm /opt/rocm/lib/libhc_am.so.2.9 /opt/rocm/lib/libhsa-runtime64.so /opt/rocm/hipblas/lib/libhipblas.so.0.1 /opt/rocm/hipsparse/lib/libhipsparse.so.0.1 -lstdc++ -lm -lgcc_s -lgcc -lc -lgcc_s -lgcc
== compile:gcc link:gcc
/usr/bin/c++  -O3 -DNDEBUG   CMakeFiles/hip_test_solver_cg_kernels.dir/cg_kernels.cpp.o  -o cg_kernels -Wl,-rpath,/opt/rocm/hip/lib:/opt/rocm/lib:/opt/rocm/hipblas/lib:/opt/rocm/hipsparse/lib ../../../core/libginkgo.a ../../../third_party/gtest/build/googlemock/gtest/./libgtest_main.a ../../../third_party/gtest/build/googlemock/gtest/./libgtest.a ../../../core/device_hooks/libginkgo_omp.a ../../../core/device_hooks/libginkgo_cuda.a ../../../reference/libginkgo_reference.a ../../libginkgo_hip.a /opt/rocm/hip/lib/libhip_hcc.so -L /opt/rocm/lib -Wl,-rpath /opt/rocm/lib -Wl,--whole-archive /opt/rocm/lib/libmcwamp.so.2.9 -Wl,--no-whole-archive -ldl -lm /opt/rocm/lib/libhc_am.so.2.9 /opt/rocm/lib/libhsa-runtime64.so /opt/rocm/hipblas/lib/libhipblas.so.0.1 /opt/rocm/hipsparse/lib/libhipsparse.so.0.1
== compile:hip link:hip
/opt/rocm/hip/bin/hipcc_cmake_linker_helper /opt/rocm/hcc    CMakeFiles/hip_test_solver_cg_kernels.dir/hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o -o cg_kernels  -L/usr/lib/gcc/x86_64-linux-gnu/7 ../../../core/libginkgo.a ../../../third_party/gtest/build/googlemock/gtest/./libgtest_main.a ../../../third_party/gtest/build/googlemock/gtest/./libgtest.a ../../../core/device_hooks/libginkgo_omp.a ../../../core/device_hooks/libginkgo_cuda.a ../../../reference/libginkgo_reference.a ../../libginkgo_hip.a /opt/rocm/hip/lib/libhip_hcc.so -L /opt/rocm/lib -Wl,-rpath /opt/rocm/lib -Wl,--whole-archive /opt/rocm/lib/libmcwamp.so.2.9 -Wl,--no-whole-archive -ldl -lm /opt/rocm/lib/libhc_am.so.2.9 /opt/rocm/lib/libhsa-runtime64.so /opt/rocm/hipblas/lib/libhipblas.so.0.1 /opt/rocm/hipsparse/lib/libhipsparse.so.0.1 -lstdc++ -lm -lgcc_s -lgcc -lc -lgcc_s -lgcc

compile

== compile:gcc link:hip
cd /root/ginkgo/build/hip/test/solver && /usr/bin/c++   -I/root/ginkgo/build -I/root/ginkgo/build/include -I/root/ginkgo/include -I/root/ginkgo -I/opt/rocm/hip/include -isystem /root/ginkgo/build/third_party/gtest/src/googletest/include  -O3 -DNDEBUG   -std=gnu++11 -o CMakeFiles/hip_test_solver_cg_kernels.dir/cg_kernels.cpp.o -c /root/ginkgo/hip/test/solver/cg_kernels.cpp
== compile:gcc link:gcc
cd /root/ginkgo/build/hip/test/solver && /usr/bin/c++   -I/root/ginkgo/build -I/root/ginkgo/build/include -I/root/ginkgo/include -I/root/ginkgo -I/opt/rocm/hip/include -isystem /root/ginkgo/build/third_party/gtest/src/googletest/include  -O3 -DNDEBUG   -std=gnu++11 -o CMakeFiles/hip_test_solver_cg_kernels.dir/cg_kernels.cpp.o -c /root/ginkgo/hip/test/solver/cg_kernels.cpp
== compile:hip link:hip
Building HIPCC object hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir/hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o
cd /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir && /usr/local/bin/cmake -E make_directory /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//.
cd /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir && /usr/local/bin/cmake -D verbose:BOOL=1 -D build_configuration:STRING=RELEASE -D generated_file:STRING=/root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//./hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o -P /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.cmake
-- Removing /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//./hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o
/usr/local/bin/cmake -E remove /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//./hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o
-- Generating dependency file: /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend.pre
/opt/rocm/hip/bin/hipcc -M /root/ginkgo/hip/test/solver/cg_kernels.hip.cpp -o /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend.pre -std=c++11 "-I/root/ginkgo/build -I/opt/rocm/include -I/opt/rocm/hipblas/include -I/opt/rocm/hipsparse/include -I/root/ginkgo/build/include -I/root/ginkgo/include -I/root/ginkgo -I/root/ginkgo/build/include -I/root/ginkgo/include -I/root/ginkgo -I/root/ginkgo/build/include -I/root/ginkgo/include -I/root/ginkgo -I/opt/rocm/hip/include -I/root/ginkgo/build/include -I/root/ginkgo/include -I/root/ginkgo -I/root/ginkgo/build/include -I/root/ginkgo/include -I/root/ginkgo -I/root/ginkgo/build/third_party/gtest/src/googletest/include -I/root/ginkgo/build/third_party/gtest/src/googletest/include"
-- Generating temporary cmake readable file: /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend.tmp
/usr/local/bin/cmake -D input_file:FILEPATH=/root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend.pre -D output_file:FILEPATH=/root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend.tmp -D verbose=1 -P /opt/rocm/hip/cmake/FindHIP/run_make2cmake.cmake
-- Copy if different /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend.tmp to /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend
/usr/local/bin/cmake -E copy_if_different /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend.tmp /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend
-- Removing /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend.tmp and /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend.pre
/usr/local/bin/cmake -E remove /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend.tmp /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o.depend.pre
-- Generating /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//./hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o
/opt/rocm/hip/bin/hipcc -c /root/ginkgo/hip/test/solver/cg_kernels.hip.cpp -o /root/ginkgo/build/hip/test/solver/CMakeFiles/hip_test_solver_cg_kernels.dir//./hip_test_solver_cg_kernels_generated_cg_kernels.hip.cpp.o -std=c++11 "-I/root/ginkgo/build -I/opt/rocm/include -I/opt/rocm/hipblas/include -I/opt/rocm/hipsparse/include -I/root/ginkgo/build/include -I/root/ginkgo/include -I/root/ginkgo -I/root/ginkgo/build/include -I/root/ginkgo/include -I/root/ginkgo -I/root/ginkgo/build/include -I/root/ginkgo/include -I/root/ginkgo -I/opt/rocm/hip/include -I/root/ginkgo/build/include -I/root/ginkgo/include -I/root/ginkgo -I/root/ginkgo/build/include -I/root/ginkgo/include -I/root/ginkgo -I/root/ginkgo/build/third_party/gtest/src/googletest/include -I/root/ginkgo/build/third_party/gtest/src/googletest/include"

@tcojean
Copy link
Member

tcojean commented Nov 26, 2019

@yhmtsai so what works for now is building with gcc and linking with HIP for these tests?

@yhmtsai
Copy link
Member Author

yhmtsai commented Nov 26, 2019

@tcojean Yes, building with gcc and linking with hip works, but I do not know how it works.
Only use this way when build ginkgo static library and the platform is hcc. Others still use gcc to compile and link.

@yhmtsai yhmtsai merged commit bb8b830 into develop Nov 27, 2019
@tcojean tcojean deleted the cuda2hip_solver_without_trs branch November 27, 2019 11:26
@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
mod:hip This is related to the HIP module. type:solver This is related to the solvers
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

4 participants