static void nrn_jacob(_NrnThread* _nt, _Memb_list* _ml, int _type){
Node *_nd; int* _ni; int _iml, _cntml;
#if CACHEVEC
_ni = _ml->_nodeindices;
#endif
_cntml = _ml->_nodecount;
for (_iml = 0; _iml < _cntml; ++_iml) {
_p = _ml->_data[_iml];
#if CACHEVEC
if (use_cachevec) {
VEC_D(_ni[_iml]) += _g;
}else
#endif
{
_nd = _ml->_nodelist[_iml];
NODED(_nd) += _g;
}
}}
/* back substitution to finish solving the matrix equations */
static void bksub(NrnThread* _nt) {
int i, i1, i2, i3;
i1 = 0;
i2 = i1 + _nt->ncell;
i3 = _nt->end;
double* vec_b = &(VEC_B(0));
double* vec_d = &(VEC_D(0));
double* vec_rhs = &(VEC_RHS(0));
int* parent_index = _nt->_v_parent_index;
#if defined(_OPENACC)
int stream_id = _nt->stream_id;
#endif
/** @todo: just for benchmarking, otherwise produces wrong results */
// clang-format off
#pragma acc parallel loop seq present( \
vec_d[0:i2], vec_rhs[0:i2]) \
async(stream_id) if (_nt->compute_gpu)
// clang-format on
for (i = i1; i < i2; ++i) {
vec_rhs[i] /= vec_d[i];
}
/** @todo: just for benchmarking, otherwise produces wrong results */
// clang-format off
#pragma acc parallel loop seq present( \
vec_b[0:i3], vec_d[0:i3], vec_rhs[0:i3], \
parent_index[0:i3]) async(stream_id) \
if (_nt->compute_gpu)
for (i = i2; i < i3; ++i) {
vec_rhs[i] -= vec_b[i] * vec_rhs[parent_index[i]];
vec_rhs[i] /= vec_d[i];
}
#pragma acc wait(stream_id)
// clang-format on
}
