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 _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_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 _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 _update_synapse(self, proj):
"""
Generate the device codes for synaptic equations. As the parallel
evaluation of local and global equations within one kernel would require
a __syncthread() call, we split up the implementation into two seperate
parts.
Return:
* a tuple contain three strings ( body, call, header )
"""
# Global variables
global_eq = generate_equation_code(proj.id, proj.synapse_type.description, 'global', 'proj', padding=1, wrap_w="plasticity")
# Semiglobal variables
semiglobal_eq = generate_equation_code(proj.id, proj.synapse_type.description, 'semiglobal', 'proj', padding=2, wrap_w="plasticity")
# Local variables
local_eq = generate_equation_code(proj.id, proj.synapse_type.description, 'local', 'proj', padding=2, wrap_w="plasticity")
# Something to do?
if global_eq.strip() == '' and semiglobal_eq.strip() == '' and local_eq.strip() == '':
return "", "", ""
# Dictionary of pre/suffixes
ids = {
'id_proj' : proj.id,
'target': proj.target,
'id_post': proj.post.id,
'id_pre': proj.pre.id,
'local_index': '[j]',
'semiglobal_index': '[rk_post]',
'global_index': '[0]',
'pre_index': '[rk_pre]',
'post_index': '[rk_post]',
'pre_prefix': 'pre_',
'post_prefix': 'post_',
'delay_nu' : '[delay[i][j]-1]', # non-uniform delay
'delay_u' : '[' + str(proj.uniform_delay-1) + ']', # uniform delay
}
body = ""
header = ""
local_call = ""
global_call = ""
semiglobal_call = ""
#
# Fill code templates for global, semiglobal and local equations
#
if global_eq.strip() != '':
global_eq, global_pre_code, kernel_args_global, kernel_args_call_global = self._process_equations( proj, global_eq, ids, 'global' )
if semiglobal_eq.strip() != '':
semiglobal_eq, semiglobal_pre_code, kernel_args_semiglobal, kernel_args_call_semiglobal = self._process_equations( proj, semiglobal_eq, ids, 'semiglobal' )
if local_eq.strip() != '':
local_eq, local_pre_code, kernel_args_local, kernel_args_call_local = self._process_equations( proj, local_eq, ids, 'local' )
#
# replace local function calls
if len(proj.synapse_type.description['functions']) > 0:
global_eq, semiglobal_eq, local_eq = self._replace_local_funcs(proj, global_eq, semiglobal_eq, local_eq)
# replace the random distributions
local_eq, global_eq = self._replace_random(local_eq, global_eq, proj.synapse_type.description['random_distributions'])
if global_eq.strip() != '':
body += self._templates['synapse_update']['global']['body'] % {
'id': proj.id,
'kernel_args': kernel_args_global,
'global_eqs': global_eq,
'target': proj.target,
'pre': proj.pre.id,
'post': proj.post.id,
'pre_loop': global_pre_code,
'float_prec': Global.config['precision']
}
header += self._templates['synapse_update']['global']['header'] % {
'id': proj.id,
'kernel_args': kernel_args_global,
'float_prec': Global.config['precision']
}
global_call = ""
target_list = proj.target if isinstance(proj.target, list) else [proj.target]
for target in target_list:
global_call += self._templates['synapse_update']['global']['call'] % {
'id_proj': proj.id,
'post': proj.post.id,
'pre': proj.pre.id,
'target': target,
'kernel_args_call': kernel_args_call_global,
'float_prec': Global.config['precision']
}
if semiglobal_eq.strip() != '':
body += self._templates['synapse_update']['semiglobal']['body'] % {
'id': proj.id,
'kernel_args': kernel_args_semiglobal,
'semiglobal_eqs': semiglobal_eq,
'target': proj.target,
'pre': proj.pre.id,
'post': proj.post.id,
'pre_loop': semiglobal_pre_code,
'float_prec': Global.config['precision']
}
header += self._templates['synapse_update']['semiglobal']['header'] % {
'id': proj.id,
'kernel_args': kernel_args_semiglobal,
'float_prec': Global.config['precision']
}
semiglobal_call = ""
target_list = proj.target if isinstance(proj.target, list) else [proj.target]
for target in target_list:
semiglobal_call += self._templates['synapse_update']['semiglobal']['call'] % {
'id_proj': proj.id,
'id_post': proj.post.id,
'id_pre': proj.pre.id,
'target': target,
'kernel_args_call': kernel_args_call_semiglobal,
'float_prec': Global.config['precision']
}
if local_eq.strip() != '':
body += self._templates['synapse_update']['local']['body'] % {
'id': proj.id,
'kernel_args': kernel_args_local,
'local_eqs': local_eq,
'target': proj.target,
'pre': proj.pre.id,
'post': proj.post.id,
'pre_loop': local_pre_code,
'float_prec': Global.config['precision']
}
header += self._templates['synapse_update']['local']['header'] % {
'id': proj.id,
'kernel_args': kernel_args_local,
'float_prec': Global.config['precision']
}
local_call = ""
target_list = proj.target if isinstance(proj.target, list) else [proj.target]
for target in target_list:
local_call += self._templates['synapse_update']['local']['call'] % {
'id_proj': proj.id,
'id_post': proj.post.id,
'id_pre': proj.pre.id,
'target': target,
'kernel_args_call': kernel_args_call_local,
'float_prec': Global.config['precision']
}
call = self._templates['synapse_update']['call'] % {
'id_proj': proj.id,
'post': proj.post.id,
'pre': proj.pre.id,
'target': proj.target,
'global_call': global_call,
'semiglobal_call': semiglobal_call,
'local_call': local_call,
'float_prec': Global.config['precision']
}
# Profiling
if self._prof_gen:
call = self._prof_gen.annotate_update_synapse(proj, call)
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 _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 _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