search

BlueBrain / nmodl

Code Generation Framework For NEURON MODeling Language
https://bluebrain.github.io/nmodl/
Apache License 2.0
54 stars 15 forks source link

Testing OpenACC GPU backend #101

Closed pramodk closed 5 years ago

pramodk commented 5 years ago

NMODL has support for OpenACC GPU backend but this has never been tested. Test coreneuron ringtest using OpenACC backend with PGI compiler (similar to MOD2C).

st4rl3ss commented 5 years ago

These are the parameters I used to compile: 

cmake ..  -DCMAKE_C_FLAGS:STRING="-O2 -ta=tesla:cuda9.0" \
          -DCMAKE_CXX_FLAGS:STRING="-O2 -ta=tesla:cuda9.0" \
          -DCOMPILE_LIBRARY_TYPE=STATIC \
          -DCUDA_HOST_COMPILER=`which gcc` \
          -DCUDA_PROPAGATE_HOST_FLAGS=OFF \
          -DENABLE_SELECTIVE_GPU_PROFILING=ON \
          -DENABLE_OPENACC=ON \
      -DENABLE_NMODL=ON \
          -DNMODL_ROOT=$HOME/nmodl \
          -DNMODL_EXTRA_FLAGS="passes --verbatim-rename --inline sympy --analytic acc --oacc" `

Issues with Coreneuron and OpenACC backend:

PGCC-W-0277-Cannot inline function _ZN111_INTERNAL_89__gpfs_bbp_cscs_ch_home_bellotta_proj_CoreNeuron_build_nmodl_openacc_coreneuron_expsyn_cpp_367bc17110coreneuron9mem_allocEmmm - data type mismatch (/gpfs/bbp.cscs.ch/home/bellotta/proj/CoreNeuron/build_nmodl_openacc/coreneuron/expsyn.cpp: 156)
PGCC-W-0277-Cannot inline function _ZN111_INTERNAL_89__gpfs_bbp_cscs_ch_home_bellotta_proj_CoreNeuron_build_nmodl_openacc_coreneuron_expsyn_cpp_367bc17110coreneuron9mem_allocEmmm - data type mismatch (/gpfs/bbp.cscs.ch/home/bellotta/proj/CoreNeuron/build_nmodl_openacc/coreneuron/expsyn.cpp: 157)
PGCC-S-0155-Compiler failed to translate accelerator region (see -Minfo messages): Could not find allocated-variable index for symbol (/gpfs/bbp.cscs.ch/home/bellotta/proj/CoreNeuron/build_nmodl_openacc/coreneuron/expsyn.cpp: 322)
ohm314 commented 5 years ago

Could you elaborate a bit? cuda_runtime_api.h needs to be included in the generated code of the MOD file? And what do you mean by "error with weight"?

st4rl3ss commented 5 years ago

Sorry if it's cryptic but it was just meant to be a reminder for me and Pramod, and not in any way an exhaustive decription of the issues I encountered.

In any case, yes the cuda header needs to be included as it's required by at least one function in the code cudaMallocManaged()

Regarding the weights I will come back at you since I don't remember the actual issue, but there was a problem with the weights variable in the generated code. In particular it wasn't defined and we had to copy the definition from another part of the code to make the compilation progress.

pramodk commented 5 years ago

@st4rl3ss : In retrospective, bullet points with code snippets is good idea (see #125).