neuronsimulator/gpuhackathon

CoreNEURON openacc to openmp offloading explorations

gpuhackathon's Issues

Create hackathon_main branches

• NMODL: https://github.com/BlueBrain/nmodl/tree/hackathon_main
• CoreNEURON: https://github.com/BlueBrain/CoreNeuron/tree/hackathon_main (protected etc.)

Investigate present(...) clauses in OpenMP, can we preserve error checking?

In OpenACC then present(...) gave us a runtime error if the data were not in fact present.
Can we keep this in OpenMP, e.g. with something like map(present, alloc: ...)?

Check nmodl for other openacc pragmas to add into overview.md

Copy channel-benchmark input data and mechanism files to shared workspace on Perlmutter and Ascent

Model size should be chosen to fit on a single GPU without being too small, i.e. O(few GB).
We should generate for at least a couple of MPI ranks so we can test asynchronous offload multiple NrnThreads in one CoreNEURON process.

Address issues ignored by the reduced hackathon_main CI

By disabling various bits of CI for hackathon work, we are postponing some issues for later.
So far the known issues are:

Code formatting

cmake-format and clang-format will have to be fixed later.

MOD2C + GPU builds

It seems that passing -mp=gpu to nvc++ breaks compilation of the OpenACC code generated by the MOD2C DSL translator. Specifically with -mp=gpu then this is an error:

NVC++-S-1053-External and Static variables are not supported in acc routine - _ZN10coreneuron11nrn_threadsE (x86_64/corenrn/mod2c/exp2syn.cpp: 385)

while without -mp=gpu it is a warning:

NVC++-W-1053-External and Static variables are not supported in acc routine - _ZN10coreneuron11nrn_threadsE (x86_64/corenrn/mod2c/exp2syn.cpp: 385)

the C++ code in question does not contain any OpenMP directives, only OpenACC. A rough outline of the usage is:

extern NrnThread* nrn_threads; // this is in a header file

static void _net_receive_kernel(Point_process* pnt) {
  NrnThread* nt = nrn_threads + pnt->_tid; // <--
  // ...
}

void _net_buf_receive(NrnThread* _nt) {
  Point_process* _pnt = _nt->pntprocs;
  int _pnt_length = _nt->n_pntproc - _nrb->_pnt_offset;
  #pragma acc parallel loop present(_pnt[0:_pnt_length], nrn_threads[0:nrn_nthread])
  for(int _j = 0; ...) {
    _net_receive_kernel(_pnt + _j);
  }
}

(internal link)

Intel compiler builds without GPU support

Adding #pragma omp target to the CoreNEURON source code breaks our CPU-only builds with the Intel compiler. Note that we already have OpenMP enabled in these builds, because it is used for higher level (CPU thread) parallelism.

