Sort EventMonitor data by index within each time step about brian2cuda HOT 2 OPEN

I can't reproduce this bug anymore and there are no more monitor tests failing in the test suite. Closing.

from brian2cuda.

Looks like this bug is still around, but only becomes apparent for large enough recordings. Because of warp-level SIMT, atomicAdd during thresholding will lead to neuron IDs being sorted if their indices are withing a single warp (group of 32 neurons). Since most tests use small neurongroups (typically less than 32 neurons), the data ends up being sorted. But the order of warps in the eventspace is random, hence for recordings of > 32 neurons, the eventspace is not sorted and the spikemonitor just copies from the eventspace is not sorted either.

Solution: Perform on copy operation per time step on the recorded data (either via thrust::sort or via custom sorting code in the spikemonitor kernel). Or find a way to sort the data in the end of the simulation only (maybe even in parallel on the GPU by storing the start indices for each time step during recording and then perform parallel sorting per time step in one kernel call at the end?).

For now this is a known issue and remains as such.

from brian2cuda.