Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Fix and improve summed variables implementation (including adding synapses to synapses support and start fixing spikegenerator) #186

Merged
merged 27 commits into from
Apr 8, 2021

Conversation

denisalevi
Copy link
Member

@denisalevi denisalevi commented Sep 3, 2020

PR for #49 [WIP]

We have three versions in the commits below that should be benchmarked against each other (see #49 discussion):

  1. Original algorithm with synapses * targets threads in total (d2ff3ff )
  2. One thread per synapse, each thread using atomicAdd on the target of its synapse in global memory to calculate summed variables (c54e64c)
  3. Same as 2. but using cudaMemset to reset summed variables instead of thrust::fill (653d487)

EDIT
This PR turned out to touch a couple of issue at the same time:

I did not do the benchmarking that I suggested above. This could be done in #197 if it turns out to be relevant.

@denisalevi denisalevi changed the base branch from master to timedarray_implementation September 3, 2020 11:34
This is needed when the index is used when setting group variables with
boolean indexing (e.g. `syn.w['k == 0'] = ...`, where `k` is the
multisynaptic index)
This reverts commit 7b764bd.
Was on wrong Brian2 branch (future release). Reapply this commit once
Brian2 is updated.
Parallelize threads over synapses, each thread does atomicAdd of the
synapse variable on the postsynaptic neuron of its synapse.
Better performance (eyeballed and assumed). Not 100% sure if this works
for all CUDA GPUs. Asked on NVIDIA forum:
https://forums.developer.nvidia.com/t/can-i-set-a-floats-to-zero-with-cudamemset/153706
This also sets CONSTANTS correctly in host templates, where previously
device pointers were used.
Fixes bug in periodicity (see #48) and makes sure the kernel works with
multiple spikespaces due to homogeneous delays. See comments in template
for potential optimizations.
@denisalevi denisalevi changed the base branch from timedarray_implementation to master March 24, 2021 12:08
@denisalevi denisalevi changed the title Fix and improve summed variables implementation Fix and improve summed variables implementation (including adding synapses to synapses support and start fixing spikegenerator) Apr 8, 2021
@denisalevi denisalevi merged commit 61e4a76 into master Apr 8, 2021
denisalevi added a commit that referenced this pull request Jun 10, 2021
Fix and improve summed variables implementation (including adding synapses to synapses support and start fixing spikegenerator)
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

1 participant