Skip to content
This repository was archived by the owner on Mar 25, 2025. It is now read-only.

Compatibility issues between Eigen and GPUs (OpenACC/CUDA)#728

Merged
pramodk merged 8 commits into
masterfrom
kotsalos/solve_eigen_GPU_issues
Sep 6, 2021
Merged

Compatibility issues between Eigen and GPUs (OpenACC/CUDA)#728
pramodk merged 8 commits into
masterfrom
kotsalos/solve_eigen_GPU_issues

Conversation

@kotsaloscv

Copy link
Copy Markdown
Contributor

This PR solves the compatibility issues between Eigen and GPUs (OpenACC/CUDA).

Two factors contribute to the above solution:

  1. New Eigen branch (version 3.5 and above). Currently we are using a mirrored version of Eigen (https://github.com/BlueBrain/eigen/releases/tag/v3.5-alpha.1).
  2. An API that makes possible to call any Eigen __device__ function from within OpenACC regions. In more details, Eigen-3.5+ provides better GPU support; however, some functions cannot be called directly from within OpenACC regions. Therefore, we need to wrap them in a special API (decorate them with __device__ & acc routine tokens), which allows us to eventually call them from OpenACC. Calling these functions from CUDA kernels presents no issue ...

This PR closes/replaces the following PRs:

  1. Eigen openacc compatibility #690
  2. Fix the issue with Eigen+OpenACC while executing on GPU #726

This PR works in combination with CoreNEURON's PR # 624 (minor changes to include the new Eigen API).

This PR closes the following issues:

  1. Eigen compatibility with OpenACC #311
  2. Issue with newton solver when used with OpenACC backend #135

@bbpbuildbot

Copy link
Copy Markdown
Collaborator

Logfiles from GitLab pipeline #15814 (:white_check_mark:) have been uploaded here!

Status and direct links:

@iomaganaris iomaganaris left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

LGTM. Just a few questions

Comment thread src/codegen/codegen_acc_visitor.cpp
Comment thread src/codegen/codegen_c_visitor.hpp Outdated
Comment thread src/codegen/codegen_c_visitor.cpp
@bbpbuildbot

Copy link
Copy Markdown
Collaborator

Logfiles from GitLab pipeline #15838 (:white_check_mark:) have been uploaded here!

Status and direct links:

@codecov-commenter

Copy link
Copy Markdown

Codecov Report

Merging #728 (20ed7d9) into master (1321c7b) will decrease coverage by 0.04%.
The diff coverage is 24.32%.

Impacted file tree graph

@@            Coverage Diff             @@
##           master     #728      +/-   ##
==========================================
- Coverage   60.71%   60.67%   -0.05%     
==========================================
  Files         200      200              
  Lines       33362    33385      +23     
==========================================
  Hits        20255    20255              
- Misses      13107    13130      +23     
Impacted Files Coverage Δ
src/codegen/codegen_acc_visitor.cpp 0.00% <0.00%> (ø)
src/codegen/codegen_acc_visitor.hpp 0.00% <ø> (ø)
src/codegen/codegen_c_visitor.hpp 91.66% <ø> (ø)
src/solver/newton/newton.hpp 94.33% <ø> (ø)
src/codegen/codegen_c_visitor.cpp 63.09% <32.14%> (-0.35%) ⬇️

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 1321c7b...20ed7d9. Read the comment docs.

@pramodk pramodk left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

Apart from minor comment, LGTM

Comment thread src/solver/newton/newton.hpp
Comment thread src/solver/newton/newton.hpp
Comment thread src/solver/newton/newton.hpp
@pramodk pramodk merged commit 5f4462a into master Sep 6, 2021
@pramodk pramodk deleted the kotsalos/solve_eigen_GPU_issues branch September 6, 2021 15:27
pramodk pushed a commit to BlueBrain/CoreNeuron that referenced this pull request Sep 7, 2021
* Resolve Eigen-OpenACC compatibility issues
   - See the changes done in NMODL: BlueBrain/nmodl/pull/728
* As part of this PR we now build `partial_piv_lu.cu` from NMODL
* Some CUDA warnings are suppressed because they are assumed
   safe and can be ignored by end users.
* Update README to fix code examples
pramodk pushed a commit to neuronsimulator/nrn that referenced this pull request Nov 2, 2022
…ueBrain/CoreNeuron#624)

* Resolve Eigen-OpenACC compatibility issues
   - See the changes done in NMODL: BlueBrain/nmodl/pull/728
* As part of this PR we now build `partial_piv_lu.cu` from NMODL
* Some CUDA warnings are suppressed because they are assumed
   safe and can be ignored by end users.
* Update README to fix code examples

CoreNEURON Repo SHA: BlueBrain/CoreNeuron@2b3e746
JCGoran pushed a commit to neuronsimulator/nrn that referenced this pull request Mar 12, 2025
…rain/nmodl#728)

* Two factors contribute to the above solution:
   - New eigen branch (version 3.5 and above). Currently we are
      using a mirrored version of Eigen in BlueBrain organisation
      https://github.com/BlueBrain/eigen/releases/tag/v3.5-alpha
   - An API that makes possible to call any Eigen `__device__` 
      function from within OpenACC regions.
* More details: Eigen-3.5+ provides better GPU support; however,
   some functions cannot be called directly from within OpenACC regions.
   Therefore, we need to wrap them in a special API (decorate them with 
    `__device__` & `acc routine` tokens), which allows us to eventually call
    them from OpenACC. Calling these functions from CUDA kernels presents
    no issue.
*  From BlueBrain/nmodl#726: Avoid use `[]` operator with eigen Matrix objects. This results
    into runtime error with OpenACC and PGI compiler.
* Note that this should works in combination with BlueBrain/CoreNeuron/pull/624

fixes BlueBrain/nmodl#311 
fixes BlueBrain/nmodl#135 

NMODL Repo SHA: BlueBrain/nmodl@5f4462a
Sign up for free to subscribe to this conversation on GitHub. Already have an account? Sign in.

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

Eigen compatibility with OpenACC Issue with newton solver when used with OpenACC backend

5 participants