From 5efbcc717dd6964e260aec3e60f28624eaa3c9d9 Mon Sep 17 00:00:00 2001 From: Thomas Nowotny Date: Wed, 9 Sep 2015 22:04:32 +0100 Subject: [PATCH] Now making compiler flags customizable through brian preferences and implemented the interface to GeNN "cpu_only" mode. This canbe useful on non-CUDA computers. --- brian2genn/correctness_testing.py | 2 +- brian2genn/device.py | 28 ++++++++++++++++++++---- brian2genn/preferences.py | 15 +++++++++++++ brian2genn/run_correctness_testing.py | 4 +++- brian2genn/templates/GNUmakefile | 11 +++++++--- brian2genn/templates/WINmakefile | 6 +++++ brian2genn/templates/engine.cc | 12 +++++++++- brian2genn/templates/runner.cc | 10 +++++++++ docs_sphinx/introduction/preferences.rst | 18 +++++++++++++++ 9 files changed, 96 insertions(+), 10 deletions(-) diff --git a/brian2genn/correctness_testing.py b/brian2genn/correctness_testing.py index d9c18083..6bbf9d89 100644 --- a/brian2genn/correctness_testing.py +++ b/brian2genn/correctness_testing.py @@ -12,7 +12,7 @@ class GeNNConfiguration(Configuration): name = 'GeNN' def before_run(self): - brian2.prefs.reset_to_defaults() +# brian2.prefs.reset_to_defaults() brian2.prefs.codegen.loop_invariant_optimisations = False brian2.set_device('genn') diff --git a/brian2genn/device.py b/brian2genn/device.py index e6708b35..50efde6d 100644 --- a/brian2genn/device.py +++ b/brian2genn/device.py @@ -877,6 +877,7 @@ def build(self, directory='output', compile=True, run=True, use_GPU=True, model_name= self.model_name, ROOTDIR= os.path.abspath(directory), source_files= self.source_files, + prefs=prefs ) open(os.path.join(directory, 'WINmakefile'), 'w').write(Makefile_tmp) else: @@ -885,6 +886,7 @@ def build(self, directory='output', compile=True, run=True, use_GPU=True, model_name= self.model_name, ROOTDIR= os.path.abspath(directory), source_files= self.source_files, + prefs= prefs ) open(os.path.join(directory, 'GNUmakefile'), 'w').write(Makefile_tmp) @@ -893,6 +895,10 @@ def build(self, directory='output', compile=True, run=True, use_GPU=True, if compile: with std_silent(debug): if os.sys.platform == 'win32': + # copy .cu file of cpu_only is desired + if prefs.devices.genn.cpu_only: + cmd= "copy runner.cu runnner_cpu_only.cc" + call(cmd, cwd=directory) bitversion= '' if os.getenv('PROCESSOR_ARCHITECTURE') == "AMD64": bitversion= 'x86_amd64' @@ -904,12 +910,23 @@ def build(self, directory='output', compile=True, run=True, use_GPU=True, # Users are required to set their path to "Visual Studio/VC", e.g. # setx VS_PATH "C:\Program Files (x86)\Microsoft Visual Studio 10.0" cmd= "\""+os.getenv('VS_PATH')+"\\VC\\vcvarsall.bat\" " + bitversion - cmd= cmd+" && buildmodel.bat "+self.model_name + " && nmake /f WINmakefile clean && nmake /f WINmakefile" + cmd= cmd+" && buildmodel.bat "+self.model_name + " 0 "; + if prefs.devices.genn.cpu_only: + cmd+= "cpu_only " + cmd+= "&& nmake /f WINmakefile clean && nmake /f WINmakefile" + if prefs.devices.genn.cpu_only: + cmd+= " CPU_ONLY=1" call(cmd, cwd=directory) else: - call(["buildmodel.sh", self.model_name], cwd=directory) - call(["make", "clean"], cwd=directory) - call(["make"], cwd=directory) + if prefs.devices.genn.cpu_only: + shutil.copy(directory+"/runner.cu", directory+"/runner_cpu_only.cc") + call(["buildmodel.sh", self.model_name, "0", "cpu_only"], cwd=directory) + call(["make", "clean"], cwd=directory) + call(["make", "cpu_only"], cwd=directory) + else: + call(["buildmodel.sh", self.model_name], cwd=directory) + call(["make", "clean"], cwd=directory) + call(["make"], cwd=directory) if run: if not with_output: @@ -919,6 +936,9 @@ def build(self, directory='output', compile=True, run=True, use_GPU=True, stdout = None stderr = None start_time = time.time() + if prefs.devices.genn.cpu_only and use_GPU: + logger.warn("Cannot use a GPU in cpu_only mode. Using CPU.") + use_GPU= False gpu_arg = "1" if use_GPU else "0" if gpu_arg == "1": where= 'on GPU' diff --git a/brian2genn/preferences.py b/brian2genn/preferences.py index 97c30d24..e1b4d782 100644 --- a/brian2genn/preferences.py +++ b/brian2genn/preferences.py @@ -37,6 +37,21 @@ def valid_connectivity_decision(expr): docs= ''' This preference allows users to set their own logical expression for the decision whether to use dense or sparse matrix methods based on Npre, Npost and Nsyn. Users should provide a valid Python expression that involves the variable sNpre (pre-synaptic neurons number), Npost (post-synaptic neuron number) and Nsyn (number of existing synapses) and return for using dense matrices and for using sparse. The default behaviour, e.g., corresponds to the logical expression 'Nsyn > 0.25*Npre*Npost'. ''', default= 'Nsyn > 0.25*Npre*Npost' + ), + unix_compiler_flags= BrianPreference( + docs= ''' + This preference allows users to set their own compiler flags for teh eventual compilation of GeNN generated code. The flags will be applied both for the nvcc compilation and C++ compiler compilation (e.g. gcc, clang). This preference is for unix-based operating systems. ''', + default= '-O3 -ffast-math' + ), + windows_compiler_flags= BrianPreference( + docs= ''' + This preference allows users to set their own compiler flags for teh eventual compilation of GeNN generated code. The flags will be applied both for the nvcc compilation and C++ compiler compilation (e.g. MCVC cl). This preference is for the windows operating system. ''', + default= '/O2' + ), + cpu_only= BrianPreference( + docs= ''' + Set this preference to True if you want to compile in GeNN's "CPU only" mode. By default this flag is False and GeNN compiles both a GPU and a CPU version. Which one is run in that case is decided by the useGPU flag in the device.buld command invocation. You might want cpu_only mode if you want to use GeNN on a system that does not have CUDA installed or does not have hardware that supports CUDA.''', + default= False ) ) diff --git a/brian2genn/run_correctness_testing.py b/brian2genn/run_correctness_testing.py index 74c4f9b7..b82c3767 100644 --- a/brian2genn/run_correctness_testing.py +++ b/brian2genn/run_correctness_testing.py @@ -14,6 +14,8 @@ prefs.codegen.loop_invariant_optimisations = False prefs.devices.genn.connectivity = 'AUTO' prefs.core.network.default_schedule= ['start', 'synapses', 'groups', 'thresholds', 'resets', 'end'] +prefs.devices.genn.unix_compiler_flags="-O1" +prefs.devices.genn.cpu_only=True #c = GeNNConfiguration() #c.before_run() @@ -31,7 +33,7 @@ #print run_feature_tests(configurations=[DefaultConfiguration, # GeNNConfiguration], # feature_tests=[NeuronGroupIntegrationLinear]).tables_and_exceptions -run_feature_tests(configurations=[DefaultConfiguration,GeNNConfiguration,CPPStandaloneConfiguration],feature_tests=[ +print run_feature_tests(configurations=[DefaultConfiguration,GeNNConfiguration,CPPStandaloneConfiguration],feature_tests=[ #run_feature_tests(configurations=[GeNNConfiguration], feature_tests=[ NeuronGroupIntegrationLinear, NeuronGroupIntegrationEuler, NeuronGroupLIF, NeuronGroupLIFRefractory, SynapsesPre, SynapsesPost, # SynapsesSTDPNoAutapse, diff --git a/brian2genn/templates/GNUmakefile b/brian2genn/templates/GNUmakefile index 3cb37f6f..3dce0a9e 100644 --- a/brian2genn/templates/GNUmakefile +++ b/brian2genn/templates/GNUmakefile @@ -1,11 +1,16 @@ ROOTDIR := {{ROOTDIR}} EXECUTABLE := runner +{% if prefs['devices.genn.cpu_only'] %} +SOURCES := runner_cpu_only.cc {% for source in source_files %} {{source}} {% endfor %} + +{% else %} SOURCES := runner.cu {% for source in source_files %} {{source}} {% endfor %} +{% endif %} + + INCLUDE_FLAGS := -I. -LINK_FLAGS := -CCFLAGS := -NVCCFLAGS := +OPTIMIZATIONFLAGS := {{prefs['devices.genn.unix_compiler_flags']}} include $(GENN_PATH)/userproject/include/makefile_common_gnu.mk diff --git a/brian2genn/templates/WINmakefile b/brian2genn/templates/WINmakefile index 003b4fe9..b786d2cf 100644 --- a/brian2genn/templates/WINmakefile +++ b/brian2genn/templates/WINmakefile @@ -1,6 +1,12 @@ EXECUTABLE = runner.exe +!IF "$(CPU_ONLY)" != "1" SOURCES = runner.cu {% for source in source_files %} {{source}} {% endfor %} +!ELSE +SOURCES = runner_cpu_only.cc {% for source in source_files %} {{source}} {% endfor %} +!ENDIF EXTRA_INCLUDE = /I %CD% +OPTIMIZATIONFLAGS = {{prefs.devices.genn.windows_compiler_flags}} + !include $(GENN_PATH)\userproject\include\makefile_common_win.mk diff --git a/brian2genn/templates/engine.cc b/brian2genn/templates/engine.cc index 20ac86a0..7a8330b8 100644 --- a/brian2genn/templates/engine.cc +++ b/brian2genn/templates/engine.cc @@ -35,9 +35,11 @@ class engine void free_device_mem(); void run(double, unsigned int); void output_state(FILE *, unsigned int); +#ifndef CPU_ONLY void getStateFromGPU(); void getSpikesFromGPU(); void getSpikeNumbersFromGPU(); +#endif void output_spikes(FILE *, unsigned int); void sum_spikes(); }; @@ -76,11 +78,13 @@ engine::engine() void engine::init(unsigned int which) { +#ifndef CPU_ONLY if (which == CPU) { } if (which == GPU) { copyStateToDevice(); } +#endif } //-------------------------------------------------------------------------- @@ -123,6 +127,7 @@ void engine::run(double runtime, //!< Duration of time to run the model for _run_{{sm.name}}_codeobject(); {% endif %} {% endfor %} +#ifndef CPU_ONLY if (which == GPU) { stepTimeGPU(); t= t-DT; @@ -144,6 +149,7 @@ void engine::run(double runtime, //!< Duration of time to run the model for pull{{sm.monitored}}StateFromDevice(); {% endfor %} } +#endif if (which == CPU) { stepTimeCPU(); t= t-DT; @@ -190,9 +196,10 @@ void engine::output_state(FILE *f, //!< File handle for a file to write the mode unsigned int which //!< Flag determining whether using GPU or CPU only ) { +#ifndef CPU_ONLY if (which == GPU) copyStateFromDevice(); - +#endif fprintf(f, "%f ", t); {% for neuron_model in neuron_models %} for (int i= 0; i < model.neuronN[{{loop.index-1}}]; i++) { @@ -213,6 +220,7 @@ void engine::output_state(FILE *f, //!< File handle for a file to write the mode fprintf(f,"\n"); } +#ifndef CPU_ONLY //-------------------------------------------------------------------------- /*! \brief Method for copying all variables of the last time step from the GPU @@ -250,6 +258,8 @@ void engine::getSpikeNumbersFromGPU() copySpikeNFromDevice(); } +#endif + //-------------------------------------------------------------------------- /*! \brief Method for writing the spikes occurred in the last time step to a file */ diff --git a/brian2genn/templates/runner.cc b/brian2genn/templates/runner.cc index 4b339dff..daf8713c 100644 --- a/brian2genn/templates/runner.cc +++ b/brian2genn/templates/runner.cc @@ -105,7 +105,9 @@ int main(int argc, char *argv[]) //----------------------------------------------------------------- eng.init(which); // this includes copying g's for the GPU version +#ifndef CPU_ONLY copyStateToDevice(); +#endif //------------------------------------------------------------------ // output general parameters to output file and start the simulation @@ -119,10 +121,12 @@ int main(int argc, char *argv[]) fprintf(timef,"%f \n", timer.getElapsedTime()); // get the final results from the GPU +#ifndef CPU_ONLY if (which == GPU) { eng.getStateFromGPU(); eng.getSpikesFromGPU(); } +#endif // translate GeNN arrays back to synaptic arrays {% for synapses in synapse_models %} {% if synapses.connectivity == 'DENSE' %} @@ -178,7 +182,13 @@ using namespace std; #include "hr_time.cpp" #include "utils.h" // for CHECK_CUDA_ERRORS +#ifndef CPU_ONLY #include +#else +#define __host__ +#define __device__ +#endif + #ifndef RAND #define RAND(Y,X) Y = Y * 1103515245 +12345;X= (unsigned int)(Y >> 16) & 32767 diff --git a/docs_sphinx/introduction/preferences.rst b/docs_sphinx/introduction/preferences.rst index b362a128..981c3e83 100644 --- a/docs_sphinx/introduction/preferences.rst +++ b/docs_sphinx/introduction/preferences.rst @@ -3,9 +3,27 @@ Brian2GeNN specific preferences Connectivity ------------------ +The preference ``prefs.devices.genn.conncetivity`` determines what +connectivity scheme is used within GeNN to represent the connections +between neurons. GeNN supports the use of full connectivity matrices +('DENSE') or a representation where connections are represented with +sparse matrix methods ('SPARSE'). When this preference is set to +'AUTO', a separately defined logic expression determines wheter +'DENSE' or 'SPARSE' are used for each individula Synapses population. +This logic expression can be set with the +``prefs.devices.genn.connectivity_decision`` preference. + Connectivity_decision ---------------------- +The ``prefs.devices.genn.connectivity_decision`` contains a logical +expression that determines whether to use 'DENSE' or 'SPARSE' matrix +methods in GeNN for a given matrix. 'DENSE' is used when the condition +evaluates to ``True`` and 'SPARSE' is used otherwise. It's standard +value is ``Nsyn > 0.25*Npre*Npost``, i.e. 'DENSE' is used whenever the +number of actually existing synaptic connections is larger than a +quarter of all possible synaptic connections between the pre- and +post-synaptic populations. Compiler options -----------------