GeNN  3.3.0
GPU enhanced Neuronal Networks (GeNN)
Best practices guide

GeNN generates code according to the network model defined by the user, and allows users to include the generated code in their programs as they want. Here we provide a guideline to setup GeNN and use generated functions. We recommend users to also have a look at the Examples, and to follow the tutorials Tutorial 1 and Tutorial 2.

Creating and simulating a network model

The user is first expected to create an object of class NNmodel by creating the function modelDefinition() which includes calls to following methods in correct order:

Then add neuron populations by:

for each neuron population. Add synapse populations by:

for each synapse population.

The modelDefinition() needs to end with calling NNmodel::finalize().

Other optional functions are explained in NNmodel class reference. At the end the function should look like this:

void modelDefinition(NNModel &model) {
model.setDT(0.5);
model.setName("YourModelName");
model.addNeuronPopulation(...);
...
model.addSynapsePopulation(...);
...
model.finalize();
}

modelSpec.h should be included in the file where this function is defined.

This function will be called by generateALL.cc to create corresponding CPU and GPU simulation codes under the <YourModelName>_CODE directory.

These functions can then be used in a .cc file which runs the simulation. This file should include <YourModelName>_CODE/definitions.h. Generated code differ from one model to the other, but core functions are the same and they should be called in correct order. First, the following variables should be defined and initialized:

  • NNmodel model // initialized by calling modelDefinition(model)
  • Array containing current input (if any)

The following are declared by GeNN but should be initialized by the user:

  • Poisson neuron offset and rates (if any)
  • Connectivity matrices (if sparse)
  • Neuron and synapse variables (if not initialising to the homogeneous initial value provided during modelDefinition)

Core functions generated by GeNN to be included in the user code include:

  • allocateMem()
  • deviceMemAllocate()
  • initialize()
  • init<model name>()
  • push<neuron or synapse name>StateToDevice()
  • pull<neuron or synapse name>StateFromDevice()
  • push<neuron name>SpikesToDevice()
  • pull<neuron name>SpikesFromDevice()
  • push<neuron name>SpikesEventsToDevice()
  • pull<neuron name>SpikesEventsFromDevice()
  • push<neuron name>CurrentSpikesToDevice()
  • pull<neuron name>CurrentSpikesFromDevice()
  • push<neuron name>CurrentSpikesEventsToDevice()
  • pull<neuron name>CurrentSpikesEventsFromDevice()
  • copyStateToDevice()
  • copyStateFromDevice()
  • copySpikesToDevice()
  • copySpikesFromDevice()
  • copySpikesEventsToDevice()
  • copySpikesEventsFromDevice()
  • copyCurrentSpikesToDevice()
  • copyCurrentSpikesFromDevice()
  • copyCurrentSpikesEventsToDevice()
  • copyCurrentSpikesEventsFromDevice()
  • stepTimeCPU()
  • stepTimeGPU()
  • freeMem()

Before calling the kernels, make sure you have copied the initial values of any neuron and synapse variables initialised on the host to the GPU. You can use the push\<neuron or synapse name\>StateToDevice() to copy from the host to the GPU. At the end of your simulation, if you want to access the variables you need to copy them back from the device using the pull\<neuron or synapse name\>StateFromDevice() function or one of the more fine-grained functions listed above. Alternatively, you can directly use the CUDA memcopy functions. Copying elements between the GPU and the host memory is very costly in terms of performance and should only be done when needed.

Floating point precision

Double precision floating point numbers are supported by devices with compute capability 1.3 or higher. If you have an older GPU, you need to use single precision floating point in your models and simulation.

GPUs are designed to work better with single precision while double precision is the standard for CPUs. This difference should be kept in mind while comparing performance.

While setting up the network for GeNN, double precision floating point numbers are used as this part is done on the CPU. For the simulation, GeNN lets users choose between single or double precision. Overall, new variables in the generated code are defined with the precision specified by NNmodel::setPrecision(unsigned int), providing GENN_FLOAT or GENN_DOUBLE as argument. GENN_FLOAT is the default value. The keyword scalar can be used in the user-defined model codes for a variable that could either be single or double precision. This keyword is detected at code generation and substituted with "float" or "double" according to the precision set by NNmodel::setPrecision(unsigned int).

There may be ambiguities in arithmetic operations using explicit numbers. Standard C compilers presume that any number defined as "X" is an integer and any number defined as "X.Y" is a double. Make sure to use the same precision in your operations in order to avoid performance loss.

Working with variables in GeNN

Model variables

User-defined model variables originate from classes derived off the NeuronModels::Base, WeightUpdateModels::Base or PostsynapticModels::Base classes. The name of model variable is defined in the model type, i.e. with a statement such as

SET_VARS({{"V", "scalar"}});

When a neuron or synapse population using this model is added to the model, the full GeNN name of the variable will be obtained by concatenating the variable name with the name of the population. For example if we a add a population called Pop using a model which contains our V variable, a variable VPop of type scalar* will be available in the global namespace of the simulation program. GeNN will pre-allocate this C array to the correct size of elements corresponding to the size of the neuron population. GeNN will also free these variables when the provided function freeMem() is called. Users can otherwise manipulate these variable arrays as they wish. For convenience, GeNN provides functions pullXXStatefromDevice() and pushXXStatetoDevice() to copy the variables associated to a neuron population XX from the device into host memory and vice versa. E.g.

