 |
User Guide
|
|
User Guide
|
Visitor for printing C++ code with OpenACC backend More...
Visitor for printing C++ code with OpenACC backend
Definition at line 30 of file codegen_acc_visitor.hpp.
#include <codegen_acc_visitor.hpp>
Protected Member Functions | |
| std::string | backend_name () const override |
| name of the code generation backend More... | |
| void | print_backend_includes () override |
| common includes : standard c++, coreneuron and backend specific More... | |
| void | print_channel_iteration_block_parallel_hint (BlockType type, const ast::Block *block) override |
| ivdep like annotation for channel iterations More... | |
| void | print_atomic_reduction_pragma () override |
| atomic update pragma for reduction statements More... | |
| void | print_memory_allocation_routine () const override |
| memory allocation routine More... | |
| void | print_abort_routine () const override |
| abort routine More... | |
| void | print_kernel_data_present_annotation_block_begin () override |
| annotations like "acc enter data present(...)" for main kernel More... | |
| void | print_kernel_data_present_annotation_block_end () override |
| end of annotation like "acc enter data" More... | |
| void | print_net_init_acc_serial_annotation_block_begin () override |
| start of annotation "acc kernels" for net_init kernel More... | |
| void | print_net_init_acc_serial_annotation_block_end () override |
| end of annotation "acc kernels" for net_init kernel More... | |
| void | print_nrn_cur_matrix_shadow_update () override |
| update to matrix elements with/without shadow vectors More... | |
| void | print_nrn_cur_matrix_shadow_reduction () override |
| reduction to matrix elements from shadow vectors More... | |
| void | print_fast_imem_calculation () override |
| fast membrane current calculation More... | |
| void | print_rhs_d_shadow_variables () override |
| setup method for setting matrix shadow vectors More... | |
| bool | nrn_cur_reduction_loop_required () override |
| if reduction block in nrn_cur required More... | |
| void | print_global_variable_device_update_annotation () override |
| update global variable from host to the device More... | |
| void | print_newtonspace_transfer_to_device () const override |
| transfer newtonspace structure to device More... | |
| void | print_instance_struct_transfer_routine_declarations () override |
| declare helper functions for copying the instance struct to the device More... | |
| void | print_instance_struct_transfer_routines (const std::vector< std::string > &ptr_members) override |
| define helper functions for copying the instance struct to the device More... | |
| void | print_instance_struct_copy_to_device () override |
| call helper function for copying the instance struct to the device More... | |
| void | print_instance_struct_delete_from_device () override |
| call helper function that deletes the instance struct from the device More... | |
| void | print_deriv_advance_flag_transfer_to_device () const override |
| update derivimplicit advance flag on the gpu device More... | |
| void | print_net_send_buf_count_update_to_host () const override |
| update NetSendBuffer_t count from device to host More... | |
| void | print_net_send_buf_update_to_host () const override |
| update NetSendBuffer_t from device to host More... | |
| virtual void | print_net_send_buf_count_update_to_device () const override |
| update NetSendBuffer_t count from host to device More... | |
| virtual void | print_dt_update_to_device () const override |
| update dt from host to device More... | |
| virtual void | print_device_stream_wait () const override |
| Print the code to synchronise/wait on stream specific to NrnThread. More... | |
| void | print_device_atomic_capture_annotation () const override |
| print atomic capture pragma More... | |
| void | print_net_send_buffering_cnt_update () const override |
| print atomic update of NetSendBuffer_t cnt More... | |
| void | print_net_send_buffering_grow () override |
| Replace default implementation by a no-op since the buffer cannot be grown up during gpu execution. More... | |
 Protected Member Functions inherited from nmodl::codegen::CodegenCoreneuronCppVisitor | |
