Avoid use of global variables in generated code

* all global variables are now wrapped into
  struct called _GlobalVars.
* previous global variables are still printed
  but their values is now copied into _GlobalVars
  during nrn_init.
* variables like slist and dlist are now removed
* we now only copy _GlobalVars into GPU; all other
  acc copyin / update for global variables is now
  removed.
* we now print all global variables as static and
  remove all their macro renaming.

Todos:

- [x] basic implementation / draft removing all global usage
- [x] run neuron test suite on CPU
- [ ] run neuron test suite on GPU
- [ ] fix TODOs
- [ ] code cleanup including memory cleanup

Related to BlueBrain/CoreNeuron/issues/141

src/mod2c_core/deriv.c

@@ -564,13 +564,6 @@ void massagederiv(q1, q2, q3, q4, sensused)
 			count++;
 		}
 	}
-	Sprintf(buf,
-	  "\n"
-	  " _slist%d = (int*)malloc(sizeof(int)*%d);\n"
-	  " _dlist%d = (int*)malloc(sizeof(int)*%d);\n"
-	  , numlist, count, numlist, count
-	);
-	Lappendstr(initlist, buf);
 
 	count = 0;
 	ITERATE(qs, deriv_used_list) {
@@ -618,31 +611,19 @@ if (s->subtype & ARRAY) { int dim = s->araydim;
 		}
 	}
 
-	Sprintf(buf,
-	  "#pragma acc enter data copyin(_slist%d[0:%d])\n"
-	  " #pragma acc enter data copyin(_dlist%d[0:%d])\n\n"
-	  , numlist, count, numlist, count);
-	Lappendstr(initlist, buf);
-
 	if (count == 0) {
 		diag("DERIVATIVE contains no derivatives", (char *)0);
 	}
 	derfun->used = count;
+	//TODO: present list could be replaced with  global_data_ptr present clause?
 	Sprintf(buf, ", _slist%d[0:%d], _dlist%d[0:%d]",
 	  numlist, count, numlist, count);
 	Lappendstr(acc_present_list, buf);
-	Sprintf(buf,
 
-	  "\n#define _slist%d _slist%d%s\n"
-	  "int* _slist%d;\n"
-	  "#pragma acc declare create(_slist%d)\n"
-	  "\n#define _dlist%d _dlist%d%s\n"
-	  "int* _dlist%d;\n"
-	  "#pragma acc declare create(_dlist%d)\n"
-	  , numlist, numlist, suffix, numlist, numlist
-	  , numlist, numlist, suffix, numlist, numlist
-	  );
-		Linsertstr(procfunc, buf);
+	Sprintf(buf, "_slist%d", numlist);
+	add_global_var("int", buf, 1, count, 1);
+	Sprintf(buf, "_dlist%d", numlist);
+	add_global_var("int", buf, 1, count, 1);
 
 #if CVODE
 	Lappendstr(procfunc, "\n/*CVODE*/\n");

src/mod2c_core/kinetic.c

@@ -458,18 +458,12 @@ Sprintf(buf, "{int _reset=0;\n");
 	Sprintf(buf, ", _slist%d[0:%d], _dlist%d[0:%d]",
 	  numlist, count, numlist, count);
 	Lappendstr(acc_present_list, buf);
-	Sprintf(buf,
-	  "\n#define _slist%d _slist%d%s\n"
-	  "int* _slist%d;\n"    
-	  "#pragma acc declare create(_slist%d)\n"
-	  "\n#define _dlist%d _dlist%d%s\n"
-	  "int* _dlist%d;\n"   
-	  "#pragma acc declare create(_dlist%d)\n"
-	  , numlist, numlist, suffix, numlist, numlist
-	  , numlist, numlist, suffix, numlist, numlist
-	  );
 
-	Linsertstr(procfunc, buf);
+	Sprintf(buf, "_slist%d", numlist);
+	add_global_var("int", buf, 1, count, 1);
+	Sprintf(buf, "_dlist%d", numlist);
+	add_global_var("int", buf, 1, count, 1);
+
 	insertstr(q4, "  } return _reset;\n");
 	movelist(q1, q4, procfunc);
 
@@ -1218,12 +1212,6 @@ static void kinlist(fun, rlst)
 			count++;
 		}
 	}
-	Sprintf(buf,
-	  "\n _slist%d = (int*)malloc(sizeof(int)*%d);\n"
-	  " _dlist%d = (int*)malloc(sizeof(int)*%d);\n"
-	  , fun->u.i, count, fun->u.i, count);
-	Lappendstr(initlist, buf);
-
 	for (i=0; i < rlst->nsym; i++) {
 		s = rlst->symorder[i];
 #if CVODE
@@ -1274,12 +1262,6 @@ if (vectorize){
 }
 		s->used = 0;
 	}
-	Sprintf(buf,
-	 "#pragma acc enter data copyin(_slist%d[0:%d])\n"
-	 " #pragma acc enter data copyin(_dlist%d[0:%d])\n\n"
-	 , fun->u.i, count, fun->u.i, count);
-	Lappendstr(initlist, buf);
-
 }
 
 /* for now we only check CONSERVE and COMPARTMENT */

