Status of the master branch is displayed at the top of the README.
The status is based on the result of the automated daily build.

Add support for continuous integration with Travis CI.
This implements bare bones support that can be extended over time.

Travis CI test environments:

    All use gcc 5.
    Test the serial distributed back end with serial and cthread threading backends.
    Test mpi with cthread.
    The tbb test failed sporadically because CMake, so it is disabled for now.

The test script:

    Builds the unit tests, global_communication tests and miniapp.
    Asserts that all unit and global_communication tests pass.
    Asserts that the miniapp runs successfully.
        does not test miniapp output for now.

There is plenty of scope for improving the tests.
A key improvement will be to use validated output for the validation and miniapp
to provide some validation.

There were some small fixes required to make the tests pass on Travis

    communication/mpi.hpp now sets default size and rank values of 1 and 0 respectively
    to allow all unit tests to pass when built with MPI.
    The wrappers around MPI API calls use const_cast to support MPI implementations that
    are not "const aware".
    A missing header was added to tests/unit/test_range to make std::unordered_multimap
    available.`

* Add `backends/fvm_types.hpp` as a single location for shared types across multicore and gpu fvm implementations.
* Use the `fvm_value_type` and `fvm_size_type` defined in `fvm_types.hpp` for the corresponding class-local `value_type` and `size_type` definitions.

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

* Move gpu-kernel code from memory to backends/gpu

A small step towards seperate back end compilation for CUDA.
Move the following code to the gpu backend:
- the memory::fill* wrappers and fill kernel
- the managed_ptr type
  - only used in backends::gpu
  - has `__device__ __host__` members for dual host-device use.

* update unit&validation tests

* Avoid format securiy warning/error in `mpi_listener.hpp`
* Update recipe classes in `test_communicator.cpp` and `test_domain_decomposition.cpp`.
* Align test names in `test_domain_decomposition.cpp` with those in unit tests.
* Fix `spike_gids` bug in `test_all2all` routine.
* Replace test assertions with `AssertionResult` returns in `test_ring` and `test_all2all`.
* Simplify no-extra-events check in `test_ring` and `test_all2all`.

Towards resolution of issue #283.

* Replace probe/sample infrastructure with new API as outlined in the `sampling_api.rst` documentation.
* Separate `cell` probe information from `cell` description object.
* Add `--list` option to `tsplot`, to summarize available time series for plotting together with their metadata.
* Add `--sample-dt` option to validation tests.
* Change validation time series comparison behaviour: linf distance metric now estimated by comparing reference data points against linearly interpreted simulation samples, rather than the other way around.
* Add utility class `any_ptr` which allows lightweight and type-checked access to a type-erased pointer.
* Replace `singleton_recipe` with two simple recipe base classes used for unit and validation tests, in `tests/simple_recipes.hpp`.
* Simplify RSS cell implementation.
* Add statistical functions for testing of probabilistic methods in `tests/unit/stats.hpp` (and .cpp):
   * simple summary stats class;
   * Kolmogorov–Smirnov one-sided statistic and CDF;
   * approximate Poisson CDF.
* Simplify and refactor miniapp `sample_trace` code.
* Add new utility class `handle_set` for managing a collection of integer handles. (Near-simplest implementation.)
* Relax const-ness constraints on `transform_iterator`.

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

Adds support for building documentation with Sphinx from reStructuredText-formatted files in the `doc` subdirectory. Automatic building has been verified with ReadTheDocs.

* Add basic documentation to the `doc` path.
* Use a git submodule and associated CMake to pull in ReadTheDocs theme at configuration time.

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

* Make `domain_decomposition` object into a flat description object.
* Put original decomposition algorithm into a new function `partitioned_load_balance` that returns a `domain_decomposition` object.

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

* Add support for CMake scripts provided by TBB.
* Update required cmake version to 3.0.

* hack to get linking to work on Cray PE

* improve comments and remove redundant include in CMakeLists

* firewall the tbb cmake files

* tbb threading back end to_string includes version number

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

Changes to `domain_decomposition`:
  * `domain_decomposition` performs two pass load balancing in constructor:
      1. first pass performs global load balance,
      2. second pass distributes cells locally between cpu and gpu cell_groups.
    The current logic for this is very simple and naive, and will be replaced with a load balancer which returns a lighter domain decomposition description in a follow up pull request.
  * Provides a simple `group_description` type that contains gid, `cell_kind` and target backend information for `cell_group_factory`.

Changes to `communicator`:
  * Constructor takes a `domain_decomposition` and recipe.
  * The interface for adding connections and constructing connection table has been removed, as this is now performed within the constructor.
  * Construction is more complicated, as connections are partitioned by source gid which requires multiple passes over the connection information in the recipe.
  * `make_event_queues` updated: spikes and connections are now partitioned by source domain, and an optimization dynamically chooses to iterate over either connection or spike list, whichever is shorter.
  * The `exchange` method now sorts `local_spikes` before global gather to facilitate the optimized spike/connection searching.

Changes to `miniapp`:
  * Automatically use gpu if available and compiled with gpu support.
  * Banner prints out useful information about number of cores, gpus and ranks.
  * Remove -g cell group size flag.

Changes to `cell_group`:
  * `cell_group` interface take a list of gid values instead of a range.
  * Updated internal `cell_group` logic to convert between gid and local indices: use a vector for local index to gid map, and a hash table for gid to local index in `cell_group` implementations that need this lookup.

changes to unit tests
  * tests for the domain decomposition
  * tests for the communicator that test ring and all2all networks

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

There was a spurious warning with gcc 7.1.x in `test_algorithms` as a result of using a Google Test macro (`EXPECT_EQ`) inside an `if` block without surrounding `{}`. This fix adds the missing braces to silence the warnings.

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

Refactor hardware-querying code.

Comprises part of the work on issue #318.

* Create `src/hardware` path for hardware-querying code.
* Create `nest::mc::hw` namespace for hardware-querying code.
* Move memory, affinity and power code from util and threading code to the new diretory and namespace.
* Add `nest::mc::threading::num_threads` function that attempts to determine the number of threads to use according to an environment variable, or by using the low level hardware querying if no environment variable is set.
* Add hardware query for counting the number of available GPUs.

Created a new cell kind that injects a user-suplied stream of spikes, along with corresponding cell_group implementation called dss_cell_group.

The miniapp is updated to test this functionality
  * An ostream spike time parser
  * A new command line argument -I with a Path.

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

Optimized reduce by key used by the GPU back end when accumulating synapse current contributions to compartment currents. This leads to significant speedup in the miniapp for cells with few compartments and many synapses.

* Implement `gpu::reduce_by_key` device function that uses warp intrinsics to perform reduction between threads in a warp before using a global atomic update to store the result.
* Add unit tests for `reduce_by_key` functionality.
* Add micro benchmarks that compare against using CUDA atomics.
* Modify `CudaPrinter` modcc class to emit `reduce_by_key` in place of `cudaAtomicAdd` functions.

Some improvements to meter reporting:
* Shorten names of metering regions in miniapp to make them easier to grep.
* JSON is no longer used as an intermediate data type when gathering distributed meters into a single report, instead conversion to JSON is performed just before writing to file.
* Add a print function for summarizing meter results to a stream.

The less than operator for spikes was not in the nest::mc namespace,
so it was not being picked up by STL algoriths and containers.
This patch makes it a friend operator of the `nest::mc::basic_spike`,
and adds some unit tests to varify that STL algorithms can use
containers of spikes.

fixes #315.

Fixes #307.

* Simplify integration complete test logic in `fvm_multicell.hpp`.
* Re-include `multi_event_stream` pretty-print output for debugging.

Replace the 20 artificial spikes we inserted until now with spikes generated with a frequency spiking cell.
This is a cell that spikes regularly for a set period.
This connections outgoing from this cell mirror all the outgoing connections of cells with gid % 20 == 0.
This means that 5% extra connections have been added to the model to implement the equivalent behaviour with an 'implemented' neuron.

To implement this behavoir a new cell_group type has been introduced.
The fs_cell is extremely small and is not implemented on the backends.

Includes:
Unit test for the rss_cell functionality

Remove improper _POSIX_C_SOURCE check.
Fixes #299.

Finalizes #184 and includes fix for #285. This PR constitutes a brutal rebase of the `devel/async-integration` branch onto current master: git patch, baby.

Implements asynchronous integration of cells within an `fvm_multicell` instance together with the code to implement efficient device-side event delivery.

Summarized patch notes from the development branch:

* Add four new views in mechanism state:
   * `vec_ci`: cv index to cell index (w.r.t. one lowered multicell instance) map.
   * `vec_t`: current or integration starting time, indexed by cell.
   * `vec_t_to`: integration stopping time, indexed by cell.
   * `vec_dt`: pre-computed delta between `vec_t` and `vec_t_to`, to address performance regression on multicore backend.
* Extend notion of indexed view in `modcc` to encompass `CellIndexedVariable` variables, which have a per-cell value.
* Add/change 'built-in' variables `t`, `t_to`, and `dt` for NMODL.
* Additional utility functions:
   * `util::minmax_value`: return one-pass reduction for minimum and maximum.
   * `util::fill`: range-compatible wrapper for `std::fill`.
   * `util::empty` now works across arbitrary sequences (defaults to calling `empty()` method if it exists).
   * Extends C++17-like functionality for path testing and manipulation (c.f. PR #233).
* Use a cached copy of the per-cell time vector for queries in `fvm_multicell`.
* Modifies back-end matrix assembly, threshold detection and stimulus implementations to be asynchronous-friendly.
* Rename some back-end mechanism variables for clarity, now that there are multiple sorts of cell index vectors present. `cell_index` for example has been renamed to `cell_cv_divs` where it acts as a partition division of cv indices (length 1+#cells), and `cell_to_cv` where it acts as a map from cell to index of first cv for that cell (length #cells).
* Move common event-delivery types (`target_handle`, `deliverable_event`) to `src/backends/event.hpp`.
* Add `multi_event_stream` data structure for `gpu` and `multicore` backends, which handles multiple streams of events, bulk loaded but individually popped.
* Add unit tests for `multi_event_stream`.
* Document event delivery process and abstraction in `src/backends/event_delivery.md`.
* Virtualize `mechanism::deliver_events`.
* Perform event delivery in `fvm_multicell::step_integration()`, using the backend-provided interfaces.
* Add zero dt check/support to matrix state implementations.
* Add zero dt unit tests for multicore and gpu back-end matrices.
* Add debugging helper `util::sepval` for printing/tracing container values.
* Add trace csv output option to miniapp.
* Add GPU kernel and unit test for end-of-integration time step test. By default, keep using the copy-to-host-and-test method, as it is faster for cell counts up to circa 10k; investigate adaptive/threshold solutions in the future.
* Explicitly compute and store per-cell and per-compartment `dt` from integration time bounds.
* Determine lower bound on number of integration steps per interval in order to avoid explicit checking of minimum cell times each step.
* Avoid any time value checking for samplers in the integration loop if none of them could be triggered in the interval.

Adds a new AVX512 target for processors supporting only the core AVX512 functionality, which currently means SkyLake Xeon processors.

Add AVX2 instrinsics back end for Haswell and Broadwell architectures.

We're still 3.5% and 5% slower than the icc `#pragma` version on Haswell and Broadwell,
respectively.