| std::string | simulator_name () override |
| Name of the simulator the code was generated for. More... | |
| int | num_thread_objects () const noexcept |
| Determine the number of threads to allocate. More... | |
| int | position_of_float_var (const std::string &name) const override |
| Determine the position in the data array for a given float variable. More... | |
| int | position_of_int_var (const std::string &name) const override |
| Determine the position in the data array for a given int variable. More... | |
| std::string | process_verbatim_token (const std::string &token) |
| Process a token in a verbatim block for possible variable renaming. More... | |
| virtual bool | is_constant_variable (const std::string &name) const |
| Check if variable is qualified as constant. More... | |
| virtual void | print_device_method_annotation () |
| Print the backend specific device method annotation. More... | |
| virtual void | print_global_method_annotation () |
| Print backend specific global method annotation. More... | |
| virtual void | print_backend_namespace_start () |
| Prints the start of namespace for the backend-specific code. More... | |
| virtual void | print_backend_namespace_stop () |
| Prints the end of namespace for the backend-specific code. More... | |
| bool | optimize_ion_variable_copies () const override |
| Check if ion variable copies should be avoided. More... | |
| void | print_top_verbatim_blocks () |
| Print top level (global scope) verbatim blocks. More... | |
| void | print_function_prototypes () override |
| Print function and procedures prototype declaration. More... | |
| std::tuple< bool, int > | check_if_var_is_array (const std::string &name) |
| Check if the given name exist in the symbol. More... | |
| void | print_table_check_function (const ast::Block &node) |
Print check_function() for functions or procedure using table. More... | |
| void | print_table_replacement_function (const ast::Block &node) |
| Print replacement function for function or procedure using table. More... | |
| void | print_check_table_thread_function () |
| Print check_table functions. More... | |
| void | print_function_or_procedure (const ast::Block &node, const std::string &name) override |
| Print nmodl function or procedure (common code) More... | |
| void | print_function_procedure_helper (const ast::Block &node) override |
| Common helper function to help printing function or procedure blocks. More... | |
| virtual void | print_procedure (const ast::ProcedureBlock &node) override |
| Print NMODL procedure in target backend code. More... | |
| void | print_function (const ast::FunctionBlock &node) override |
| Print NMODL function in target backend code. More... | |
| void | print_function_tables (const ast::FunctionTableBlock &node) |
| Print NMODL function_table in target backend code. More... | |
| bool | is_functor_const (const ast::StatementBlock &variable_block, const ast::StatementBlock &functor_block) |
| Checks whether the functor_block generated by sympy solver modifies any variable outside its scope. More... | |
| void | print_functor_definition (const ast::EigenNewtonSolverBlock &node) |
Based on the EigenNewtonSolverBlock passed print the definition needed for its functor. More... | |
| virtual void | print_eigen_linear_solver (const std::string &float_type, int N) |
| std::string | internal_method_arguments () override |
| Arguments for functions that are defined and used internally. More... | |
| ParamVector | internal_method_parameters () override |
| Parameters for internally defined functions. More... | |
| const char * | external_method_arguments () noexcept override |
| Arguments for external functions called from generated code. More... | |
| const char * | external_method_parameters (bool table=false) noexcept override |
| Parameters for functions in generated code that are called back from external code. More... | |
| std::string | nrn_thread_arguments () const override |
| Arguments for "_threadargs_" macro in neuron implementation. More... | |
| std::string | nrn_thread_internal_arguments () override |
| Arguments for "_threadargs_" macro in neuron implementation. More... | |
| std::string | replace_if_verbatim_variable (std::string name) |
| Replace commonly used verbatim variables. More... | |
| std::string | process_verbatim_text (std::string const &text) override |
| Process a verbatim block for possible variable renaming. More... | |
| std::string | register_mechanism_arguments () const override |
| Arguments for register_mech or point_register_mech function. More... | |
| std::string | conc_write_statement (const std::string &ion_name, const std::string &concentration, int index) override |
| Generate Function call statement for nrn_wrote_conc. More... | |
| void | print_first_pointer_var_index_getter () |
| Print the getter method for index position of first pointer variable. More... | |
| void | print_first_random_var_index_getter () |
| Print the getter method for index position of first RANDOM variable. More... | |
| void | print_num_variable_getter () |
| Print the getter methods for float and integer variables count. More... | |
| void | print_net_receive_arg_size_getter () |
| Print the getter method for getting number of arguments for net_receive. More... | |
| void | print_mech_type_getter () |
| Print the getter method for returning mechtype. More... | |
| void | print_memb_list_getter () |
| Print the getter method for returning membrane list from NrnThread. More... | |
| void | print_namespace_start () override |
Prints the start of the coreneuron namespace. More... | |
| void | print_namespace_stop () override |
Prints the end of the coreneuron namespace. More... | |
| void | print_thread_getters () |
| Print the getter method for thread variables and ids. More... | |
| std::string | float_variable_name (const SymbolType &symbol, bool use_instance) const override |
Determine the name of a float variable given its symbol. More... | |
| std::string | int_variable_name (const IndexVariableInfo &symbol, const std::string &name, bool use_instance) const override |
Determine the name of an int variable given its symbol. More... | |
| std::string | global_variable_name (const SymbolType &symbol, bool use_instance=true) const override |
| Determine the variable name for a global variable given its symbol. More... | |
| std::string | get_variable_name (const std::string &name, bool use_instance=true) const override |
| Determine variable name in the structure of mechanism properties. More... | |
| void | print_backend_info () override |
| Print top file header printed in generated code. More... | |
| void | print_standard_includes () override |
| Print standard C/C++ includes. More... | |
| void | print_coreneuron_includes () |
| Print includes from coreneuron. More... | |
| void | print_sdlists_init (bool print_initializers) override |
| void | print_mechanism_global_var_structure (bool print_initializers) override |
| Print the structure that wraps all global variables used in the NMODL. More... | |
| void | print_global_variables_for_hoc () override |
| Print byte arrays that register scalar and vector variables for hoc interface. More... | |
| void | print_mechanism_register () override |
| Print the mechanism registration function. More... | |
| void | print_thread_memory_callbacks () |
| Print thread related memory allocation and deallocation callbacks. More... | |
| void | print_ion_var_structure () |
| Print structure of ion variables used for local copies. More... | |
| virtual void | print_ion_var_constructor (const std::vector< std::string > &members) |
| Print constructor of ion variables. More... | |
| virtual void | print_ion_variable () |
| Print the ion variable struct. More... | |
| void | print_setup_range_variable () |
| Print the function that initialize range variable with different data type. More... | |
| std::string | get_range_var_float_type (const SymbolType &symbol) |
| Returns floating point type for given range variable symbol. More... | |
| void | print_initial_block (const ast::InitialBlock *node) |
| Print initial block statements. More... | |
| virtual void | print_global_function_common_code (BlockType type, const std::string &function_name="") override |
| Print common code for global functions like nrn_init, nrn_cur and nrn_state. More... | |
| void | print_nrn_init (bool skip_init_check=true) |
Print the nrn_init function definition. More... | |
| virtual void | print_before_after_block (const ast::Block *node, size_t block_id) |
| Print NMODL before / after block in target backend code. More... | |
| void | print_nrn_constructor () override |
| Print nrn_constructor function definition. More... | |
| void | print_nrn_destructor () override |
| Print nrn_destructor function definition. More... | |
| void | print_functors_definitions () |
Go through the map of EigenNewtonSolverBlock s and their corresponding functor names and print the functor definitions before the definitions of the functions of the generated file. More... | |
| void | print_nrn_alloc () override |
| Print nrn_alloc function definition. More... | |
| void | print_watch_activate () |
| Print watch activate function. More... | |
| void | print_watch_check () |
| Print watch activate function. More... | |
| void | print_net_receive_common_code (const ast::Block &node, bool need_mech_inst=true) |
| Print the common code section for net receive related methods. More... | |
| void | print_net_send_call (const ast::FunctionCall &node) override |
Print call to net_send. More... | |
| void | print_net_move_call (const ast::FunctionCall &node) override |
| Print call to net_move. More... | |
| void | print_net_event_call (const ast::FunctionCall &node) override |
| Print call to net_event. More... | |
| void | print_net_init () |
| Print initial block in the net receive block. More... | |
| void | print_send_event_move () |
| Print send event move block used in net receive as well as watch. More... | |
| virtual std::string | net_receive_buffering_declaration () |
Generate the target backend code for the net_receive_buffering function delcaration. More... | |
| virtual void | print_get_memb_list () |
Print the target backend code for defining and checking a local Memb_list variable. More... | |
| virtual void | print_net_receive_loop_begin () |
Print the code for the main net_receive loop. More... | |
| virtual void | print_net_receive_loop_end () |
Print the code for closing the main net_receive loop. More... | |
| void | print_net_receive_buffering (bool need_mech_inst=true) |
| Print kernel for buffering net_receive events. More... | |
| void | print_net_send_buffering () |
| Print kernel for buffering net_send events. More... | |
| void | print_net_receive_kernel () |
Print net_receive kernel function definition. More... | |
| void | print_net_receive () |
Print net_receive function definition. More... | |
| void | print_derivimplicit_kernel (const ast::Block &block) |
Print derivative kernel when derivimplicit method is used. More... | |
| void | print_nrn_state () override |
| Print nrn_state / state update function definition. More... | |
| void | print_nrn_current (const ast::BreakpointBlock &node) override |
Print the nrn_current kernel. More... | |
| void | print_nrn_cur_conductance_kernel (const ast::BreakpointBlock &node) override |
Print the nrn_cur kernel with NMODL conductance keyword provisions. More... | |
| void | print_nrn_cur_non_conductance_kernel () override |
Print the nrn_cur kernel without NMODL conductance keyword provisions. More... | |
| void | print_nrn_cur_kernel (const ast::BreakpointBlock &node) override |
| Print main body of nrn_cur function. More... | |
| void | print_nrn_cur () override |
| Print nrn_cur / current update function definition. More... | |
| void | print_headers_include () override |
| Print all includes. More... | |
| void | print_namespace_begin () override |
| Print start of namespaces. More... | |
| void | print_namespace_end () override |
| Print end of namespaces. More... | |
| void | print_common_getters () |
| Print common getters. More... | |
| void | print_data_structures (bool print_initializers) override |
| Print all classes. More... | |
| void | print_v_unused () const override |
| Set v_unused (voltage) for NRN_PRCELLSTATE feature. More... | |
| void | print_g_unused () const override |
| Set g_unused (conductance) for NRN_PRCELLSTATE feature. More... | |
| virtual void | print_compute_functions () override |
| Print all compute functions for every backend. More... | |
| virtual void | print_codegen_routines () override |
| Print entry point to code generation. More... | |
| void | visit_derivimplicit_callback (const ast::DerivimplicitCallback &node) override |
| visit node of type ast::DerivimplicitCallback More... | |
| void | visit_eigen_newton_solver_block (const ast::EigenNewtonSolverBlock &node) override |
| visit node of type ast::EigenNewtonSolverBlock More... | |
| void | visit_eigen_linear_solver_block (const ast::EigenLinearSolverBlock &node) override |
| visit node of type ast::EigenLinearSolverBlock More... | |
| void | visit_for_netcon (const ast::ForNetcon &node) override |
| visit node of type ast::ForNetcon More... | |
| virtual void | visit_watch_statement (const ast::WatchStatement &node) override |
| visit node of type ast::WatchStatement More... | |
| template<typename T > | |
| void | print_function_declaration (const T &node, const std::string &name) |
| Print prototype declarations of functions or procedures. More... | |
| Protected Member Functions inherited from nmodl::codegen::CodegenCppVisitor | |
| std::string | nmodl_version () const noexcept |
| Return Nmodl language version. More... | |
| std::string | instance_struct () const |
| Name of structure that wraps range variables. More... | |
| std::string | node_data_struct () const |
| Name of structure that wraps node variables. More... | |
| std::string | global_struct () const |
| Name of structure that wraps global variables. More... | |
| std::string | global_struct_instance () const |
Name of the (host-only) global instance of global_struct More... | |
| const char * | local_var_type () const noexcept |
| Data type for the local variables. More... | |
| const char * | default_float_data_type () const noexcept |
| Default data type for floating point elements. More... | |
| const std::string & | float_data_type () const noexcept |
| Data type for floating point elements specified on command line. More... | |
| const char * | default_int_data_type () const noexcept |
| Default data type for integer (offset) elements. More... | |
| const char * | operator_for_rhs () const noexcept |
| Operator for rhs vector update (matrix update) More... | |
| const char * | operator_for_d () const noexcept |
| Operator for diagonal vector update (matrix update) More... | |
| std::string | get_channel_info_var_name () const noexcept |
| Name of channel info variable. More... | |
| bool | ion_variable_struct_required () const |
| Check if a structure for ion variables is required. More... | |
| template<typename T > | |
| bool | has_parameter_of_name (const T &node, const std::string &name) |
| Check if function or procedure node has parameter with given name. More... | |
| bool | net_send_buffer_required () const noexcept |
| Check if net_send_buffer is required. More... | |
| bool | net_receive_buffering_required () const noexcept |
| Check if net receive/send buffering kernels required. More... | |
| bool | nrn_state_required () const noexcept |
| Check if nrn_state function is required. More... | |
| bool | nrn_cur_required () const noexcept |
| Check if nrn_cur function is required. More... | |
| bool | net_receive_required () const noexcept |
| Check if net_receive function is required. More... | |
| bool | range_variable_setup_required () const noexcept |
| Check if setup_range_variable function is required. More... | |
| bool | net_receive_exist () const noexcept |
| Check if net_receive node exist. More... | |
| bool | breakpoint_exist () const noexcept |
| Check if breakpoint node exist. More... | |
| bool | defined_method (const std::string &name) const |
| Check if given method is defined in this model. More... | |
| bool | is_net_send (const std::string &name) const noexcept |
Checks if given function name is net_send. More... | |
| bool | is_net_move (const std::string &name) const noexcept |
Checks if given function name is net_move. More... | |
| bool | is_net_event (const std::string &name) const noexcept |
Checks if given function name is net_event. More... | |
| int | float_variables_size () const |
| Number of float variables in the model. More... | |
| int | int_variables_size () const |
| Number of integer variables in the model. More... | |
| std::string | format_double_string (const std::string &value) |
Convert a given double value to its string representation. More... | |
| std::string | format_float_string (const std::string &value) |
Convert a given float value to its string representation. More... | |
| void | update_index_semantics () |
| populate all index semantics needed for registration with coreneuron More... | |
| std::vector< SymbolType > | get_float_variables () const |
Determine all float variables required during code generation. More... | |
| std::vector< IndexVariableInfo > | get_int_variables () |
Determine all int variables required during code generation. More... | |
| std::vector< std::string > | ion_read_statements (BlockType type) const |
| For a given output block type, return statements for all read ion variables. More... | |
| std::vector< std::string > | ion_read_statements_optimized (BlockType type) const |
| For a given output block type, return minimal statements for all read ion variables. More... | |
| std::vector< ShadowUseStatement > | ion_write_statements (BlockType type) |
| For a given output block type, return statements for writing back ion variables. More... | |
| std::string | process_shadow_update_statement (const ShadowUseStatement &statement, BlockType type) |
| Process shadow update statement. More... | |
| std::string | breakpoint_current (std::string current) const |
| Determine the variable name for the "current" used in breakpoint block taking into account intermediate code transformations. More... | |
| virtual void | print_global_var_struct_decl () |
| Instantiate global var instance. More... | |
| void | print_statement_block (const ast::StatementBlock &node, bool open_brace=true, bool close_brace=true) |
| Print any statement block in nmodl with option to (not) print braces. More... | |
| virtual void | print_function_call (const ast::FunctionCall &node) |
| Print call to internal or external function. More... | |
| void | rename_function_arguments () |
| Rename function/procedure arguments that conflict with default arguments. More... | |
| template<typename T > | |
| void | print_vector_elements (const std::vector< T > &elements, const std::string &separator, const std::string &prefix="") |
| Print the items in a vector as a list. More... | |
| std::string | add_escape_quote (const std::string &text) const |
| Add quotes to string to be output. More... | |
| std::string | method_name (const std::string &name) const |
| Constructs the name of a function or procedure. More... | |
| SymbolType | make_symbol (const std::string &name) const |
| Creates a temporary symbol. More... | |
| std::string | update_if_ion_variable_name (const std::string &name) const |
| Determine the updated name if the ion variable has been optimized. More... | |
| virtual void | print_global_var_struct_assertions () const |
| Print static assertions about the global variable struct. More... | |
| void | print_prcellstate_macros () const |
| Print declaration of macro NRN_PRCELLSTATE for debugging. More... | |
| void | print_mechanism_info () |
| Print backend code for byte array that has mechanism information (to be registered with NEURON/CoreNEURON) More... | |
| void | print_nmodl_constants () |
| Print the nmodl constants used in backend code. More... | |
| void | visit_binary_expression (const ast::BinaryExpression &node) override |
| visit node of type ast::BinaryExpression More... | |
| void | visit_binary_operator (const ast::BinaryOperator &node) override |
| visit node of type ast::BinaryOperator More... | |
| void | visit_boolean (const ast::Boolean &node) override |
| visit node of type ast::Boolean More... | |
| void | visit_double (const ast::Double &node) override |
| visit node of type ast::Double More... | |
| void | visit_else_if_statement (const ast::ElseIfStatement &node) override |
| visit node of type ast::ElseIfStatement More... | |
| void | visit_else_statement (const ast::ElseStatement &node) override |
| visit node of type ast::ElseStatement More... | |
| void | visit_float (const ast::Float &node) override |
| visit node of type ast::Float More... | |
| void | visit_from_statement (const ast::FromStatement &node) override |
| visit node of type ast::FromStatement More... | |
| void | visit_function_call (const ast::FunctionCall &node) override |
| visit node of type ast::FunctionCall More... | |
| void | visit_if_statement (const ast::IfStatement &node) override |
| visit node of type ast::IfStatement More... | |
| void | visit_indexed_name (const ast::IndexedName &node) override |
| visit node of type ast::IndexedName More... | |
| void | visit_integer (const ast::Integer &node) override |
| visit node of type ast::Integer More... | |
| void | visit_local_list_statement (const ast::LocalListStatement &node) override |
| visit node of type ast::LocalListStatement More... | |
| void | visit_name (const ast::Name &node) override |
| visit node of type ast::Name More... | |
| void | visit_paren_expression (const ast::ParenExpression &node) override |
| visit node of type ast::ParenExpression More... | |
| void | visit_prime_name (const ast::PrimeName &node) override |
| visit node of type ast::PrimeName More... | |
| void | visit_statement_block (const ast::StatementBlock &node) override |
| void | visit_string (const ast::String &node) override |
| visit node of type ast::String More... | |
| void | visit_unary_operator (const ast::UnaryOperator &node) override |
| visit node of type ast::UnaryOperator More... | |
| void | visit_unit (const ast::Unit &node) override |
| visit node of type ast::Unit More... | |
| void | visit_var_name (const ast::VarName &node) override |
| void | visit_verbatim (const ast::Verbatim &node) override |
| visit node of type ast::Verbatim More... | |
| void | visit_while_statement (const ast::WhileStatement &node) override |
| visit node of type ast::WhileStatement More... | |
| void | visit_update_dt (const ast::UpdateDt &node) override |
| visit node of type ast::UpdateDt More... | |
| void | visit_protect_statement (const ast::ProtectStatement &node) override |
| visit node of type ast::ProtectStatement More... | |
| void | visit_mutex_lock (const ast::MutexLock &node) override |
| visit node of type ast::MutexLock More... | |
| void | visit_mutex_unlock (const ast::MutexUnlock &node) override |
| visit node of type ast::MutexUnlock More... | |
| void | visit_solution_expression (const ast::SolutionExpression &node) override |
| visit node of type ast::SolutionExpression More... | |
| std::string | compute_method_name (BlockType type) const |
Additional Inherited Members | |
| Public Member Functions inherited from nmodl::codegen::CodegenCoreneuronCppVisitor | |
| void | print_instance_variable_setup () |
| Print the function that initialize instance structure. More... | |
| void | print_mechanism_range_var_structure (bool print_initializers) override |
| Print the structure that wraps all range and int variables required for the NMODL. More... | |
| CodegenCppVisitor (std::string mod_filename, const std::string &output_dir, std::string float_type, const bool optimize_ionvar_copies, size_t blame_line=0) | |
| Constructs the C++ code generator visitor. More... | |
| CodegenCppVisitor (std::string mod_filename, std::ostream &stream, std::string float_type, const bool optimize_ionvar_copies, size_t blame_line=0) | |
| Constructs the C++ code generator visitor. More... | |
| Public Member Functions inherited from nmodl::codegen::CodegenCppVisitor | |
| CodegenCppVisitor (std::string mod_filename, const std::string &output_dir, std::string float_type, const bool optimize_ionvar_copies, size_t blame_line=0) | |
| Constructs the C++ code generator visitor. More... | |
| CodegenCppVisitor (std::string mod_filename, std::ostream &stream, std::string float_type, const bool optimize_ionvar_copies, size_t blame_line=0) | |
| Constructs the C++ code generator visitor. More... | |
| void | setup (const ast::Program &node) |
| Setup the target backend code generator. More... | |
| void | visit_program (const ast::Program &program) override |
| Main and only member function to call after creating an instance of this class. More... | |
| Public Member Functions inherited from nmodl::visitor::ConstAstVisitor | |
| void | visit_node (const ast::Node &node) override |
| visit node of type ast::Node More... | |
| void | visit_statement (const ast::Statement &node) override |
| visit node of type ast::Statement More... | |
| void | visit_expression (const ast::Expression &node) override |
| visit node of type ast::Expression More... | |
| void | visit_block (const ast::Block &node) override |
| visit node of type ast::Block More... | |
| void | visit_identifier (const ast::Identifier &node) override |
| visit node of type ast::Identifier More... | |
| void | visit_number (const ast::Number &node) override |
| visit node of type ast::Number More... | |
| void | visit_string (const ast::String &node) override |
| visit node of type ast::String More... | |
| void | visit_integer (const ast::Integer &node) override |
| visit node of type ast::Integer More... | |
| void | visit_float (const ast::Float &node) override |
| visit node of type ast::Float More... | |
| void | visit_double (const ast::Double &node) override |
| visit node of type ast::Double More... | |
| void | visit_boolean (const ast::Boolean &node) override |
| visit node of type ast::Boolean More... | |
| void | visit_name (const ast::Name &node) override |
| visit node of type ast::Name More... | |
| void | visit_prime_name (const ast::PrimeName &node) override |
| visit node of type ast::PrimeName More... | |
| void | visit_indexed_name (const ast::IndexedName &node) override |
| visit node of type ast::IndexedName More... | |
| void | visit_var_name (const ast::VarName &node) override |
| visit node of type ast::VarName More... | |
| void | visit_argument (const ast::Argument &node) override |
| visit node of type ast::Argument More... | |
| void | visit_react_var_name (const ast::ReactVarName &node) override |
| visit node of type ast::ReactVarName More... | |
| void | visit_read_ion_var (const ast::ReadIonVar &node) override |
| visit node of type ast::ReadIonVar More... | |
| void | visit_write_ion_var (const ast::WriteIonVar &node) override |
| visit node of type ast::WriteIonVar More... | |
| void | visit_nonspecific_cur_var (const ast::NonspecificCurVar &node) override |
| visit node of type ast::NonspecificCurVar More... | |
| void | visit_electrode_cur_var (const ast::ElectrodeCurVar &node) override |
| visit node of type ast::ElectrodeCurVar More... | |
| void | visit_range_var (const ast::RangeVar &node) override |
| visit node of type ast::RangeVar More... | |
| void | visit_global_var (const ast::GlobalVar &node) override |
| visit node of type ast::GlobalVar More... | |
| void | visit_pointer_var (const ast::PointerVar &node) override |
| visit node of type ast::PointerVar More... | |
| void | visit_random_var (const ast::RandomVar &node) override |
| visit node of type ast::RandomVar More... | |
| void | visit_bbcore_pointer_var (const ast::BbcorePointerVar &node) override |
| visit node of type ast::BbcorePointerVar More... | |
| void | visit_extern_var (const ast::ExternVar &node) override |
| visit node of type ast::ExternVar More... | |
| void | visit_param_block (const ast::ParamBlock &node) override |
| visit node of type ast::ParamBlock More... | |
| void | visit_independent_block (const ast::IndependentBlock &node) override |
| visit node of type ast::IndependentBlock More... | |
| void | visit_assigned_block (const ast::AssignedBlock &node) override |
| visit node of type ast::AssignedBlock More... | |
| void | visit_state_block (const ast::StateBlock &node) override |
| visit node of type ast::StateBlock More... | |
| void | visit_initial_block (const ast::InitialBlock &node) override |
| visit node of type ast::InitialBlock More... | |
| void | visit_constructor_block (const ast::ConstructorBlock &node) override |
| visit node of type ast::ConstructorBlock More... | |
| void | visit_destructor_block (const ast::DestructorBlock &node) override |
| visit node of type ast::DestructorBlock More... | |
| void | visit_statement_block (const ast::StatementBlock &node) override |
| visit node of type ast::StatementBlock More... | |
| void | visit_derivative_block (const ast::DerivativeBlock &node) override |
| visit node of type ast::DerivativeBlock More... | |
| void | visit_linear_block (const ast::LinearBlock &node) override |
| visit node of type ast::LinearBlock More... | |
| void | visit_non_linear_block (const ast::NonLinearBlock &node) override |
| visit node of type ast::NonLinearBlock More... | |
| void | visit_discrete_block (const ast::DiscreteBlock &node) override |
| visit node of type ast::DiscreteBlock More... | |
| void | visit_function_table_block (const ast::FunctionTableBlock &node) override |
| visit node of type ast::FunctionTableBlock More... | |
| void | visit_function_block (const ast::FunctionBlock &node) override |
| visit node of type ast::FunctionBlock More... | |
| void | visit_procedure_block (const ast::ProcedureBlock &node) override |
| visit node of type ast::ProcedureBlock More... | |
| void | visit_net_receive_block (const ast::NetReceiveBlock &node) override |
| visit node of type ast::NetReceiveBlock More... | |
| void | visit_solve_block (const ast::SolveBlock &node) override |
| visit node of type ast::SolveBlock More... | |
| void | visit_breakpoint_block (const ast::BreakpointBlock &node) override |
| visit node of type ast::BreakpointBlock More... | |
| void | visit_before_block (const ast::BeforeBlock &node) override |
| visit node of type ast::BeforeBlock More... | |
| void | visit_after_block (const ast::AfterBlock &node) override |
| visit node of type ast::AfterBlock More... | |
| void | visit_ba_block (const ast::BABlock &node) override |
| visit node of type ast::BABlock More... | |
| void | visit_for_netcon (const ast::ForNetcon &node) override |
| visit node of type ast::ForNetcon More... | |
| void | visit_kinetic_block (const ast::KineticBlock &node) override |
| visit node of type ast::KineticBlock More... | |
| void | visit_unit_block (const ast::UnitBlock &node) override |
| visit node of type ast::UnitBlock More... | |
| void | visit_constant_block (const ast::ConstantBlock &node) override |
| visit node of type ast::ConstantBlock More... | |
| void | visit_neuron_block (const ast::NeuronBlock &node) override |
| visit node of type ast::NeuronBlock More... | |
| void | visit_unit (const ast::Unit &node) override |
| visit node of type ast::Unit More... | |
| void | visit_double_unit (const ast::DoubleUnit &node) override |
| visit node of type ast::DoubleUnit More... | |
| void | visit_local_var (const ast::LocalVar &node) override |
| visit node of type ast::LocalVar More... | |
| void | visit_limits (const ast::Limits &node) override |
| visit node of type ast::Limits More... | |
| void | visit_number_range (const ast::NumberRange &node) override |
| visit node of type ast::NumberRange More... | |
| void | visit_constant_var (const ast::ConstantVar &node) override |
| visit node of type ast::ConstantVar More... | |
| void | visit_binary_operator (const ast::BinaryOperator &node) override |
| visit node of type ast::BinaryOperator More... | |
| void | visit_unary_operator (const ast::UnaryOperator &node) override |
| visit node of type ast::UnaryOperator More... | |
| void | visit_reaction_operator (const ast::ReactionOperator &node) override |
| visit node of type ast::ReactionOperator More... | |
| void | visit_paren_expression (const ast::ParenExpression &node) override |
| visit node of type ast::ParenExpression More... | |
| void | visit_binary_expression (const ast::BinaryExpression &node) override |
| visit node of type ast::BinaryExpression More... | |
| void | visit_diff_eq_expression (const ast::DiffEqExpression &node) override |
| visit node of type ast::DiffEqExpression More... | |
| void | visit_unary_expression (const ast::UnaryExpression &node) override |
| visit node of type ast::UnaryExpression More... | |
| void | visit_non_lin_equation (const ast::NonLinEquation &node) override |
| visit node of type ast::NonLinEquation More... | |
| void | visit_lin_equation (const ast::LinEquation &node) override |
| visit node of type ast::LinEquation More... | |
| void | visit_function_call (const ast::FunctionCall &node) override |
| visit node of type ast::FunctionCall More... | |
| void | visit_watch (const ast::Watch &node) override |
| visit node of type ast::Watch More... | |
| void | visit_ba_block_type (const ast::BABlockType &node) override |
| visit node of type ast::BABlockType More... | |
| void | visit_unit_def (const ast::UnitDef &node) override |
| visit node of type ast::UnitDef More... | |
| void | visit_factor_def (const ast::FactorDef &node) override |
| visit node of type ast::FactorDef More... | |
| void | visit_valence (const ast::Valence &node) override |
| visit node of type ast::Valence More... | |
| void | visit_unit_state (const ast::UnitState &node) override |
| visit node of type ast::UnitState More... | |
| void | visit_local_list_statement (const ast::LocalListStatement &node) override |
| visit node of type ast::LocalListStatement More... | |
| void | visit_model (const ast::Model &node) override |
| visit node of type ast::Model More... | |
| void | visit_define (const ast::Define &node) override |
| visit node of type ast::Define More... | |
| void | visit_include (const ast::Include &node) override |
| visit node of type ast::Include More... | |
| void | visit_param_assign (const ast::ParamAssign &node) override |
| visit node of type ast::ParamAssign More... | |
| void | visit_assigned_definition (const ast::AssignedDefinition &node) override |
| visit node of type ast::AssignedDefinition More... | |
| void | visit_conductance_hint (const ast::ConductanceHint &node) override |
| visit node of type ast::ConductanceHint More... | |
| void | visit_expression_statement (const ast::ExpressionStatement &node) override |
| visit node of type ast::ExpressionStatement More... | |
| void | visit_protect_statement (const ast::ProtectStatement &node) override |
| visit node of type ast::ProtectStatement More... | |
| void | visit_from_statement (const ast::FromStatement &node) override |
| visit node of type ast::FromStatement More... | |
| void | visit_while_statement (const ast::WhileStatement &node) override |
| visit node of type ast::WhileStatement More... | |
| void | visit_if_statement (const ast::IfStatement &node) override |
| visit node of type ast::IfStatement More... | |
| void | visit_else_if_statement (const ast::ElseIfStatement &node) override |
| visit node of type ast::ElseIfStatement More... | |
| void | visit_else_statement (const ast::ElseStatement &node) override |
| visit node of type ast::ElseStatement More... | |
| void | visit_watch_statement (const ast::WatchStatement &node) override |
| visit node of type ast::WatchStatement More... | |
| void | visit_mutex_lock (const ast::MutexLock &node) override |
| visit node of type ast::MutexLock More... | |
| void | visit_mutex_unlock (const ast::MutexUnlock &node) override |
| visit node of type ast::MutexUnlock More... | |
| void | visit_conserve (const ast::Conserve &node) override |
| visit node of type ast::Conserve More... | |
| void | visit_compartment (const ast::Compartment &node) override |
| visit node of type ast::Compartment More... | |
| void | visit_lon_difuse (const ast::LonDifuse &node) override |
| visit node of type ast::LonDifuse More... | |
| void | visit_reaction_statement (const ast::ReactionStatement &node) override |
| visit node of type ast::ReactionStatement More... | |
| void | visit_lag_statement (const ast::LagStatement &node) override |
| visit node of type ast::LagStatement More... | |
| void | visit_constant_statement (const ast::ConstantStatement &node) override |
| visit node of type ast::ConstantStatement More... | |
| void | visit_table_statement (const ast::TableStatement &node) override |
| visit node of type ast::TableStatement More... | |
| void | visit_suffix (const ast::Suffix &node) override |
| visit node of type ast::Suffix More... | |
| void | visit_useion (const ast::Useion &node) override |
| visit node of type ast::Useion More... | |
| void | visit_nonspecific (const ast::Nonspecific &node) override |
| visit node of type ast::Nonspecific More... | |
| void | visit_electrode_current (const ast::ElectrodeCurrent &node) override |
| visit node of type ast::ElectrodeCurrent More... | |
| void | visit_range (const ast::Range &node) override |
| visit node of type ast::Range More... | |
| void | visit_global (const ast::Global &node) override |
| visit node of type ast::Global More... | |
| void | visit_random_var_list (const ast::RandomVarList &node) override |
| visit node of type ast::RandomVarList More... | |
| void | visit_pointer (const ast::Pointer &node) override |
| visit node of type ast::Pointer More... | |
| void | visit_bbcore_pointer (const ast::BbcorePointer &node) override |
| visit node of type ast::BbcorePointer More... | |
| void | visit_external (const ast::External &node) override |
| visit node of type ast::External More... | |
| void | visit_thread_safe (const ast::ThreadSafe &node) override |
| visit node of type ast::ThreadSafe More... | |
| void | visit_verbatim (const ast::Verbatim &node) override |
| visit node of type ast::Verbatim More... | |
| void | visit_line_comment (const ast::LineComment &node) override |
| visit node of type ast::LineComment More... | |
| void | visit_block_comment (const ast::BlockComment &node) override |
| visit node of type ast::BlockComment More... | |
| void | visit_ontology_statement (const ast::OntologyStatement &node) override |
| visit node of type ast::OntologyStatement More... | |
| void | visit_program (const ast::Program &node) override |
| visit node of type ast::Program More... | |
| void | visit_nrn_state_block (const ast::NrnStateBlock &node) override |
| visit node of type ast::NrnStateBlock More... | |
| void | visit_eigen_newton_solver_block (const ast::EigenNewtonSolverBlock &node) override |
| visit node of type ast::EigenNewtonSolverBlock More... | |
| void | visit_eigen_linear_solver_block (const ast::EigenLinearSolverBlock &node) override |
| visit node of type ast::EigenLinearSolverBlock More... | |
| void | visit_wrapped_expression (const ast::WrappedExpression &node) override |
| visit node of type ast::WrappedExpression More... | |
| void | visit_derivimplicit_callback (const ast::DerivimplicitCallback &node) override |
| visit node of type ast::DerivimplicitCallback More... | |
| void | visit_solution_expression (const ast::SolutionExpression &node) override |
| visit node of type ast::SolutionExpression More... | |
| void | visit_update_dt (const ast::UpdateDt &node) override |
| visit node of type ast::UpdateDt More... | |
| Public Member Functions inherited from nmodl::visitor::ConstVisitor | |
| virtual | ~ConstVisitor ()=default |
| Protected Types inherited from nmodl::codegen::CodegenCppVisitor | |
| using | SymbolType = std::shared_ptr< symtab::Symbol > |
| using | ParamVector = std::vector< std::tuple< std::string, std::string, std::string, std::string > > |
| A vector of parameters represented by a 4-tuple of strings: More... | |
| Static Protected Member Functions inherited from nmodl::codegen::CodegenCppVisitor | |
| static bool | need_semicolon (const ast::Statement &node) |
| Check if a semicolon is required at the end of given statement. More... | |
| static std::string | get_parameter_str (const ParamVector ¶ms) |
| Generate the string representing the procedure parameter declaration. More... | |
| static bool | statement_to_skip (const ast::Statement &node) |
| Check if given statement should be skipped during code generation. More... | |
| static std::pair< std::string, std::string > | read_ion_variable_name (const std::string &name) |
| Return ion variable name and corresponding ion read variable name. More... | |
| static std::pair< std::string, std::string > | write_ion_variable_name (const std::string &name) |
| Return ion variable name and corresponding ion write variable name. More... | |
| Protected Attributes inherited from nmodl::codegen::CodegenCppVisitor | |
| std::unique_ptr< CodePrinter > | printer |
| Code printer object for target (C++) More... | |
| std::string | mod_filename |
| Name of mod file (without .mod suffix) More... | |
| std::string | float_type = codegen::naming::DEFAULT_FLOAT_TYPE |
| Data type of floating point variables. More... | |
| bool | optimize_ionvar_copies = true |
| Flag to indicate if visitor should avoid ion variable copies. More... | |
| codegen::CodegenInfo | info |
| All ast information for code generation. More... | |
| symtab::SymbolTable * | program_symtab = nullptr |
| Symbol table for the program. More... | |
| std::vector< SymbolType > | codegen_float_variables |
| All float variables for the model. More... | |
| std::vector< IndexVariableInfo > | codegen_int_variables |
| All int variables for the model. More... | |
| std::vector< SymbolType > | codegen_global_variables |
| All global variables for the model. More... | |
| bool | enable_variable_name_lookup = true |
| Variable name should be converted to instance name (but not for function arguments) More... | |
| bool | printing_net_receive = false |
true if currently net_receive block being printed More... | |
| bool | printing_net_init = false |
true if currently initial block of net_receive being printed More... | |
| bool | printing_top_verbatim_blocks = false |
true if currently printing top level verbatim blocks More... | |
| bool | internal_method_call_encountered = false |
true if internal method call was encountered while processing verbatim block More... | |
| int | current_watch_statement = 0 |
| Index of watch statement being printed. More... | |
|
overrideprotectedvirtual |
name of the code generation backend
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 70 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
if reduction block in nrn_cur required
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 226 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
abort routine
OpenACC kernels running on GPU doesn't support abort().
CUDA/OpenACC supports assert() in device kernel that can be used for similar purpose. Also, printf is supported on device.
assert(0==1) pattern which is used for OpenACC. Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 108 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
atomic update pragma for reduction statements
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 58 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
common includes : standard c++, coreneuron and backend specific
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 64 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
ivdep like annotation for channel iterations
Depending programming model and compiler, we print compiler hint for parallelization.
For example:
#pragma ivdep
for(int id=0; id<nodecount; id++) {
#pragma acc parallel loop
for(int id=0; id<nodecount; id++) {
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 34 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
update derivimplicit advance flag on the gpu device
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 323 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
print atomic capture pragma
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 329 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
Print the code to synchronise/wait on stream specific to NrnThread.
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 335 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
update dt from host to device
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 365 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
fast membrane current calculation
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 187 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
update global variable from host to the device
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 231 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
call helper function for copying the instance struct to the device
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 307 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
call helper function that deletes the instance struct from the device
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 315 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
declare helper functions for copying the instance struct to the device
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 257 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
define helper functions for copying the instance struct to the device
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 269 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
annotations like "acc enter data present(...)" for main kernel
Each kernel like nrn_init, nrn_state and nrn_cur could be offloaded to accelerator.
In this case, at very top level, we print pragma for data present. For example:
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 143 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
end of annotation like "acc enter data"
End of print_kernel_enter_data_begin.
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 214 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
memory allocation routine
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 75 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
start of annotation "acc kernels" for net_init kernel
INITIAL block from NET_RECEIVE generates net_init function.
The net_init function pointer is registered with the coreneuron and called from the CPU. As the data is on GPU, we need to launch net_init on the GPU.
serial construct to launch serial kernels. This is during initialization but still inefficient. This should be improved when we drop MOD2C. Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 160 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
end of annotation "acc kernels" for net_init kernel
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 168 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
update NetSendBuffer_t count from host to device
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 357 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
update NetSendBuffer_t count from device to host
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 342 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
update NetSendBuffer_t from device to host
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 348 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
print atomic update of NetSendBuffer_t cnt
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 116 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
Replace default implementation by a no-op since the buffer cannot be grown up during gpu execution.
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 125 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
transfer newtonspace structure to device
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 241 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
reduction to matrix elements from shadow vectors
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 206 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
update to matrix elements with/without shadow vectors
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 174 of file codegen_acc_visitor.cpp.
|
overrideprotectedvirtual |
setup method for setting matrix shadow vectors
Reimplemented from nmodl::codegen::CodegenCoreneuronCppVisitor.
Definition at line 221 of file codegen_acc_visitor.cpp.