src/mod2c_core/modl.c

@@ -78,6 +78,44 @@ extern char* nmodl_version_;
 extern int   usederivstatearray;
 #endif
 
+// global variables, current count and current capacity for array to store
+global_variable_t* global_variables = NULL;
+int global_variables_count = 0;
+int global_variables_capacity = 0;
+
+// add new global variable that needs to be printed
+// TODO: add function for deallocation/cleanup of global_variables
+void add_global_var(const char* type,
+        const char* name,
+        int is_array,
+        int array_length,
+        int skip_initialisation) {
+
+    // if there is no capacity to store variable, reallocate array
+    // note that typically there are small number of global variables
+    // that need to be printed. So ~64 is more than enough in most
+    // of the cases.
+    if (global_variables_capacity <= global_variables_count) {
+        int new_size = sizeof(global_variable_t) * (global_variables_capacity + 64);
+        global_variables = (global_variable_t*) realloc(global_variables, new_size);
+
+        if(global_variables == NULL) {
+            diag("Error while memory allocation, realloc failed!");
+        }
+    }
+
+    // add new variable at the end
+    int i = global_variables_count;
+    strncpy(global_variables[i].type, type, NRN_VARTYPE_BUFSIZE);
+    strncpy(global_variables[i].name, name, NRN_BUFSIZE);
+    global_variables[i].is_array = is_array;
+    global_variables[i].array_length = array_length;
+    global_variables[i].skip_initialisation = skip_initialisation;
+
+    global_variables_count++;
+}
+
+
 extern int yyparse();
 extern int mkdir_p();
 

src/mod2c_core/modl.h

@@ -76,6 +76,9 @@
 #define NRN_BUFSIZE MAX_PATH
 #endif
 
+// buffer to store variable type like int, float or double
+#define NRN_VARTYPE_BUFSIZE 16
+
 typedef struct Item    List;		/* list of mixed items */
 typedef struct Item {
 	short           itemtype;
@@ -310,3 +313,38 @@ extern Item    *qlint;
 #define IGNORE(arg)	arg
 #define Free(arg)	free((void *)(arg))
 #endif
+
+
+// represent a global variable at file scope level
+typedef struct global_variable_t {
+
+    // type is base type like "int" or "double"
+    char type[NRN_VARTYPE_BUFSIZE];
+
+    // name is name of the variable
+    char name[NRN_BUFSIZE];
+
+    // > 0 if variable is of type array
+    int is_array;
+
+    // length of array if is_array > 0
+    int array_length;
+
+    // > 0 if this variable doesn't need to be
+    // initialized from another global variable
+    // in the file.
+    int skip_initialisation;
+
+} global_variable_t;
+
+// all global variables in the program
+extern struct global_variable_t* global_variables;
+
+// count of global variables
+extern int global_variables_count;
+
+// current capacity of global variable array
+extern int global_variables_capacity;
+
+// add new global variable
+void add_global_var(const char* type, const char* name, int is_array, int array_length, int skip_initialisation);

src/mod2c_core/noccout.c

@@ -50,6 +50,7 @@ extern List	*begin_dion_stmt(), *end_dion_stmt();
 extern List* conductance_;
 extern List* breakpoint_local_current_;
 extern List* newtonspace_list;
+extern List* globals_update_list;
 static void conductance_cout();
 extern int net_send_buffer_in_initial;
 #endif
@@ -751,6 +752,7 @@ void c_out_vectorize()
 	P("#undef PI\n");
 	printlist(defs_list);
 	printlist(firstlist);
+	printlist(globals_update_list);
 	if (modelline) {
 		Fprintf(fcout, "static const char *modelname = \"%s\";\n\n", modelline);
 	} else {
@@ -806,9 +808,6 @@ void c_out_vectorize()
 	  , derivimplic_listnum, derivimplic_listnum);
  	  P(buf);
 	}
-	  if (net_send_seen_ && !artificial_cell) {
-	    P("  #pragma acc update device (_mechtype) if(_nt->compute_gpu)\n");
-	  }
 	ITERATE(q, newtonspace_list) {
 	  P(STR(q));
 	}
@@ -828,7 +827,7 @@ void c_out_vectorize()
        With c++ files we declare routine seq in which case static variables
        can't be used. So, introducing _celsius_ as default implementation.
 	*/
-	P("_acc_globals_update();\n");
+	P("_update_global_variables();\n");
 
 	 pr_layout_for_p(1, NRN_INIT);
 
@@ -1074,6 +1073,10 @@ void c_out_vectorize()
 	*/
 	/* initlists() is called once to setup slist and dlist pointers */
 	P("\nstatic void _initlists(){\n");
+	// _initlists is called first from the _reg function
+	// and hence setup global variable pointer here so
+	// that it's valid to begin with
+	P(" _global_variables_ptr = &_global_variables;\n");
 	P(" double _x; double* _p = &_x;\n");
 	P(" int _i; static int _first = 1;\n");
 	P(" int _cntml_actual=1;\n");