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

* Rename `miniapps` subdirectory to `example`.
* Have all example executables be built under `example` in the build directory.
* Update Travis CI to run miniapp from new path.

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

* Add three new mechanisms: `nax.mod`, `kdrmt.mod` and `kamt.mod`.
* Add new built-in math operators to `modcc`: `min`, `max`, `abs` and `exprelr`. `exprelr` is defined as the reciprocal of the 'exprel' function, exprel(x)=x/(exp(x)-1), exprel(1)=1. This function occurs frequently in HH-style mechanisms, and having a built-in operator avoids the ad hoc `vtrap` functions found in NMODL files in the wild.
* Split Arbor SIMD intrinsics support into AVX2- and AVX512-specific files.
* Add unit tests for new maths operators for C++, SIMD and CUDA implementations.

Reduce differences between `util::optional` and `std::optional`.

* Rename `util::nothing` to `util::nullopt`.
* Replace `util::get()` by`util::optional::value()`.
* Add correct move semantics to `util::optional::value()`.
* Add `util::optional::value_or()` method.
* Remove unused monadic functionality from `util::optional` and `util::uninitialized`.
* Update code formatting to closer adhere to coding guidelines.
* Add convenience ""_s string constructor for unit tests.
* Update unit tests accordingly.

Fixes #419.

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

* Fix various compiler warnings in `modcc` when using Apple Clang.

Implement generic `event_generator`, and some common generators that should be useful to users out of the box:
* poisson sequence                                                                                       
* uniform regular sequence
* from a user-prescribed sequence
* from a `std::vector` of post synaptic events.

Extend `recipe` API to provide an `recipe::event_generators(gid)` that returns a list of `event_generator`s attached to a cell with `gid`.

Update the `model` class and `merge_events` method to support event generators.

Fixes #401 

Fix bug in `rss_cell_group`s with more than one cell with different `dt` or `start_time` values.
Fixes #410.

Vasileios Karakasis authored 7 years ago and Sam Yates committed 7 years ago

Also removes vestigial `intitial` procedure generated from a MOD file with no INITIAL block.
Fixes #407. 

Restructure the miniapp in such a way that we have to option to have multiple parallel mini-applications.

Move the original miniapp directory to a miniapps directory.
Output executable also in a nested miniapps directory.
Update the Travis to point to the new executable location.

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

* Add querying methods to modcc ion species type to simplify code generation.
* Add `WriteBack` type that records per-mechanism ion usage metadata.
* Add `write_back` implementation for C and CUDA printers that adds weighted concentration
  contributions to global concentration fields.
* Extend `uses_ion` interface to provide information about which concentration fields are modified by mechanisms.
* Update `update_ion_symbols` lambda that is responsible for adding metadata about interactions between mechanism fields and external ion species fields, creating a write-back when the mechanism field is a state variable or an indexed variable if it is a current/reversal potential, or else an error.
* Add `test_ca.mod` mechanism that writes calcium ion concentration for testing purposes.
* Add back-end callbacks that initialize concentration values.
* Update `mechanism::uses_ion()` to return information regarding mechanism ion concentration updates.
* Add `mechanism::write_back()` method.
* Update stimulus mechanism specialization to use new mechanism interface.
* Update `ion` type to calculate default concentration contributions in CVs where the concentration is only partially determined by mechanisms.
* Update `fvm_multicell` to calculate the default ion concentration weights.
* Add unit test for ion concentration weight determination.

Fixes #373 

* Remove optimization option (use SIMD options for vectorization).
* Remove arbor utility library dependencies from modcc (pending separation of utility lib from arbor lib source).
* Split target (cpu, gpu) specification from vectorization architecture (avx2, avx512).
* Remove `Options` singleton; replace with structure local to `modcc.cpp`.
* Tidy `modcc` option parsing and main function; allow a single invocation of `modcc` to generate code for multiple backends.
* Rename generated sources to include backend target in filename.
* Always run a constant simplification pass on generated procedures.
* Remove file i/o code from `Module` and `modcc` main function; move functionality to new functions in `io` namespace. (Note: in on-going mechanism revamp, other i/o utility code will reside in the `io` namespace and subdirectory.)
* Remove classes `ConstantFolderVisitor` and `ExpressionClassifierVisitor` that are no longer used.
* Modify CMakeLists.txt files, `backends/*/fvm.cpp` to reflect the new filenames of generated sources.
* Small formatting changes in `modcc` source to reflect coding guidelines (incomplete).

Fixes #399.

Update `test_printers.cpp`:
* Fix compilation and scope management issues, and re-include in CMakeLists.txt.
* Update to use `verbose_print` (run-time verbosity).
* Apply expected output tests to `CPrinter` and `CUDAPrinter` outputs.

