Parallelize `source` mode of synaptic effects application about brian2cuda HOT 4 OPEN

Copied relevant part for source mode optimization from my comment in #12:

No serializing needed for no_or_const_delay_mode when presynaptic variables are modified
In no_or_const_delay_mode we don't need to serialize when the on_pre pathway modifies presynaptic neuron variables as long as we loop through the spikespace (as we are currently doing), since we only apply synaptic effects for a single spiking neuron at a time. Currently this would get us in our 'pre' serializing mode (only a single thread does everything). So whenever we get into our 'pre' mode, we only need to modify a single presynaptic neurons state variables, so all we need to do, is to make sure that only a single thread does that. E.g. adding something like

if (tid==0 && bid==0)
   // presynaptic code
// postsynaptic code (one thread one postsynaptic neuron)

This would mean that in no_or_const_delay_mode we could - while applying postsynaptic effects in parallel anyways - also modify presynaptic ones on the fly (and synaptic variables also of course). Assuming only one synapse per pre/post neuron.

EDIT: In the 'pre' mode (now called 'source' mode) the presynaptic code has to be applied once per synapse. Therefore in aboves snippet thread 0 of block 0 would loop through all synapses and each time apply the same effect to the same source neuron. But maybe we could simplify this by just multiplying the effect times the number of synapses. We would have to catch different effect cases though: additive effects (instead of adding N times +v, do once +N*v); multiplicative effects (instead of N times *M, do once *(M^N)); setting a variable to a constant (which we would only need to apply once).

EDIT 2: We can only aggregate the effect if it does not depend on a postsynaptic variable. If it does, we are back to serializing. But maybe instead of having one thread load postsynaptic variables serialized, we could have multiple threads do coalesced memory reads and then let just one thread at the time apply the presynaptic effect, e.g. (needs profiling)

for (i=0; i<num_synapses; i++)
    if (tid == i)
       // apply presynaptic effect

from brian2cuda.

Came across an example using STDP with neuronal trace, which results in source mode pathways. So this can be relevant. @moritzaugustin

from brian2cuda.

good point

let's plan this for the release sprint (at least via atomics should not be too hard)

from brian2cuda.

Additional comment for my pre serialization comment above:

For the case where a source update depends on target state variables, my comment states that we need to serialize the summation of the postsynaptic variables. Instead, we could also use some reduction scheme, which can be easily performed via thrust::reduce, which can now be called form device code using dynamic parallelism. This could maybe even be performed concurrently while other stateupdated (from target or synapse variables) are performed and only in the end of the kernel one source variable update is performed.

from brian2cuda.