Eigen openacc compatibility by kotsaloscv · Pull Request #690 · BlueBrain/nmodl

kotsaloscv · 2021-06-07T06:19:26Z

This PR solves the compatibility issues of Eigen with GPUs (I fixed both the OpenACC & CUDA backends).

The Eigen::PartialPivLU (LU decomposition to solve linear systems) is not compatible with GPUs (no device tokens).

For matrices up to 4x4, the Eigen inverse() has template specializations decorated with host & device tokens. Therefore, we use the inverse method instead of the PartialPivLU (requires an invertible matrix) which supports both CPUs & GPUs.

For matrices 5x5 and above, Eigen does not provide GPU-enabled methods to solve small linear systems. For this reason, we use the Crout LU decomposition (Legacy code : coreneuron/sim/scopmath/crout_thread.cpp). For the CPU exectutions, we use the standard PartialPivLU from Eigen.

For the Crout LU-decomposition, I have included a unit test to compare its result with Eigen::PartialPivLU.

bbpbuildbot · 2021-06-07T06:19:29Z

Can one of the admins verify this patch?

bbpbuildbot · 2021-06-07T06:52:24Z

Logfiles from GitLab pipeline #8037 (:no_entry:) have been uploaded here!

Status and direct links:

bbp-hpcteam

Just a quick note for failing clang-format: note that we use clang-format v11 for formatting check under CI. I have added an item for this Thursday meeting to discuss project specific clang-format version.

bbpbuildbot · 2021-06-07T11:49:13Z

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

Status and direct links:

bbpbuildbot · 2021-06-07T11:50:06Z

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

Status and direct links:

cattabiani · 2021-06-07T12:48:12Z

+    static int count_length(const std::vector<SymbolType>& variables) {
+        int length = 0;
+        for (const auto& variable: variables) {
+            length += variable->get_length();
+        }
+        return length;
+    };


this function is not really needed since it can be replaced with a one-liner:

std::accumulate(v.begin(), v.end(), 0, [](int l, const SymbolType& variable) {return a+= variable->get_length(); })

