Repository: myokit/myokit (language: python)

Issue:
OpenCL sim: Don't create and pass around idiff vector when diffusion is not enabled
Slight optimisation. Currently size 1 so shouldn't matter too much.


Discussion / hints:

Fix patch:
diff --git a/myokit/_sim/fiber_tissue.c b/myokit/_sim/fiber_tissue.c
index 220b43b48..2acb1d5df 100644
--- a/myokit/_sim/fiber_tissue.c
+++ b/myokit/_sim/fiber_tissue.c
@@ -77,7 +77,7 @@ void pyprint(const char* text)
  *  var      : The variable to add to the logs, if its name is present
  * Returns 0 if not added, 1 if added.
  */
-static int log_add(PyObject* log_dict, PyObject** logs, Real** vars, int i, char* name, const Real* var)
+static int log_add(PyObject* log_dict, PyObject** logs, Real** vars, unsigned long i, char* name, const Real* var)
 {
     int added = 0;
     PyObject* key = PyUnicode_FromString(name);
@@ -102,10 +102,10 @@ PyObject *platform_name;// A python string specifying the platform to use
 PyObject *device_name;  // A python string specifying the device to use
 char* kernel_source_f;  // The kernel code for the fiber model
 char* kernel_source_t;  // The kernel code for the tissue model
-int nfx;                // The number of cells in the x direction (fiber)
-int nfy;                // The number of cells in the y direction (fiber)
-int ntx;                // The number of cells in the x direction (tissue)
-int nty;                // The number of cells in the y direction (tissue)
+unsigned long nfx;      // The number of cells in the x direction (fiber)
+unsigned long nfy;      // The number of cells in the y direction (fiber)
+unsigned long ntx;      // The number of cells in the x direction (tissue)
+unsigned long nty;      // The number of cells in the y direction (tissue)
 int i_vm_f;             // The state index of the membrane potential (fiber)
 int i_vm_t;             // The state index of the membrane potential (tissue)
 double gfx;             // The cell-to-cell conductance in the fiber direction
@@ -121,9 +121,9 @@ PyObject* state_in_t;   // The initial state (tissue)
 PyObject* state_out_f;  // The final state (fiber)
 PyObject* state_out_t;  // The final state (tissue)
 PyObject *protocol;     // A pacing protocol
-int cfx;                // The x-coord. on the fiber where the tissue connects
-int ctx;                // The x-coord. on the tissue where the fiber connects
-int cty;                // The first connected y-coord. on the tissue
+unsigned long cfx;      // The x-coord. on the fiber where the tissue connects
+unsigned long ctx;      // The x-coord. on the tissue where the fiber connects
+unsigned long cty;      // The first connected y-coord. on the tissue
 PyObject *log_dict_f;   // A logging dict for the fiber
 PyObject *log_dict_t;   // A logging dict for the tissue
 double log_interval;    // The time between log writes
@@ -194,23 +194,23 @@ Real arg_gtx;
 Real arg_gty;
 Real arg_gft;
 
-// Logging
-PyObject** logs_f;      // An array of pointers to a PyObject (fiber)
-PyObject** logs_t;      // An array of pointers to a PyObject (tissue)
-Real** vars_f;          // An array of pointers to values to log (fiber)
-Real** vars_t;          // An array of pointers to values to log (tissue)
-int n_vars_f;           // Number of logging variables (fiber)
-int n_vars_t;           // Number of logging variables (tissue)
-double tnext_log;       /* The next logging point */
-unsigned long inext_log;// The number of logged steps
-int logging_states_f;   // True if any states are being logged
-int logging_states_t;   // True if any states are being logged
-int logging_diffusion_f;// True if diffusion current is being logged.
-int logging_diffusion_t;// True if diffusion current is being logged.
-int logging_inters_f;   // True if any intermediary vars are being logged
-int logging_inters_t;   // True if any intermediary vars are being logged
-int n_inter_f;          // The number of unique intermediary variables logged
-int n_inter_t;          // The number of unique intermediary variables logged
+/* Logging */
+PyObject** logs_f;       // An array of pointers to a PyObject (fiber)
+PyObject** logs_t;       // An array of pointers to a PyObject (tissue)
+Real** vars_f;           // An array of pointers to values to log (fiber)
+Real** vars_t;           // An array of pointers to values to log (tissue)
+unsigned long n_vars_f;  // Number of logging variables (fiber)
+unsigned long n_vars_t;  // Number of logging variables (tissue)
+double tnext_log;        /* The next logging point */
+unsigned long inext_log; // The number of logged steps
+int logging_states_f;    // True if any states are being logged
+int logging_states_t;    // True if any states are being logged
+int logging_diffusion_f; // True if diffusion current is being logged.
+int logging_diffusion_t; // True if diffusion current is being logged.
+int logging_inters_f;    // True if any intermediary vars are being logged
+int logging_inters_t;    // True if any intermediary vars are being logged
+unsigned long n_inter_f; // The number of unique intermediary variables logged
+unsigned long n_inter_t; // The number of unique intermediary variables logged
 
 // Temporary objects: decref before re-using for another var
 // (Unless you got it through PyList_GetItem or PyTuble_GetItem)
@@ -310,7 +310,7 @@ sim_init(PyObject* self, PyObject* args)
     cl_int flag;
 
     // Iteration
-    int i, j, k;
+    unsigned long i, j, k;
 
     // Platform and device id
     cl_platform_id platform_id;
@@ -318,14 +318,14 @@ sim_init(PyObject* self, PyObject* args)
 
     // Variable names
     char log_var_name[1023];
-    int k_vars;
+    unsigned long k_vars;
 
     // Compilation error message
     size_t blog_size;
     char *blog;
 
     // Cell coupling
-    int nsf, nst;
+    unsigned long nsf, nst;
 
     #ifdef MYOKIT_DEBUG
     printf("Starting initialization.\n");
@@ -367,12 +367,12 @@ sim_init(PyObject* self, PyObject* args)
     vars_t = NULL;
 
     // Check input arguments
-    if(!PyArg_ParseTuple(args, "OOssiiiiiidddddiiidddOOOOOOOdOO",
+    if(!PyArg_ParseTuple(args, "OOsskkkkiidddddkkkdddOOOOOOOdOO",
             &platform_name,
             &device_name,
             &kernel_source_f,
             &kernel_source_t,
-            &nfx,
+            &nfx,       // Small 'k' = unsigned long
             &nfy,
             &ntx,
             &nty,
@@ -456,19 +456,19 @@ sim_init(PyObject* self, PyObject* args)
         PyErr_SetString(PyExc_Exception, "'state_out_t' must be a list.");
         return sim_clean();
     }
-    if(PyList_Size(state_in_f) != nfx * nfy * n_state_f) {
+    if((unsigned long)PyList_Size(state_in_f) != nfx * nfy * n_state_f) {
         PyErr_SetString(PyExc_Exception, "'state_in_f' must have size nfx * nfy * n_states_f.");
         return sim_clean();
     }
-    if(PyList_Size(state_in_t) != ntx * nty * n_state_t) {
+    if((unsigned long)PyList_Size(state_in_t) != ntx * nty * n_state_t) {
         PyErr_SetString(PyExc_Exception, "'state_in_t' must have size ntx * nty * n_states_t.");
         return sim_clean();
     }
-    if(PyList_Size(state_out_f) != nfx * nfy * n_state_f) {
+    if((unsigned long)PyList_Size(state_out_f) != nfx * nfy * n_state_f) {
         PyErr_SetString(PyExc_Exception, "'state_out_f' must have size nfx * nfy * n_states_f.");
         return sim_clean();
     }
-    if(PyList_Size(state_out_t) != ntx * nty * n_state_t) {
+    if((unsigned long)PyList_Size(state_out_t) != ntx * nty * n_state_t) {
         PyErr_SetString(PyExc_Exception, "'state_out_t' must have size ntx * nty * n_states_t.");
         return sim_clean();
     }
@@ -494,7 +494,6 @@ sim_init(PyObject* self, PyObject* args)
     //
     // Set up pacing system
     //
-    flag_pacing;
     pacing = ESys_Create(&flag_pacing);
     if(flag_pacing!=ESys_OK) { ESys_SetPyErr(flag_pacing); return sim_clean(); }
     flag_pacing = ESys_Populate(pacing, protocol);
@@ -538,7 +537,7 @@ sim_init(PyObject* self, PyObject* args)
         flt = PyList_GetItem(state_in_f, i);    // Don't decref!
         if(!PyFloat_Check(flt)) {
             char errstr[200];
-            sprintf(errstr, "Item %d in fiber state vector is not a float.", i);
+            sprintf(errstr, "Item %u in fiber state vector is not a float.", (unsigned int)i);
             PyErr_SetString(PyExc_Exception, errstr);
             return sim_clean();
         }
@@ -554,7 +553,7 @@ sim_init(PyObject* self, PyObject* args)
         flt = PyList_GetItem(state_in_t, i);
         if(!PyFloat_Check(flt)) {
             char errstr[200];
-            sprintf(errstr, "Item %d in tissue state vector is not a float.", i);
+            sprintf(errstr, "Item %u in tissue state vector is not a float.", (unsigned int)i);
             PyErr_SetString(PyExc_Exception, errstr);
             return sim_clean();
         }
@@ -715,7 +714,7 @@ sim_init(PyObject* self, PyObject* args)
         clGetProgramBuildInfo(program_t, device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &blog_size);
         blog = (char*)malloc(blog_size);
         clGetProgramBuildInfo(program_t, device_id, CL_PROGRAM_BUILD_LOG, blog_size, blog, NULL);
-        fprintf(stderr, "OpenCL Error: Fiber kernel failed to compile.\n");
+        fprintf(stderr, "OpenCL Error: Tissue kernel failed to compile.\n");
         fprintf(stderr, "----------------------------------------");
         fprintf(stderr, "---------------------------------------\n");
         fprintf(stderr, "%s\n", blog);
@@ -754,6 +753,9 @@ sim_init(PyObject* self, PyObject* args)
     if(mcl_flag(clSetKernelArg(kernel_cell_f, i++, sizeof(mbuf_idiff_f), &mbuf_idiff_f))) return sim_clean();
     if(mcl_flag(clSetKernelArg(kernel_cell_f, i++, sizeof(mbuf_inter_log_f), &mbuf_inter_log_f))) return sim_clean();
     if(mcl_flag(clSetKernelArg(kernel_cell_f, i++, sizeof(mbuf_field_data), &mbuf_field_data))) return sim_clean();
+    #ifdef MYOKIT_DEBUG
+    printf("Set up cell kernel for fiber model.\n");
+    #endif
 
     i = 0;
     if(mcl_flag(clSetKernelArg(kernel_diff_f, i++, sizeof(nfx), &nfx))) return sim_clean();
@@ -762,6 +764,9 @@ sim_init(PyObject* self, PyObject* args)
     if(mcl_flag(clSetKernelArg(kernel_diff_f, i++, sizeof(arg_gfy), &arg_gfy))) return sim_clean();
     if(mcl_flag(clSetKernelArg(kernel_diff_f, i++, sizeof(mbuf_state_f), &mbuf_state_f))) return sim_clean();
     if(mcl_flag(clSetKernelArg(kernel_diff_f, i++, sizeof(mbuf_idiff_f), &mbuf_idiff_f))) return sim_clean();
+    #ifdef MYOKIT_DEBUG
+    printf("Set up diffusion kernel for fiber model.\n");
+    #endif
 
     i = 0;
     if(mcl_flag(clSetKernelArg(kernel_cell_t, i++, sizeof(ntx), &ntx))) return sim_clean();
@@ -773,6 +778,9 @@ sim_init(PyObject* self, PyObject* args)
     if(mcl_flag(clSetKernelArg(kernel_cell_t, i++, sizeof(mbuf_idiff_t), &mbuf_idiff_t))) return sim_clean();
     if(mcl_flag(clSetKernelArg(kernel_cell_t, i++, sizeof(mbuf_inter_log_t), &mbuf_inter_log_t))) return sim_clean();
     if(mcl_flag(clSetKernelArg(kernel_cell_t, i++, sizeof(mbuf_field_data), &mbuf_field_data))) return sim_clean();
+    #ifdef MYOKIT_DEBUG
+    printf("Set up cell kernel for tissue model.\n");
+    #endif
 
     i = 0;
     if(mcl_flag(clSetKernelArg(kernel_diff_t, i++, sizeof(ntx), &ntx))) return sim_clean();
@@ -781,6 +789,9 @@ sim_init(PyObject* self, PyObject* args)
     if(mcl_flag(clSetKernelArg(kernel_diff_t, i++, sizeof(arg_gty), &arg_gty))) return sim_clean();
     if(mcl_flag(clSetKernelArg(kernel_diff_t, i++, sizeof(mbuf_state_t), &mbuf_state_t))) return sim_clean();
     if(mcl_flag(clSetKernelArg(kernel_diff_t, i++, sizeof(mbuf_idiff_t), &mbuf_idiff_t))) return sim_clean();
+    #ifdef MYOKIT_DEBUG
+    printf("Set up diffusion kernel for tissue model.\n");
+    #endif
 
     /* Set indices of coupled cell indexes in state and diff vectors */
     nsf = n_state_f;
@@ -799,9 +810,8 @@ sim_init(PyObject* self, PyObject* args)
     if(mcl_flag(clSetKernelArg(kernel_diff_ft, 11, sizeof(mbuf_state_t), &mbuf_state_t))) return sim_clean();
     if(mcl_flag(clSetKernelArg(kernel_diff_ft, 12, sizeof(mbuf_idiff_f), &mbuf_idiff_f))) return sim_clean();
     if(mcl_flag(clSetKernelArg(kernel_diff_ft, 13, sizeof(mbuf_idiff_t), &mbuf_idiff_t))) return sim_clean();
-
     #ifdef MYOKIT_DEBUG
-    printf("Arguments passed into kernels.\n");
+    printf("Set up fiber-tissue diffusion kernel.\n");
     #endif
 
     /*
@@ -842,7 +852,7 @@ if var is not None:
 <?
 var = modelf.binding('diffusion_current')
 if var is not None:
-    print(3*tab + 'sprintf(log_var_name, "%d.%d.' + var.qname() + '", j, i);')
+    print(3*tab + 'sprintf(log_var_name, "%u.%u.' + var.qname() + '", (unsigned int)j, (unsigned int)i);')
     print(3*tab + 'if(log_add(log_dict_f, logs_f, vars_f, k_vars, log_var_name, &rvec_idiff_f[i*nfx+j])) {')
     print(4*tab + 'logging_diffusion_f = 1;')
     print(4*tab + 'k_vars++;')
@@ -857,7 +867,7 @@ if var is not None:
         for(j=0; j<nfx; j++) {
 <?
 for var in modelf.states():
-    print(3*tab + 'sprintf(log_var_name, "%d.%d.' + var.qname() + '", j, i);' )
+    print(3*tab + 'sprintf(log_var_name, "%u.%u.' + var.qname() + '", (unsigned int)j, (unsigned int)i);' )
     print(3*tab + 'if(log_add(log_dict_f, logs_f, vars_f, k_vars, log_var_name, &rvec_state_f[(i*nfx+j)*n_state_f+' + str(var.indice()) + '])) {')
     print(4*tab + 'logging_states_f = 1;')
     print(4*tab + 'k_vars++;')
@@ -873,7 +883,7 @@ for var in modelf.states():
             for(k=0; k<n_inter_f; k++) {
                 ret = PyList_GetItem(inter_log_f, k); /* Don't decref */
 <?
-print(4*tab + 'sprintf(log_var_name, "%d.%d.%s", j, i, PyBytes_AsString(ret));')
+print(4*tab + 'sprintf(log_var_name, "%u.%u.%s", (unsigned int)j, (unsigned int)i, PyBytes_AsString(ret));')
 print(4*tab + 'if(log_add(log_dict_f, logs_f, vars_f, k_vars, log_var_name, &rvec_inter_log_f[(i*nfx+j)*n_inter_f+k])) {')
 print(5*tab + 'logging_inters_f = 1;')
 print(5*tab + 'k_vars++;')
@@ -891,7 +901,7 @@ print(4*tab + '}')
     }
 
     #ifdef MYOKIT_DEBUG
-    printf("Created log for %d fiber variables.\n", n_vars_f);
+    printf("Created log for %u fiber variables.\n", (unsigned int)n_vars_f);
     #endif
 
     /* Now set up tissue logging */
@@ -913,7 +923,7 @@ if var is not None:
 <?
 var = modelt.binding('diffusion_current')
 if var is not None:
-    print(3*tab + 'sprintf(log_var_name, "%d.%d.' + var.qname() + '", j, i);')
+    print(3*tab + 'sprintf(log_var_name, "%u.%u.' + var.qname() + '", (unsigned int)j, (unsigned int)i);')
     print(3*tab + 'if(log_add(log_dict_t, logs_t, vars_t, k_vars, log_var_name, &rvec_idiff_t[i*ntx+j])) {')
     print(4*tab + 'logging_diffusion_t = 1;')
     print(4*tab + 'k_vars++;')
@@ -928,7 +938,7 @@ if var is not None:
         for(j=0; j<ntx; j++) {
 <?
 for var in modelt.states():
-    print(3*tab + 'sprintf(log_var_name, "%d.%d.' + var.qname() + '", j, i);' )
+    print(3*tab + 'sprintf(log_var_name, "%u.%u.' + var.qname() + '", (unsigned int)j, (unsigned int)i);' )
     print(3*tab + 'if(log_add(log_dict_t, logs_t, vars_t, k_vars, log_var_name, &rvec_state_t[(i*ntx+j)*n_state_t+' + str(var.indice()) + '])) {')
     print(4*tab + 'logging_states_t = 1;')
     print(4*tab + 'k_vars++;')
@@ -944,7 +954,7 @@ for var in modelt.states():
             for(k=0; k<n_inter_t; k++) {
                 ret = PyList_GetItem(inter_log_t, k); /* Don't decref */
 <?
-print(4*tab + 'sprintf(log_var_name, "%d.%d.%s", j, i, PyBytes_AsString(ret));')
+print(4*tab + 'sprintf(log_var_name, "%u.%u.%s", (unsigned int)j, (unsigned int)i, PyBytes_AsString(ret));')
 print(4*tab + 'if(log_add(log_dict_t, logs_t, vars_t, k_vars, log_var_name, &rvec_inter_log_t[(i*ntx+j)*n_inter_t+k])) {')
 print(5*tab + 'logging_inters_t = 1;')
 print(5*tab + 'k_vars++;')
@@ -962,7 +972,7 @@ print(4*tab + '}')
     }
 
     #ifdef MYOKIT_DEBUG
-    printf("Created log for %d tissue variables.\n", n_vars_t);
+    printf("Created log for %u tissue variables.\n", (unsigned int)n_vars_t);
     #endif
 
     /* Log update method: */
@@ -991,9 +1001,9 @@ static PyObject*
 sim_step(PyObject *self, PyObject *args)
 {
     ESys_Flag flag_pacing;
-    long steps_left_in_run;
+    unsigned long steps_left_in_run;
     cl_int flag;
-    int i;
+    unsigned long i;
     double d;
     int logging_condition;
 
@@ -1185,8 +1195,8 @@ sim_step(PyObject *self, PyObject *args)
  */
 static PyMethodDef SimMethods[] = {
     {"sim_init", sim_init, METH_VARARGS, "Initialize the simulation."},
-    {"sim_step", sim_step, METH_VARARGS, "Perform the next step in the simulation."},
-    {"sim_clean", py_sim_clean, METH_VARARGS, "Clean up after an aborted simulation."},
+    {"sim_step", sim_step, METH_NOARGS, "Perform the next step in the simulation."},
+    {"sim_clean", py_sim_clean, METH_NOARGS, "Clean up after an aborted simulation."},
     {NULL},
 };
 
diff --git a/myokit/_sim/fiber_tissue.py b/myokit/_sim/fiber_tissue.py
index 5cbc58e26..7c2d2814a 100644
--- a/myokit/_sim/fiber_tissue.py
+++ b/myokit/_sim/fiber_tissue.py
@@ -11,9 +11,6 @@
 import myokit
 import platform
 
-# OpenCLSim keywords
-from .openclsim import KEYWORDS
-
 # Location of C and OpenCL sources
 SOURCE_FILE = 'fiber_tissue.c'
 KERNEL_FILE = 'openclsim.cl'
@@ -327,16 +324,7 @@ def __init__(
             'diffusion_current': 'idiff',
         })
 
-        # Reserve keywords
-        from myokit.formats import opencl
-        self._modelf.reserve_unique_names(*opencl.keywords)
-        self._modelt.reserve_unique_names(*opencl.keywords)
-        self._modelf.reserve_unique_names(
-            *['calc_' + c.name() for c in self._modelf.components()])
-        self._modelt.reserve_unique_names(
-            *['calc_' + c.name() for c in self._modelt.components()])
-        self._modelf.reserve_unique_names(*KEYWORDS)
-        self._modelt.reserve_unique_names(*KEYWORDS)
+        # Create unique names
         self._modelf.create_unique_names()
         self._modelt.create_unique_names()
 
diff --git a/myokit/_sim/openclsim.c b/myokit/_sim/openclsim.c
index 8e5f66bda..f096c7943 100644
--- a/myokit/_sim/openclsim.c
+++ b/myokit/_sim/openclsim.c
@@ -56,7 +56,7 @@ if precision == myokit.DOUBLE_PRECISION:
  *  var      : The variable to add to the logs, if its name is present
  * Returns 0 if not added, 1 if added.
  */
-static int log_add(PyObject* log_dict, PyObject** logs, Real** vars, int i, char* name, const Real* var)
+static int log_add(PyObject* log_dict, PyObject** logs, Real** vars, unsigned long i, char* name, const Real* var)
 {
     int added = 0;
     PyObject* key = PyUnicode_FromString(name);
@@ -76,12 +76,16 @@ static int log_add(PyObject* log_dict, PyObject** logs, Real** vars, int i, char
 // Simulation state
 int running = 0;    // 1 if a simulation has been initialized, 0 if it's clean
 
+// Note: OpenCL functions such as get_global_id() all return a size_t defined
+// on the device (so different from the .C size_t). This is either a 32 or a
+// 64 bit int. To be safe, we just use a ulong everywhere.
+
 // Input arguments
 PyObject *platform_name;// A python string specifying the platform to use
 PyObject *device_name;  // A python string specifying the device to use
 char* kernel_source;    // The kernel code
-int nx;                 // The number of cells in the x direction
-int ny;                 // The number of cells in the y direction
+unsigned long nx;       // The number of cells in the x direction
+unsigned long ny;       // The number of cells in the y direction
 double gx;              // The cell-to-cell conductance in the x direction
 double gy;              // The cell-to-cell conductance in the y direction
 double tmin;            // The initial simulation time
@@ -116,8 +120,8 @@ Real *rvec_state = NULL;
 Real *rvec_idiff = NULL;
 Real *rvec_inter_log = NULL;
 Real *rvec_field_data = NULL;
-int *rvec_conn1 = NULL;
-int *rvec_conn2 = NULL;
+unsigned long *rvec_conn1 = NULL;
+unsigned long *rvec_conn2 = NULL;
 Real *rvec_conn3 = NULL;
 size_t dsize_state;
 size_t dsize_idiff;
@@ -147,7 +151,7 @@ int diffusion;
 
 // Arbitrary geometry diffusion
 PyObject* connections;  // List of connection tuples
-int n_connections;
+unsigned long n_connections;
 
 // OpenCL work group sizes
 size_t global_work_size[2];
@@ -162,20 +166,20 @@ Real arg_gx;
 Real arg_gy;
 
 /* Logging */
-PyObject** logs = NULL; /* An array of pointers to a PyObject */
-Real** vars = NULL;     /* An array of pointers to values to log */
-int n_vars;             /* Number of logging variables */
-double tnext_log;       /* The next logging point */
-unsigned long inext_log;/* The number of logged steps */
-int logging_diffusion;  /* True if diffusion current is being logged. */
-int logging_states;     /* True if any states are being logged */
-int logging_inters;     /* True if any intermediary variables are being logged. */
-int n_inter;            /* The number of intermediary variables to log */
+PyObject** logs = NULL;     /* An array of pointers to a PyObject */
+Real** vars = NULL;         /* An array of pointers to values to log */
+unsigned long n_vars;       /* Number of logging variables */
+double tnext_log;           /* The next logging point */
+unsigned long inext_log;    /* The number of logged steps */
+int logging_diffusion;      /* True if diffusion current is being logged. */
+int logging_states;         /* True if any states are being logged */
+int logging_inters;         /* True if any intermediary variables are being logged. */
+unsigned long n_inter;      /* The number of intermediary variables to log */
 /* The relationship between n_inter and logging_inters isn't straightforward: */
 /* n_inter is the number of different model variables (so membrane.V, not */
 /* 1.2.membrane.V) being logged, while logging_inters is 1 if at least one */
 /* simulation variable (1.2.membrane.V) is listed in the given log. */
-int n_field_data;       /* The number of floats in the field data */
+unsigned long n_field_data;       /* The number of floats in the field data */
 
 /* Temporary objects: decref before re-using for another var */
 /* (Unless you got it through PyList_GetItem or PyTuble_GetItem) */
@@ -260,7 +264,7 @@ sim_clean()
     return 0;
 }
 static PyObject*
-py_sim_clean()
+py_sim_clean(PyObject *self, PyObject *args)
 {
     #ifdef MYOKIT_DEBUG
     printf("Python py_sim_clean called.\n");
@@ -285,7 +289,7 @@ sim_init(PyObject* self, PyObject* args)
     cl_int flag;
 
     // Iteration
-    int i, j, k;
+    unsigned long i, j, k;
 
     // Platform and device id
     cl_platform_id platform_id;
@@ -296,7 +300,7 @@ sim_init(PyObject* self, PyObject* args)
 
     // Variable names
     char log_var_name[1023];
-    int k_vars;
+    unsigned long k_vars;
 
     // Compilation error message
     size_t blog_size;
@@ -342,11 +346,12 @@ sim_init(PyObject* self, PyObject* args)
     list_update_str = NULL;
 
     // Check input arguments
-    if(!PyArg_ParseTuple(args, "OOsiibddOdddOOOOdOO",
+    // https://docs.python.org/3.8/c-api/arg.html#c.PyArg_ParseTuple
+    if(!PyArg_ParseTuple(args, "OOskkbddOdddOOOOdOO",
             &platform_name,     // Must be bytes
             &device_name,       // Must be bytes
             &kernel_source,
-            &nx,
+            &nx,                // Small 'k' = unsigned long
             &ny,
             &diffusion,
             &gx,
@@ -368,7 +373,7 @@ sim_init(PyObject* self, PyObject* args)
         return 0;
     }
     dt = default_dt;
-    dt_min = 0; /*dt * 1e-2; Not sure why this was >0. Causes issues with timing sometimes. */
+    dt_min = 0;
     arg_dt = (Real)dt;
     arg_gx = (Real)gx;
     arg_gy = (Real)gy;
@@ -399,7 +404,7 @@ sim_init(PyObject* self, PyObject* args)
         PyErr_SetString(PyExc_Exception, "'state_in' must be a list.");
         return sim_clean();
     }
-    if(PyList_Size(state_in) != nx * ny * n_state) {
+    if((unsigned long)PyList_Size(state_in) != nx * ny * n_state) {
         PyErr_SetString(PyExc_Exception, "'state_in' must have size nx * ny * n_states.");
         return sim_clean();
     }
@@ -407,7 +412,7 @@ sim_init(PyObject* self, PyObject* args)
         PyErr_SetString(PyExc_Exception, "'state_out' must be a list.");
         return sim_clean();
     }
-    if(PyList_Size(state_out) != nx * ny * n_state) {
+    if((unsigned long)PyList_Size(state_out) != nx * ny * n_state) {
         PyErr_SetString(PyExc_Exception, "'state_out' must have size nx * ny * n_states.");
         return sim_clean();
     }
@@ -464,7 +469,7 @@ sim_init(PyObject* self, PyObject* args)
         flt = PyList_GetItem(state_in, i);    // Don't decref!
         if(!PyFloat_Check(flt)) {
             char errstr[200];
-            sprintf(errstr, "Item %d in state vector is not a float.", i);
+            sprintf(errstr, "Item %u in state vector is not a float.", (unsigned int)i);
             PyErr_SetString(PyExc_Exception, errstr);
             return sim_clean();
         }
@@ -499,7 +504,7 @@ sim_init(PyObject* self, PyObject* args)
             flt = PyList_GetItem(field_data, i);    // No need to decref
             if(!PyFloat_Check(flt)) {
                 char errstr[200];
-                sprintf(errstr, "Item %d in field data is not a float.", i);
+                sprintf(errstr, "Item %u in field data is not a float.", (unsigned int)i);
                 PyErr_SetString(PyExc_Exception, errstr);
                 return sim_clean();
             }
@@ -517,11 +522,11 @@ sim_init(PyObject* self, PyObject* args)
             return sim_clean();
         }
         n_connections = PyList_Size(connections);
-        dsize_conn1 = n_connections * sizeof(int);  // Use int; same as nx, ny, etc.
-        dsize_conn2 = n_connections * sizeof(int);
+        dsize_conn1 = n_connections * sizeof(unsigned long);  // Same type as nx, ny, etc.
+        dsize_conn2 = n_connections * sizeof(unsigned long);
         dsize_conn3 = n_connections * sizeof(Real);
-        rvec_conn1 = (int*)malloc(dsize_conn1);
-        rvec_conn2 = (int*)malloc(dsize_conn2);
+        rvec_conn1 = (unsigned long*)malloc(dsize_conn1);
+        rvec_conn2 = (unsigned long*)malloc(dsize_conn2);
         rvec_conn3 = (Real*)malloc(dsize_conn3);
         for(i=0; i<n_connections; i++) {
             flt = PyList_GetItem(connections, i);   // Borrowed reference
@@ -536,11 +541,11 @@ sim_init(PyObject* self, PyObject* args)
 
             ret = PyTuple_GetItem(flt, 0);  // Borrowed reference
             if(PyLong_Check(ret)) {
-                rvec_conn1[i] = (int)PyLong_AsLong(ret);
-#if PY_MAJOR_VERSION < 3
+                rvec_conn1[i] = (unsigned long)PyLong_AsLong(ret);
+            #if PY_MAJOR_VERSION < 3
             } else if (PyInt_Check(ret)) {
-                rvec_conn1[i] = (int)PyInt_AsLong(ret);
-#endif
+                rvec_conn1[i] = (unsigned long)PyInt_AsLong(ret);
+            #endif
             } else {
                 PyErr_SetString(PyExc_Exception, "First item in each connection tuple must be int");
                 return sim_clean();
@@ -548,11 +553,11 @@ sim_init(PyObject* self, PyObject* args)
 
             ret = PyTuple_GetItem(flt, 1);  // Borrowed reference
             if(PyLong_Check(ret)) {
-                rvec_conn2[i] = (int)PyLong_AsLong(ret);
-#if PY_MAJOR_VERSION < 3
+                rvec_conn2[i] = (unsigned long)PyLong_AsLong(ret);
+            #if PY_MAJOR_VERSION < 3
             } else if(PyInt_Check(ret)) {
-                rvec_conn2[i] = (int)PyInt_AsLong(ret);
-#endif
+                rvec_conn2[i] = (unsigned long)PyInt_AsLong(ret);
+            #endif
             } else {
                 PyErr_SetString(PyExc_Exception, "Second item in each connection tuple must be int");
                 return sim_clean();
@@ -643,13 +648,13 @@ sim_init(PyObject* self, PyObject* args)
 
     #ifdef MYOKIT_DEBUG
     printf("Created buffers.\n");
-    printf("State buffer size:%d.\n", dsize_state);
-    printf("Idiff buffer size:%d.\n", dsize_idiff);
-    printf("Inter-log buffer size:%d.\n", dsize_inter_log);
-    printf("Field-data buffer size:%d.\n", dsize_field_data);
-    printf("Connections-1 buffer size:%d.\n", dsize_conn1);
-    printf("Connections-2 buffer size:%d.\n", dsize_conn2);
-    printf("Connections-3 buffer size:%d.\n", dsize_conn3);
+    printf("State buffer size: %d.\n", (int)dsize_state);
+    printf("Idiff buffer size: %d.\n", (int)dsize_idiff);
+    printf("Inter-log buffer size: %d.\n", (int)dsize_inter_log);
+    printf("Field-data buffer size: %d.\n", (int)dsize_field_data);
+    printf("Connections-1 buffer size: %d.\n", (int)dsize_conn1);
+    printf("Connections-2 buffer size: %d.\n", (int)dsize_conn2);
+    printf("Connections-3 buffer size: %d.\n", (int)dsize_conn3);
     #endif
 
     /* Copy data into buffers */
@@ -780,7 +785,7 @@ sim_init(PyObject* self, PyObject* args)
     }
     n_vars = PyDict_Size(log_dict);
     #ifdef MYOKIT_DEBUG
-    printf("Number of variables to log:%d.\n", n_vars);
+    printf("Number of variables to log:%u.\n", (unsigned int)n_vars);
     #endif
     logs = (PyObject**)malloc(sizeof(PyObject*)*n_vars); // Pointers to logging lists
     #ifdef MYOKIT_DEBUG
@@ -811,9 +816,9 @@ if var is not None:
 var = model.binding('diffusion_current')
 if var is not None:
     if dims == 1:
-        print(3*tab + 'sprintf(log_var_name, "%d.' + var.qname() + '", j);')
+        print(3*tab + 'sprintf(log_var_name, "%u.' + var.qname() + '", (unsigned int)j);')
     else:
-        print(3*tab + 'sprintf(log_var_name, "%d.%d.' + var.qname() + '", j, i);')
+        print(3*tab + 'sprintf(log_var_name, "%u.%u.' + var.qname() + '", (unsigned int)j, (unsigned int)i);')
     print(3*tab + 'if(log_add(log_dict, logs, vars, k_vars, log_var_name, &rvec_idiff[i*nx+j])) {')
     print(4*tab + 'logging_diffusion = 1;')
     print(4*tab + 'k_vars++;')
@@ -829,9 +834,9 @@ if var is not None:
 <?
 for var in model.states():
     if dims == 1:
-        print(3*tab + 'sprintf(log_var_name, "%d.' + var.qname() + '", j);')
+        print(3*tab + 'sprintf(log_var_name, "%u.' + var.qname() + '", (unsigned int)j);')
     else:
-        print(3*tab + 'sprintf(log_var_name, "%d.%d.' + var.qname() + '", j, i);' )
+        print(3*tab + 'sprintf(log_var_name, "%u.%u.' + var.qname() + '", (unsigned int)j, (unsigned int)i);' )
     print(3*tab + 'if(log_add(log_dict, logs, vars, k_vars, log_var_name, &rvec_state[(i*nx+j)*n_state+' + str(var.indice()) + '])) {')
     print(4*tab + 'logging_states = 1;')
     print(4*tab + 'k_vars++;')
@@ -848,9 +853,9 @@ for var in model.states():
                 ret = PyList_GetItem(inter_log, k); // Don't decref
 <?
 if dims == 1:
-    print(4*tab + 'sprintf(log_var_name, "%d.%s", j, PyBytes_AsString(ret));')
+    print(4*tab + 'sprintf(log_var_name, "%u.%s", (unsigned int)j, PyBytes_AsString(ret));')
 else:
-    print(4*tab + 'sprintf(log_var_name, "%d.%d.%s", j, i, PyBytes_AsString(ret));')
+    print(4*tab + 'sprintf(log_var_name, "%u.%u.%s", (unsigned int)j, (unsigned int)i, PyBytes_AsString(ret));')
 
 print(4*tab + 'if(log_add(log_dict, logs, vars, k_vars, log_var_name, &rvec_inter_log[(i*nx+j)*n_inter+k])) {')
 print(5*tab + 'logging_inters = 1;')
@@ -869,7 +874,7 @@ print(4*tab + '}')
     }
 
     #ifdef MYOKIT_DEBUG
-    printf("Created log for %d variables.\n", n_vars);
+    printf("Created log for %u variables.\n", (unsigned int)n_vars);
     #endif
 
     /* Log update method: */
@@ -900,7 +905,7 @@ sim_step(PyObject *self, PyObject *args)
     ESys_Flag flag_pacing;
     long steps_left_in_run;
     cl_int flag;
-    int i;
+    unsigned long i;
     double d;
     int logging_condition;
 
@@ -1075,8 +1080,8 @@ sim_step(PyObject *self, PyObject *args)
  */
 static PyMethodDef SimMethods[] = {
     {"sim_init", sim_init, METH_VARARGS, "Initialize the simulation."},
-    {"sim_step", sim_step, METH_VARARGS, "Perform the next step in the simulation."},
-    {"sim_clean", py_sim_clean, METH_VARARGS, "Clean up after an aborted simulation."},
+    {"sim_step", sim_step, METH_NOARGS, "Perform the next step in the simulation."},
+    {"sim_clean", py_sim_clean, METH_NOARGS, "Clean up after an aborted simulation."},
     {NULL},
 };
 
diff --git a/myokit/_sim/openclsim.cl b/myokit/_sim/openclsim.cl
index b7702710a..235312b97 100644
--- a/myokit/_sim/openclsim.cl
+++ b/myokit/_sim/openclsim.cl
@@ -148,10 +148,11 @@ if precision == myokit.DOUBLE_PRECISION:
 /* Number of scalar fields */
 #define n_field <?=str(len(fields))?>
 
-/* Indice of membrane potential in state vector */
-#define i_vm <?= model.label('membrane_potential').indice() ?>
-
 <?
+if diffusion:
+    print('/* Indice of membrane potential in state vector */')
+    print('#define i_vm ' + str(model.label('membrane_potential').indice()))
+
 if precision == myokit.SINGLE_PRECISION:
     print('/* Using single precision floats */')
     print('typedef float Real;')
@@ -220,9 +221,9 @@ for comp, ilist in comp_in.items():
     # Function header
     olist = comp_out[comp]
     args = [
-        'const uint of1',
-        'const uint of2',
-        'const uint of3',
+        'const unsigned long of1',
+        'const unsigned long of2',
+        'const unsigned long of3',
         '__global Real *state',
         '__global Real *inter_log',
         'const __global Real *field_data',
@@ -254,9 +255,9 @@ if diffusion and paced_cells:
     print(' * Calculate pacing current per cell')
     print(' */')
     print('inline Real calculate_pacing(')
-    print(tab + 'const uint cid,')
-    print(tab + 'const uint ix,')
-    print(tab + 'const uint iy,')
+    print(tab + 'const unsigned long cid,')
+    print(tab + 'const unsigned long ix,')
+    print(tab + 'const unsigned long iy,')
     print(tab + 'const Real pace)')
     print('{')
 
@@ -265,18 +266,16 @@ if diffusion and paced_cells:
         nx, ny, x, y = paced_cells
         xlo, ylo = str(x), str(y)
         xhi, yhi = str(x + nx), str(y + ny)
-        # Code below to be added in after unit tests are up and running
-        #cond = []
-        #if x > 0:
-        #   cond.append('ix >= ' + str(x))
-        #cond.append('ix < ' + str(x + nx))
-        #if y > 1:
-        #    bits.append('iy >= ' + str(y))
-        #cond.append('iy < ' + str(y + ny))
-        #cond = ' && '.join(cond)
-        #print(tab + 'return (' + cond + ') ? pace : 0;')
-        print(tab + 'return (ix >= ' + xlo + ' && ix < ' + xhi + ' && iy >= '
-            + ylo + ' && iy < ' + yhi + ') ? pace : 0;')
+
+        cond = []
+        if x > 0:
+           cond.append('ix >= ' + str(x))
+        if y > 0:
+            cond.append('iy >= ' + str(y))
+        cond.append('ix < ' + str(x + nx))
+        cond.append('iy < ' + str(y + ny))
+        cond = ' && '.join(cond)
+        print(tab + 'return (' + cond + ') ? pace : 0;')
     else:
         # Explicit cell selection
         for id in paced_cells:
@@ -301,8 +300,8 @@ if diffusion and paced_cells:
  *  field_data : A vector containing all field data
  */
 __kernel void cell_step(
-    const uint nx,
-    const uint ny,
+    const unsigned long nx,
+    const unsigned long ny,
     const Real time,
     const Real dt,
     const Real pace_in,
@@ -312,16 +311,16 @@ __kernel void cell_step(
     const __global Real* field_data
     )
 {
-    const uint ix = get_global_id(0);
-    const uint iy = get_global_id(1);
+    const unsigned long ix = get_global_id(0);
+    const unsigned long iy = get_global_id(1);
     if(ix >= nx) return;
     if(iy >= ny) return;
 
     // Offset of this cell's state in the state vector
-    const uint cid = ix + iy * nx;
-    const uint of1 = cid * n_state;
-    const uint of2 = cid * n_inter;
-    const uint of3 = cid * n_field;
+    const unsigned long cid = ix + iy * nx;
+    const unsigned long of1 = cid * n_state;
+    const unsigned long of2 = cid * n_inter;
+    const unsigned long of3 = cid * n_field;
 
     // Pacing
 <?
@@ -371,6 +370,9 @@ for var in model.states():
 ?>
 }
 
+<?
+if diffusion and not connections:
+    print("""
 /*
  * Performs a single diffusion step
  *
@@ -383,24 +385,24 @@ for var in model.states():
  *  idiff : The diffusion current vector
  */
 __kernel void diff_step(
-    const uint nx,
-    const uint ny,
+    const unsigned long nx,
+    const unsigned long ny,
     const Real gx,
     const Real gy,
     __global Real *state,
     __global Real *idiff)
 {
-    const uint ix = get_global_id(0);
-    const uint iy = get_global_id(1);
+    const unsigned long ix = get_global_id(0);
+    const unsigned long iy = get_global_id(1);
     if(ix >= nx) return;
-    if (iy >= ny) return;
+    if(iy >= ny) return;
 
     // Offset of this cell's Vm in the state vector
-    const uint cid = ix + iy * nx;
-    const uint of1 = cid * n_state + i_vm;
+    const unsigned long cid = ix + iy * nx;
+    const unsigned long of1 = cid * n_state + i_vm;
 
     // Diffusion, x-direction
-    uint ofp, ofm;
+    unsigned long ofp, ofm;
     if(nx > 1) {
         ofp = of1 + n_state;
         ofm = of1 - n_state;
@@ -434,8 +436,8 @@ __kernel void diff_step(
         }
     }
 }
+    """)
 
-<?
 if connections:
     print("""
 /*
@@ -484,19 +486,19 @@ inline void AtomicAdd(volatile __global Real *var, const Real operand)
  *  idiff : The diffusion current vector
  */
 __kernel void diff_arb_step(
-    const uint count,
-    __global int *cell1,
-    __global int *cell2,
+    const unsigned long count,
+    __global unsigned long *cell1,
+    __global unsigned long *cell2,
     __global Real *conductance,
     __global Real *state,
     __global Real *idiff)
 {
-    const uint ix = get_global_id(0);
+    const unsigned long ix = get_global_id(0);
     if(ix >= count) return;
 
     // Cell indices and conductance
-    int i1 = cell1[ix];
-    int i2 = cell2[ix];
+    unsigned long i1 = cell1[ix];
+    unsigned long i2 = cell2[ix];
     Real i12 = conductance[ix] * (state[i1 * n_state + i_vm] - state[i2 * n_state + i_vm]);
 
     // Diffusion
@@ -513,10 +515,10 @@ __kernel void diff_arb_step(
  *  idiff : The diffusion current vector
  */
 __kernel void diff_arb_reset(
-    const uint count,
+    const unsigned long count,
     __global Real *idiff)
 {
-    const uint ix = get_global_id(0);
+    const unsigned long ix = get_global_id(0);
     if(ix >= count) return;
     idiff[ix] = 0;
 }
@@ -547,26 +549,26 @@ __kernel void diff_arb_reset(
  *  idiff_t : The diffusion current vector for the tissue model
  */
 __kernel void diff_step_fiber_tissue(
-    const uint nfx,
-    const uint nfy,
-    const uint ntx,
-    const uint ctx,
-    const uint cty,
-    const uint nsf,
-    const uint nst,
-    const uint ivf,
-    const uint ivt,
+    const unsigned long nfx,
+    const unsigned long nfy,
+    const unsigned long ntx,
+    const unsigned long ctx,
+    const unsigned long cty,
+    const unsigned long nsf,
+    const unsigned long nst,
+    const int ivf,
+    const int ivt,
     const Real gft,
     __global Real *state_f,
     __global Real *state_t,
     __global Real *idiff_f,
     __global Real *idiff_t)
 {
-    const uint cid = get_global_id(0);
-    const uint iff = (nfx - 1) + cid * nfx;
-    const uint ift = ctx + (cty + cid) * ntx;
-    const uint off = iff * nsf + ivf;
-    const uint oft = ift * nst + ivt;
+    const unsigned long cid = get_global_id(0);
+    const unsigned long iff = (nfx - 1) + cid * nfx;
+    const unsigned long ift = ctx + (cty + cid) * ntx;
+    const unsigned long off = iff * nsf + ivf;
+    const unsigned long oft = ift * nst + ivt;
     if (cid < nfy) {
         // Connection
         Real idiff = gft * (state_f[off] - state_t[oft]);
diff --git a/myokit/_sim/openclsim.py b/myokit/_sim/openclsim.py
index b25c3904c..d7d455067 100644
--- a/myokit/_sim/openclsim.py
+++ b/myokit/_sim/openclsim.py
@@ -202,14 +202,15 @@ def __init__(
 
         # Get membrane potential variable (from pre-cloned model!)
         vm = model.label('membrane_potential')
-        if vm is None:
-            raise ValueError(
-                'This simulation requires the membrane potential'
-                ' variable to be labelled as "membrane_potential".')
-        if not vm.is_state():
-            raise ValueError(
-                'The variable labelled as membrane potential must'
-                ' be a state variable.')
+        if self._diffusion_enabled or self._rl:
+            if vm is None:
+                raise ValueError(
+                    'This simulation requires the membrane potential'
+                    ' variable to be labelled as "membrane_potential".')
+            if not vm.is_state():
+                raise ValueError(
+                    'The variable labelled as membrane potential must'
+                    ' be a state variable.')
 
         #if vm.is_referenced():
         #  raise ValueError('This simulation requires that no other variables'
@@ -234,7 +235,10 @@ def __init__(
         else:
             # Clone model, store
             self._model = model.clone()
-            self._vm = self._model.get(vm.qname())
+            if vm is None:
+                self._vm = None
+            else:
+                self._vm = self._model.get(vm.qname())
             del(model, vm)
 
         # Set default conductance values
@@ -276,14 +280,7 @@ def __init__(
             inputs['diffusion_current'] = 'idiff'
         self._bound_variables = self._model.prepare_bindings(inputs)
 
-        # Reserve keywords
-        from myokit.formats import opencl
-        self._model.reserve_unique_names(*opencl.keywords)
-        self._model.reserve_unique_names(
-            *['calc_' + c.name() for c in self._model.components()])
-        self._model.reserve_unique_names(
-            *['D_' + c.uname() for c in self._model.states()])
-        self._model.reserve_unique_names(*KEYWORDS)
+        # Create unique names
         self._model.create_unique_names()
 
         # Create back-end
@@ -1139,12 +1136,12 @@ def _run(self, duration, log, log_interval, report_nan, progress, msg):
                 txt.append('Simulation variables during:')
                 txt.append('  Time: ' + myokit.float.str(
                     bounds[0][vtime], precision=self._precision))
-                if vpace:
+                if vpace is not None:
                     txt.append(
                         '  Pacing variable: ' + myokit.float.str(
                             bounds[0][vpace], precision=self._precision)
                         + is_paced_str)
-                if vdiff:
+                if vdiff is not None:
                     txt.append(
                         '  Diffusion current: ' + myokit.float.str(
                             bounds[0][vdiff], precision=self._precision))
@@ -1186,12 +1183,12 @@ def _run(self, duration, log, log_interval, report_nan, progress, msg):
                     txt.append('Simulation variables before:')
                     txt.append('  Time: ' + myokit.float.str(
                         bounds[1][vtime], precision=self._precision))
-                    if vpace:
+                    if vpace is not None:
                         txt.append(
                             '  Pacing variable: ' + myokit.float.str(
                                 bounds[1][vpace], precision=self._precision)
                             + is_paced_str)
-                    if vdiff:
+                    if vdiff is not None:
                         txt.append(
                             '  Diffusion current: ' + myokit.float.str(
                                 bounds[1][vdiff], precision=self._precision))
@@ -1645,68 +1642,3 @@ def time(self):
         """
         return self._time
 
-
-KEYWORDS = [
-    'AtomicAdd',
-    'calculate_pacing',
-    'cell1',
-    'cell2',
-    'cell_step',
-    'cid',
-    'conductance',
-    'ctx',
-    'cty',
-    'diff_step',
-    'diff_arb_step',
-    'diff_arb_reset',
-    'diff_step_fiber_tissue',
-    'dt',
-    'floatVal',
-    'gft',
-    'gx',
-    'gy',
-    'i1',
-    'i2',
-    'i12',
-    'idiff',
-    'idiff_f',
-    'idiff_in',
-    'idiff_t',
-    'iff',
-    'ift',
-    'inter_log',
-    'intVal',
-    'ivf',
-    'ivt',
-    'ix',
-    'iy',
-    'i_vm',
-    'newVal',
-    'nfx',
-    'nfy',
-    'nsf',
-    'nst',
-    'ntx',
-    'nx',
-    'nx_paced',
-    'ny',
-    'ny_paced',
-    'n_inter',
-    'n_state',
-    'of1',
-    'of2',
-    'off',
-    'ofm',
-    'ofp',
-    'oft',
-    'operand',
-    'pace',
-    'pace_in',
-    'prevVal',
-    'Real',
-    'source',
-    'state',
-    'state_f',
-    'state_t',
-    'time',
-]


Test patch:
diff --git a/myokit/tests/test_simulation_fiber_tissue.py b/myokit/tests/test_simulation_fiber_tissue.py
old mode 100644
new mode 100755
index 3cfd35204..54948cc96
--- a/myokit/tests/test_simulation_fiber_tissue.py
+++ b/myokit/tests/test_simulation_fiber_tissue.py
@@ -14,7 +14,7 @@
 
 import myokit
 
-from shared import OpenCL_FOUND, DIR_DATA
+from shared import OpenCL_FOUND, DIR_DATA, WarningCollector
 
 # Unit testing in Python 2 and 3
 try:
@@ -52,17 +52,20 @@ def test_basic(self):
         nty = 6
 
         # Create simulation
-        s = myokit.FiberTissueSimulation(
-            mf,
-            mt,
-            p,
-            ncells_fiber=(nfx, nfy),
-            ncells_tissue=(ntx, nty),
-            nx_paced=10,
-            g_fiber=(235, 100),
-            g_tissue=(9, 5),
-            g_fiber_tissue=9
-        )
+        with WarningCollector() as w:
+            s = myokit.FiberTissueSimulation(
+                mf,
+                mt,
+                p,
+                ncells_fiber=(nfx, nfy),
+                ncells_tissue=(ntx, nty),
+                nx_paced=10,
+                g_fiber=(235, 100),
+                g_tissue=(9, 5),
+                g_fiber_tissue=9
+            )
+            self.assertIn('deprecated', w.text())
+
         s.set_step_size(0.0012)
         # Set up logging
         logf = [
@@ -104,18 +107,19 @@ def test_against_cvode(self):
         p = myokit.pacing.blocktrain(1000, 2.0, offset=0)
 
         # Create simulation
-        s1 = myokit.FiberTissueSimulation(
-            m,
-            m,
-            p,
-            ncells_fiber=(1, 1),
-            ncells_tissue=(1, 1),
-            nx_paced=1,
-            g_fiber=(0, 0),
-            g_tissue=(0, 0),
-            g_fiber_tissue=0,
-            precision=myokit.DOUBLE_PRECISION,
-        )
+        with WarningCollector():
+            s1 = myokit.FiberTissueSimulation(
+                m,
+                m,
+                p,
+                ncells_fiber=(1, 1),
+                ncells_tissue=(1, 1),
+                nx_paced=1,
+                g_fiber=(0, 0),
+                g_tissue=(0, 0),
+                g_fiber_tissue=0,
+                precision=myokit.DOUBLE_PRECISION,
+            )
         s1.set_step_size(0.01)
 
         # Set up logging
@@ -200,3 +204,17 @@ def test_against_cvode(self):
         self.assertLess(e2, 0.05)
         self.assertLess(e3, 0.01)
 
+
+if __name__ == '__main__':
+
+    import sys
+    if '-v' in sys.argv:
+        print('Running in debug/verbose mode')
+        debug = True
+    else:
+        print('Add -v for more debug output')
+
+    import warnings
+    warnings.simplefilter('always')
+
+    unittest.main()
diff --git a/myokit/tests/test_simulation_opencl.py b/myokit/tests/test_simulation_opencl.py
index 2973660ca..ff3d3131a 100755
--- a/myokit/tests/test_simulation_opencl.py
+++ b/myokit/tests/test_simulation_opencl.py
@@ -478,6 +478,19 @@ def test_diffusion_free(self):
         self.assertFalse(np.all(d['membrane.V', 0] == d['membrane.V', 2]))
         self.assertFalse(np.all(d['membrane.V', 1] == d['membrane.V', 2]))
 
+        # Create without diffusion, 1d, no membrane_potential label
+        m = self.m.clone()
+        m.label('membrane_potential').set_label(None)
+        s = myokit.SimulationOpenCL(m, self.p, ncells=3, diffusion=False)
+
+        # Test that all cells depolarise and are the same
+        d = s.run(10, log=['membrane.V']).npview()
+        self.assertGreater(np.max(d['membrane.V', 0]), 0)
+        self.assertGreater(np.max(d['membrane.V', 1]), 0)
+        self.assertGreater(np.max(d['membrane.V', 2]), 0)
+        self.assertTrue(np.all(d['membrane.V', 0] == d['membrane.V', 1]))
+        self.assertTrue(np.all(d['membrane.V', 0] == d['membrane.V', 2]))
+
     def test_neighbours_0d(self):
         # Test listing neighbours in a 0d simulation
 
diff --git a/myokit/tests/test_simulation_opencl_log_interval.py b/myokit/tests/test_simulation_opencl_log_interval.py
index d6c4d7f79..4cb7db9aa 100755
--- a/myokit/tests/test_simulation_opencl_log_interval.py
+++ b/myokit/tests/test_simulation_opencl_log_interval.py
@@ -14,7 +14,7 @@
 
 import myokit
 
-from shared import OpenCL_FOUND, DIR_DATA
+from shared import OpenCL_FOUND, DIR_DATA, WarningCollector
 from test_simulation_log_interval import PeriodicTest
 
 debug = False
@@ -41,8 +41,9 @@ def test_periodic(self):
         Test periodic logging.
         """
         m, p, x = myokit.load(os.path.join(DIR_DATA, 'lr-1991.mmt'))
-        s = myokit.FiberTissueSimulation(
-            m, m, p, ncells_fiber=(1, 1), ncells_tissue=(1, 1))
+        with WarningCollector():
+            s = myokit.FiberTissueSimulation(
+                m, m, p, ncells_fiber=(1, 1), ncells_tissue=(1, 1))
 
         # Set tolerance for equality testing
         emax = 1e-2  # Time steps for logging are approximate