CMake build options refactoring for Makefile and removing PG_ACC_BUGS (…

…#455)

* CMake build options refactoring for Makefile
   * move all makefile related options into single file MakefileBuildOptions.cmake
   * remove all option processing from extra/CMakeLists.txt
   * fix with neuron linking : type CN_PGI_ACC_FLAGS -> PGI_ACC_FLAGS
* Use PGI archive module because of issue with latest NVHPC modules
* fix ISPC build issue and nrnivmodl-core failure on OSX
* Avoid P2 UC2 ndoes as GPFS is not mounted (HELP-13390)
* Use spack branch only if specified
* Remove use of PG_ACC_BUGS
   * PG_ACC_BUGS was added when we saw bugs with PGI compiler
     version 16.3
   * Except the case of union, no more real buggy scenario exists
   * So remove all usages of PG_ACC_BUGS
   * Remove some hardcoded ion array size by const variable
   * update mod2c to master
* mods files and cpp translated files are in same x86_64/corenrn/mod2c directory

CMake/MakefileBuildOptions.cmake

@@ -0,0 +1,73 @@
+# =============================================================================
+# Common CXX and ISPC flags
+# =============================================================================
+
+# ISPC should compile with --pic by default
+set(CMAKE_ISPC_FLAGS "${CMAKE_ISPC_FLAGS} --pic")
+
+# =============================================================================
+# NMODL CLI options : common and backend specific
+# =============================================================================
+# if user pass arguments then use those as common arguments
+if ("${CORENRN_NMODL_FLAGS}" STREQUAL "")
+  set(NMODL_COMMON_ARGS "passes --inline")
+else()
+  set(NMODL_COMMON_ARGS ${CORENRN_NMODL_FLAGS})
+endif()
+
+set(NMODL_CPU_BACKEND_ARGS "host --c")
+set(NMODL_ISPC_BACKEND_ARGS "host --ispc")
+set(NMODL_ACC_BACKEND_ARGS "host --c acc --oacc")
+
+# =============================================================================
+# Extract Compile definitions : common to all backend
+# =============================================================================
+get_directory_property(COMPILE_DEFS COMPILE_DEFINITIONS)
+if(COMPILE_DEFS)
+    set(CORENRN_COMMON_COMPILE_DEFS "")
+    foreach(flag ${COMPILE_DEFS})
+        set(CORENRN_COMMON_COMPILE_DEFS "${CORENRN_COMMON_COMPILE_DEFS} -D${flag}")
+    endforeach()
+endif()
+
+# =============================================================================
+# link flags : common to all backend
+# =============================================================================
+# ~~~
+# find_cuda uses FindThreads that adds below imported target we
+# shouldn't add imported target to link line
+# ~~~
+list(REMOVE_ITEM CORENRN_LINK_LIBS "Threads::Threads")
+
+# replicate CMake magic to transform system libs to -l<libname>
+foreach(link_lib ${CORENRN_LINK_LIBS})
+    if(${link_lib} MATCHES "\-l.*")
+        string(APPEND CORENRN_COMMON_LDFLAGS " ${link_lib}")
+        continue()
+    endif()
+    get_filename_component(path ${link_lib} DIRECTORY)
+    if(NOT path)
+        string(APPEND CORENRN_COMMON_LDFLAGS " -l${link_lib}")
+    elseif("${path}" MATCHES "^(/lib|/lib64|/usr/lib|/usr/lib64)$")
+        get_filename_component(libname ${link_lib} NAME_WE)
+        string(REGEX REPLACE "^lib" "" libname ${libname})
+        string(APPEND CORENRN_COMMON_LDFLAGS " -l${libname}")
+    else()
+        string(APPEND CORENRN_COMMON_LDFLAGS " ${link_lib}")
+    endif()
+endforeach()
+
+# =============================================================================
+# compile flags : common to all backend
+# =============================================================================
+# PGI compiler adds --c++14;-A option for C++14, remove ";"
+string(REPLACE ";" " " CXX14_STD_FLAGS "${CMAKE_CXX14_STANDARD_COMPILE_OPTION}")
+string(TOUPPER "${CMAKE_BUILD_TYPE}" _BUILD_TYPE)
+set(CORENRN_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${CMAKE_CXX_FLAGS_${_BUILD_TYPE}} ${CXX14_STD_FLAGS}")
+
+# =============================================================================
+# nmodl/mod2c related options : TODO
+# =============================================================================
+# name of nmodl/mod2c binary
+get_filename_component(nmodl_name ${CORENRN_MOD2CPP_BINARY} NAME)
+set(nmodl_binary_name ${nmodl_name})

CMake/OpenAccHelper.cmake

@@ -4,10 +4,14 @@
 # See top-level LICENSE file for details.
 # =============================================================================
 
+# =============================================================================
+# Prepare compiler flags for GPU target
+# =============================================================================
 if(CORENRN_ENABLE_GPU)
+
   # cuda unified memory support
   if(CORENRN_ENABLE_CUDA_UNIFIED_MEMORY)
-    add_definitions(-DUNIFIED_MEMORY)
+    set(UNIFIED_MEMORY_DEF -DUNIFIED_MEMORY)
   endif()
 
   # if user don't specify host compiler, use gcc from $PATH
@@ -20,44 +24,50 @@ if(CORENRN_ENABLE_GPU)
 
   # various flags for PGI compiler with GPU build
   if(${CMAKE_C_COMPILER_ID} STREQUAL "PGI")
+
     # workaround for old PGI version
-    add_definitions(-DPG_ACC_BUGS)
-    set(ACC_FLAGS "-acc")
+    set(PGI_ACC_FLAGS "-acc")
     # 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
+    # inlining of large functions for OpenACC
     set(PGI_INLINE_FLAGS "-Minline=size:200,levels:10")
-    # C/C++ compiler flags
-    set(CMAKE_C_FLAGS "${ACC_FLAGS} ${CMAKE_C_FLAGS}")
-    set(CMAKE_CXX_FLAGS "${ACC_FLAGS} ${CMAKE_CXX_FLAGS} ${PGI_DIAG_FLAGS}")
+
     # avoid PGI adding standard compliant "-A" flags
     set(CMAKE_CXX11_STANDARD_COMPILE_OPTION --c++11)
     set(CMAKE_CXX14_STANDARD_COMPILE_OPTION --c++14)
+
   else()
     message(FATAL_ERROR "GPU support is available via OpenACC using PGI/NVIDIA compilers."
                         " Use NVIDIA HPC SDK with -DCMAKE_C_COMPILER=nvc -DCMAKE_CXX_COMPILER=nvc++")
   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)
   endif()
   find_package(CUDA 9.0 REQUIRED)
   set(CUDA_SEPARABLE_COMPILATION ON)
   set(CUDA_PROPAGATE_HOST_FLAGS OFF)
-  add_definitions(-DCUDA_PROFILING)
-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(CUDA_PROFILING_DEF -DCUDA_PROFILING)
+
+  set(CORENRN_ACC_GPU_DEFS "${UNIFIED_MEMORY_DEF} ${CUDA_PROFILING_DEF}")
+  set(CORENRN_ACC_GPU_FLAGS "${PGI_ACC_FLAGS} ${PGI_DIAG_FLAGS} ${PGI_INLINE_FLAGS}")
+
+  add_definitions(${CORENRN_ACC_GPU_DEFS})
+  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${CORENRN_ACC_GPU_FLAGS}")
+endif()
+
+# =============================================================================
+# Set global property that will be used by NEURON to link with CoreNEURON
+# =============================================================================
+if(CORENRN_ENABLE_GPU)
+  set_property(
+    GLOBAL
+    PROPERTY
+      CORENEURON_LIB_LINK_FLAGS
+      "${PGI_ACC_FLAGS} -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}"
+  )
+else()
   set_property(GLOBAL PROPERTY CORENEURON_LIB_LINK_FLAGS
                                "-L${CMAKE_HOST_SYSTEM_PROCESSOR} -lcorenrnmech")
 endif(CORENRN_ENABLE_GPU)

CMakeLists.txt

@@ -95,13 +95,6 @@ set(LIKWID_DIR
     ""
     CACHE PATH "Path to likwid performance analysis suite")
 
-set(CORENRN_FRONTEND_C_COMPILER
-    gcc
-    CACHE FILEPATH "C compiler for building mod2c [frontend]")
-set(CORENRN_FRONTEND_CXX_COMPILER
-    g++
-    CACHE FILEPATH "C++ compiler for building mod2c [frontend]")
-
 if(CORENEURON_AS_SUBPROJECT)
   set(CORENRN_ENABLE_UNIT_TESTS OFF)
 endif()
@@ -126,11 +119,6 @@ include(OpenAccHelper)
 find_package(PythonInterp REQUIRED)
 find_package(Perl REQUIRED)
 
-# =============================================================================
-# ISPC should compile with --pic by default
-# =============================================================================
-set(CMAKE_ISPC_FLAGS "--pic ${CMAKE_ISPC_FLAGS}")
-
 # =============================================================================
 # Common build options
 # =============================================================================
@@ -147,7 +135,6 @@ endif()
 # Build option specific compiler flags
 # =============================================================================
 if(${CMAKE_CXX_COMPILER_ID} STREQUAL "PGI")
-  add_definitions(-DSWAP_ENDIAN_DISABLE_ASM)
   # PGI with llvm code generation doesn't have necessary assembly intrinsic headers
   add_definitions(-DEIGEN_DONT_VECTORIZE=1)
 endif()
@@ -168,7 +155,6 @@ endif()
 
 if(CORENRN_ENABLE_ISPC)
   enable_language(ISPC)
-  add_definitions("-DISPC_INTEROP=1")
   set(CORENRN_ENABLE_NMODL ON)
 endif()
 
@@ -283,7 +269,6 @@ if(CORENRN_ENABLE_NMODL)
   if(CORENRN_ENABLE_GPU)
     string(APPEND CORENRN_NMODL_FLAGS " acc --oacc")
   endif()
-  separate_arguments(NMODL_EXTRA_FLAGS_LIST UNIX_COMMAND "${CORENRN_NMODL_FLAGS}")
 else()
   include(AddMod2cSubmodule)
   set(CORENRN_MOD2CPP_BINARY ${CMAKE_BINARY_DIR}/bin/mod2c_core${CMAKE_EXECUTABLE_SUFFIX})
@@ -309,10 +294,17 @@ if(CORENRN_ENABLE_LIKWID_PROFILING)
   add_definitions("-DLIKWID_PERFMON")
 endif()
 
+# =============================================================================
+# Common CXX flags : ignore unknown pragma warnings
+# =============================================================================
+set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${IGNORE_UNKNOWN_PRAGMA_FLAGS}")
+
 # =============================================================================
 # Add main directories
 # =============================================================================
 add_subdirectory(coreneuron)
+
+include(MakefileBuildOptions)
 add_subdirectory(extra)
 
 if(CORENRN_ENABLE_UNIT_TESTS)
@@ -325,14 +317,6 @@ endif()
 install(FILES CMake/coreneuron-config.cmake DESTINATION share/cmake)
 install(EXPORT coreneuron DESTINATION share/cmake)
 
-# just for printing the compiler flags in the build status
-string(TOUPPER ${CMAKE_BUILD_TYPE} BUILD_TYPE_UPPER)
-if(BUILD_TYPE_UPPER MATCHES "CUSTOM")
-  set(COMPILER_FLAGS "${CMAKE_CXX_FLAGS}")
-else()
-  set(COMPILER_FLAGS "${CMAKE_CXX_FLAGS_${BUILD_TYPE_UPPER}}")
-endif()
-
 if(NOT CORENEURON_AS_SUBPROJECT)
   # =============================================================================
   # Setup Doxygen documentation
@@ -402,7 +386,7 @@ if(cmake_generator_tolower MATCHES "makefile")
 
   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 "COMPILE FLAGS       | ${CORENRN_CXX_FLAGS}")
   message(STATUS "Build Type          | ${COMPILE_LIBRARY_TYPE}")
   message(STATUS "MPI                 | ${CORENRN_ENABLE_MPI}")
   if(CORENRN_ENABLE_MPI)

