def annotate_update_neuron(self, pop, code):
"""
annotate the update neuron code
"""
if Global.config["num_threads"] == 1:
prof_begin = cpp11_profile_template['update_neuron']['before'] % {
'name': pop.name
}
prof_end = cpp11_profile_template['update_neuron']['after'] % {
'name': pop.name
}
else:
prof_begin = cpp11_omp_profile_template['update_neuron'][
'before'] % {
'name': pop.name
}
prof_end = cpp11_omp_profile_template['update_neuron']['after'] % {
'name': pop.name
}
prof_code = """
// first run, measuring average time
%(prof_begin)s
%(code)s
%(prof_end)s
""" % {
'code': code,
'prof_begin': tabify(prof_begin, 2),
'prof_end': tabify(prof_end, 2)
}
return prof_code
def annotate_computesum_spiking(self, proj, code):
"""
annotate the computesum compuation code
"""
if Global.config["num_threads"] == 1:
prof_begin = cpp11_profile_template['compute_psp']['before'] % {
'name': 'proj' + str(proj.id)
}
prof_end = cpp11_profile_template['compute_psp']['after'] % {
'name': 'proj' + str(proj.id)
}
else:
prof_begin = cpp11_omp_profile_template['compute_psp']['before'] % {
'name': 'proj' + str(proj.id)
}
prof_end = cpp11_omp_profile_template['compute_psp']['after'] % {
'name': 'proj' + str(proj.id)
}
prof_code = """
// first run, measuring average time
%(prof_begin)s
%(code)s
%(prof_end)s
""" % {
'code': code,
'prof_begin': tabify(prof_begin, 2),
'prof_end': tabify(prof_end, 2)
}
return prof_code
def annotate_update_rng(self, pop, code):
"""
annotate update rng kernel (only for CPUs available)
"""
if Global.config["num_threads"] == 1:
prof_begin = cpp11_profile_template['update_rng']['before'] % {
'name': pop.name
}
prof_end = cpp11_profile_template['update_rng']['after'] % {
'name': pop.name
}
else:
prof_begin = cpp11_omp_profile_template['update_rng']['before'] % {
'name': pop.name
}
prof_end = cpp11_omp_profile_template['update_rng']['after'] % {
'name': pop.name
}
prof_dict = {
'code': code,
'prof_begin': tabify(prof_begin, 2),
'prof_end': tabify(prof_end, 2)
}
prof_code = """
%(prof_begin)s
%(code)s
%(prof_end)s
"""
return prof_code % prof_dict
def annotate_spike_cond(self, pop, code):
"""
annotate the spike condition code
"""
if Global.config["num_threads"] == 1:
prof_begin = cpp11_profile_template['spike_gather']['before'] % {
'name': pop.name
}
prof_end = cpp11_profile_template['spike_gather']['after'] % {
'name': pop.name
}
else:
prof_begin = cpp11_omp_profile_template['spike_gather'][
'before'] % {
'name': pop.name
}
prof_end = cpp11_omp_profile_template['spike_gather']['after'] % {
'name': pop.name
}
prof_dict = {
'code': code,
'prof_begin': tabify(prof_begin, 2),
'prof_end': tabify(prof_end, 2)
}
prof_code = """
%(prof_begin)s
%(code)s
%(prof_end)s
""" % prof_dict
return prof_code
def _generate_cuda(self, convolve_code, sum_code):
"""
Update the ProjectionGenerator._specific_template structure and bypass the standard CUDA code generation.
"""
pool_operation = self.synapse_type.operation
# default value for sum in code depends on operation
sum_default = "0.0"
if pool_operation == "min":
sum_default = "FLT_MAX"
elif pool_operation == "max":
sum_default = "FLT_MIN"
pool_template = ""
pool_op_code = cuda_op_code[pool_operation]
pool_dict = {
'sum_default': sum_default,
'float_prec': Global.config['precision'],
}
if len(self.pre.geometry) == 2:
pool_template = cuda_pooling_code_2d
pool_dict.update({
'row_extent': self.extent[0],
'col_extent': self.extent[1],
'col_size': self.pre.geometry[1],
'operation': tabify(pool_op_code, 3)
})
pooling_code = pool_template % pool_dict
elif len(self.pre.geometry) == 3:
pool_template = cuda_pooling_code_3d
pool_dict.update({
'row_extent': self.extent[0],
'col_extent': self.extent[1],
'plane_extent': self.extent[2],
'row_size': self.pre.geometry[0],
'col_size': self.pre.geometry[1],
'plane_size': self.pre.geometry[2],
'operation': tabify(pool_op_code, 4)
})
pooling_code = pool_template % pool_dict
else:
raise NotImplementedError
# Specific template for generation
pool_dict = deepcopy(pooling_template_cuda)
for key, value in pool_dict.iteritems():
value = value % {
'id_proj': self.id,
'id_pre': self.pre.id,
'id_post': self.post.id,
'size_post': self.post.size,
'target': self.target,
'float_prec': Global.config['precision'],
'pooling_code': pooling_code
}
pool_dict[key] = value
self._specific_template.update(pool_dict)
def _generate_cuda(self, convolve_code, sum_code):
"""
Update the ProjectionGenerator._specific_template structure and bypass the standard CUDA code generation.
"""
pool_operation = self.synapse_type.operation
# default value for sum in code depends on operation
sum_default = "0.0"
if pool_operation == "min":
sum_default = "FLT_MAX"
elif pool_operation == "max":
sum_default = "FLT_MIN"
pool_template = ""
pool_op_code = cuda_op_code[pool_operation]
pool_dict = {
'sum_default': sum_default,
'float_prec': Global.config['precision'],
}
if len(self.pre.geometry) == 2:
pool_template = cuda_pooling_code_2d
pool_dict.update({
'row_extent': self.extent[0],
'col_extent': self.extent[1],
'col_size': self.pre.geometry[1],
'operation': tabify(pool_op_code, 3)
})
pooling_code = pool_template % pool_dict
elif len(self.pre.geometry) == 3:
pool_template = cuda_pooling_code_3d
pool_dict.update({
'row_extent': self.extent[0],
'col_extent': self.extent[1],
'plane_extent': self.extent[2],
'row_size': self.pre.geometry[0],
'col_size': self.pre.geometry[1],
'plane_size': self.pre.geometry[2],
'operation': tabify(pool_op_code, 4)
})
pooling_code = pool_template % pool_dict
else:
raise NotImplementedError
# Specific template for generation
pool_dict = deepcopy(pooling_template_cuda)
for key, value in pool_dict.iteritems():
value = value % {
'id_proj': self.id,
'id_pre': self.pre.id,
'id_post': self.post.id,
'size_post': self.post.size,
'target': self.target,
'float_prec': Global.config['precision'],
'pooling_code': pooling_code
}
pool_dict[key] = value
self._specific_template.update(pool_dict)
def _determine_size_in_bytes(self, pop):
"""
Generate code template to determine size in bytes for the C++ object *pop*. Please note, that this contain only
default elementes (parameters, variables).
User defined elements, parallelization support data structures or similar are not considered. Consequently
implementing generators should extent the resulting code template.
"""
from ANNarchy.generator.Utils import tabify
code = ""
# Parameters
code += "// Parameters\n"
for attr in pop.neuron_type.description['parameters']:
ids = {'ctype': attr['ctype'], 'name': attr['name']}
if attr['locality'] == "global":
code += "size_in_bytes += sizeof(%(ctype)s);\t// %(name)s\n" % ids
else:
code += "size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s\n" % ids
# Variables
code += "// Variables\n"
for attr in pop.neuron_type.description['variables']:
ids = {'ctype': attr['ctype'], 'name': attr['name']}
if attr['locality'] == "global":
code += "size_in_bytes += sizeof(%(ctype)s);\t// %(name)s\n" % ids
else:
code += "size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s\n" % ids
code = tabify(code, 2)
return code
def _clear_container(self, pop):
"""
Generate code template to destroy allocated container of the C++ object *pop*.
User defined elements, parallelization support data structures or similar are not considered. Consequently
implementing generators should extent the resulting code template.
"""
from ANNarchy.generator.Utils import tabify
code = ""
# Variables
code += "// Variables\n"
for attr in pop.neuron_type.description['variables']:
# HD: we need to clear only local variables, the others are no vectors
if attr['locality'] == "local":
# HD: clear alone does not deallocate, it only resets size.
# So we need to call shrink_to_fit afterwards.
ids = {'ctype': attr['ctype'], 'name': attr['name']}
code += "%(name)s.clear();\n" % ids
code += "%(name)s.shrink_to_fit();\n" % ids
# Mean-FR
if pop.neuron_type.description['type'] == 'spike':
code += """
// Mean Firing Rate
for (auto it = _spike_history.begin(); it != _spike_history.end(); it++) {
while(!it->empty())
it->pop();
}
_spike_history.clear();
_spike_history.shrink_to_fit();
"""
code = tabify(code, 2)
return code
def _determine_size_in_bytes(self, pop):
"""
Generate code template to determine size in bytes for the C++ object *pop*. Please note, that this contain only
default elementes (parameters, variables).
User defined elements, parallelization support data structures or similar are not considered. Consequently
implementing generators should extent the resulting code template.
"""
from ANNarchy.generator.Utils import tabify
code = ""
# Parameters
code += "// Parameters\n"
for attr in pop.neuron_type.description['parameters']:
ids = {'ctype': attr['ctype'], 'name': attr['name']}
if attr['locality'] == "global":
code += "size_in_bytes += sizeof(%(ctype)s);\t// %(name)s\n" % ids
else:
code += "size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s\n" % ids
# Variables
code += "// Variables\n"
for attr in pop.neuron_type.description['variables']:
ids = {'ctype': attr['ctype'], 'name': attr['name']}
if attr['locality'] == "global":
code += "size_in_bytes += sizeof(%(ctype)s);\t// %(name)s\n" % ids
else:
code += "size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s\n" % ids
code = tabify(code, 2)
return code
def _init_globalops(self, pop):
"""
Generate the C++ codes for the initialization of global operations
within the Population::init_population() method.
"""
if len(pop.global_operations) == 0:
return ""
code = "\n// Initialize global operations\n"
for op in pop.global_operations:
ids = {
'op': op['function'],
'var': op['variable'],
'type': Global.config['precision']
}
if Global._check_paradigm("openmp"):
code += """_%(op)s_%(var)s = 0.0;
""" % ids
elif Global._check_paradigm("cuda"):
code += """_%(op)s_%(var)s = 0.0;
cudaMalloc((void**)&_gpu_%(op)s_%(var)s, sizeof(%(type)s));
""" % ids
else:
raise NotImplementedError
return tabify(code, 2)
def _body_structural_plasticity(self):
"""
Call of pruning or creating methods if necessary.
"""
# Pruning if any
pruning = ""
creating = ""
if Global.config['structural_plasticity']:
for proj in self._projections:
if 'pruning' in proj.synapse_type.description.keys():
pruning += tabify(
"proj%(id)s.pruning();" % {'id': proj.id}, 1)
if 'creating' in proj.synapse_type.description.keys():
creating += tabify(
"proj%(id)s.creating();" % {'id': proj.id}, 1)
return creating + pruning
def annotate_update_synapse(self, proj, code):
"""
annotate the update synapse code, generated by ProjectionGenerator.update_synapse()
"""
if Global.config["num_threads"] == 1:
prof_begin = cpp11_profile_template['update_synapse']['before']
prof_end = cpp11_profile_template['update_synapse']['after']
else:
prof_begin = cpp11_omp_profile_template['update_synapse']['before']
prof_end = cpp11_omp_profile_template['update_synapse']['after']
prof_code = """
// first run, measuring average time
%(prof_begin)s
%(code)s
%(prof_end)s
""" % {
'code': code,
'prof_begin': tabify(prof_begin, 2),
'prof_end': tabify(prof_end, 2)
}
return prof_code
def _determine_size_in_bytes(self, pop):
"""
Generate code template to determine size in bytes for the C++ object *pop*. Please note, that this contain only
default elements (parameters, variables). User defined elements, parallelization support data structures or
similar are not considered.
Consequently implementing generators should extent the resulting code template. This is done by filling the
'size_in_bytes' field in the _specific_template.
"""
if 'size_in_bytes' in pop._specific_template.keys():
return pop._specific_template['size_in_bytes']
from ANNarchy.generator.Utils import tabify
code = ""
# Parameters
code += "// Parameters\n"
for attr in pop.neuron_type.description['parameters']:
ids = {'ctype': attr['ctype'], 'name': attr['name']}
if attr['locality'] == "global":
code += "size_in_bytes += sizeof(%(ctype)s);\t// %(name)s\n" % ids
else:
code += "size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s\n" % ids
# Variables
code += "// Variables\n"
for attr in pop.neuron_type.description['variables']:
ids = {'ctype': attr['ctype'], 'name': attr['name']}
if attr['locality'] == "global":
code += "size_in_bytes += sizeof(%(ctype)s);\t// %(name)s\n" % ids
else:
code += "size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s\n" % ids
# Random variables
code += "// RNGs\n"
if Global._check_paradigm("openmp"):
for dist in pop.neuron_type.description['random_distributions']:
ids = {'ctype': dist['ctype'], 'name': dist['name']}
if dist['locality'] == "local":
code += "size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s\n" % ids
else:
code += "size_in_bytes += sizeof(%(ctype)s);\t// %(name)s\n" % ids
else:
for dist in pop.neuron_type.description['random_distributions']:
code += "size_in_bytes += sizeof(curandState*);\t// gpu_%(name)s\n" % {
'name': dist['name']
}
code = tabify(code, 2)
return code
def _update_globalops(self, pop):
"""
Update of global functions is a call of pre-implemented functions defined in GlobalOperationTemplate. In case
of openMP this calls will take place in the population header.
We consider two cases:
a) the number of neurons is small, then we compute the operation thread-wise with openMP tasks (TODO: overhead??)
b) the number of neurons is high enough for a parallel implementation, where we first compute the local result
and then reduce over all available threads.
"""
if len(pop.global_operations) == 0:
return ""
if True: # parallel reduction is currently disabled
from ANNarchy.generator.Template.GlobalOperationTemplate import global_operation_templates_st_call as call_template
code = ""
for op in pop.global_operations:
code += call_template[op['function']] % {'var': op['variable']}
return """
if ( _active ){
// register tasks
#pragma omp single nowait
{
%(code)s
}
#pragma omp taskwait
}""" % {
'code': code
}
else:
from ANNarchy.generator.Template.GlobalOperationTemplate import global_operation_templates_omp_call as call_template
from ANNarchy.generator.Template.GlobalOperationTemplate import global_operation_templates_omp_reduce as red_template
code = ""
for op in pop.global_operations:
code += call_template[op['function']] % {'var': op['variable']}
code += red_template[op['function']] % {'var': op['variable']}
return """
if ( _active ){
%(code)s
}""" % {
'code': tabify(code, 3)
}
def _clear_container(self, pop):
"""
Generate code template to destroy allocated container of the C++ object *pop*.
"""
from ANNarchy.generator.Utils import tabify
code = ""
# Variables
code += "// Variables\n"
for attr in pop.neuron_type.description['variables']:
# HD: we need to clear only local variables, the others are no vectors
if attr['locality'] == "local":
# HD: clear alone does not deallocate, it only resets size.
# So we need to call shrink_to_fit afterwards.
ids = {'ctype': attr['ctype'], 'name': attr['name']}
code += "%(name)s.clear();\n" % ids
code += "%(name)s.shrink_to_fit();\n" % ids
# Mean-FR
if pop.neuron_type.description['type'] == 'spike':
code += """
// Mean Firing Rate
for (auto it = _spike_history.begin(); it != _spike_history.end(); it++) {
while(!it->empty())
it->pop();
}
_spike_history.clear();
_spike_history.shrink_to_fit();
"""
# Random variables
code += "\n// RNGs\n"
for dist in pop.neuron_type.description['random_distributions']:
rng_ids = {
'id': pop.id,
'rd_name': dist['name'],
}
code += self._templates['rng'][dist['locality']]['clear'] % rng_ids
code = tabify(code, 2)
return code
def _clear_container(self, proj):
"""
Generate code template to destroy allocated container of the C++ object *proj*.
User defined elements, parallelization support data structures or similar are not considered. Consequently
implementing generators should extent the resulting code template.
"""
from ANNarchy.generator.Utils import tabify
code = ""
# Variables
code += "// Variables\n"
for attr in proj.synapse_type.description['variables']:
# HD: clear alone does not deallocate, it only resets size.
# So we need to call shrink_to_fit afterwards.
ids = {'ctype': attr['ctype'], 'name': attr['name']}
code += "%(name)s.clear();\n" % ids
code += "%(name)s.shrink_to_fit();\n" % ids
code = tabify(code, 2)
return code
def _clear_container(self, proj):
"""
Generate code template to destroy allocated container of the C++ object *proj*.
User defined elements, parallelization support data structures or similar are not considered. Consequently
implementing generators should extent the resulting code template.
"""
from ANNarchy.generator.Utils import tabify
code = ""
# Variables
code += "// Variables\n"
for attr in proj.synapse_type.description['variables']:
# HD: clear alone does not deallocate, it only resets size.
# So we need to call shrink_to_fit afterwards.
ids = {'ctype': attr['ctype'], 'name': attr['name']}
code += "%(name)s.clear();\n" % ids
code += "%(name)s.shrink_to_fit();\n" % ids
code = tabify(code, 2)
return code
def _determine_size_in_bytes(self, proj):
"""
Generate code template to determine size in bytes for the C++ object *proj*. Please note, that this contain only
default elementes (parameters, variables).
User defined elements, parallelization support data structures or similar are not considered. Consequently
implementing generators should extent the resulting code template.
"""
from ANNarchy.generator.Utils import tabify
code = ""
# Parameters
code += "// Parameters\n"
for attr in proj.synapse_type.description['parameters']:
ids = {'ctype': attr['ctype'], 'name': attr['name']}
code += "size_in_bytes += sizeof(%(ctype)s);\t// %(name)s\n" % ids
# Variables
code += "// Variables\n"
for attr in proj.synapse_type.description['variables']:
ids = {'ctype': attr['ctype'], 'name': attr['name']}
if attr['locality'] == "global":
code += "size_in_bytes += sizeof(%(ctype)s);\t// %(name)s\n" % ids
elif attr['locality'] == "semiglobal":
code += "size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s\n" % ids
else:
if proj._storage_format == "lil":
code += """size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s
for(auto it = %(name)s.begin(); it != %(name)s.end(); it++)
size_in_bytes += (it->capacity()) * sizeof(%(ctype)s);\n""" % ids
elif proj._storage_format == "csr":
code += """size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s""" % ids
else:
# TODO: sanity check???
pass
code = tabify(code, 2)
return code
def _determine_size_in_bytes(self, proj):
"""
Generate code template to determine size in bytes for the C++ object *proj*. Please note, that this contain only
default elementes (parameters, variables).
User defined elements, parallelization support data structures or similar are not considered. Consequently
implementing generators should extent the resulting code template.
"""
from ANNarchy.generator.Utils import tabify
code = ""
# Parameters
code += "// Parameters\n"
for attr in proj.synapse_type.description['parameters']:
ids = {'ctype': attr['ctype'], 'name': attr['name']}
code += "size_in_bytes += sizeof(%(ctype)s);\t// %(name)s\n" % ids
# Variables
code += "// Variables\n"
for attr in proj.synapse_type.description['variables']:
ids = {'ctype': attr['ctype'], 'name': attr['name']}
if attr['locality'] == "global":
code += "size_in_bytes += sizeof(%(ctype)s);\t// %(name)s\n" % ids
elif attr['locality'] == "semiglobal":
code += "size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s\n" % ids
else:
if proj._storage_format == "lil":
code += """size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s
for(auto it = %(name)s.begin(); it != %(name)s.end(); it++)
size_in_bytes += (it->capacity()) * sizeof(%(ctype)s);\n""" % ids
elif proj._storage_format == "csr":
code += """size_in_bytes += sizeof(%(ctype)s) * %(name)s.capacity();\t// %(name)s""" % ids
else:
# TODO: sanity check???
pass
code = tabify(code, 2)
return code
def reset_computesum(self, pop):
"""
For rate-coded neurons each step the weighted sum of inputs is computed. The implementation
codes of the computes_rate kernel expect zeroed arrays. The same applies for the AccProjections
used for the computation of the BOLD signal.
Hint: this method is called directly by CodeGenerator.
"""
code = ""
# HD: use set to remove doublons
for target in sorted(set(pop.targets)):
code += self._templates['rate_psp']['reset'] % {
'id': pop.id,
'name': pop.name,
'target': target,
'float_prec': Global.config['precision']
}
# we need to sync the memsets
if len(code) > 1:
code += tabify("#pragma omp barrier\n", 1)
return code
def _generate_cuda(self, convolve_code, sum_code):
"""
Update the ProjectionGenerator._specific_template structure and bypass the standard CUDA code generation.
"""
pool_operation = self.synapse_type.operation
# default value for sum in code depends on operation
sum_default = "0.0"
if pool_operation == "min":
sum_default = "FLT_MAX"
elif pool_operation == "max":
sum_default = "FLT_MIN"
# operation to perform
pool_op_code = cuda_op_code[pool_operation] % {
'float_prec': Global.config['precision']
}
# result dictionary with code for
# body, call and header
pool_template = {}
base_ids = {
'id_proj': self.id,
'id_pre': self.pre.id,
'id_post': self.post.id,
'target': self.target,
'float_prec': Global.config['precision'],
'size_post': self.post.size # TODO: population views?
}
# The correct templates depends on both
# kernel-geometry and extent
if len(self.pre.geometry) == 2:
# For small extents, we compute multiple coords within one warp. If one extent can fill alone
# a half-warp we switch to the other implementation.
if self.extent[0] < 6:
pool_op_reduce_code = cuda_pooling_code_2d_small_extent[
'reduce_code'][pool_operation] % {
'float_prec': Global.config['precision'],
'row_extent': int(self.extent[0]),
'col_extent': int(self.extent[1])
}
pool_dict = deepcopy(base_ids)
pool_dict.update({
'sum_default': sum_default,
'row_extent': int(self.extent[0]),
'col_extent': int(self.extent[1]),
'row_size': int(self.pre.geometry[0]),
'col_size': int(self.pre.geometry[1]),
'operation': tabify(pool_op_code, 3),
'operation_reduce': pool_op_reduce_code
})
pool_template['psp_body'] = cuda_pooling_code_2d_small_extent[
'psp_body'] % pool_dict
pool_template[
'psp_header'] = cuda_pooling_code_2d_small_extent[
'psp_header'] % pool_dict
pool_template['psp_call'] = cuda_pooling_code_2d_small_extent[
'psp_call'] % pool_dict
else:
pool_op_reduce_code = cuda_pooling_code_2d['reduce_code'][
pool_operation] % {
'float_prec': Global.config['precision'],
'row_extent': int(self.extent[0]),
'col_extent': int(self.extent[1])
}
pool_dict = deepcopy(base_ids)
pool_dict.update({
'sum_default':
sum_default,
'row_extent':
int(self.extent[0]),
'col_extent':
int(self.extent[1]),
'row_size':
int(self.pre.geometry[0]),
'col_size':
int(self.pre.geometry[1]),
'operation':
tabify(pool_op_code, 3),
'operation_reduce':
tabify(pool_op_reduce_code, 2)
})
pool_template['psp_body'] = remove_trailing_spaces(
cuda_pooling_code_2d['psp_body'] % pool_dict)
pool_template['psp_header'] = cuda_pooling_code_2d[
'psp_header'] % pool_dict
pool_template[
'psp_call'] = cuda_pooling_code_2d['psp_call'] % pool_dict
elif len(self.pre.geometry) == 3:
pool_dict = deepcopy(base_ids)
pool_dict.update({
'sum_default': sum_default,
'row_extent': self.extent[0],
'col_extent': self.extent[1],
'plane_extent': self.extent[2],
'row_size': self.pre.geometry[0],
'col_size': self.pre.geometry[1],
'plane_size': self.pre.geometry[2],
'operation': tabify(pool_op_code, 4)
})
pool_template['psp_body'] = remove_trailing_spaces(
cuda_pooling_code_3d['psp_body'] % pool_dict)
pool_template[
'psp_header'] = cuda_pooling_code_3d['psp_header'] % pool_dict
pool_template[
'psp_call'] = cuda_pooling_code_3d['psp_header'] % pool_dict
else:
raise NotImplementedError
# Update psp fields
self._specific_template.update(pool_template)
# Specific template for generation (wrapper, etc)
pool_dict = deepcopy(pooling_template_cuda)
for key, value in pool_dict.items():
value = value % base_ids
pool_dict[key] = value
self._specific_template.update(pool_dict)
self._specific_template['wrapper_connector_call'] = ""
self._specific_template['access_parameters_variables'] = ""
self._specific_template['size_in_bytes'] = "//TODO:\n"
def _pop_recorder_class(self, pop):
"""
Creates population recording class code.
Returns:
* complete code as string
Templates:
omp_population, cuda_population
"""
if Global.config['paradigm'] == "openmp":
template = RecTemplate.omp_population
elif Global.config['paradigm'] == "cuda":
template = RecTemplate.cuda_population
else:
raise NotImplementedError
tpl_code = template['template']
init_code = ""
recording_code = ""
recording_target_code = ""
struct_code = ""
determine_size = ""
# The post-synaptic potential for rate-code (weighted sum) as well
# as the conductance variables are handled seperatly.
target_list = []
targets = []
for t in pop.neuron_type.description['targets']:
if isinstance(t, list):
for t2 in t:
targets.append(t2)
else:
targets.append(t)
for t in pop.targets:
if isinstance(t, list):
for t2 in t:
targets.append(t2)
else:
targets.append(t)
targets = sorted(list(set(targets)))
if pop.neuron_type.type == 'rate':
for target in targets:
tar_dict = {'id': pop.id, 'type' : Global.config['precision'], 'name': '_sum_'+target}
struct_code += template['local']['struct'] % tar_dict
init_code += template['local']['init'] % tar_dict
recording_target_code += template['local']['recording'] % tar_dict
else:
for target in targets:
tar_dict = {'id': pop.id, 'type' : Global.config['precision'], 'name': 'g_'+target}
struct_code += template['local']['struct'] % tar_dict
init_code += template['local']['init'] % tar_dict
recording_target_code += template['local']['recording'] % tar_dict
# to skip this entry in the following loop
target_list.append('g_'+target)
# Record global and local attributes
attributes = []
for var in pop.neuron_type.description['parameters'] + pop.neuron_type.description['variables']:
# Skip targets
if var['name'] in target_list:
continue
# Avoid doublons
if var['name'] in attributes:
continue
attributes.append(var['name'])
ids = {
'id': pop.id,
'name': var['name'],
'type': var['ctype']
}
struct_code += template[var['locality']]['struct'] % ids
init_code += template[var['locality']]['init'] % ids
recording_code += template[var['locality']]['recording'] % ids
# Memory management
if var['locality'] == "global":
determine_size += "size_in_bytes += sizeof(%(type)s);\t//%(name)s\n" % ids
else:
determine_size += "size_in_bytes += sizeof(%(type)s) * %(name)s.capacity();\t//%(name)s\n" % ids
# Spike events
if pop.neuron_type.type == 'spike':
struct_code += """
// Local variable %(name)s
std::map<int, std::vector< %(type)s > > %(name)s ;
bool record_%(name)s ;
void clear_spike() {
for ( auto it = spike.begin(); it != spike.end(); it++ ) {
it->second.clear();
}
}
""" % {'type' : 'long int', 'name': 'spike'}
init_code += """
this->%(name)s = std::map<int, std::vector< %(type)s > >();
if(!this->partial){
for(int i=0; i<pop%(id)s.size; i++) {
this->%(name)s[i]=std::vector<%(type)s>();
}
}
else{
for(int i=0; i<this->ranks.size(); i++) {
this->%(name)s[this->ranks[i]]=std::vector<%(type)s>();
}
}
this->record_%(name)s = false; """ % {'id': pop.id, 'type' : 'long int', 'name': 'spike'}
recording_code += RecTemplate.recording_spike_tpl[Global.config['paradigm']] % {'id': pop.id, 'type' : 'int', 'name': 'spike'}
ids = {
'id': pop.id,
'init_code': init_code,
'struct_code': struct_code,
'recording_code': recording_code,
'recording_target_code': recording_target_code,
'determine_size': tabify(determine_size, 2)
}
return tpl_code % ids
def _generate_bank_code(self):
# Operation to be performed: sum, max, min, mean
operation = self.synapse_type.operation
# Main code
code = tabify("sum = 0.0;", 3)
# Generate for loops
for dim in range(self.dim_kernel-1):
code += tabify("""
for(int %(index)s_w = 0; %(index)s_w < %(size)s;%(index)s_w++) {
""" % { 'index': indices[dim], 'size': self.weights.shape[dim+1]}, dim)
# Compute indices
if dim < self.dim_kernel:
code += tabify("""int %(index)s_pre = coord[%(dim)s] %(operator)s (%(index)s_w - %(center)s);""" % { 'id_proj': self.id, 'index': indices[dim], 'dim': dim, 'operator': '-' if self.method=='convolution' else '+', 'center': self._center_filter(self.weights.shape[dim+1])}, 1)
else:
code += tabify("""int %(index)s_pre = coord[%(dim)s];""" % { 'id_proj': self.id, 'index': indices[dim], 'dim': dim}, 1)
# Check indices
if operation in ['sum', 'mean']:
if isinstance(self.padding, str): # 'border'
code += tabify("""
if (%(index)s_pre < 0) %(index)s_pre = 0 ;
if (%(index)s_pre > %(max_size)s) %(index)s_pre = %(max_size)s ;
""" % { 'index': indices[dim], 'dim': dim, 'max_size': self.pre.geometry[dim] -1}, 1+dim)
else:
code += tabify("""
if ((%(index)s_pre < 0) || (%(index)s_pre > %(max_size)s)) {
sum += %(padding)s;
continue;
}
""" % { 'index': indices[dim], 'padding': self.padding, 'max_size': self.pre.geometry[dim] -1}, 1+dim)
else: # min, max
code += tabify("""
if ((%(index)s_pre < 0) || (%(index)s_pre > %(max_size)s)){
continue;
}
""" % { 'index': indices[dim], 'max_size': self.pre.geometry[dim] -1}, 1+dim)
# Compute pre-synaptic rank
code +=tabify("""
rk_pre = %(value)s;""" % {'value': self._coordinates_to_rank('pre', self.pre.geometry)}, 1+dim)
# Compute the increment
index = "[coord["+str(self.dim_pre)+"]]"
for dim in range(self.dim_kernel-1):
index += '[' + indices[dim] + '_w]'
increment = self.synapse_type.description['psp']['cpp'] % {
'id_pre': self.pre.id,
'id_post': self.post.id,
'local_index': index,
'global_index': '[i]',
'pre_index': '[rk_pre]',
'post_index': '[rk_post]',
'pre_prefix': 'pop'+str(self.pre.id)+'.',
'post_prefix': 'pop'+str(self.post.id)+'.'}
# Delays
if self.delays > Global.config['dt']:
increment = increment.replace(
'pop%(id_pre)s.r[rk_pre]' % {'id_pre': self.pre.id},
'delayed_r[rk_pre]'
)
# Apply the operation
if operation == "sum":
code += tabify("""
sum += %(increment)s""" % {'increment': increment}, 1+dim)
elif operation == "max":
code += tabify("""
double _psp = %(increment)s
if(_psp > sum) sum = _psp;""" % {'increment': increment}, 1+dim)
elif operation == "min":
code += tabify("""
double _psp = %(increment)s
if(_psp < sum) sum = _psp;""" % {'increment': increment}, 1+dim)
elif operation == "mean":
code += tabify("""
sum += %(increment)s""" % {'increment': increment}, 1+dim)
else:
Global._error('Operation', operation, 'is not implemented yet for shared projections.')
# Close for loops
for dim in range(self.dim_kernel-1):
code += tabify("""
}""", self.dim_kernel-1-dim)
impl_code = code % {'id_proj': self.id,
'target': self.target,
'id_pre': self.pre.id, 'name_pre': self.pre.name, 'size_pre': self.pre.size,
'id_post': self.post.id, 'name_post': self.post.name, 'size_post': self.post.size
}
# sum code
if operation == "mean":
sum_code = """sum/%(filter_size)s""" % {'filter_size': self.weights.size}
else:
sum_code = "sum"
return impl_code, sum_code
def _pop_recorder_class(self, pop):
"""
Creates population recording class code.
Returns:
* complete code as string
Templates:
omp_population, cuda_population
"""
if Global.config['paradigm'] == "openmp":
template = RecTemplate.omp_population
elif Global.config['paradigm'] == "cuda":
template = RecTemplate.cuda_population
else:
raise NotImplementedError
tpl_code = template['template']
init_code = ""
recording_code = ""
recording_target_code = ""
struct_code = ""
determine_size = ""
# The post-synaptic potential for rate-code (weighted sum) as well
# as the conductance variables are handled seperatly.
target_list = []
targets = []
for t in pop.neuron_type.description['targets']:
if isinstance(t, list):
for t2 in t:
targets.append(t2)
else:
targets.append(t)
for t in pop.targets:
if isinstance(t, list):
for t2 in t:
targets.append(t2)
else:
targets.append(t)
targets = sorted(list(set(targets)))
if pop.neuron_type.type == 'rate':
for target in targets:
tar_dict = {
'id': pop.id,
'type': Global.config['precision'],
'name': '_sum_' + target
}
struct_code += template['local']['struct'] % tar_dict
init_code += template['local']['init'] % tar_dict
recording_target_code += template['local'][
'recording'] % tar_dict
else:
for target in targets:
tar_dict = {
'id': pop.id,
'type': Global.config['precision'],
'name': 'g_' + target
}
struct_code += template['local']['struct'] % tar_dict
init_code += template['local']['init'] % tar_dict
recording_target_code += template['local'][
'recording'] % tar_dict
# to skip this entry in the following loop
target_list.append('g_' + target)
# Record global and local attributes
attributes = []
for var in pop.neuron_type.description[
'parameters'] + pop.neuron_type.description['variables']:
# Skip targets
if var['name'] in target_list:
continue
# Avoid doublons
if var['name'] in attributes:
continue
attributes.append(var['name'])
ids = {'id': pop.id, 'name': var['name'], 'type': var['ctype']}
struct_code += template[var['locality']]['struct'] % ids
init_code += template[var['locality']]['init'] % ids
recording_code += template[var['locality']]['recording'] % ids
# Memory management
if var['locality'] == "global":
determine_size += "size_in_bytes += sizeof(%(type)s);\t//%(name)s\n" % ids
else:
determine_size += "size_in_bytes += sizeof(%(type)s) * %(name)s.capacity();\t//%(name)s\n" % ids
# Spike events
if pop.neuron_type.type == 'spike':
struct_code += """
// Local variable %(name)s
std::map<int, std::vector< %(type)s > > %(name)s ;
bool record_%(name)s ;
void clear_spike() {
for ( auto it = spike.begin(); it != spike.end(); it++ ) {
it->second.clear();
}
}
""" % {
'type': 'long int',
'name': 'spike'
}
init_code += """
this->%(name)s = std::map<int, std::vector< %(type)s > >();
if(!this->partial){
for(int i=0; i<pop%(id)s.size; i++) {
this->%(name)s[i]=std::vector<%(type)s>();
}
}
else{
for(int i=0; i<this->ranks.size(); i++) {
this->%(name)s[this->ranks[i]]=std::vector<%(type)s>();
}
}
this->record_%(name)s = false; """ % {
'id': pop.id,
'type': 'long int',
'name': 'spike'
}
recording_code += RecTemplate.recording_spike_tpl[
Global.config['paradigm']] % {
'id': pop.id,
'type': 'int',
'name': 'spike'
}
ids = {
'id': pop.id,
'init_code': init_code,
'struct_code': struct_code,
'recording_code': recording_code,
'recording_target_code': recording_target_code,
'determine_size': tabify(determine_size, 2)
}
return tpl_code % ids
def _computesum_spiking(self, proj):
"""
Generate code for the spike propagation. As ANNarchy supports a set of
different data structures, this method split in up into several sub
functions.
In contrast to _computsum_rate() the spike propagation kernel need
to implement the signal transmission (event- as well as continous)
and also the equations filled in the 'pre-spike' field of synapse
desctiption.
"""
# Specific template ?
if 'header' in proj._specific_template.keys() and \
'body' in proj._specific_template.keys() and \
'call' in proj._specific_template.keys():
try:
header = proj._specific_template['header']
body = proj._specific_template['body']
call = proj._specific_template['call']
except KeyError:
Global._error('header,spike_count body and call should be overwritten')
return header, body, call
# some variables needed for the final templates
psp_code = ""
kernel_args = ""
kernel_args_call = ""
pre_spike_code = ""
kernel_deps = []
if proj.max_delay > 1 and proj.uniform_delay == -1:
Global._error("Non-uniform delays are not supported yet on GPUs.")
# some basic definitions
ids = {
# identifiers
'id_proj' : proj.id,
'id_post': proj.post.id,
'id_pre': proj.pre.id,
# common for all equations
'local_index': "[syn_idx]",
'semiglobal_index': '[post_rank]',
'global_index': '[0]',
'float_prec': Global.config['precision'],
# psp specific
'pre_prefix': 'pre_',
'post_prefix': 'post_',
'pre_index': '[col_idx[syn_idx]]',
'post_index': '[post_rank]'
}
#
# All statements in the 'pre_spike' field of synapse description
#
for var in proj.synapse_type.description['pre_spike']:
if var['name'] == "g_target": # synaptic transmission
psp_dict = {
'psp': var['cpp'].split('=')[1] % ids,
'float_prec': Global.config['precision']
}
psp_code += "%(float_prec)s tmp = %(psp)s\natomicAdd(&g_target[post_rank], tmp);" % psp_dict
else:
condition = ""
# Check conditions to update the variable
if var['name'] == 'w': # Surround it by the learning flag
condition = "plasticity"
# Flags avoids pre-spike evaluation when post fires at the same time
if 'unless_post' in var['flags']:
simultaneous = "pop%(id_pre)s_last_spike[_pr] != pop%(id_post)s_last_spike[%(semiglobal_index)s]" % ids
if condition == "":
condition = simultaneous
else:
condition += "&&(" + simultaneous + ")"
eq_dict = {
'eq': var['eq'],
'cpp': var['cpp'] % ids,
'bounds': get_bounds(var) % ids,
'condition': condition,
}
# Generate the code
if condition != "":
pre_spike_code += """
// unless_post can prevent evaluation of presynaptic variables
if(%(condition)s){
// %(eq)s
%(cpp)s
%(bounds)s
}
""" % eq_dict
else: # Normal synaptic variable
pre_spike_code += """
// %(eq)s
%(cpp)s
%(bounds)s""" % eq_dict
# Update the dependencies
kernel_deps.append(var['name']) # right side
for dep in var['dependencies']: # left side
kernel_deps.append(dep)
#
# Event-driven integration of synaptic variables
#
has_exact = False
event_driven_code = ''
for var in proj.synapse_type.description['variables']:
if var['method'] == 'event-driven':
has_exact = True
event_dict = {
'eq': var['eq'],
'exact': var['cpp'].replace('(t)', '(t-1)') % ids
}
event_driven_code += """
// %(eq)s
%(exact)s
""" % event_dict
# add the dependencies to kernel dependencies
for dep in var['dependencies']:
kernel_deps.append(dep)
# Does an event-driven variable occur?
if has_exact:
event_driven_code += """
// Update the last event for the synapse
_last_event%(local_index)s = t;
""" % ids
# event-driven requires access to last event variable
kernel_args += ", long* _last_event"
kernel_args_call += ", proj%(id_proj)s._gpu_last_event" % ids
# Add pre- and post-synaptic population dependencies
pre_post_deps = list(set(proj.synapse_type.description['dependencies']['pre'] + proj.synapse_type.description['dependencies']['post']))
pre_post_args = self._gen_kernel_args(proj, pre_post_deps, pre_post_deps)
kernel_args += pre_post_args[0]
kernel_args_call += pre_post_args[1]
# Add synaptic variables to kernel arguments
kernel_deps = list(set(kernel_deps)) # sort out doublings
for dep in kernel_deps:
if dep == "w" or dep == "g_target":
# already contained
continue
_, attr = self._get_attr_and_type(proj, dep)
attr_ids = {
'id_proj': proj.id,
'name': attr['name'],
'type': attr['ctype']
}
kernel_args += ", %(type)s* %(name)s" % attr_ids
kernel_args_call += ", proj%(id_proj)s.gpu_%(name)s" % attr_ids
#
# Finally, fill the templates
#
if 'psp' in proj.synapse_type.description.keys(): # not event-based
# transfrom psp equation
psp_code = proj.synapse_type.description['psp']['cpp']
# update dependencies
for dep in proj.synapse_type.description['psp']['dependencies']:
if dep == "w":
continue
_, attr = self._get_attr_and_type(proj, dep)
attr_ids = {
'id_proj': proj.id,
'type': attr['ctype'],
'name': attr['name']
}
kernel_args += ", %(type)s* %(name)s" % attr_ids
kernel_args_call += ", proj%(id_proj)s.gpu_%(name)s" % attr_ids
psp_code = proj.synapse_type.description['psp']['cpp'] % ids
# select the correct template
template = self._templates['spike_transmission']['continous'][proj._storage_order]
call = ""
target_list = proj.target if isinstance(proj.target, list) else [proj.target]
for target in target_list:
call += template['call'] % {
'id_proj': proj.id,
'id_pre': proj.pre.id,
'id_post': proj.post.id,
'target_arg': ', pop%(id_post)s.gpu_g_%(target)s' % {'id_post': proj.post.id, 'target': target},
'target': target,
'kernel_args': kernel_args_call,
'float_prec': Global.config['precision']
}
body = template['body'] % {
'id_proj': proj.id,
'target_arg': proj.target,
'kernel_args': kernel_args,
'psp': psp_code,
'pre_code': tabify(pre_spike_code, 3),
'float_prec': Global.config['precision']
}
header = template['header'] % {
'id': proj.id,
'kernel_args': kernel_args,
'target_arg': 'g_'+proj.target,
'float_prec': Global.config['precision']
}
else: # event-based
# select the correct template
template = self._templates['spike_transmission']['event_driven'][proj._storage_order]
# Connectivity description
if proj._storage_order == "post_to_pre":
conn_header = "int* col_ptr, int* row_idx, int* inv_idx, %(float_prec)s *w, %(float_prec)s* g_target" % ids
conn_call = "proj%(id_proj)s.gpu_col_ptr, proj%(id_proj)s.gpu_row_idx, proj%(id_proj)s.gpu_inv_idx, proj%(id_proj)s.gpu_w" % ids
else:
conn_call = "proj%(id_proj)s._gpu_row_ptr, proj%(id_proj)s._gpu_col_idx, proj%(id_proj)s.gpu_w" % ids
conn_body = "int* row_ptr, int* col_idx, %(float_prec)s* w, %(float_prec)s* g_target" %ids
conn_header = "int* row_ptr, int* col_idx, %(float_prec)s *w, %(float_prec)s* g_target" %ids
# Population sizes
pre_size = proj.pre.size if isinstance(proj.pre, Population) else proj.pre.population.size
post_size = proj.post.size if isinstance(proj.post, Population) else proj.post.population.size
call = ""
target_list = proj.target if isinstance(proj.target, list) else [proj.target]
for target in target_list:
call += template['call'] % {
'id_proj': proj.id,
'id_pre': proj.pre.id,
'id_post': proj.post.id,
'target': target,
'kernel_args': kernel_args_call % {'id_post': proj.post.id, 'target': target},
'conn_args': conn_call + ", pop%(id_post)s.gpu_g_%(target)s" % {'id_post': proj.post.id, 'target': target}
}
body = template['body'] % {
'id': proj.id,
'float_prec': Global.config['precision'],
'conn_arg': conn_header,
'kernel_args': kernel_args,
'event_driven': tabify(event_driven_code, 2),
'psp': tabify(psp_code, 4),
'pre_event': tabify(pre_spike_code, 4),
'pre_size': pre_size,
'post_size': post_size,
}
header = template['header'] % {
'id': proj.id,
'float_prec': Global.config['precision'],
'conn_header': conn_header,
'kernel_args': kernel_args
}
return header, body, call
def _update_spiking_neuron(self, pop):
"""
Generate the neural update code for GPU devices. We split up the
calculation into three parts:
* evolvement of global differential equations
* evolvement of local differential equations
* spike gathering
Return:
* tuple of three code snippets (body, header, call)
"""
# Is there any variable?
if len(pop.neuron_type.description['variables']) == 0:
return "", "", ""
# The purpose of this lines is explained in _update_rate_neuron
# HD: 19. May 2017
if 'update_variables' in pop._specific_template.keys():
call = """
// host side update of neurons
pop%(id)s.update();
""" % {'id': pop.id}
return "", "", call
header = ""
body = ""
local_call = ""
global_call = ""
# some defaults
ids = {
'id': pop.id,
'local_index': "[i]",
'global_index': "[0]"
}
#
# Process the global and local equations and generate first set of kernels
#
# parse the equations
glob_eqs = generate_equation_code(pop.id, pop.neuron_type.description, locality='global', padding=1) % ids
loc_eqs = generate_equation_code(pop.id, pop.neuron_type.description, locality='local', padding=2) % ids
# Gather pre-loop declaration (dt/tau for ODEs) and
# update the related kernels
pre_code = ""
for var in pop.neuron_type.description['variables']:
if 'pre_loop' in var.keys() and len(var['pre_loop']) > 0:
pre_code += Global.config['precision'] + ' ' + var['pre_loop']['name'] + ' = ' + var['pre_loop']['value'] + ';\n'
if pre_code.strip() != '':
pre_code = """
// Updating the step sizes
""" + tabify(pre_code, 1) % ids
# replace pow() for SDK < 6.5
loc_eqs = check_and_apply_pow_fix(loc_eqs)
glob_eqs = check_and_apply_pow_fix(glob_eqs)
# replace local function calls
if len(pop.neuron_type.description['functions']) > 0:
glob_eqs, loc_eqs = self._replace_local_funcs(pop, glob_eqs, loc_eqs)
# replace the random distributions
loc_eqs, glob_eqs = self._replace_random(loc_eqs, glob_eqs, pop.neuron_type.description['random_distributions'])
# Global variables
if glob_eqs.strip() != '':
add_args_header, add_args_call = self._gen_kernel_args(pop, 'global')
# global operations
for op in pop.global_operations:
ids = {
'id': pop.id,
'type': Global.config['precision'],
'op': op['function'],
'var': op['variable']
}
add_args_header += """, %(type)s _%(op)s_%(var)s """ % ids
add_args_call += """, pop%(id)s._%(op)s_%(var)s""" % ids
# finalize code templates
body += CUDATemplates.population_update_kernel['global']['body'] % {
'id': pop.id,
'add_args': add_args_header,
'pre_loop': pre_code,
'global_eqs':glob_eqs
}
header += CUDATemplates.population_update_kernel['global']['header'] % {
'id': pop.id, 'add_args': add_args_header
}
global_call = CUDATemplates.population_update_kernel['global']['call'] % {
'id': pop.id, 'add_args': add_args_call, 'stream_id': pop.id
}
# Local variables
if loc_eqs.strip() != '':
add_args_header, add_args_call = self._gen_kernel_args(pop, 'local')
# global operations
for op in pop.global_operations:
ids = {
'id': pop.id,
'type': Global.config['precision'],
'op': op['function'],
'var': op['variable']
}
add_args_header += """, %(type)s _%(op)s_%(var)s """ % ids
add_args_call += """, pop%(id)s._%(op)s_%(var)s""" % ids
# Is there a refractory period?
if pop.neuron_type.refractory or pop.refractory:
# within refractory perid, only conductance variables
refr_eqs = generate_equation_code(pop.id, pop.neuron_type.description, 'local', conductance_only=True, padding=3) % ids
# 'old' loc_eqs is now only executed ouside refractory period
loc_eqs = """
// Refractory period
if( refractory_remaining[i] > 0){
%(eqs)s
// Decrement the refractory period
refractory_remaining[i]--;
} else{
%(loc_eqs)s
}
""" % {'eqs': refr_eqs, 'loc_eqs': loc_eqs}
add_args_header += ", int *refractory, int* refractory_remaining"
add_args_call += """, pop%(id)s.gpu_refractory, pop%(id)s.gpu_refractory_remaining""" %{'id':pop.id}
# finalize code templates
body += CUDATemplates.population_update_kernel['local']['body'] % {
'id': pop.id, 'add_args': add_args_header, 'pop_size': pop.size, 'pre_loop': pre_code, 'local_eqs': loc_eqs
}
header += CUDATemplates.population_update_kernel['local']['header'] % {
'id': pop.id, 'add_args': add_args_header
}
local_call = CUDATemplates.population_update_kernel['local']['call'] % {
'id': pop.id, 'add_args': add_args_call, 'stream_id': pop.id
}
# Call statement consists of two parts
call = """
// Updating the local and global variables of population %(id)s
if ( pop%(id)s._active ) {
%(global_call)s
%(local_call)s
}
""" % {'id':pop.id, 'global_call': global_call, 'local_call': local_call}
if self._prof_gen:
call = self._prof_gen.annotate_update_neuron(pop, call)
#
# Process the spike condition and generate 2nd set of kernels
#
cond = pop.neuron_type.description['spike']['spike_cond'] % ids
reset = ""
for eq in pop.neuron_type.description['spike']['spike_reset']:
reset += """
%(reset)s
""" % {'reset': eq['cpp'] % ids}
# arguments
header_args = ""
call_args = ""
# gather all attributes required by this kernel
kernel_deps = []
for var in pop.neuron_type.description['spike']['spike_cond_dependencies']:
kernel_deps.append(var)
for reset_eq in pop.neuron_type.description['spike']['spike_reset']:
kernel_deps.append(reset_eq['name'])
for var in reset_eq['dependencies']:
kernel_deps.append(var)
kernel_deps = list(set(kernel_deps)) # remove doubled entries
# generate header, call and body args
for var in kernel_deps:
attr = self._get_attr(pop, var)
header_args += ", "+attr['ctype']+"* " + var
call_args += ", pop"+str(pop.id)+".gpu_"+var
# Is there a refractory period?
if pop.neuron_type.refractory or pop.refractory:
refrac_inc = "refractory_remaining[i] = refractory[i];"
header_args += ", int *refractory, int* refractory_remaining"
call_args += """, pop%(id)s.gpu_refractory, pop%(id)s.gpu_refractory_remaining""" %{'id':pop.id}
else:
refrac_inc = ""
# dependencies of CSR storage_order
if pop._storage_order == 'pre_to_post':
header_args += ", unsigned int* spike_count"
call_args += ", pop"+str(pop.id)+".gpu_spike_count"
spike_gather_decl = """volatile int pos = 0;
*spike_count = 0;"""
spike_count = """
// transfer back the spike counter (needed by record)
cudaMemcpyAsync( &pop%(id)s.spike_count, pop%(id)s.gpu_spike_count, sizeof(unsigned int), cudaMemcpyDeviceToHost, streams[%(stream_id)s]);
#ifdef _DEBUG
cudaError_t err = cudaGetLastError();
if ( err != cudaSuccess )
std::cout << "record_spike_count: " << cudaGetErrorString(err) << std::endl;
#endif""" %{'id':pop.id, 'stream_id':pop.id}
spike_count_cpy = """pop%(id)s.spike_count"""%{'id':pop.id}
else:
spike_gather_decl = ""
spike_count = ""
spike_count_cpy = """pop%(id)s.size"""%{'id':pop.id}
spike_gather = """
if ( %(cond)s ) {
%(reset)s
// store spike event
int pos = atomicAdd ( num_events, 1);
spiked[pos] = i;
last_spike[i] = t;
// refractory
%(refrac_inc)s
}
""" % {'cond': cond, 'reset': reset, 'refrac_inc': refrac_inc}
body += CUDATemplates.spike_gather_kernel['body'] % {
'id': pop.id,
'pop_size': str(pop.size),
'default': Global.config['precision'] + ' dt, int* spiked, long int* last_spike',
'args': header_args,
'decl': spike_gather_decl,
'spike_gather': spike_gather
}
header += CUDATemplates.spike_gather_kernel['header'] % {
'id': pop.id,
'default': Global.config['precision'] + ' dt, int* spiked, long int* last_spike',
'args': header_args
}
if pop.max_delay > 1:
default_args = 'dt, pop%(id)s.gpu_delayed_spiked.front(), pop%(id)s.gpu_last_spike' % {'id': pop.id}
else: # no_delay
default_args = 'dt, pop%(id)s.gpu_spiked, pop%(id)s.gpu_last_spike' % {'id': pop.id}
spike_gather = CUDATemplates.spike_gather_kernel['call'] % {
'id': pop.id,
'default': default_args,
'args': call_args % {'id': pop.id},
'stream_id': pop.id,
'spike_count': spike_count,
'spike_count_cpy': spike_count_cpy
}
if self._prof_gen:
spike_gather = self._prof_gen.annotate_spike_gather(pop, spike_gather)
call += spike_gather
return body, header, call
def _update_spiking_neuron(self, pop):
# Neural update
from ANNarchy.generator.Utils import generate_equation_code, tabify
# Is there a refractory period?
if pop.neuron_type.refractory or pop.refractory:
# Get the equations
eqs = generate_equation_code(
pop.id,
pop.neuron_type.description,
'local',
conductance_only=True,
padding=4) % { 'id': pop.id,
'local_index': "[i]",
'semiglobal_index': '',
'global_index': ''}
# Generate the code snippet
code = """
// Refractory period
if( refractory_remaining[i] > 0){
%(eqs)s
// Decrement the refractory period
refractory_remaining[i]--;
continue;
}
""" % {'eqs': eqs}
refrac_inc = "refractory_remaining[i] = refractory[i];"
else:
code = ""
refrac_inc = ""
# Global variables
global_code = ""
eqs = generate_equation_code(pop.id, pop.neuron_type.description, 'global', padding=3) % {'id': pop.id, 'local_index': "[i]", 'semiglobal_index': '', 'global_index': ''}
if eqs.strip() != "":
global_code += """
// Updating the global variables
%(eqs)s
""" % {'eqs': eqs}
# Gather pre-loop declaration (dt/tau for ODEs)
pre_code = ""
for var in pop.neuron_type.description['variables']:
if 'pre_loop' in var.keys() and len(var['pre_loop']) > 0:
pre_code += var['ctype'] + ' ' + var['pre_loop']['name'] + ' = ' + var['pre_loop']['value'] + ';\n'
if len(pre_code) > 0:
pre_code = """
// Updating the step sizes
""" + tabify(pre_code, 3)
global_code = pre_code % {'id': pop.id, 'local_index': "[i]", 'semiglobal_index': '', 'global_index': ''} + global_code
# OMP code
omp_code = "#pragma omp parallel for" if (Global.config['num_threads'] > 1 and pop.size > Global.OMP_MIN_NB_NEURONS) else ""
omp_critical_code = "#pragma omp critical" if (Global.config['num_threads'] > 1 and pop.size > Global.OMP_MIN_NB_NEURONS) else ""
# Local variables, evaluated in parallel
code += generate_equation_code(pop.id, pop.neuron_type.description, 'local', padding=4) % {'id': pop.id, 'local_index': "[i]", 'semiglobal_index': '', 'global_index': ''}
# Process the condition
cond = pop.neuron_type.description['spike']['spike_cond'] % {'id': pop.id, 'local_index': "[i]", 'semiglobal_index': '', 'global_index': ''}
# Reset equations
reset = ""
for eq in pop.neuron_type.description['spike']['spike_reset']:
reset += """
%(reset)s
""" % {'reset': eq['cpp'] % {'id': pop.id, 'local_index': "[i]", 'semiglobal_index': '', 'global_index': ''}}
# Mean Firing rate
mean_FR_push, mean_FR_update = self._update_fr(pop)
# Gather code
spike_gather = """
// Spike emission
if(%(condition)s){ // Condition is met
// Reset variables
%(reset)s
// Store the spike
%(omp_critical_code)s
{
spiked.push_back(i);
}
last_spike[i] = t;
// Refractory period
%(refrac_inc)s
%(mean_FR_push)s
}
%(mean_FR_update)s
"""% {'condition' : cond,
'reset': reset,
'refrac_inc': refrac_inc,
'mean_FR_push': mean_FR_push,
'mean_FR_update': mean_FR_update,
'omp_critical_code': omp_critical_code}
code += spike_gather
# finish code
final_code = """
if( _active ) {
spiked.clear();
%(global_code)s
// Updating local variables
%(omp_code)s
for(int i = 0; i < size; i++){
%(code)s
}
} // active
""" % {
'code': code,
'global_code': global_code,
'omp_code': omp_code
}
# if profiling enabled, annotate with profiling code
if self._prof_gen:
final_code = self._prof_gen.annotate_update_neuron(pop, final_code)
return final_code
def header_struct(self, pop, annarchy_dir):
"""
Specialized implementation of PopulationGenerator.header_struct() for
generation of an openMP header.
"""
# Generate declaration and accessors of all parameters and variables
declaration_parameters_variables, access_parameters_variables = self._generate_decl_and_acc(pop)
# Additional includes and structures
include_additional = ""
struct_additional = ""
declare_additional = ""
init_additional = ""
reset_additional = ""
access_additional = ""
# Declare global operations as extern at the beginning of the file
extern_global_operations = ""
for op in pop.global_operations:
extern_global_operations += global_op_extern_dict[op['function']] % {'type': Global.config['precision']}
# Initialize parameters and variables
init_parameters_variables = self._init_population(pop)
# Spike-specific stuff
reset_spike = ""
declare_spike = ""
init_spike = ""
if pop.neuron_type.description['type'] == 'spike':
spike_tpl = self._templates['spike_specific']
# Main data for spiking pops
declare_spike += spike_tpl['declare_spike'] % {'id': pop.id}
init_spike += spike_tpl['init_spike'] % {'id': pop.id}
reset_spike += spike_tpl['reset_spike'] % {'id': pop.id}
# If there is a refractory period
if pop.neuron_type.refractory or pop.refractory:
declare_spike += spike_tpl['declare_refractory'] % {'id': pop.id}
init_spike += spike_tpl['init_refractory'] % {'id': pop.id}
reset_spike += spike_tpl['reset_refractory'] % {'id': pop.id}
# Process eventual delay
declare_delay = ""; init_delay = ""; update_delay = ""; reset_delay = ""
if pop.max_delay > 1:
declare_delay, init_delay, update_delay, reset_delay = self._delay_code(pop)
# Process mean FR computations
declare_FR, init_FR = self._init_fr(pop)
update_FR = self._update_fr(pop)
# Update random distributions
update_rng = self._update_random_distributions(pop)
# Update global operations
update_global_ops = self._update_globalops(pop)
# Update the neural variables
if pop.neuron_type.type == "rate":
body, header, update_call = self._update_rate_neuron(pop)
else:
body, header, update_call = self._update_spiking_neuron(pop)
update_variables = ""
# Memory transfers
host_to_device, device_to_host = self._memory_transfers(pop)
# Stop condition
stop_condition = self._stop_condition(pop)
# Local functions
host_local_func, device_local_func = self._local_functions(pop)
declaration_parameters_variables += host_local_func
# Profiling
if self._prof_gen:
include_profile = """#include "Profiling.h"\n"""
declare_profile, init_profile = self._prof_gen.generate_init_population(pop)
else:
include_profile = ""
init_profile = ""
declare_profile = ""
## When everything is generated, we override the fields defined by the specific population
if 'include_additional' in pop._specific_template.keys():
include_additional = pop._specific_template['include_additional']
if 'struct_additional' in pop._specific_template.keys():
struct_additional = pop._specific_template['struct_additional']
if 'extern_global_operations' in pop._specific_template.keys():
extern_global_operations = pop._specific_template['extern_global_operations']
if 'declare_spike_arrays' in pop._specific_template.keys():
declare_spike = pop._specific_template['declare_spike_arrays']
if 'declare_parameters_variables' in pop._specific_template.keys():
declaration_parameters_variables = pop._specific_template['declare_parameters_variables']
if 'declare_additional' in pop._specific_template.keys():
declare_additional = pop._specific_template['declare_additional']
if 'declare_FR' in pop._specific_template.keys():
declare_FR = pop._specific_template['declare_FR']
if 'declare_delay' in pop._specific_template.keys() and pop.max_delay > 1:
declare_delay = pop._specific_template['declare_delay']
if 'access_parameters_variables' in pop._specific_template.keys():
access_parameters_variables = pop._specific_template['access_parameters_variables']
if 'access_additional' in pop._specific_template.keys():
access_additional = pop._specific_template['access_additional']
if 'init_parameters_variables' in pop._specific_template.keys():
init_parameters_variables = pop._specific_template['init_parameters_variables']
if 'init_spike' in pop._specific_template.keys():
init_spike = pop._specific_template['init_spike']
if 'init_delay' in pop._specific_template.keys() and pop.max_delay > 1:
init_delay = pop._specific_template['init_delay']
if 'init_FR' in pop._specific_template.keys():
init_FR = pop._specific_template['init_FR']
if 'init_additional' in pop._specific_template.keys():
init_additional = pop._specific_template['init_additional']
if 'reset_spike' in pop._specific_template.keys():
reset_spike = pop._specific_template['reset_spike']
if 'reset_delay' in pop._specific_template.keys() and pop.max_delay > 1:
reset_delay = pop._specific_template['reset_delay']
if 'reset_additional' in pop._specific_template.keys():
reset_additional = pop._specific_template['reset_additional']
if 'update_variables' in pop._specific_template.keys():
update_variables = pop._specific_template['update_variables']
if 'update_rng' in pop._specific_template.keys():
update_rng = pop._specific_template['update_rng']
if 'update_FR' in pop._specific_template.keys():
update_FR = pop._specific_template['update_FR']
if 'update_delay' in pop._specific_template.keys() and pop.max_delay > 1:
update_delay = pop._specific_template['update_delay']
if 'update_global_ops' in pop._specific_template.keys():
update_global_ops = pop._specific_template['update_global_ops']
# Fill the template
code = self._templates['population_header'] % {
'float_prec': Global.config['precision'],
'id': pop.id,
'name': pop.name,
'size': pop.size,
'include_additional': include_additional,
'include_profile': include_profile,
'struct_additional': struct_additional,
'extern_global_operations': extern_global_operations,
'declare_spike_arrays': declare_spike,
'declare_parameters_variables': declaration_parameters_variables,
'declare_additional': declare_additional,
'declare_delay': declare_delay,
'declare_FR': declare_FR,
'declare_profile': declare_profile,
'access_parameters_variables': access_parameters_variables,
'access_additional': access_additional,
'init_parameters_variables': init_parameters_variables,
'init_spike': init_spike,
'init_delay': init_delay,
'init_FR': init_FR,
'init_additional': init_additional,
'init_profile': init_profile,
'reset_spike': reset_spike,
'reset_delay': reset_delay,
'reset_additional': reset_additional,
'update_FR': update_FR,
'update_variables': update_variables,
'update_rng': update_rng,
'update_delay': update_delay,
'update_global_ops': update_global_ops,
'stop_condition': stop_condition,
'host_to_device': host_to_device,
'device_to_host': device_to_host
}
# Store the complete header definition in a single file
with open(annarchy_dir+'/generate/net'+str(self._net_id)+'/pop'+str(pop.id)+'.hpp', 'w') as ofile:
ofile.write(code)
# Basic informations common to all populations
pop_desc = {
'include': """#include "pop%(id)s.hpp"\n""" % {'id': pop.id},
'extern': """extern PopStruct%(id)s pop%(id)s;\n"""% {'id': pop.id},
'instance': """PopStruct%(id)s pop%(id)s;\n"""% {'id': pop.id},
'init': """ pop%(id)s.init_population();\n""" % {'id': pop.id}
}
pop_desc['custom_func'] = device_local_func
pop_desc['update'] = update_call
pop_desc['update_body'] = body
pop_desc['update_header'] = header
pop_desc['update_delay'] = """ pop%(id)s.update_delay();\n""" % {'id': pop.id} if pop.max_delay > 1 else ""
pop_desc['update_FR'] = """ pop%(id)s.update_FR();\n""" % {'id': pop.id} if pop.neuron_type.type == "spike" else ""
if len(pop.global_operations) > 0:
pop_desc['gops_update'] = self._update_globalops(pop) % {'id': pop.id}
pop_desc['host_to_device'] = tabify("pop%(id)s.host_to_device();" % {'id':pop.id}, 1)+"\n"
pop_desc['device_to_host'] = tabify("pop%(id)s.device_to_host();" % {'id':pop.id}, 1)+"\n"
return pop_desc
def header_struct(self, proj, annarchy_dir):
"""
Generate the codes for the pop[id].hpp file. This file contains
the c-style structure with all data members and equation codes (in
case of openMP).
"""
# configure Connectivity base class
self.configure(proj)
# Generate declarations and accessors for the variables
decl, accessor = self._declaration_accessors(proj)
# concurrent streams
decl['cuda_stream'] = cuda_templates['cuda_stream']
# Initiliaze the projection
init_parameters_variables = self._init_parameters_variables(proj)
variables_body, variables_header, variables_call = self._update_synapse(proj)
# Update the random distributions
init_rng = self._init_random_distributions(proj)
# Post event
post_event_body, post_event_header, post_event_call = self._post_event(proj)
# Compute sum is the trickiest part
psp_header, psp_body, psp_call = self._computesum_rate(proj) if proj.synapse_type.type == 'rate' else self._computesum_spiking(proj)
# Detect event-driven variables
has_event_driven = False
for var in proj.synapse_type.description['variables']:
if var['method'] == 'event-driven':
has_event_driven = True
# Detect delays to eventually generate the code
has_delay = proj.max_delay > 1
# Connectivity matrix
connectivity_matrix = self._connectivity(proj)
# Memory transfers
host_device_transfer, device_host_transfer = self._memory_transfers(proj)
# Local functions
host_local_func, device_local_func = self._local_functions(proj)
decl['parameters_variables'] += host_local_func
# Profiling
if self._prof_gen:
include_profile = """#include "Profiling.h"\n"""
declare_profile, init_profile = self._prof_gen.generate_init_projection(proj)
else:
include_profile = ""
init_profile = ""
declare_profile = ""
# Additional info (overwritten)
include_additional = ""
struct_additional = ""
init_additional = ""
access_additional = ""
cuda_flattening = ""
if 'include_additional' in proj._specific_template.keys():
include_additional = proj._specific_template['include_additional']
if 'struct_additional' in proj._specific_template.keys():
struct_additional = proj._specific_template['struct_additional']
if 'init_additional' in proj._specific_template.keys():
init_additional = proj._specific_template['init_additional']
if 'access_additional' in proj._specific_template.keys():
access_additional = proj._specific_template['access_additional']
inverse_connectivity_matrix = connectivity_matrix['init_inverse'] % {
'id_proj': proj.id,
'id_pre': proj.pre.id,
'id_post': proj.post.id,
'post_size': proj.post.size # only needed by CSR
}
if 'cuda_flattening' in proj._specific_template.keys():
cuda_flattening = proj._specific_template['cuda_flattening']
else:
if proj._storage_format == "lil":
cuda_flattening = self._templates['flattening'] % {
'id_post':proj.post.id
}
final_code = self._templates['projection_header'] % {
'id_pre': proj.pre.id,
'id_post': proj.post.id,
'id_proj': proj.id,
'name_pre': proj.pre.name,
'name_post': proj.post.name,
'target': proj.target,
'include_additional': include_additional,
'include_profile': include_profile,
'struct_additional': struct_additional,
'declare_connectivity_matrix': connectivity_matrix['declare'],
'declare_inverse_connectivity_matrix': connectivity_matrix['declare_inverse'],
'declare_delay': decl['declare_delay'] if has_delay else "",
'declare_event_driven': decl['event_driven'] if has_event_driven else "",
'declare_rng': decl['rng'],
'declare_parameters_variables': decl['parameters_variables'],
'declare_cuda_stream': decl['cuda_stream'],
'declare_additional': decl['additional'],
'declare_profile': declare_profile,
'init_connectivity_matrix': connectivity_matrix['init'] % {'float_prec': Global.config['precision']},
'init_inverse_connectivity_matrix': inverse_connectivity_matrix,
'init_event_driven': "",
'init_rng': init_rng,
'init_parameters_variables': init_parameters_variables,
'init_additional': init_additional,
'init_profile': init_profile,
'access_connectivity_matrix': connectivity_matrix['accessor'],
'access_parameters_variables': accessor,
'access_additional': access_additional,
'host_to_device': host_device_transfer,
'device_to_host': device_host_transfer,
'cuda_flattening': cuda_flattening
}
# Store the file in generate ( will be compared with files contained
# in build/ later on )
with open(annarchy_dir+'/generate/net'+str(self._net_id)+'/proj'+str(proj.id)+'.hpp', 'w') as ofile:
ofile.write(final_code)
# Build dictionary for inclusions in ANNarchy.cu
proj_desc = {
'include': """#include "proj%(id)s.hpp"\n""" % {'id': proj.id},
'extern': """extern ProjStruct%(id)s proj%(id)s;\n"""% {'id': proj.id},
'instance': """ProjStruct%(id)s proj%(id)s;\n"""% {'id': proj.id},
'init': """ proj%(id)s.init_projection();\n""" % {'id' : proj.id}
}
proj_desc['psp_header'] = psp_header
proj_desc['psp_body'] = psp_body
proj_desc['psp_call'] = psp_call
proj_desc['custom_func'] = device_local_func
proj_desc['update_synapse_header'] = variables_header
proj_desc['update_synapse_body'] = variables_body
proj_desc['update_synapse_call'] = variables_call
proj_desc['postevent_header'] = post_event_header
proj_desc['postevent_body'] = post_event_body
proj_desc['postevent_call'] = post_event_call
proj_desc['host_to_device'] = tabify("proj%(id)s.host_to_device();" % {'id':proj.id}, 1)+"\n"
proj_desc['device_to_host'] = tabify("proj%(id)s.device_to_host();" % {'id':proj.id}, 1)+"\n"
return proj_desc
def _update_rate_neuron(self, pop):
"""
Generate the execution code for updating neural variables, more precise local and global ones.
The resulting code comprise of several parts: creating of local and global update code, generating
function prototype and finally calling statement.
Returns:
a tuple of three strings, comprising of:
* body: kernel implementation
* header: kernel prototypes
* call: kernel call
"""
# HD ( 18. Nov. 2016 )
#
# In some user-defined cases the host and device side need something to do to
# in order to realize specific functionality. Yet I simply add a update()
# call, if update_variables was set.
if 'update_variables' in pop._specific_template.keys():
call = """
// host side update of neurons
pop%(id)s.update();
""" % {'id': pop.id}
return "", "", call
# Is there any variable?
if len(pop.neuron_type.description['variables']) == 0:
return "", "", ""
header = ""
body = ""
local_call = ""
global_call = ""
# some defaults
ids = {
'id': pop.id,
'local_index': "[i]",
'global_index': '[0]'
}
# parse the equations
glob_eqs = generate_equation_code(pop.id, pop.neuron_type.description, locality='global', padding=1) % ids
loc_eqs = generate_equation_code(pop.id, pop.neuron_type.description, locality='local', padding=2) % ids
# Gather pre-loop declaration (dt/tau for ODEs)
pre_code = ""
for var in pop.neuron_type.description['variables']:
if 'pre_loop' in var.keys() and len(var['pre_loop']) > 0:
pre_code += Global.config['precision'] + ' ' + var['pre_loop']['name'] + ' = ' + var['pre_loop']['value'] + ';\n'
if pre_code.strip() != '':
pre_code = """
// Updating the step sizes
""" + tabify(pre_code, 1) % ids
# sum() must generate _sum___all__[i] = _sum_exc[i] + sum_inh[i] + ... at the beginning of local equations
if '__all__' in pop.neuron_type.description['targets']:
eqs = " "*8 + "// Sum over all targets\n"
eqs += " "*8 + "_sum___all__[i] = "
for target in pop.targets:
eqs += "_sum_" + target + '[i] + '
eqs = eqs[:-2]
eqs += ';\n\n'
loc_eqs = eqs + loc_eqs
# replace pow() for SDK < 6.5
loc_eqs = check_and_apply_pow_fix(loc_eqs)
glob_eqs = check_and_apply_pow_fix(glob_eqs)
# replace local function calls
if len(pop.neuron_type.description['functions']) > 0:
glob_eqs, loc_eqs = self._replace_local_funcs(pop, glob_eqs, loc_eqs)
# replace the random distributions
loc_eqs, glob_eqs = self._replace_random(loc_eqs, glob_eqs, pop.neuron_type.description['random_distributions'])
# Global variables
if glob_eqs.strip() != '':
add_args_header, add_args_call = self._gen_kernel_args(pop, 'global')
# global operations
for op in pop.global_operations:
ids = {
'id': pop.id,
'type': Global.config['precision'],
'op': op['function'],
'var': op['variable']
}
add_args_header += """, %(type)s _%(op)s_%(var)s """ % ids
add_args_call += """, pop%(id)s._%(op)s_%(var)s""" % ids
# finalize code templates
body += CUDATemplates.population_update_kernel['global']['body'] % {
'id': pop.id,
'add_args': add_args_header,
'global_eqs':glob_eqs,
'pre_loop': pre_code
}
header += CUDATemplates.population_update_kernel['global']['header'] % {
'id': pop.id, 'add_args': add_args_header
}
global_call = CUDATemplates.population_update_kernel['global']['call'] % {
'id': pop.id, 'add_args': add_args_call
}
# Local variables
if loc_eqs.strip() != '':
add_args_header, add_args_call = self._gen_kernel_args(pop, 'local')
# targets
for target in sorted(list(set(pop.neuron_type.description['targets'] + pop.targets))):
add_args_header += """, %(type)s* _sum_%(target)s""" % {'type': Global.config['precision'], 'target' : target}
add_args_call += """, pop%(id)s.gpu__sum_%(target)s""" % {'id': pop.id, 'target' : target}
# global operations
for op in pop.global_operations:
ids = {
'id': pop.id,
'type': Global.config['precision'],
'op': op['function'],
'var': op['variable']
}
add_args_header += """, %(type)s _%(op)s_%(var)s """ % ids
add_args_call += """, pop%(id)s._%(op)s_%(var)s""" % ids
# finalize code templates
body += CUDATemplates.population_update_kernel['local']['body'] % {
'id': pop.id,
'add_args': add_args_header,
'pop_size': pop.size,
'local_eqs': loc_eqs,
'pre_loop': tabify(pre_code % ids, 1)
}
header += CUDATemplates.population_update_kernel['local']['header'] % {
'id': pop.id,
'add_args': add_args_header
}
local_call = CUDATemplates.population_update_kernel['local']['call'] % {
'id': pop.id,
'add_args': add_args_call
}
# Call statement consists of two parts
call = """
// Updating the local and global variables of population %(id)s
if ( pop%(id)s._active ) {
%(global_call)s
%(local_call)s
}
""" % {'id':pop.id, 'global_call': global_call, 'local_call': local_call}
if self._prof_gen:
call = self._prof_gen.annotate_update_neuron(pop, call)
return body, header, call
def _post_event(self, proj):
"""
Post-synaptic event kernel for CUDA devices
"""
if proj.synapse_type.type == "rate":
return "", "", ""
if proj.synapse_type.description['post_spike'] == []:
return "", "", ""
if proj._storage_format == "lil":
ids = {
'id_proj' : proj.id,
'target': proj.target,
'id_post': proj.post.id,
'id_pre': proj.pre.id,
'local_index': "[j]",
'semiglobal_index': '[i]',
'global_index': '[0]',
'pre_index': '[pre_rank[j]]',
'post_index': '[post_rank[i]]',
'pre_prefix': 'pre_',
'post_prefix': ''
}
elif proj._storage_format == "csr":
ids = {
'id_proj' : proj.id,
'target': proj.target,
'id_post': proj.post.id,
'id_pre': proj.pre.id,
'local_index': "[col_idx[j]]",
'semiglobal_index': '[i]',
'global_index': '',
'pre_index': '[pre_rank[j]]',
'post_index': '[post_rank[i]]',
'pre_prefix': 'pop'+ str(proj.pre.id) + '.',
'post_prefix': 'pop'+ str(proj.post.id) + '.'
}
else:
raise NotImplementedError
add_args_header = ""
add_args_call = ""
# Event-driven integration
has_event_driven = False
for var in proj.synapse_type.description['variables']:
if var['method'] == 'event-driven':
has_event_driven = True
# Generate event-driven code
event_driven_code = ""
event_deps = []
if has_event_driven:
# event-driven rely on last pre-synaptic event
add_args_header += ", long* _last_event"
add_args_call += ", proj%(id_proj)s._gpu_last_event" % {'id_proj': proj.id}
for var in proj.synapse_type.description['variables']:
if var['method'] == 'event-driven':
event_driven_code += '// ' + var['eq'] + '\n'
event_driven_code += var['cpp'] % ids + '\n'
for deps in var['dependencies']:
event_deps.append(deps)
event_driven_code += """
// Update the last event for the synapse
_last_event%(local_index)s = t;
""" % {'local_index' : '[j]'}
# Gather the equations
post_code = ""
post_deps = []
for post_eq in proj.synapse_type.description['post_spike']:
post_code += '// ' + post_eq['eq'] + '\n'
if post_eq['name'] == 'w':
post_code += "if(plasticity)\n"
post_code += post_eq['cpp'] % ids + '\n'
post_code += get_bounds(post_eq) % ids + '\n'
# add dependencies, only right side!
for deps in post_eq['dependencies']:
post_deps.append(deps)
# left side of equations is not part of dependencies
post_deps.append(post_eq['name'])
post_code = tabify(post_code, 2)
# Create add_args for event-driven eqs and post_event
kernel_deps = list(set(post_deps+event_deps)) # variables can occur in several eqs
for dep in kernel_deps:
if dep == "w":
continue
_, attr = self._get_attr_and_type(proj, dep)
add_args_header += ', %(type)s* %(name)s' % {'type': attr['ctype'], 'name': attr['name']}
add_args_call += ', proj%(id)s.gpu_%(name)s' % {'id': proj.id, 'name': attr['name']}
if proj._storage_format == "lil":
conn_header = "int* row_ptr, int* pre_ranks,"
conn_call = ", proj%(id_proj)s.gpu_row_ptr, proj%(id_proj)s.gpu_pre_rank"
templates = self._templates['post_event']['post_to_pre']
elif proj._storage_format == "csr":
conn_header = "int* row_ptr, int *col_idx, "
conn_call = ", proj%(id_proj)s._gpu_row_ptr, proj%(id_proj)s._gpu_col_idx"
templates = self._templates['post_event']['pre_to_post']
else:
raise NotImplementedError
postevent_header = templates['header'] % {
'id_proj': proj.id,
'conn_args': conn_header,
'add_args': add_args_header,
'float_prec': Global.config['precision']
}
postevent_body = templates['body'] % {
'id_proj': proj.id,
'conn_args': conn_header,
'add_args': add_args_header,
'event_driven': tabify(event_driven_code, 2),
'post_code': post_code,
'float_prec': Global.config['precision']
}
postevent_call = ""
target_list = proj.target if isinstance(proj.target, list) else [proj.target]
for target in target_list:
postevent_call += templates['call'] % {
'id_proj': proj.id,
'id_pre': proj.pre.id,
'id_post': proj.post.id,
'target': target,
'conn_args': conn_call % ids,
'add_args': add_args_call
}
return postevent_body, postevent_header, postevent_call
def _update_rate_neuron(self, pop):
"""
Generate the code template for neural update step, more precise updating of variables.
The code comprise of two major parts: global and local update, second one parallelized
with an openmp for construct, if number of threads is greater than one and the number
of neurons exceed a minimum amount of neurons ( defined as Global.OMP_MIN_NB_NEURONS)
"""
from ANNarchy.generator.Utils import generate_equation_code, tabify
code = ""
# Random distributions
deps =[]
for rd in pop.neuron_type.description['random_distributions']:
for dep in rd['dependencies']:
deps += dep
# Global variables
eqs = generate_equation_code(pop.id, pop.neuron_type.description, 'global', padding=3) % {'id': pop.id, 'local_index': "[i]", 'semiglobal_index': '', 'global_index': ''}
if eqs.strip() != "":
code += """
// Updating the global variables
%(eqs)s
""" % {'eqs': eqs}
# Gather pre-loop declaration (dt/tau for ODEs)
pre_code =""
for var in pop.neuron_type.description['variables']:
if 'pre_loop' in var.keys() and len(var['pre_loop']) > 0:
pre_code += var['ctype'] + ' ' + var['pre_loop']['name'] + ' = ' + var['pre_loop']['value'] + ';\n'
code = tabify(pre_code, 3) % {'id': pop.id, 'local_index': "[i]", 'semiglobal_index': '', 'global_index': ''} + code
eqs = ""
# sum() must generate _sum___all__[i] = _sum_exc[i] + sum_inh[i] + ... at the beginning
if '__all__' in pop.neuron_type.description['targets']:
eqs += " "*16 + "// Sum over all targets\n"
eqs += " "*16 + "_sum___all__[i] = "
for target in pop.targets:
eqs += "_sum_" + target + '[i] + '
eqs = eqs[:-2]
eqs += ';\n\n'
# Local variables, evaluated in parallel
eqs += generate_equation_code(pop.id, pop.neuron_type.description, 'local', padding=4) % {
'id': pop.id,
'local_index': "[i]",
'semiglobal_index': '',
'global_index': ''}
if eqs.strip() != "":
omp_code = "#pragma omp parallel for" if (Global.config['num_threads'] > 1 and pop.size > Global.OMP_MIN_NB_NEURONS) else ""
code += """
// Updating the local variables
%(omp_code)s
for(int i = 0; i < size; i++){
%(eqs)s
}
""" % {'eqs': eqs, 'omp_code': omp_code}
# finish code
final_code = """
if( _active ) {
%(code)s
} // active
""" % {'code': code}
# if profiling enabled, annotate with profiling code
if self._prof_gen:
final_code = self._prof_gen.annotate_update_neuron(pop, final_code)
return final_code
def _generate_omp(self,
filter_definition,
filter_pyx_definition,
convolve_code,
sum_code,
kernel=True):
"""
OpenMP code generation.
"""
# Basic ids
base_ids = {
'id_proj': self.id,
'size_post': self.post.size,
'float_prec': Global.config['precision']
}
# Fill the basic definitions
conv_dict = deepcopy(convole_template_omp)
for key, value in conv_dict.items():
value = value % base_ids
conv_dict[key] = value
self._specific_template.update(conv_dict)
# Kernel-based method: specify w with the correct dimension
if kernel:
self._specific_template['declare_parameters_variables'] = tabify(
filter_definition.strip(), 1)
self._specific_template['export_parameters_variables'] = ""
self._specific_template['access_parameters_variables'] = """
// Local parameter w
%(type_w)s get_w() { return w; }
void set_w(%(type_w)s value) { w = value; }
""" % {
'type_w': filter_definition.replace(' w;', '')
}
self._specific_template['export_connectivity'] += """
# Local variable w
%(type_w)s get_w()
void set_w(%(type_w)s)
""" % {
'type_w': filter_pyx_definition.replace(' w', '')
}
self._specific_template['wrapper_init_connectivity'] += """
proj%(id_proj)s.set_w(weights)
""" % {
'id_proj': self.id
}
self._specific_template['wrapper_access_connectivity'] += """
# Local variable w
def get_w(self):
return proj%(id_proj)s.get_w()
def set_w(self, value):
proj%(id_proj)s.set_w( value )
def get_dendrite_w(self, int rank):
return proj%(id_proj)s.get_w()
def set_dendrite_w(self, int rank, value):
proj%(id_proj)s.set_w(value)
def get_synapse_w(self, int rank_post, int rank_pre):
return 0.0
def set_synapse_w(self, int rank_post, int rank_pre, %(float_prec)s value):
pass
""" % {
'id_proj': self.id,
'float_prec': Global.config['precision']
}
# Override the monitor to avoid recording the weights
self._specific_template['monitor_class'] = ""
self._specific_template['monitor_export'] = ""
self._specific_template['monitor_wrapper'] = ""
# OMP code
omp_code = ""
if Global.config['num_threads'] > 1:
omp_code = """
#pragma omp for private(sum, rk_pre, coord) %(psp_schedule)s""" % {
'psp_schedule':
"" if not 'psp_schedule' in self._omp_config.keys() else
self._omp_config['psp_schedule']
}
# HD ( 16.10.2015 ):
# pre-load delayed firing rate in a local array, so we
# prevent multiple accesses to pop%(id_pre)s._delayed_r[delay-1]
# wheareas delay is set available as variable
# TODO HD: wouldn't it be much better to reduce delay globaly, instead of the substraction here???
if self.delays > Global.config['dt']:
pre_load_r = """
// pre-load delayed firing rate
auto delayed_r = pop%(id_pre)s._delayed_r[delay-1];
""" % {
'id_pre': self.pre.id
}
else:
pre_load_r = ""
# Compute sum
wsum = """
if ( _transmission && pop%(id_pre)s._active ) {
int* coord;
""" + pre_load_r + """
%(omp_code)s
for(int i = 0; i < %(size_post)s; i++){
coord = pre_coords[i].data();
// perform the convolution
""" + tabify(convolve_code, 1) + """
// store result
pop%(id_post)s._sum_%(target)s[i] += """ + sum_code + """;
} // for
} // if
"""
self._specific_template['psp_code'] = wsum % \
{ 'id_proj': self.id,
'target': self.target,
'id_pre': self.pre.id, 'name_pre': self.pre.name, 'size_pre': self.pre.size,
'id_post': self.post.id, 'name_post': self.post.name, 'size_post': self.post.size,
'omp_code': omp_code,
'convolve_code': convolve_code
}
self._specific_template['size_in_bytes'] = """
// post-ranks
size_in_bytes += sizeof(std::vector<int>);
size_in_bytes += post_rank.capacity() * sizeof(int);
// pre-coords
size_in_bytes += sizeof(std::vector<std::vector<int>>);
size_in_bytes += pre_coords.capacity() * sizeof(std::vector<int>);
for (auto it = pre_coords.begin(); it != pre_coords.end(); it++) {
size_in_bytes += it->capacity() * sizeof(int);
}
// filter
// TODO:
"""
self._specific_template['clear'] = """
def _update_rate_neuron(self, pop):
"""
Generate the code template for neural update step, more precise updating of variables.
The code comprise of two major parts: global and local update, second one parallelized
with an openmp for construct, if number of threads is greater than one and the number
of neurons exceed a minimum amount of neurons ( defined as Global.OMP_MIN_NB_NEURONS)
"""
code = ""
id_dict = {
'id': pop.id,
'local_index': "[i]",
'semiglobal_index': '',
'global_index': ''
}
# Random distributions
deps = []
for rd in pop.neuron_type.description['random_distributions']:
for dep in rd['dependencies']:
deps += dep
# Global variables
eqs = generate_equation_code(pop.id,
pop.neuron_type.description,
locality='global',
padding=3)
eqs = eqs % id_dict
if eqs.strip() != "":
code += """
// Updating the global variables
#pragma omp single
{
%(eqs)s
}
""" % {
'eqs': eqs
}
# Gather pre-loop declaration (dt/tau for ODEs)
pre_code = ""
for var in pop.neuron_type.description['variables']:
if 'pre_loop' in var.keys() and len(var['pre_loop']) > 0:
pre_code += var['ctype'] + ' ' + var['pre_loop'][
'name'] + ' = ' + var['pre_loop']['value'] + ';\n'
code = tabify(pre_code, 3) % {
'id': pop.id,
'local_index': "[i]",
'semiglobal_index': '',
'global_index': ''
} + code
eqs = ""
# sum() must generate _sum___all__[i] = _sum_exc[i] + sum_inh[i] + ... at the beginning
if '__all__' in pop.neuron_type.description['targets']:
eqs += " " * 16 + "// Sum over all targets\n"
eqs += " " * 16 + "_sum___all__[i] = "
for target in pop.targets:
eqs += "_sum_" + target + '[i] + '
eqs = eqs[:-2]
eqs += ';\n\n'
# Local variables, evaluated in parallel
eqs = generate_equation_code(pop.id,
pop.neuron_type.description,
'local',
padding=4)
eqs = eqs % id_dict
if eqs.strip() != "":
code += """
// Updating the local variables
#pragma omp for simd
for (int i = 0; i < size; i++) {
%(eqs)s
}
""" % {
'eqs': eqs
}
# finish code
final_code = """
if( _active ) {
%(code)s
} // active
""" % {
'code': code
}
# if profiling enabled, annotate with profiling code
if self._prof_gen:
final_code = self._prof_gen.annotate_update_neuron(pop, final_code)
return final_code
def _generate_bank_code(self):
# Operation to be performed: sum, max, min, mean
operation = self.synapse_type.operation
# Main code
code = tabify("sum = 0.0;\n", 3)
# Generate for loops
for dim in range(self.dim_kernel - 1):
code += tabify(
"""
for(int %(index)s_w = 0; %(index)s_w < %(size)s;%(index)s_w++) {
""" % {
'index': indices[dim],
'size': self.weights.shape[dim + 1]
}, dim)
# Compute indices
if dim < self.dim_kernel:
code += tabify(
"""int %(index)s_pre = coord[%(dim)s] %(operator)s (%(index)s_w - %(center)s);"""
% {
'id_proj': self.id,
'index': indices[dim],
'dim': dim,
'operator': '+',
'center': self._center_filter(
self.weights.shape[dim + 1])
}, 1)
else:
code += tabify(
"""int %(index)s_pre = coord[%(dim)s];""" % {
'id_proj': self.id,
'index': indices[dim],
'dim': dim
}, 1)
# Check indices
if operation in ['sum', 'mean']:
if isinstance(self.padding, str): # 'border'
code += tabify(
"""
if (%(index)s_pre < 0) %(index)s_pre = 0 ;
if (%(index)s_pre > %(max_size)s) %(index)s_pre = %(max_size)s ;
""" % {
'index': indices[dim],
'dim': dim,
'max_size': self.pre.geometry[dim] - 1
}, 1 + dim)
else:
code += tabify(
"""
if ((%(index)s_pre < 0) || (%(index)s_pre > %(max_size)s)) {
sum += %(padding)s;
continue;
}
""" % {
'index': indices[dim],
'padding': self.padding,
'max_size': self.pre.geometry[dim] - 1
}, 1 + dim)
else: # min, max
code += tabify(
"""
if ((%(index)s_pre < 0) || (%(index)s_pre > %(max_size)s)){
continue;
}
""" % {
'index': indices[dim],
'max_size': self.pre.geometry[dim] - 1
}, 1 + dim)
# Compute pre-synaptic rank
code += tabify(
"""
rk_pre = %(value)s;""" %
{'value': self._coordinates_to_rank('pre', self.pre.geometry)},
1 + dim)
# Compute the increment
index = "[coord[" + str(self.dim_pre) + "]]"
for dim in range(self.dim_kernel - 1):
index += '[' + indices[dim] + '_w]'
increment = self.synapse_type.description['psp']['cpp'] % {
'id_pre': self.pre.id,
'id_post': self.post.id,
'local_index': index,
'global_index': '[i]',
'pre_index': '[rk_pre]',
'post_index': '[rk_post]',
'pre_prefix': 'pop' + str(self.pre.id) + '.',
'post_prefix': 'pop' + str(self.post.id) + '.'
}
# Delays
if self.delays > Global.config['dt']:
increment = increment.replace(
'pop%(id_pre)s.r[rk_pre]' % {'id_pre': self.pre.id},
'delayed_r[rk_pre]')
# Apply the operation
if operation == "sum":
code += tabify(
"""
sum += %(increment)s""" % {'increment': increment}, 1 + dim)
elif operation == "max":
code += tabify(
"""
%(float_prec)s _psp = %(increment)s
if(_psp > sum) sum = _psp;""" % {
'increment': increment,
'float_prec': Global.config['precision']
}, 1 + dim)
elif operation == "min":
code += tabify(
"""
%(float_prec)s _psp = %(increment)s
if(_psp < sum) sum = _psp;""" % {
'increment': increment,
'float_prec': Global.config['precision']
}, 1 + dim)
elif operation == "mean":
code += tabify(
"""
sum += %(increment)s""" % {'increment': increment}, 1 + dim)
else:
Global._error('SharedProjection: Operation', operation,
'is not implemented yet for shared projections.')
# Close for loops
for dim in range(self.dim_kernel - 1):
code += tabify("""
}""", self.dim_kernel - 1 - dim)
impl_code = code % {
'id_proj': self.id,
'target': self.target,
'id_pre': self.pre.id,
'name_pre': self.pre.name,
'size_pre': self.pre.size,
'id_post': self.post.id,
'name_post': self.post.name,
'size_post': self.post.size
}
# sum code
if operation == "mean":
sum_code = """sum/%(filter_size)s""" % {
'filter_size': self.weights.size
}
else:
sum_code = "sum"
return impl_code, sum_code
def _update_spiking_neuron(self, pop):
"""
Update code for the spiking neurons comprise of two parts, update of
ODE and test spiking condition.
"""
id_dict = {
'id': pop.id,
'local_index': "[i]",
'semiglobal_index': '',
'global_index': ''
}
# Is there an axonal spike condition?
if pop.neuron_type.axon_spike:
# get the conditions for axonal and neural spike event
axon_cond = pop.neuron_type.description['axon_spike'][
'spike_cond'] % id_dict
neur_cond = pop.neuron_type.description['spike'][
'spike_cond'] % id_dict
# state changes if axonal spike occur
axon_reset = ""
for eq in pop.neuron_type.description['axon_spike']['spike_reset']:
axon_reset += """
%(reset)s
""" % {
'reset': eq['cpp'] % id_dict
}
# Simply extent the spiking vector, as the axonal spike
# either manipulate neuron state nor consider refractoriness.
# TODO: what happens for plastic networks? As it seems to be unclear,
# after discussion with JB it is currently disabled (HD: 21.01.2019)
axon_spike_code = """
// Axon Spike Event, only if there was not already an event
if( (%(axon_condition)s) && !(%(neur_condition)s) ) {
axonal.push_back(i);
%(axon_reset)s
}
""" % {
'axon_condition': axon_cond,
'neur_condition': neur_cond,
'axon_reset': axon_reset
}
else:
axon_spike_code = ""
# Global variables
global_code = ""
eqs = generate_equation_code(pop.id,
pop.neuron_type.description,
locality='global',
with_refractory=False,
padding=3) % id_dict
if eqs.strip() != "":
global_code += """
// Updating the global variables
#pragma omp single
{
%(eqs)s
}
""" % {
'eqs': eqs
}
# Is there a refractory period?
has_refractory = True if (pop.neuron_type.refractory
or pop.refractory) else False
# Gather pre-loop declaration (dt/tau for ODEs)
pre_code = ""
for var in pop.neuron_type.description['variables']:
if 'pre_loop' in var.keys() and len(var['pre_loop']) > 0:
pre_code += var['ctype'] + ' ' + var['pre_loop'][
'name'] + ' = ' + var['pre_loop']['value'] + ';\n'
if len(pre_code) > 0:
pre_code = """
// Updating the step sizes
""" + tabify(pre_code, 3)
global_code = pre_code % id_dict + global_code
# Local variables, evaluated in parallel
local_code = generate_equation_code(pop.id,
pop.neuron_type.description,
locality='local',
with_refractory=has_refractory,
padding=4) % id_dict
# Decrement of refractoriness
if has_refractory:
local_code += tabify(
"""
// Decrement the refractory period
refractory_remaining[i] -= (1 - in_ref[i]);
""", 4)
# Process the condition
cond = pop.neuron_type.description['spike']['spike_cond'] % id_dict
# Reset equations
reset = ""
for eq in pop.neuron_type.description['spike']['spike_reset']:
reset += """
%(reset)s
""" % {
'reset': eq['cpp'] % id_dict
}
# Mean Firing rate
mean_FR_push, mean_FR_update = self._update_fr(pop)
# Increment of the refractory variable
if has_refractory:
# By default, it is refractory, but users can specify another one
refrac_var = "refractory[i]"
if isinstance(pop.neuron_type.refractory, str):
found = False
for param in pop.neuron_type.description[
"parameters"] + pop.neuron_type.description[
"variables"]:
if param["name"] == pop.neuron_type.refractory:
if param['locality'] == 'local':
refrac_var = "int(" + pop.neuron_type.refractory + "[i]/dt)"
else:
refrac_var = "int(" + pop.neuron_type.refractory + "/dt)"
found = True
break
if not found:
Global._error("refractory = " +
pop.neuron_type.refractory +
": parameter or variable does not exist.")
refrac_inc = "refractory_remaining[i] = %(refrac_var)s;" % {
'refrac_var': refrac_var
}
omp_code = "#pragma omp for" if pop.size > Global.OMP_MIN_NB_NEURONS else ""
comp_inref = """
// compute which neurons are in refractory
#pragma omp for
for (int i = 0; i < size; i++) {
in_ref[i] = (refractory_remaining[i] > 0) ? 0 : 1;
}
""" % {
'omp_code': omp_code
}
else:
refrac_inc = ""
comp_inref = ""
# Gather code
omp_critical_code = "#pragma omp critical" if pop.size > Global.OMP_MIN_NB_NEURONS else ""
refrac_check = tabify(
"""
// check if neuron is in refractory
if (in_ref[i]==0)
continue;
""", 3)
if pop.size > Global.OMP_MIN_NB_NEURONS:
gather_code = """
if( _active ) {
#pragma omp master
{
spiked.clear();
}
auto local_spikes = std::vector<int>();
#pragma omp barrier
#pragma omp for nowait
for (int i = 0; i < size; i++) {
%(refrac_check)s
// Spike emission
if( %(condition)s ) { // Condition is met
// Reset variables
%(reset)s
// Store the spike
local_spikes.push_back(i);
last_spike[i] = t;
// Refractory period
%(refrac_inc)s
%(mean_FR_push)s
}
%(mean_FR_update)s
%(axon_spike_code)s
}
local_spiked_sizes[tid+1] = local_spikes.size();
#pragma omp barrier
#pragma omp single
{
for (int i = 1; i < (num_threads+1); i++) {
local_spiked_sizes[i] += local_spiked_sizes[i-1];
}
spiked.resize(spiked.size()+local_spiked_sizes[num_threads]);
}
std::copy(local_spikes.begin(), local_spikes.end(), spiked.begin() + local_spiked_sizes[tid]);
} // active
"""
else:
gather_code = """
if( _active ) {
#pragma omp barrier
#pragma omp single
{
spiked.clear();
for (int i = 0; i < size; i++) {
%(refrac_check)s
// Spike emission
if( %(condition)s ) { // Condition is met
// Reset variables
%(reset)s
// Store the spike
spiked.push_back(i);
last_spike[i] = t;
// Refractory period
%(refrac_inc)s
%(mean_FR_push)s
}
%(mean_FR_update)s
%(axon_spike_code)s
}
}
} // active
"""
final_spike_gather = gather_code % {
'condition': cond,
'refrac_check': refrac_check if has_refractory else "",
'reset': reset,
'refrac_inc': refrac_inc,
'mean_FR_push': mean_FR_push,
'mean_FR_update': mean_FR_update,
'omp_critical_code': omp_critical_code,
'axon_spike_code': axon_spike_code
}
# If axonal events are defined
if pop.neuron_type.axon_spike:
global_code = "axonal.clear();\n" + global_code
# finish code
final_eq = """
if( _active ) {
%(comp_inref)s
%(global_code)s
// Updating local variables
#pragma omp for simd
for (int i = 0; i < size; i++) {
%(local_code)s
}
} // active
""" % {
'comp_inref': comp_inref,
'local_code': local_code,
'global_code': global_code,
}
# if profiling enabled, annotate with profiling code
if self._prof_gen:
final_eq = self._prof_gen.annotate_update_neuron(pop, final_eq)
final_spike_gather = self._prof_gen.annotate_spike_cond(
pop, final_spike_gather)
return final_eq, final_spike_gather
def header_struct(self, pop, annarchy_dir):
"""
Specialized implementation of PopulationGenerator.header_struct() for
generation of an openMP header.
two passes:
* generate the codes for population header
* fill the dictionary with call codes (return)
"""
self._templates = deepcopy(openmp_templates)
# Generate declaration and accessors of all parameters and variables
declaration_parameters_variables, access_parameters_variables = self._generate_decl_and_acc(
pop)
# Additional includes and structures
include_additional = ""
access_additional = ""
struct_additional = ""
declare_additional = ""
init_additional = ""
reset_additional = ""
# Declare global operations as extern at the beginning of the file
extern_global_operations = ""
for op in pop.global_operations:
extern_global_operations += global_op_extern_dict[
op['function']] % {
'type': Global.config['precision']
}
# Initialize parameters and variables
init_parameters_variables = self._init_population(pop)
# Spike-specific stuff
reset_spike = ""
declare_spike = ""
init_spike = ""
if pop.neuron_type.description['type'] == 'spike':
spike_specific_tpl = self._templates['spike_specific']
# Main data for spiking pops
declare_spike += spike_specific_tpl['spike']['declare'] % {
'id': pop.id
}
init_spike += spike_specific_tpl['spike']['init'] % {'id': pop.id}
reset_spike += spike_specific_tpl['spike']['reset'] % {
'id': pop.id
}
# If there is a refractory period
if pop.neuron_type.refractory or pop.refractory:
declare_spike += spike_specific_tpl['refractory']['declare'] % {
'id': pop.id
}
if isinstance(pop.neuron_type.description['refractory'],
str): # no need to instantiate refractory
init_spike += spike_specific_tpl['refractory'][
'init_extern'] % {
'id': pop.id
}
else:
init_spike += spike_specific_tpl['refractory']['init'] % {
'id': pop.id
}
reset_spike += spike_specific_tpl['refractory']['reset'] % {
'id': pop.id
}
# If axonal spike condition was defined
if pop.neuron_type.axon_spike:
declare_spike += spike_specific_tpl['axon_spike']['declare']
init_spike += spike_specific_tpl['axon_spike']['init']
reset_spike += spike_specific_tpl['axon_spike']['reset']
# Process eventual delay
declare_delay = ""
init_delay = ""
update_delay = ""
update_max_delay = ""
reset_delay = ""
if pop.max_delay > 1:
declare_delay, init_delay, update_delay, update_max_delay, reset_delay = self._delay_code(
pop)
# Process mean FR computations
declare_FR, init_FR, reset_FR = self._init_fr(pop)
reset_spike += reset_FR
# Init rng_dist
init_rng_dist, _ = self._init_random_dist(pop)
# Update random distributions
update_rng = self._update_random_distributions(pop)
# Update global operations
update_global_ops = self._update_globalops(pop)
# Defintion of local functions
declaration_parameters_variables += self._local_functions(pop)
# Update the neural variables
if pop.neuron_type.type == 'rate':
update_variables = self._update_rate_neuron(pop)
test_spike_cond = ""
else:
update_variables, test_spike_cond = self._update_spiking_neuron(
pop)
# Stop condition
stop_condition = self._stop_condition(pop)
# Memory management
determine_size_in_bytes = self._determine_size_in_bytes(pop)
clear_container = self._clear_container(pop)
# Profiling
if self._prof_gen:
include_profile = """#include "Profiling.h"\n"""
declare_profile, init_profile = self._prof_gen.generate_init_population(
pop)
else:
include_profile = ""
init_profile = ""
declare_profile = ""
## When everything is generated, we override the fields defined by the specific population
if 'include_additional' in pop._specific_template.keys():
include_additional = pop._specific_template['include_additional']
if 'struct_additional' in pop._specific_template.keys():
struct_additional = pop._specific_template['struct_additional']
if 'extern_global_operations' in pop._specific_template.keys():
extern_global_operations = pop._specific_template[
'extern_global_operations']
if 'declare_spike_arrays' in pop._specific_template.keys():
declare_spike = pop._specific_template['declare_spike_arrays']
if 'declare_parameters_variables' in pop._specific_template.keys():
declaration_parameters_variables = pop._specific_template[
'declare_parameters_variables']
if 'declare_additional' in pop._specific_template.keys():
declare_additional = pop._specific_template['declare_additional']
if 'declare_FR' in pop._specific_template.keys():
declare_FR = pop._specific_template['declare_FR']
if 'declare_delay' in pop._specific_template.keys(
) and pop.max_delay > 1:
declare_delay = pop._specific_template['declare_delay']
if 'access_parameters_variables' in pop._specific_template.keys():
access_parameters_variables = pop._specific_template[
'access_parameters_variables']
if 'access_additional' in pop._specific_template.keys():
access_additional = pop._specific_template['access_additional']
if 'init_parameters_variables' in pop._specific_template.keys():
init_parameters_variables = pop._specific_template[
'init_parameters_variables']
if 'init_spike' in pop._specific_template.keys():
init_spike = pop._specific_template['init_spike']
if 'init_delay' in pop._specific_template.keys() and pop.max_delay > 1:
init_delay = pop._specific_template['init_delay']
if 'init_FR' in pop._specific_template.keys():
init_FR = pop._specific_template['init_FR']
if 'init_additional' in pop._specific_template.keys():
init_additional = pop._specific_template['init_additional']
if 'reset_spike' in pop._specific_template.keys():
reset_spike = pop._specific_template['reset_spike']
if 'reset_delay' in pop._specific_template.keys(
) and pop.max_delay > 1:
reset_delay = pop._specific_template['reset_delay']
if 'reset_additional' in pop._specific_template.keys():
reset_additional = pop._specific_template['reset_additional']
if 'update_variables' in pop._specific_template.keys():
update_variables = pop._specific_template['update_variables']
if 'test_spike_cond' in pop._specific_template.keys():
test_spike_cond = pop._specific_template['test_spike_cond']
if 'update_rng' in pop._specific_template.keys():
update_rng = pop._specific_template['update_rng']
if 'update_delay' in pop._specific_template.keys(
) and pop.max_delay > 1:
update_delay = pop._specific_template['update_delay']
if 'update_max_delay' in pop._specific_template.keys(
) and pop.max_delay > 1:
update_max_delay = pop._specific_template['update_max_delay']
if 'update_global_ops' in pop._specific_template.keys():
update_global_ops = pop._specific_template['update_global_ops']
# Fill the template
code = self._templates['population_header'] % {
# version tag
'annarchy_version': ANNarchy.__release__,
#'time_stamp': '{:%Y-%b-%d %H:%M:%S}'.format(datetime.datetime.now()),
# fill code templates
'float_prec': Global.config['precision'],
'id': pop.id,
'name': pop.name,
'size': pop.size,
'include_additional': include_additional,
'include_profile': include_profile,
'struct_additional': struct_additional,
'extern_global_operations': extern_global_operations,
'declare_spike_arrays': declare_spike,
'declare_parameters_variables': declaration_parameters_variables,
'declare_additional': declare_additional,
'declare_delay': declare_delay,
'declare_FR': declare_FR,
'declare_profile': declare_profile,
'access_parameters_variables': access_parameters_variables,
'access_additional': access_additional,
'init_parameters_variables': init_parameters_variables,
'init_spike': init_spike,
'init_delay': init_delay,
'init_FR': init_FR,
'init_additional': init_additional,
'init_rng_dist': init_rng_dist,
'init_profile': init_profile,
'reset_spike': reset_spike,
'reset_delay': reset_delay,
'reset_additional': reset_additional,
'update_variables': update_variables,
'test_spike_cond': test_spike_cond,
'update_rng': update_rng,
'update_delay': update_delay,
'update_max_delay': update_max_delay,
'update_global_ops': update_global_ops,
'stop_condition': stop_condition,
'determine_size': determine_size_in_bytes,
'clear_container': clear_container
}
# remove right-trailing spaces
code = remove_trailing_spaces(code)
# Store the complete header definition in a single file
with open(
annarchy_dir + '/generate/net' + str(self._net_id) + '/pop' +
str(pop.id) + '.hpp', 'w') as ofile:
ofile.write(code)
# Basic informations common to all populations
pop_desc = {
'include': """#include "pop%(id)s.hpp"\n""" % {
'id': pop.id
},
'extern': """extern PopStruct%(id)s pop%(id)s;\n""" % {
'id': pop.id
},
'instance': """PopStruct%(id)s pop%(id)s;\n""" % {
'id': pop.id
},
'init': """ pop%(id)s.init_population();\n""" % {
'id': pop.id
}
}
# Generate the calls to be made in the main ANNarchy.cpp
if len(pop.neuron_type.description['variables']
) > 0 or 'update_variables' in pop._specific_template.keys():
if update_variables != "":
pop_desc['update'] = """ pop%(id)s.update(tid);\n""" % {
'id': pop.id
}
if pop.neuron_type.type == "spike":
pop_desc[
'update'] += """ pop%(id)s.spike_gather(tid, nt);\n""" % {
'id': pop.id
}
if len(pop.neuron_type.description['random_distributions']) > 0:
pop_desc['rng_update'] = """ pop%(id)s.update_rng(tid);\n""" % {
'id': pop.id
}
if pop.max_delay > 1:
pop_desc['delay_update'] = tabify(
"""pop%(id)s.update_delay();\n""" % {'id': pop.id}, 1)
if len(pop.global_operations) > 0:
pop_desc[
'gops_update'] = """ pop%(id)s.update_global_ops(tid, nt);\n""" % {
'id': pop.id
}
return pop_desc
def _init_parameters_variables(self, proj, single_spmv_matrix):
"""
Generate initialization code for variables / parameters of the
projection *proj*.
Returns 3 values:
ret1 (str): weight initialization
ret2 (str): delay initialization
ret3 (str): other initializations (e. g. event-driven)
"""
# Is it a specific projection?
if 'init_parameters_variables' in proj._specific_template.keys():
return proj._specific_template['init_parameters_variables']
# Learning by default
code = ""
weight_code = ""
# choose initialization templates based on chosen paradigm
attr_init_tpl = self._templates['attribute_cpp_init']
attributes = []
# Initialize parameters
for var in proj.synapse_type.description['parameters'] + proj.synapse_type.description['variables']:
# Avoid doublons
if var['name'] in attributes:
continue
# Important to select which template
locality = var['locality']
attr_type = 'parameter' if var in proj.synapse_type.description['parameters'] else 'variable'
# The synaptic weight
if var['name'] == 'w':
if var['locality'] == "global" or proj._has_single_weight():
if cpp_connector_available(proj.connector_name, proj._storage_format, proj._storage_order):
weight_code = tabify("w = w_dist_arg1;", 2)
else:
weight_code = tabify("w = values[0][0];", 2)
elif var['locality'] == "local":
if cpp_connector_available(proj.connector_name, proj._storage_format, proj._storage_order): # Init weights in CPP
if proj.connector_weight_dist == None:
init_code = self._templates['attribute_cpp_init']['local'] % {
'init': 'w_dist_arg1',
'type': var['ctype'],
'attr_type': 'parameter' if var in proj.synapse_type.description['parameters'] else 'variable',
'name': var['name']
}
elif isinstance(proj.connector_weight_dist, ANNRandom.Uniform):
if single_spmv_matrix:
init_code = "w = init_matrix_variable_uniform<%(float_prec)s>(w_dist_arg1, w_dist_arg2, rng[0]);"
else:
init_code = "w = init_matrix_variable_uniform<%(float_prec)s>(w_dist_arg1, w_dist_arg2, rng);"
elif isinstance(proj.connector_weight_dist, ANNRandom.Normal):
if single_spmv_matrix:
init_code = "w = init_matrix_variable_normal<%(float_prec)s>(w_dist_arg1, w_dist_arg2, rng[0]);"
else:
init_code = "w = init_matrix_variable_normal<%(float_prec)s>(w_dist_arg1, w_dist_arg2, rng);"
elif isinstance(proj.connector_weight_dist, ANNRandom.LogNormal):
if proj.connector_weight_dist.min==None and proj.connector_weight_dist.max==None:
if single_spmv_matrix:
init_code = "w = init_matrix_variable_log_normal<%(float_prec)s>(w_dist_arg1, w_dist_arg2, rng[0]);"
else:
init_code = "w = init_matrix_variable_log_normal<%(float_prec)s>(w_dist_arg1, w_dist_arg2, rng);"
else:
min_code = "std::numeric_limits<%(float_prec)s>::min()" if proj.connector_weight_dist.min==None else str(proj.connector_weight_dist.min)
max_code = "std::numeric_limits<%(float_prec)s>::max()" if proj.connector_weight_dist.max==None else str(proj.connector_weight_dist.max)
if single_spmv_matrix:
init_code = "w = init_matrix_variable_log_normal_clip<%(float_prec)s>(w_dist_arg1, w_dist_arg2, rng[0], "+min_code+", "+max_code+");"
else:
init_code = "w = init_matrix_variable_log_normal_clip<%(float_prec)s>(w_dist_arg1, w_dist_arg2, rng, "+min_code+", "+max_code+");"
else:
raise NotImplementedError( str(type(proj.connector_weight_dist)) + " is not available for CPP-side connection patterns.")
if Global._check_paradigm("cuda"):
init_code += "\ngpu_w = init_matrix_variable_gpu<%(float_prec)s>(w);"
weight_code = tabify(init_code % {'float_prec': Global.config['precision']}, 2)
# Init_from_lil
else:
init = 'false' if var['ctype'] == 'bool' else ('0' if var['ctype'] == 'int' else '0.0')
weight_code = attr_init_tpl[locality] % {
'id': proj.id,
'id_post': proj.post.id,
'name': var['name'],
'type': var['ctype'],
'init': init,
'attr_type': attr_type,
'float_prec': Global.config['precision']
}
weight_code += tabify("update_matrix_variable_all<%(float_prec)s>(w, values);" % {'float_prec': Global.config['precision']}, 2)
if Global._check_paradigm("cuda"):
weight_code += tabify("\nw_host_to_device = true;", 2)
else:
raise NotImplementedError
# All other variables
else:
init = 'false' if var['ctype'] == 'bool' else ('0' if var['ctype'] == 'int' else '0.0')
var_ids = {
'id': proj.id,
'id_post': proj.post.id,
'name': var['name'],
'type': var['ctype'],
'init': init,
'attr_type': attr_type,
'float_prec': Global.config['precision']
}
if Global._check_paradigm("cuda") and locality == "global":
code += attr_init_tpl[locality][attr_type] % var_ids
else:
code += attr_init_tpl[locality] % var_ids
attributes.append(var['name'])
# Initialize delays differs for construction from LIL or CPP inited patterns
if proj.max_delay > 1:
# Special case: we have non-uniform delays, but not determined by a RandomDistribution
# This will caused most likely by custom connectivity pattern
if proj.connector_delay_dist == None and proj.uniform_delay==-1:
id_pre = proj.pre.id if not isinstance(proj.pre, PopulationView) else proj.pre.population.id
if proj.synapse_type.type == "rate":
delay_code = self._templates['delay']['nonuniform_rate_coded']['init'] % {'id_pre': id_pre}
else:
delay_code = self._templates['delay']['nonuniform_spiking']['init'] % {'id_pre': id_pre}
#
# uniform delay
elif proj.connector_delay_dist == None:
if cpp_connector_available(proj.connector_name, proj._storage_format, proj._storage_order):
delay_code = tabify("delay = d_dist_arg1;", 2)
else:
delay_code = self._templates['delay']['uniform']['init']
#
# non-uniform delay drawn from distribution
elif isinstance(proj.connector_delay_dist, ANNRandom.RandomDistribution):
if cpp_connector_available(proj.connector_name, proj._storage_format, proj._storage_order):
rng_init = "rng[0]" if single_spmv_matrix else "rng"
delay_code = tabify("""
delay = init_matrix_variable_discrete_uniform<int>(d_dist_arg1, d_dist_arg2, %(rng_init)s);
max_delay = -1;""" % {'id_pre': proj.pre.id, 'rng_init': rng_init}, 2)
else:
id_pre = proj.pre.id if not isinstance(proj.pre, PopulationView) else proj.pre.population.id
if proj.synapse_type.type == "rate":
delay_code = self._templates['delay']['nonuniform_rate_coded']['init'] % {'id_pre': id_pre}
else:
delay_code = self._templates['delay']['nonuniform_spiking']['init'] % {'id_pre': id_pre}
else:
raise NotImplementedError( str(type(proj.connector_weight_dist)) + " is not available.")
else:
delay_code = ""
# If no psp is defined, it's event-driven
has_event_driven = False
for var in proj.synapse_type.description['variables']:
if var['method'] == 'event-driven':
has_event_driven = True
break
if has_event_driven:
code += self._templates['event_driven']['cpp_init']
# Pruning
if Global.config['structural_plasticity']:
if 'pruning' in proj.synapse_type.description.keys():
code += """
// Pruning
_pruning = false;
_pruning_period = 1;
_pruning_offset = 0;
"""
if 'creating' in proj.synapse_type.description.keys():
code += """
// Creating
_creating = false;
_creating_period = 1;
_creating_offset = 0;
"""
return weight_code, delay_code, code
def _pop_recorder_class(self, pop):
"""
Creates population recording class code.
Returns:
* complete code as string
Templates:
omp_population, cuda_population
"""
if Global.config['paradigm'] == "openmp":
template = RecTemplate.omp_population
elif Global.config['paradigm'] == "cuda":
template = RecTemplate.cuda_population
else:
raise NotImplementedError
tpl_code = template['template']
init_code = ""
recording_code = ""
recording_target_code = ""
struct_code = ""
determine_size = ""
clear_code = ""
# The post-synaptic potential for rate-code (weighted sum) as well
# as the conductance variables are handled seperatly.
target_list = []
targets = []
for t in pop.neuron_type.description['targets']:
if isinstance(t, list):
for t2 in t:
targets.append(t2)
else:
targets.append(t)
for t in pop.targets:
if isinstance(t, list):
for t2 in t:
targets.append(t2)
else:
targets.append(t)
targets = sorted(list(set(targets)))
if pop.neuron_type.type == 'rate':
for target in targets:
tar_dict = {'id': pop.id, 'type' : Global.config['precision'], 'name': '_sum_'+target}
struct_code += template['local']['struct'] % tar_dict
init_code += template['local']['init'] % tar_dict
recording_target_code += template['local']['recording'] % tar_dict
else:
for target in targets:
tar_dict = {'id': pop.id, 'type' : Global.config['precision'], 'name': 'g_'+target}
struct_code += template['local']['struct'] % tar_dict
init_code += template['local']['init'] % tar_dict
recording_target_code += template['local']['recording'] % tar_dict
# to skip this entry in the following loop
target_list.append('g_'+target)
# Record global and local attributes
attributes = []
for var in pop.neuron_type.description['variables']:
# Skip targets
if var['name'] in target_list:
continue
# Avoid doublons
if var['name'] in attributes:
continue
attributes.append(var['name'])
ids = {
'id': pop.id,
'name': var['name'],
'type': var['ctype']
}
struct_code += template[var['locality']]['struct'] % ids
init_code += template[var['locality']]['init'] % ids
recording_code += template[var['locality']]['recording'] % ids
clear_code += template[var['locality']]['clear'] % ids
# Memory management
if var['locality'] == "global":
determine_size += """
// global variable %(name)s
size_in_bytes += sizeof(%(type)s);
""" % ids
elif var['locality'] == "semiglobal":
determine_size += """
// semiglobal variable %(name)s
size_in_bytes += sizeof(%(type)s) * %(name)s.capacity();
""" % ids
else:
determine_size += """
// local variable %(name)s
size_in_bytes += sizeof(std::vector<%(type)s>) * %(name)s.capacity();
for(auto it=%(name)s.begin(); it!= %(name)s.end(); it++) {
size_in_bytes += it->capacity() * sizeof(%(type)s);
}
""" % ids
# Record spike events
if pop.neuron_type.type == 'spike':
base_tpl = RecTemplate.recording_spike_tpl
rec_dict = {
'id': pop.id,
'type' : 'long int',
'name': 'spike',
'rec_target': 'spiked'
}
struct_code += base_tpl['struct'] % rec_dict
init_code += base_tpl['init'] % rec_dict
recording_code += base_tpl['record'][Global.config['paradigm']] % rec_dict
determine_size += base_tpl['size_in_bytes'][Global.config['paradigm']] % rec_dict
clear_code += base_tpl['clear'][Global.config['paradigm']] % rec_dict
# Record axon spike events
if pop.neuron_type.axon_spike:
rec_dict = {
'id': pop.id,
'type' : 'long int',
'name': 'axon_spike',
'rec_target': 'axonal'
}
struct_code += base_tpl['struct'] % rec_dict
init_code += base_tpl['init'] % rec_dict
recording_code += base_tpl['record'][Global.config['paradigm']] % rec_dict
ids = {
'id': pop.id,
'init_code': init_code,
'struct_code': struct_code,
'recording_code': recording_code,
'recording_target_code': recording_target_code,
'determine_size': tabify(determine_size, 2),
'clear_monitor_code': clear_code
}
return tpl_code % ids
def _update_spiking_neuron(self, pop):
# Neural update
from ANNarchy.generator.Utils import generate_equation_code, tabify
# Is there a refractory period?
if pop.neuron_type.refractory or pop.refractory:
# Get the equations
eqs = generate_equation_code(pop.id,
pop.neuron_type.description,
'local',
conductance_only=True,
padding=4) % {
'id': pop.id,
'local_index': "[i]",
'semiglobal_index': '',
'global_index': ''
}
# Generate the code snippet
code = """
// Refractory period
if( refractory_remaining[i] > 0){
%(eqs)s
// Decrement the refractory period
refractory_remaining[i]--;
continue;
}
""" % {
'eqs': eqs
}
refrac_inc = "refractory_remaining[i] = refractory[i];"
else:
code = ""
refrac_inc = ""
# Global variables
global_code = ""
eqs = generate_equation_code(
pop.id, pop.neuron_type.description, 'global', padding=3) % {
'id': pop.id,
'local_index': "[i]",
'semiglobal_index': '',
'global_index': ''
}
if eqs.strip() != "":
global_code += """
// Updating the global variables
%(eqs)s
""" % {
'eqs': eqs
}
# Gather pre-loop declaration (dt/tau for ODEs)
pre_code = ""
for var in pop.neuron_type.description['variables']:
if 'pre_loop' in var.keys() and len(var['pre_loop']) > 0:
pre_code += var['ctype'] + ' ' + var['pre_loop'][
'name'] + ' = ' + var['pre_loop']['value'] + ';\n'
if len(pre_code) > 0:
pre_code = """
// Updating the step sizes
""" + tabify(pre_code, 3)
global_code = pre_code % {
'id': pop.id,
'local_index': "[i]",
'semiglobal_index': '',
'global_index': ''
} + global_code
# OMP code
omp_code = "#pragma omp parallel for" if (
Global.config['num_threads'] > 1
and pop.size > Global.OMP_MIN_NB_NEURONS) else ""
omp_critical_code = "#pragma omp critical" if (
Global.config['num_threads'] > 1
and pop.size > Global.OMP_MIN_NB_NEURONS) else ""
# Local variables, evaluated in parallel
code += generate_equation_code(
pop.id, pop.neuron_type.description, 'local', padding=4) % {
'id': pop.id,
'local_index': "[i]",
'semiglobal_index': '',
'global_index': ''
}
# Process the condition
cond = pop.neuron_type.description['spike']['spike_cond'] % {
'id': pop.id,
'local_index': "[i]",
'semiglobal_index': '',
'global_index': ''
}
# Reset equations
reset = ""
for eq in pop.neuron_type.description['spike']['spike_reset']:
reset += """
%(reset)s
""" % {
'reset': eq['cpp'] % {
'id': pop.id,
'local_index': "[i]",
'semiglobal_index': '',
'global_index': ''
}
}
# Mean Firing rate
mean_FR_push, mean_FR_update = self._update_fr(pop)
# Gather code
spike_gather = """
// Spike emission
if(%(condition)s){ // Condition is met
// Reset variables
%(reset)s
// Store the spike
%(omp_critical_code)s
{
spiked.push_back(i);
}
last_spike[i] = t;
// Refractory period
%(refrac_inc)s
%(mean_FR_push)s
}
%(mean_FR_update)s
""" % {
'condition': cond,
'reset': reset,
'refrac_inc': refrac_inc,
'mean_FR_push': mean_FR_push,
'mean_FR_update': mean_FR_update,
'omp_critical_code': omp_critical_code
}
code += spike_gather
# finish code
final_code = """
if( _active ) {
spiked.clear();
%(global_code)s
// Updating local variables
%(omp_code)s
for(int i = 0; i < size; i++){
%(code)s
}
} // active
""" % {
'code': code,
'global_code': global_code,
'omp_code': omp_code
}
# if profiling enabled, annotate with profiling code
if self._prof_gen:
final_code = self._prof_gen.annotate_update_neuron(pop, final_code)
return final_code
