Skip to content
This repository was archived by the owner on Mar 20, 2023. It is now read-only.
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
73 changes: 73 additions & 0 deletions CMake/MakefileBuildOptions.cmake
Original file line number Diff line number Diff line change
@@ -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})
50 changes: 30 additions & 20 deletions CMake/OpenAccHelper.cmake
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand All @@ -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)
32 changes: 8 additions & 24 deletions CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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()
Expand All @@ -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
# =============================================================================
Expand All @@ -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()
Expand All @@ -168,7 +155,6 @@ endif()

if(CORENRN_ENABLE_ISPC)
enable_language(ISPC)
add_definitions("-DISPC_INTEROP=1")
set(CORENRN_ENABLE_NMODL ON)
endif()

Expand Down Expand Up @@ -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})
Expand All @@ -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)
Expand All @@ -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
Expand Down Expand Up @@ -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)
Expand Down
11 changes: 0 additions & 11 deletions coreneuron/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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
# =============================================================================
Expand All @@ -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}
Expand Down
6 changes: 2 additions & 4 deletions coreneuron/apps/main1.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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();

Expand Down
3 changes: 1 addition & 2 deletions coreneuron/gpu/nrn_acc_manager.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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*));
}
}
Expand Down
19 changes: 7 additions & 12 deletions coreneuron/mechanism/eion.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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)")
Expand All @@ -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};
Expand Down Expand Up @@ -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);
Expand Down
4 changes: 3 additions & 1 deletion coreneuron/mechanism/mechanism.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 2 additions & 0 deletions coreneuron/mechanism/membfunc.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -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 \
Expand All @@ -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,
Expand Down
9 changes: 3 additions & 6 deletions coreneuron/mechanism/register_mech.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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 (*)();
Expand Down
Loading