Skip to content

Commit

Permalink
Fix wrong bundle parallelisation
Browse files Browse the repository at this point in the history
  • Loading branch information
denisalevi committed Mar 16, 2018
1 parent 40fcf22 commit 215fd13
Show file tree
Hide file tree
Showing 2 changed files with 12 additions and 16 deletions.
24 changes: 9 additions & 15 deletions brian2cuda/templates/synapses.cu
Original file line number Diff line number Diff line change
Expand Up @@ -130,24 +130,18 @@ num_loops = 1;
{% elif synaptic_effects == "target" %}
// Synaptic effects modify target group variables but NO source group variables.
num_blocks = num_parallel_blocks;
// if no bundles: num_threads = 1
num_threads_per_bundle = 1;
num_loops = 1;
if ({{pathway.name}}_scalar_delay)
{
if ({{owner.name}}_multiple_pre_post)
num_threads = 1;
else
num_threads = max_threads_per_block;
}
else // heterogeneous delays
{
num_threads = {{pathway.name}}_max_num_bundles;
if (num_threads > max_threads_per_block)
{
num_threads = 1;
if (!{{owner.name}}_multiple_pre_post){
if ({{pathway.name}}_scalar_delay)
num_threads = max_threads_per_block;
}
else // heterogeneous delays
num_threads = {{pathway.name}}_max_bundle_size;
}
if (num_threads > max_threads_per_block)
num_threads = max_threads_per_block;
// num_threads_per_bundle only used for heterogeneous delays
num_threads_per_bundle = num_threads;
{% elif synaptic_effects == "source" %}
// Synaptic effects modify source group variables.
num_blocks = 1;
Expand Down
4 changes: 3 additions & 1 deletion brian2cuda/templates/synapses_initialise_queue.cu
Original file line number Diff line number Diff line change
Expand Up @@ -359,6 +359,8 @@ void _run_{{pathobj}}_initialise_queue()
num_synapses = h_unique_delay_start_idx_by_pre_id[i][bundle_idx + 1] - synapses_start_idx;
//h_size_by_bundle_id_by_pre[i].push_back(num_synapses);
h_size_by_bundle_id.push_back(num_synapses);
if (num_synapses > {{pathobj}}_max_bundle_size)
{{pathobj}}_max_bundle_size = num_synapses;
int32_t* synapse_bundle = new int32_t[num_synapses];
// TODO: don't copy synapses to synapse_bundle, just cudaMemcpy it directly to the device with
// CudaSafeCall( cudaMemcpy(d_..., h_synapses_by_pre_id[i] + synapses_start_idx, ...)
Expand Down Expand Up @@ -443,7 +445,7 @@ void _run_{{pathobj}}_initialise_queue()
sum_num_unique_elements += num_unique_elements;
} // end if(num_elements < 0)
}
long unsigned int num_bundle_ids = sum_num_unique_elements;
num_bundle_ids = sum_num_unique_elements;
// floor(mean(h_size_by_bundle_id))
{{pathobj}}_mean_bundle_size = sum_num_synapses / num_bundle_ids;
//delete [] h_size_by_bundle_id_by_pre;
Expand Down

0 comments on commit 215fd13

Please sign in to comment.