/gpfs/bbp.cscs.ch/ssd/apps/hpc/jenkins/deploy/externals/2021-01-06/linux-rhel7-x86_64/gcc-9.3.0/hpe-mpi-2.25.hmpt-l6dob7/bin/mpicxx -fPIC  -qopenmp  -Wno-unknown-pragmas -g  -O0   -shared -Wl,-soname,libcoreneuron.so -o ../lib/libcoreneuron.so CMakeFiles/coreneuron.dir/apps/corenrn_parameters.cpp.o CMakeFiles/coreneuron.dir/apps/main1.cpp.o CMakeFiles/coreneuron.dir/gpu/nrn_acc_manager.cpp.o CMakeFiles/coreneuron.dir/io/core2nrn_data_return.cpp.o CMakeFiles/coreneuron.dir/io/file_utils.cpp.o CMakeFiles/coreneuron.dir/io/global_vars.cpp.o CMakeFiles/coreneuron.dir/io/lfp.cpp.o CMakeFiles/coreneuron.dir/io/mech_report.cpp.o CMakeFiles/coreneuron.dir/io/mem_layout_util.cpp.o CMakeFiles/coreneuron.dir/io/mk_mech.cpp.o CMakeFiles/coreneuron.dir/io/nrn2core_data_init.cpp.o CMakeFiles/coreneuron.dir/io/nrn_checkpoint.cpp.o CMakeFiles/coreneuron.dir/io/nrn_filehandler.cpp.o CMakeFiles/coreneuron.dir/io/nrn_setup.cpp.o CMakeFiles/coreneuron.dir/io/output_spikes.cpp.o CMakeFiles/coreneuron.dir/io/phase1.cpp.o CMakeFiles/coreneuron.dir/io/phase2.cpp.o CMakeFiles/coreneuron.dir/io/prcellstate.cpp.o CMakeFiles/coreneuron.dir/io/reports/binary_report_handler.cpp.o CMakeFiles/coreneuron.dir/io/reports/nrnreport.cpp.o CMakeFiles/coreneuron.dir/io/reports/report_configuration_parser.cpp.o CMakeFiles/coreneuron.dir/io/reports/report_event.cpp.o CMakeFiles/coreneuron.dir/io/reports/report_handler.cpp.o CMakeFiles/coreneuron.dir/io/reports/sonata_report_handler.cpp.o CMakeFiles/coreneuron.dir/io/setup_fornetcon.cpp.o CMakeFiles/coreneuron.dir/mechanism/capac.cpp.o CMakeFiles/coreneuron.dir/mechanism/eion.cpp.o CMakeFiles/coreneuron.dir/mechanism/mech_mapping.cpp.o CMakeFiles/coreneuron.dir/mechanism/patternstim.cpp.o CMakeFiles/coreneuron.dir/mechanism/register_mech.cpp.o CMakeFiles/coreneuron.dir/network/cvodestb.cpp.o CMakeFiles/coreneuron.dir/network/multisend.cpp.o CMakeFiles/coreneuron.dir/network/multisend_setup.cpp.o CMakeFiles/coreneuron.dir/network/netcvode.cpp.o CMakeFiles/coreneuron.dir/network/netpar.cpp.o CMakeFiles/coreneuron.dir/network/partrans.cpp.o CMakeFiles/coreneuron.dir/network/partrans_setup.cpp.o CMakeFiles/coreneuron.dir/network/tqueue.cpp.o CMakeFiles/coreneuron.dir/permute/balance.cpp.o CMakeFiles/coreneuron.dir/permute/cellorder.cpp.o CMakeFiles/coreneuron.dir/permute/cellorder1.cpp.o CMakeFiles/coreneuron.dir/permute/cellorder2.cpp.o CMakeFiles/coreneuron.dir/permute/data_layout.cpp.o CMakeFiles/coreneuron.dir/permute/node_permute.cpp.o CMakeFiles/coreneuron.dir/sim/fadvance_core.cpp.o CMakeFiles/coreneuron.dir/sim/fast_imem.cpp.o CMakeFiles/coreneuron.dir/sim/finitialize.cpp.o CMakeFiles/coreneuron.dir/sim/multicore.cpp.o CMakeFiles/coreneuron.dir/sim/solve_core.cpp.o CMakeFiles/coreneuron.dir/sim/treeset_core.cpp.o CMakeFiles/coreneuron.dir/utils/ispc/globals.cpp.o CMakeFiles/coreneuron.dir/utils/ivocvect.cpp.o CMakeFiles/coreneuron.dir/utils/lpt.cpp.o CMakeFiles/coreneuron.dir/utils/memory.cpp.o CMakeFiles/coreneuron.dir/utils/memory_utils.cpp.o CMakeFiles/coreneuron.dir/utils/nrn_stats.cpp.o CMakeFiles/coreneuron.dir/utils/nrnoc_aux.cpp.o CMakeFiles/coreneuron.dir/utils/nrntimeout.cpp.o CMakeFiles/coreneuron.dir/utils/progressbar/progressbar.cpp.o CMakeFiles/coreneuron.dir/utils/randoms/nrnran123.cpp.o CMakeFiles/coreneuron.dir/utils/string_utils.cpp.o CMakeFiles/coreneuron.dir/utils/utils.cpp.o CMakeFiles/coreneuron.dir/utils/vrecord.cpp.o CMakeFiles/coreneuron.dir/config/config.cpp.o CMakeFiles/coreneuron.dir/mpi/core/nrnmpi_def_cinc.cpp.o CMakeFiles/coreneuron.dir/mpi/core/nrnmpi.cpp.o CMakeFiles/coreneuron.dir/mpi/core/nrnmpidec.cpp.o CMakeFiles/corenrn_mpi.dir/mpi/lib/mpispike.cpp.o CMakeFiles/corenrn_mpi.dir/mpi/lib/nrnmpi.cpp.o -Wl,-rpath,/gpfs/bbp.cscs.ch/ssd/apps/hpc/jenkins/deploy/libraries/2021-01-06/linux-rhel7-x86_64/intel-19.1.2.254/reportinglib-2.5.6-7eg3tq/lib64:/gpfs/bbp.cscs.ch/ssd/apps/hpc/jenkins/deploy/libraries/2021-01-06/linux-rhel7-x86_64/intel-19.1.2.254/libsonata-report-1.0.0.20211007-2gn7wj/lib:::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::::: /gpfs/bbp.cscs.ch/ssd/apps/hpc/jenkins/deploy/libraries/2021-01-06/linux-rhel7-x86_64/intel-19.1.2.254/reportinglib-2.5.6-7eg3tq/lib64/libreportinglib.so /gpfs/bbp.cscs.ch/ssd/apps/hpc/jenkins/deploy/libraries/2021-01-06/linux-rhel7-x86_64/intel-19.1.2.254/libsonata-report-1.0.0.20211007-2gn7wj/lib/libsonatareport.so 
ld: cannot find /gpfs/bbp.cscs.ch/ssd/apps/hpc/jenkins/deploy/compilers/2021-01-06/linux-rhel7-x86_64/gcc-4.8.5/intel-20.0.2-ilowey/compilers_and_libraries_2020.2.254/linux/compiler/lib/intel64_lin/i_ofldbegin_target.o: No such file or directory
ld: cannot find -lioffload_target

(internal link)