There is a potential data race in the `assemble_matrix_interleaved` kernel, where threads in a different warp can overwrite the `buffer_v` and `buffer_i` values before they are used to update the `d` and `rhs` vectors.

This race has been exercised in the asynchronous event delivery branch.

* Add `__syncthreads()` to assemble matrix interleaved kernel after `d` and `rhs` update.

The `emit_gather()` function emitted the "wrong" instruction in terms of its
arguments but the instruction actually generated was correct, because
the `simd_printer` was passing the arguments to `emit_gather()` in a
different order, which was though the correct order for the finally emitted
instruction. Complicated? This commit cleans this up.

Fixes #273.

Note: this incorporates the unit test patch in PR #275.

* Store index into targets collection with synapse CV before sorting, and use this target index to store the permuted instance index in the correct target slot.

Unit test did not take into account that the target handles set by the lowered cell will be grouped by cell.

Demonstrates the problem in #273.

* Test correct assignment of handles to synapse mechanisms and handles under a variety of circumstances.

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

GCC has a power-pc specific bug

https://gcc.gnu.org/bugzilla/show_bug.cgi?id=26374

This is simple to workaround: remove `l` from a double literal in a `constexpr` expression.

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

Fixes #265.

In the `modcc`-generated mechanism code, the `cnexp` solver method uses an expensive call to `exp` to integrate dependent variables over one time step. This commit replaces the exponential with a second-order Padé approximation.

  * Modify `modcc` to insert `exp_pade_11` and `exp_pade_22` functions into every module, which define Padé approximations of second and fourth order respectively (m=n=1 and m=n=2).
  * Have `cnexp` solver use `exp_pade_11` instead of the built in `exp` unary operator.

