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