First try to implement synapse to synapse support

Problem is, that we need access to target variables, which gives an
error when the target is a `Synapse`, see #133

brian2cuda/templates/objects.cu

@@ -110,18 +110,18 @@ int brian::num_threads_per_warp;
 {% for S in synapses | sort(attribute='name') %}
 {% for path in S._pathways | sort(attribute='name') %}
 __global__ void {{path.name}}_init(
-                int Nsource,
-                int Ntarget,
                 int32_t* sources,
                 int32_t* targets,
                 double dt,
-                int32_t start,
-                int32_t stop
+                int32_t source_start,
+                int32_t source_stop,
+                int32_t target_start,
+                int32_t target_stop
         )
 {
     using namespace brian;
 
-    {{path.name}}.init(Nsource, Ntarget, sources, targets, dt, start, stop);
+    {{path.name}}.init(sources, targets, dt, start, stop);
 }
 {% endfor %}
 {% endfor %}
@@ -188,7 +188,9 @@ void _init_arrays()
             thrust::raw_pointer_cast(&dev{{dynamic_array_specs[path.synapse_targets]}}[0]),
             0,  //was dt, maybe irrelevant?
             {{path.source.start}},
-            {{path.source.stop}}
+            {{path.source.stop}},
+            {{path.target.start}},
+            {{path.target.stop}}
             );
     CUDA_CHECK_ERROR("{{path.name}}_init");
     {% endfor %}

brian2cuda/templates/synapses_classes.cu

@@ -14,33 +14,40 @@
 class SynapticPathway
 {
 public:
-    // total number of neurons in source and target NeuronGroup / Subgroup
-    int Nsource;
-    int Ntarget;
-
     int32_t* dev_sources;
     int32_t* dev_targets;
 
     // first and last index in source NeuronGroup corresponding to Subgroup in SynapticPathway
     // important for Subgroups created with syntax: NeuronGroup(N=4000,...)[:3200]
-    int spikes_start;
-    int spikes_stop;
+    int32_t spikes_start;
+    int32_t spikes_stop;
+
+    // first and last index in target Group corresponding to Subgroup in SynapticPathway
+    // important for Subgroups created with syntax: NeuronGroup(N=4000,...)[:3200]
+    // We need these to be initialised in `objects.cu` for synapses to synapses to work,
+    // since our connectivity matrix (`synapses_initialise_queue` tempalte)
+    // needs access to the target group IDs, which we can't access at python
+    // runtime (before the network is run)
+    int32_t targets_start;
+    int32_t targets_stop;
+
 
     double dt;
     CudaSpikeQueue* queue;
     bool no_or_const_delay_mode;
 
     //our real constructor
-    __device__ void init(int _Nsource, int _Ntarget, int32_t* _sources,
-                int32_t* _targets, double _dt, int _spikes_start, int _spikes_stop)
+    __device__ void init(int32_t* _sources, int32_t* _targets, double _dt,
+            int32_t _spikes_start, int32_t _spikes_stop, int32_t _targets_start,
+            int32_t _targets_stop)
     {
-        Nsource = _Nsource;
-        Ntarget = _Ntarget;
         dev_sources = _sources;
         dev_targets = _targets;
         dt = _dt;
         spikes_start = _spikes_start;
         spikes_stop = _spikes_stop;
+        targets_start = _targets_start;
+        targets_stop = _targets_stop;
         queue = new CudaSpikeQueue;
     };
 

brian2cuda/templates/synapses_initialise_queue.cu

@@ -29,7 +29,7 @@
  #}
 
 {% macro cu_file() %}
-{# USES_VARIABLES { delay } #}
+{# USES_VARIABLES { delay, _n_sources, _n_targets } #}
 #include <thrust/sort.h>
 #include <thrust/reduce.h>
 #include <iostream>
@@ -167,9 +167,9 @@ void _run_{{pathobj}}_initialise_queue()
     // simulation time step
     double dt = {{owner.clock.name}}.dt[0];
     // number of neurons in source group
-    int source_N = {{owner.source.N}};
+    int source_N = {{constant_or_scalar('_n_sources', variables['_n_sources'])}};
     // number of neurons in target group
-    int target_N = {{owner.target.N}};
+    int target_N = {{constant_or_scalar('_n_targets', variables['_n_targets'])}};
 
     //TODO: for multiple SynapticPathways for the same Synapses object (on_pre and on_post) the following copy is identical in both pathways initialise templates
     {% if not no_or_const_delay_mode %}
@@ -287,10 +287,10 @@ void _run_{{pathobj}}_initialise_queue()
     {% endif %}
     for(int syn_id = 0; syn_id < syn_N; syn_id++)  // loop through all synapses
     {
-        // pre/post_neuron_id are integers from 0 to Nsource/Ntarget (from corresponding SynapticPathway)
+        // pre/post_neuron_id are integers from 0 to the number of neurons in source/target group (from corresponding SynapticPathway)
         // this is relevant only when using Subgroups where they might be NOT equal to the idx in their NeuronGroup
-        int32_t pre_neuron_id = {{get_array_name(owner.synapse_sources, access_data=False)}}[syn_id] - {{owner.source.start}};
-        int32_t post_neuron_id = {{get_array_name(owner.synapse_targets, access_data=False)}}[syn_id] - {{owner.target.start}};
+        int32_t pre_neuron_id = {{get_array_name(owner.synapse_sources, access_data=False)}}[syn_id] - {{owner}}.spikes_start;
+        int32_t post_neuron_id = {{get_array_name(owner.synapse_targets, access_data=False)}}[syn_id] - {{owner}}.targets_start;
 
         {% if not no_or_const_delay_mode %}
         int delay = (int)({{_dynamic_delay}}[syn_id] / dt + 0.5);