Move at_time inside coreneuron namespace and redefine at_time in ispc (#403)

* Move at_time inside coreneuron namespace and redefine at_time in ispc
* Explicitely set uniform and varying in at_time ispc function
* Added documentation for the at_time function in C++ and ISPC backends

Co-authored-by: Pramod Kumbhar <[email protected]>

Author: iomaganaris

coreneuron/mechanism/membfunc.hpp

@@ -188,6 +188,8 @@ extern void artcell_net_move(void**, Point_process*, double);
 extern void nrn2ncs_outputevent(int netcon_output_index, double firetime);
 extern bool nrn_use_localgid_;
 extern void net_sem_from_gpu(int sendtype, int i_vdata, int, int ith, int ipnt, double, double);
+#pragma acc routine seq
+extern int at_time(NrnThread*, double);
 
 // _OPENACC and/or NET_RECEIVE_BUFFERING
 extern void net_sem_from_gpu(int, int, int, int, int, double, double);
@@ -196,11 +198,3 @@ extern void hoc_malchk(void); /* just a stub */
 extern void* hoc_Emalloc(size_t);
 
 }  // namespace coreneuron
-
-/*
- * at_time function declared outside of coreneuron namespace
- * because it can be also called from ISPC code. That's also
- * the reason for extern "C".
- */
-#pragma acc routine seq
-extern "C" int at_time(coreneuron::NrnThread*, double);

coreneuron/mechanism/nrnoc_ml.ispc

@@ -168,5 +168,12 @@ struct NrnThread {
 
 extern uniform double ispc_celsius;
 
-extern "C"
-int at_time(NrnThread* nt, varying double te);
+// NOTE : this implementation is duplicated from "coreneuron/network/cvodestb.cpp"
+// If changes are required, make sure to change CPP as well.
+static inline int at_time(uniform NrnThread* nt, varying double te) {
+    varying double x = te - 1e-11;
+    if (x <= nt->_t && x > (nt->_t - nt->_dt)) {
+        return 1;
+    }
+    return 0;
+}

coreneuron/network/cvodestb.cpp

@@ -105,12 +105,14 @@ void fixed_play_continuous(NrnThread* nt) {
     }
 }
 
-}  // namespace coreneuron
-
-extern "C" int at_time(coreneuron::NrnThread* nt, double te) {
+// NOTE : this implementation is duplicated in "coreneuron/mechanism/nrnoc_ml.ispc"
+// for the ISPC backend. If changes are required, make sure to change ISPC as well.
+int at_time(NrnThread* nt, double te) {
     double x = te - 1e-11;
     if (x <= nt->_t && x > (nt->_t - nt->_dt)) {
         return 1;
     }
     return 0;
 }
+
+}  // namespace coreneuron