coreneuron/CMakeLists.txt

@@ -56,11 +56,6 @@ set(NMODL_UNITS_FILE "${CMAKE_BINARY_DIR}/share/mod2c/nrnunits.lib")
 file(COPY ${CORENEURON_PROJECT_SOURCE_DIR}/coreneuron/mechanism/mech/modfile
      DESTINATION ${CMAKE_BINARY_DIR}/share)
 
-# eion.cpp depends on CORENRN_USE_LEGACY_UNITS
-set(LegacyFR_FILES mechanism/eion.cpp apps/main1.cpp io/global_vars.cpp)
-set_source_files_properties(${LegacyFR_FILES} PROPERTIES COMPILE_FLAGS
-                              "-DCORENRN_USE_LEGACY_UNITS=${CORENRN_USE_LEGACY_UNITS}")
-
 # =============================================================================
 # coreneuron GPU library
 # =============================================================================
@@ -85,12 +80,6 @@ if(CORENRN_ENABLE_GPU)
   set_source_files_properties(${OPENACC_EXCLUDED_FILES} PROPERTIES COMPILE_FLAGS
                               "-DDISABLE_OPENACC")
 
-  # TODO : only older PGI versions?
-  if(${CMAKE_C_COMPILER_ID} STREQUAL "PGI")
-    set_source_files_properties(${CMAKE_CURRENT_SOURCE_DIR}/scopmath_core/sparse_thread.c
-                                PROPERTIES COMPILE_FLAGS "-ta=tesla:nollvm")
-  endif()
-
   # compile cuda files for multiple architecture
   cuda_add_library(
     "cudacoreneuron" ${CORENEURON_CUDA_FILES}

coreneuron/apps/main1.cpp

@@ -77,10 +77,8 @@ bool corenrn_units_use_legacy() {
 
 void (*nrn2core_part2_clean_)();
 
-#ifdef ISPC_INTEROP
 // cf. utils/ispc_globals.c
 extern double ispc_celsius;
-#endif
 
 /**
  * If "export OMP_NUM_THREADS=n" is not set then omp by default sets
@@ -213,9 +211,9 @@ void nrn_init_and_load_data(int argc,
 
     corenrn_param.celsius = celsius;
 
-#ifdef ISPC_INTEROP
+    // for ispc backend
     ispc_celsius = celsius;
-#endif
+
     // create net_cvode instance
     mk_netcvode();
 

coreneuron/gpu/nrn_acc_manager.cpp

@@ -969,8 +969,7 @@ void nrn_ion_global_map_copyto_device() {
             (double**)acc_copyin(nrn_ion_global_map, sizeof(double*) * nrn_ion_global_map_size);
         for (int j = 0; j < nrn_ion_global_map_size; j++) {
             if (nrn_ion_global_map[j]) {
-                /* @todo: fix this constant size 3 :( */
-                double* d_mechmap = (double*)acc_copyin(nrn_ion_global_map[j], 3 * sizeof(double));
+                double* d_mechmap = (double*)acc_copyin(nrn_ion_global_map[j], ion_global_map_member_size * sizeof(double));
                 acc_memcpy_to_device(&(d_data[j]), &d_mechmap, sizeof(double*));
             }
         }

