Skip to content

Commit

Permalink
Copy variables to host after group_variable_set
Browse files Browse the repository at this point in the history
  • Loading branch information
denisalevi committed Jun 10, 2022
1 parent de69f0f commit 3d89de9
Show file tree
Hide file tree
Showing 2 changed files with 27 additions and 1 deletion.
22 changes: 22 additions & 0 deletions brian2cuda/templates/group_variable_set.cu
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,28 @@
///// endblock kernel_maincode /////
{% endblock %}

{% block extra_kernel_call_post %}
{# We need to copy modifed variables back to host in case they are used in
codeobjects that run on the host, which are synapse connect calls (e.g. in the
connect condition) and before run synapses push spikes, which initialized
synaptic variables.
#}
{% for var, varname in written_variables.items() %}
{% if var.dynamic %}
{{varname}} = dev{{varname}};
{% else %}
CUDA_SAFE_CALL(
cudaMemcpy(
{{varname}},
dev{{varname}},
sizeof({{c_data_type(var.dtype)}})*_num_{{varname}},
cudaMemcpyDeviceToHost
)
);
{% endif %}
{% endfor %}
{% endblock %}

{# _num_group_idx is defined in HOST_CONSTANTS, so we can't set _N before #}
{% block define_N %}
{% endblock %}
Expand Down
6 changes: 5 additions & 1 deletion brian2cuda/templates/group_variable_set_conditional.cu
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,11 @@


{% block extra_kernel_call_post %}
{# We want to copy only those variables that were potentially modified in aboves kernel call. #}
{# We need to copy modifed variables back to host in case they are used in
codeobjects that run on the host, which are synapse connect calls (e.g. in the
connect condition) and before run synapses push spikes, which initialized
synaptic variables.
#}
{% for var, varname in written_variables.items() %}
{% if var.dynamic %}
{{varname}} = dev{{varname}};
Expand Down

0 comments on commit 3d89de9

Please sign in to comment.