GPU build improvements for NEURON integration (BlueBrain/CoreNeuron#373)

* Update README with latest workflow

* Avoid linking to C++ MPI libraries
  - MPI C++ bindings are removed from MPI
  - avoid linking mpi c++ library for simplicity
  - print mpi c/c++ compiler

* Improvemenets for GPU build integration with NEURON
  - create property CORENEURON_LIB_LINK_FLAGS with necessary LDFLAGS
    to link coreneuron with neuron via nrnivmodl
  - suppress additional PGI warnings which are not necessary to emit
    with old mod2c
  - minor change in nrn_setup.hpp : functions declarated as static
    in the header result into warning as definitions are in .cpp file

* Update weights from GPU before copying to NEURON

CoreNEURON Repo SHA: BlueBrain/CoreNeuron@dca1dc0

Author: pramodk

cmake/coreneuron/CMakeLists.txt

@@ -159,6 +159,10 @@ endif()
 if(CORENRN_ENABLE_MPI)
   find_package(MPI REQUIRED)
   add_definitions("-DNRNMPI=1")
+  # avoid linking to C++ bindings
+  add_definitions("-DMPI_NO_CPPBIND=1")
+  add_definitions("-DOMPI_SKIP_MPICXX=1")
+  add_definitions("-DMPICH_SKIP_MPICXX=1")
 else()
   add_definitions("-DNRNMPI=0")
   add_definitions("-DNRN_MULTISEND=0")
@@ -369,6 +373,8 @@ if(cmake_generator_tolower MATCHES "makefile")
   message(STATUS " Build option       | Status")
   message(STATUS "--------------------+--------------------------------------------------------")
 
+  message(STATUS "C COMPILER          | ${CMAKE_C_COMPILER}")
+  message(STATUS "CXX COMPILER        | ${CMAKE_CXX_COMPILER}")
   message(STATUS "COMPILE FLAGS       | ${COMPILER_FLAGS} ${CMAKE_CXX_FLAGS}")
   message(STATUS "Build Type          | ${COMPILE_LIBRARY_TYPE}")
   message(STATUS "MPI                 | ${CORENRN_ENABLE_MPI}")

cmake/coreneuron/OpenAccHelper.cmake

@@ -23,8 +23,8 @@ if(CORENRN_ENABLE_GPU)
     # workaround for old PGI version
     add_definitions(-DPG_ACC_BUGS)
     set(ACC_FLAGS "-acc")
-    # disable very verbose diagnosis messages
-    set(PGI_DIAG_FLAGS "--diag_suppress 177")
+    # disable very verbose diagnosis messages and obvious warnings for mod2c
+    set(PGI_DIAG_FLAGS "--diag_suppress 161,177,550")
     # some of the mod files can have too many functions, increase inline level
     set(PGI_INLINE_FLAGS "-Minline=size:200,levels:10")
     # C/C++ compiler flags
@@ -37,6 +37,14 @@ if(CORENRN_ENABLE_GPU)
     message(WARNING "Non-PGI compiler : make sure to add required compiler flags to enable OpenACC")
   endif()
 
+  # set property for neuron to link with coreneuron libraries
+  set_property(
+    GLOBAL
+    PROPERTY
+      CORENEURON_LIB_LINK_FLAGS
+      "-acc -rdynamic -lrt -Wl,--whole-archive -L${CMAKE_HOST_SYSTEM_PROCESSOR} -lcorenrnmech -L${CMAKE_INSTALL_PREFIX}/lib -lcoreneuron -lcudacoreneuron -Wl,--no-whole-archive ${CUDA_cudart_static_LIBRARY}"
+  )
+
   # find_cuda produce verbose messages : use new behavior to use _ROOT variables
   if(POLICY CMP0074)
     cmake_policy(SET CMP0074 NEW)
@@ -49,4 +57,6 @@ else(CORENRN_ENABLE_GPU)
   # OpenACC pragmas are not guarded, disable all unknown pragm warnings
   set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${IGNORE_UNKNOWN_PRAGMA_FLAGS}")
   set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${IGNORE_UNKNOWN_PRAGMA_FLAGS}")
+  set_property(GLOBAL PROPERTY CORENEURON_LIB_LINK_FLAGS
+                               "-L${CMAKE_HOST_SYSTEM_PROCESSOR} -lcorenrnmech")
 endif(CORENRN_ENABLE_GPU)

external/mod2c

@@ -1 +1 @@
-Subproject commit 2b3ff22c4a785db174e915dc36a9ead5543ba673
+Subproject commit 7b0623a2e70dc0ab0f143719f725efb2526964b5

src/coreneuron/CMakeLists.txt

@@ -168,7 +168,7 @@ add_library(corenrnmech ${COMPILE_LIBRARY_TYPE} ${CORENEURON_HEADER_FILES} ${ENG
 
 target_link_libraries(
   coreneuron
-  ${MPI_CXX_LIBRARIES}
+  ${MPI_C_LIBRARIES}
   ${reportinglib_LIBRARY}
   ${sonatareport_LIBRARY}
   ${link_cudacoreneuron}
@@ -207,8 +207,7 @@ add_dependencies(coreneuron scopmath)
 include_directories(${CORENEURON_PROJECT_SOURCE_DIR})
 add_executable(nrniv-core "apps/coreneuron.cpp")
 
-target_link_libraries(nrniv-core corenrnmech ${reportinglib_LIBRARY} ${sonatareport_LIBRARY}
-                      ${MPI_C_LIBRARIES} ${MPI_CXX_LIBRARIES})
+target_link_libraries(nrniv-core corenrnmech ${reportinglib_LIBRARY} ${sonatareport_LIBRARY})
 
 set_target_properties(nrniv-core PROPERTIES RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/bin)
 

src/coreneuron/apps/main1.cpp

@@ -595,7 +595,13 @@ extern "C" int run_solve_core(int argc, char** argv) {
 
     // copy weights back to NEURON NetCon
     if (nrn2core_all_weights_return_) {
+
+        // first update weights from gpu
+        update_weights_from_gpu(nrn_threads, nrn_nthread);
+
+        // store weight pointers
         std::vector<double*> weights(nrn_nthread, NULL);
+
         // could be one thread more (empty) than in NEURON but does not matter
         for (int i=0; i < nrn_nthread; ++i) {
           weights[i] = nrn_threads[i].weights;

src/coreneuron/gpu/nrn_acc_manager.cpp

@@ -765,6 +765,24 @@ void update_voltage_from_gpu(NrnThread* nt) {
     }
 }
 
+/**
+ * Copy weights from GPU to CPU
+ *
+ * User may record NetCon weights at the end of simulation.
+ * For this purpose update weights of all NrnThread objects
+ * from GPU to CPU.
+ */
+void update_weights_from_gpu(NrnThread* threads, int nthreads) {
+    for (int i = 0; i < nthreads; i++) {
+        NrnThread* nt = threads + i;
+        size_t n_weight = nt->n_weight;
+        if (nt->compute_gpu && n_weight > 0) {
+            double* weights = nt->weights;
+            #pragma acc update host(weights[0 : n_weight])
+        }
+    }
+}
+
 void update_matrix_from_gpu(NrnThread* _nt) {
 #ifdef _OPENACC
     if (_nt->compute_gpu && (_nt->end > 0)) {

src/coreneuron/gpu/nrn_acc_manager.hpp

@@ -21,6 +21,7 @@ void update_net_receive_buffer(NrnThread* _nt);
 void realloc_net_receive_buffer(NrnThread* nt, Memb_list* ml);
 void update_net_send_buffer_on_host(NrnThread* nt, NetSendBuffer_t* nsb);
 void update_voltage_from_gpu(NrnThread* nt);
+void update_weights_from_gpu(NrnThread* threads, int nthreads);
 void init_gpu();
 
 }  // namespace coreneuron

src/coreneuron/io/nrn_setup.hpp

@@ -37,10 +37,10 @@ THE POSSIBILITY OF SUCH DAMAGE.
 #include "coreneuron/io/mem_layout_util.hpp"
 
 namespace coreneuron {
-static void read_phase1(NrnThread& nt, UserParams& userParams);
-static void read_phase2(NrnThread& nt, UserParams& userParams);
-static void read_phase3(NrnThread& nt, UserParams& userParams);
-static void read_phasegap(NrnThread& nt, UserParams& userParams);
+void read_phase1(NrnThread& nt, UserParams& userParams);
+void read_phase2(NrnThread& nt, UserParams& userParams);
+void read_phase3(NrnThread& nt, UserParams& userParams);
+void read_phasegap(NrnThread& nt, UserParams& userParams);
 static void setup_ThreadData(NrnThread& nt);
 
 // Functions to load and clean data;