coreneuron/mechanism/eion.cpp

@@ -48,21 +48,12 @@ THE POSSIBILITY OF SUCH DAMAGE.
 #endif
 
 #if defined(_OPENACC)
-#if defined(PG_ACC_BUGS)
 #define _PRAGMA_FOR_INIT_ACC_LOOP_ \
     _Pragma(                       \
-        "acc parallel loop present(pd[0:_cntml_padded*5], ppd[0:1], nrn_ion_global_map[0:nrn_ion_global_map_size][0:3]) if(nt->compute_gpu)")
+        "acc parallel loop present(pd[0:_cntml_padded*5], ppd[0:1], nrn_ion_global_map[0:nrn_ion_global_map_size][0:ion_global_map_member_size]) if(nt->compute_gpu)")
 #define _PRAGMA_FOR_CUR_ACC_LOOP_ \
     _Pragma(                      \
-        "acc parallel loop present(pd[0:_cntml_padded*5], nrn_ion_global_map[0:nrn_ion_global_map_size][0:3]) if(nt->compute_gpu) async(stream_id)")
-#else
-#define _PRAGMA_FOR_INIT_ACC_LOOP_ \
-    _Pragma(                       \
-        "acc parallel loop present(pd[0:_cntml_padded*5], ppd[0:1], nrn_ion_global_map[0:nrn_ion_global_map_size]) if(nt->compute_gpu)")
-#define _PRAGMA_FOR_CUR_ACC_LOOP_ \
-    _Pragma(                      \
-        "acc parallel loop present(pd[0:_cntml_padded*5], nrn_ion_global_map[0:nrn_ion_global_map_size]) if(nt->compute_gpu) async(stream_id)")
-#endif
+        "acc parallel loop present(pd[0:_cntml_padded*5], nrn_ion_global_map[0:nrn_ion_global_map_size][0:ion_global_map_member_size]) if(nt->compute_gpu) async(stream_id)")
 #define _PRAGMA_FOR_SEC_ORDER_CUR_ACC_LOOP_ \
     _Pragma(                                \
         "acc parallel loop present(pd[0:_cntml_padded*5], ni[0:_cntml_actual], _vec_rhs[0:_nt->end]) if(_nt->compute_gpu) async(stream_id)")
