Update Brian2CUDA to work with new run_args feature (brian-team/brian…

…#1429)

brian2cuda/device.py

@@ -433,7 +433,15 @@ def check_openmp_compatible(self, nb_threads):
         if nb_threads > 0:
             raise NotImplementedError("Using OpenMP in a CUDA standalone project is not supported")
 
-    def generate_objects_source(self, writer, arange_arrays, synapses, static_array_specs, networks):
+    def generate_objects_source(
+        self,
+        writer,
+        arange_arrays,
+        synapses,
+        static_array_specs,
+        networks,
+        timed_arrays,
+    ):
         sm_multiplier = prefs.devices.cuda_standalone.SM_multiplier
         num_parallel_blocks = prefs.devices.cuda_standalone.parallel_blocks
         curand_generator_type = prefs.devices.cuda_standalone.random_number_generator_type
@@ -443,7 +451,7 @@ def generate_objects_source(self, writer, arange_arrays, synapses, static_array_
         for var, varname in self.arrays.items():
             if var.name.endswith('space'):  # get all eventspace variables
                 self.eventspace_arrays[var] = varname
-                #if hasattr(var, 'owner') and isinstance(v.owner, Clock):
+                # if hasattr(var, 'owner') and isinstance(v.owner, Clock):
                 if isinstance(var.owner, SpikeGeneratorGroup):
                     self.spikegenerator_eventspaces.append(varname)
         for var in self.eventspace_arrays.keys():
@@ -474,6 +482,7 @@ def generate_objects_source(self, writer, arange_arrays, synapses, static_array_
                         networks=networks,
                         code_objects=self.code_objects.values(),
                         get_array_filename=self.get_array_filename,
+                        get_array_name=self.get_array_name,
                         all_codeobj_with_host_rng=self.codeobjects_with_rng["host_api"]["all_runs"],
                         sm_multiplier=sm_multiplier,
                         num_parallel_blocks=num_parallel_blocks,
@@ -487,6 +496,7 @@ def generate_objects_source(self, writer, arange_arrays, synapses, static_array_
                         profile_statemonitor_copy_to_host=profile_statemonitor_copy_to_host,
                         profile_statemonitor_vars=profile_statemonitor_vars,
                         subgroups_with_spikemonitor=sorted(subgroups_with_spikemonitor),
+                        timed_arrays=timed_arrays,
                         variables_on_host_only=self.variables_on_host_only)
         # Reinsert deleted entries, in case we use self.arrays later? maybe unnecassary...
         self.arrays.update(self.eventspace_arrays)
@@ -1227,7 +1237,7 @@ def generate_makefile(self, writer, cpp_compiler, cpp_compiler_flags, cpp_linker
             )
             writer.write('makefile', makefile_tmp)
 
-    def build(self, directory='output',
+    def build(self, directory='output', results_directory="results",
               compile=True, run=True, debug=False, clean=False,
               with_output=True, disable_asserts=False,
               additional_source_files=None,
@@ -1307,6 +1317,15 @@ def build(self, directory='output',
             directory = tempfile.mkdtemp(prefix='brian_standalone_')
         self.project_dir = directory
         ensure_directory(directory)
+        if os.path.isabs(results_directory):
+            raise TypeError(
+                "The 'results_directory' argument needs to be a relative path but was "
+                f"'{results_directory}'."
+            )
+        # Translate path to absolute path which ends with /
+        self.results_dir = os.path.join(
+            os.path.abspath(os.path.join(directory, results_directory)), ""
+        )
 
         # Determine compiler flags and directories
         cpp_compiler, cpp_default_extra_compile_args = get_compiler_and_args()
@@ -1447,9 +1466,14 @@ def build(self, directory='output',
 
         self.generate_codeobj_source(self.writer)
 
-        self.generate_objects_source(self.writer, self.arange_arrays,
-                                     self.synapses, self.static_array_specs,
-                                     self.networks)
+        self.generate_objects_source(
+            self.writer,
+            self.arange_arrays,
+            self.synapses,
+            self.static_array_specs,
+            self.networks,
+            self.timed_arrays,
+        )
         self.generate_network_source(self.writer)
         self.generate_synapses_classes_source(self.writer)
         self.generate_run_source(self.writer)
@@ -1483,13 +1507,18 @@ def build(self, directory='output',
         if compile:
             self.compile_source(directory, cpp_compiler, debug, clean)
             if run:
-                self.run(directory, with_output, run_args)
+                self.run(
+                    directory=directory,
+                    results_directory=results_directory,
+                    with_output=with_output,
+                    run_args=run_args,
+                )
 
     def network_run(self, net, duration, report=None, report_period=10*second,
                     namespace=None, profile=False, level=0, **kwds):
-        ###################################################
-        ### This part is copied from CPPStandaoneDevice ###
-        ###################################################
+        ####################################################
+        ### This part is copied from CPPStandaloneDevice ###
+        ####################################################
         self.networks.add(net)
         if kwds:
             logger.warn(('Unsupported keyword argument(s) provided for run: '
@@ -1671,12 +1700,15 @@ def network_run(self, net, duration, report=None, report_period=10*second,
             if clock not in all_clocks:
                 run_lines.append(f'{net.name}.add(&{clock.name}, NULL);')
 
-        run_lines.extend(self.code_lines['before_network_run'])
-        # run everything that is run on a clock
+        run_lines.extend(self.code_lines["before_network_run"])
+        if not self.run_args_applied:
+            run_lines.append("set_from_command_line(args);")
+            self.run_args_applied = True
         run_lines.append(
-            f'{net.name}.run({float(duration)!r}, {report_call}, {float(report_period)!r});'
+            f"{net.name}.run({float(duration)!r}, {report_call},"
+            f" {float(report_period)!r});"
         )
-        run_lines.extend(self.code_lines['after_network_run'])
+        run_lines.extend(self.code_lines["after_network_run"])
         # for multiple runs, the random number buffer needs to be reset
         run_lines.append('random_number_buffer.run_finished();')
         # nvprof stuff
@@ -1758,7 +1790,6 @@ def network_restore(self, net, *args, **kwds):
                                    'supported in CUDA standalone'))
 
 
-
 def prepare_codeobj_code_for_rng(codeobj):
     '''
     Prepare a CodeObject for random number generation (RNG).

brian2cuda/templates/main.cu

@@ -25,8 +25,28 @@
 
 {{report_func|autoindent}}
 
+void set_from_command_line(const std::vector<std::string> args)
+{
+    for (const auto& arg : args) {
+		// Split into two parts
+		size_t equal_sign = arg.find("=");
+		auto name = arg.substr(0, equal_sign);
+		auto value = arg.substr(equal_sign + 1, arg.length());
+		brian::set_variable_by_name(name, value);
+	}
+}
+
 int main(int argc, char **argv)
 {
+    std::vector<std::string> args(argv + 1, argv + argc);
+    if (args.size() >=2 && args[0] == "--results_dir")
+    {
+        brian::results_dir = args[1];
+        #ifdef DEBUG
+        std::cout << "Setting results dir to '" << brian::results_dir << "'" << std::endl;
+        #endif
+        args.erase(args.begin(), args.begin()+2);
+    }
     {{'\n'.join(code_lines['before_start'])|autoindent}}
 
     // seed variable set in Python through brian2.seed() calls can use this

brian2cuda/templates/objects.cu

@@ -1,5 +1,19 @@
 {% macro cu_file() %}
 
+{% macro set_from_value(var_dtype, array_name) %}
+{% if c_data_type(var_dtype) == 'double' %}
+set_variable_from_value<double>(name, {{array_name}}, var_size, (double)atof(s_value.c_str()));
+{% elif c_data_type(var_dtype) == 'float' %}
+set_variable_from_value<float>(name, {{array_name}}, var_size, (float)atof(s_value.c_str()));
+{% elif c_data_type(var_dtype) == 'int32_t' %}
+set_variable_from_value<int32_t>(name, {{array_name}}, var_size, (int32_t)atoi(s_value.c_str()));
+{% elif c_data_type(var_dtype) == 'int64_t' %}
+set_variable_from_value<int64_t>(name, {{array_name}}, var_size, (int64_t)atol(s_value.c_str()));
+{% elif c_data_type(var_dtype) == 'char' %}
+set_variable_from_value(name, {{array_name}}, var_size, (char)atoi(s_value.c_str()));
+{% endif %}
+{%- endmacro %}
+
 #include "objects.h"
 #include "synapses_classes.h"
 #include "brianlib/clocks.h"
@@ -18,6 +32,7 @@
 #include <curand_kernel.h>
 
 size_t brian::used_device_memory = 0;
+std::string brian::results_dir = "results/";  // can be overwritten by --results_dir command line arg
 
 //////////////// clocks ///////////////////
 {% for clock in clocks | sort(attribute='name') %}
@@ -29,6 +44,138 @@ Clock brian::{{clock.name}};
 Network brian::{{net.name}};
 {% endfor %}
 
+void set_variable_from_value(std::string varname, char* var_pointer, size_t size, char value) {
+    #ifdef DEBUG
+    std::cout << "Setting '" << varname << "' to " << (value == 1 ? "True" : "False") << std::endl;
+    #endif
+    std::fill(var_pointer, var_pointer+size, value);
+}
+
+template<class T> void set_variable_from_value(std::string varname, T* var_pointer, size_t size, T value) {
+    #ifdef DEBUG
+    std::cout << "Setting '" << varname << "' to " << value << std::endl;
+    #endif
+    std::fill(var_pointer, var_pointer+size, value);
+}
+
+template<class T> void set_variable_from_file(std::string varname, T* var_pointer, size_t data_size, std::string filename) {
+    ifstream f;
+    streampos size;
+    #ifdef DEBUG
+    std::cout << "Setting '" << varname << "' from file '" << filename << "'" << std::endl;
+    #endif
+    f.open(filename, ios::in | ios::binary | ios::ate);
+    size = f.tellg();
+    if (size != data_size) {
+        std::cerr << "Error reading '" << filename << "': file size " << size << " does not match expected size " << data_size << std::endl;
+        return;
+    }
+    f.seekg(0, ios::beg);
+    if (f.is_open())
+        f.read(reinterpret_cast<char *>(var_pointer), data_size);
+    else
+        std::cerr << "Could not read '" << filename << "'" << std::endl;
+    if (f.fail())
+        std::cerr << "Error reading '" << filename << "'" << std::endl;
+}
+
+//////////////// set arrays by name ///////
+void brian::set_variable_by_name(std::string name, std::string s_value) {
+	size_t var_size;
+	size_t data_size;
+	std::for_each(s_value.begin(), s_value.end(), [](char& c) // modify in-place
+    {
+        c = std::tolower(static_cast<unsigned char>(c));
+    });
+    if (s_value == "true")
+        s_value = "1";
+    else if (s_value == "false")
+        s_value = "0";
+	// non-dynamic arrays
+	{% for var, varname in array_specs | dictsort(by='value') %}
+    {% if not var in dynamic_array_specs and not var.read_only %}
+    if (name == "{{var.owner.name}}.{{var.name}}") {
+        var_size = {{var.size}};
+        data_size = {{var.size}}*sizeof({{c_data_type(var.dtype)}});
+        if (s_value[0] == '-' || (s_value[0] >= '0' && s_value[0] <= '9')) {
+            // set from single value
+            {{ set_from_value(var.dtype, "brian::" + get_array_name(var)) }}
+        } else {
+            // set from file
+            set_variable_from_file(name, brian::{{get_array_name(var)}}, data_size, s_value);
+        }
+        {% if get_array_name(var) not in variables_on_host_only %}
+        // copy to device
+        CUDA_SAFE_CALL(
+            cudaMemcpy(
+                brian::dev{{get_array_name(var)}},
+                &brian::{{get_array_name(var)}}[0],
+                sizeof(brian::{{get_array_name(var)}}[0])*brian::_num_{{get_array_name(var)}},
+                cudaMemcpyHostToDevice
+            )
+        );
+        {% endif %}
+        return;
+    }
+    {% endif %}
+    {% endfor %}
+    // dynamic arrays (1d)
+    {% for var, varname in dynamic_array_specs | dictsort(by='value') %}
+    {% if not var.read_only %}
+    if (name == "{{var.owner.name}}.{{var.name}}") {
+        var_size = brian::{{get_array_name(var, access_data=False)}}.size();
+        data_size = var_size*sizeof({{c_data_type(var.dtype)}});
+        if (s_value[0] == '-' || (s_value[0] >= '0' && s_value[0] <= '9')) {
+            // set from single value
+            {{ set_from_value(var.dtype, "&brian::" + get_array_name(var, False) + "[0]") }}
+        } else {
+            // set from file
+            set_variable_from_file(name, &brian::{{get_array_name(var, False)}}[0], data_size, s_value);
+        }
+        {% if get_array_name(var) not in variables_on_host_only %}
+        // copy to device
+        CUDA_SAFE_CALL(
+            cudaMemcpy(
+                thrust::raw_pointer_cast(&brian::dev{{get_array_name(var, False)}}[0]),
+                &brian::{{get_array_name(var, False)}}[0],
+                sizeof(brian::{{get_array_name(var, False)}}[0])*brian::{{get_array_name(var, False)}}.size(),
+                cudaMemcpyHostToDevice
+            )
+        );
+        {% endif %}
+        return;
+    }
+    {% endif %}
+    {% endfor %}
+    {% for var, varname in timed_arrays | dictsort(by='value') %}
+    if (name == "{{varname}}.values") {
+        var_size = {{var.values.size}};
+        data_size = var_size*sizeof({{c_data_type(var.values.dtype)}});
+        if (s_value[0] == '-' || (s_value[0] >= '0' && s_value[0] <= '9')) {
+            // set from single value
+            {{ set_from_value(var.values.dtype, "brian::" + varname + "_values") }}
+
+        } else {
+            // set from file
+            set_variable_from_file(name, brian::{{varname}}_values, data_size, s_value);
+        }
+        {% if varname + "_values" not in variables_on_host_only %}
+        // copy to device
+        CUDA_SAFE_CALL(
+            cudaMemcpy(
+                brian::dev{{varname}}_values,
+                &brian::{{varname}}_values[0],
+                data_size,
+                cudaMemcpyHostToDevice
+            )
+        );
+        {% endif %}
+        return;
+    }
+    {% endfor %}
+    std::cerr << "Cannot set unknown variable '" << name << "'." << std::endl;
+    exit(1);
+}
 //////////////// arrays ///////////////////
 {% for var, varname in array_specs | dictsort(by='value') %}
 {% if not var in dynamic_array_specs %}
@@ -382,7 +529,7 @@ void _write_arrays()
             );
     {% endif %}
     ofstream outfile_{{varname}};
-    outfile_{{varname}}.open("{{get_array_filename(var) | replace('\\', '\\\\')}}", ios::binary | ios::out);
+    outfile_{{varname}}.open(results_dir + "{{get_array_filename(var) | replace('\\', '\\\\')}}", ios::binary | ios::out);
     if(outfile_{{varname}}.is_open())
     {
         outfile_{{varname}}.write(reinterpret_cast<char*>({{varname}}), {{var.size}}*sizeof({{c_data_type(var.dtype)}}));
@@ -399,7 +546,7 @@ void _write_arrays()
     {{varname}} = dev{{varname}};
     {% endif %}
     ofstream outfile_{{varname}};
-    outfile_{{varname}}.open("{{get_array_filename(var) | replace('\\', '\\\\')}}", ios::binary | ios::out);
+    outfile_{{varname}}.open(results_dir + "{{get_array_filename(var) | replace('\\', '\\\\')}}", ios::binary | ios::out);
     if(outfile_{{varname}}.is_open())
     {
         outfile_{{varname}}.write(reinterpret_cast<char*>(thrust::raw_pointer_cast(&{{varname}}[0])), {{varname}}.size()*sizeof({{c_data_type(var.dtype)}}));
@@ -418,7 +565,7 @@ void _write_arrays()
         double copy_time_statemon;
         {% endif %}
         ofstream outfile_{{varname}};
-        outfile_{{varname}}.open("{{get_array_filename(var) | replace('\\', '\\\\')}}", ios::binary | ios::out);
+        outfile_{{varname}}.open(results_dir + "{{get_array_filename(var) | replace('\\', '\\\\')}}", ios::binary | ios::out);
         if(outfile_{{varname}}.is_open())
         {
             {% if var in profile_statemonitor_vars %}
@@ -450,7 +597,7 @@ void _write_arrays()
     {% if profiled_codeobjects is defined and profiled_codeobjects %}
     // Write profiling info to disk
     ofstream outfile_profiling_info;
-    outfile_profiling_info.open("results/profiling_info.txt", ios::out);
+    outfile_profiling_info.open(results_dir + "profiling_info.txt", ios::out);
     if(outfile_profiling_info.is_open())
     {
     {% for codeobj in profiled_codeobjects | sort %}
@@ -476,7 +623,7 @@ void _write_arrays()
     {% endif %}
     // Write last run info to disk
     ofstream outfile_last_run_info;
-    outfile_last_run_info.open("results/last_run_info.txt", ios::out);
+    outfile_last_run_info.open(results_dir + "last_run_info.txt", ios::out);
     if(outfile_last_run_info.is_open())
     {
         outfile_last_run_info << (Network::_last_run_time) << " " << (Network::_last_run_completed_fraction) << std::endl;
@@ -600,6 +747,7 @@ typedef {{curand_float_type}} randomNumber_t;  // random number type
 namespace brian {
 
 extern size_t used_device_memory;
+extern std::string results_dir;
 
 //////////////// clocks ///////////////////
 {% for clock in clocks %}
@@ -611,6 +759,8 @@ extern Clock {{clock.name}};
 extern Network {{net.name}};
 {% endfor %}
 
+extern void set_variable_by_name(std::string, std::string);
+
 //////////////// dynamic arrays 1d ///////////
 {% for var, varname in dynamic_array_specs | dictsort(by='value') %}
 extern thrust::host_vector<{{c_data_type(var.dtype)}}> {{varname}};