diff --git a/CMakeLists.txt b/CMakeLists.txt index 5b7a8beba558ddb122fee99e6f381aeff96ebe1d..86aff17b95f532bbddc5beb17f02cce3ff363fff 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -93,7 +93,7 @@ include_directories(${CMAKE_SOURCE_DIR}/vector) include_directories(${CMAKE_SOURCE_DIR}/include) include_directories(${CMAKE_SOURCE_DIR}/src) include_directories(${CMAKE_SOURCE_DIR}/miniapp) -include_directories(${CMAKE_SOURCE_DIR}/modcc/src) +include_directories(${CMAKE_SOURCE_DIR}/modcc) include_directories(${CMAKE_SOURCE_DIR}) if( "${WITH_TBB}" STREQUAL "ON" ) include_directories(${TBB_INCLUDE_DIRS}) diff --git a/modcc/CMakeLists.txt b/modcc/CMakeLists.txt index 508202cb0c157e867449aa0142efee551c538ff7..aa98b8617c6581eee3fe386eaaf168b964c56d96 100644 --- a/modcc/CMakeLists.txt +++ b/modcc/CMakeLists.txt @@ -1,6 +1,27 @@ -# generated .a and .so go into /lib -set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib) -set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib) +set(MODCC_SOURCES + token.cpp + lexer.cpp + expression.cpp + parser.cpp + textbuffer.cpp + cprinter.cpp + functionexpander.cpp + functioninliner.cpp + cudaprinter.cpp + expressionclassifier.cpp + constantfolder.cpp + errorvisitor.cpp + module.cpp +) -add_subdirectory(src) +add_library(compiler ${MODCC_SOURCES}) + +add_executable(modcc modcc.cpp) + +target_link_libraries(modcc LINK_PUBLIC compiler) + +set_target_properties(modcc + PROPERTIES + RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/modcc" +) diff --git a/modcc/src/blocks.hpp b/modcc/blocks.hpp similarity index 99% rename from modcc/src/blocks.hpp rename to modcc/blocks.hpp index 58db08e155343039f03b41577b7887f05480e647..b9e41bfa3d2fb1963691abae539f4ad94d45de2c 100644 --- a/modcc/src/blocks.hpp +++ b/modcc/blocks.hpp @@ -164,3 +164,4 @@ inline std::ostream& operator<< (std::ostream& os, AssignedBlock const& A) { return os; } + diff --git a/modcc/src/constantfolder.cpp b/modcc/constantfolder.cpp similarity index 98% rename from modcc/src/constantfolder.cpp rename to modcc/constantfolder.cpp index 3dfb609f7c06f16972633cc3c5f3d5edbfe7b2a2..8f04714fd2670f82a32459e8b18bc8f24df8c52d 100644 --- a/modcc/src/constantfolder.cpp +++ b/modcc/constantfolder.cpp @@ -31,7 +31,8 @@ void ConstantFolderVisitor::visit(UnaryExpression *e) { e->expression()->accept(this); if(is_number) { if(!e->is_number()) { - e->replace_expression(make_expression<NumberExpression>(e->location(), value)); + e->replace_expression( + make_expression<NumberExpression>(e->location(), value)); } switch(e->op()) { case tok::minus : @@ -173,4 +174,3 @@ void ConstantFolderVisitor::visit(IfExpression *e) { e->false_branch()->accept(this); } } - diff --git a/modcc/src/constantfolder.hpp b/modcc/constantfolder.hpp similarity index 96% rename from modcc/src/constantfolder.hpp rename to modcc/constantfolder.hpp index 1fc5f507a4bbdd3d58622e7772944bd1d7b25b3a..ede3fefd68d3f732f4f804c88fb06b704d3d201d 100644 --- a/modcc/src/constantfolder.hpp +++ b/modcc/constantfolder.hpp @@ -22,6 +22,5 @@ public: // store intermediate results as long double, i.e. 80-bit precision long double value = 0.; - bool is_number = false; + bool is_number = false; }; - diff --git a/modcc/src/cprinter.cpp b/modcc/cprinter.cpp similarity index 99% rename from modcc/src/cprinter.cpp rename to modcc/cprinter.cpp index b1991ea9d52c4f2e227033650bff22f36cb72c48..cf4b3d93afbcab6847732cea15629c47642de29f 100644 --- a/modcc/src/cprinter.cpp +++ b/modcc/cprinter.cpp @@ -82,6 +82,7 @@ CPrinter::CPrinter(Module &m, bool o) text_.add_line(); ////////////////////////////////////////////// + // constructor ////////////////////////////////////////////// int num_vars = array_variables.size(); text_.add_line(class_name + "(view_type vec_v, view_type vec_i, const_index_view node_index)"); @@ -141,7 +142,6 @@ CPrinter::CPrinter(Module &m, bool o) } text_.add_line(); - //text_.add_line("INIT_PROFILE"); text_.decrease_indentation(); text_.add_line("}"); @@ -891,3 +891,4 @@ void CPrinter::visit(BinaryExpression *e) { // reset parent precedence parent_op_ = pop; } + diff --git a/modcc/src/cprinter.hpp b/modcc/cprinter.hpp similarity index 100% rename from modcc/src/cprinter.hpp rename to modcc/cprinter.hpp diff --git a/modcc/src/cudaprinter.cpp b/modcc/cudaprinter.cpp similarity index 55% rename from modcc/src/cudaprinter.cpp rename to modcc/cudaprinter.cpp index 24433d2be230ea0d083e4223698cecec501577fd..3dad1b5a33e6410deb3e4f1400fd8feb1c34a834 100644 --- a/modcc/src/cudaprinter.cpp +++ b/modcc/cudaprinter.cpp @@ -1,3 +1,5 @@ +#include <algorithm> + #include "cudaprinter.hpp" #include "lexer.hpp" @@ -31,11 +33,16 @@ CUDAPrinter::CUDAPrinter(Module &m, bool o) text_.add_line("#include <cmath>"); text_.add_line("#include <limits>"); text_.add_line(); - text_.add_line("#include <indexedview.hpp>"); text_.add_line("#include <mechanism.hpp>"); - text_.add_line("#include <target.hpp>"); + text_.add_line("#include <mechanism_interface.hpp>"); + //text_.add_line("#include <gpu/util.hpp>"); text_.add_line(); + + text_.add_line("namespace nest{ namespace mc{ namespace mechanisms{ namespace gpu{ namespace " + m.name() + "{"); + text_.add_line(); + increase_indentation(); + //////////////////////////////////////////////////////////// // generate the parameter pack //////////////////////////////////////////////////////////// @@ -69,32 +76,62 @@ CUDAPrinter::CUDAPrinter(Module &m, bool o) param_pack.push_back(tname + ".index.data()"); } - text_.add_line("// matrix"); - text_.add_line("T* vec_rhs;"); - text_.add_line("T* vec_d;"); + text_.add_line("// voltage and current state within the cell"); text_.add_line("T* vec_v;"); - param_pack.push_back("matrix_.vec_rhs().data()"); - param_pack.push_back("matrix_.vec_d().data()"); - param_pack.push_back("matrix_.vec_v().data()"); + text_.add_line("T* vec_i;"); + param_pack.push_back("vec_v_.data()"); + param_pack.push_back("vec_i_.data()"); + + text_.add_line("T* vec_area;"); + param_pack.push_back("vec_area_.data()"); text_.add_line("// node index information"); text_.add_line("I* ni;"); - text_.add_line("unsigned long n;"); + text_.add_line("unsigned long n_;"); text_.decrease_indentation(); text_.add_line("};"); text_.add_line(); - param_pack.push_back("node_indices_.data()"); - param_pack.push_back("node_indices_.size()"); - + param_pack.push_back("node_index_.data()"); + param_pack.push_back("node_index_.size()"); //////////////////////////////////////////////////////// // write the CUDA kernels //////////////////////////////////////////////////////// - text_.add_line("namespace impl {"); - text_.add_line("namespace " + m.name() + " {"); - text_.add_line(); + text_.add_line("namespace kernels {"); { increase_indentation(); + + text_.add_line("__device__"); + text_.add_line("inline double atomicAdd(double* address, double val) {"); + text_.increase_indentation(); + text_.add_line("using I = unsigned long long int;"); + text_.add_line("I* address_as_ull = (I*)address;"); + text_.add_line("I old = *address_as_ull, assumed;"); + text_.add_line("do {"); + text_.add_line("assumed = old;"); + text_.add_line("old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val+__longlong_as_double(assumed)));"); + text_.add_line("} while (assumed != old);"); + text_.add_line("return __longlong_as_double(old);"); + text_.decrease_indentation(); + text_.add_line("}"); + text_.add_line(); + /* + text_.add_line("__device__"); + text_.add_line("inline double atomicSub(double* address, double val) {"); + text_.increase_indentation(); + text_.add_line("return atomicAdd(address, -val);"); + text_.decrease_indentation(); + text_.add_line("}"); + text_.add_line(); + text_.add_line("__device__"); + text_.add_line("inline float atomicSub(float* address, float val) {"); + text_.increase_indentation(); + text_.add_line("return atomicAdd(address, -val);"); + text_.decrease_indentation(); + text_.add_line("}"); + text_.add_line(); + */ + // forward declarations of procedures for(auto const &var : m.symbols()) { if( var.second->kind()==symbolKind::procedure @@ -107,122 +144,287 @@ CUDAPrinter::CUDAPrinter(Module &m, bool o) } // print stubs that call API method kernels that are defined in the - // imp::name namespace + // kernels::name namespace auto proctest = [] (procedureKind k) {return k == procedureKind::normal || k == procedureKind::api; }; for(auto const &var : m.symbols()) { - if( var.second->kind()==symbolKind::procedure - && proctest(var.second->is_procedure()->kind())) + if (var.second->kind()==symbolKind::procedure && + proctest(var.second->is_procedure()->kind())) { var.second->accept(this); } } decrease_indentation(); } - text_.add_line("} // namespace " + m.name()); - text_.add_line("} // namespace impl"); + text_.add_line("} // namespace kernels"); text_.add_line(); ////////////////////////////////////////////// ////////////////////////////////////////////// - std::string class_name = "Mechanism_" + m.name(); - - text_ << "template<typename T, typename I>\n"; - text_ << "class " + class_name + " : public Mechanism<T, I, targetKind::gpu> {\n"; - text_ << "public:\n\n"; - text_ << " using base = Mechanism<T, I, targetKind::gpu>;\n"; - text_ << " using value_type = typename base::value_type;\n"; - text_ << " using size_type = typename base::size_type;\n"; - text_ << " using vector_type = typename base::vector_type;\n"; - text_ << " using view_type = typename base::view_type;\n"; - text_ << " using index_type = typename base::index_type;\n"; - text_ << " using index_view = typename index_type::view_type;\n"; - text_ << " using indexed_view= typename base::indexed_view;\n\n"; - text_ << " using matrix_type = typename base::matrix_type;\n\n"; - text_ << " using param_pack_type = " << m.name() << "_ParamPack<T,I>;\n\n"; + std::string class_name = "mechanism_" + m.name(); + + text_.add_line("template<typename T, typename I>"); + text_.add_line("class " + class_name + " : public ::nest::mc::mechanisms::gpu::mechanism<T, I> {"); + text_.add_line("public:"); + text_.increase_indentation(); + text_.add_line("using base = ::nest::mc::mechanisms::gpu::mechanism<T, I>;"); + text_.add_line("using value_type = typename base::value_type;"); + text_.add_line("using size_type = typename base::size_type;"); + text_.add_line("using vector_type = typename base::vector_type;"); + text_.add_line("using view_type = typename base::view_type;"); + text_.add_line("using index_type = typename base::index_type;"); + text_.add_line("using index_view = typename base::index_view;"); + text_.add_line("using const_index_view = typename base::const_index_view;"); + text_.add_line("using indexed_view_type= typename base::indexed_view_type;"); + text_.add_line("using ion_type = typename base::ion_type;"); + text_.add_line("using param_pack_type = " + m.name() + "_ParamPack<T,I>;"); ////////////////////////////////////////////// ////////////////////////////////////////////// for(auto& ion: m.neuron_block().ions) { auto tname = "Ion" + ion.name; - text_ << " struct " + tname + " {\n"; + text_.add_line("struct " + tname + " {"); + text_.increase_indentation(); for(auto& field : ion.read) { - text_ << " view_type " + field.spelling + ";\n"; + text_.add_line("view_type " + field.spelling + ";"); } for(auto& field : ion.write) { - text_ << " view_type " + field.spelling + ";\n"; + text_.add_line("view_type " + field.spelling + ";"); } - text_ << " index_type index;\n"; - text_ << " std::size_t memory() const { return sizeof(size_type)*index.size(); }\n"; - text_ << " std::size_t size() const { return index.size(); }\n"; - text_ << " };\n"; - text_ << " " + tname + " ion_" + ion.name + ";\n\n"; + text_.add_line("index_type index;"); + text_.add_line("std::size_t memory() const { return sizeof(size_type)*index.size(); }"); + text_.add_line("std::size_t size() const { return index.size(); }"); + text_.decrease_indentation(); + text_.add_line("};"); + text_.add_line(tname + " ion_" + ion.name + ";"); + text_.add_line(); } ////////////////////////////////////////////// + // constructor ////////////////////////////////////////////// + int num_vars = array_variables.size(); - text_ << " " + class_name + "(\n"; - text_ << " matrix_type* matrix,\n"; - text_ << " index_view node_indices)\n"; - text_ << " : base(matrix, node_indices)\n"; - text_ << " {\n"; - text_ << " size_type num_fields = " << num_vars << ";\n"; - text_ << " size_type n = size();\n"; - text_ << " data_ = vector_type(n * num_fields);\n"; - text_ << " data_(memory::all) = std::numeric_limits<value_type>::quiet_NaN();\n"; + text_.add_line(); + text_.add_line("template <typename IVT>"); + text_.add_line(class_name + "(view_type vec_v, view_type vec_i, IVT node_index) :"); + text_.add_line(" base(vec_v, vec_i, node_index)"); + text_.add_line("{"); + text_.increase_indentation(); + text_.add_gutter() << "size_type num_fields = " << num_vars << ";"; + text_.end_line(); + + text_.add_line(); + text_.add_line("// calculate the padding required to maintain proper alignment of sub arrays"); + text_.add_line("auto alignment = data_.alignment();"); + text_.add_line("auto field_size_in_bytes = sizeof(value_type)*size();"); + text_.add_line("auto remainder = field_size_in_bytes % alignment;"); + text_.add_line("auto padding = remainder ? (alignment - remainder)/sizeof(value_type) : 0;"); + text_.add_line("auto field_size = size()+padding;"); + + text_.add_line(); + text_.add_line("// allocate memory"); + text_.add_line("data_ = vector_type(field_size * num_fields);"); + text_.add_line("data_(memory::all) = std::numeric_limits<value_type>::quiet_NaN();"); + + // assign the sub-arrays + // replace this : data_(1*n, 2*n); + // with this : data_(1*field_size, 1*field_size+n); + + text_.add_line(); + text_.add_line("// asign the sub-arrays"); for(int i=0; i<num_vars; ++i) { char namestr[128]; sprintf(namestr, "%-15s", array_variables[i]->name().c_str()); - text_ << " " << namestr << " = data_(" << i << "*n, " << i+1 << "*n);\n"; + text_.add_line( + array_variables[i]->name() + " = data_(" + + std::to_string(i) + "*field_size, " + std::to_string(i+1) + "*field_size);"); } + for(auto const& var : array_variables) { double val = var->value(); // only non-NaN fields need to be initialized, because data_ // is NaN by default if(val == val) { - text_ << " " << var->name() << "(memory::all) = " << val << ";\n"; + text_.add_line(var->name() + "(memory::all) = " + std::to_string(val) + ";"); } } - text_ << " INIT_PROFILE\n"; - text_ << " }\n\n"; + text_.add_line(); + text_.decrease_indentation(); + text_.add_line("}"); ////////////////////////////////////////////// ////////////////////////////////////////////// - text_ << " using base::size;\n\n"; + text_.add_line("using base::size;"); + text_.add_line(); - text_ << " std::size_t memory() const override {\n"; - text_ << " auto s = std::size_t{0};\n"; - text_ << " s += data_.size()*sizeof(value_type);\n"; + text_.add_line("std::size_t memory() const override {"); + text_.increase_indentation(); + text_.add_line("auto s = std::size_t{0};"); + text_.add_line("s += data_.size()*sizeof(value_type);"); for(auto& ion: m.neuron_block().ions) { - text_ << " s += ion_" + ion.name + ".memory();\n"; + text_.add_line("s += ion_" + ion.name + ".memory();"); } - text_ << " return s;\n"; - text_ << " }\n\n"; + text_.add_line("return s;"); + text_.decrease_indentation(); + text_.add_line("}"); + text_.add_line(); - text_ << " void set_params(value_type t_, value_type dt_) override {\n"; - text_ << " t = t_;\n"; - text_ << " dt = dt_;\n"; - text_ << " param_pack_ = param_pack_type{\n"; - //for(auto i=0; i<param_pack.size(); ++i) - for(auto &str: param_pack) { - text_ << " " << str << ",\n"; + // print the member funtion that + // * sets time step parameters + // * packs up the parameters for use on the GPU + text_.add_line("void set_params(value_type t_, value_type dt_) override {"); + text_.increase_indentation(); + text_.add_line("t = t_;"); + text_.add_line("dt = dt_;"); + text_.add_line("param_pack_ ="); + text_.increase_indentation(); + text_.add_line("param_pack_type {"); + text_.increase_indentation(); + for(auto& str: param_pack) { + text_.add_line(str + ","); } - text_ << " };\n"; - text_ << " }\n\n"; + text_.decrease_indentation(); + text_.add_line("};"); + text_.decrease_indentation(); + text_.decrease_indentation(); + text_.add_line("}"); + text_.add_line(); - text_ << " std::string name() const override {\n"; - text_ << " return \"" << m.name() << "\";\n"; - text_ << " }\n\n"; + // name member function + text_.add_line("std::string name() const override {"); + text_.increase_indentation(); + text_.add_line("return \"" + m.name() + "\";"); + text_.decrease_indentation(); + text_.add_line("}"); + text_.add_line(); std::string kind_str = m.kind() == moduleKind::density ? "mechanismKind::density" - : "mechanismKind::point_process"; - text_ << " mechanismKind kind() const override {\n"; - text_ << " return " << kind_str << ";\n"; - text_ << " }\n\n"; + : "mechanismKind::point"; + text_.add_line("mechanismKind kind() const override {"); + text_.increase_indentation(); + text_.add_line("return " + kind_str + ";"); + text_.decrease_indentation(); + text_.add_line("}"); + text_.add_line(); + + ////////////////////////////////////////////// + // print ion channel interface + ////////////////////////////////////////////// + // return true/false indicating if cell has dependency on k + auto const& ions = m.neuron_block().ions; + auto find_ion = [&ions] (ionKind k) { + return std::find_if( + ions.begin(), ions.end(), + [k](IonDep const& d) {return d.kind()==k;} + ); + }; + auto has_ion = [&ions, find_ion] (ionKind k) { + return find_ion(k) != ions.end(); + }; + + // bool uses_ion(ionKind k) const override + text_.add_line("bool uses_ion(ionKind k) const override {"); + text_.increase_indentation(); + text_.add_line("switch(k) {"); + text_.increase_indentation(); + text_.add_gutter() + << "case ionKind::na : return " + << (has_ion(ionKind::Na) ? "true" : "false") << ";"; + text_.end_line(); + text_.add_gutter() + << "case ionKind::ca : return " + << (has_ion(ionKind::Ca) ? "true" : "false") << ";"; + text_.end_line(); + text_.add_gutter() + << "case ionKind::k : return " + << (has_ion(ionKind::K) ? "true" : "false") << ";"; + text_.end_line(); + text_.decrease_indentation(); + text_.add_line("}"); + text_.add_line("return false;"); + text_.decrease_indentation(); + text_.add_line("}"); + text_.add_line(); + + /*************************************************************************** + * + * ion channels have the following fields : + * + * --------------------------------------------------- + * label Ca Na K name + * --------------------------------------------------- + * iX ica ina ik current + * eX eca ena ek reversal_potential + * Xi cai nai ki internal_concentration + * Xo cao nao ko external_concentration + * gX gca gna gk conductance + * --------------------------------------------------- + * + **************************************************************************/ + + // void set_ion(ionKind k, ion_type& i) override + // TODO: this is done manually, which isn't going to scale + auto has_variable = [] (IonDep const& ion, std::string const& name) { + if( std::find_if(ion.read.begin(), ion.read.end(), + [&name] (Token const& t) {return t.spelling==name;} + ) != ion.read.end() + ) return true; + if( std::find_if(ion.write.begin(), ion.write.end(), + [&name] (Token const& t) {return t.spelling==name;} + ) != ion.write.end() + ) return true; + return false; + }; + text_.add_line("void set_ion(ionKind k, ion_type& i) override {"); + text_.increase_indentation(); + text_.add_line("using nest::mc::algorithms::index_into;"); + if(has_ion(ionKind::Na)) { + auto ion = find_ion(ionKind::Na); + text_.add_line("if(k==ionKind::na) {"); + text_.increase_indentation(); + text_.add_line("ion_na.index = index_into(i.node_index(), node_index_);"); + if(has_variable(*ion, "ina")) text_.add_line("ion_na.ina = i.current();"); + if(has_variable(*ion, "ena")) text_.add_line("ion_na.ena = i.reversal_potential();"); + if(has_variable(*ion, "nai")) text_.add_line("ion_na.nai = i.internal_concentration();"); + if(has_variable(*ion, "nao")) text_.add_line("ion_na.nao = i.external_concentration();"); + text_.add_line("return;"); + text_.decrease_indentation(); + text_.add_line("}"); + } + if(has_ion(ionKind::Ca)) { + auto ion = find_ion(ionKind::Ca); + text_.add_line("if(k==ionKind::ca) {"); + text_.increase_indentation(); + text_.add_line("ion_ca.index = index_into(i.node_index(), node_index_);"); + if(has_variable(*ion, "ica")) text_.add_line("ion_ca.ica = i.current();"); + if(has_variable(*ion, "eca")) text_.add_line("ion_ca.eca = i.reversal_potential();"); + if(has_variable(*ion, "cai")) text_.add_line("ion_ca.cai = i.internal_concentration();"); + if(has_variable(*ion, "cao")) text_.add_line("ion_ca.cao = i.external_concentration();"); + text_.add_line("return;"); + text_.decrease_indentation(); + text_.add_line("}"); + } + if(has_ion(ionKind::K)) { + auto ion = find_ion(ionKind::K); + text_.add_line("if(k==ionKind::k) {"); + text_.increase_indentation(); + text_.add_line("ion_k.index = index_into(i.node_index(), node_index_);"); + if(has_variable(*ion, "ik")) text_.add_line("ion_k.ik = i.current();"); + if(has_variable(*ion, "ek")) text_.add_line("ion_k.ek = i.reversal_potential();"); + if(has_variable(*ion, "ki")) text_.add_line("ion_k.ki = i.internal_concentration();"); + if(has_variable(*ion, "ko")) text_.add_line("ion_k.ko = i.external_concentration();"); + text_.add_line("return;"); + text_.decrease_indentation(); + text_.add_line("}"); + } + text_.add_line("throw std::domain_error(nest::mc::util::pprintf(\"mechanism % does not support ion type\\n\", name()));"); + text_.decrease_indentation(); + text_.add_line("}"); + text_.add_line(); ////////////////////////////////////////////// @@ -235,47 +437,53 @@ CUDAPrinter::CUDAPrinter(Module &m, bool o) { auto proc = var.second->is_api_method(); auto name = proc->name(); - text_ << " void " << name << "() {\n"; - text_ << " auto n = size();\n"; - text_ << " auto thread_dim = 192;\n"; - text_ << " dim3 dim_block(thread_dim);\n"; - text_ << " dim3 dim_grid(n/dim_block.x + (n%dim_block.x ? 1 : 0) );\n\n"; - text_ << " START_PROFILE\n"; - text_ << " impl::" << m.name() << "::" << name << "<T,I>" - << "<<<dim_grid, dim_block>>>(param_pack_);\n"; - text_ << " STOP_PROFILE\n"; - text_ << " }\n"; + text_.add_line("void " + name + "() {"); + text_.increase_indentation(); + text_.add_line("auto n = size();"); + text_.add_line("auto thread_dim = 192;"); + text_.add_line("dim3 dim_block(thread_dim);"); + text_.add_line("dim3 dim_grid(n/dim_block.x + (n%dim_block.x ? 1 : 0) );"); + text_.add_line(); + text_.add_line( + "kernels::" + name + "<T,I>" + + "<<<dim_grid, dim_block>>>(param_pack_);"); + text_.decrease_indentation(); + text_.add_line("}"); + text_.add_line(); } } ////////////////////////////////////////////// ////////////////////////////////////////////// - //text_ << "private:\n\n"; - text_ << " vector_type data_;\n\n"; + text_.add_line("vector_type data_;"); for(auto var: array_variables) { - text_ << " view_type " << var->name() << ";\n"; + text_.add_line("view_type " + var->name() + ";"); } for(auto var: scalar_variables) { double val = var->value(); // test the default value for NaN // useful for error propogation from bad initial conditions if(val==val) { - text_ << " value_type " << var->name() << " = " << val << ";\n"; + text_.add_line("value_type " + var->name() + " = " + std::to_string(val) + ";"); } else { // the cuda compiler has a bug that doesn't allow initialization of // class members with std::numer_limites<>. So simply set to zero. - text_ << " value_type " << var->name() - << " = value_type{0};\n"; + text_.add_line("value_type " + var->name() + " = value_type{0};"); } } - text_ << " using base::matrix_;\n"; - text_ << " using base::node_indices_;\n\n"; - text_ << " param_pack_type param_pack_;\n\n"; - text_ << " DATA_PROFILE\n"; - text_ << "};\n"; + text_.add_line("using base::vec_v_;"); + text_.add_line("using base::vec_i_;"); + text_.add_line("using base::vec_area_;"); + text_.add_line("using base::node_index_;"); + text_.add_line(); + text_.add_line("param_pack_type param_pack_;"); + decrease_indentation(); + text_.add_line("};"); + decrease_indentation(); + text_.add_line("}}}}} // namespaces"); } void CUDAPrinter::visit(Expression *e) { @@ -482,7 +690,7 @@ void CUDAPrinter::visit(APIMethod *e) { text_.add_line(); text_.add_line("auto tid_ = threadIdx.x + blockDim.x*blockIdx.x;"); - text_.add_line("auto const n_ = params_.n;"); + text_.add_line("auto const n_ = params_.n_;"); text_.add_line(); text_.add_line("if(tid_<n_) {"); increase_indentation(); diff --git a/modcc/src/cudaprinter.hpp b/modcc/cudaprinter.hpp similarity index 100% rename from modcc/src/cudaprinter.hpp rename to modcc/cudaprinter.hpp diff --git a/modcc/src/error.hpp b/modcc/error.hpp similarity index 100% rename from modcc/src/error.hpp rename to modcc/error.hpp diff --git a/modcc/src/errorvisitor.cpp b/modcc/errorvisitor.cpp similarity index 100% rename from modcc/src/errorvisitor.cpp rename to modcc/errorvisitor.cpp diff --git a/modcc/src/errorvisitor.hpp b/modcc/errorvisitor.hpp similarity index 100% rename from modcc/src/errorvisitor.hpp rename to modcc/errorvisitor.hpp diff --git a/modcc/src/expression.cpp b/modcc/expression.cpp similarity index 99% rename from modcc/src/expression.cpp rename to modcc/expression.cpp index 96fc173d7576084f1bdb1547d7c8b4e78f94d9bb..06a70e0c99bb7a9d9c38d472f2b2382e0f44371d 100644 --- a/modcc/src/expression.cpp +++ b/modcc/expression.cpp @@ -22,9 +22,9 @@ inline std::string to_string(symbolKind k) { inline std::string to_string(procedureKind k) { switch(k) { - case procedureKind::normal : + case procedureKind::normal : return "procedure"; - case procedureKind::api : + case procedureKind::api : return "APIprocedure"; case procedureKind::initial : return "initial"; @@ -702,12 +702,14 @@ expression_ptr IfExpression::clone() const { } #include "visitor.hpp" + /* Visitor hooks */ void Expression::accept(Visitor *v) { v->visit(this); } + void Symbol::accept(Visitor *v) { v->visit(this); } diff --git a/modcc/src/expression.hpp b/modcc/expression.hpp similarity index 100% rename from modcc/src/expression.hpp rename to modcc/expression.hpp diff --git a/modcc/src/expressionclassifier.cpp b/modcc/expressionclassifier.cpp similarity index 99% rename from modcc/src/expressionclassifier.cpp rename to modcc/expressionclassifier.cpp index 9b06f14fb1bc836feb9a7e0a15fd16d651537b97..29008ec7674cf166d20ec2f005c127b7d7c176d4 100644 --- a/modcc/src/expressionclassifier.cpp +++ b/modcc/expressionclassifier.cpp @@ -320,3 +320,4 @@ void ExpressionClassifierVisitor::visit(CallExpression *e) { } } } + diff --git a/modcc/src/expressionclassifier.hpp b/modcc/expressionclassifier.hpp similarity index 100% rename from modcc/src/expressionclassifier.hpp rename to modcc/expressionclassifier.hpp diff --git a/modcc/src/functionexpander.cpp b/modcc/functionexpander.cpp similarity index 99% rename from modcc/src/functionexpander.cpp rename to modcc/functionexpander.cpp index dd32ba7662e952799ec9697c4d3407b853456e62..1ebd9d4908824ab85dc122f283679b3d52fb3672 100644 --- a/modcc/src/functionexpander.cpp +++ b/modcc/functionexpander.cpp @@ -162,3 +162,4 @@ lower_function_arguments(std::vector<expression_ptr>& args) return new_statements; } + diff --git a/modcc/src/functionexpander.hpp b/modcc/functionexpander.hpp similarity index 100% rename from modcc/src/functionexpander.hpp rename to modcc/functionexpander.hpp diff --git a/modcc/src/functioninliner.cpp b/modcc/functioninliner.cpp similarity index 100% rename from modcc/src/functioninliner.cpp rename to modcc/functioninliner.cpp diff --git a/modcc/src/functioninliner.hpp b/modcc/functioninliner.hpp similarity index 100% rename from modcc/src/functioninliner.hpp rename to modcc/functioninliner.hpp diff --git a/modcc/src/identifier.hpp b/modcc/identifier.hpp similarity index 100% rename from modcc/src/identifier.hpp rename to modcc/identifier.hpp diff --git a/modcc/src/lexer.cpp b/modcc/lexer.cpp similarity index 93% rename from modcc/src/lexer.cpp rename to modcc/lexer.cpp index 8846da84aefad764976342a269081bcfd7fe7748..b4b8ccf936b7d028f79353afbbb4c05864c4478b 100644 --- a/modcc/src/lexer.cpp +++ b/modcc/lexer.cpp @@ -87,8 +87,7 @@ Token Lexer::parse() { continue; // number - case '0': case '1' : case '2' : case '3' : case '4': - case '5': case '6' : case '7' : case '8' : case '9': + case '0' ... '9': case '.': t.spelling = number(); @@ -97,14 +96,8 @@ Token Lexer::parse() { return t; // identifier or keyword - case 'a': case 'b': case 'c': case 'd': case 'e': case 'f': case 'g': - case 'h': case 'i': case 'j': case 'k': case 'l': case 'm': case 'n': - case 'o': case 'p': case 'q': case 'r': case 's': case 't': case 'u': - case 'v': case 'w': case 'x': case 'y': case 'z': - case 'A': case 'B': case 'C': case 'D': case 'E': case 'F': case 'G': - case 'H': case 'I': case 'J': case 'K': case 'L': case 'M': case 'N': - case 'O': case 'P': case 'Q': case 'R': case 'S': case 'T': case 'U': - case 'V': case 'W': case 'X': case 'Y': case 'Z': + case 'a' ... 'z': + case 'A' ... 'Z': case '_': // get std::string of the identifier t.spelling = identifier(); diff --git a/modcc/src/lexer.hpp b/modcc/lexer.hpp similarity index 100% rename from modcc/src/lexer.hpp rename to modcc/lexer.hpp diff --git a/modcc/src/location.hpp b/modcc/location.hpp similarity index 100% rename from modcc/src/location.hpp rename to modcc/location.hpp diff --git a/modcc/src/mechanism.hpp b/modcc/mechanism.hpp similarity index 100% rename from modcc/src/mechanism.hpp rename to modcc/mechanism.hpp diff --git a/modcc/src/memop.hpp b/modcc/memop.hpp similarity index 99% rename from modcc/src/memop.hpp rename to modcc/memop.hpp index a2819cad103a2ef4b93105eab66335f888b0f6c8..e962a78548f1e3b20c745bbf69c3b716fe36fdc8 100644 --- a/modcc/src/memop.hpp +++ b/modcc/memop.hpp @@ -35,3 +35,4 @@ struct MemOp { } } }; + diff --git a/modcc/src/modcc.cpp b/modcc/modcc.cpp similarity index 98% rename from modcc/src/modcc.cpp rename to modcc/modcc.cpp index 423f855893a8b8fbb606676c58b1a8eb4a57ab41..eb2b29e9351e927b7c32d0a24ffd0ba464d1a547 100644 --- a/modcc/src/modcc.cpp +++ b/modcc/modcc.cpp @@ -123,12 +123,15 @@ int main(int argc, char **argv) { //////////////////////////////////////////////////////////// // semantic analysis //////////////////////////////////////////////////////////// - if(options.verbose) std::cout << green("[") + "semantic analysis" + green("]") << std::endl; + if(options.verbose) + std::cout << green("[") + "semantic analysis" + green("]") << "\n"; + m.semantic(); if( m.has_error() || m.has_warning() ) { std::cout << m.error_string() << std::endl; } + if(m.status() == lexerStatus::error) { return 1; } diff --git a/modcc/src/modccutil.hpp b/modcc/modccutil.hpp similarity index 100% rename from modcc/src/modccutil.hpp rename to modcc/modccutil.hpp diff --git a/modcc/src/module.cpp b/modcc/module.cpp similarity index 98% rename from modcc/src/module.cpp rename to modcc/module.cpp index d00d90d27c190c0fb74165fe5644dd8d0c15ea7a..6001e1b458f8cf16ffa488ca10ceb4e46268787e 100644 --- a/modcc/src/module.cpp +++ b/modcc/module.cpp @@ -106,7 +106,7 @@ bool Module::semantic() { // Helper which iterates over a vector of Symbols, moving them into the // symbol table. - // Returns false if a symbol name clases with the name of a symbol that + // Returns false if a symbol name clashes with the name of a symbol that // is already in the symbol table. auto move_symbols = [this] (std::vector<symbol_ptr>& symbol_list) { for(auto& symbol: symbol_list) { @@ -149,8 +149,9 @@ bool Module::semantic() { || s->kind() == symbolKind::procedure) { #ifdef LOGGING - std::cout << "\nfunction inlining for " << s->location() << "\n" << s->to_string() << "\n"; - std::cout << green("\n-call site lowering-\n\n"); + std::cout << "\nfunction inlining for " << s->location() << "\n" + << s->to_string() << "\n" + << green("\n-call site lowering-\n\n"); #endif // first perform semantic analysis s->semantic(symbols_); @@ -338,6 +339,7 @@ bool Module::semantic() { } } } + a = b = nullptr; return std::make_pair(a, b); }; @@ -347,7 +349,7 @@ bool Module::semantic() { // There are two APIMethods generated from BREAKPOINT. // The first is nrn_state, which is the first case handled below. // The second is nrn_current, which is handled after this block - auto state_api = make_empty_api_method("nrn_state", "breakpoint"); + auto state_api = make_empty_api_method("nrn_state", "breakpoint"); auto api_state = state_api.first; auto breakpoint = state_api.second; @@ -774,7 +776,7 @@ bool Module::optimize() { auto kind = symbol.second->kind(); BlockExpression* body; if(kind == symbolKind::procedure) { - // we are only interested in true procedurs and APIMethods + // we are only interested in true procedures and APIMethods auto proc = symbol.second->is_procedure(); auto pkind = proc->kind(); if(pkind == procedureKind::normal || pkind == procedureKind::api ) @@ -811,4 +813,3 @@ bool Module::optimize() { return true; } - diff --git a/modcc/src/module.hpp b/modcc/module.hpp similarity index 100% rename from modcc/src/module.hpp rename to modcc/module.hpp diff --git a/modcc/src/parser.cpp b/modcc/parser.cpp similarity index 97% rename from modcc/src/parser.cpp rename to modcc/parser.cpp index a0884900f2dee855a2ddbc12c47d7306ad278383..0b4b81068ed22111b6c28ed04313d245ce613f4a 100644 --- a/modcc/src/parser.cpp +++ b/modcc/parser.cpp @@ -449,6 +449,7 @@ void Parser::parse_parameter_block() { return; } + int success = 1; // there are no use cases for curly brace in a UNITS block, so we don't have to count them get_token(); while(token_.type!=tok::rbrace && token_.type!=tok::eof) { @@ -457,7 +458,8 @@ void Parser::parse_parameter_block() { // read the parameter name if(token_.type != tok::identifier) { - goto parm_error; + success = 0; + goto parm_exit; } parm.token = token_; // save full token @@ -471,7 +473,8 @@ void Parser::parse_parameter_block() { get_token(); } if(token_.type != tok::number) { - goto parm_error; + success = 0; + goto parm_exit; } parm.value += token_.spelling; // store value as a string get_token(); @@ -481,27 +484,27 @@ void Parser::parse_parameter_block() { if(line==location_.line && token_.type == tok::lparen) { parm.units = unit_description(); if(status_ == lexerStatus::error) { - goto parm_error; + success = 0; + goto parm_exit; } } block.parameters.push_back(parm); } - // errer if EOF before closeing curly brace + // error if EOF before closing curly brace if(token_.type==tok::eof) { error("PARAMETER block must have closing '}'"); - goto parm_error; + goto parm_exit; } get_token(); // consume closing brace module_->parameter_block(block); - return; -parm_error: +parm_exit: // only write error message if one hasn't already been logged by the lexer - if(status_==lexerStatus::happy) { + if(!success && status_==lexerStatus::happy) { error(pprintf("PARAMETER block unexpected symbol '%'", token_.spelling)); } return; @@ -518,6 +521,8 @@ void Parser::parse_assigned_block() { return; } + int success = 1; + // there are no use cases for curly brace in an ASSIGNED block, so we don't have to count them get_token(); while(token_.type!=tok::rbrace && token_.type!=tok::eof) { @@ -526,7 +531,8 @@ void Parser::parse_assigned_block() { // the first token must be ... if(token_.type != tok::identifier) { - goto ass_error; + success = 0; + goto ass_exit; } // read all of the identifiers until we run out of identifiers or reach a new line while(token_.type == tok::identifier && line == location_.line) { @@ -538,7 +544,8 @@ void Parser::parse_assigned_block() { if(line==location_.line && token_.type == tok::lparen) { auto u = unit_description(); if(status_ == lexerStatus::error) { - goto ass_error; + success = 0; + goto ass_exit; } for(auto const& t : variables) { block.parameters.push_back(Id(t, "", u)); @@ -551,20 +558,19 @@ void Parser::parse_assigned_block() { } } - // errer if EOF before closeing curly brace + // error if EOF before closing curly brace if(token_.type==tok::eof) { error("ASSIGNED block must have closing '}'"); - goto ass_error; + goto ass_exit; } get_token(); // consume closing brace module_->assigned_block(block); - return; -ass_error: +ass_exit: // only write error message if one hasn't already been logged by the lexer - if(status_==lexerStatus::happy) { + if(!success && status_==lexerStatus::happy) { error(pprintf("ASSIGNED block unexpected symbol '%'", token_.spelling)); } return; @@ -575,15 +581,20 @@ std::vector<Token> Parser::unit_description() { int startline = location_.line; std::vector<Token> tokens; - // chec that we start with a left parenthesis - if(token_.type != tok::lparen) - goto unit_error; + // check that we start with a left parenthesis + if(token_.type != tok::lparen) { + error(pprintf("unit description must start with a parenthesis '%'", tokens)); + goto unit_exit; + } + get_token(); while(token_.type != tok::rparen) { // check for illegal tokens or a new line - if( !is_in(token_.type,legal_tokens) || startline < location_.line ) - goto unit_error; + if( !is_in(token_.type,legal_tokens) || startline < location_.line ) { + error(pprintf("incorrect unit description '%'", tokens)); + goto unit_exit; + } // add this token to the set tokens.push_back(token_); @@ -592,10 +603,7 @@ std::vector<Token> Parser::unit_description() { // remove trailing right parenthesis ')' get_token(); - return tokens; - -unit_error: - error(pprintf("incorrect unit description '%'", tokens)); +unit_exit: return tokens; } diff --git a/modcc/src/parser.hpp b/modcc/parser.hpp similarity index 100% rename from modcc/src/parser.hpp rename to modcc/parser.hpp diff --git a/modcc/src/perfvisitor.hpp b/modcc/perfvisitor.hpp similarity index 99% rename from modcc/src/perfvisitor.hpp rename to modcc/perfvisitor.hpp index 67ed51afd057cc81ce9c5f7642a645da3ced74a9..694a9284dfe977e02bd1d1cbb3b5935600223905 100644 --- a/modcc/src/perfvisitor.hpp +++ b/modcc/perfvisitor.hpp @@ -67,6 +67,10 @@ public: // leave UnaryExpression to throw, to catch // any missed specializations //////////////////////////////////////////////////// + void visit(UnaryExpression *e) override { + // do nothing + } + void visit(NegUnaryExpression *e) override { // this is a simplification // we would have to perform analysis of parent nodes to ensure that @@ -259,4 +263,3 @@ private: std::set<Symbol*> indexed_writes_; std::set<Symbol*> vector_writes_; }; - diff --git a/modcc/src/scope.hpp b/modcc/scope.hpp similarity index 99% rename from modcc/src/scope.hpp rename to modcc/scope.hpp index 6a55d6ea982dc6436a9938c60984e46b92347fcf..ef70040073d84adae6c33e07f0565180c18a9f8e 100644 --- a/modcc/src/scope.hpp +++ b/modcc/scope.hpp @@ -124,3 +124,4 @@ typename Scope<Symbol>::symbol_map* Scope<Symbol>::globals() { return global_symbols_; } + diff --git a/modcc/src/CMakeLists.txt b/modcc/src/CMakeLists.txt deleted file mode 100644 index aa98b8617c6581eee3fe386eaaf168b964c56d96..0000000000000000000000000000000000000000 --- a/modcc/src/CMakeLists.txt +++ /dev/null @@ -1,27 +0,0 @@ -set(MODCC_SOURCES - token.cpp - lexer.cpp - expression.cpp - parser.cpp - textbuffer.cpp - cprinter.cpp - functionexpander.cpp - functioninliner.cpp - cudaprinter.cpp - expressionclassifier.cpp - constantfolder.cpp - errorvisitor.cpp - module.cpp -) - -add_library(compiler ${MODCC_SOURCES}) - -add_executable(modcc modcc.cpp) - -target_link_libraries(modcc LINK_PUBLIC compiler) - -set_target_properties(modcc - PROPERTIES - RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/modcc" -) - diff --git a/modcc/src/textbuffer.cpp b/modcc/textbuffer.cpp similarity index 100% rename from modcc/src/textbuffer.cpp rename to modcc/textbuffer.cpp diff --git a/modcc/src/textbuffer.hpp b/modcc/textbuffer.hpp similarity index 100% rename from modcc/src/textbuffer.hpp rename to modcc/textbuffer.hpp diff --git a/modcc/src/token.cpp b/modcc/token.cpp similarity index 100% rename from modcc/src/token.cpp rename to modcc/token.cpp diff --git a/modcc/src/token.hpp b/modcc/token.hpp similarity index 100% rename from modcc/src/token.hpp rename to modcc/token.hpp diff --git a/modcc/src/visitor.hpp b/modcc/visitor.hpp similarity index 76% rename from modcc/src/visitor.hpp rename to modcc/visitor.hpp index de90fdbc6271bcacef545d5d3f4f635b78c27d22..faab5cca2500e7a144263a3402364284860bc227 100644 --- a/modcc/src/visitor.hpp +++ b/modcc/visitor.hpp @@ -15,43 +15,35 @@ /// heavily inspired by the DMD D compiler : github.com/D-Programming-Language/dmd class Visitor { public: - virtual void visit(Expression *e) { - throw compiler_exception("unimplemented visitor", Location()); - } - virtual void visit(Symbol *e) { visit((Expression*) e); } - virtual void visit(LocalVariable *e) { visit((Expression*) e); } - virtual void visit(IdentifierExpression *e) { visit((Expression*) e); } - virtual void visit(NumberExpression *e) { visit((Expression*) e); } - virtual void visit(LocalDeclaration *e) { visit((Expression*) e); } - virtual void visit(ArgumentExpression *e) { visit((Expression*) e); } - virtual void visit(PrototypeExpression *e) { visit((Expression*) e); } - virtual void visit(CallExpression *e) { visit((Expression*) e); } - virtual void visit(VariableExpression *e) { visit((Expression*) e); } - virtual void visit(IndexedVariable *e) { visit((Expression*) e); } - virtual void visit(FunctionExpression *e) { visit((Expression*) e); } - virtual void visit(IfExpression *e) { visit((Expression*) e); } - virtual void visit(SolveExpression *e) { visit((Expression*) e); } - virtual void visit(DerivativeExpression *e) { visit((Expression*) e); } - - virtual void visit(ProcedureExpression *e) { visit((Expression*) e); } + virtual void visit(Expression *e) = 0; + virtual void visit(Symbol *e) { visit((Expression*) e); } + virtual void visit(LocalVariable *e) { visit((Expression*) e); } + virtual void visit(IdentifierExpression *e) { visit((Expression*) e); } + virtual void visit(NumberExpression *e) { visit((Expression*) e); } + virtual void visit(LocalDeclaration *e) { visit((Expression*) e); } + virtual void visit(ArgumentExpression *e) { visit((Expression*) e); } + virtual void visit(PrototypeExpression *e) { visit((Expression*) e); } + virtual void visit(CallExpression *e) { visit((Expression*) e); } + virtual void visit(VariableExpression *e) { visit((Expression*) e); } + virtual void visit(IndexedVariable *e) { visit((Expression*) e); } + virtual void visit(FunctionExpression *e) { visit((Expression*) e); } + virtual void visit(IfExpression *e) { visit((Expression*) e); } + virtual void visit(SolveExpression *e) { visit((Expression*) e); } + virtual void visit(DerivativeExpression *e) { visit((Expression*) e); } + virtual void visit(ProcedureExpression *e) { visit((Expression*) e); } virtual void visit(NetReceiveExpression *e) { visit((ProcedureExpression*) e); } - virtual void visit(APIMethod *e) { visit((Expression*) e); } - - virtual void visit(BlockExpression *e) { visit((Expression*) e); } - virtual void visit(InitialBlock *e) { visit((BlockExpression*) e); } + virtual void visit(APIMethod *e) { visit((Expression*) e); } + virtual void visit(BlockExpression *e) { visit((Expression*) e); } + virtual void visit(InitialBlock *e) { visit((BlockExpression*) e); } - virtual void visit(UnaryExpression *e) { - throw compiler_exception("unimplemented visitor (UnaryExpression)", Location()); - } - virtual void visit(NegUnaryExpression *e) { visit((UnaryExpression*) e); } - virtual void visit(ExpUnaryExpression *e) { visit((UnaryExpression*) e); } - virtual void visit(LogUnaryExpression *e) { visit((UnaryExpression*) e); } - virtual void visit(CosUnaryExpression *e) { visit((UnaryExpression*) e); } - virtual void visit(SinUnaryExpression *e) { visit((UnaryExpression*) e); } + virtual void visit(UnaryExpression *e) = 0; + virtual void visit(NegUnaryExpression *e) { visit((UnaryExpression*) e); } + virtual void visit(ExpUnaryExpression *e) { visit((UnaryExpression*) e); } + virtual void visit(LogUnaryExpression *e) { visit((UnaryExpression*) e); } + virtual void visit(CosUnaryExpression *e) { visit((UnaryExpression*) e); } + virtual void visit(SinUnaryExpression *e) { visit((UnaryExpression*) e); } - virtual void visit(BinaryExpression *e) { - throw compiler_exception("unimplemented visitor (BinaryExpression)", Location()); - } + virtual void visit(BinaryExpression *e) = 0; virtual void visit(AssignmentExpression *e) { visit((BinaryExpression*) e); } virtual void visit(AddBinaryExpression *e) { visit((BinaryExpression*) e); } virtual void visit(SubBinaryExpression *e) { visit((BinaryExpression*) e); } diff --git a/tests/modcc/CMakeLists.txt b/tests/modcc/CMakeLists.txt index d17002a6d7cd7e08e2a1cd9a921462d220660b16..044eec0a09c3ebdb530fd7b8196a52024da6127e 100644 --- a/tests/modcc/CMakeLists.txt +++ b/tests/modcc/CMakeLists.txt @@ -1,36 +1,22 @@ -set(HEADERS - ${PROJECT_SOURCE_DIR}/modcc/src/constantfolder.hpp - ${PROJECT_SOURCE_DIR}/modcc/src/cprinter.hpp - ${PROJECT_SOURCE_DIR}/modcc/src/expressionclassifier.hpp - ${PROJECT_SOURCE_DIR}/modcc/src/lexer.hpp - ${PROJECT_SOURCE_DIR}/modcc/src/module.hpp - ${PROJECT_SOURCE_DIR}/modcc/src/perfvisitor.hpp - ${PROJECT_SOURCE_DIR}/modcc/src/parser.hpp - ${PROJECT_SOURCE_DIR}/src/util.hpp - ${PROJECT_SOURCE_DIR}/modcc/src/modccutil.hpp -# ${PROJECT_SOURCE_DIR}/modcc/src/variablerenamer.hpp - test.hpp -) - set(MODCC_TEST_SOURCES - # unit tests - test_lexer.cpp - test_module.cpp - test_optimization.cpp - test_parser.cpp - #test_printers.cpp - test_visitors.cpp + # unit tests + test_lexer.cpp + test_module.cpp + test_optimization.cpp + test_parser.cpp + #test_printers.cpp + test_visitors.cpp - # unit test driver - driver.cpp + # unit test driver + driver.cpp ) add_definitions("-DDATADIR=\"${CMAKE_SOURCE_DIR}/data\"") -add_executable(test_modcc ${MODCC_TEST_SOURCES} ${HEADERS}) +add_executable(test_modcc ${MODCC_TEST_SOURCES}) target_link_libraries(test_modcc LINK_PUBLIC compiler gtest) set_target_properties(test_modcc - PROPERTIES - RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/tests" + PROPERTIES + RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/tests" ) diff --git a/tests/modcc/test.hpp b/tests/modcc/test.hpp index e0b99219420ce4a36e744d4d2c9ce101e9d79d2f..588b75a36afcb09107adc65b862fd65eceb5a9c6 100644 --- a/tests/modcc/test.hpp +++ b/tests/modcc/test.hpp @@ -12,18 +12,18 @@ #define VERBOSE_PRINT(x) #endif -static expression_ptr parse_line_expression(std::string const& s) { +inline expression_ptr parse_line_expression(std::string const& s) { return Parser(s).parse_line_expression(); } -static expression_ptr parse_expression(std::string const& s) { +inline expression_ptr parse_expression(std::string const& s) { return Parser(s).parse_expression(); } -static expression_ptr parse_function(std::string const& s) { +inline expression_ptr parse_function(std::string const& s) { return Parser(s).parse_function(); } -static expression_ptr parse_procedure(std::string const& s) { +inline expression_ptr parse_procedure(std::string const& s) { return Parser(s).parse_procedure(); }