The validation tests pass for both the 2nd and 4th order approximations; the second order approximation will suffices.

Remove some old compiler directives and properties from modcc-generated kernels that were not used by the intel compiler, and were not recognised by gcc.

Now the optimized x86 kernels can be compiled by both gcc and intel compilers.

* Add micro-benchmark `cuda_compare_and_reduce`.
* Add support for `.cu` cuda benchmarks in `tests/ubench/CMakeLists.txt`.
* Update `tests/ubench/README.md` with benchmark summary and results.
* Updates the version of google benchmark library, for fixed benchmark iteration support.

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

Fixes #266.

Use CUDA to compile the `cell_group_factory` so that the CUDA back end is compiled correctly, instead of the null back end proxy.
  * Added bonus: the miniapp is now compiled using host C++ compiler instead of `nvcc`.

This is a little bit hacky, because this is a stop gap until we have separate compilation of CUDA code.

Fixes #263.

* Add 'verbose' field to command line option struct, set with '-v'.
* If verbose flag is true, emit option summary to stdout.

Refactor model and recipe to build models that have different cell types.

    Refactor recipe::get_cell to return unique_any so that.
        All recipe definitions in tests and miniap had to be updated to use
        the new interface.
    Make a cell_group_factory that forwards arguments for building a
    cell group to the appropriate cell_group constructor.
    Refactor model to use generic cell types
        Constructor now delegates cell_group generation to the
        cell_group_factory.
        Add an implementation file model.cpp for model to reduce compilation
        times (by 2-7 seconds on my desktop).
    Refactor probe enumeration code in model and cell_group
        add interface to cell_group for querying enumeration of probes
        in a cell_group
        use this interface instead of directly computing enumeration in
        model constructor, which no longer has easy access to probe
        information.

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

