diff --git a/mechanisms/default/exp2syn.mod b/mechanisms/default/exp2syn.mod index 225b162a6da5e446c3b1fc0f4c9455f9e6afddf3..3be17ba2676b7fb548bd2a7797e98d7f88403fcf 100644 --- a/mechanisms/default/exp2syn.mod +++ b/mechanisms/default/exp2syn.mod @@ -28,8 +28,6 @@ INITIAL { BREAKPOINT { SOLVE state METHOD cnexp - LOCAL g - g = B - A i = (B - A)*(v - e) } diff --git a/modcc/printer/cprinter.cpp b/modcc/printer/cprinter.cpp index 02afcf13e0ab04d78dc59f5a4aea43fcf07fb854..8eb2fc2186ff30586789020f48e748e6b2767153 100644 --- a/modcc/printer/cprinter.cpp +++ b/modcc/printer/cprinter.cpp @@ -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. diff --git a/modcc/printer/cprinter.hpp b/modcc/printer/cprinter.hpp index 57f94d5957bd3f1d334fc83da6efe6b90c9574b9..872de3ec97860a5fc4ba58d593304bb7d25eae04 100644 --- a/modcc/printer/cprinter.hpp +++ b/modcc/printer/cprinter.hpp @@ -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) {} diff --git a/modcc/printer/gpuprinter.cpp b/modcc/printer/gpuprinter.cpp index 9696f383faf4347a8f043ac66daf2496fb65fab3..ac3fe69c05e526d6fd491c2f900e10556283a6f1 100644 --- a/modcc/printer/gpuprinter.cpp +++ b/modcc/printer/gpuprinter.cpp @@ -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 { diff --git a/modcc/printer/printerutil.cpp b/modcc/printer/printerutil.cpp index af94fdc8bd007a052d11d898aa2b522f3203c4fc..86a983a7688bc16a71c5079a7bb2a38c98e44e30 100644 --- a/modcc/printer/printerutil.cpp +++ b/modcc/printer/printerutil.cpp @@ -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"; diff --git a/modcc/printer/printerutil.hpp b/modcc/printer/printerutil.hpp index bf3e406c877c4f53db77757339c4708ce33a937a..5addde91b244df8116e42d9312dfd72a12034be3 100644 --- a/modcc/printer/printerutil.hpp +++ b/modcc/printer/printerutil.hpp @@ -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.