From faf19b152ff0ac9764fd79eecdb61550e7dda1c3 Mon Sep 17 00:00:00 2001 From: pramodk Date: Tue, 18 Aug 2020 21:11:12 +0200 Subject: [PATCH 1/6] Add SEClamp support in CoreNEURON - add svclmp.mod mod file from NEURON - update registeration callbacks - expose at_time in the header file --- coreneuron/CMakeLists.txt | 3 +- coreneuron/mechanism/mech/cfile/cabvars.h | 4 +- coreneuron/mechanism/mech/modfile/svclmp.mod | 107 +++++++++++++++++++ coreneuron/mechanism/membfunc.hpp | 1 + 4 files changed, 112 insertions(+), 3 deletions(-) create mode 100755 coreneuron/mechanism/mech/modfile/svclmp.mod diff --git a/coreneuron/CMakeLists.txt b/coreneuron/CMakeLists.txt index 810bc57f8..dd64ffa83 100644 --- a/coreneuron/CMakeLists.txt +++ b/coreneuron/CMakeLists.txt @@ -58,7 +58,8 @@ set(INBUILT_MOD_FILES netstim.mod passive.mod pattern.mod - stim.mod) + stim.mod + svclmp.mod) # Has helper macro for generating cpp file from mod2c include(NmodlHelper) diff --git a/coreneuron/mechanism/mech/cfile/cabvars.h b/coreneuron/mechanism/mech/cfile/cabvars.h index ce6a511ca..fa1ee2c05 100644 --- a/coreneuron/mechanism/mech/cfile/cabvars.h +++ b/coreneuron/mechanism/mech/cfile/cabvars.h @@ -32,7 +32,7 @@ extern void capacitance_reg(void), _passive_reg(void), #if EXTRACELLULAR extracell_reg_(void), #endif - _stim_reg(void), _hh_reg(void), _netstim_reg(void), _expsyn_reg(void), _exp2syn_reg(void); + _stim_reg(void), _hh_reg(void), _netstim_reg(void), _expsyn_reg(void), _exp2syn_reg(void), _svclmp_reg(void); static void (*mechanism[])(void) = {/* type will start at 3 */ capacitance_reg, _passive_reg, @@ -40,6 +40,6 @@ static void (*mechanism[])(void) = {/* type will start at 3 */ /* extracellular requires special handling and must be type 5 */ extracell_reg_, #endif - _stim_reg, _hh_reg, _expsyn_reg, _netstim_reg, _exp2syn_reg, 0}; + _stim_reg, _hh_reg, _expsyn_reg, _netstim_reg, _exp2syn_reg, _svclmp_reg, 0}; } // namespace coreneuron diff --git a/coreneuron/mechanism/mech/modfile/svclmp.mod b/coreneuron/mechanism/mech/modfile/svclmp.mod new file mode 100755 index 000000000..daa5de2d4 --- /dev/null +++ b/coreneuron/mechanism/mech/modfile/svclmp.mod @@ -0,0 +1,107 @@ +TITLE svclmp.mod +COMMENT +Single electrode Voltage clamp with three levels. +Clamp is on at time 0, and off at time +dur1+dur2+dur3. When clamp is off the injected current is 0. +The clamp levels are amp1, amp2, amp3. +i is the injected current, vc measures the control voltage) +Do not insert several instances of this model at the same location in order to +make level changes. That is equivalent to independent clamps and they will +have incompatible internal state values. +The electrical circuit for the clamp is exceedingly simple: +vc ---'\/\/`--- cell + rs + +Note that since this is an electrode current model v refers to the +internal potential which is equivalent to the membrane potential v when +there is no extracellular membrane mechanism present but is v+vext when +one is present. +Also since i is an electrode current, +positive values of i depolarize the cell. (Normally, positive membrane currents +are outward and thus hyperpolarize the cell) +ENDCOMMENT + +INDEPENDENT {t FROM 0 TO 1 WITH 1 (ms)} + +DEFINE NSTEP 3 + +NEURON { + POINT_PROCESS SEClamp + ELECTRODE_CURRENT i + RANGE dur1, amp1, dur2, amp2, dur3, amp3, rs, vc, i +} + +UNITS { + (nA) = (nanoamp) + (mV) = (millivolt) + (uS) = (microsiemens) +} + + +PARAMETER { + rs = 1 (megohm) <1e-9, 1e9> + dur1 (ms) amp1 (mV) + dur2 (ms) <0,1e9> amp2 (mV) + dur3 (ms) <0,1e9> amp3 (mV) +} + +ASSIGNED { + v (mV) : automatically v + vext when extracellular is present + i (nA) + vc (mV) + tc2 (ms) + tc3 (ms) + on +} + +INITIAL { + tc2 = dur1 + dur2 + tc3 = tc2 + dur3 + on = 0 +} + +BREAKPOINT { + SOLVE icur METHOD after_cvode + vstim() +} + +PROCEDURE icur() { + if (on) { + i = (vc - v)/rs + }else{ + i = 0 + } +} + +COMMENT +The SOLVE of icur() in the BREAKPOINT block is necessary to compute +i=(vc - v(t))/rs instead of i=(vc - v(t-dt))/rs +This is important for time varying vc because the actual i used in +the implicit method is equivalent to (vc - v(t)/rs due to the +calculation of di/dv from the BREAKPOINT block. +The reason this works is because the SOLVE statement in the BREAKPOINT block +is executed after the membrane potential is advanced. + +It is a shame that vstim has to be called twice but putting the call +in a SOLVE block would cause playing a Vector into vc to be off by one +time step. +ENDCOMMENT + +PROCEDURE vstim() { + on = 1 + if (dur1) {at_time(dur1)} + if (dur2) {at_time(tc2)} + if (dur3) {at_time(tc3)} + if (t < dur1) { + vc = amp1 + }else if (t < tc2) { + vc = amp2 + }else if (t < tc3) { + vc = amp3 + }else { + vc = 0 + on = 0 + } + icur() +} + diff --git a/coreneuron/mechanism/membfunc.hpp b/coreneuron/mechanism/membfunc.hpp index 36204bd15..10f9e1126 100644 --- a/coreneuron/mechanism/membfunc.hpp +++ b/coreneuron/mechanism/membfunc.hpp @@ -188,6 +188,7 @@ extern void artcell_net_move(void**, Point_process*, double); extern void nrn2ncs_outputevent(int netcon_output_index, double firetime); extern bool nrn_use_localgid_; extern void net_sem_from_gpu(int sendtype, int i_vdata, int, int ith, int ipnt, double, double); +extern int at_time(NrnThread*, double); // _OPENACC and/or NET_RECEIVE_BUFFERING extern void net_sem_from_gpu(int, int, int, int, int, double, double); From a6ff6296344304ca81569d9c4091feb9afe2cfe5 Mon Sep 17 00:00:00 2001 From: Ioannis Magkanaris Date: Mon, 5 Oct 2020 18:29:15 +0200 Subject: [PATCH 2/6] Fixed compilation issue for GPU with at_time and updated NMODL commit --- coreneuron/mechanism/membfunc.hpp | 1 + external/nmodl | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/coreneuron/mechanism/membfunc.hpp b/coreneuron/mechanism/membfunc.hpp index 10f9e1126..7ec078612 100644 --- a/coreneuron/mechanism/membfunc.hpp +++ b/coreneuron/mechanism/membfunc.hpp @@ -188,6 +188,7 @@ extern void artcell_net_move(void**, Point_process*, double); extern void nrn2ncs_outputevent(int netcon_output_index, double firetime); extern bool nrn_use_localgid_; extern void net_sem_from_gpu(int sendtype, int i_vdata, int, int ith, int ipnt, double, double); +#pragma acc routine seq extern int at_time(NrnThread*, double); // _OPENACC and/or NET_RECEIVE_BUFFERING diff --git a/external/nmodl b/external/nmodl index f90f671e0..528043dde 160000 --- a/external/nmodl +++ b/external/nmodl @@ -1 +1 @@ -Subproject commit f90f671e03fd537540de359e52e21852a335f2b1 +Subproject commit 528043dde011fac17e0e1d7222a9cf5371c0fcab From 301033f8a0a73238d3d980423900fc7a47ea42e8 Mon Sep 17 00:00:00 2001 From: Ioannis Magkanaris Date: Tue, 6 Oct 2020 18:48:07 +0200 Subject: [PATCH 3/6] Added definition of at_time in nrnoc_ml.ispc --- coreneuron/mechanism/nrnoc_ml.ispc | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/coreneuron/mechanism/nrnoc_ml.ispc b/coreneuron/mechanism/nrnoc_ml.ispc index a2956cb80..da8839086 100644 --- a/coreneuron/mechanism/nrnoc_ml.ispc +++ b/coreneuron/mechanism/nrnoc_ml.ispc @@ -166,4 +166,12 @@ struct NrnThread { void* mapping; }; -extern uniform double ispc_celsius; \ No newline at end of file +extern uniform double ispc_celsius; + +static inline int at_time(NrnThread* nt, double te) { + double x = te - 1e-11; + if (x <= nt->_t && x > (nt->_t - nt->_dt)) { + return 1; + } + return 0; +} From b9428c4b3ad7e341af888377258c4bb79c91b26b Mon Sep 17 00:00:00 2001 From: Ioannis Magkanaris Date: Fri, 9 Oct 2020 18:49:28 +0200 Subject: [PATCH 4/6] Second running version for definition of at_time in ispc --- coreneuron/mechanism/nrnoc_ml.ispc | 9 ++------- coreneuron/network/cvodestb.cpp | 5 +++-- 2 files changed, 5 insertions(+), 9 deletions(-) diff --git a/coreneuron/mechanism/nrnoc_ml.ispc b/coreneuron/mechanism/nrnoc_ml.ispc index da8839086..21fcaa73c 100644 --- a/coreneuron/mechanism/nrnoc_ml.ispc +++ b/coreneuron/mechanism/nrnoc_ml.ispc @@ -168,10 +168,5 @@ struct NrnThread { extern uniform double ispc_celsius; -static inline int at_time(NrnThread* nt, double te) { - double x = te - 1e-11; - if (x <= nt->_t && x > (nt->_t - nt->_dt)) { - return 1; - } - return 0; -} +extern "C" +unmasked int at_time(NrnThread* nt, varying double te); diff --git a/coreneuron/network/cvodestb.cpp b/coreneuron/network/cvodestb.cpp index d84515e58..372bbea89 100644 --- a/coreneuron/network/cvodestb.cpp +++ b/coreneuron/network/cvodestb.cpp @@ -105,11 +105,12 @@ void fixed_play_continuous(NrnThread* nt) { } } -int at_time(NrnThread* nt, double te) { +} // namespace coreneuron + +extern "C" int at_time(coreneuron::NrnThread* nt, double te) { double x = te - 1e-11; if (x <= nt->_t && x > (nt->_t - nt->_dt)) { return 1; } return 0; } -} // namespace coreneuron From 87c9c70047a21cba9e00f29b3316aeb9e4d368a2 Mon Sep 17 00:00:00 2001 From: Ioannis Magkanaris Date: Fri, 9 Oct 2020 19:33:46 +0200 Subject: [PATCH 5/6] Small fix for at_time declaration outside ispc --- coreneuron/mechanism/membfunc.hpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/coreneuron/mechanism/membfunc.hpp b/coreneuron/mechanism/membfunc.hpp index 7ec078612..f64b98853 100644 --- a/coreneuron/mechanism/membfunc.hpp +++ b/coreneuron/mechanism/membfunc.hpp @@ -188,8 +188,6 @@ extern void artcell_net_move(void**, Point_process*, double); extern void nrn2ncs_outputevent(int netcon_output_index, double firetime); extern bool nrn_use_localgid_; extern void net_sem_from_gpu(int sendtype, int i_vdata, int, int ith, int ipnt, double, double); -#pragma acc routine seq -extern int at_time(NrnThread*, double); // _OPENACC and/or NET_RECEIVE_BUFFERING extern void net_sem_from_gpu(int, int, int, int, int, double, double); @@ -198,3 +196,6 @@ extern void hoc_malchk(void); /* just a stub */ extern void* hoc_Emalloc(size_t); } // namespace coreneuron + +#pragma acc routine seq +extern "C" int at_time(coreneuron::NrnThread*, double); \ No newline at end of file From de938556367ed1da0561b58d47e7c63ca1e6c81e Mon Sep 17 00:00:00 2001 From: Ioannis Magkanaris Date: Fri, 9 Oct 2020 22:00:05 +0200 Subject: [PATCH 6/6] Removed unmasked from ispc declaration and added documentation for at_time --- coreneuron/mechanism/membfunc.hpp | 7 ++++++- coreneuron/mechanism/nrnoc_ml.ispc | 2 +- 2 files changed, 7 insertions(+), 2 deletions(-) diff --git a/coreneuron/mechanism/membfunc.hpp b/coreneuron/mechanism/membfunc.hpp index f64b98853..ea073e59b 100644 --- a/coreneuron/mechanism/membfunc.hpp +++ b/coreneuron/mechanism/membfunc.hpp @@ -197,5 +197,10 @@ extern void* hoc_Emalloc(size_t); } // namespace coreneuron +/* + * at_time function declared outside of coreneuron namespace + * because it can be also called from ISPC code. That's also + * the reason for extern "C". + */ #pragma acc routine seq -extern "C" int at_time(coreneuron::NrnThread*, double); \ No newline at end of file +extern "C" int at_time(coreneuron::NrnThread*, double); diff --git a/coreneuron/mechanism/nrnoc_ml.ispc b/coreneuron/mechanism/nrnoc_ml.ispc index 21fcaa73c..1036dc935 100644 --- a/coreneuron/mechanism/nrnoc_ml.ispc +++ b/coreneuron/mechanism/nrnoc_ml.ispc @@ -169,4 +169,4 @@ struct NrnThread { extern uniform double ispc_celsius; extern "C" -unmasked int at_time(NrnThread* nt, varying double te); +int at_time(NrnThread* nt, varying double te);