Support for shared libraries in GPU execution (python launch support) (#795)

* coreneuron and mechanism library can be built as shared and it
  enables launching coreneuron on GPU via python
* update MOD2C and NMODL fixes to handle GLOBAL variables
      See BlueBrain/mod2c/pull/78
      See BlueBrain/nmodl/pull/904
* removed acc/openmp global annotations for celsius, pi and secondorder
  and they don't need to be copied on GPU
* Pass Memb_list* as an argument for all common prototypes in order
   to support global variables via argument
* free ml->instance if not empty
* add link to libscopmath in neuron as well
* nrn_ghk is now declared inline.
* homegrown present table to avoid dynamic loading + acc_deviceptr limitations
* use -gpu=nordc and make #pragma acc routine seq functions inline
* drop -lscopmath as its folded in elsewhere
* random123 header reorganisation
* try and cleanup CLI11 handling.
* try and consolidate build logic
* some CORENEURON_ -> CORENRN_ for consistency.
* export OpenACC flags to NEURON separately as well as part
     of the whole ... -lcoreneuron ... link line.
* libcoreneuron.so -> libcorenrnmech.so, try and fix static builds
* do not enable OpenMP in shared/OpenACC builds.
* add rpaths inside nrnivmodl-core.
* accept a private destructor function pointer from generated mechanisms
* drop ${TEST_EXEC_PREFIX} that was causing simple tests to be executed on many ranks.
* CORENEURON_GPU_DEBUG: add environment variable that enables cnrn_target_* debug messages.

fixes #141

Co-authored-by: Olli Lupton <[email protected]>

Author: pramodk

.gitlab-ci.yml

@@ -81,127 +81,179 @@ build:nmodl:
     SPACK_PACKAGE_COMPILER: gcc
 
 # Build CoreNEURON
-build:coreneuron:mod2c:nvhpc:acc:
-  extends: [.build, .spack_nvhpc]
-  variables:
-    SPACK_PACKAGE: coreneuron
-    # See https://github.com/BlueBrain/CoreNeuron/issues/518 re: build_type
-    SPACK_PACKAGE_SPEC: +gpu+openmp+tests~legacy-unit build_type=RelWithDebInfo
-
-# Build CoreNEURON with Unified Memory on GPU
-build:coreneuron:mod2c:nvhpc:acc:unified:
-  extends: [.build, .spack_nvhpc]
-  variables:
-    SPACK_PACKAGE: coreneuron
-    # See https://github.com/BlueBrain/CoreNeuron/issues/518 re: build_type
-    SPACK_PACKAGE_SPEC: +gpu+unified+openmp+tests~legacy-unit build_type=RelWithDebInfo
-
-.build_coreneuron_nmodl:
+.build_coreneuron:
   extends: [.build]
   variables:
+    SPACK_PACKAGE: coreneuron
     # NEURON depends on py-mpi4py, most of whose dependencies are pulled in by
     # nmodl%gcc, with the exception of MPI, which is pulled in by
     # coreneuron%{nvhpc,intel}. hpe-mpi is an external package anyway, so
     # setting its compiler is just changing how it is labelled in the
     # dependency graph and not changing which installation is used, but this
     # means that in the NEURON step an existing py-mpi4py%gcc can be used.
     # Otherwise a new py-mpi4py with hpe-mpi%{nvhpc,intel} will be built.
+    # caliper: papi%nvhpc does not build; use the caliper from the deployment
     # TODO: fix this more robustly so we don't have to play so many games.
-    SPACK_PACKAGE_DEPENDENCIES: ^hpe-mpi%gcc
+    SPACK_PACKAGE_DEPENDENCIES: ^hpe-mpi%gcc ^caliper%gcc+cuda cuda_arch=70
 
-build:coreneuron:nmodl:nvhpc:omp:
-  extends: [.build_coreneuron_nmodl, .spack_nvhpc]
+# TODO: improve coverage by switching an Intel build to be statically linked
+# TODO: improve coverage by switching an Intel build to RelWithDebInfo
+# TODO: improve coverage by enabling +openmp on an Intel build
+build:coreneuron:mod2c:intel:shared:debug:
+  extends: [.build_coreneuron, .spack_intel]
   variables:
-    SPACK_PACKAGE: coreneuron
-    # See https://github.com/BlueBrain/CoreNeuron/issues/518 re: build_type
-    SPACK_PACKAGE_SPEC: +nmodl+openmp+gpu+tests~legacy-unit~sympy build_type=RelWithDebInfo
+    SPACK_PACKAGE_SPEC: +caliper~gpu~legacy-unit~nmodl~openmp+shared+tests~unified build_type=Debug
+
+build:coreneuron:nmodl:intel:debug:legacy:
+  extends: [.build_coreneuron, .spack_intel]
   needs: ["build:nmodl"]
+  variables:
+    SPACK_PACKAGE_SPEC: +caliper~gpu~legacy-unit+nmodl~openmp~shared~sympy+tests~unified build_type=Debug
 
-build:coreneuron:nmodl:nvhpc:acc:
-  extends: [.build_coreneuron_nmodl, .spack_nvhpc]
+# Disable caliper to improve coverage
+build:coreneuron:nmodl:intel:shared:debug:
+  extends: [.build_coreneuron, .spack_intel]
+  needs: ["build:nmodl"]
   variables:
-    SPACK_PACKAGE: coreneuron
-    # See https://github.com/BlueBrain/CoreNeuron/issues/518 re: build_type
-    # Sympy + OpenMP target offload does not currently work with NVHPC
-    SPACK_PACKAGE_SPEC: +nmodl~openmp+gpu+tests~legacy-unit+sympy build_type=RelWithDebInfo
+    SPACK_PACKAGE_DEPENDENCIES: ^hpe-mpi%gcc
+    SPACK_PACKAGE_SPEC: ~caliper~gpu~legacy-unit+nmodl~openmp+shared+sympy+tests~unified build_type=Debug
+
+# Not linked to a NEURON build+test job, see
+# https://github.com/BlueBrain/CoreNeuron/issues/594
+build:coreneuron:mod2c:nvhpc:acc:debug:unified:
+  extends: [.build_coreneuron, .spack_nvhpc]
+  variables:
+    SPACK_PACKAGE_SPEC: +caliper+gpu~legacy-unit~nmodl+openmp~shared+tests+unified build_type=Debug
+
+# Shared + OpenACC + OpenMP host threading has problems
+build:coreneuron:mod2c:nvhpc:acc:shared:
+  extends: [.build_coreneuron, .spack_nvhpc]
+  variables:
+    SPACK_PACKAGE_SPEC: +caliper+gpu~legacy-unit~nmodl~openmp+shared+tests~unified build_type=RelWithDebInfo
+
+build:coreneuron:nmodl:nvhpc:acc:debug:legacy:
+  extends: [.build_coreneuron, .spack_nvhpc]
   needs: ["build:nmodl"]
+  variables:
+    SPACK_PACKAGE_SPEC: +caliper+gpu~legacy-unit+nmodl~openmp~shared~sympy+tests~unified build_type=Debug
 
-build:coreneuron:mod2c:intel:
-  extends: [.build, .spack_intel]
+build:coreneuron:nmodl:nvhpc:acc:shared:
+  extends: [.build_coreneuron, .spack_nvhpc]
+  needs: ["build:nmodl"]
   variables:
-    SPACK_PACKAGE: coreneuron
-    SPACK_PACKAGE_SPEC: +tests~legacy-unit build_type=Debug
+    SPACK_PACKAGE_SPEC: +caliper+gpu~legacy-unit+nmodl~openmp+shared+sympy+tests~unified build_type=RelWithDebInfo
 
-build:coreneuron:nmodl:intel:
-  extends: [.build_coreneuron_nmodl, .spack_intel]
+build:coreneuron:nmodl:nvhpc:omp:legacy:
+  extends: [.build_coreneuron, .spack_nvhpc]
+  needs: ["build:nmodl"]
   variables:
-    SPACK_PACKAGE: coreneuron
-    SPACK_PACKAGE_SPEC: +nmodl+tests~legacy-unit build_type=Debug
+    SPACK_PACKAGE_SPEC: +caliper+gpu~legacy-unit+nmodl+openmp~shared~sympy+tests~unified build_type=RelWithDebInfo
+
+build:coreneuron:nmodl:nvhpc:omp:debug:
+  extends: [.build_coreneuron, .spack_nvhpc]
   needs: ["build:nmodl"]
+  variables:
+    SPACK_PACKAGE_SPEC: +caliper+gpu~legacy-unit+nmodl+openmp~shared+sympy+tests~unified build_type=Debug
 
 # Build NEURON
-build:neuron:mod2c:nvhpc:acc:
+build:neuron:mod2c:intel:shared:debug:
+  extends: [.build_neuron, .spack_intel]
+  needs: ["build:coreneuron:mod2c:intel:shared:debug"]
+
+build:neuron:nmodl:intel:debug:legacy:
+  extends: [.build_neuron, .spack_intel]
+  needs: ["build:coreneuron:nmodl:intel:debug:legacy"]
+
+build:neuron:nmodl:intel:shared:debug:
+  extends: [.build_neuron, .spack_intel]
+  needs: ["build:coreneuron:nmodl:intel:shared:debug"]
+
+build:neuron:mod2c:nvhpc:acc:shared:
   extends: [.build_neuron, .spack_nvhpc]
-  needs: ["build:coreneuron:mod2c:nvhpc:acc"]
+  needs: ["build:coreneuron:mod2c:nvhpc:acc:shared"]
 
-build:neuron:nmodl:nvhpc:omp:
+build:neuron:nmodl:nvhpc:acc:debug:legacy:
   extends: [.build_neuron, .spack_nvhpc]
-  needs: ["build:coreneuron:nmodl:nvhpc:omp"]
+  needs: ["build:coreneuron:nmodl:nvhpc:acc:debug:legacy"]
 
-build:neuron:nmodl:nvhpc:acc:
+build:neuron:nmodl:nvhpc:acc:shared:
   extends: [.build_neuron, .spack_nvhpc]
-  needs: ["build:coreneuron:nmodl:nvhpc:acc"]
+  needs: ["build:coreneuron:nmodl:nvhpc:acc:shared"]
 
-build:neuron:mod2c:intel:
-  extends: [.build_neuron, .spack_intel]
-  needs: ["build:coreneuron:mod2c:intel"]
+build:neuron:nmodl:nvhpc:omp:legacy:
+  extends: [.build_neuron, .spack_nvhpc]
+  needs: ["build:coreneuron:nmodl:nvhpc:omp:legacy"]
 
-build:neuron:nmodl:intel:
-  extends: [.build_neuron, .spack_intel]
-  needs: ["build:coreneuron:nmodl:intel"]
+build:neuron:nmodl:nvhpc:omp:debug:
+  extends: [.build_neuron, .spack_nvhpc]
+  needs: ["build:coreneuron:nmodl:nvhpc:omp:debug"]
 
 # Test CoreNEURON
-test:coreneuron:mod2c:nvhpc:acc:
+test:coreneuron:mod2c:intel:shared:debug:
+  extends: [.ctest]
+  needs: ["build:coreneuron:mod2c:intel:shared:debug"]
+
+test:coreneuron:nmodl:intel:debug:legacy:
+  extends: [.ctest]
+  needs: ["build:coreneuron:nmodl:intel:debug:legacy"]
+
+test:coreneuron:nmodl:intel:shared:debug:
+  extends: [.ctest]
+  needs: ["build:coreneuron:nmodl:intel:shared:debug"]
+
+test:coreneuron:mod2c:nvhpc:acc:debug:unified:
   extends: [.ctest, .gpu_node]
-  needs: ["build:coreneuron:mod2c:nvhpc:acc"]
+  needs: ["build:coreneuron:mod2c:nvhpc:acc:debug:unified"]
 
-test:coreneuron:mod2c:nvhpc:acc:unified:
+test:coreneuron:mod2c:nvhpc:acc:shared:
   extends: [.ctest, .gpu_node]
-  needs: ["build:coreneuron:mod2c:nvhpc:acc:unified"]
+  needs: ["build:coreneuron:mod2c:nvhpc:acc:shared"]
 
-test:coreneuron:nmodl:nvhpc:omp:
+test:coreneuron:nmodl:nvhpc:acc:debug:legacy:
   extends: [.ctest, .gpu_node]
-  needs: ["build:coreneuron:nmodl:nvhpc:omp"]
+  needs: ["build:coreneuron:nmodl:nvhpc:acc:debug:legacy"]
 
-test:coreneuron:nmodl:nvhpc:acc:
+test:coreneuron:nmodl:nvhpc:acc:shared:
   extends: [.ctest, .gpu_node]
-  needs: ["build:coreneuron:nmodl:nvhpc:acc"]
+  needs: ["build:coreneuron:nmodl:nvhpc:acc:shared"]
 
-test:coreneuron:mod2c:intel:
-  extends: [.ctest]
-  needs: ["build:coreneuron:mod2c:intel"]
+test:coreneuron:nmodl:nvhpc:omp:legacy:
+  extends: [.ctest, .gpu_node]
+  needs: ["build:coreneuron:nmodl:nvhpc:omp:legacy"]
 
-test:coreneuron:nmodl:intel:
-  extends: [.ctest]
-  needs: ["build:coreneuron:nmodl:intel"]
+test:coreneuron:nmodl:nvhpc:omp:debug:
+  extends: [.ctest, .gpu_node]
+  needs: ["build:coreneuron:nmodl:nvhpc:omp:debug"]
 
 # Test NEURON
-test:neuron:mod2c:nvhpc:acc:
+test:neuron:mod2c:intel:shared:debug:
+  extends: [.test_neuron]
+  needs: ["build:neuron:mod2c:intel:shared:debug"]
+
+test:neuron:nmodl:intel:debug:legacy:
+  extends: [.test_neuron]
+  needs: ["build:neuron:nmodl:intel:debug:legacy"]
+
+test:neuron:nmodl:intel:shared:debug:
+  extends: [.test_neuron]
+  needs: ["build:neuron:nmodl:intel:shared:debug"]
+
+test:neuron:mod2c:nvhpc:acc:shared:
   extends: [.test_neuron, .gpu_node]
-  needs: ["build:neuron:mod2c:nvhpc:acc"]
+  needs: ["build:neuron:mod2c:nvhpc:acc:shared"]
 
-test:neuron:nmodl:nvhpc:omp:
+test:neuron:nmodl:nvhpc:acc:debug:legacy:
   extends: [.test_neuron, .gpu_node]
-  needs: ["build:neuron:nmodl:nvhpc:omp"]
+  needs: ["build:neuron:nmodl:nvhpc:acc:debug:legacy"]
 
-test:neuron:nmodl:nvhpc:acc:
+test:neuron:nmodl:nvhpc:acc:shared:
   extends: [.test_neuron, .gpu_node]
-  needs: ["build:neuron:nmodl:nvhpc:acc"]
+  needs: ["build:neuron:nmodl:nvhpc:acc:shared"]
 
-test:neuron:mod2c:intel:
-  extends: [.test_neuron]
-  needs: ["build:neuron:mod2c:intel"]
+test:neuron:nmodl:nvhpc:omp:legacy:
+  extends: [.test_neuron, .gpu_node]
+  needs: ["build:neuron:nmodl:nvhpc:omp:legacy"]
 
-test:neuron:nmodl:intel:
-  extends: [.test_neuron]
-  needs: ["build:neuron:nmodl:intel"]
+test:neuron:nmodl:nvhpc:omp:debug:
+  extends: [.test_neuron, .gpu_node]
+  needs: ["build:neuron:nmodl:nvhpc:omp:debug"]