Add a non-copyable variant of `util::any`.

The two main use cases for such a container are:
 1. Storing types that are not copyable.
 2. Ensuring that no copies are made of copyable types that have to be stored in a type-erased container.                                                                           
                                                                                                          
`unique_any` has the same semantics as `any` with the exception of copy construction and copy assignment, which are explicitly forbidden. The requirement that the contained type be copy-constructable has also been relaxed.

The `any_cast` non-member functions have been overridden for `unique_any`, with the same semantics as for `any`.

Ben Cumming authored 7 years ago and Sam Yates committed 7 years ago

Partial implementation of `std::any` from C++17 standard. See: http://en.cppreference.com/w/cpp/utility/any

The implementation is in the `util` library as `util::any`.

Deviations from the standards description of `std::any`:
  * Does not avoid dynamic allocation of small objects.
  * Does not implement the `in_place_type<T>` constructors from the standard.
  * Does not implement the `emplace` modifier from the standard.

The domain decomposition, whereby cells were partitioned across domains/MPI
ranks and then grouped together, was performed in an ad-hoc manner.
This PR modularizes the domain decomposition.

* A `domain_decomposition` class performs the cell partitioning given the recipe
  and some parameters.
* This implementation has a single `domain_decomposition` class and flags are 
  provided in a `group_rules` struct
  * The `domain_decomposition` could be specialized, or the simple rules struct
    could provide some sort of policy implementation in the future when the 
    need arises
* The domain_decomposition class is initialized inside the model constructor
  and is maintained as state of the model.
* The cell model constructor interface has been simplified to a single constructor
  that takes a recipe and a reference to a `domain_decomposition`.
  * In the future we might pass in only a moveable subset of the 
    `domain_decomposition` information required to build the model.                                      

fixes #242.

Refer to issue #247 and pr #248.

* Replace `compat::sink()` usage with a simpler `compat::barrier_if_icc_leq`, as the icc bug also could be tickled depending on how the `sink()` code was structured.

w-klijn authored 7 years ago and Sam Yates committed 7 years ago

Fixes #249.

* Use CMake source directory instead of current source directory as the working directory in git submodule update for `google-benchmark`, in order to support git versions prior to 1.8.4.

Fixes #247.

Forces evaluation of partition bounds call within the `div_component_sampler` constructor, which for subtle and obscure reasons is mis- or un-computed with icpc and the `-xMIC-AVX512` target option.

* Add `compat::sink()` function that forces evaluation of the argument.
* Add `compat::sink_if_icc_leq(version, const X&)` function that performs this only when run with the Intel compiler subject to version bound.
* Use the `compat::sink_if_icc_leq` function in the `div_component_sampler` constructor.