Skip to content
Snippets Groups Projects
Unverified Commit 8f525914 authored by Thorsten Hater's avatar Thorsten Hater Committed by GitHub
Browse files

Bug fix: point mechs applying weights to ionic concentrations (#1960) 

Until now, point mechanisms applied a weighting factor when updating internal/external concentrations.
This was an oversight and is only correct in density mechanisms. This has been rectified, but might change
point mechanism semantics for injecting/modifying Xi/Xo values.
parent 8494607e
No related branches found
No related tags found
No related merge requests found
Hide whitespace changes
Inline Side-by-side
mechanisms/default/exp2syn.mod
+ 0
− 2
View file @ 8f525914
......@@ -28,8 +28,6 @@ INITIAL {
BREAKPOINT {
SOLVE state METHOD cnexp
LOCAL g
g = B - A
i = (B - A)*(v - e)
}
......
modcc/printer/cprinter.cpp
+ 84
− 79
View file @ 8f525914
......@@ -52,8 +52,8 @@ struct index_prop {
void emit_procedure_proto(std::ostream&, ProcedureExpression*, const std::string&, const std::string& qualified = "");
void emit_simd_procedure_proto(std::ostream&, ProcedureExpression*, const std::string&, const std::string& qualified = "");
void emit_masked_simd_procedure_proto(std::ostream&, ProcedureExpression*, const std::string&, const std::string& qualified = "");
void emit_api_body(std::ostream&, APIMethod*, bool cv_loop = true, bool ppack_iface=true, bool use_additive=false);
void emit_simd_api_body(std::ostream&, APIMethod*, const std::vector<VariableExpression*>& scalars, bool use_additive);
void emit_api_body(std::ostream&, APIMethod*, const ApiFlags& flags={});
void emit_simd_api_body(std::ostream&, APIMethod*, const std::vector<VariableExpression*>& scalars, const ApiFlags&);
void emit_simd_index_initialize(std::ostream& out, const std::list<index_prop>& indices, simd_expr_constraint constraint);
void emit_simd_body_for_loop(std::ostream& out,
......@@ -61,14 +61,14 @@ void emit_simd_body_for_loop(std::ostream& out,
const std::vector<LocalVariable*>& indexed_vars,
const std::list<index_prop>& indices,
const simd_expr_constraint& constraint,
bool use_additive);
const ApiFlags&);
void emit_simd_for_loop_per_constraint(std::ostream& out, BlockExpression* body,
const std::vector<LocalVariable*>& indexed_vars,
const std::list<index_prop>& indices,
const simd_expr_constraint& constraint,
std::string constraint_name,
bool use_additive);
const ApiFlags&);
struct cprint {
Expression* expr_;
......@@ -233,10 +233,13 @@ ARB_LIBMODCC_API std::string emit_cpp_source(const Module& module_, const printe
// Make implementations
auto emit_body = [&](APIMethod *p, bool add=false) {
auto flags = ApiFlags{}
.additive(add)
.point(moduleKind::point == module_.kind());
if (with_simd) {
emit_simd_api_body(out, p, vars.scalars, add);
emit_simd_api_body(out, p, vars.scalars, flags);
} else {
emit_api_body(out, p, true, true, add);
emit_api_body(out, p, flags);
}
};
......@@ -344,7 +347,7 @@ ARB_LIBMODCC_API std::string emit_cpp_source(const Module& module_, const printe
pp_var_pfx,
net_receive_api->args().empty() ? "weight" : net_receive_api->args().front()->is_argument()->name());
out << indent << indent << indent << indent;
emit_api_body(out, net_receive_api, false, false, true);
emit_api_body(out, net_receive_api, net_recv_flags);
out << popindent << "}\n" << popindent << "}\n" << popindent << "}\n" << popindent << "}\n\n";
} else {
out << "static void apply_events(arb_mechanism_ppack*, arb_deliverable_event_stream*) {}\n\n";
......@@ -364,7 +367,7 @@ ARB_LIBMODCC_API std::string emit_cpp_source(const Module& module_, const printe
pp_var_pfx,
time_arg);
out << indent << indent << indent << indent;
emit_api_body(out, post_event_api, false, false);
emit_api_body(out, post_event_api, post_evt_flags);
out << popindent << "}\n" << popindent << "}\n" << popindent << "}\n" << popindent << "}\n";
} else {
out << "static void post_event(arb_mechanism_ppack*) {}\n";
......@@ -560,70 +563,65 @@ std::list<index_prop> gather_indexed_vars(const std::vector<LocalVariable*>& ind
return indices;
}
void emit_state_read(std::ostream& out, LocalVariable* local) {
ENTER(out);
out << "arb_value_type " << cprint(local) << " = ";
if (local->is_read() || (local->is_write() && decode_indexed_variable(local->external_variable()).additive)) {
auto d = decode_indexed_variable(local->external_variable());
out << scaled(d.scale) << deref(d) << ";\n";
}
else {
out << "0;\n";
}
EXIT(out);
}
void emit_state_update(std::ostream& out, Symbol* from, IndexedVariable* external, bool use_additive) {
if (!external->is_write()) return;
ENTER(out);
auto d = decode_indexed_variable(external);
if (d.readonly) throw compiler_exception("Cannot assign to read-only external state: "+external->to_string());
std::string var, weight = pp_var_pfx + "weight[i_]", scale = scaled(1.0/d.scale), name = from->name();
double coeff = 1.0/d.scale;
{
std::stringstream v, s, w;
v << deref(d); var = v.str();
}
if (d.additive && use_additive) {
out << fmt::format("{3} -= {0};\n"
"{0} = fma({1}{2}, {3}, {0});\n",
var, scale, weight, name);
}
else if (d.accumulate) {
out << deref(d) << " = fma("
<< scaled(coeff) << pp_var_pfx << "weight[i_], "
<< from->name() << ", " << deref(d) << ");\n";
}
else {
out << deref(d) << " = " << scaled(coeff) << from->name() << ";\n";
}
EXIT(out);
}
void emit_api_body(std::ostream& out, APIMethod* method, bool cv_loop, bool ppack_iface, bool use_additive) {
void emit_api_body(std::ostream& out, APIMethod* method, const ApiFlags& flags) {
ENTER(out);
auto body = method->body();
auto indexed_vars = indexed_locals(method->scope());
std::list<index_prop> indices = gather_indexed_vars(indexed_vars, "i_");
if (!body->statements().empty()) {
ppack_iface && out << "PPACK_IFACE_BLOCK;\n";
cv_loop && out << fmt::format("for (arb_size_type i_ = 0; i_ < {}width; ++i_) {{\n", pp_var_pfx)
<< indent;
if (flags.ppack_iface) out << "PPACK_IFACE_BLOCK;\n";
if (flags.cv_loop) {
out << fmt::format("for (arb_size_type i_ = 0; i_ < {}width; ++i_) {{\n",
pp_var_pfx)
<< indent;
}
for (auto index: indices) {
out << "auto " << source_index_i_name(index) << " = " << source_var(index) << "[" << index.index_name << "];\n";
out << fmt::format("auto {} = {}[{}];\n",
source_index_i_name(index),
source_var(index),
index.index_name);
}
for (auto& sym: indexed_vars) {
emit_state_read(out, sym);
auto d = decode_indexed_variable(sym->external_variable());
out << "arb_value_type " << cprint(sym) << " = ";
if (sym->is_read() || (sym->is_write() && d.additive)) {
out << scaled(d.scale) << deref(d) << ";\n";
}
else {
out << "0;\n";
}
}
out << cprint(body);
for (auto& sym: indexed_vars) {
emit_state_update(out, sym, sym->external_variable(), use_additive);
if (!sym->external_variable()->is_write()) continue;
auto d = decode_indexed_variable(sym->external_variable());
bool use_weight = d.always_use_weight || !flags.is_point;
if (d.readonly) throw compiler_exception("Cannot assign to read-only external state: "+sym->to_string());
std::string
var,
weight = use_weight ? pp_var_pfx + "weight[i_]" : "1.0",
scale = scaled(1.0/d.scale),
name = sym->name();
{
std::stringstream v; v << deref(d); var = v.str();
}
if (d.additive && flags.use_additive) {
out << fmt::format("{3} -= {0};\n"
"{0} = fma({1}{2}, {3}, {0});\n",
var, scale, weight, name);
}
else if (d.accumulate) {
out << fmt::format("{} = fma({}{}, {}, {});\n",
var, scale, weight, name, var);
}
else {
out << var << " = " << scale << name << ";\n";
}
}
cv_loop && out << popindent << "}\n";
if (flags.cv_loop) out << popindent << "}\n";
}
EXIT(out);
}
......@@ -820,11 +818,15 @@ void emit_simd_state_read(std::ostream& out, LocalVariable* local, simd_expr_con
EXIT(out);
}
void emit_simd_state_update(std::ostream& out, Symbol* from, IndexedVariable* external, simd_expr_constraint constraint, bool use_additive) {
void emit_simd_state_update(std::ostream& out,
Symbol* from, IndexedVariable* external,
simd_expr_constraint constraint,
const ApiFlags& flags) {
if (!external->is_write()) return;
ENTER(out);
auto d = decode_indexed_variable(external);
if (d.readonly) {
throw compiler_exception("Cannot assign to read-only external state: "+external->to_string());
}
......@@ -842,11 +844,13 @@ void emit_simd_state_update(std::ostream& out, Symbol* from, IndexedVariable* ex
scaled = ss.str();
}
if (d.additive && use_additive) {
std::string weight = (d.always_use_weight || !flags.is_point) ? "w_" : "simd_cast<simd_value>(1.0)";
if (d.additive && flags.use_additive) {
if (d.index_var_kind == index_kind::node) {
if (constraint == simd_expr_constraint::contiguous) {
out << fmt::format("indirect({} + {}, simd_width_) = S::mul(w_, {});\n",
data, node, scaled);
out << fmt::format("indirect({} + {}, simd_width_) = S::mul({}, {});\n",
data, node, weight, scaled);
}
else {
// We need this instead of simple assignment!
......@@ -854,17 +858,17 @@ void emit_simd_state_update(std::ostream& out, Symbol* from, IndexedVariable* ex
" simd_value t_{}0_ = simd_cast<simd_value>(0.0);\n"
" assign(t_{}0_, indirect({}, simd_cast<simd_index>({}), simd_width_, constraint_category_));\n"
" {} -= t_{}0_;\n"
" indirect({}, simd_cast<simd_index>({}), simd_width_, constraint_category_) += S::mul(w_, {});\n"
" indirect({}, simd_cast<simd_index>({}), simd_width_, constraint_category_) += S::mul({}, {});\n"
"}}\n",
name,
name, data, node,
scaled, name,
data, node, scaled);
data, node, weight, scaled);
}
}
else {
out << fmt::format("indirect({}, {}, simd_width_, index_constraint::none) = S::mul(w_, {});\n",
data, index, scaled);
out << fmt::format("indirect({}, {}, simd_width_, index_constraint::none) = S::mul({}, {});\n",
data, index, weight, scaled);
}
}
else if (d.accumulate) {
......@@ -874,17 +878,17 @@ void emit_simd_state_update(std::ostream& out, Symbol* from, IndexedVariable* ex
case simd_expr_constraint::contiguous:
out << "simd_value " << tempvar << ";\n"
<< "assign(" << tempvar << ", indirect(" << data << " + " << node << ", simd_width_));\n"
<< tempvar << " = S::fma(w_, " << scaled << ", " << tempvar << ");\n"
<< tempvar << " = S::fma(" << weight << ", " << scaled << ", " << tempvar << ");\n"
<< "indirect(" << data << " + " << node << ", simd_width_) = " << tempvar << ";\n";
break;
case simd_expr_constraint::constant:
out << "indirect(" << data << ", simd_cast<simd_index>(" << node << "), simd_width_, constraint_category_) += S::mul(w_, " << scaled << ");\n";
out << "indirect(" << data << ", simd_cast<simd_index>(" << node << "), simd_width_, constraint_category_) += S::mul(" << weight << ", " << scaled << ");\n";
break;
default:
out << "indirect(" << data << ", " << node << ", simd_width_, constraint_category_) += S::mul(w_, " << scaled << ");\n";
out << "indirect(" << data << ", " << node << ", simd_width_, constraint_category_) += S::mul(" << weight << ", " << scaled << ");\n";
}
} else {
out << "indirect(" << data << ", " << index << ", simd_width_, index_constraint::none) += S::mul(w_, " << scaled << ");\n";
out << "indirect(" << data << ", " << index << ", simd_width_, index_constraint::none) += S::mul(" << weight << ", " << scaled << ");\n";
}
}
else if (d.index_var_kind == index_kind::node) {
......@@ -958,7 +962,7 @@ void emit_simd_body_for_loop(
const std::vector<VariableExpression*>& scalars,
const std::list<index_prop>& indices,
const simd_expr_constraint& constraint,
bool use_additive) {
const ApiFlags& flags) {
ENTER(out);
emit_simd_index_initialize(out, indices, constraint);
......@@ -972,7 +976,7 @@ void emit_simd_body_for_loop(
out << printer;
for (auto& sym: indexed_vars) {
emit_simd_state_update(out, sym, sym->external_variable(), constraint, use_additive);
emit_simd_state_update(out, sym, sym->external_variable(), constraint, flags);
}
EXIT(out);
}
......@@ -983,7 +987,7 @@ void emit_simd_for_loop_per_constraint(std::ostream& out, BlockExpression* body,
const std::list<index_prop>& indices,
const simd_expr_constraint& constraint,
std::string underlying_constraint_name,
bool use_additive) {
const ApiFlags& flags) {
ENTER(out);
out << fmt::format("constraint_category_ = index_constraint::{1};\n"
"for (auto i_ = 0ul; i_ < {0}index_constraints.n_{1}; i_++) {{\n"
......@@ -995,13 +999,15 @@ void emit_simd_for_loop_per_constraint(std::ostream& out, BlockExpression* body,
"assign(w_, indirect(({}weight+index_), simd_width_));\n",
pp_var_pfx);
emit_simd_body_for_loop(out, body, indexed_vars, scalars, indices, constraint, use_additive);
emit_simd_body_for_loop(out, body, indexed_vars, scalars, indices, constraint, flags);
out << popindent << "}\n";
EXIT(out);
}
void emit_simd_api_body(std::ostream& out, APIMethod* method, const std::vector<VariableExpression*>& scalars, bool use_additive) {
void emit_simd_api_body(std::ostream& out, APIMethod* method,
const std::vector<VariableExpression*>& scalars,
const ApiFlags& flags) {
auto body = method->body();
auto indexed_vars = indexed_locals(method->scope());
......@@ -1023,26 +1029,25 @@ void emit_simd_api_body(std::ostream& out, APIMethod* method, const std::vector<
simd_expr_constraint constraint = simd_expr_constraint::contiguous;
std::string underlying_constraint = "contiguous";
emit_simd_for_loop_per_constraint(out, body, indexed_vars, scalars, indices, constraint, underlying_constraint, use_additive);
emit_simd_for_loop_per_constraint(out, body, indexed_vars, scalars, indices, constraint, underlying_constraint, flags);
//Generate for loop for all independent simd_vectors
constraint = simd_expr_constraint::other;
underlying_constraint = "independent";
emit_simd_for_loop_per_constraint(out, body, indexed_vars, scalars, indices, constraint, underlying_constraint, use_additive);
emit_simd_for_loop_per_constraint(out, body, indexed_vars, scalars, indices, constraint, underlying_constraint, flags);
//Generate for loop for all simd_vectors that have no optimizing constraints
constraint = simd_expr_constraint::other;
underlying_constraint = "none";
emit_simd_for_loop_per_constraint(out, body, indexed_vars, scalars, indices, constraint, underlying_constraint, use_additive);
emit_simd_for_loop_per_constraint(out, body, indexed_vars, scalars, indices, constraint, underlying_constraint, flags);
//Generate for loop for all constant simd_vectors
constraint = simd_expr_constraint::constant;
underlying_constraint = "constant";
emit_simd_for_loop_per_constraint(out, body, indexed_vars, scalars, indices, constraint, underlying_constraint, use_additive);
emit_simd_for_loop_per_constraint(out, body, indexed_vars, scalars, indices, constraint, underlying_constraint, flags);
}
else {
// We may nonetheless need to read a global scalar indexed variable.
......
modcc/printer/cprinter.hpp
+ 15
− 0
View file @ 8f525914
......@@ -45,6 +45,21 @@ enum class simd_expr_constraint{
other
};
struct ApiFlags {
bool cv_loop = true;
bool ppack_iface=true;
bool use_additive=false;
bool is_point=false;
ApiFlags& loop(bool v) { cv_loop = v; return *this; }
ApiFlags& iface(bool v) { ppack_iface = v; return *this; }
ApiFlags& additive(bool v) { use_additive = v; return *this; }
ApiFlags& point(bool v) { is_point = v; return *this; }
};
const ApiFlags net_recv_flags = {false, false, true, false};
const ApiFlags post_evt_flags = {false, false, false, false};
class ARB_LIBMODCC_API SimdPrinter: public Visitor {
public:
SimdPrinter(std::ostream& out): out_(out) {}
......
modcc/printer/gpuprinter.cpp
+ 19
− 19
View file @ 8f525914
......@@ -28,10 +28,10 @@ static std::string scaled(double coeff) {
}
void emit_api_body_cu(std::ostream& out, APIMethod* method, bool is_point_proc, bool cv_loop = true, bool ppack=true, bool additive=false);
void emit_api_body_cu(std::ostream& out, APIMethod* method, const ApiFlags&);
void emit_procedure_body_cu(std::ostream& out, ProcedureExpression* proc);
void emit_state_read_cu(std::ostream& out, LocalVariable* local);
void emit_state_update_cu(std::ostream& out, Symbol* from, IndexedVariable* external, bool is_point_proc, bool use_additive);
void emit_state_update_cu(std::ostream& out, Symbol* from, IndexedVariable* external, const ApiFlags&);
const char* index_id(Symbol *s);
......@@ -226,7 +226,7 @@ ARB_LIBMODCC_API std::string emit_gpu_cu_source(const Module& module_, const pri
<< "void " << e->name() << "(arb_mechanism_ppack params_) {\n" << indent
<< "int n_ = params_.width;\n"
<< "int tid_ = threadIdx.x + blockDim.x*blockIdx.x;\n";
emit_api_body_cu(out, e, is_point_proc, true, true, additive);
emit_api_body_cu(out, e, ApiFlags{}.point(is_point_proc).additive(additive));
out << popindent << "}\n\n";
}
};
......@@ -264,7 +264,7 @@ ARB_LIBMODCC_API std::string emit_gpu_cu_source(const Module& module_, const pri
net_receive_api->args().empty() ? "weight" : net_receive_api->args().front()->is_argument()->name(),
pp_var_pfx);
out << indent << indent << indent << indent;
emit_api_body_cu(out, net_receive_api, is_point_proc, false, false, false);
emit_api_body_cu(out, net_receive_api, ApiFlags{}.point(is_point_proc).loop(false).iface(false));
out << popindent << "}\n" << popindent << "}\n" << popindent << "}\n" << popindent << "}\n";
}
......@@ -285,7 +285,7 @@ ARB_LIBMODCC_API std::string emit_gpu_cu_source(const Module& module_, const pri
time_arg,
pp_var_pfx);
out << indent << indent << indent << indent;
emit_api_body_cu(out, post_event_api, is_point_proc, false, false);
emit_api_body_cu(out, post_event_api, ApiFlags{}.point(is_point_proc).loop(false).iface(false));
out << popindent << "}\n" << popindent << "}\n" << popindent << "}\n" << popindent << "}\n";
}
......@@ -365,7 +365,7 @@ static std::string index_i_name(const std::string& index_var) {
return index_var+"i_";
}
void emit_api_body_cu(std::ostream& out, APIMethod* e, bool is_point_proc, bool cv_loop, bool ppack, bool additive) {
void emit_api_body_cu(std::ostream& out, APIMethod* e, const ApiFlags& flags) {
auto body = e->body();
auto indexed_vars = indexed_locals(e->scope());
......@@ -418,7 +418,7 @@ void emit_api_body_cu(std::ostream& out, APIMethod* e, bool is_point_proc, bool
}
if (!body->statements().empty()) {
if (is_point_proc) {
if (flags.is_point) {
// The run length information is only required if this method will
// update an indexed variable, like current or conductance.
// This is the case if one of the external variables "is_write".
......@@ -428,8 +428,8 @@ void emit_api_body_cu(std::ostream& out, APIMethod* e, bool is_point_proc, bool
out << "unsigned lane_mask_ = arb::gpu::ballot(0xffffffff, tid_<n_);\n";
}
}
ppack && out << "PPACK_IFACE_BLOCK;\n";
cv_loop && out << "if (tid_<n_) {\n" << indent;
if (flags.ppack_iface) out << "PPACK_IFACE_BLOCK;\n";
if (flags.cv_loop) out << "if (tid_<n_) {\n" << indent;
for (auto& index: indices) {
out << "auto " << index_i_name(index.source_var)
......@@ -443,9 +443,9 @@ void emit_api_body_cu(std::ostream& out, APIMethod* e, bool is_point_proc, bool
out << cuprint(body);
for (auto& sym: indexed_vars) {
emit_state_update_cu(out, sym, sym->external_variable(), is_point_proc, additive);
emit_state_update_cu(out, sym, sym->external_variable(), flags);
}
cv_loop && out << popindent << "}\n";
if (flags.cv_loop) out << popindent << "}\n";
}
}
......@@ -470,9 +470,8 @@ namespace {
void emit_state_read_cu(std::ostream& out, LocalVariable* local) {
out << "arb_value_type " << cuprint(local) << " = ";
if (local->is_read() || (local->is_write() && decode_indexed_variable(local->external_variable()).additive)) {
auto d = decode_indexed_variable(local->external_variable());
auto d = decode_indexed_variable(local->external_variable());
if (local->is_read() || (local->is_write() && d.additive)) {
if (d.scale != 1) {
out << as_c_double(d.scale) << "*";
}
......@@ -485,7 +484,7 @@ void emit_state_read_cu(std::ostream& out, LocalVariable* local) {
void emit_state_update_cu(std::ostream& out, Symbol* from,
IndexedVariable* external, bool is_point_proc, bool use_additive) {
IndexedVariable* external, const ApiFlags& flags) {
if (!external->is_write()) return;
auto d = decode_indexed_variable(external);
if (d.readonly) {
......@@ -497,11 +496,12 @@ void emit_state_update_cu(std::ostream& out, Symbol* from,
auto data = pp_var_pfx + d.data_var;
auto index = index_i_name(d.outer_index_var());
auto var = deref(d);
auto weight = scale + pp_var_pfx + "weight[tid_]";
std::string weight = (d.always_use_weight || !flags.is_point) ? pp_var_pfx + "weight[tid_]" : "1.0";
weight = scale + weight;
if (d.additive && use_additive) {
if (d.additive && flags.use_additive) {
out << name << " -= " << var << ";\n";
if (is_point_proc) {
if (flags.is_point) {
out << fmt::format("::arb::gpu::reduce_by_key({}*{}, {}, {}, lane_mask_);\n", weight, name, data, index);
}
else {
......@@ -509,7 +509,7 @@ void emit_state_update_cu(std::ostream& out, Symbol* from,
}
}
else if (d.accumulate) {
if (is_point_proc) {
if (flags.is_point) {
out << "::arb::gpu::reduce_by_key(" << weight << "*" << name << ',' << data << ", " << index << ", lane_mask_);\n";
}
else {
......
modcc/printer/printerutil.cpp
+ 3
− 0
View file @ 8f525914
......@@ -143,6 +143,7 @@ ARB_LIBMODCC_API indexed_variable_info decode_indexed_variable(IndexedVariable*
v.accumulate = true;
v.additive = false;
v.readonly = true;
v.always_use_weight = true;
std::string ion_pfx;
if (sym->is_ion()) {
......@@ -209,6 +210,7 @@ ARB_LIBMODCC_API indexed_variable_info decode_indexed_variable(IndexedVariable*
case sourceKind::ion_iconc:
v.data_var = ion_pfx+".internal_concentration";
v.readonly = false;
v.always_use_weight = false;
break;
case sourceKind::ion_diffusive:
v.data_var = ion_pfx+".diffusive_concentration";
......@@ -219,6 +221,7 @@ ARB_LIBMODCC_API indexed_variable_info decode_indexed_variable(IndexedVariable*
case sourceKind::ion_econc:
v.data_var = ion_pfx+".external_concentration";
v.readonly = false;
v.always_use_weight = false;
break;
case sourceKind::ion_valence:
v.data_var = ion_pfx+".ionic_charge";
......
modcc/printer/printerutil.hpp
+ 1
− 0
View file @ 8f525914
......@@ -150,6 +150,7 @@ struct ARB_LIBMODCC_API indexed_variable_info {
bool accumulate = true; // true => add with weight_ factor on assignment
bool readonly = false; // true => can never be assigned to by a mechanism
bool additive = false; // only additive contributions allowed?
bool always_use_weight = false; // can disable weighting?
// Scale is the conversion factor from the data variable
// to the NMODL value.
......
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment