Include Eigen API from NMODL : Eigen OpenACC compatibility issues (#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

CMakeLists.txt

@@ -176,6 +176,15 @@ if(CORENRN_ENABLE_GPU)
         " --generate-code=arch=compute_${cuda_arch},code=[compute_${cuda_arch},sm_${cuda_arch}]")
     endforeach()
   endif()
+
+  # ~~~
+  # Needed for the Eigen GPU support Warning suppression (Eigen GPU-related):
+  # 3057 : Warning on ignoring __host__ annotation in some functions
+  # 3085 : Warning on redeclaring a __host__ function as __host__ __device__
+  # ~~~
+  set(CMAKE_CUDA_FLAGS
+      "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr -Xcudafe --diag_suppress=3057,--diag_suppress=3085"
+  )
 endif()
 
 # =============================================================================

README.md

@@ -133,7 +133,7 @@ With CoreNEURON, existing NEURON models can be run with minimal changes. For a g
 3. If GPU support is enabled during build, enable GPU execution using :
 	```
 	coreneuron.gpu = True
-  ```
+	```
 
 4. Use `psolve` to run simulation after initialization :
 

coreneuron/CMakeLists.txt

@@ -128,6 +128,14 @@ if(CORENRN_ENABLE_GPU)
   # setting LANGUAGE=CXX in non-GPU builds does not work.
   list(REMOVE_ITEM CORENEURON_CODE_FILES "${CMAKE_CURRENT_SOURCE_DIR}/utils/randoms/nrnran123.cpp")
   list(APPEND CORENEURON_CODE_FILES ${CORENEURON_CUDA_FILES})
+
+  # Eigen-3.5+ provides better GPU support. However, some functions cannot be called directly from
+  # within an OpenACC region. 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 ...
+  if(CORENRN_ENABLE_NMODL AND EXISTS ${CORENRN_MOD2CPP_INCLUDE}/partial_piv_lu/partial_piv_lu.cu)
+    list(APPEND CORENEURON_CODE_FILES ${CORENRN_MOD2CPP_INCLUDE}/partial_piv_lu/partial_piv_lu.cu)
+  endif()
 endif()
 
 # =============================================================================