From fa0d7aef8ec7469fd48cd90135790d4ca4af7297 Mon Sep 17 00:00:00 2001
From: Sam Yates <yates@cscs.ch>
Date: Mon, 20 Aug 2018 10:57:22 +0200
Subject: [PATCH] Global temperature for NMODL mechanisms. (#565)
Global temperature for mechanisms.
* Make 'celsius' magic in modcc: now an indexed variable.
* Add a new temperature data source for indexed variables.
* Add support to printers for indexed variables that reference a scalar.
* Check that indexed variables aren't used in PROCEDURE blocks (this is a problem not just for 'celsius').
* Modify built-in mod files to pass celsius as a parameter to rates() procedures.
* Add global temperature to shared_state classes, and initialize through backend mechanism superclasses.
* Add some infrastructure for unit-test only mechanisms.
* Set modcc flags globally in top level CMakeLists.txt.
* Add test mechanism/module for checking celsius setting.
* Add unit test for multicore and gpu mechanism celsius setting.
* Make common mechanism private field data access helper for unit tests.
* Use helper in temperature, synapses tests.
* Fix warning in `distribued_context.hpp` about errant semicolon.
* Fix global scalar ref for SIMD printing.
* Use correct ARB_CXXOPT_ARCH instead of incorrect CXXOPT_ARCH in various CMakeLists.txt files.
* Add special case for no-non scalar indexed variables in API loop in SIMD printing.
Fixes #386
---
CMakeLists.txt | 7 ++
arbor/backends/gpu/mechanism.cpp | 2 +
arbor/backends/gpu/mechanism_ppack_base.hpp | 1 +
arbor/backends/gpu/shared_state.cpp | 2 +
arbor/backends/gpu/shared_state.hpp | 1 +
arbor/backends/multicore/mechanism.cpp | 2 +
arbor/backends/multicore/mechanism.hpp | 3 +-
arbor/backends/multicore/shared_state.cpp | 3 +
arbor/backends/multicore/shared_state.hpp | 1 +
aux/CMakeLists.txt | 1 +
include/arbor/distributed_context.hpp | 2 +-
mechanisms/BuildModules.cmake | 9 +++
mechanisms/CMakeLists.txt | 10 +--
mechanisms/mod/hh.mod | 8 +--
mechanisms/mod/kamt.mod | 6 +-
mechanisms/mod/kdrmt.mod | 6 +-
mechanisms/mod/nax.mod | 6 +-
modcc/expression.cpp | 61 ++++++++++++----
modcc/expression.hpp | 3 +
modcc/identifier.hpp | 1 +
modcc/module.cpp | 30 ++++++--
modcc/printer/cprinter.cpp | 80 +++++++++++++++------
modcc/printer/cudaprinter.cpp | 16 ++++-
modcc/printer/printerutil.cpp | 4 ++
modcc/printer/printerutil.hpp | 1 +
modcc/scope.hpp | 9 +++
modcc/solvers.hpp | 3 +
test/ubench/CMakeLists.txt | 1 +
test/unit-distributed/CMakeLists.txt | 4 +-
test/unit/CMakeLists.txt | 48 +++++++++++--
test/unit/mech_private_field_access.cpp | 64 +++++++++++++++++
test/unit/mech_private_field_access.hpp | 16 +++++
test/unit/mod/celsius_test.mod | 27 +++++++
test/unit/test_mech_temperature.cpp | 80 +++++++++++++++++++++
test/unit/test_synapses.cpp | 28 +++-----
test/unit/unit_test_catalogue.cpp | 28 ++++++++
test/unit/unit_test_catalogue.hpp | 5 ++
test/validation/CMakeLists.txt | 1 +
38 files changed, 485 insertions(+), 95 deletions(-)
create mode 100644 test/unit/mech_private_field_access.cpp
create mode 100644 test/unit/mech_private_field_access.hpp
create mode 100644 test/unit/mod/celsius_test.mod
create mode 100644 test/unit/test_mech_temperature.cpp
create mode 100644 test/unit/unit_test_catalogue.cpp
create mode 100644 test/unit/unit_test_catalogue.hpp
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 6ff82e3a..b4fc22b8 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -218,6 +218,13 @@ else()
set(ARB_WITH_EXTERNAL_MODCC FALSE)
endif()
+set(ARB_MODCC_FLAGS)
+if(ARB_VECTORIZE)
+ list(APPEND ARB_MODCC_FLAGS "--simd")
+endif()
+if(ARB_WITH_PROFILING)
+ list(APPEND ARB_MODCC_FLAGS "--profile")
+endif()
#----------------------------------------------------------
# Configure targets in sub-directories.
diff --git a/arbor/backends/gpu/mechanism.cpp b/arbor/backends/gpu/mechanism.cpp
index e545e791..9a70b562 100644
--- a/arbor/backends/gpu/mechanism.cpp
+++ b/arbor/backends/gpu/mechanism.cpp
@@ -76,6 +76,8 @@ void mechanism::instantiate(unsigned id,
pp->vec_v_ = shared.voltage.data();
pp->vec_i_ = shared.current_density.data();
+ pp->temperature_degC_ = shared.temperature_degC.data();
+
auto ion_state_tbl = ion_state_table();
num_ions_ = ion_state_tbl.size();
diff --git a/arbor/backends/gpu/mechanism_ppack_base.hpp b/arbor/backends/gpu/mechanism_ppack_base.hpp
index a867f69d..fb256579 100644
--- a/arbor/backends/gpu/mechanism_ppack_base.hpp
+++ b/arbor/backends/gpu/mechanism_ppack_base.hpp
@@ -35,6 +35,7 @@ struct mechanism_ppack_base {
const value_type* vec_dt_;
const value_type* vec_v_;
value_type* vec_i_;
+ const value_type* temperature_degC_;
const index_type* node_index_;
const value_type* weight_;
diff --git a/arbor/backends/gpu/shared_state.cpp b/arbor/backends/gpu/shared_state.cpp
index 9ab08608..5cd8972c 100644
--- a/arbor/backends/gpu/shared_state.cpp
+++ b/arbor/backends/gpu/shared_state.cpp
@@ -120,6 +120,7 @@ shared_state::shared_state(
dt_cv(n_cv),
voltage(n_cv),
current_density(n_cv),
+ temperature_degC(1),
deliverable_events(n_cell)
{}
@@ -139,6 +140,7 @@ void shared_state::reset(fvm_value_type initial_voltage, fvm_value_type temperat
memory::fill(current_density, 0);
memory::fill(time, 0);
memory::fill(time_to, 0);
+ memory::fill(temperature_degC, temperature_K - 273.15);
for (auto& i: ion_data) {
i.second.reset(temperature_K);
diff --git a/arbor/backends/gpu/shared_state.hpp b/arbor/backends/gpu/shared_state.hpp
index da1ae4ef..81174715 100644
--- a/arbor/backends/gpu/shared_state.hpp
+++ b/arbor/backends/gpu/shared_state.hpp
@@ -75,6 +75,7 @@ struct shared_state {
array dt_cv; // Maps CV index to dt [ms].
array voltage; // Maps CV index to membrane voltage [mV].
array current_density; // Maps CV index to current density [A/m²].
+ array temperature_degC; // Global temperature [°C] (length 1 array).
std::unordered_map<ionKind, ion_state> ion_data;
diff --git a/arbor/backends/multicore/mechanism.cpp b/arbor/backends/multicore/mechanism.cpp
index 81c48e8f..63bd3449 100644
--- a/arbor/backends/multicore/mechanism.cpp
+++ b/arbor/backends/multicore/mechanism.cpp
@@ -78,6 +78,8 @@ void mechanism::instantiate(unsigned id, backend::shared_state& shared, const la
vec_v_ = shared.voltage.data();
vec_i_ = shared.current_density.data();
+ temperature_degC_ = &shared.temperature_degC;
+
auto ion_state_tbl = ion_state_table();
n_ion_ = ion_state_tbl.size();
for (auto i: ion_state_tbl) {
diff --git a/arbor/backends/multicore/mechanism.hpp b/arbor/backends/multicore/mechanism.hpp
index 8724d8aa..af08a931 100644
--- a/arbor/backends/multicore/mechanism.hpp
+++ b/arbor/backends/multicore/mechanism.hpp
@@ -72,12 +72,13 @@ protected:
// Non-owning views onto shared cell state, excepting ion state.
- const index_type* vec_ci_; // CV to cell index.
+ const index_type* vec_ci_; // CV to cell index.
const value_type* vec_t_; // Cell index to cell-local time.
const value_type* vec_t_to_; // Cell index to cell-local integration step time end.
const value_type* vec_dt_; // CV to integration time step.
const value_type* vec_v_; // CV to cell membrane voltage.
value_type* vec_i_; // CV to cell membrane current density.
+ const value_type* temperature_degC_; // Pointer to global temperature scalar.
deliverable_event_stream* event_stream_ptr_;
// Per-mechanism index and weight data, excepting ion indices.
diff --git a/arbor/backends/multicore/shared_state.cpp b/arbor/backends/multicore/shared_state.cpp
index 6cb520ef..38120f4a 100644
--- a/arbor/backends/multicore/shared_state.cpp
+++ b/arbor/backends/multicore/shared_state.cpp
@@ -1,3 +1,4 @@
+#include <cfloat>
#include <cmath>
#include <iostream>
#include <string>
@@ -129,6 +130,7 @@ shared_state::shared_state(
dt_cv(n_cv, pad(alignment)),
voltage(n_cv, pad(alignment)),
current_density(n_cv, pad(alignment)),
+ temperature_degC(NAN),
deliverable_events(n_cell)
{
// For indices in the padded tail of cv_to_cell, set index to last valid cell index.
@@ -155,6 +157,7 @@ void shared_state::reset(fvm_value_type initial_voltage, fvm_value_type temperat
util::fill(current_density, 0);
util::fill(time, 0);
util::fill(time_to, 0);
+ temperature_degC = temperature_K - 273.15;
for (auto& i: ion_data) {
i.second.reset(temperature_K);
diff --git a/arbor/backends/multicore/shared_state.hpp b/arbor/backends/multicore/shared_state.hpp
index 294f4534..25247f14 100644
--- a/arbor/backends/multicore/shared_state.hpp
+++ b/arbor/backends/multicore/shared_state.hpp
@@ -94,6 +94,7 @@ struct shared_state {
array dt_cv; // Maps CV index to dt [ms].
array voltage; // Maps CV index to membrane voltage [mV].
array current_density; // Maps CV index to current density [A/m²].
+ fvm_value_type temperature_degC; // Global temperature [°C].
std::unordered_map<ionKind, ion_state> ion_data;
diff --git a/aux/CMakeLists.txt b/aux/CMakeLists.txt
index d3030154..2a23eb46 100644
--- a/aux/CMakeLists.txt
+++ b/aux/CMakeLists.txt
@@ -8,6 +8,7 @@ set(aux-sources
)
add_library(arbor-aux ${aux-sources})
+target_compile_options(arbor-aux PRIVATE ${ARB_CXXOPT_ARCH})
target_link_libraries(arbor-aux PUBLIC ext-json arbor)
target_include_directories(arbor-aux PUBLIC include)
set_target_properties(arbor-aux PROPERTIES OUTPUT_NAME arboraux)
diff --git a/include/arbor/distributed_context.hpp b/include/arbor/distributed_context.hpp
index 5f50baa9..4c95292a 100644
--- a/include/arbor/distributed_context.hpp
+++ b/include/arbor/distributed_context.hpp
@@ -88,7 +88,7 @@ private:
virtual void barrier() const = 0;
virtual std::string name() const = 0;
- ARB_PP_FOREACH(ARB_INTERFACE_COLLECTIVES_, ARB_COLLECTIVE_TYPES_);
+ ARB_PP_FOREACH(ARB_INTERFACE_COLLECTIVES_, ARB_COLLECTIVE_TYPES_)
virtual std::vector<std::string> gather(std::string value, int root) const = 0;
virtual ~interface() {}
diff --git a/mechanisms/BuildModules.cmake b/mechanisms/BuildModules.cmake
index a10850f6..945671f2 100644
--- a/mechanisms/BuildModules.cmake
+++ b/mechanisms/BuildModules.cmake
@@ -5,6 +5,15 @@ include(CMakeParseArguments)
function(build_modules)
cmake_parse_arguments(build_modules "" "MODCC;TARGET;SOURCE_DIR;DEST_DIR;MECH_SUFFIX" "MODCC_FLAGS;GENERATES" ${ARGN})
+ if("${build_modules_SOURCE_DIR}" STREQUAL "")
+ set(build_modules_SOURCE_DIR "${CMAKE_CURRENT_SOURCE_DIR}")
+ endif()
+
+ if("${build_modules_DEST_DIR}" STREQUAL "")
+ set(build_modules_SOURCE_DIR "${CMAKE_CURRENT_BINARY_DIR}")
+ endif()
+ file(MAKE_DIRECTORY "${build_modules_DEST_DIR}")
+
set(all_generated)
foreach(mech ${build_modules_UNPARSED_ARGUMENTS})
set(mod "${build_modules_SOURCE_DIR}/${mech}.mod")
diff --git a/mechanisms/CMakeLists.txt b/mechanisms/CMakeLists.txt
index 2ad38b11..3597672e 100644
--- a/mechanisms/CMakeLists.txt
+++ b/mechanisms/CMakeLists.txt
@@ -10,14 +10,6 @@ set(mod_srcdir "${CMAKE_CURRENT_SOURCE_DIR}/mod")
set(mech_dir "${CMAKE_CURRENT_BINARY_DIR}/generated")
file(MAKE_DIRECTORY "${mech_dir}")
-set(modcc_flags)
-if(ARB_VECTORIZE)
- list(APPEND modcc_flags "--simd")
-endif()
-if(ARB_WITH_PROFILING)
- list(APPEND modcc_flags "--profile")
-endif()
-
set(external_modcc)
if(ARB_WITH_EXTERNAL_MODCC)
set(external_modcc MODCC ${modcc})
@@ -28,7 +20,7 @@ build_modules(
SOURCE_DIR "${mod_srcdir}"
DEST_DIR "${mech_dir}"
${external_modcc}
- MODCC_FLAGS -t cpu -t gpu ${modcc_flags}
+ MODCC_FLAGS -t cpu -t gpu ${ARB_MODCC_FLAGS}
GENERATES .hpp _cpu.cpp _gpu.cpp _gpu.cu
TARGET build_all_mods
)
diff --git a/mechanisms/mod/hh.mod b/mechanisms/mod/hh.mod
index 00554aec..3e7a145c 100644
--- a/mechanisms/mod/hh.mod
+++ b/mechanisms/mod/hh.mod
@@ -17,7 +17,7 @@ PARAMETER {
gkbar = .036 (S/cm2)
gl = .0003 (S/cm2)
el = -54.3 (mV)
- celsius = 6.3 (degC)
+ celsius
}
STATE {
@@ -47,20 +47,20 @@ BREAKPOINT {
}
INITIAL {
- rates(v)
+ rates(v, celsius)
m = minf
h = hinf
n = ninf
}
DERIVATIVE states {
- rates(v)
+ rates(v, celsius)
m' = (minf-m)/mtau
h' = (hinf-h)/htau
n' = (ninf-n)/ntau
}
-PROCEDURE rates(v)
+PROCEDURE rates(v, celsius)
{
LOCAL alpha, beta, sum, q10
diff --git a/mechanisms/mod/kamt.mod b/mechanisms/mod/kamt.mod
index d897163b..e218e8b3 100644
--- a/mechanisms/mod/kamt.mod
+++ b/mechanisms/mod/kamt.mod
@@ -59,18 +59,18 @@ BREAKPOINT {
}
INITIAL {
- trates(v)
+ trates(v,celsius)
m=minf
h=hinf
}
DERIVATIVE states {
- trates(v)
+ trates(v,celsius)
m' = (minf-m)/mtau
h' = (hinf-h)/htau
}
-PROCEDURE trates(v) {
+PROCEDURE trates(v,celsius) {
LOCAL qt
qt=q10^((celsius-24)/10)
diff --git a/mechanisms/mod/kdrmt.mod b/mechanisms/mod/kdrmt.mod
index c61ea5e9..54c07044 100644
--- a/mechanisms/mod/kdrmt.mod
+++ b/mechanisms/mod/kdrmt.mod
@@ -50,16 +50,16 @@ BREAKPOINT {
}
INITIAL {
- trates(v)
+ trates(v,celsius)
m=minf
}
DERIVATIVE states {
- trates(v)
+ trates(v,celsius)
m' = (minf-m)/mtau
}
-PROCEDURE trates(v) {
+PROCEDURE trates(v,celsius) {
LOCAL qt
LOCAL alpm, betm
LOCAL tmp
diff --git a/mechanisms/mod/nax.mod b/mechanisms/mod/nax.mod
index 2feddda2..931fe82f 100644
--- a/mechanisms/mod/nax.mod
+++ b/mechanisms/mod/nax.mod
@@ -63,18 +63,18 @@ BREAKPOINT {
}
INITIAL {
- trates(v,sh)
+ trates(v,sh,celsius)
m=minf
h=hinf
}
DERIVATIVE states {
- trates(v,sh)
+ trates(v,sh,celsius)
m' = (minf-m)/mtau
h' = (hinf-h)/htau
}
-PROCEDURE trates(vm,sh2) {
+PROCEDURE trates(vm,sh2,celsius) {
LOCAL a, b, qt
qt=q10^((celsius-24)/10)
a = trap0(vm,tha+sh2,Ra,qa)
diff --git a/modcc/expression.cpp b/modcc/expression.cpp
index 5689fe55..960339f8 100644
--- a/modcc/expression.cpp
+++ b/modcc/expression.cpp
@@ -92,14 +92,27 @@ void IdentifierExpression::semantic(scope_ptr scp) {
yellow(spelling_)));
return;
}
- // if the symbol is an indexed variable, this is the first time that the
- // indexed variable is used in this procedure. In which case, we create
- // a local variable which refers to the indexed variable, which will be
- // found for any subsequent variable lookup inside the procedure
+
+ // If the symbol is an indexed variable, and we're in an API block,
+ // create a local variable which refers to the indexed variable,
+ // which will be found for any subsequent variable lookup inside the
+ // procedure.
+ //
+ // If, however, we are in a PROCEDURE or FUNCTION block, we do not
+ // have access to indexed variables and this constitutes an error.
+
if(auto sym = s->is_indexed_variable()) {
- auto var = new LocalVariable(location_, spelling_);
- var->external_variable(sym);
- s = scope_->add_local_symbol(spelling_, scope_type::symbol_ptr{var});
+ if (scope_->in_api_context()) {
+ auto var = new LocalVariable(location_, spelling_);
+ var->external_variable(sym);
+ s = scope_->add_local_symbol(spelling_, scope_type::symbol_ptr{var});
+ }
+ else {
+ error( pprintf("the symbol '%' refers to an external quantity "
+ "and is unavailable in a function or procedure",
+ yellow(spelling_)));
+ return;
+ }
}
// save the symbol
@@ -445,9 +458,11 @@ std::string ProcedureExpression::to_string() const {
return str;
}
-void ProcedureExpression::semantic(scope_type::symbol_map &global_symbols) {
+void ProcedureExpression::semantic(scope_ptr scp) {
+ scope_ = scp;
+
// assert that the symbol is already visible in the global_symbols
- if(global_symbols.find(name()) == global_symbols.end()) {
+ if(scope_->find_global(name()) == nullptr) {
throw compiler_exception(
"attempt to perform semantic analysis for procedure '"
+ yellow(name())
@@ -455,9 +470,6 @@ void ProcedureExpression::semantic(scope_type::symbol_map &global_symbols) {
location_);
}
- // create the scope for this procedure
- scope_ = std::make_shared<scope_type>(global_symbols);
-
// add the argumemts to the list of local variables
for(auto& a : args_) {
a->semantic(scope_);
@@ -476,6 +488,23 @@ void ProcedureExpression::semantic(scope_type::symbol_map &global_symbols) {
symbol_ = scope_->find_global(name());
}
+void ProcedureExpression::semantic(scope_type::symbol_map &global_symbols) {
+ // create the scope for this procedure and run semantic pass on it
+ scope_ptr scp = std::make_shared<scope_type>(global_symbols);
+ switch (kind_) {
+ case procedureKind::derivative:
+ case procedureKind::kinetic:
+ case procedureKind::initial:
+ case procedureKind::breakpoint:
+ scp->in_api_context(true);
+ break;
+ default:
+ scp->in_api_context(false);
+ break;
+ }
+ semantic(scp);
+}
+
/*******************************************************************************
APIMethod
*******************************************************************************/
@@ -500,6 +529,14 @@ std::string APIMethod::to_string() const {
return str;
}
+void APIMethod::semantic(scope_type::symbol_map &global_symbols) {
+ // create the scope for this procedure, marking it as an API context,
+ // and run semantic pass on it
+ scope_ptr scp = std::make_shared<scope_type>(global_symbols);
+ scp->in_api_context(true);
+ semantic(scp);
+}
+
/*******************************************************************************
InitialBlock
*******************************************************************************/
diff --git a/modcc/expression.hpp b/modcc/expression.hpp
index 1b793591..12b9fc5f 100644
--- a/modcc/expression.hpp
+++ b/modcc/expression.hpp
@@ -1021,6 +1021,7 @@ public:
body_ = std::move(new_body);
}
+ void semantic(scope_ptr scp) override;
void semantic(scope_type::symbol_map &scp) override;
ProcedureExpression* is_procedure() override {return this;}
std::string to_string() const override;
@@ -1049,6 +1050,8 @@ public:
: ProcedureExpression(loc, std::move(name), std::move(args), std::move(body), procedureKind::api)
{}
+ using ProcedureExpression::semantic;
+ void semantic(scope_type::symbol_map &scp) override;
APIMethod* is_api_method() override {return this;}
void accept(Visitor *v) override;
diff --git a/modcc/identifier.hpp b/modcc/identifier.hpp
index 6642caac..3ddd3c56 100644
--- a/modcc/identifier.hpp
+++ b/modcc/identifier.hpp
@@ -51,6 +51,7 @@ enum class sourceKind {
ion_revpot,
ion_iconc,
ion_econc,
+ temperature,
no_source
};
diff --git a/modcc/module.cpp b/modcc/module.cpp
index dbabd329..81312784 100644
--- a/modcc/module.cpp
+++ b/modcc/module.cpp
@@ -4,6 +4,7 @@
#include <iostream>
#include <memory>
#include <set>
+#include <string>
#include <unordered_set>
#include "errorvisitor.hpp"
@@ -467,17 +468,32 @@ void Module::add_variables_to_symbols() {
continue;
}
- // Parameters are scalar by default, but may later be changed to range.
- linkageKind linkage = linkageKind::local;
- auto& sym = create_variable(id.token,
- accessKind::read, visibilityKind::global, linkage, rangeKind::scalar);
+ // Special case: 'celsius' is an external indexed-variable with a special
+ // data source. Retrieval of value is handled especially by printers.
- // set default value if one was specified
- if (id.has_value()) {
- sym->is_variable()->value(std::stod(id.value));
+ if (id.name() == "celsius") {
+ create_indexed_variable("celsius", "celsius",
+ sourceKind::temperature, tok::eq, accessKind::read, ionKind::none, Location());
+ }
+ else {
+ // Parameters are scalar by default, but may later be changed to range.
+ auto& sym = create_variable(id.token,
+ accessKind::read, visibilityKind::global, linkageKind::local, rangeKind::scalar);
+
+ // Set default value if one was specified.
+ if (id.has_value()) {
+ sym->is_variable()->value(std::stod(id.value));
+ }
}
}
+ // Remove `celsius` from the parameter block, as it is not a true parameter anymore.
+ parameter_block_.parameters.erase(
+ std::remove_if(parameter_block_.begin(), parameter_block_.end(),
+ [](const Id& id) { return id.name() == "celsius"; }),
+ parameter_block_.end()
+ );
+
// Add 'assigned' variables, ignoring built-in voltage variable "v".
for (const Id& id: assigned_block_) {
if (id.name() == "v") {
diff --git a/modcc/printer/cprinter.cpp b/modcc/printer/cprinter.cpp
index 0c7a6376..99cc3dc9 100644
--- a/modcc/printer/cprinter.cpp
+++ b/modcc/printer/cprinter.cpp
@@ -394,7 +394,12 @@ void CPrinter::visit(VariableExpression *sym) {
void CPrinter::visit(IndexedVariable *sym) {
indexed_variable_info v = decode_indexed_variable(sym);
- out_ << v.data_var << "[" << v.index_var << "[i_]]";
+ if (v.scalar()) {
+ out_ << v.data_var << "[0]";
+ }
+ else {
+ out_ << v.data_var << "[" << v.index_var << "[i_]]";
+ }
}
void CPrinter::visit(CallExpression* e) {
@@ -450,6 +455,10 @@ void emit_state_read(std::ostream& out, LocalVariable* local) {
void emit_state_update(std::ostream& out, Symbol* from, IndexedVariable* external) {
if (!external->is_write()) return;
+ if (decode_indexed_variable(external).scalar()) {
+ throw compiler_exception("Cannot assign to global scalar: "+external->to_string());
+ }
+
const char* op = external->op()==tok::plus? " += ": " -= ";
out << cprint(external) << op << from->name() << ";\n";
}
@@ -524,9 +533,14 @@ void SimdPrinter::visit(AssignmentExpression* e) {
void SimdPrinter::visit(IndexedVariable *sym) {
indexed_variable_info v = decode_indexed_variable(sym);
- out_ << "S::indirect(" << v.data_var
- << ", " << index_i_name(v.index_var)
- << ", constraint_category_)";
+ if (v.scalar()) {
+ out_ << v.data_var << "[0]";
+ }
+ else {
+ out_ << "S::indirect(" << v.data_var
+ << ", " << index_i_name(v.index_var)
+ << ", constraint_category_)";
+ }
}
void SimdPrinter::visit(CallExpression* e) {
@@ -576,18 +590,23 @@ void emit_simd_state_read(std::ostream& out, LocalVariable* local, simd_expr_con
if (local->is_read()) {
indexed_variable_info v = decode_indexed_variable(local->external_variable());
- if(constraint == simd_expr_constraint::contiguous) {
+ if (v.scalar()) {
+ out << "(" << v.data_var
+ << "[0]);\n";
+ }
+ else if (constraint == simd_expr_constraint::contiguous) {
out << "(" << v.data_var
- << " + " << v.index_var
- << "[index_]);\n";
+ << " + " << v.index_var
+ << "[index_]);\n";
}
- else if(constraint == simd_expr_constraint::constant){
+ else if (constraint == simd_expr_constraint::constant) {
out << "(" << v.data_var
- << "[" << v.index_var
- << "element0]);\n";
+ << "[" << v.index_var
+ << "element0]);\n";
}
- else
+ else {
out << "(" << simdprint(local->external_variable()) << ");\n";
+ }
}
else {
out << " = 0;\n";
@@ -600,14 +619,19 @@ void emit_simd_state_update(std::ostream& out, Symbol* from, IndexedVariable* ex
const char* op = external->op()==tok::plus? " += ": " -= ";
indexed_variable_info v = decode_indexed_variable(external);
- if(constraint == simd_expr_constraint::contiguous) {
- out << "simd_value t_"<< external->name() <<"(" << v.data_var << " + " << v.index_var << "[index_]);\n";
- out << "t_" << external->name() << op << from->name() << ";\n";
- out << "t_" << external->name() << ".copy_to(" << v.data_var << " + " << v.index_var << "[index_]);\n";
-
+ if (v.scalar()) {
+ throw compiler_exception("Cannot assign to global scalar: "+external->to_string());
}
else {
- out << simdprint(external) << op << from->name() << ";\n";
+ if (constraint == simd_expr_constraint::contiguous) {
+ out << "simd_value t_"<< external->name() <<"(" << v.data_var << " + " << v.index_var << "[index_]);\n";
+ out << "t_" << external->name() << op << from->name() << ";\n";
+ out << "t_" << external->name() << ".copy_to(" << v.data_var << " + " << v.index_var << "[index_]);\n";
+
+ }
+ else {
+ out << simdprint(external) << op << from->name() << ";\n";
+ }
}
}
@@ -674,9 +698,16 @@ void emit_simd_api_body(std::ostream& out, APIMethod* method, moduleKind module_
auto body = method->body();
auto indexed_vars = indexed_locals(method->scope());
+ std::vector<LocalVariable*> scalar_indexed_vars;
std::unordered_set<std::string> indices;
for (auto& sym: indexed_vars) {
- indices.insert(decode_indexed_variable(sym->external_variable()).index_var);
+ auto info = decode_indexed_variable(sym->external_variable());
+ if (!info.scalar()) {
+ indices.insert(info.index_var);
+ }
+ else {
+ scalar_indexed_vars.push_back(sym);
+ }
}
if (!body->statements().empty()) {
@@ -718,11 +749,16 @@ void emit_simd_api_body(std::ostream& out, APIMethod* method, moduleKind module_
}
else {
- out << "unsigned n_ = width_;\n\n";
+ // We may nonetheless need to read a global scalar indexed variable.
+ for (auto& sym: scalar_indexed_vars) {
+ emit_simd_state_read(out, sym, simd_expr_constraint::other);
+ }
+
out <<
- "for (unsigned i_ = 0; i_ < n_; i_ += simd_width_) {\n" << indent;
- out << simdprint(body);
- out << popindent << "}\n";
+ "unsigned n_ = width_;\n\n"
+ "for (unsigned i_ = 0; i_ < n_; i_ += simd_width_) {\n" << indent <<
+ simdprint(body) << popindent <<
+ "}\n";
}
}
}
diff --git a/modcc/printer/cudaprinter.cpp b/modcc/printer/cudaprinter.cpp
index 86849396..1013d5a8 100644
--- a/modcc/printer/cudaprinter.cpp
+++ b/modcc/printer/cudaprinter.cpp
@@ -372,7 +372,10 @@ void emit_api_body_cu(std::ostream& out, APIMethod* e, bool is_point_proc) {
std::unordered_set<std::string> indices;
for (auto& sym: indexed_vars) {
- indices.insert(decode_indexed_variable(sym->external_variable()).index_var);
+ auto d = decode_indexed_variable(sym->external_variable());
+ if (!d.scalar()) {
+ indices.insert(d.index_var);
+ }
}
if (!body->statements().empty()) {
@@ -420,6 +423,10 @@ void emit_state_update_cu(std::ostream& out, Symbol* from,
const bool is_minus = external->op()==tok::minus;
auto d = decode_indexed_variable(external);
+ if (d.scalar()) {
+ throw compiler_exception("Cannot assign to global scalar: "+external->to_string());
+ }
+
if (is_point_proc) {
out << "arb::gpu::reduce_by_key(";
is_minus && out << "-";
@@ -439,7 +446,12 @@ void CudaPrinter::visit(VariableExpression *sym) {
void CudaPrinter::visit(IndexedVariable *e) {
auto d = decode_indexed_variable(e);
- out_ << "params_." << d.data_var << "[" << index_i_name(d.index_var) << "]";
+ if (d.scalar()) {
+ out_ << "params_." << d.data_var << "[0]";
+ }
+ else {
+ out_ << "params_." << d.data_var << "[" << index_i_name(d.index_var) << "]";
+ }
}
void CudaPrinter::visit(CallExpression* e) {
diff --git a/modcc/printer/printerutil.cpp b/modcc/printer/printerutil.cpp
index 6050fa06..09bc4c73 100644
--- a/modcc/printer/printerutil.cpp
+++ b/modcc/printer/printerutil.cpp
@@ -144,6 +144,10 @@ indexed_variable_info decode_indexed_variable(IndexedVariable* sym) {
case sourceKind::ion_econc:
data_var=ion_pfx+".external_concentration";
break;
+ case sourceKind::temperature:
+ data_var="temperature_degC_";
+ index_var=""; // scalar global
+ break;
default:
throw compiler_exception(pprintf("unrecognized indexed data source: %", sym), sym->location());
}
diff --git a/modcc/printer/printerutil.hpp b/modcc/printer/printerutil.hpp
index 6881d2cd..208c577a 100644
--- a/modcc/printer/printerutil.hpp
+++ b/modcc/printer/printerutil.hpp
@@ -117,6 +117,7 @@ NetReceiveExpression* find_net_receive(const Module& m);
struct indexed_variable_info {
std::string data_var;
std::string index_var;
+ bool scalar() const { return index_var.empty(); }
};
indexed_variable_info decode_indexed_variable(IndexedVariable* sym);
diff --git a/modcc/scope.hpp b/modcc/scope.hpp
index be2eee29..d7aca81e 100644
--- a/modcc/scope.hpp
+++ b/modcc/scope.hpp
@@ -31,9 +31,18 @@ public:
symbol_map& locals();
symbol_map* globals();
+ bool in_api_context() const {
+ return api_context_;
+ }
+
+ void in_api_context(bool flag) {
+ api_context_ = flag;
+ }
+
private:
symbol_map* global_symbols_=nullptr;
symbol_map local_symbols_;
+ bool api_context_ = false;
};
template<typename Symbol>
diff --git a/modcc/solvers.hpp b/modcc/solvers.hpp
index e2bbc353..33f5b00b 100644
--- a/modcc/solvers.hpp
+++ b/modcc/solvers.hpp
@@ -49,6 +49,9 @@ public:
" but no METHOD was specified in the SOLVE statement",
deriv->location()});
}
+ else {
+ visit((Expression*)e);
+ }
}
};
diff --git a/test/ubench/CMakeLists.txt b/test/ubench/CMakeLists.txt
index a63aee73..a0fb0040 100644
--- a/test/ubench/CMakeLists.txt
+++ b/test/ubench/CMakeLists.txt
@@ -24,6 +24,7 @@ foreach(bench_src ${bench_sources})
string(REGEX REPLACE "\\.[^.]*$" "" bench_exe ${bench_src})
add_executable(${bench_exe} EXCLUDE_FROM_ALL "${bench_src}")
target_link_libraries(${bench_exe} arbor arbor-private-headers ext-benchmark)
+ target_compile_options(${bench_exe} PRIVATE ${ARB_CXXOPT_ARCH})
list(APPEND bench_exe_list ${bench_exe})
endforeach()
diff --git a/test/unit-distributed/CMakeLists.txt b/test/unit-distributed/CMakeLists.txt
index 6e5030f0..fa7d15f9 100644
--- a/test/unit-distributed/CMakeLists.txt
+++ b/test/unit-distributed/CMakeLists.txt
@@ -9,13 +9,13 @@ set(unit-distributed_sources
)
add_executable(unit-local ${unit-distributed_sources})
-target_compile_options(unit-local PRIVATE ${CXXOPT_ARCH})
+target_compile_options(unit-local PRIVATE ${ARB_CXXOPT_ARCH})
target_compile_definitions(unit-local PRIVATE TEST_LOCAL)
target_link_libraries(unit-local PRIVATE gtest arbor arbor-aux arbor-private-headers)
if(ARB_WITH_MPI)
add_executable(unit-mpi ${unit-distributed_sources})
- target_compile_options(unit-mpi PRIVATE ${CXXOPT_ARCH})
+ target_compile_options(unit-mpi PRIVATE ${ARB_CXXOPT_ARCH})
target_compile_definitions(unit-mpi PRIVATE TEST_MPI)
target_link_libraries(unit-mpi PRIVATE gtest arbor arbor-aux arbor-private-headers)
endif()
diff --git a/test/unit/CMakeLists.txt b/test/unit/CMakeLists.txt
index 42205d7d..2945330a 100644
--- a/test/unit/CMakeLists.txt
+++ b/test/unit/CMakeLists.txt
@@ -1,10 +1,40 @@
+# Build mechanisms used solely in unit tests.
+
+set(test_mechanisms
+ celsius_test
+)
+
+include(${PROJECT_SOURCE_DIR}/mechanisms/BuildModules.cmake)
+
+set(external_modcc)
+if(ARB_WITH_EXTERNAL_MODCC)
+ set(external_modcc MODCC ${modcc})
+endif()
+set(test_mech_dir ${CMAKE_CURRENT_BINARY_DIR}/mechanisms)
+
+build_modules(
+ ${test_mechanisms}
+ SOURCE_DIR mod
+ DEST_DIR "${test_mech_dir}"
+ ${external_modcc}
+ MODCC_FLAGS -t cpu -t gpu ${ARB_MODCC_FLAGS}
+ GENERATES .hpp _cpu.cpp _gpu.cpp _gpu.cu
+ TARGET build_test_mods
+)
+
+set(test_mech_sources)
+foreach(mech ${test_mechanisms})
+ list(APPEND test_mech_sources ${test_mech_dir}/${mech}_cpu.cpp)
+ if(ARB_WITH_CUDA)
+ list(APPEND test_mech_sources ${test_mech_dir}/${mech}_gpu.cpp)
+ list(APPEND test_mech_sources ${test_mech_dir}/${mech}_gpu.cu)
+ endif()
+endforeach()
+
+
# TODO: test_mechanism and mechanism prototype comparisons must
# be re-jigged.
-# Build prototype mechanisms for testing in test_mechanisms.
-#
-#include(${PROJECT_SOURCE_DIR}/mechanisms/BuildModules.cmake)
-#
# set(proto_mechanisms pas hh expsyn exp2syn test_kin1 test_kinlva test_ca)
# set(mech_proto_dir "${CMAKE_CURRENT_BINARY_DIR}/mech_proto")
# file(MAKE_DIRECTORY "${mech_proto_dir}")
@@ -46,6 +76,7 @@ set(unit_sources
test_matrix.cpp
test_mc_cell.cpp
test_mechanisms.cpp
+ test_mech_temperature.cpp
test_mechcat.cpp
test_merge_events.cpp
test_multi_event_stream.cpp
@@ -83,7 +114,9 @@ set(unit_sources
test.cpp
# common routines
+ mech_private_field_access.cpp
stats.cpp
+ unit_test_catalogue.cpp
)
if(ARB_WITH_CUDA)
@@ -103,8 +136,9 @@ if(ARB_WITH_CUDA)
)
endif()
-add_executable(unit ${unit_sources})
-target_compile_options(unit PRIVATE ${CXXOPT_ARCH})
+add_executable(unit ${unit_sources} ${test_mech_sources})
+add_dependencies(unit build_test_mods)
+target_compile_options(unit PRIVATE ${ARB_CXXOPT_ARCH})
target_compile_definitions(unit PRIVATE "-DDATADIR=\"${CMAKE_CURRENT_SOURCE_DIR}/swc\"")
+target_include_directories(unit PRIVATE "${CMAKE_CURRENT_BINARY_DIR}")
target_link_libraries(unit PRIVATE gtest arbor arbor-private-headers arbor-aux)
-
diff --git a/test/unit/mech_private_field_access.cpp b/test/unit/mech_private_field_access.cpp
new file mode 100644
index 00000000..3653e87b
--- /dev/null
+++ b/test/unit/mech_private_field_access.cpp
@@ -0,0 +1,64 @@
+#include <arbor/version.hpp>
+
+#include "backends/multicore/fvm.hpp"
+#include "backends/multicore/mechanism.hpp"
+#include "util/maputil.hpp"
+
+#ifdef ARB_GPU_ENABLED
+#include "backends/gpu/fvm.hpp"
+#include "backends/gpu/mechanism.hpp"
+#include "memory/cuda_wrappers.hpp"
+#endif
+
+#include "common.hpp"
+#include "mech_private_field_access.hpp"
+
+using namespace arb;
+using field_table_type = std::vector<std::pair<const char*, fvm_value_type**>>;
+
+// Multicore mechanisms:
+
+ACCESS_BIND(field_table_type (multicore::mechanism::*)(), multicore_field_table_ptr, &multicore::mechanism::field_table)
+
+std::vector<fvm_value_type> mechanism_field(multicore::mechanism* m, const std::string& key) {
+ auto opt_ptr = util::value_by_key((m->*multicore_field_table_ptr)(), key);
+ if (!opt_ptr) throw std::logic_error("internal error: no such field in mechanism");
+
+ const fvm_value_type* field_data = *opt_ptr.value();
+ return std::vector<fvm_value_type>(field_data, field_data+m->size());
+}
+
+// GPU mechanisms:
+
+#ifdef ARB_GPU_ENABLED
+ACCESS_BIND(field_table_type (gpu::mechanism::*)(), gpu_field_table_ptr, &gpu::mechanism::field_table)
+
+std::vector<fvm_value_type> mechanism_field(gpu::mechanism* m, const std::string& key) {
+ auto opt_ptr = util::value_by_key((m->*gpu_field_table_ptr)(), key);
+ if (!opt_ptr) throw std::logic_error("internal error: no such field in mechanism");
+
+ const fvm_value_type* field_data = *opt_ptr.value();
+ std::vector<fvm_value_type> values(m->size());
+
+ cudaDeviceSynchronize();
+ memory::cuda_memcpy_d2h(values.data(), field_data, sizeof(fvm_value_type)*m->size());
+ return values;
+}
+#endif
+
+// Generic access:
+
+std::vector<fvm_value_type> mechanism_field(mechanism* m, const std::string& key) {
+ if (auto p = dynamic_cast<multicore::mechanism*>(m)) {
+ return mechanism_field(p, key);
+ }
+
+#ifdef ARB_GPU_ENABLED
+ if (auto p = dynamic_cast<gpu::mechanism*>(m)) {
+ return mechanism_field(p, key);
+ }
+#endif
+
+ throw std::logic_error("internal error: mechanism instantiated on unknown backend");
+}
+
diff --git a/test/unit/mech_private_field_access.hpp b/test/unit/mech_private_field_access.hpp
new file mode 100644
index 00000000..50c7b26b
--- /dev/null
+++ b/test/unit/mech_private_field_access.hpp
@@ -0,0 +1,16 @@
+#pragma once
+
+#include <memory>
+#include <string>
+#include <vector>
+
+#include <arbor/mechanism.hpp>
+
+// Get a copy of the data within a mechanisms's (private) named field.
+
+std::vector<arb::fvm_value_type> mechanism_field(arb::mechanism* m, const std::string& key);
+
+template <typename DerivedMechPtr>
+inline std::vector<arb::fvm_value_type> mechanism_field(const std::unique_ptr<DerivedMechPtr>& m, const std::string& key) {
+ return mechanism_field(static_cast<arb::mechanism*>(m.get()), key);
+}
diff --git a/test/unit/mod/celsius_test.mod b/test/unit/mod/celsius_test.mod
new file mode 100644
index 00000000..838ef243
--- /dev/null
+++ b/test/unit/mod/celsius_test.mod
@@ -0,0 +1,27 @@
+NEURON {
+ SUFFIX celsius_test
+}
+
+PARAMETER {
+ celsius
+}
+
+STATE {
+ c
+}
+
+ASSIGNED {
+}
+
+BREAKPOINT {
+ SOLVE states
+}
+
+DERIVATIVE states {
+ c = celsius
+}
+
+INITIAL {
+ c = 0
+}
+
diff --git a/test/unit/test_mech_temperature.cpp b/test/unit/test_mech_temperature.cpp
new file mode 100644
index 00000000..8d152d99
--- /dev/null
+++ b/test/unit/test_mech_temperature.cpp
@@ -0,0 +1,80 @@
+#include <vector>
+
+#include <arbor/mechanism.hpp>
+#include <arbor/version.hpp>
+
+#include "backends/multicore/fvm.hpp"
+#ifdef ARB_GPU_ENABLED
+#include "backends/gpu/fvm.hpp"
+#endif
+
+#include "common.hpp"
+#include "mech_private_field_access.hpp"
+#include "unit_test_catalogue.hpp"
+
+using namespace arb;
+
+template <typename backend>
+void run_celsius_test() {
+ auto cat = make_unit_test_catalogue();
+
+ // one cell, three CVs:
+
+ fvm_size_type ncell = 1;
+ fvm_size_type ncv = 3;
+ std::vector<fvm_index_type> cv_to_cell(ncv, 0);
+
+ auto celsius_test = cat.instance<backend>("celsius_test");
+ auto shared_state = std::make_unique<typename backend::shared_state>(
+ ncell, cv_to_cell, celsius_test->data_alignment());
+
+ mechanism::layout layout;
+ layout.weight.assign(ncv, 1.);
+ for (fvm_size_type i = 0; i<ncv; ++i) {
+ layout.cv.push_back(i);
+ }
+
+ celsius_test->instantiate(0, *shared_state, layout);
+
+ double temperature_K = 300.;
+ double temperature_C = temperature_K-273.15;
+
+ shared_state->reset(-65., temperature_K);
+
+ // expect 0 value in state 'c' after init:
+
+ celsius_test->nrn_init();
+ std::vector<fvm_value_type> expected_c_values(ncv, 0.);
+
+ EXPECT_EQ(expected_c_values, mechanism_field(celsius_test.get(), "c"));
+
+ // expect temperature_C value in state 'c' after state update:
+
+ celsius_test->nrn_state();
+ expected_c_values.assign(ncv, temperature_C);
+
+ EXPECT_EQ(expected_c_values, mechanism_field(celsius_test.get(), "c"));
+
+ // reset with new temperature and repeat test:
+
+ temperature_K = 290.;
+ temperature_C = temperature_K-273.15;
+
+ shared_state->reset(-65., temperature_K);
+ celsius_test->nrn_init();
+
+ celsius_test->nrn_state();
+ expected_c_values.assign(ncv, temperature_C);
+
+ EXPECT_EQ(expected_c_values, mechanism_field(celsius_test.get(), "c"));
+}
+
+TEST(mech_temperature, celsius) {
+ run_celsius_test<multicore::backend>();
+}
+
+#ifdef ARB_GPU_ENABLED
+TEST(mech_temperature_gpu, celsius) {
+ run_celsius_test<gpu::backend>();
+}
+#endif
diff --git a/test/unit/test_synapses.cpp b/test/unit/test_synapses.cpp
index 0c5464e6..1428455d 100644
--- a/test/unit/test_synapses.cpp
+++ b/test/unit/test_synapses.cpp
@@ -15,6 +15,7 @@
#include "util/range.hpp"
#include "common.hpp"
+#include "mech_private_field_access.hpp"
using namespace arb;
@@ -23,17 +24,7 @@ using shared_state = backend::shared_state;
using value_type = backend::value_type;
using size_type = backend::size_type;
-// Access to mechanisms protected data:
-using field_table_type = std::vector<std::pair<const char*, value_type**>>;
-ACCESS_BIND(field_table_type (multicore::mechanism::*)(), field_table_ptr, &multicore::mechanism::field_table)
-
-util::range<const value_type*> mechanism_field(std::unique_ptr<multicore::mechanism>& m, const std::string& key) {
- if (auto opt_ptr = util::value_by_key((m.get()->*field_table_ptr)(), key)) {
- const value_type* field = *opt_ptr.value();
- return util::make_range(field, field+m->size());
- }
- throw std::logic_error("internal error: no such field in mechanism");
-}
+// Access to more mechanism protected data:
ACCESS_BIND(const value_type* multicore::mechanism::*, vec_v_ptr, &multicore::mechanism::vec_v_)
ACCESS_BIND(value_type* multicore::mechanism::*, vec_i_ptr, &multicore::mechanism::vec_i_)
@@ -74,6 +65,11 @@ static bool all_equal_to(const Seq& s, double v) {
});
}
+template <typename A, typename B>
+auto unique_cast(std::unique_ptr<B> p) {
+ return std::unique_ptr<A>(dynamic_cast<A*>(p.release()));
+}
+
TEST(synapses, syn_basic_state) {
using util::fill;
using value_type = multicore::backend::value_type;
@@ -83,16 +79,10 @@ TEST(synapses, syn_basic_state) {
int num_comp = 4;
int num_cell = 1;
- auto multicore_mechanism_instance = [](const char* name) {
- return std::unique_ptr<multicore::mechanism>(
- dynamic_cast<multicore::mechanism*>(
- global_default_catalogue().instance<backend>(name).release()));
- };
-
- auto expsyn = multicore_mechanism_instance("expsyn");
+ auto expsyn = unique_cast<multicore::mechanism>(global_default_catalogue().instance<backend>("expsyn"));
ASSERT_TRUE(expsyn);
- auto exp2syn = multicore_mechanism_instance("exp2syn");
+ auto exp2syn = unique_cast<multicore::mechanism>(global_default_catalogue().instance<backend>("exp2syn"));
ASSERT_TRUE(exp2syn);
auto align = std::max(expsyn->data_alignment(), exp2syn->data_alignment());
diff --git a/test/unit/unit_test_catalogue.cpp b/test/unit/unit_test_catalogue.cpp
new file mode 100644
index 00000000..52a2f72c
--- /dev/null
+++ b/test/unit/unit_test_catalogue.cpp
@@ -0,0 +1,28 @@
+#include <arbor/mechcat.hpp>
+#include <arbor/version.hpp>
+
+#ifdef ARB_GPU_ENABLED
+#include "backends/gpu/fvm.hpp"
+#endif
+#include "backends/multicore/fvm.hpp"
+
+#include "unit_test_catalogue.hpp"
+#include "mechanisms/celsius_test.hpp"
+
+#include "../gtest.h"
+
+using namespace arb;
+
+mechanism_catalogue make_unit_test_catalogue() {
+ mechanism_catalogue cat;
+
+ cat.add("celsius_test", mechanism_celsius_test_info());
+
+ cat.register_implementation("celsius_test", make_mechanism_celsius_test<multicore::backend>());
+#ifdef ARB_GPU_ENABLED
+ cat.register_implementation("celsius_test", make_mechanism_celsius_test<gpu::backend>());
+#endif
+
+ return cat;
+}
+
diff --git a/test/unit/unit_test_catalogue.hpp b/test/unit/unit_test_catalogue.hpp
new file mode 100644
index 00000000..6516ffa2
--- /dev/null
+++ b/test/unit/unit_test_catalogue.hpp
@@ -0,0 +1,5 @@
+#pragma once
+
+#include <arbor/mechcat.hpp>
+
+arb::mechanism_catalogue make_unit_test_catalogue();
diff --git a/test/validation/CMakeLists.txt b/test/validation/CMakeLists.txt
index 8826d8e6..6257a38d 100644
--- a/test/validation/CMakeLists.txt
+++ b/test/validation/CMakeLists.txt
@@ -14,6 +14,7 @@ set(validation_sources
)
add_executable(validate ${validation_sources})
+target_compile_options(validate PRIVATE ${ARB_CXXOPT_ARCH})
target_compile_definitions(validate PRIVATE "ARB_DATADIR=\"${ARB_VALIDATION_DATA_DIR}\"")
target_link_libraries(validate PRIVATE gtest arbor arbor-aux ext-json)
--
GitLab