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

pragma unroll warning in hip #492

Merged
merged 6 commits into from
Apr 23, 2020
Merged

pragma unroll warning in hip #492

merged 6 commits into from
Apr 23, 2020

Conversation

yhmtsai
Copy link
Member

@yhmtsai yhmtsai commented Mar 31, 2020

This PR removes the pragma unroll warning in hip.
In general, the warning in ginkgo is caused by using break in the loop.

for () {
  if (condition) break;
  // other commands
}

->

for () {
  if (!condition) {
    // other commands
  }
}

to make pragma unroll work. ( I am not sure whether it affect the performance.)

copy_matrix unroll problem:
When using clang and hip to compile, copy_matrix gives the wrong result on custom precision.
I need to use #pragma unroll 1 to disable it.

I add local one variable to store source_row value to make everything work as expected.
I only go through some different part of fix_13.isa and unfix_13.isa. The assembly looks similar except the number of VGPR.
For some reasons, hip use many VGPR (registers) in the original version (without using the local variable)
unfix_13.isa contains NumSpilledVGPRs and use full VGPR (256), but fix_13.isa does not need to use full VGPR.
unroll_isa.zip
If someone is interested in the assembly code, I use # main - start and # main - end around

const auto val = source_row[i * increment];
destination[get_row_major_index<mod>(idx, row_perm, stride)] = static_cast<ResultValueType>(val);

or destination[get_row_major_index<mod>(idx, row_perm, stride)] = static_cast<ResultValueType>(source_row[i * increment]);
This issue ROCm/aomp#24 also mentions some weird things about NumSpilledVGPRs although the reason might be different from this PR.

Note. The original version gives wrong answer when casting to custom precision format and the wrong result are stable.

@yhmtsai yhmtsai added the mod:hip This is related to the HIP module. label Mar 31, 2020
@yhmtsai yhmtsai self-assigned this Mar 31, 2020
@codecov
Copy link

codecov bot commented Mar 31, 2020

Codecov Report

Merging #492 into develop will decrease coverage by 0.18%.
The diff coverage is n/a.

Impacted file tree graph

@@             Coverage Diff             @@
##           develop     #492      +/-   ##
===========================================
- Coverage    88.64%   88.45%   -0.19%     
===========================================
  Files          268      268              
  Lines        16946    16894      -52     
===========================================
- Hits         15021    14944      -77     
- Misses        1925     1950      +25     
Impacted Files Coverage Δ
omp/components/format_conversion.hpp 40.00% <0.00%> (-60.00%) ⬇️
include/ginkgo/core/matrix/ell.hpp 68.75% <0.00%> (-31.25%) ⬇️
include/ginkgo/core/base/lin_op.hpp 58.33% <0.00%> (-11.91%) ⬇️
include/ginkgo/core/matrix/hybrid.hpp 73.83% <0.00%> (-10.38%) ⬇️
include/ginkgo/core/matrix/sellp.hpp 81.81% <0.00%> (-8.43%) ⬇️
include/ginkgo/core/base/name_demangling.hpp 85.71% <0.00%> (-3.18%) ⬇️
reference/components/format_conversion.hpp 100.00% <0.00%> (ø)
include/ginkgo/core/base/utils.hpp 87.23% <0.00%> (+1.81%) ⬆️
include/ginkgo/core/matrix/csr.hpp 70.70% <0.00%> (+2.07%) ⬆️
include/ginkgo/core/matrix/dense.hpp 97.67% <0.00%> (+2.32%) ⬆️
... and 2 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 8ec56dd...2ec6529. Read the comment docs.

@pratikvn pratikvn added the is:affects-performance This is related to something which affects performance. label Apr 1, 2020
@@ -109,7 +109,6 @@ __device__ void reduce(const Group &__restrict__ group,
{
const auto local_id = group.thread_rank();

#pragma unroll
for (int k = group.size() / 2; k >= config::warp_size; k /= 2) {
Copy link
Member Author

Choose a reason for hiding this comment

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

This group is larger than warp_size, so we can not unroll this loop in the compile time.

Copy link
Member

Choose a reason for hiding this comment

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

Makes sense, thanks for clearing that up.

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.

I am perfectly fine with changing the breaks, especially if it improves the unrolling. However, I think there are some minor errors in the code currently.

common/components/reduction.hpp.inc Show resolved Hide resolved
common/components/warp_blas.hpp.inc Outdated Show resolved Hide resolved
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 except for Thomas' __GNUC__ comment.

@yhmtsai
Copy link
Member Author

yhmtsai commented Apr 2, 2020

__GNUC__, which stands for supporting GNU C, is also defined in clang compiler.
hcc/hipcc also define clang, so I can not use clang to identify clang.
does anyone have an idea for identify clang/gcc only?

@upsj
Copy link
Member

upsj commented Apr 2, 2020

This is only device code, right? Wouldn't it be sufficient to differentiate between nvcc and the ROCm compiler? Why do we need to check for gcc?

@yhmtsai
Copy link
Member Author

yhmtsai commented Apr 2, 2020

Yes, it is only on device.
For some reasons, using gcc in hip works well but using clang gives wrong answer.
Thus, I also wonder whether there are some potential bug in the code.

@tcojean
Copy link
Member

tcojean commented Apr 2, 2020

For gcc:

#define GCC_COMPILER (defined(__GNUC__) && !defined(__clang__) && !defined(__INTEL_COMPILER))

For clang (maybe?):

#define CLANG_COMPILER (defined(__clang__) && !defined(__HIPCC__))

I did not check it but probably something like this.

@upsj
Copy link
Member

upsj commented Apr 2, 2020

Is there any way to look at the intermediate representations generated by ROCm? I guess they separate their code just like nvcc does, right? Because they should hopefully give an insight into the differences between the compilers.

@yhmtsai yhmtsai added the 1:ST:do-not-merge Please do not merge PR this yet. label Apr 3, 2020
@yhmtsai yhmtsai requested review from thoasm and tcojean April 15, 2020 14:51
@yhmtsai yhmtsai added 1:ST:ready-for-review This PR is ready for review and removed 1:ST:do-not-merge Please do not merge PR this yet. labels Apr 15, 2020
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!

common/components/warp_blas.hpp.inc Outdated 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!

@yhmtsai yhmtsai 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 Apr 22, 2020
@sonarcloud
Copy link

sonarcloud bot commented Apr 22, 2020

Kudos, SonarCloud Quality Gate passed!

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

No Coverage information No Coverage information
0.0% 0.0% Duplication

@yhmtsai yhmtsai merged commit b4feb96 into develop Apr 23, 2020
@yhmtsai yhmtsai deleted the pragma_unroll branch April 23, 2020 06:43
@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. is:affects-performance This is related to something which affects performance. mod:hip This is related to the HIP module.
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

6 participants