Address binary operation rendering issue:
* Move C-style expression rendering common to `CPrinter` and `CUDAPrinter` to a specialized `Visitor`, `CExprRenderer`.
* Use operator associativity to determine need for parentheses around sub-expressions of a binary expression.

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

* Properly iterate over all event queues in the case where a cell group has more than once cell (triggered in e.g. GPU miniapp).
* Update GPU unit test to use the new `cell_group::advance()` interface.

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

Derivatives should only appear on the left hand side of expressions that describe the time evolution of state variables. Without this check `modcc` segfaulted when processing a derivative expression of a non-state variable.

* Specialize the semantic analysis of `DerivativeExpression` to enforce that derivatives are only applied to state variables.

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

* Remove `cell_group::enqueue_events` interface: `cell_group::advance` is given the events to deliver in the next epoch directly.
* Implement event lane merging in `model` with stand-alone `merge_events` function.
* Ensure event-delivery order is preserved across differing domain decompositions by performing a lexicographic sort over all of the `postsynaptic_spike_event` fields: fixes #385. 
                                                                                                         
Fixes #378.

These classes aim to provide an alternative to `TextBuffer` in `modcc` printers.

* Implement a streambuf wrapper `prefixbuf` that prepends a prefix string to each line.
* Adds stream manipulators `setprefix`, `indent`, `popindent`, `settab` for manipulating the prefix in a `prefixbuf`.
* Provide a wrapper stream `pfxstringstream` around `std::stringbuf` with `std::ostringstream` semantics.

As implemented, a prefix string set with the `setprefix` is applied to the underlying streambuf, and is not a property of the stream. This behaviour may be surprising, so consider storing this prefix with the stream as a possible enhancement.

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

Replace constant values for ion species reversal potentials taken from HH model with values that are updated on each time step via the Nernst equation.

* Implement Nernst equation in multicore and gpu back-ends.
* Extend interface of `ion` type to proved a method `update_reversal_potential` which calls the back-end nernst routine.
* Add `valency` and default concentration values to the ion species.
* Add interface for resetting ion species state (for restarting simulations).

Fixes #376 

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

Update the FVM formulation to use current densities instead of currents.

Modifications to modcc:
* Update printers to store and use weights for point process mechanisms,
* Scale ion species current contributions by area proportion, similarly to contributions to the accumulated current.

Changes to FVM code:
* Update weights calculation for density and point processes mechanisms:
    * density channels use relative proportion of CV area, i.e. "density",
    * point processes use the reciprocal of the CV area to convert to a density.
* Add `cv_area` parameter for matrix constructor, which is used by matrix assembly to convert current densities to currents.
* Update stimulus implementations (gpu and cpu backends) to contribute current densities.

Other changes:
* Update unit tests to use new interfaces.
* Update units section in LaTeX docs.

Fixes #374.

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

Fixes #383 (segfault owing to empty vector of binners).

* Add call to `set_binning_policy` in `mc_cell_group` constructor with default `binning_kind::none` policy.

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

Optimize host-side postsynaptic event queues:

* Have communicator store stage local postsynaptic events per cell instead of per cell group.
* Pass (by reference) the relevant subrange of vectors of per-cell events to the cell group `enqueue_events` method.
* Identify calls to cell group `enqueue_events` and `advance` with an integration epoch identification allowing the use of an efficient thread-safe data structure for the cell group event stores.
* Perform merge of new events and currently stored events not for delivery in current epoch in parallel across cells in `mc_cell_group::enqueue_events`.
* Store one event binner object per cell in `mc_cell_group` to avoid hash table lookup overheads in the `binning_kind::following` case, and simplifying the `event_binner` class.
* Add convenience time comparison functional object `event_time_less`.

Fixes #368.

Fixes #350 

* Replace parameter_list with mechanism_spec.
* Add prototype for mechanism parameter schema checking.
* Allow mechanism weights to be set after construction.
* Combine range parameters on density mechanisms by linear contribution in CVs.
* Cable segment electrical parameters are now member variables.
* Publish mechanism parameter information through new method `mechanism::field_info`; note this will be replaced/improved in upcoming dynamic mechanism catalog work.
* Access mechanism parameter scalars and range data via `mechanism::field_view_ptr` and `mechanism::field_value_ptr` methods.
* Allow mechanism 'global' parameters to be set via a method of specializing mechanisms (and giving them corresponding aliases).
* Extend recipe interface to allow querying of per-cell-kind global information for use by cell group implementations.
* Add unit tests for above - note that linear density mechanism parameter test is tightly coup...

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

Fixes #367.

* Add empty INITIAL block if NMODL file does not supply one.

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

Add a micro-benchmark that illustrates scaling issues with event-setup on cell groups with many cells.

The benchmark also illustrates some alternative approaches that are significant optimizations over the current approach.

