diff --git a/ionic/coreneuron/ionic.cpp b/ionic/coreneuron/ionic.cpp
new file mode 100644
index 0000000..8d25d1b
--- /dev/null
+++ b/ionic/coreneuron/ionic.cpp
@@ -0,0 +1,289 @@
+/*********************************************************
+Model Name      : ionic
+Filename        : ionic.mod
+NMODL Version   : 7.7.0
+Vectorized      : true
+Threadsafe      : true
+Created         : DATE
+Simulator       : CoreNEURON
+Backend         : C++ (api-compatibility)
+NMODL Compiler  : VERSION
+*********************************************************/
+
+#include <math.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+#include <coreneuron/gpu/nrn_acc_manager.hpp>
+#include <coreneuron/mechanism/mech/mod2c_core_thread.hpp>
+#include <coreneuron/mechanism/register_mech.hpp>
+#include <coreneuron/nrnconf.h>
+#include <coreneuron/nrniv/nrniv_decl.h>
+#include <coreneuron/sim/multicore.hpp>
+#include <coreneuron/sim/scopmath/newton_thread.hpp>
+#include <coreneuron/utils/ivocvect.hpp>
+#include <coreneuron/utils/nrnoc_aux.hpp>
+#include <coreneuron/utils/randoms/nrnran123.h>
+
+
+namespace coreneuron {
+    #ifndef NRN_PRCELLSTATE
+    #define NRN_PRCELLSTATE 0
+    #endif
+
+
+    /** channel information */
+    static const char *mechanism_info[] = {
+        "7.7.0",
+        "ionic",
+        0,
+        0,
+        0,
+        0
+    };
+
+
+    /** all global variables */
+    struct ionic_Store {
+        int na_type{};
+        int reset{};
+        int mech_type{};
+    };
+    static_assert(std::is_trivially_copy_constructible_v<ionic_Store>);
+    static_assert(std::is_trivially_move_constructible_v<ionic_Store>);
+    static_assert(std::is_trivially_copy_assignable_v<ionic_Store>);
+    static_assert(std::is_trivially_move_assignable_v<ionic_Store>);
+    static_assert(std::is_trivially_destructible_v<ionic_Store>);
+    ionic_Store ionic_global;
+
+
+    /** all mechanism instance variables and global variables */
+    struct ionic_Instance  {
+        double* ina{};
+        double* ena{};
+        double* v_unused{};
+        double* g_unused{};
+        const double* ion_ina{};
+        double* ion_ena{};
+        ionic_Store* global{&ionic_global};
+    };
+
+
+    /** connect global (scalar) variables to hoc -- */
+    static DoubScal hoc_scalar_double[] = {
+        {nullptr, nullptr}
+    };
+
+
+    /** connect global (array) variables to hoc -- */
+    static DoubVec hoc_vector_double[] = {
+        {nullptr, nullptr, 0}
+    };
+
+
+    static inline int first_pointer_var_index() {
+        return -1;
+    }
+
+
+    static inline int first_random_var_index() {
+        return -1;
+    }
+
+
+    static inline int float_variables_size() {
+        return 4;
+    }
+
+
+    static inline int int_variables_size() {
+        return 2;
+    }
+
+
+    static inline int get_mech_type() {
+        return ionic_global.mech_type;
+    }
+
+
+    static inline Memb_list* get_memb_list(NrnThread* nt) {
+        if (!nt->_ml_list) {
+            return nullptr;
+        }
+        return nt->_ml_list[get_mech_type()];
+    }
+
+
+    static inline void* mem_alloc(size_t num, size_t size, size_t alignment = 16) {
+        void* ptr;
+        posix_memalign(&ptr, alignment, num*size);
+        memset(ptr, 0, size);
+        return ptr;
+    }
+
+
+    static inline void mem_free(void* ptr) {
+        free(ptr);
+    }
+
+
+    static inline void coreneuron_abort() {
+        abort();
+    }
+
+    // Allocate instance structure
+    static void nrn_private_constructor_ionic(NrnThread* nt, Memb_list* ml, int type) {
+        assert(!ml->instance);
+        assert(!ml->global_variables);
+        assert(ml->global_variables_size == 0);
+        auto* const inst = new ionic_Instance{};
+        assert(inst->global == &ionic_global);
+        ml->instance = inst;
+        ml->global_variables = inst->global;
+        ml->global_variables_size = sizeof(ionic_Store);
+    }
+
+    // Deallocate the instance structure
+    static void nrn_private_destructor_ionic(NrnThread* nt, Memb_list* ml, int type) {
+        auto* const inst = static_cast<ionic_Instance*>(ml->instance);
+        assert(inst);
+        assert(inst->global);
+        assert(inst->global == &ionic_global);
+        assert(inst->global == ml->global_variables);
+        assert(ml->global_variables_size == sizeof(ionic_Store));
+        delete inst;
+        ml->instance = nullptr;
+        ml->global_variables = nullptr;
+        ml->global_variables_size = 0;
+    }
+
+    /** initialize mechanism instance variables */
+    static inline void setup_instance(NrnThread* nt, Memb_list* ml) {
+        auto* const inst = static_cast<ionic_Instance*>(ml->instance);
+        assert(inst);
+        assert(inst->global);
+        assert(inst->global == &ionic_global);
+        assert(inst->global == ml->global_variables);
+        assert(ml->global_variables_size == sizeof(ionic_Store));
+        int pnodecount = ml->_nodecount_padded;
+        Datum* indexes = ml->pdata;
+        inst->ina = ml->data+0*pnodecount;
+        inst->ena = ml->data+1*pnodecount;
+        inst->v_unused = ml->data+2*pnodecount;
+        inst->g_unused = ml->data+3*pnodecount;
+        inst->ion_ina = nt->_data;
+        inst->ion_ena = nt->_data;
+    }
+
+
+
+    static void nrn_alloc_ionic(double* data, Datum* indexes, int type) {
+        // do nothing
+    }
+
+
+    void nrn_constructor_ionic(NrnThread* nt, Memb_list* ml, int type) {
+        #ifndef CORENEURON_BUILD
+        int nodecount = ml->nodecount;
+        int pnodecount = ml->_nodecount_padded;
+        const int* node_index = ml->nodeindices;
+        double* data = ml->data;
+        const double* voltage = nt->_actual_v;
+        Datum* indexes = ml->pdata;
+        ThreadDatum* thread = ml->_thread;
+        auto* const inst = static_cast<ionic_Instance*>(ml->instance);
+
+        #endif
+    }
+
+
+    void nrn_destructor_ionic(NrnThread* nt, Memb_list* ml, int type) {
+        #ifndef CORENEURON_BUILD
+        int nodecount = ml->nodecount;
+        int pnodecount = ml->_nodecount_padded;
+        const int* node_index = ml->nodeindices;
+        double* data = ml->data;
+        const double* voltage = nt->_actual_v;
+        Datum* indexes = ml->pdata;
+        ThreadDatum* thread = ml->_thread;
+        auto* const inst = static_cast<ionic_Instance*>(ml->instance);
+
+        #endif
+    }
+
+
+    /** initialize channel */
+    void nrn_init_ionic(NrnThread* nt, Memb_list* ml, int type) {
+        int nodecount = ml->nodecount;
+        int pnodecount = ml->_nodecount_padded;
+        const int* node_index = ml->nodeindices;
+        double* data = ml->data;
+        const double* voltage = nt->_actual_v;
+        Datum* indexes = ml->pdata;
+        ThreadDatum* thread = ml->_thread;
+
+        setup_instance(nt, ml);
+        auto* const inst = static_cast<ionic_Instance*>(ml->instance);
+
+        if (_nrn_skip_initmodel == 0) {
+            #pragma omp simd
+            #pragma ivdep
+            for (int id = 0; id < nodecount; id++) {
+                int node_id = node_index[id];
+                double v = voltage[node_id];
+                #if NRN_PRCELLSTATE
+                inst->v_unused[id] = v;
+                #endif
+                inst->ina[id] = inst->ion_ina[indexes[0*pnodecount + id]];
+                inst->ion_ena[indexes[1*pnodecount + id]] = inst->ena[id];
+            }
+        }
+    }
+
+
+    /** update state */
+    void nrn_state_ionic(NrnThread* nt, Memb_list* ml, int type) {
+        int nodecount = ml->nodecount;
+        int pnodecount = ml->_nodecount_padded;
+        const int* node_index = ml->nodeindices;
+        double* data = ml->data;
+        const double* voltage = nt->_actual_v;
+        Datum* indexes = ml->pdata;
+        ThreadDatum* thread = ml->_thread;
+        auto* const inst = static_cast<ionic_Instance*>(ml->instance);
+
+        #pragma omp simd
+        #pragma ivdep
+        for (int id = 0; id < nodecount; id++) {
+            int node_id = node_index[id];
+            double v = voltage[node_id];
+            #if NRN_PRCELLSTATE
+            inst->v_unused[id] = v;
+            #endif
+            inst->ina[id] = inst->ion_ina[indexes[0*pnodecount + id]];
+            inst->ena[id] = 42.0;
+            inst->ion_ena[indexes[1*pnodecount + id]] = inst->ena[id];
+        }
+    }
+
+
+    /** register channel with the simulator */
+    void _ionic_reg() {
+
+        int mech_type = nrn_get_mechtype("ionic");
+        ionic_global.mech_type = mech_type;
+        if (mech_type == -1) {
+            return;
+        }
+
+        _nrn_layout_reg(mech_type, 0);
+        register_mech(mechanism_info, nrn_alloc_ionic, nullptr, nullptr, nrn_state_ionic, nrn_init_ionic, nrn_private_constructor_ionic, nrn_private_destructor_ionic, first_pointer_var_index(), 1);
+        ionic_global.na_type = nrn_get_mechtype("na_ion");
+
+        hoc_register_prop_size(mech_type, float_variables_size(), int_variables_size());
+        hoc_register_dparam_semantics(mech_type, 0, "na_ion");
+        hoc_register_dparam_semantics(mech_type, 1, "na_ion");
+        hoc_register_var(hoc_scalar_double, hoc_vector_double, NULL);
+    }
+}
diff --git a/ionic/neuron/ionic.cpp b/ionic/neuron/ionic.cpp
new file mode 100644
index 0000000..bd16be9
--- /dev/null
+++ b/ionic/neuron/ionic.cpp
@@ -0,0 +1,267 @@
+/*********************************************************
+Model Name      : ionic
+Filename        : ionic.mod
+NMODL Version   : 7.7.0
+Vectorized      : true
+Threadsafe      : true
+Created         : DATE
+Simulator       : NEURON
+Backend         : C++ (api-compatibility)
+NMODL Compiler  : VERSION
+*********************************************************/
+
+#include <math.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#include "mech_api.h"
+#include "neuron/cache/mechanism_range.hpp"
+#include "nrniv_mf.h"
+#include "section_fwd.hpp"
+
+/* NEURON global macro definitions */
+/* VECTORIZED */
+#define NRN_VECTORIZED 1
+
+static constexpr auto number_of_datum_variables = 2;
+static constexpr auto number_of_floating_point_variables = 4;
+
+namespace {
+template <typename T>
+using _nrn_mechanism_std_vector = std::vector<T>;
+using _nrn_model_sorted_token = neuron::model_sorted_token;
+using _nrn_mechanism_cache_range = neuron::cache::MechanismRange<number_of_floating_point_variables, number_of_datum_variables>;
+using _nrn_mechanism_cache_instance = neuron::cache::MechanismInstance<number_of_floating_point_variables, number_of_datum_variables>;
+using _nrn_non_owning_id_without_container = neuron::container::non_owning_identifier_without_container;
+template <typename T>
+using _nrn_mechanism_field = neuron::mechanism::field<T>;
+template <typename... Args>
+void _nrn_mechanism_register_data_fields(Args&&... args) {
+    neuron::mechanism::register_data_fields(std::forward<Args>(args)...);
+}
+}  // namespace
+
+Prop* hoc_getdata_range(int type);
+
+
+namespace neuron {
+    #ifndef NRN_PRCELLSTATE
+    #define NRN_PRCELLSTATE 0
+    #endif
+
+
+    /** channel information */
+    static const char *mechanism_info[] = {
+        "7.7.0",
+        "ionic",
+        0,
+        0,
+        0,
+        0
+    };
+
+
+    /* NEURON global variables */
+    static Symbol* _na_sym;
+    static int mech_type;
+    static Prop* _extcall_prop;
+    /* _prop_id kind of shadows _extcall_prop to allow validity checking. */
+    static _nrn_non_owning_id_without_container _prop_id{};
+    static int hoc_nrnpointerindex = -1;
+    static _nrn_mechanism_std_vector<Datum> _extcall_thread;
+
+
+    /** all global variables */
+    struct ionic_Store {
+    };
+    static_assert(std::is_trivially_copy_constructible_v<ionic_Store>);
+    static_assert(std::is_trivially_move_constructible_v<ionic_Store>);
+    static_assert(std::is_trivially_copy_assignable_v<ionic_Store>);
+    static_assert(std::is_trivially_move_assignable_v<ionic_Store>);
+    static_assert(std::is_trivially_destructible_v<ionic_Store>);
+    ionic_Store ionic_global;
+
+
+    /** all mechanism instance variables and global variables */
+    struct ionic_Instance  {
+        double* ina{};
+        double* ena{};
+        double* v_unused{};
+        double* g_unused{};
+        const double* const* ion_ina{};
+        double* const* ion_ena{};
+        ionic_Store* global{&ionic_global};
+    };
+
+
+    struct ionic_NodeData  {
+        int const * nodeindices;
+        double const * node_voltages;
+        double * node_diagonal;
+        double * node_rhs;
+        int nodecount;
+    };
+
+
+    static ionic_Instance make_instance_ionic(_nrn_mechanism_cache_range& _ml) {
+        return ionic_Instance {
+            _ml.template fpfield_ptr<0>(),
+            _ml.template fpfield_ptr<1>(),
+            _ml.template fpfield_ptr<2>(),
+            _ml.template fpfield_ptr<3>(),
+            _ml.template dptr_field_ptr<0>(),
+            _ml.template dptr_field_ptr<1>()
+        };
+    }
+
+
+    static ionic_NodeData make_node_data_ionic(NrnThread& _nt, Memb_list& _ml_arg) {
+        return ionic_NodeData {
+            _ml_arg.nodeindices,
+            _nt.node_voltage_storage(),
+            _nt.node_d_storage(),
+            _nt.node_rhs_storage(),
+            _ml_arg.nodecount
+        };
+    }
+
+
+    static void nrn_alloc_ionic(Prop* _prop) {
+        Prop *prop_ion{};
+        Datum *_ppvar{};
+        _ppvar = nrn_prop_datum_alloc(mech_type, 2, _prop);
+        _nrn_mechanism_access_dparam(_prop) = _ppvar;
+        _nrn_mechanism_cache_instance _ml_real{_prop};
+        auto* const _ml = &_ml_real;
+        size_t const _iml{};
+        assert(_nrn_mechanism_get_num_vars(_prop) == 4);
+        /*initialize range parameters*/
+        _nrn_mechanism_access_dparam(_prop) = _ppvar;
+        Symbol * na_sym = hoc_lookup("na_ion");
+        Prop * na_prop = need_memb(na_sym);
+        _ppvar[0] = _nrn_mechanism_get_param_handle(na_prop, 3);
+        _ppvar[1] = _nrn_mechanism_get_param_handle(na_prop, 0);
+    }
+
+
+    /* Neuron setdata functions */
+    extern void _nrn_setdata_reg(int, void(*)(Prop*));
+    static void _setdata(Prop* _prop) {
+        _extcall_prop = _prop;
+        _prop_id = _nrn_get_prop_id(_prop);
+    }
+    static void _hoc_setdata() {
+        Prop *_prop = hoc_getdata_range(mech_type);
+        _setdata(_prop);
+        hoc_retpushx(1.);
+    }
+    /* Mechanism procedures and functions */
+
+
+    /** connect global (scalar) variables to hoc -- */
+    static DoubScal hoc_scalar_double[] = {
+        {nullptr, nullptr}
+    };
+
+
+    /** connect global (array) variables to hoc -- */
+    static DoubVec hoc_vector_double[] = {
+        {nullptr, nullptr, 0}
+    };
+
+
+    /* declaration of user functions */
+
+
+    /* connect user functions to hoc names */
+    static VoidFunc hoc_intfunc[] = {
+        {"setdata_ionic", _hoc_setdata},
+        {0, 0}
+    };
+    static NPyDirectMechFunc npy_direct_func_proc[] = {
+    };
+
+
+    void nrn_init_ionic(_nrn_model_sorted_token const& _sorted_token, NrnThread* _nt, Memb_list* _ml_arg, int _type) {
+        _nrn_mechanism_cache_range _lmr{_sorted_token, *_nt, *_ml_arg, _type};
+        auto inst = make_instance_ionic(_lmr);
+        auto node_data = make_node_data_ionic(*_nt, *_ml_arg);
+        auto nodecount = _ml_arg->nodecount;
+        auto* const _ml = &_lmr;
+        auto* _thread = _ml_arg->_thread;
+        for (int id = 0; id < nodecount; id++) {
+            
+            int node_id = node_data.nodeindices[id];
+            auto* _ppvar = _ml_arg->pdata[id];
+            auto v = node_data.node_voltages[node_id];
+            inst.v_unused[id] = v;
+            inst.ina[id] = (*inst.ion_ina[id]);
+            (*inst.ion_ena[id]) = inst.ena[id];
+        }
+    }
+
+
+    void nrn_state_ionic(_nrn_model_sorted_token const& _sorted_token, NrnThread* _nt, Memb_list* _ml_arg, int _type) {
+        _nrn_mechanism_cache_range _lmr{_sorted_token, *_nt, *_ml_arg, _type};
+        auto inst = make_instance_ionic(_lmr);
+        auto node_data = make_node_data_ionic(*_nt, *_ml_arg);
+        auto nodecount = _ml_arg->nodecount;
+        auto* const _ml = &_lmr;
+        auto* _thread = _ml_arg->_thread;
+        for (int id = 0; id < nodecount; id++) {
+            
+            int node_id = node_data.nodeindices[id];
+            auto* _ppvar = _ml_arg->pdata[id];
+            auto v = node_data.node_voltages[node_id];
+            inst.ina[id] = (*inst.ion_ina[id]);
+            inst.ena[id] = 42.0;
+            (*inst.ion_ena[id]) = inst.ena[id];
+        }
+    }
+
+
+    /** nrn_jacob function */
+    static void nrn_jacob_ionic(_nrn_model_sorted_token const& _sorted_token, NrnThread* _nt, Memb_list* _ml_arg, int _type) {
+        _nrn_mechanism_cache_range _lmr{_sorted_token, *_nt, *_ml_arg, _type};
+        auto inst = make_instance_ionic(_lmr);
+        auto node_data = make_node_data_ionic(*_nt, *_ml_arg);
+        auto nodecount = _ml_arg->nodecount;
+        for (int id = 0; id < nodecount; id++) {
+            // set conductances properly
+            int node_id = node_data.nodeindices[id];
+            node_data.node_diagonal[node_id] += inst.g_unused[id];
+        }
+    }
+
+
+    static void _initlists() {
+    }
+
+
+    /** register channel with the simulator */
+    extern "C" void _ionic_reg() {
+        _initlists();
+
+        ion_reg("na", -10000.);
+
+        _na_sym = hoc_lookup("na_ion");
+
+        register_mech(mechanism_info, nrn_alloc_ionic, nullptr, nrn_jacob_ionic, nrn_state_ionic, nrn_init_ionic, hoc_nrnpointerindex, 1);
+
+        mech_type = nrn_get_mechtype(mechanism_info[1]);
+        _nrn_mechanism_register_data_fields(mech_type,
+            _nrn_mechanism_field<double>{"ina"} /* 0 */,
+            _nrn_mechanism_field<double>{"ena"} /* 1 */,
+            _nrn_mechanism_field<double>{"v_unused"} /* 2 */,
+            _nrn_mechanism_field<double>{"g_unused"} /* 3 */,
+            _nrn_mechanism_field<double*>{"ion_ina", "na_ion"} /* 0 */,
+            _nrn_mechanism_field<double*>{"ion_ena", "na_ion"} /* 1 */
+        );
+
+        hoc_register_prop_size(mech_type, 4, 2);
+        hoc_register_dparam_semantics(mech_type, 0, "na_ion");
+        hoc_register_dparam_semantics(mech_type, 1, "na_ion");
+        hoc_register_var(hoc_scalar_double, hoc_vector_double, hoc_intfunc);
+        hoc_register_npy_direct(mech_type, npy_direct_func_proc);
+    }
+}