pullPopStateFromDevice();

would copy the C array VPop from device memory into host memory (and any other variables that the population Pop may have).

The user can also directly use CUDA memory copy commands independent of the provided convenience functions. The relevant device pointers for all variables that exist in host memory have the same name as the host variable but are prefixed with d_. For example, the copy command that would be contained in pullPopStateFromDevice() will look like

unsigned int size = sizeof(scalar) * nPop;
cudaMemcpy(VPop, d_VPop, size, cudaMemcpyDeviceToHost);

where nPop is an integer containing the population size of the Pop population.

Thes conventions also apply to the the variables of postsynaptic and weight update models.

Note
Be aware that the above naming conventions do assume that variables from the weightupdate models and the postSynModels that are used together in a synapse population are unique. If both the weightupdate model and the postSynModel have a variable of the same name, the behaviour is undefined.

Built-in Variables in GeNN

GeNN has no explicitly hard-coded synapse and neuron variables. Users are free to name the variable of their models as they want. However, there are some reserved variables that are used for intermediary calculations and communication between different parts of the generated code. They can be used in the user defined code but no other variables should be defined with these names.

  • DT : Time step (typically in ms) for simulation; Neuron integration can be done in multiple sub-steps inside the neuron model for numerical stability (see Traub-Miles and Izhikevich neuron model variations in Neuron models).
  • addtoinSyn : This variable is used by WeightUpdateModels::Base for updating synaptic input. The way it is modified is defined using the SET_SIM_CODE or SET_EVENT_CODE macros, therefore if a user defines her own model she should update this variable to contain the input to the post-synaptic model.
  • updatelinsyn : At the end of the synaptic update by addtoinSyn, final values are copied back to the d_inSyn<synapsePopulation> variables which will be used in the next step of the neuron update to provide the input to the postsynaptic neurons. This keyword designated where the changes to addtoinSyn have been completed and it is safe to update the summed synaptic input and write back to d_inSyn<synapsePopulation> in device memory.
  • inSyn: This is an intermediary synapse variable which contains the summed input into a postsynaptic neuron (originating from the addtoinSyn variables of the incoming synapses) .
  • Isyn : This is a local variable which contains the (summed) input current to a neuron. It is typically the sum of any explicit current input and all synaptic inputs. The way its value is calculated during the update of the postsynaptic neuron is defined by the code provided in the postsynaptic model. For example, the standard PostsynapticModels::ExpCond postsynaptic model defines
    SET_APPLY_INPUT_CODE("$(Isyn) += $(inSyn)*($(E)-$(V))");
    which implements a conductance based synapse in which the postsynaptic current is given by $I_{\rm syn}= g*s*(V_{\rm rev}-V_{\rm post})$.
    Note
    The addtoinSyn variables from all incoming synapses are automatically summed and added to the current value of inSyn.
    The value resulting from the current converter code is assigned to Isyn and can then be used in neuron sim code like so:
    $(V)+= (-$(V)+$(Isyn))*DT
  • sT : This is a neuron variable containing the last spike time of each neuron and is automatically generated for pre and postsynaptic neuron groups if they are connected using a synapse population with a weight update model that has SET_NEEDS_PRE_SPIKE_TIME(true) or SET_NEEDS_POST_SPIKE_TIME(true) set.

In addition to these variables, neuron variables can be referred to in the synapse models by calling $(<neuronVarName>_pre) for the presynaptic neuron population, and $(<neuronVarName>_post) for the postsynaptic population. For example, $(sT_pre), $(sT_post), $(V_pre), etc.

Debugging suggestions

In Linux, users can call cuda-gdb to debug on the GPU. Example projects in the userproject directory come with a flag to enable debugging (DEBUG=1). genn-buildmodel.sh has a debug flag (-d) to generate debugging data. If you are executing a project with debugging on, the code will be compiled with -g -G flags. In CPU mode the executable will be run in gdb, and in GPU mode it will be run in cuda-gdb in tui mode.

Note
Do not forget to switch debugging flags -g and -G off after debugging is complete as they may negatively affect performance.

On Mac, some versions of clang aren't supported by the CUDA toolkit. This is a recurring problem on Fedora as well, where CUDA doesn't keep up with GCC releases. You can either hack the CUDA header which checks compiler versions - cuda/include/host_config.h - or just use an older XCode version (6.4 works fine).

On Windows models can also be debugged and developed by opening the vcxproj file used to build the model in Visual Studio. From here files can be added to the project, build settings can be adjusted and the full suite of Visual Studio debugging and profiling tools can be used.

Note
When opening the models in the userproject directory in Visual Studio, right-click on the project in the solution explorer, select 'Properties'. Then, making sure the desired configuration is selected, navigate to 'Debugging' under 'Configuration Properties', set the 'Working Directory' to '..' and the 'Command Arguments' to match those passed to genn-buildmodel e.g. 'outdir 1' to use an output directory called outdir and to run the model on the GPU.

Previous | Top | Next