Issue with membrane pointer list in OpenACC backend

[pramodk]

Currently CoreNEURON copied membrane list to GPU as:

        /* nt._ml_list is used in NET_RECEIVE block and should have valid membrane list id*/
        Memb_list** d_ml_list =
            (Memb_list**)acc_copyin(nt->_ml_list, n_memb_func * sizeof(Memb_list*));
        acc_memcpy_to_device(&(d_nt->_ml_list), &(d_ml_list), sizeof(Memb_list**));

Some mechanisms kernel use those in compute block as:

    void nrn_init_CaDynamics_DC0(NrnThread* nt, Memb_list* ml, int type) {
        #pragma acc data present(nt, ml, CaDynamics_DC0_global)
        {
            int nodecount = ml->nodecount;
            int pnodecount = ml->_nodecount_padded;
            const int* __restrict__ node_index = ml->nodeindices;
            double* __restrict__ data = ml->data;
            const double* __restrict__ voltage = nt->_actual_v;
            Datum* __restrict__ indexes = ml->pdata;
            ThreadDatum* __restrict__ thread = ml->_thread;

            setup_instance(nt, ml);
            CaDynamics_DC0_Instance* __restrict__ inst = (CaDynamics_DC0_Instance*) ml->instance;
            if (_nrn_skip_initmodel == 0) {
                int start = 0;
                int end = nodecount;
                #pragma acc parallel loop present(inst, node_index, data, voltage, indexes, thread)
                for (int id = start; id < end; id++)  {
                    int node_id = node_index[id];
                    double v = voltage[node_id];
                    IonCurVar ionvar;
                    ionvar.cai = inst->ion_cai[indexes[1*pnodecount+id]];
                    inst->surftovol[id] = 4.0 / inst->diam[indexes[3*pnodecount+id]];
                    ionvar.cai = CaDynamics_DC0_global.minCai;
                    inst->ion_cai[indexes[1*pnodecount+id]] = ionvar.cai;
                    int ca_type = CaDynamics_DC0_global.ca_type;
                    nrn_wrote_conc(ca_type, &(inst->ion_cai[indexes[1*pnodecount+id]]), 1, inst->style_ca[2], nrn_ion_global_map, celsius, nt->_ml_list[ca_type]->_nodecount_padded);
                }
            }
        }
    }

this sgfaults whenever nt->_ml_list[ca_type] is accessed. If I try to add nt->_ml_list in present clause as:

            Memb_list** ml_list = nt->_ml_list;
            if (_nrn_skip_initmodel == 0) {
                int start = 0;
                int end = nodecount;
                #pragma acc parallel loop present(inst, node_index, data, voltage, indexes, thread, ml_list)

then we get :

FATAL ERROR: data in PRESENT clause was not found on device 1: name=ml_list host:0x1d60470

Looking at CPU/GPU pi=pointers using PGI_ACC_DEBUG=1 tells that the data is present on GPU.

I am missing something but it's not clear at the moment :)

[pramodk]

Oops..the ca_type is not initialised when mechanism is registered:

    /** register channel with the simulator */
    void _CaDynamics_DC0_reg()  {
        setup_global_variables();
        int mech_type = nrn_get_mechtype("CaDynamics_DC0");
        CaDynamics_DC0_global.mech_type = mech_type;
        if (mech_type == -1) {
            return;
        }

        _nrn_layout_reg(mech_type, get_memory_layout());
        register_mech(mechanism, nrn_alloc_CaDynamics_DC0, NULL, NULL, nrn_state_CaDynamics_DC0, nrn_init_CaDynamics_DC0, first_pointer_var_index(), 1);
        CaDynamics_DC0_global.ca_type = nrn_get_mechtype("ca_ion");

setup_global_variables should be called after ca_type is set.

Note : the error about ml_list in PRESENT clause was not found on device is still bit puzzle. But this also happens with mod2c generated code even though the code works perfectly fine. Most likely the issue with semantic of array of pointers. Will take care of this later.

[pramodk]

Addressed in #146, specifically in the commit 49ea279164edd16ef28c8b1c54c5b1672616d7ca.