Fixes #357.

* Use ARB_ and arb_ as variable prefixes in place of NMC_ and nmc_.
* Replace references to 'NestMC' and 'NEST MC' to refer instead to Arbor.
* Use 'arbor' as the sim name in generated validation data.
* Reflow long-line paragraphs in `tests/ubench/README.md`.
* Change names of CUDA mechanism and CUDA kernel libraries to include arbor name.

Change nest::mc namespaces to arb

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

Update README.md to use badge from updated Travis CI page, and updated references to the old nestmc-proto repository.

* Fix attribution for the Cephes library.

* Small language fix

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

Separate compilation for all CUDA code.

* Move all CUDA kernels to their own .cu files, together with C++ function wrappers.
* Compile all CUDA .cu files to a single static library.
* Merge gpu and multicore backend validation tests.
* Simply and clean up cruft from CMakeLists.txt files.

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

Separable compilation of the CUDA kernels generated by modcc from NMODL files.

CMake scripts:
* Update the `build_modules()` helper function to cleanly handle calls to modcc that generate multiple output files.
* Add a new library target `gpu_mechanisms` for the separately compiled CUDA kernels and the implementation of their C wrappers.
* Reduce verbosity of compilation messages.

* Simplify mechanism C++ namespace use: move everything in nest::mc::mechanisms::gpu::_mechanism-name_ into `nest::mc::gpu`, and similarly for multicore mechanism implementations, ions.
* Remove template parameters for `value_type` and `size_type` from all of the 
mechanism implementations, and use `fvm_value_type` and `fvm_size_type` everywhere instead.

modcc changes:
* Modify `CUDAPrinter` to keep track of 3 text buffers, one each for 
  "implementation", "interface" and "implementation interface":
* Write the CUDA implementation interface to `X_impl.hpp`, comprising the definition of the mechanism-specific 'X_ParamParck' struct used to pass function arguments to the CUDA kernels.
* Write the CUDA kernels and C wrappers to `X_impl.cu`.
* Write the public C++ mechanism interface (with calls to implementation wrappers) to `X.hpp`.
* Modify modcc driver to support multiple generated output files.

Only compile AVX2 intrinsics unit tests if NMC_VECTORIZE_TARGET=AVX2 is defined.

Fixes #351.

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

Refactor the threshold_watcher and stack data structures in the gpu backend so that they are amenable to separable compilation.

* Make `gpu::stack<T>` have a host-only interface that wraps a POD type `gpu::stack_base<T>`.
* Implement a `push_back(stack_base, value)` method in `backends/gpu/kernels/stack.hpp` that is visible only to device code.
* Move `test_thresholds` kernel to a .cu file, replacing template parameters with types provided by `backends/fvm_types.hpp`.
* Add a simple C function interface, callable from host side code, defined in `backends/gpu/threshold_common.hpp`.
* Simplify the `gpu::impl::padded_size` function (both to read and in terms of efficiency).
* Use `typeid` as the default for pretty-printing types in the memory back end.
* Update the `test_gpu_stack` unit test to support new gpu stack interface.
* Fix bug in the `test_spikes` unit test, which was not running the GPU back end in the cuda unit tests.

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

The unit test for AVX2 intrinsics does not compile with gcc.
This is a quick fix to get master to compile, while the test is fixed.

Provides the following transcendentals intrinsics:

* `nmc_mm256_exp_pd`
* `nmc_mm256_log_pd`
* `nmc_mm256_pow_pd`
* `nmc_mm256_frexp_pd`

The first three are the equivalent of the corresponding SVML intrinsics without the `nmc` prefix.

The last one is used by the `log` function, but I decided it's nice to have it public.

All results are tested against full precision standard library implementation and provide equal results (except for the `pow` case). `NaN`s and infinities are treated according to the standard.

Limitations:

* Subnormals are treated as zeros by `frexp` and as a result by `log`.

Collect samples in FVM lowered cells across an integration period, prior to delivery to sampler callbacks.

Fixes #283.

* Make `multi_event_stream` a generic data structure, based on the event type. These are now used for both `deliverable_event` objects and `sample_event` objects.
* Add generic accessors for event data, separating access to event metadata (cell index, time) from payload (weight, sampling info etc.).
* Make a flat, CUDA-friendly data structure describing the `multi_event_queue` marked event state, for passing to mechanisms and kernels. This also brings the underlying representations of the `multicore` and `gpu` versions of `multi_event_queue` closer.
* Implement batched sampling kernel for GPU.
* Additional utility function: `util::is_sorted_by`.

Overhead of implementation and sampling process is negligible for CPU, and with no samples less than 1% on GPU. Running the miniapp with 10'000 cells and sampling 1% of the somata at 0.01ms adds approximately 10% to simulation time. 

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 p...

* 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`.