Add 'net_send'. (#11)

net_send/coreneuron/toggle.cpp

@@ -0,0 +1,379 @@
+/*********************************************************
+Model Name      : toggle
+Filename        : toggle.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",
+        "toggle",
+        0,
+        "y",
+        0,
+        0,
+        0
+    };
+
+
+    /** all global variables */
+    struct toggle_Store {
+        int point_type{};
+        int reset{};
+        int mech_type{};
+    };
+    static_assert(std::is_trivially_copy_constructible_v<toggle_Store>);
+    static_assert(std::is_trivially_move_constructible_v<toggle_Store>);
+    static_assert(std::is_trivially_copy_assignable_v<toggle_Store>);
+    static_assert(std::is_trivially_move_assignable_v<toggle_Store>);
+    static_assert(std::is_trivially_destructible_v<toggle_Store>);
+    toggle_Store toggle_global;
+
+
+    /** all mechanism instance variables and global variables */
+    struct toggle_Instance  {
+        double* y{};
+        double* v_unused{};
+        double* tsave{};
+        const double* node_area{};
+        const int* point_process{};
+        const int* tqitem{};
+        toggle_Store* global{&toggle_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 num_net_receive_args() {
+        return 1;
+    }
+
+
+    static inline int float_variables_size() {
+        return 3;
+    }
+
+
+    static inline int int_variables_size() {
+        return 3;
+    }
+
+
+    static inline int get_mech_type() {
+        return toggle_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_toggle(NrnThread* nt, Memb_list* ml, int type) {
+        assert(!ml->instance);
+        assert(!ml->global_variables);
+        assert(ml->global_variables_size == 0);
+        auto* const inst = new toggle_Instance{};
+        assert(inst->global == &toggle_global);
+        ml->instance = inst;
+        ml->global_variables = inst->global;
+        ml->global_variables_size = sizeof(toggle_Store);
+    }
+
+    // Deallocate the instance structure
+    static void nrn_private_destructor_toggle(NrnThread* nt, Memb_list* ml, int type) {
+        auto* const inst = static_cast<toggle_Instance*>(ml->instance);
+        assert(inst);
+        assert(inst->global);
+        assert(inst->global == &toggle_global);
+        assert(inst->global == ml->global_variables);
+        assert(ml->global_variables_size == sizeof(toggle_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<toggle_Instance*>(ml->instance);
+        assert(inst);
+        assert(inst->global);
+        assert(inst->global == &toggle_global);
+        assert(inst->global == ml->global_variables);
+        assert(ml->global_variables_size == sizeof(toggle_Store));
+        int pnodecount = ml->_nodecount_padded;
+        Datum* indexes = ml->pdata;
+        inst->y = ml->data+0*pnodecount;
+        inst->v_unused = ml->data+1*pnodecount;
+        inst->tsave = ml->data+2*pnodecount;
+        inst->node_area = nt->_data;
+        inst->point_process = ml->pdata;
+        inst->tqitem = ml->pdata;
+    }
+
+
+
+    static void nrn_alloc_toggle(double* data, Datum* indexes, int type) {
+        // do nothing
+    }
+
+
+    void nrn_constructor_toggle(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<toggle_Instance*>(ml->instance);
+
+        #endif
+    }
+
+
+    void nrn_destructor_toggle(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<toggle_Instance*>(ml->instance);
+
+        #endif
+    }
+
+
+    static inline void net_send_buffering(const NrnThread* nt, NetSendBuffer_t* nsb, int type, int vdata_index, int weight_index, int point_index, double t, double flag) {
+        int i = 0;
+        i = nsb->_cnt++;
+        if (i >= nsb->_size) {
+            nsb->grow();
+        }
+        if (i < nsb->_size) {
+            nsb->_sendtype[i] = type;
+            nsb->_vdata_index[i] = vdata_index;
+            nsb->_weight_index[i] = weight_index;
+            nsb->_pnt_index[i] = point_index;
+            nsb->_nsb_t[i] = t;
+            nsb->_nsb_flag[i] = flag;
+        }
+    }
+
+
+    static inline void net_receive_kernel_toggle(double t, Point_process* pnt, toggle_Instance* inst, NrnThread* nt, Memb_list* ml, int weight_index, double flag) {
+        int tid = pnt->_tid;
+        int id = pnt->_i_instance;
+        double v = 0;
+        int nodecount = ml->nodecount;
+        int pnodecount = ml->_nodecount_padded;
+        double* data = ml->data;
+        double* weights = nt->weights;
+        Datum* indexes = ml->pdata;
+        ThreadDatum* thread = ml->_thread;
+
+        inst->tsave[id] = t;
+        {
+            inst->y[id] = 1.0;
+        }
+    }
+
+
+    static void net_receive_toggle(Point_process* pnt, int weight_index, double flag) {
+        NrnThread* nt = nrn_threads + pnt->_tid;
+        Memb_list* ml = get_memb_list(nt);
+        NetReceiveBuffer_t* nrb = ml->_net_receive_buffer;
+        if (nrb->_cnt >= nrb->_size) {
+            realloc_net_receive_buffer(nt, ml);
+        }
+        int id = nrb->_cnt;
+        nrb->_pnt_index[id] = pnt-nt->pntprocs;
+        nrb->_weight_index[id] = weight_index;
+        nrb->_nrb_t[id] = nt->_t;
+        nrb->_nrb_flag[id] = flag;
+        nrb->_cnt++;
+    }
+
+
+    void net_buf_receive_toggle(NrnThread* nt) {
+        Memb_list* ml = get_memb_list(nt);
+        if (!ml) {
+            return;
+        }
+
+        NetReceiveBuffer_t* nrb = ml->_net_receive_buffer;
+        auto* const inst = static_cast<toggle_Instance*>(ml->instance);
+        int count = nrb->_displ_cnt;
+        #pragma omp simd
+        #pragma ivdep
+        for (int i = 0; i < count; i++) {
+            int start = nrb->_displ[i];
+            int end = nrb->_displ[i+1];
+            for (int j = start; j < end; j++) {
+                int index = nrb->_nrb_index[j];
+                int offset = nrb->_pnt_index[index];
+                double t = nrb->_nrb_t[index];
+                int weight_index = nrb->_weight_index[index];
+                double flag = nrb->_nrb_flag[index];
+                Point_process* point_process = nt->pntprocs + offset;
+                net_receive_kernel_toggle(t, point_process, inst, nt, ml, weight_index, flag);
+            }
+        }
+        nrb->_displ_cnt = 0;
+        nrb->_cnt = 0;
+
+        NetSendBuffer_t* nsb = ml->_net_send_buffer;
+        for (int i=0; i < nsb->_cnt; i++) {
+            int type = nsb->_sendtype[i];
+            int tid = nt->id;
+            double t = nsb->_nsb_t[i];
+            double flag = nsb->_nsb_flag[i];
+            int vdata_index = nsb->_vdata_index[i];
+            int weight_index = nsb->_weight_index[i];
+            int point_index = nsb->_pnt_index[i];
+            net_sem_from_gpu(type, vdata_index, weight_index, tid, point_index, t, flag);
+        }
+        nsb->_cnt = 0;
+    }
+
+
+    /** initialize channel */
+    void nrn_init_toggle(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<toggle_Instance*>(ml->instance);
+
+        if (_nrn_skip_initmodel == 0) {
+            #pragma omp simd
+            #pragma ivdep
+            for (int id = 0; id < nodecount; id++) {
+                inst->tsave[id] = -1e20;
+                int node_id = node_index[id];
+                double v = voltage[node_id];
+                #if NRN_PRCELLSTATE
+                inst->v_unused[id] = v;
+                #endif
+                inst->y[id] = 0.0;
+                net_send_buffering(nt, ml->_net_send_buffer, 0, inst->tqitem[2*pnodecount+id], 0, inst->point_process[1*pnodecount+id], nt->_t+2.0, 1.0);
+            }
+        }
+
+        NetSendBuffer_t* nsb = ml->_net_send_buffer;
+        for (int i=0; i < nsb->_cnt; i++) {
+            int type = nsb->_sendtype[i];
+            int tid = nt->id;
+            double t = nsb->_nsb_t[i];
+            double flag = nsb->_nsb_flag[i];
+            int vdata_index = nsb->_vdata_index[i];
+            int weight_index = nsb->_weight_index[i];
+            int point_index = nsb->_pnt_index[i];
+            net_sem_from_gpu(type, vdata_index, weight_index, tid, point_index, t, flag);
+        }
+        nsb->_cnt = 0;
+    }
+
+
+    /** register channel with the simulator */
+    void _toggle_reg() {
+
+        int mech_type = nrn_get_mechtype("toggle");
+        toggle_global.mech_type = mech_type;
+        if (mech_type == -1) {
+            return;
+        }
+
+        _nrn_layout_reg(mech_type, 0);
+        point_register_mech(mechanism_info, nrn_alloc_toggle, nullptr, nullptr, nullptr, nrn_init_toggle, nrn_private_constructor_toggle, nrn_private_destructor_toggle, first_pointer_var_index(), nullptr, nullptr, 1);
+
+        hoc_register_prop_size(mech_type, float_variables_size(), int_variables_size());
+        hoc_register_dparam_semantics(mech_type, 0, "area");
+        hoc_register_dparam_semantics(mech_type, 1, "pntproc");
+        hoc_register_dparam_semantics(mech_type, 2, "netsend");
+        hoc_register_net_receive_buffering(net_buf_receive_toggle, mech_type);
+        set_pnt_receive(mech_type, net_receive_toggle, nullptr, num_net_receive_args());
+        hoc_register_net_send_buffering(mech_type);
+        hoc_register_var(hoc_scalar_double, hoc_vector_double, NULL);
+    }
+}