@@ -74,6 +65,10 @@ THE POSSIBILITY OF SUCH DAMAGE.
 
 namespace coreneuron {
 
+// for each ion it refers to internal concentration, external concentration, and charge,
+const int ion_global_map_member_size = 3;
+
+
 #define nparm 5
 static const char* mechanism[] = {/*just a template*/
                                   "0", "na_ion", "ena", "nao", "nai", 0, "ina", "dina_dv_", 0, 0};
@@ -130,7 +125,7 @@ void ion_reg(const char* name, double valence) {
             }
             nrn_ion_global_map_size = mechtype + 1;
         }
-        nrn_ion_global_map[mechtype] = (double*)emalloc(3 * sizeof(double));
+        nrn_ion_global_map[mechtype] = (double*)emalloc(ion_global_map_member_size * sizeof(double));
 
         register_mech((const char**)mechanism, nrn_alloc_ion, nrn_cur_ion, (mod_f_t)0, (mod_f_t)0,
                       (mod_f_t)nrn_init_ion, -1, 1);

coreneuron/mechanism/mechanism.hpp

@@ -33,7 +33,9 @@ THE POSSIBILITY OF SUCH DAMAGE.
 #include "coreneuron/utils/memory.h"
 
 namespace coreneuron {
-#if PG_ACC_BUGS
+// OpenACC with PGI compiler has issue when union is used and hence use struct
+// \todo check if newer PGI versions has resolved this issue
+#if defined(_OPENACC)
 struct ThreadDatum {
     int i;
     double* pval;

coreneuron/mechanism/membfunc.hpp

@@ -86,6 +86,7 @@ struct BAMech {
 
 extern int nrn_ion_global_map_size;
 extern double** nrn_ion_global_map;
+extern const int ion_global_map_member_size;
 
 #define NRNPOINTER                                                            \
     4 /* added on to list of mechanism variables.These are                    \
@@ -94,6 +95,7 @@ pointers which connect variables  from other mechanisms via the _ppval array. \
 
 #define _AMBIGUOUS 5
 
+
 extern int nrn_get_mechtype(const char*);
 extern const char* nrn_get_mechname(int);  // slow. use memb_func[i].sym if posible
 extern int register_mech(const char** m,

coreneuron/mechanism/register_mech.cpp

@@ -40,12 +40,9 @@ THE POSSIBILITY OF SUCH DAMAGE.
 namespace coreneuron {
 int secondorder = 0;
 double t, dt, celsius;
-#if defined(PG_ACC_BUGS)
-// clang-format off
-    #pragma acc declare copyin(secondorder)
-    #pragma acc declare copyin(celsius)
-// clang-format on
-#endif
+// declare copyin required for correct initialization
+#pragma acc declare copyin(secondorder)
+#pragma acc declare copyin(celsius)
 int rev_dt;
 
 using Pfrv = void (*)();