partialPivLu: use nvc++ -cuda

BlueBrain · Aug 16, 2022 · d44340f · d44340f
1 parent 16b918c
commit d44340f
Show file tree

Hide file tree

Showing 4 changed files with 80 additions and 70 deletions.
diff --git a/ext/eigen b/ext/eigen
diff --git a/src/solver/partial_piv_lu/partial_piv_lu.cpp b/src/solver/partial_piv_lu/partial_piv_lu.cpp
@@ -0,0 +1,61 @@
+/*************************************************************************
+ * Copyright (C) 2018-2022 Blue Brain Project
+ *
+ * This file is part of NMODL distributed under the terms of the GNU
+ * Lesser General Public License. See top-level LICENSE file for details.
+ *************************************************************************/
+#define NMODL_EIGEN_NO_OPENACC  // disable OpenACC/OpenMP annotations
+#include "partial_piv_lu/partial_piv_lu.h"
+
+// See explanation in partial_piv_lu.h
+NMODL_EIGEN_ATTR VecType<1> partialPivLu1(const MatType<1>& A, const VecType<1>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<2> partialPivLu2(const MatType<2>& A, const VecType<2>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<3> partialPivLu3(const MatType<3>& A, const VecType<3>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<4> partialPivLu4(const MatType<4>& A, const VecType<4>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<5> partialPivLu5(const MatType<5>& A, const VecType<5>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<6> partialPivLu6(const MatType<6>& A, const VecType<6>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<7> partialPivLu7(const MatType<7>& A, const VecType<7>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<8> partialPivLu8(const MatType<8>& A, const VecType<8>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<9> partialPivLu9(const MatType<9>& A, const VecType<9>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<10> partialPivLu10(const MatType<10>& A, const VecType<10>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<11> partialPivLu11(const MatType<11>& A, const VecType<11>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<12> partialPivLu12(const MatType<12>& A, const VecType<12>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<13> partialPivLu13(const MatType<13>& A, const VecType<13>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<14> partialPivLu14(const MatType<14>& A, const VecType<14>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<15> partialPivLu15(const MatType<15>& A, const VecType<15>& b) {
+    return A.partialPivLu().solve(b);
+}
+NMODL_EIGEN_ATTR VecType<16> partialPivLu16(const MatType<16>& A, const VecType<16>& b) {
+    return A.partialPivLu().solve(b);
+}
+
+// Currently there is an issue in Eigen (GPU-branch) for matrices 17x17 and above.
+// ToDo: Check in a future release if this issue is resolved!
diff --git a/src/solver/partial_piv_lu/partial_piv_lu.cu b/src/solver/partial_piv_lu/partial_piv_lu.cu
diff --git a/src/solver/partial_piv_lu/partial_piv_lu.h b/src/solver/partial_piv_lu/partial_piv_lu.h
@@ -34,17 +34,24 @@ using VecType = Eigen::Matrix<double, dim, 1, Eigen::ColMajor, dim, 1>;
 //    instantiations visible in this header file that call:
 //  - partialPivLuN(...), functions that are declared in this header but defined
 //    in the CUDA file partial_piv_lu.cu.
+#ifdef NMODL_EIGEN_NO_OPENACC
+#define NMODL_EIGEN_ATTR __host__ __device__
+#define NMODL_EIGEN_ROUTINE_SEQ
+#else
 nrn_pragma_omp(declare target)
-nrn_pragma_acc(routine seq)
+#define NMODL_EIGEN_ATTR
+#define NMODL_EIGEN_ROUTINE_SEQ nrn_pragma_acc(routine seq)
+#endif
+NMODL_EIGEN_ROUTINE_SEQ
 template <int dim>
-EIGEN_DEVICE_FUNC VecType<dim> partialPivLu(const MatType<dim>&, const VecType<dim>&);
-#define InstantiatePartialPivLu(N)                                                               \
-    nrn_pragma_acc(routine seq)                                                                  \
-    EIGEN_DEVICE_FUNC VecType<N> partialPivLu##N(const MatType<N>&, const VecType<N>&);          \
-    nrn_pragma_acc(routine seq)                                                                  \
-    template <>                                                                                  \
-    EIGEN_DEVICE_FUNC inline VecType<N> partialPivLu(const MatType<N>& A, const VecType<N>& b) { \
-        return partialPivLu##N(A, b);                                                            \
+NMODL_EIGEN_ATTR VecType<dim> partialPivLu(const MatType<dim>&, const VecType<dim>&);
+#define InstantiatePartialPivLu(N)                                                              \
+    NMODL_EIGEN_ROUTINE_SEQ                                                                     \
+    NMODL_EIGEN_ATTR VecType<N> partialPivLu##N(const MatType<N>&, const VecType<N>&);          \
+    NMODL_EIGEN_ROUTINE_SEQ                                                                     \
+    template <>                                                                                 \
+    NMODL_EIGEN_ATTR inline VecType<N> partialPivLu(const MatType<N>& A, const VecType<N>& b) { \
+        return partialPivLu##N(A, b);                                                           \
     }
 InstantiatePartialPivLu(1)
 InstantiatePartialPivLu(2)
@@ -63,4 +70,6 @@ InstantiatePartialPivLu(14)
 InstantiatePartialPivLu(15)
 InstantiatePartialPivLu(16)
 #undef InstantiatePartialPivLu
+#ifndef NMODL_EIGEN_NO_OPENACC
 nrn_pragma_omp(end declare target)
+#endif