diff --git a/CMakeLists.txt b/CMakeLists.txt
index 6ff82e3a4657bd1d4889842a759e7c3597f38a44..b4fc22b8ac4eb769bb83077efdd4b9a80bc8717e 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 e545e791806f8bef9e21a8d7c48e2fcf8a962d87..9a70b562210256d012b278179af85aa6ff71f822 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 a867f69da375e985d55eda099d8801df56cb5adc..fb256579dbacd3325094a014c98c1e85f754d643 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 9ab086088dbe6a973ab44602521209b17ad8cc3c..5cd8972c45ab44913977612d2ae7710e846c1300 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 da1ae4effe82c7969fc548960d1b7aeb3377b82d..81174715d60c62e074da9db07cf4a507e2e07245 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 81c48e8f890c1ec92d2ac7f41d57e8a728267e5c..63bd34493edb51f690e34ce1b53d6d294c96783b 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 8724d8aa03704f1ebea7f506a8659b9eb17c86c7..af08a9318ee8af6e1d6f650d0c12bab3833b89f8 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 6cb520ef3e0cd4843bd2eb23f92a350a1e5c3631..38120f4a15fecb7725943d5a286b304669c954c9 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 294f453430373fbbec62be5dd0e5a8e2798679a0..25247f1435d377eb8a01a107738ca49d656f0907 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 d30301542bc64f71178b77dbc3ac525e18f9ecf3..2a23eb46bfbdc36d13659ad45d97bb435201854e 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 5f50baa99397d055e70b0bb287c6f1ae77fad8ef..4c95292ab127851e4a4420efa3de355d4ba66e9c 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 a10850f66eed3fbe72afc797942b52d53d459add..945671f21d565ad2ca487979f56d74d9baf95915 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 2ad38b1126ea4025fee0d9a4685b2c937c916c1b..3597672e3af91f650dbaaa468a40317c735fe023 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 00554aecfba7a35679c90b60c255ae2ee14e5b6f..3e7a145ced520677f472632d8a773e6b6ac6e730 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 d897163b24634d650b969ad28a3665038de7fec3..e218e8b3f8ec2a1facabbb76b56cdb130467121b 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 c61ea5e92ff04265f3bc4631189fdfb30fd8b857..54c070441deb82925b7468c0e04d62de9fe4a205 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 2feddda26aca3a6c3f01a06855cec3e11df448c2..931fe82fb08d7f4bee717e9d49290f5c2299bc87 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 5689fe5599ad82c2855bf888d2c1d52755fff142..960339f80043842f6643de94c6e563e5679047cf 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 1b793591a3e73f099a2314953d5dcc686ecb7ca9..12b9fc5f174aede12fbe6c67f59c1f235db044c7 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 6642caacf668b12f17981ac2b05be0509e1b480c..3ddd3c56ea34dcd853b0d3a6d0d1984be3c509a8 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 dbabd32989770f0bd0e7d4bcd4c4df7ef6182be5..81312784081f2bf80f15b78d191eb8b8db502966 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 0c7a6376ffd031ec1214f84801e0b7c3f61046b5..99cc3dc9d09367b238234250134613aad691b60d 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 86849396a4bc524919e226e2f2f518bf0f0692bb..1013d5a80cd6efc59d84133e672c677505283c61 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 6050fa06bf730bd497a63fa4f5f05bdc95da6de2..09bc4c7371384d6e21368aeb87af8f898d3d8567 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 6881d2cd4a2da21764caa0065dbcfb29333e0ea7..208c577af369c07705909afcd880ec7e5517c31b 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 be2eee29450104c0e85d36e75004e9ee3bca2b92..d7aca81e4e0a952513d6c8162f1f5b9e1bfc253e 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 e2bbc35323cda815a5f1918d9a16b7e00189d2b9..33f5b00bb50fe9b94e0ff17651151d5e6d145fbe 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 a63aee730f4527a4275479037c436d35c81fa75d..a0fb0040738736606125f3016c110db7b47d8167 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 6e5030f05dd3eaa19f97d8c15519063bce4bfa88..fa7d15f98b71c891a9d6ab2f9daec87d4f849375 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 42205d7da6e082249b2a2bfb97e1a2c08b339170..2945330af8a482cdf0ebbb28d995d940fd684803 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 0000000000000000000000000000000000000000..3653e87b64546251cdde57128f249043c1b7e610
--- /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 0000000000000000000000000000000000000000..50c7b26bc8a54190920ecb2b963af94ed9eea4f7
--- /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 0000000000000000000000000000000000000000..838ef243055f0fe9feb995b701a2e2cd6ac61f1b
--- /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 0000000000000000000000000000000000000000..8d152d9990fe131ff4fc91113721d163cb973911
--- /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 0c5464e698680d5375593857758975bd830ce05f..1428455d82203481e8376d225a99cde1fec7e080 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 0000000000000000000000000000000000000000..52a2f72c50adea56d7a97e9d7086b2d0ecfa3259
--- /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 0000000000000000000000000000000000000000..6516ffa23bcc03b6351041bd527a9b0f4abe5384
--- /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 8826d8e611e952f3ddae3fe154308c8fc1e18386..6257a38dccbf35aef0329bb1eb21523ed4c1d9b2 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)