cattabiani · 2021-06-07T12:52:14Z

 void CodegenCudaVisitor::print_backend_includes() {
    printer->add_line("#include <cuda.h>");
+
+    if (info.eigen_linear_solver_exist && count_length(info.state_vars) > 4) {


not blocking but a one-liner with accumulate here is better than a new static function since the code remains cleaner. As you prefer

cattabiani · 2021-06-07T12:55:26Z

+
+
+template <typename T>
+bool test_Crout_correctness(T rtol = 1e-6, T atol = 1e-6) {


these tolerances are quite big. why?

@cattabiani you mean that we can relax them even more (like 1e-3)?

I think he meant the inverse - why not something much closer to eps ?

Yes, I meant what Ohm said

@ohm314 & @cattabiani : I have modified Crout solver and I have resolved an accuracy issue that I had. I computed the relative error (A*x-b)/b for both eigen and crout, and crout's relative error was consistently two orders of magnitude higher than eigen one (before the modification). Therefore, for small tolerances my tests where failing. Now (after the modification), the relative error for both solvers has the same order of magnitude.

ohm314 · 2021-06-07T14:05:43Z

+
+
+template <typename T>
+bool test_Crout_correctness(T rtol = 1e-6, T atol = 1e-6) {


I think he meant the inverse - why not something much closer to eps ?

ohm314 · 2021-06-07T14:55:45Z

+#pragma acc routine seq
+#endif
+template <typename T>
+EIGEN_DEVICE_FUNC inline void Crout(int d, T* S, T* D) {


can this not be done with anything better than C-style arrays?

I prefered not to change the initial solver (C-style arrays) to avoid bugs. The C-style arrays are readily combined with the .data() of Eigen and therefore, I think that after an extensive testing, we could consider its modernization (if everything passes).

…resolve accuracy issues

bbpbuildbot · 2021-06-08T13:09:26Z

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

Status and direct links:

…wMajor storage order

bbpbuildbot · 2021-06-09T07:26:42Z

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

Status and direct links:

…wMajor storage order (improvement with in-place transposition)

bbpbuildbot · 2021-06-09T09:39:18Z

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

Status and direct links:

…r and its unit test (full compatibility with OpenACC/CUDA backends)

bbpbuildbot · 2021-06-09T18:15:44Z

Logfiles from GitLab pipeline #8532 (:no_entry:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-06-09T18:15:58Z

Logfiles from GitLab pipeline #8531 (:no_entry:) have been uploaded here!

Status and direct links:

…ag for PGI - resolve failing CI

bbpbuildbot · 2021-06-10T08:23:16Z

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

Status and direct links:

bbpbuildbot · 2021-07-20T09:35:12Z

Logfiles from GitLab pipeline #11210 (:no_entry:) have been uploaded here!

Status and direct links:

bbpbuildbot · 2021-08-27T10:10:45Z

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

Status and direct links:

olupton

Looking good, let's test it again after a merge/rebase with the latest master.

olupton · 2021-08-27T10:52:53Z

 #include <algorithm>
 #include <cmath>
 #include <ctime>
+#include <numeric>


Why is this needed?

olupton · 2021-08-27T10:53:56Z

+                                                   const std::string& X,
+                                                   const std::string& Jm,
+                                                   const std::string& F) {
+    // The Eigen::PartialPivLU is not compatible with GPUs (no __device__ tokens).


Is this still true with the latest Eigen changes we plan to use? (https://github.com/BlueBrain/eigen/)

olupton · 2021-08-27T10:57:37Z

  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} --diag_suppress 1,82,111,115,177,186,611,997,1097,1625")
+
+  # Needed for Eigen
+  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wc,--pending_instantiations=0")


The surrounding code has changed a bit in master, it would be good to merge/rebase before final review/testing.

olupton · 2021-08-27T11:01:00Z

+    using VecType = Matrix<T, Dynamic, 1>;
+
+    std::random_device rd;  // seeding
+    std::mt19937 mt(rd());


If we're going to use a pseudorandom seed then we should at least print what it is, in case we need to reproduce/debug an issue.

bbpbuildbot · 2021-09-02T11:59:33Z

Logfiles from GitLab pipeline #15548 (:no_entry:) have been uploaded here!

Status and direct links:

kotsaloscv · 2021-09-02T12:05:02Z

For the build PGI pipeline we will need to bring back in CMake the following:
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wc,--pending_instantiations=0")
That's why it is failing for now

pramodk · 2021-09-06T15:11:43Z

Superseded by #728

Christos Kotsalos added 9 commits May 31, 2021 08:53

Eigen compatibility with OpenACC #311 : WIP

13eb7c9

Merge branch 'master' into eigen_openacc_compatibility

16afc0d

Eigen compatibility with OpenACC #311 : WIP

f3734ea

hpc-coding-conventions update

a1a95f1

WIP

887141a

Eigen compatibility with OpenACC #311 : WIP

d057bdb

Eigen compatibility with OpenACC #311 : WIP

dda0ff0

Eigen compatibility with OpenACC #311 : WIP

62bdd32

Merge branch 'master' into eigen_openacc_compatibility

79a96e5

kotsaloscv requested review from iomaganaris and pramodk June 7, 2021 06:19

kotsaloscv linked an issue Jun 7, 2021 that may be closed by this pull request

Eigen compatibility with OpenACC #311

Closed

kotsaloscv requested review from cattabiani and ohm314 June 7, 2021 08:20

bbp-hpcteam reviewed Jun 7, 2021

View reviewed changes

Eigen compatibility with OpenACC #311 : clang-format v11

dd4b984

cattabiani reviewed Jun 7, 2021

View reviewed changes

ohm314 reviewed Jun 7, 2021

View reviewed changes

Eigen compatibility with OpenACC #311 : Crout solver modification to …

3466882

…resolve accuracy issues

Eigen compatibility with OpenACC #311 : Resolved a bug with Eigen::Ro…

808d06e

…wMajor storage order

Eigen compatibility with OpenACC #311 : Resolved a bug with Eigen::Ro…

41fc88c

…wMajor storage order (improvement with in-place transposition)

Eigen compatibility with OpenACC #311 / #135 : Updated Newton's solve…

a94d025

…r and its unit test (full compatibility with OpenACC/CUDA backends)

kotsaloscv linked an issue Jun 9, 2021 that may be closed by this pull request

Issue with newton solver when used with OpenACC backend #135

Closed

Eigen compatibility with OpenACC #311 / #135 : Added a compilation fl…

c4a0e21

…ag for PGI - resolve failing CI

Eigen compatibility with OpenACC #311 / #135 : Merge with master

a369989

kotsaloscv closed this Aug 27, 2021

kotsaloscv reopened this Aug 27, 2021

olupton reviewed Aug 27, 2021

View reviewed changes

kotsaloscv removed request for iomaganaris and pramodk August 27, 2021 12:24

Merge branch 'master' into eigen_openacc_compatibility

6d32376

kotsaloscv mentioned this pull request Sep 6, 2021

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

Merged

pramodk closed this Sep 6, 2021

pramodk mentioned this pull request Oct 5, 2022

Remove dependency with Eigen (and use internal Crout solver implementation instead) #943

Closed

1uc deleted the eigen_openacc_compatibility branch July 12, 2024 14:03



		template <typename T>
		bool test_Crout_correctness(T rtol = 1e-6, T atol = 1e-6) {

Uh oh!

Conversation

kotsaloscv commented Jun 7, 2021

Uh oh!

bbpbuildbot commented Jun 7, 2021

Uh oh!

bbpbuildbot commented Jun 7, 2021

Uh oh!

bbp-hpcteam left a comment

Choose a reason for hiding this comment

Uh oh!

bbpbuildbot commented Jun 7, 2021

Uh oh!

bbpbuildbot commented Jun 7, 2021

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

bbpbuildbot commented Jun 8, 2021

Uh oh!

bbpbuildbot commented Jun 9, 2021

Uh oh!

bbpbuildbot commented Jun 9, 2021

Uh oh!

bbpbuildbot commented Jun 9, 2021

Uh oh!

bbpbuildbot commented Jun 9, 2021

Uh oh!

bbpbuildbot commented Jun 10, 2021

Uh oh!

bbpbuildbot commented Jul 20, 2021

Uh oh!

bbpbuildbot commented Aug 27, 2021

Uh oh!

olupton left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

bbpbuildbot commented Sep 2, 2021

Uh oh!

kotsaloscv commented Sep 2, 2021

Uh oh!

pramodk commented Sep 6, 2021

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

7 participants