update rtd build image

- sim: added LogicSim2V, LogicSim4V, LogicSim6V (faster and simpler API)
- sim: fixed model for: many-input gates in bench; dangling nodes; POs that drive further logic
- Circuit: dot output fix and better style, fixed equivalence check of lines and nodes.
- Techlib: Added KYUPY library with supported simulation primitives for debug and testing
- tests: fixed and added more for better coverage
- more type annotations
- updated project config and dependencies for better uv experience, newer lark, newer numpy

.gitignore

@@ -1,12 +1,14 @@
-__pycache__
+__pycache__/
 .ipynb_checkpoints
 .pytest_cache
 .DS_Store
-*.pyc
-docs/_build
-build
-dist
 .idea
 .vscode
-src/kyupy.egg-info
+.venv
+*.py[oc]
+docs/_build
+build/
+dist/
+wheels/
+*.egg-info
 *nogit*

.python-version

@@ -0,0 +1 @@
+3.13

.readthedocs.yaml

@@ -1,15 +1,20 @@
 version: 2
 
 build:
-  os: "ubuntu-20.04"
+  os: "ubuntu-24.04"
   tools:
-    python: "3.8"
+    python: "3.12"
+  jobs:
+    post_create_environment:
+      - python -m pip install sphinx_rtd_theme myst-parser
 
 sphinx:
   fail_on_warning: true
+  configuration: docs/conf.py
 
 python:
   # Install our python package before building the docs
   install:
     - method: pip
       path: .
+

LICENSE.txt

@@ -1,6 +1,6 @@
 MIT License
 
-Copyright (c) 2020-2023 Stefan Holst
+Copyright (c) 2020-2025 Stefan Holst
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal

README.md

@@ -0,0 +1,32 @@
+KyuPy - Pythonic Processing of VLSI Circuits
+============================================
+
+KyuPy is a Python package for processing and analysis of non-hierarchical gate-level VLSI designs.
+It contains fundamental building blocks for research software in the fields of VLSI test, diagnosis and reliability:
+
+* Efficient data structures for gate-level circuits and related design data.
+* Partial [lark](https://github.com/lark-parser/lark) parsers for common design files like
+  bench, gate-level Verilog, standard delay format (SDF), standard test interface language (STIL), design exchange format (DEF).
+* Bit-parallel gate-level 2-, 4-, and 8-valued logic simulation.
+* GPU-accelerated high-throughput gate-level timing simulation.
+* High-performance through the use of [numpy](https://numpy.org) and [numba](https://numba.pydata.org).
+
+
+Getting Started
+---------------
+
+KyuPy is available in [PyPI](https://pypi.org/project/kyupy).
+It requires Python 3.10 or newer, [lark](https://pypi.org/project/lark), and [numpy](https://numpy.org).
+Although optional, [numba](https://numba.pydata.org) should be installed for best performance. [numba-cuda](https://nvidia.github.io/numba-cuda/) is required for GPU/CUDA acceleration.
+If numba is not available, KyuPy will automatically fall back to slow, pure Python execution.
+
+The Jupyter Notebook [Introduction.ipynb](https://github.com/s-holst/kyupy/blob/main/examples/Introduction.ipynb) contains some useful examples to get familiar with the API.
+
+
+Development
+-----------
+
+To work with the latest pre-release source code, clone the [KyuPy GitHub repository](https://github.com/s-holst/kyupy).
+
+* Using ``pip``: Run ``pip install -e .`` within your local checkout to make the package available in your Python environment. The source code comes with tests that can be run with ``pytest``.
+* Using ``uv``: Run ``uv run pytest`` within your local checkout to get started.

README.rst

@@ -1,32 +0,0 @@
-KyuPy - Pythonic Processing of VLSI Circuits
-============================================
-
-KyuPy is a Python package for processing and analysis of non-hierarchical gate-level VLSI designs.
-It contains fundamental building blocks for research software in the fields of VLSI test, diagnosis and reliability:
-
-* Efficient data structures for gate-level circuits and related design data.
-* Partial `lark <https://github.com/lark-parser/lark>`_ parsers for common design files like
-  bench, gate-level Verilog, standard delay format (SDF), standard test interface language (STIL), design exchange format (DEF).
-* Bit-parallel gate-level 2-, 4-, and 8-valued logic simulation.
-* GPU-accelerated high-throughput gate-level timing simulation.
-* High-performance through the use of `numpy <https://numpy.org>`_ and `numba <https://numba.pydata.org>`_.
-
-
-Getting Started
----------------
-
-KyuPy is available in `PyPI <https://pypi.org/project/kyupy>`_.
-It requires Python 3.8 or newer, `lark-parser <https://pypi.org/project/lark-parser>`_, and `numpy`_.
-Although optional, `numba`_ should be installed for best performance.
-GPU/CUDA support in numba may `require some additional setup <https://numba.readthedocs.io/en/stable/cuda/index.html>`_.
-If numba is not available, KyuPy will automatically fall back to slow, pure Python execution.
-
-The Jupyter Notebook `Introduction.ipynb <https://github.com/s-holst/kyupy/blob/main/examples/Introduction.ipynb>`_ contains some useful examples to get familiar with the API.
-
-
-Development
------------
-
-To work with the latest pre-release source code, clone the `KyuPy GitHub repository <https://github.com/s-holst/kyupy>`_.
-Run ``pip install -e .`` within your local checkout to make the package available in your Python environment.
-The source code comes with tests that can be run with ``pytest``.

docs/Makefile

@@ -1,4 +1,4 @@
-# pip install sphinx sphinx-rtd-theme
+# pip install sphinx sphinx-rtd-theme myst-parser
 #
 # Minimal makefile for Sphinx documentation
 #

docs/conf.py

@@ -20,11 +20,11 @@ sys.path.insert(0, os.path.abspath('../src'))
 # -- Project information -----------------------------------------------------
 
 project = 'KyuPy'
-copyright = '2020-2023, Stefan Holst'
+copyright = '2020-2026, Stefan Holst'
 author = 'Stefan Holst'
 
 # The full version, including alpha/beta/rc tags
-release = '0.0.4'
+release = '0.0.7'
 
 
 # -- General configuration ---------------------------------------------------
@@ -35,6 +35,7 @@ release = '0.0.4'
 extensions = [
     'sphinx.ext.autodoc',
     'sphinx_rtd_theme',
+    'myst_parser',
 ]
 
 # Add any paths that contain templates here, relative to this directory.

docs/index.rst

@@ -1,4 +1,5 @@
-.. include:: ../README.rst
+.. include:: ../README.md
+   :parser: myst_parser.sphinx_
 
 API Reference
 -------------

docs/simulators.rst

@@ -14,6 +14,15 @@ Logic Simulation - :mod:`kyupy.logic_sim`
 .. autoclass:: kyupy.logic_sim.LogicSim
    :members:
 
+.. autoclass:: kyupy.logic_sim.LogicSim2V
+   :members:
+
+.. autoclass:: kyupy.logic_sim.LogicSim4V
+   :members:
+
+.. autoclass:: kyupy.logic_sim.LogicSim6V
+   :members:
+
 
 Timing Simulation - :mod:`kyupy.wave_sim`
 -----------------------------------------

pyproject.toml

@@ -0,0 +1,42 @@
+[project]
+name = "kyupy"
+version = "0.0.7"
+authors = [
+  { name="Stefan Holst", email="mail@s-holst.de" },
+]
+description = 'High-performance processing and analysis of non-hierarchical VLSI designs'
+readme = "README.md"
+requires-python = ">=3.12,<3.14"
+dependencies = [
+    "numpy>=1.17.0",
+    "lark>=1.3.0",
+]
+classifiers = [
+        "Development Status :: 3 - Alpha",
+        "Environment :: GPU :: NVIDIA CUDA",
+        "Intended Audience :: Science/Research",
+        "Topic :: Scientific/Engineering :: Electronic Design Automation (EDA)",
+        "License :: OSI Approved :: MIT License",
+        "Operating System :: OS Independent",
+        "Programming Language :: Python :: 3",
+]
+license = "MIT"
+license-files = ["LICENSE.txt"]
+
+[project.urls]
+homepage = "https://github.com/s-holst/kyupy"
+
+[build-system]
+requires = ["hatchling"]
+build-backend = "hatchling.build"
+
+[dependency-groups]
+dev = [
+    "myst-parser>=4.0.1",
+    "sphinx>=8.1.3",
+    "sphinx-rtd-theme>=3.0.2",
+    "pytest>=9.0.1",
+    "ipython>=9.14.1",
+    "jupyter>=1.1.1",
+]
+

setup.py

@@ -1,40 +0,0 @@
-from setuptools import setup, find_packages
-
-with open('README.rst', 'r') as f:
-    long_description = f.read()
-
-setup(
-    name='kyupy',
-    version='0.0.4',
-    description='High-performance processing and analysis of non-hierarchical VLSI designs',
-    long_description=long_description,
-    long_description_content_type='text/x-rst',
-    packages=find_packages(where='src'),
-    package_dir={'': 'src'},
-    url='https://github.com/s-holst/kyupy',
-    author='Stefan Holst',
-    author_email='mail@s-holst.de',
-    python_requires='>=3.8',
-    install_requires=[
-        'numpy>=1.17.0',
-        'lark-parser>=0.8.0'
-    ],
-    extras_requires={
-        'dev': [
-            'pytest>=6.1',
-        ],
-    },
-    classifiers=[
-        'Development Status :: 3 - Alpha',
-        'Environment :: GPU :: NVIDIA CUDA',
-        'Intended Audience :: Science/Research',
-        'Topic :: Scientific/Engineering :: Electronic Design Automation (EDA)',
-        'License :: OSI Approved :: MIT License',
-        'Operating System :: OS Independent',
-        'Programming Language :: Python :: 3',
-        'Programming Language :: Python :: 3 :: Only',
-        'Programming Language :: Python :: 3.8',
-        'Programming Language :: Python :: 3.9',
-        'Programming Language :: Python :: 3.10',
-    ],
-)

src/kyupy/__init__.py

@@ -12,6 +12,8 @@ import gzip
 
 import numpy as np
 
+type NestedNumericList = list['int|NestedNumericList']
+type NestedStrIntDict = dict[str,'int|NestedStrIntDict']
 
 _pop_count_lut = np.asarray([bin(x).count('1') for x in range(256)])
 
@@ -57,6 +59,18 @@ def hr_bytes(nbytes):
         multiplier += 1
     return f'{nbytes:.1f}{["", "ki", "Mi", "Gi", "Ti", "Pi"][multiplier]}B'
 
+def eng(number):
+    """Formats a given number using engineering notation."""
+    exponent = 0
+    if abs(number) < 1:
+        while abs(number) >= 1000:
+            number *= 1000
+            exponent -= 3
+    else:
+        while abs(number) >= 1000:
+            number /= 1000
+            exponent += 3
+    return f'{number:.0f}' + (f'e{exponent}' if exponent != 0 else '')
 
 def hr_time(seconds):
     """Formats a given time interval for human readability."""
@@ -99,9 +113,9 @@ class Timer:
 
 
 class Timers:
-    def __init__(self, t={}): self.timers = defaultdict(Timer) | t
+    def __init__(self, t={}): self.timers: dict[str, Timer] = defaultdict(Timer) | t
     def __getitem__(self, name): return self.timers[name]
-    def __repr__(self): return '{' + ', '.join([f'{k}: {v}' for k, v in self.timers.items()]) + '}'
+    def __repr__(self): return '{' + ', '.join([f'{k}: {v:.2f}' for k, v in self.dict().items()]) + '}'
     def __add__(self, t):
         tmr = Timers(self.timers)
         for k, v in t.timers.items(): tmr.timers[k] += v
@@ -109,7 +123,7 @@ class Timers:
     def sum(self):
         return sum([v.s for v in self.timers.values()])
     def dict(self):
-        return dict([(k, v.s) for k, v in self.timers.items()])
+        return dict(sorted([(k, v.s) for k, v in self.timers.items()], key=lambda item: -item[1]))
 
 
 class Log:
@@ -138,10 +152,10 @@ class Log:
         self._limit = limit
 
     def stop_limit(self):
-        if self.filtered > 0:
-            log.info(f'{self.filtered} more messages (filtered).')
-            self.filtered = 0
         self._limit = -1
+        if self.filtered > 0:
+            self.info(f'{self.filtered} more messages (filtered).')
+            self.filtered = 0
 
     def __getstate__(self):
         return {'elapsed': time.perf_counter() - self.start}
@@ -149,6 +163,8 @@ class Log:
     def __setstate__(self, state):
         self.logfile = sys.stdout
         self.indent = 0
+        self._limit = -1
+        self.filtered = 0
         self.start = time.perf_counter() - state['elapsed']
 
     def write(self, s, indent=0):
@@ -169,7 +185,7 @@ class Log:
             return
         t = time.perf_counter() - self.start
         self.logfile.write(f'# {t:011.3f} {level} {message}\n')
-        self.logfile.flush()
+        #self.logfile.flush()
         self._limit -= 1
 
     def info(self, message):
@@ -184,25 +200,35 @@ class Log:
         """Log an error message."""
         self.log('E', message)
 
+    def progress(self):
+        class Progress:
+            def __init__(self, l): self.l = l
+            def __enter__(self):
+                self.start_time = time.perf_counter()
+                self.lastlog_time = self.start_time
+                self.log_interval = 5
+                return self
+            def update(self, completion: float, msg: str = ''):
+                assert completion > 0 and completion <= 1
+                current_time = time.perf_counter()
+                if current_time > self.lastlog_time + self.log_interval:
+                    elapsed_time = current_time - self.start_time
+                    total_time = elapsed_time / completion
+                    rem_time = total_time - elapsed_time
+                    self.l.log(':', f'{msg} - {completion*100:.0f}% done {hr_time(elapsed_time)} elapsed {hr_time(rem_time)} remaining')
+                    self.log_interval = min(600, int(self.log_interval*1.5))
+                    self.lastlog_time = current_time
+            def __exit__(self, *args): pass
+        return Progress(self)
+
     def range(self, *args):
         """A generator that operates just like the ``range()`` built-in, and also occasionally logs the progress
         and compute time estimates."""
         elems = len(range(*args))
-        start_time = time.perf_counter()
+        with self.progress() as p:
-        lastlog_time = start_time
+            for elem, i in enumerate(range(*args)):
-        log_interval = 5
+                yield i
-        for elem, i in enumerate(range(*args)):
+                p.update((elem + 1) / elems)
-            yield i
-            current_time = time.perf_counter()
-            if current_time > lastlog_time + log_interval:
-                done = (elem + 1) / elems
-                elapsed_time = current_time - start_time
-                total_time = elapsed_time / done
-                rem_time = total_time - elapsed_time
-                self.log(
-                    ':', f'{done*100:.0f}% done {hr_time(elapsed_time)} elapsed {hr_time(rem_time)} remaining')
-                log_interval = min(600, int(log_interval*1.5))
-                lastlog_time = current_time
 
 
 log = Log()
@@ -279,8 +305,6 @@ if importlib.util.find_spec('numba') is not None:
     try:
         list(numba.cuda.gpus)
         from numba import cuda
-        from numba.core import config
-        config.CUDA_LOW_OCCUPANCY_WARNINGS = False
     except CudaSupportError:
         log.warn('Cuda unavailable. Falling back to pure Python.')
         cuda = MockCuda()

src/kyupy/atalanta.py

@@ -0,0 +1,64 @@
+"""A parser and dumper for Atalanta's plain-text test pattern format (*.test).
+
+A *.test file contains one test pattern per line in the form ``<index>: <bits>``,
+where ``<bits>`` is a string of ``'0'`` and ``'1'`` characters giving the values of
+the primary inputs (in circuit order). Lines starting with ``'*'`` are comments.
+
+This parser also accepts files without ``<index>:`` and patterns elements in 8-valued
+logic such as ``'R'``, ``'F'``, ``'X'``.
+"""
+
+import re
+
+import numpy as np
+
+from . import readtext, logic
+from .circuit import Circuit
+
+_pattern_re = re.compile(r'^\s*(?:\d+:\s*)?([01HLRFX-]+)\s*$')
+
+
+class TestFile:
+    """Intermediate representation of a *.test pattern file.
+
+    :param text: The contents of a *.test file.
+    """
+
+    def __init__(self, text: str):
+        patterns = []
+        for line in text.splitlines():
+            if (m := _pattern_re.match(line)) is not None:
+                patterns.append(logic.mvarray(m.group(1)))
+
+        self.patterns = np.stack(patterns, axis=-1) if patterns \
+            else np.zeros((0, 0), dtype=np.uint8)
+        """The parsed patterns as a mvarray (see :py:mod:`~kyupy.logic`).
+
+        The second-to-last axis goes along primary inputs, the last axis goes along patterns.
+        """
+
+    def tests(self, circuit: Circuit):
+        """Assembles and returns the test pattern set for the given circuit.
+
+        :param circuit: The circuit to assemble the patterns for. The patterns will follow the
+            :py:attr:`~kyupy.circuit.Circuit.io_nodes` ordering of the circuit.
+        :return: A logic array (see :py:mod:`~kyupy.logic`). The values for primary inputs are
+            filled, all other values are left unassigned.
+        """
+        pi_locs = [i for i, n in enumerate(circuit.io_nodes) if len(n.ins) == 0]
+        tests = np.full((len(circuit.io_nodes), self.patterns.shape[-1]), logic.UNASSIGNED, dtype=np.uint8)
+        tests[pi_locs] = self.patterns
+        return tests
+
+
+def parse(text) -> TestFile:
+    """Parses the given ``text`` and returns a :class:`TestFile` object."""
+    return TestFile(text)
+
+
+def load(file) -> TestFile:
+    """Parses the contents of ``file`` and returns a :class:`TestFile` object.
+
+    Files with `.gz`-suffix are decompressed on-the-fly.
+    """
+    return parse(readtext(file))

src/kyupy/bench.py

@@ -5,12 +5,20 @@ Historically it was first used in the
 `ISCAS89 benchmark set <https://people.engr.ncsu.edu/brglez/CBL/benchmarks/ISCAS89/>`_.
 Besides loading these benchmarks, this module is also useful for easily constructing simple circuits:
 ``c = bench.parse('input(x, y) output(a, o, n) a=and(x,y) o=or(x,y) n=not(x)')``.
+
+This parser automatically translates gate to simulation primitives supported by kyupy.
+In most cases this means just adding the number of inputs to the name: e.g. NAND -> NAND2, OR -> OR4, INV -> INV1.
+Since simulators only support at most four-input-gates, gates with more than four inputs are implemented as tree.
 """
 
 from lark import Lark, Transformer
 
 from .circuit import Circuit, Node, Line
-from . import readtext
+from . import readtext, batchrange, log
+
+def treeify(l, max_degree=4):
+    if len(l) <= max_degree: return l
+    return treeify([l[bo:bo+bs] for bo, bs in batchrange(len(l), max_degree)])
 
 
 class BenchTransformer(Transformer):
@@ -18,16 +26,62 @@ class BenchTransformer(Transformer):
     def __init__(self, name):
         super().__init__()
         self.c = Circuit(name)
+        self.inputs = set()
+        self.outputs = set()
 
     def start(self, _): return self.c
 
-    def parameters(self, args): return [self.c.get_or_add_fork(str(name)) for name in args]
+    def parameters(self, args): return [self.c.get_or_add_fork(str(name)) for name in args if name is not None]
 
-    def interface(self, args): self.c.io_nodes.extend(args[0])
+    def input(self, args):
+        for input in args[0]:
+            if input in self.outputs:
+                new_input = self.c.get_or_add_fork(input.name + '~i')
+                Line(self.c, new_input, input)
+                log.warn(f'input-output passthrough, renaming input: {input.name} -> {new_input.name}')
+                input = new_input
+            self.inputs.add(input)
+            self.c.io_nodes.append(input)
+
+    def output(self, args):
+        for output in args[0]:
+            if output in self.inputs:
+                new_output = self.c.get_or_add_fork(output.name + '~o')
+                Line(self.c, output, new_output)
+                log.warn(f'input-output passthrough, renaming output: {output.name} -> {new_output.name}')
+                output = new_output
+            self.outputs.add(output)
+            self.c.io_nodes.append(output)
+
+    def _cell_tree_inner(self, name, kind, inner_kind, drivers):
+        cell = Node(self.c, name, f'{kind}{len(drivers)}')
+        fork = self.c.get_or_add_fork(name)
+        Line(self.c, cell, fork)
+        for i, d in enumerate(drivers):
+            while isinstance(d, list) and len(d) == 1: d = d[0]
+            if isinstance(d, list):
+                d = self._cell_tree_inner(f'{name}~{i}', inner_kind, inner_kind, d)
+            Line(self.c, d, cell)
+        return fork
+
+    def cell_tree(self, name, kind, drivers):
+        root_kind = kind.upper()
+        inner_kind = root_kind
+        if root_kind == 'NAND': inner_kind = 'AND'
+        if root_kind == 'NOR': inner_kind = 'OR'
+        if root_kind == 'XNOR': inner_kind = 'XOR'
+        return self._cell_tree_inner(name, root_kind, inner_kind, treeify(drivers))
 
     def assignment(self, args):
-        name, cell_type, drivers = args
+        name, kind, drivers = args
-        cell = Node(self.c, str(name), str(cell_type))
+        if kind.upper() in ('AND', 'NAND', 'OR', 'NOR', 'XOR', 'XNOR'):
+            self.cell_tree(name, kind, drivers)
+            return
+        if kind.upper().startswith('BUF'):
+            kind = 'BUF1'
+        elif kind.upper().startswith('INV') or kind.upper().startswith('NOT'):
+            kind = 'INV1'
+        cell = Node(self.c, str(name), str(kind))
         Line(self.c, cell, self.c.get_or_add_fork(str(name)))
         for d in drivers: Line(self.c, d, cell)
 
@@ -35,8 +89,8 @@ class BenchTransformer(Transformer):
 GRAMMAR = r"""
     start: (statement)*
     statement: input | output | assignment
-    input: ("INPUT" | "input") parameters -> interface
+    input: ("INPUT" | "input") parameters
-    output: ("OUTPUT" | "output") parameters -> interface
+    output: ("OUTPUT" | "output") parameters
     assignment: NAME "=" NAME parameters
     parameters: "(" [ NAME ( "," NAME )* ] ")"
     NAME: /[-_a-z0-9]+/i
@@ -44,14 +98,14 @@ GRAMMAR = r"""
     """
 
 
-def parse(text, name=None):
+def parse(text, name=None) -> Circuit:
     """Parses the given ``text`` as ISCAS89 bench code.
 
     :param text: A string with bench code.
     :param name: The name of the circuit. Circuit names are not included in bench descriptions.
     :return: A :class:`Circuit` object.
     """
-    return Lark(GRAMMAR, parser="lalr", transformer=BenchTransformer(name)).parse(text)
+    return Lark(GRAMMAR, parser="lalr", transformer=BenchTransformer(name)).parse(text) # type: ignore
 
 
 def load(file, name=None):

src/kyupy/circuit.py

@@ -10,20 +10,44 @@ Circuit graphs also define an ordering of inputs, outputs and other nodes to eas
 
 """
 
+from __future__ import annotations
+
 from collections import deque, defaultdict
 import re
 
 import numpy as np
 
+from . import NestedNumericList, NestedStrIntDict
 
-class GrowingList(list):
+class GrowingList[T](list[T]):
     def __setitem__(self, index, value):
-        if index >= len(self):
+        if value is None: self.has_nones = True
-            self.extend([None] * (index + 1 - len(self)))
+        if index == len(self): return super().append(value)
+        if index > len(self):
+            super().extend([None] * (index + 1 - len(self)))
+            self.has_nones = True
         super().__setitem__(index, value)
 
-    def free_index(self):
+    # Override __getitem__ to return None when reading beyond the list
-        return next((i for i, x in enumerate(self) if x is None), len(self))
+    # instead of throwing an exception. Type checker complains about the None return
+    # type, though. Probably not needed anyways.
+
+    #def __getitem__(self, index) -> list[T] | T | None:
+    #    if isinstance(index, slice): return super().__getitem__(index)
+    #    return super().__getitem__(index) if index < len(self) else None
+
+    @property
+    def free_idx(self):
+        fi = len(self)
+        if hasattr(self, 'has_nones') and self.has_nones:
+            fi = next((i for i, x in enumerate(self) if x is None), len(self))
+            self.has_nones = fi < len(self)
+        return fi
+
+    def without_nones(self):
+        for item in self:
+            if item is not None:
+                yield item
 
 
 class IndexList(list):
@@ -76,10 +100,10 @@ class Node:
         by allocating an array or list :code:`my_data` of length :code:`len(n.circuit.nodes)` and
         accessing it by :code:`my_data[n.index]` or simply by :code:`my_data[n]`.
         """
-        self.ins = GrowingList()
+        self.ins: GrowingList[Line] = GrowingList()
         """A list of input connections (:class:`Line` objects).
         """
-        self.outs = GrowingList()
+        self.outs: GrowingList[Line] = GrowingList()
         """A list of output connections (:class:`Line` objects).
         """
 
@@ -114,10 +138,10 @@ class Node:
 
         This is ok, because (name, kind) is unique within a circuit.
         """
-        return self.name == other.name and self.kind == other.kind
+        return self.name == other.name and self.kind == other.kind and self.circuit == other.circuit
 
     def __hash__(self):
-        return hash((self.name, self.kind))
+        return hash((self.name, self.kind, id(self.circuit)))
 
 
 class Line:
@@ -135,7 +159,7 @@ class Line:
     Use the explicit case only if connections to specific pins are required.
     It may overwrite any previous line references in the connection list of the nodes.
     """
-    def __init__(self, circuit, driver, reader):
+    def __init__(self, circuit: Circuit, driver: Node | tuple[Node, None|int], reader: Node | tuple[Node, None|int]):
         self.circuit = circuit
         """The :class:`Circuit` object the line is part of.
         """
@@ -147,21 +171,21 @@ class Line:
         by allocating an array or list :code:`my_data` of length :code:`len(l.circuit.lines)` and
         accessing it by :code:`my_data[l.index]` or simply by :code:`my_data[l]`.
         """
-        if not isinstance(driver, tuple): driver = (driver, driver.outs.free_index())
+        if not isinstance(driver, tuple): driver = (driver, driver.outs.free_idx)
-        self.driver = driver[0]
+        self.driver: Node = driver[0]
         """The :class:`Node` object that drives this line.
         """
-        self.driver_pin = driver[1]
+        self.driver_pin = driver[1] if driver[1] is not None else self.driver.outs.free_idx
         """The output pin position of the driver node this line is connected to.
 
         This is the position in the list :py:attr:`Node.outs` of the driving node this line referenced from:
         :code:`self.driver.outs[self.driver_pin] == self`.
         """
-        if not isinstance(reader, tuple): reader = (reader, reader.ins.free_index())
+        if not isinstance(reader, tuple): reader = (reader, reader.ins.free_idx)
-        self.reader = reader[0]
+        self.reader: Node = reader[0]
         """The :class:`Node` object that reads this line.
         """
-        self.reader_pin = reader[1]
+        self.reader_pin = reader[1] if reader[1] is not None else self.reader.ins.free_idx
         """The input pin position of the reader node this line is connected to.
 
         This is the position in the list :py:attr:`Node.ins` of the reader node this line referenced from:
@@ -183,8 +207,6 @@ class Line:
                 for i, l in enumerate(self.driver.outs): l.driver_pin = i
         if self.reader is not None: self.reader.ins[self.reader_pin] = None
         if self.circuit is not None: del self.circuit.lines[self.index]
-        self.driver = None
-        self.reader = None
         self.circuit = None
 
     def __index__(self):
@@ -258,20 +280,19 @@ class Circuit:
         Use the :class:`Node` constructor and :py:attr:`Node.remove()` to add and remove nodes.
         """
 
-    @property
+    def s_nodes(self, tlib: 'TechLib') -> list: # type: ignore
-    def s_nodes(self):
         """A list of all primary I/Os as well as all flip-flops and latches in the circuit (in that order).
 
         The s_nodes list defines the order of all ports and all sequential elements in the circuit.
         This list is constructed on-the-fly. If used in some inner toop, consider caching the list for better performance.
         """
-        return list(self.io_nodes) + [n for n in self.nodes if 'dff' in n.kind.lower()] + [n for n in self.nodes if 'latch' in n.kind.lower()]
+        return list(self.io_nodes) + [n for n in self.nodes if tlib.is_dff(n.kind)] + [n for n in self.nodes if tlib.is_latch(n.kind)]
 
     def io_locs(self, prefix):
-        """Returns the indices of primary I/Os that start with given name prefix.
+        """Returns a list of indices of primary I/Os that start with given name prefix.
 
         The returned values are used to index into the :py:attr:`io_nodes` array.
-        If only one I/O cell matches the given prefix, a single integer is returned.
+        If only one I/O cell matches the given prefix, a list with a single integer is returned.
         If a bus matches the given prefix, a sorted list of indices is returned.
         Busses are identified by integers in the cell names following the given prefix.
         Lists for bus indices are sorted from LSB (e.g. :code:`data[0]`) to MSB (e.g. :code:`data[31]`).
@@ -281,33 +302,33 @@ class Circuit:
         """
         return self._locs(prefix, list(self.io_nodes))
 
-    def s_locs(self, prefix):
+    def s_locs(self, prefix, tlib: 'TechLib'): # type: ignore
         """Returns the indices of I/Os and sequential elements that start with given name prefix.
 
         The returned values are used to index into the :py:attr:`s_nodes` list.
         It works the same as :py:attr:`io_locs`. See there for more details.
         """
-        return self._locs(prefix, self.s_nodes)
+        return self._locs(prefix, self.s_nodes(tlib))
 
-    def _locs(self, prefix, nodes):
+    def _locs(self, prefix, nodes:list[Node]) -> NestedNumericList:  # can return list[list[...]]
-        d_top = dict()
+        d_top: NestedStrIntDict = dict()
         for i, n in enumerate(nodes):
-            if m := re.match(fr'({prefix}.*?)((?:[\d_\[\]])*$)', n.name):
+            if m := re.match(fr'({re.escape(prefix)}.*?)((?:[\d_\[\]])*$)', n.name):
                 path = [m[1]] + [int(v) for v in re.split(r'[_\[\]]+', m[2]) if len(v) > 0]
                 d = d_top
                 for j in path[:-1]:
                     d[j] = d.get(j, dict())
                     d = d[j]
+                    assert isinstance(d, dict)
                 d[path[-1]] = i
 
         # sort recursively for multi-dimensional lists.
-        def sorted_values(d): return [sorted_values(v) for k, v in sorted(d.items())] if isinstance(d, dict) else d
+        def sorted_values(d) -> NestedNumericList: return [sorted_values(v) for k, v in sorted(d.items())] if isinstance(d, dict) else d
         l = sorted_values(d_top)
-        while isinstance(l, list) and len(l) == 1: l = l[0]
+        while isinstance(l, list) and len(l) == 1 and isinstance(l[0], list): l = l[0]
-        return None if isinstance(l, list) and len(l) == 0 else l
+        return l
 
-    @property
+    def stats(self, tlib: 'TechLib'): # type: ignore
-    def stats(self):
         """A dictionary with the counts of all different elements in the circuit.
 
         The dictionary contains the number of all different kinds of nodes, the number
@@ -325,24 +346,25 @@ class Circuit:
         stats['__line__'] = len(self.lines)
         for n in self.cells.values():
             stats[n.kind] += 1
-            if 'dff' in n.kind.lower(): stats['__dff__'] += 1
+            if tlib.is_dff(n.kind): stats['__dff__'] += 1
-            elif 'latch' in n.kind.lower(): stats['__latch__'] += 1
+            elif tlib.is_latch(n.kind): stats['__latch__'] += 1
-            elif 'put' not in n.kind.lower(): stats['__comb__'] += 1 # no input or output
+            elif 'put' not in n.kind.lower(): stats['__comb__'] += 1 # don't count input or output towards comb
         stats['__seq__'] = stats['__dff__'] + stats['__latch__']
         return dict(stats)
 
     def get_or_add_fork(self, name):
         return self.forks[name] if name in self.forks else Node(self, name)
 
-    def remove_dangling_nodes(self, root_node:Node):
+    def remove_dangling_nodes(self, root_node:Node, keep=[]):
         if len([l for l in root_node.outs if l is not None]) > 0: return
         lines = [l for l in root_node.ins if l is not None]
         drivers = [l.driver for l in lines]
+        if root_node in keep: return
         root_node.remove()
         for l in lines:
             l.remove()
         for d in drivers:
-            self.remove_dangling_nodes(d)
+            self.remove_dangling_nodes(d, keep=keep)
 
     def eliminate_1to1_forks(self):
         """Removes all forks that drive only one node.
@@ -370,6 +392,21 @@ class Circuit:
             in_line.reader_pin = out_reader_pin
             in_line.reader.ins[in_line.reader_pin] = in_line
 
+    def remove_forks(self):
+        ios = set(self.io_nodes)
+        for n in list(self.forks.values()):
+            if n in ios: continue
+            d = None
+            if (l := n.ins[0]) is not None:
+                d = l.driver
+                l.remove()
+            for l in list(n.outs):
+                if l is None: continue
+                r, rp = l.reader, l.reader_pin
+                l.remove()
+                if d is not None: Line(self, d, (r, rp))
+            n.remove()
+
     def substitute(self, node, impl):
         """Replaces a given node with the given implementation circuit.
 
@@ -410,7 +447,7 @@ class Circuit:
                     node_map[n] = Node(self, f'{node.name}~{n.name}', n.kind)
             elif len(n.outs) > 0 and len(n.ins) > 0:  # output is also read by impl. circuit, need to add a fork.
                 node_map[n] = Node(self, f'{node.name}~{n.name}')
-            elif len(n.ins) == 0 and len(n.outs) > 1:  # input is read by multiple nodes, need to add fork.
+            elif len(n.ins) == 0 and len(n.outs) != 1:  # input is read by multiple nodes (or no nodes), need to add fork.
                 node_map[n] = Node(self, f'{node.name}~{n.name}')
         for l in impl.lines:  # add all internal lines to main circuit
             if l.reader in node_map and l.driver in node_map:
@@ -428,7 +465,7 @@ class Circuit:
         for l, ll in zip(impl_out_lines, node_out_lines):  # connect outputs
             if ll is None:
                 if l.driver in node_map:
-                    self.remove_dangling_nodes(node_map[l.driver])
+                    self.remove_dangling_nodes(node_map[l.driver], keep=ios)
                 continue
             if len(l.reader.outs) > 0:  # output is also read by impl. circuit, connect to fork.
                 ll.driver = node_map[l.reader]
@@ -447,6 +484,21 @@ class Circuit:
             if n.kind in tlib.cells:
                 self.substitute(n, tlib.cells[n.kind][0])
 
+    def remove_constants(self):
+        c1gen = None
+        for n in self.nodes:
+            if n.kind == '__const0__':  # just remove, unconnected inputs are defined 0.
+                for l in n.outs:
+                    l.remove()
+                n.remove()
+            elif n.kind == '__const1__':
+                if c1gen is None: c1gen = Node(self, '__const1gen__', 'INV1')  # one unique const 1 generator
+                for l in n.outs:
+                    r, rp = l.reader, l.reader_pin
+                    l.remove()
+                    Line(self, c1gen, (r, rp))
+                n.remove()
+
     def copy(self):
         """Returns a deep copy of the circuit.
         """
@@ -488,9 +540,6 @@ class Circuit:
         for n in state['io_nodes']:
             self.io_nodes.append(self.nodes[n])
 
-    def __eq__(self, other):
-        return self.nodes == other.nodes and self.lines == other.lines and self.io_nodes == other.io_nodes
-
     def __repr__(self):
         return f'{{name: "{self.name}", cells: {len(self.cells)}, forks: {len(self.forks)}, lines: {len(self.lines)}, io_nodes: {len(self.io_nodes)}}}'
 
@@ -501,14 +550,15 @@ class Circuit:
         substrings 'dff' or 'latch' are yielded first.
         """
         visit_count = np.zeros(len(self.nodes), dtype=np.uint32)
-        queue = deque(n for n in self.nodes if len(n.ins) == 0 or 'dff' in n.kind.lower() or 'latch' in n.kind.lower())
+        start = set(n for n in self.nodes if len(n.ins) == 0 or 'dff' in n.kind.lower() or 'latch' in n.kind.lower())
+        queue = deque(start)
         while len(queue) > 0:
             n = queue.popleft()
             for line in n.outs:
                 if line is None: continue
                 succ = line.reader
                 visit_count[succ] += 1
-                if visit_count[succ] == len(succ.ins) and 'dff' not in succ.kind.lower() and 'latch' not in succ.kind.lower():
+                if visit_count[succ] == len(succ.ins) and succ not in start:
                     queue.append(succ)
             yield n
 
@@ -530,24 +580,23 @@ class Circuit:
                 if line is not None:
                     yield line
 
-    def reversed_topological_order(self):
+    def reversed_topological_order(self, tlib: 'TechLib'): # type: ignore
         """Generator function to iterate over all nodes in reversed topological order.
 
-        Nodes without output lines and nodes whose :py:attr:`Node.kind` contains the
+        Nodes without output lines and sequential nodes (flip-flops, latches) are yielded first.
-        substrings 'dff' or 'latch' are yielded first.
         """
         visit_count = [0] * len(self.nodes)
-        queue = deque(n for n in self.nodes if len(n.outs) == 0 or 'dff' in n.kind.lower() or 'latch' in n.kind.lower())
+        queue = deque(n for n in self.nodes if len(n.outs) == 0 or tlib.is_dff(n.kind) or tlib.is_latch(n.kind))
         while len(queue) > 0:
             n = queue.popleft()
             for line in n.ins:
                 pred = line.driver
                 visit_count[pred] += 1
-                if visit_count[pred] == len(pred.outs) and 'dff' not in pred.kind.lower() and 'latch' not in pred.kind.lower():
+                if visit_count[pred] == len(pred.outs) and not tlib.is_dff(pred.kind) and not tlib.is_latch(pred.kind):
                     queue.append(pred)
             yield n
 
-    def fanin(self, origin_nodes):
+    def fanin(self, origin_nodes, tlib: 'TechLib'): # type: ignore
         """Generator function to iterate over the fan-in cone of a given list of origin nodes.
 
         Nodes are yielded in reversed topological order.
@@ -555,16 +604,34 @@ class Circuit:
         marks = [False] * len(self.nodes)
         for n in origin_nodes:
             marks[n] = True
-        for n in self.reversed_topological_order():
+        for n in self.reversed_topological_order(tlib):
             if not marks[n]:
-                for line in n.outs:
+                for line in n.outs.without_nones():
-                    if line is not None:
+                    marks[n] |= marks[line.reader]
-                        marks[n] |= marks[line.reader]
             if marks[n]:
                 yield n
 
-    def fanout_free_regions(self):
+    def fanout(self, origin_nodes: list[Node], node_filter = lambda n: 'dff' not in n.name.lower()):
-        for stem in self.reversed_topological_order():
+        """Generator function to iterate over the fan-out cone of a given list of origin nodes.
+
+        origin_nodes are yielded first, followed by nodes driven by them in a breadth-first manner.
+        The search stops at nodes for which node_filter returns False.
+        Only origin_nodes and nodes for which node_filter returned True are yielded.
+        By default, search stops at flip-flops.
+        """
+        queue = deque(origin_nodes)
+        yielded = set()
+        while len(queue) > 0:
+            n = queue.popleft()
+            for line in n.outs.without_nones():
+                succ = line.reader
+                if succ not in yielded and node_filter(succ):
+                    yielded.add(succ)
+                    queue.append(succ)
+            yield n
+
+    def fanout_free_regions(self, tlib: 'TechLib'): # type: ignore
+        for stem in self.reversed_topological_order(tlib):
             if len(stem.outs) == 1 and 'dff' not in stem.kind.lower(): continue
             region = []
             if 'dff' in stem.kind.lower():
@@ -584,11 +651,11 @@ class Circuit:
                 region.append(n)
             yield stem, region
 
-    def dot(self, format='svg'):
+    def dot(self, tlib: 'TechLib', format='svg', graph_attr={}, line_labels={}):
         from graphviz import Digraph
-        dot = Digraph(format=format, graph_attr={'rankdir': 'LR', 'splines': 'true'})
+        dot = Digraph(format=format, graph_attr={'rankdir': 'LR', 'splines': 'true', 'size': '10', 'ranksep': '0.1'} | graph_attr)
 
-        s_dict = dict((n, i) for i, n in enumerate(self.s_nodes))
+        s_dict = dict((n, i) for i, n in enumerate(self.s_nodes(tlib)))
         node_level = np.zeros(len(self.nodes), dtype=np.uint32)
         level_nodes = defaultdict(list)
         for n, lv in self.topological_order_with_level():
@@ -599,17 +666,24 @@ class Circuit:
             with dot.subgraph() as s:
                 s.attr(rank='same')
                 for n in level_nodes[lv]:
-                    ins = '|'.join([f'<i{i}>{i}' for i in range(len(n.ins))])
+                    ins = '{' + '|'.join([f'<i{i}>{i}' for i in range(len(n.ins))]) + '}|' if len(n.ins) > 1 else ''
-                    outs = '|'.join([f'<o{i}>{i}' for i in range(len(n.outs))])
+                    outs = '|{' + '|'.join([f'<o{i}>{i}' for i in range(len(n.outs))]) + '}' if len(n.outs) > 1 else ''
                     io = f' [{s_dict[n]}]' if n in s_dict else ''
-                    s.node(name=str(n.index), label = f'{{{{{ins}}}|{n.index}{io}\n{n.kind}\n{n.name}|{{{outs}}}}}', shape='record')
+                    color = '#f5f5f5' if n.kind == '__fork__' else '#cccccc'
+                    kind = '' if n.kind == '__fork__' else fr'\n{n.kind}'
+                    s.node(name=str(n.index), label = fr'{{{ins}{n.index}{io}{kind}\n{n.name}{outs}}}', shape='record', style='filled', fillcolor=color)
 
         for l in self.lines:
-            driver, reader = f'{l.driver.index}:o{l.driver_pin}', f'{l.reader.index}:i{l.reader_pin}'
+            driver = f'{l.driver.index}:o{l.driver_pin}' if len(l.driver.outs)>1 else f'{l.driver.index}'
-            if node_level[l.driver] >= node_level[l.reader]:
+            reader = f'{l.reader.index}:i{l.reader_pin}' if len(l.reader.ins)>1 else f'{l.reader.index}'
-                dot.edge(driver, reader, style='dotted', label=str(l.index))
+            label = str(line_labels.get(l, l.index))
-                pass
+            if node_level[l.driver] == node_level[l.reader]:
+                dot.node(f'_{l.index}_')
+                dot.edge(driver, f'_{l.index}_', style='dotted', label=label)
+                dot.edge(f'_{l.index}_', reader, style='dotted', label=label)
+            elif node_level[l.driver] > node_level[l.reader]:
+                dot.edge(driver, reader, style='dotted', label=label)
             else:
-                dot.edge(driver, reader, label=str(l.index))
+                dot.edge(driver, reader, label=label)
 
         return dot

src/kyupy/def_file.py

@@ -6,7 +6,7 @@ as an intermediate representation (:class:`DefFile` object).
 
 from collections import defaultdict
 
-from lark import Lark, Transformer, Tree
+from lark import Lark, Transformer
 
 from kyupy import readtext
 
@@ -85,6 +85,8 @@ class DefFile:
         self.pins = {}
         self.specialnets = {}
         self.nets = {}
+        self.design = ""
+        self.diearea = 0
 
 
 class DefTransformer(Transformer):
@@ -140,9 +142,11 @@ class DefTransformer(Transformer):
     def comp_stmt(self, args):
         name = args[0].value
         kind = args[1].value
-        point = args[2]
+        for aidx in range(2, len(args)):
-        orientation = args[3].value
+            if args[aidx] in ('PLACED', 'FIXED'):
-        self.def_file.components[name] = (kind, point, orientation)
+                point = args[aidx+1]
+                orientation = args[aidx+2].value
+                self.def_file.components[name] = (kind, point, orientation)
 
     def pins_stmt(self, args):
         pin = DefPin(args[0].value)
@@ -207,6 +211,7 @@ GRAMMAR = r"""
                 | /DIEAREA/ point+ ";"
                 | /ROW/ ID ID NUMBER NUMBER ID do_step ";"
                 | /TRACKS/ /[XY]/ NUMBER "DO" NUMBER "STEP" NUMBER "LAYER" ID ";"
+                | /GCELLGRID/ /[XY]/ NUMBER "DO" NUMBER "STEP" NUMBER ";"
                 | propdef | vias | nondef | comp | pins | pinprop | spnets | nets
 
     propdef: "PROPERTYDEFINITIONS" propdef_stmt* "END" "PROPERTYDEFINITIONS"
@@ -228,7 +233,9 @@ GRAMMAR = r"""
                         | "+" /VIA/ ID )* ";"
 
     comp: "COMPONENTS" NUMBER ";" comp_stmt* "END" "COMPONENTS"
-    comp_stmt: "-" ID ID "+" "PLACED" point ID ";"
+    comp_stmt: "-" ID ID ( "+" /PLACED/ point ID
+                         | "+" /FIXED/ point ID
+                         | "+" /SOURCE/ ID )* ";"
 
     pins: "PINS" NUMBER ";" pins_stmt* "END" "PINS"
     pins_stmt: "-" ID pins_opt* ";"
@@ -239,6 +246,8 @@ GRAMMAR = r"""
             | "+" /PORT/
             | "+" /LAYER/ ID point point
             | "+" /PLACED/ point ID
+            | "+" /FIXED/ point ID
+
 
     pinprop: "PINPROPERTIES" NUMBER ";" pinprop_stmt* "END" "PINPROPERTIES"
     pinprop_stmt: "-" "PIN" ID "+" "PROPERTY" ID STRING ";"
@@ -266,10 +275,10 @@ GRAMMAR = r"""
     wire: ID wire_opt points
     wire_opt: ( "TAPER" | "TAPERRULE" ID )? ("STYLE" ID)?
 
-    points: point ( point | points_via )+
+    points: point ( point | points_via_or_rect )+
-    points_via: ID ORIENTATION?
+    points_via_or_rect: ID ( ORIENTATION? | "(" (NUMBER|SIGNED_NUMBER|/\*/) (NUMBER|SIGNED_NUMBER|/\*/) (NUMBER|SIGNED_NUMBER|/\*/) (NUMBER|SIGNED_NUMBER|/\*/) ")" )
 
-    point: "(" (NUMBER|/\*/) (NUMBER|/\*/) NUMBER? ")"
+    point: "(" (NUMBER|SIGNED_NUMBER|/\*/) (NUMBER|SIGNED_NUMBER|/\*/) NUMBER? ")"
 
     do_step: "DO" NUMBER "BY" NUMBER "STEP" (NUMBER|SIGNED_NUMBER) (NUMBER|SIGNED_NUMBER)
 
@@ -284,9 +293,9 @@ GRAMMAR = r"""
     """
 
 
-def parse(text):
+def parse(text) -> DefFile:
     """Parses the given ``text`` and returns a :class:`DefFile` object."""
-    return Lark(GRAMMAR, parser="lalr", transformer=DefTransformer()).parse(text)
+    return Lark(GRAMMAR, parser="lalr", transformer=DefTransformer()).parse(text) # type: ignore
 
 
 def load(file):

src/kyupy/logic.py

@@ -47,8 +47,9 @@ The functions in this module use the ``mv...`` and ``bp...`` prefixes to signify
 from collections.abc import Iterable
 
 import numpy as np
+from numpy.typing import DTypeLike
 
-from . import numba, hr_bytes
+from . import numba, hr_bytes, NestedNumericList
 
 
 ZERO = 0b000
@@ -80,7 +81,30 @@ on a signal. ``'N'``, ``'n'``, and ``'v'`` are interpreted as ``NPULSE``.
 """
 
 
-def interpret(value):
+#                      0 X - 1 P R F N
+_mv_xor   = np.array([[0,1,1,3,4,5,6,7], # 0
+                      [1,1,1,1,1,1,1,1], # X
+                      [1,1,1,1,1,1,1,1], # -
+                      [3,1,1,0,7,6,5,4], # 1
+                      [4,1,1,7,4,5,6,7], # P
+                      [5,1,1,6,5,4,7,6], # R
+                      [6,1,1,5,6,7,4,5], # F
+                      [7,1,1,4,7,6,5,4], # N
+                      ], dtype=np.uint8)
+
+#                      0 X - 1 P R F N
+_mv_merge = np.array([[0,0,0,1,1,1,1,1], # 0
+                      [0,1,1,3,4,5,6,7], # X
+                      [0,1,2,3,4,5,6,7], # -
+                      [1,3,3,3,1,1,1,1], # 1
+                      [1,4,4,1,4,1,1,1], # P
+                      [1,5,5,1,1,5,1,1], # R
+                      [1,6,6,1,1,1,6,1], # F
+                      [1,7,7,1,1,1,1,7], # N
+                      ], dtype=np.uint8)
+
+
+def interpret(value) -> int|NestedNumericList:
     """Converts characters, strings, and lists of them to lists of logic constants defined above.
 
     :param value: A character (string of length 1), Boolean, Integer, None, or Iterable.
@@ -99,11 +123,12 @@ def interpret(value):
     return UNKNOWN
 
 
-def mvarray(*a):
+def mvarray(*a) -> np.ndarray[tuple[int]|tuple[int,int],np.dtype[np.uint8]]:
     """Converts (lists of) Boolean values or strings into a multi-valued array.
 
-    The given values are interpreted and the axes are arranged as per KyuPy's convention.
+    Iterable over characters, Booleans, Integers (0,1), None are interpreted as a single pattern.
-    Use this function to convert strings into multi-valued arrays.
+    Multi-character strings are Iterables over characters and are interpreted as a single pattern.
+    Lists of said patterns form a pattern set as ndarray with the axes are arranged as per KyuPy's convention.
     """
     mva = np.array(interpret(a), dtype=np.uint8)
     if mva.ndim < 2: return mva
@@ -114,7 +139,7 @@ def mvarray(*a):
 def mv_str(mva, delim='\n'):
     """Renders a given multi-valued array into a string.
     """
-    sa = np.choose(mva, np.array([*'0X-1PRFN'], dtype=np.unicode_))
+    sa = np.choose(mva, np.array([*'0X-1PRFN'], dtype=np.str_))
     if not hasattr(mva, 'ndim') or mva.ndim == 0: return sa
     if mva.ndim == 1: return ''.join(sa)
     return delim.join([''.join(c) for c in sa.swapaxes(-1,-2)])
@@ -190,17 +215,6 @@ def mv_and(x1, x2, out=None):
     return out
 
 
-def _mv_xor(out, *ins):
-    any_unknown = (ins[0] == UNKNOWN) | (ins[0] == UNASSIGNED)
-    for inp in ins[1:]: any_unknown |= (inp == UNKNOWN) | (inp == UNASSIGNED)
-
-    out[...] = ZERO
-    for inp in ins:
-        np.bitwise_xor(out, inp & 0b011, out=out)
-        np.bitwise_or(out, inp & 0b100, out=out)
-    np.putmask(out, any_unknown, UNKNOWN)
-
-
 def mv_xor(x1, x2, out=None):
     """A multi-valued XOR operator.
 
@@ -210,7 +224,25 @@ def mv_xor(x1, x2, out=None):
     :return: A multi-valued array with the result.
     """
     out = out or np.empty(np.broadcast(x1, x2).shape, dtype=np.uint8)
-    _mv_xor(out, x1, x2)
+    out[...] = _mv_xor[x1, x2]
+    return out
+
+
+def mv_merge(x1, x2, out=None):
+    """A multi-valued merge operator.
+
+    Merging identical values ``a`` and ``a`` yields ``a``.
+    Merging ``X`` and ``-`` yields ``X``.
+    Merging any value ``a`` with ``X`` or ``-`` yields ``a``.
+    Merging any other values ``a != b`` yields ``X``.
+
+    :param x1: A multi-valued array.
+    :param x2: A multi-valued array.
+    :param out: An optional storage destination. If None, a new multi-valued array is returned.
+    :return: A multi-valued array with the result.
+    """
+    out = out or np.empty(np.broadcast(x1, x2).shape, dtype=np.uint8)
+    out[...] = _mv_merge[x1, x2]
     return out
 
 
@@ -241,6 +273,8 @@ def mv_latch(d, t, q_prev, out=None):
 def mv_transition(init, final, out=None):
     """Computes the logic transitions from the initial values of ``init`` to the final values of ``final``.
     Pulses in the input data are ignored. If any of the inputs are ``UNKNOWN``, the result is ``UNKNOWN``.
+    If init is ``UNASSIGNED``, the result is the final value of ``final``.
+    If final is ``UNASSIGNED``, the result is the initial value of ``init``.
     If both inputs are ``UNASSIGNED``, the result is ``UNASSIGNED``.
 
     :param init: A multi-valued array.
@@ -251,7 +285,9 @@ def mv_transition(init, final, out=None):
     out = out or np.empty(np.broadcast(init, final).shape, dtype=np.uint8)
     out[...] = (init & 0b010) | (final & 0b001)
     out[...] |= ((out << 1) ^ (out << 2)) & 0b100
-    unknown = (init == UNKNOWN) | (init == UNASSIGNED) | (final == UNKNOWN) | (final == UNASSIGNED)
+    out[...] = np.choose(init == UNASSIGNED, [out, (final & 0b001) * ONE])
+    out[...] = np.choose(final == UNASSIGNED, [out, ((init & 0b010) >> 1) * ONE])
+    unknown = (init == UNKNOWN) | (final == UNKNOWN)
     unassigned = (init == UNASSIGNED) & (final == UNASSIGNED)
     np.putmask(out, unknown, UNKNOWN)
     np.putmask(out, unassigned, UNASSIGNED)
@@ -265,6 +301,18 @@ def mv_to_bp(mva):
     return np.packbits(unpackbits(mva)[...,:3], axis=-2, bitorder='little').swapaxes(-1,-2)
 
 
+def mv_init(mva):
+    """Returns the initial binary values for mva.
+    """
+    return (mva>>1) & ((mva>>2)|mva) & 1
+
+
+def mv_final(mva):
+    """Returns the final binary value of mva.
+    """
+    return mva & ((mva>>2)|(mva>>1)) & 1
+
+
 def bparray(*a):
     """Converts (lists of) Boolean values or strings into a bit-parallel array.
 
@@ -417,7 +465,7 @@ def unpackbits(a : np.ndarray):
     return np.unpackbits(a.view(np.uint8), bitorder='little').reshape(*a.shape, 8*a.itemsize)
 
 
-def packbits(a, dtype=np.uint8):
+def packbits(a, dtype:DTypeLike=np.uint8):
     """Packs the values of a boolean-valued array ``a`` along its last axis into bits.
 
     Similar to ``np.packbits``, but returns an array of given dtype and the shape of ``a`` with the last axis removed.

src/kyupy/logic_sim.py

@@ -1,21 +1,38 @@
-"""A high-throughput combinational logic simulator.
+"""High-throughput combinational logic simulators.
 
-The class :py:class:`~kyupy.logic_sim.LogicSim` performs parallel simulations of the combinational part of a circuit.
+These simulators take batches of input assignments to a logic circuit and compute batches of output values.
-The logic operations are performed bit-parallel on packed numpy arrays (see bit-parallel (bp) array description in :py:mod:`~kyupy.logic`).
+Each batch is simulated in data-parallel fashion using packed numpy arrays (see bit-parallel (bp) array description in :py:mod:`~kyupy.logic`) for each signal in the circuit.
-Simple sequential circuits can be simulated by repeated assignments and propagations.
+All simulations within a batch are mutually independent.
-However, this simulator ignores the clock network and simply assumes that all state-elements are clocked all the time.
+Inner simulation loops (that iterate over the topologically sorted circuit model) are just-in-time (JIT) compiled using numba.
+JIT compiling triggered by the first simulation may take several seconds.
+If numba is not available, inner loop is executed in the python interpreter and is significantly slower.
+
+Circuits passed to these simulators must contain only cells defined in the :py:class:`~kyupy.techlib.KYUPY` cell library.
+Other circuits must be mapped to that library first using :py:method:`~kyupy.circuit.resolve_tlib_cells()`.
+
+Sequential circuits are supported with some caveats.
+The simulation model is built by stripping out all flip-flops, clock, and asynchronous set/reset logic.
+Values stored in the flip-flops become part of the input assignments and output values (pseudo-primary I/O). 
+Only the combinational portion is actually simulated.
+Simulation results remain correct under the assumption that all flip-clops are clocked all the time and set/reset logic is passive.
 """
 
 import math
+import warnings
 
 import numpy as np
 
-from . import numba, logic, hr_bytes, sim
+from . import numba, logic, hr_bytes, sim, eng, cdiv, batchrange
-from .circuit import Circuit
+from .circuit import Circuit, Line
+
 
 class LogicSim(sim.SimOps):
     """A bit-parallel naïve combinational simulator for 2-, 4-, or 8-valued logic.
 
+    .. deprecated::
+        Use the specialized :py:class:`LogicSim2V`, :py:class:`LogicSim4V`, or :py:class:`LogicSim6V`
+        simulators instead.
+
     :param circuit: The circuit to simulate.
     :param sims: The number of parallel logic simulations to perform.
     :param m: The arity of the logic, must be 2, 4, or 8.
@@ -23,14 +40,18 @@ class LogicSim(sim.SimOps):
     :param strip_forks: If True, forks are not included in the simulation model to save memory and simulation time.
     """
     def __init__(self, circuit: Circuit, sims: int = 8, m: int = 8, c_reuse: bool = False, strip_forks: bool = False):
+        warnings.warn(
+            'LogicSim is deprecated; use LogicSim2V, LogicSim4V, or LogicSim6V instead.',
+            DeprecationWarning, stacklevel=2)
         assert m in [2, 4, 8]
         super().__init__(circuit, c_reuse=c_reuse, strip_forks=strip_forks)
         self.m = m
         self.mdim = math.ceil(math.log2(m))
         self.sims = sims
-        nbytes = (sims - 1) // 8 + 1
+        nbytes = cdiv(sims, 8)
 
         self.c = np.zeros((self.c_len, self.mdim, nbytes), dtype=np.uint8)
+        self.c_dirty = np.full(self.c_len, 1, dtype=np.uint8)
         self.s = np.zeros((2, self.s_len, 3, nbytes), dtype=np.uint8)
         """Logic values of the sequential elements (flip-flops) and ports.
 
@@ -42,16 +63,17 @@ class LogicSim(sim.SimOps):
         Access this array to assign new values to the (P)PIs or read values from the (P)POs.
         """
         self.s[:,:,1,:] = 255  # unassigned
+        self._full_mask = np.full(self.c.shape[-1], 255, dtype=np.uint8)
 
     def __repr__(self):
-        return f'{{name: "{self.circuit.name}", sims: {self.sims}, m: {self.m}, c_bytes: {self.c.nbytes}}}'
+        return f'{{name: "{self.circuit.name}", sims: {self.sims}, m: {self.m}, c_bytes: {eng(self.c.nbytes)}}}'
 
     def s_to_c(self):
         """Copies the values from ``s[0]`` the inputs of the combinational portion.
         """
         self.c[self.pippi_c_locs] = self.s[0, self.pippi_s_locs, :self.mdim]
 
-    def c_prop(self, inject_cb=None):
+    def c_prop(self, sims=None, inject_cb=None, fault_line:Line|int=-1, fault_mask=None, fault_model=2):
         """Propagate the input values through the combinational circuit towards the outputs.
 
         Performs all logic operations in topological order.
@@ -60,17 +82,26 @@ class LogicSim(sim.SimOps):
         :param inject_cb: A callback function for manipulating intermediate signal values.
             This function is called with a line and its new logic values (in bit-parallel format) after
             evaluation of a node. The callback may manipulate the given values in-place, the simulation
-            resumes with the manipulated values after the callback returns.
+            resumes with the manipulated values after the callback returns. Specifying this callback
+            may reduce performance as it disables jit compilation.
         :type inject_cb: ``f(Line, ndarray)``
         """
+        fault_line = int(fault_line)
+        if fault_mask is None:
+            fault_mask = self._full_mask  # default: full mask
+        else:
+            if len(fault_mask) < self.c.shape[-1]:  # pad mask with 0's if necessary
+                fault_mask2 = np.full(self.c.shape[-1], 0, dtype=np.uint8)
+                fault_mask2[:len(fault_mask)] = fault_mask
+                fault_mask = fault_mask2
         t0 = self.c_locs[self.tmp_idx]
         t1 = self.c_locs[self.tmp2_idx]
         if self.m == 2:
             if inject_cb is None:
-                _prop_cpu(self.ops, self.c_locs, self.c)
+                c_prop_2v_cpu(self.ops, self.c_locs, self.c, self.c_dirty, int(fault_line), fault_mask, int(fault_model))
             else:
-                for op, o0, i0, i1, i2, i3 in self.ops[:,:6]:
+                for op, o0l, i0l, i1l, i2l, i3l in self.ops[:,:6]:
-                    o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0, i0, i1, i2, i3)]
+                    o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0l, i0l, i1l, i2l, i3l)]
                     if op == sim.BUF1: self.c[o0]=self.c[i0]
                     elif op == sim.INV1: self.c[o0] = ~self.c[i0]
                     elif op == sim.AND2: self.c[o0] = self.c[i0] & self.c[i1]
@@ -105,10 +136,10 @@ class LogicSim(sim.SimOps):
                     elif op == sim.OAI211:self.c[o0] = ~((self.c[i0] | self.c[i1]) & self.c[i2] & self.c[i3])
                     elif op == sim.MUX21: self.c[o0] = (self.c[i0] & ~self.c[i2]) | (self.c[i1] & self.c[i2])
                     else: print(f'unknown op {op}')
-                    inject_cb(o0, self.s[o0])
+                    inject_cb(o0l, self.c[o0])
         elif self.m == 4:
-            for op, o0, i0, i1, i2, i3 in self.ops[:,:6]:
+            for op, o0l, i0l, i1l, i2l, i3l in self.ops[:,:6]:
-                o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0, i0, i1, i2, i3)]
+                o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0l, i0l, i1l, i2l, i3l)]
                 if op == sim.BUF1: self.c[o0]=self.c[i0]
                 elif op == sim.INV1: logic.bp4v_not(self.c[o0], self.c[i0])
                 elif op == sim.AND2: logic.bp4v_and(self.c[o0], self.c[i0], self.c[i1])
@@ -181,9 +212,19 @@ class LogicSim(sim.SimOps):
                     logic.bp4v_and(self.c[t1], self.c[i1], self.c[i2])
                     logic.bp4v_or(self.c[o0], self.c[t0], self.c[t1])
                 else: print(f'unknown op {op}')
+                if inject_cb is not None: inject_cb(o0l, self.c[o0])
+                if fault_line >= 0 and o0l == fault_line:
+                    if fault_model == 0:
+                        self.c[o0] = self.c[o0] & ~fault_mask[np.newaxis]
+                    elif fault_model == 1:
+                        self.c[o0] = self.c[o0] | fault_mask[np.newaxis]
+                    else:
+                        self.c[t0, 0] = ~(self.c[o0, 0] & self.c[o0, 1] & fault_mask)
+                        self.c[o0, 1] = ~self.c[o0, 0] & ~self.c[o0, 1] & fault_mask
+                        self.c[o0, 0] = self.c[t0, 0]
         else:
-            for op, o0, i0, i1, i2, i3 in self.ops[:,:6]:
+            for op, o0l, i0l, i1l, i2l, i3l in self.ops[:,:6]:
-                o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0, i0, i1, i2, i3)]
+                o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0l, i0l, i1l, i2l, i3l)]
                 if op == sim.BUF1: self.c[o0]=self.c[i0]
                 elif op == sim.INV1: logic.bp8v_not(self.c[o0], self.c[i0])
                 elif op == sim.AND2: logic.bp8v_and(self.c[o0], self.c[i0], self.c[i1])
@@ -256,7 +297,7 @@ class LogicSim(sim.SimOps):
                     logic.bp8v_and(self.c[t1], self.c[i1], self.c[i2])
                     logic.bp8v_or(self.c[o0], self.c[t0], self.c[t1])
                 else: print(f'unknown op {op}')
-                if inject_cb is not None: inject_cb(o0, self.s[o0])
+                if inject_cb is not None: inject_cb(o0l, self.c[o0])
 
     def c_to_s(self):
         """Copies (captures) the results of the combinational portion to ``s[1]``.
@@ -295,10 +336,113 @@ class LogicSim(sim.SimOps):
             self.s_ppo_to_ppi()
 
 
+class LogicSim2V(sim.SimOps):
+    """A bit-parallel naïve combinational simulator for 2-valued logic.
+
+    :param circuit: The circuit to simulate.
+    :param sims: The number of parallel logic simulations to perform.
+    :param c_reuse: If True, intermediate signal values may get overwritten when not needed anymore to save memory.
+    :param strip_forks: If True, forks are not included in the simulation model to save memory and simulation time.
+           Caveat: faults on fanout branches will not be injected if forks are stripped.
+    """
+    def __init__(self, circuit: Circuit, sims: int = 8, c_reuse: bool = False, strip_forks: bool = False):
+        super().__init__(circuit, c_reuse=c_reuse, strip_forks=strip_forks)
+        self.sims = sims
+        nbytes = cdiv(sims, 8)
+
+        self.c = np.zeros((self.c_len, 1, nbytes), dtype=np.uint8)
+        """Logic values within the combinational portion of the circuit.
+
+        In bit-parallel (bp) storage format.
+        Storage locations are indirectly addressed.
+        Data for line `l` is in `self.c[self.c_locs[l]]`.
+        """
+        self.c_dirty = np.full(self.c_len, 1, dtype=np.uint8)
+        """Marker for logic values that have changed recently for use in incremental simulation.
+
+        A node n will be evaluated by `c_prop()` if and only if at least one of its inputs is marked dirty (`c_dirty[c_locs[n.ins[i]]] == 1`).
+        If a node is evaluated, its output is marked dirty.
+        By default, this array is all-1, therefore every node is evaluated every time by `c_prop()`.
+        Enable incremental simulation by setting this array to all-0, change signal values in `c` and set `c_dirty` to 1 for all changed signals.
+        Next call to `c_prop()` will only evaluate nodes that read from the changed signals.
+        To restore the original simulation state, revert the signal value changes in `c`, and call `c_prop()` again.
+        Caveats:
+        - Injected changes that are located downstream of other changes will have no effect.
+        - Only works with `c_reuse=False` (default).
+        """
+        self.s_assign = np.zeros((self.s_len, self.sims), dtype=np.uint8)
+        """Logic values assigned to the ports and flip-flops of the circuit.
+
+        The simulator reads (P)PI values from here.
+        Values assigned to PO positions are ignored.
+        This field is a 2-dimensional array and expects values in the mv storage format.
+        First index is the port position (defined by `self.circuit.s_nodes`), second index is the pattern index (0 ... `self.sims-1`).
+        """
+        self.s_result = np.zeros((self.s_len, self.sims), dtype=np.uint8)
+        """Logic values at the ports and flip-flops of the circuit after simulation.
+
+        The simulator writes (P)PO values here.
+        Values assigned to PI positions are ignored.
+        This field is a 2-dimensional array and expects values in the mv storage format.
+        First index is the port position (defined by `self.circuit.s_nodes`), second index is the pattern index (0 ... `self.sims-1`).
+        """
+        self._full_mask = np.full(self.c.shape[-1], 255, dtype=np.uint8)
+
+    def __repr__(self):
+        return f'{{name: "{self.circuit.name}", sims: {self.sims}, c_bytes: {eng(self.c.nbytes)}}}'
+
+    def s_to_c(self):
+        """Assigns the values from ``self.s_assign`` to the inputs of the combinational portion.
+        """
+        self.c[self.pippi_c_locs] = logic.mv_to_bp(self.s_assign[self.pippi_s_locs])[:,:1,:]
+
+    def c_prop(self, fault_line=-1, fault_model=2, fault_mask=None):
+        if fault_mask is None:
+            fault_mask = self._full_mask  # default: full mask
+        else:
+            if len(fault_mask) < self.c.shape[-1]:  # pad mask with 0's if necessary
+                fault_mask2 = np.full(self.c.shape[-1], 0, dtype=np.uint8)
+                fault_mask2[:len(fault_mask)] = fault_mask
+                fault_mask = fault_mask2
+        c_prop_2v_cpu(self.ops, self.c_locs, self.c, self.c_dirty, int(fault_line), fault_mask, int(fault_model))
+
+    def c_to_s(self):
+        """Captures the results of the combinational portion into ``self.s_result``.
+        """
+        self.s_result[self.poppo_s_locs] = logic.bp_to_mv(self.c[self.poppo_c_locs])[:,:self.sims] * logic.ONE
+
+    def c_ppo_to_ppi(self):
+        """Copies the result data for all PPOs (flip-flops) to PPIs for a next simulation cycle.
+        """
+        self.c[self.ppi_c_locs] = self.c[self.ppo_c_locs].copy()  # copy prevents undefined behavior if (P)PIs are connected directly to (P)POs
+
+    def allocate(self):
+        """Allocates a new pattern array with appropriate dimensions ``(self.s_len, self.sims)`` for one-pass simulation.
+        """
+        return np.full((self.s_len, self.sims), logic.ZERO, dtype=np.uint8)
+
+    def simulate(self, patterns, responses = None, cycles:int = 1, fault_line=-1, fault_model=2, fault_mask=None):
+        assert cycles >= 1
+        for bo, bs in batchrange(patterns.shape[-1], self.sims):
+            self.s_assign[self.pippi_s_locs, :bs] = patterns[self.pippi_s_locs, bo:bo+bs]
+            self.s_to_c()
+            for cycle in range(cycles):
+                self.c_prop(fault_line=fault_line, fault_model=fault_model, fault_mask=fault_mask)
+                if cycle < (cycles-1): self.c_ppo_to_ppi()
+            self.c_to_s()
+            if responses is None:
+                patterns[self.poppo_s_locs, bo:bo+bs] = self.s_result[self.poppo_s_locs, :bs]
+            else:
+                responses[self.poppo_s_locs, bo:bo+bs] = self.s_result[self.poppo_s_locs, :bs]
+        return patterns if responses is None else responses
+
+
 @numba.njit
-def _prop_cpu(ops, c_locs, c):
+def c_prop_2v_cpu(ops, c_locs, c, c_dirty, fault_line, fault_mask, fault_model):
-    for op, o0, i0, i1, i2, i3 in ops[:,:6]:
+    for op, o0l, i0l, i1l, i2l, i3l in ops[:,:6]:
-        o0, i0, i1, i2, i3 = [c_locs[x] for x in (o0, i0, i1, i2, i3)]
+        o0, i0, i1, i2, i3 = [c_locs[x] for x in (o0l, i0l, i1l, i2l, i3l)]
+        if fault_line < 0 or o0l != fault_line:  # fault injection forces node evaluation
+            if not (c_dirty[i0] | c_dirty[i1] | c_dirty[i2] | c_dirty[i3]): continue
         if op == sim.BUF1: c[o0]=c[i0]
         elif op == sim.INV1: c[o0] = ~c[i0]
         elif op == sim.AND2: c[o0] = c[i0] & c[i1]
@@ -333,3 +477,519 @@ def _prop_cpu(ops, c_locs, c):
         elif op == sim.OAI211: c[o0] = ~((c[i0] | c[i1]) & c[i2] & c[i3])
         elif op == sim.MUX21: c[o0] = (c[i0] & ~c[i2]) | (c[i1] & c[i2])
         else: print(f'unknown op {op}')
+        if fault_line >= 0 and o0l == fault_line:
+            if fault_model == 0:
+                c[o0] = c[o0] & ~fault_mask
+            elif fault_model == 1:
+                c[o0] = c[o0] | fault_mask
+            elif fault_model == 2:
+                c[o0] = c[o0] ^ fault_mask
+        c_dirty[o0] = 1
+
+
+class LogicSim4V(sim.SimOps):
+    """A bit-parallel naïve combinational simulator for 4-valued logic.
+
+    Logic values: 0, 1, -, X
+
+    :param circuit: The circuit to simulate.
+    :param sims: The number of parallel logic simulations to perform.
+    :param c_reuse: If True, intermediate signal values may get overwritten when not needed anymore to save memory.
+    :param strip_forks: If True, forks are not included in the simulation model to save memory and simulation time.
+    """
+    def __init__(self, circuit: Circuit, sims: int = 8, c_reuse: bool = False, strip_forks: bool = False):
+        super().__init__(circuit, c_reuse=c_reuse, strip_forks=strip_forks)
+        self.sims = sims
+        nbytes = cdiv(sims, 8)
+
+        self.c = np.zeros((self.c_len, 2, nbytes), dtype=np.uint8)
+        """Logic values within the combinational portion of the circuit.
+
+        In bit-parallel (bp) storage format.
+        Storage locations are indirectly addressed.
+        Data for line `l` is in `self.c[self.c_locs[l]]`.
+        """
+        self.s_assign = np.full((self.s_len, self.sims), logic.UNASSIGNED, dtype=np.uint8)
+        """Logic values assigned to the ports and flip-flops of the circuit.
+        The simulator reads (P)PI values from here.
+        Values assigned to PO positions are ignored.
+        This field is a 2-dimensional array and expects values in the mv storage format.
+        First index is the port position (defined by `self.circuit.s_nodes`), second
+        index is the pattern index (0 ... `self.sims-1`).
+        """
+        self.s_result = np.full((self.s_len, self.sims), logic.UNASSIGNED, dtype=np.uint8)
+        """Logic values at the ports and flip-flops of the circuit after simulation.
+        The simulator writes (P)PO values here.
+        Values assigned to PI positions are ignored.
+        This field is a 2-dimensional array and expects values in the mv storage format.
+        First index is the port position (defined by `self.circuit.s_nodes`), second
+        index is the pattern index (0 ... `self.sims-1`).
+        """
+        self._full_mask = np.full(self.c.shape[-1], 255, dtype=np.uint8)
+
+    def __repr__(self):
+        return f'{{name: "{self.circuit.name}", sims: {self.sims}, s_len: {self.s_len}, c_bytes: {eng(self.c.nbytes)}}}'
+
+    def s_to_c(self):
+        """Assigns the values from ``self.s_assign`` to the inputs of the combinational portion.
+        """
+        self.c[self.pippi_c_locs] = logic.mv_to_bp(self.s_assign[self.pippi_s_locs])[:,:2,:]
+
+    def c_prop(self, fault_line=-1, fault_model=2, fault_mask=None):
+        if fault_mask is None:
+            fault_mask = self._full_mask  # default: full mask
+        else:
+            if len(fault_mask) < self.c.shape[-1]:  # pad mask with 0's if necessary
+                fault_mask2 = np.full(self.c.shape[-1], 0, dtype=np.uint8)
+                fault_mask2[:len(fault_mask)] = fault_mask
+                fault_mask = fault_mask2
+        c_prop_4v_cpu(self.ops, self.c_locs, self.c, int(fault_line), fault_mask, int(fault_model), self.tmp_idx, self.tmp2_idx)
+
+    def c_to_s(self):
+        """Captures the results of the combinational portion into ``self.s_result``.
+        """
+        self.s_result[self.poppo_s_locs] = logic.bp_to_mv(self.c[self.poppo_c_locs])[:,:self.sims]
+
+    def c_ppo_to_ppi(self):
+        """Copies the result data for all PPOs (flip-flops) to PPIs for a next simulation cycle.
+        """
+        self.c[self.ppi_c_locs] = self.c[self.ppo_c_locs].copy()  # copy prevents undefined behavior if (P)PIs are connected directly to (P)POs
+
+    def allocate(self):
+        """Allocates a new pattern array with appropriate dimensions ``(self.s_len, self.sims)`` for one-pass simulation.
+        """
+        return np.full((self.s_len, self.sims), logic.ZERO, dtype=np.uint8)
+
+    def simulate(self, patterns, responses = None, cycles:int = 1, fault_line=-1, fault_model=2, fault_mask=None):
+        assert cycles >= 1
+        for bo, bs in batchrange(patterns.shape[-1], self.sims):
+            self.s_assign[self.pippi_s_locs, :bs] = patterns[self.pippi_s_locs, bo:bo+bs]
+            self.s_to_c()
+            for cycle in range(cycles):
+                self.c_prop(fault_line=fault_line, fault_model=fault_model, fault_mask=fault_mask)
+                if cycle < (cycles-1): self.c_ppo_to_ppi()
+            self.c_to_s()
+            if responses is None:
+                patterns[self.poppo_s_locs, bo:bo+bs] = self.s_result[self.poppo_s_locs, :bs]
+            else:
+                responses[self.poppo_s_locs, bo:bo+bs] = self.s_result[self.poppo_s_locs, :bs]
+        return patterns if responses is None else responses
+
+
+@numba.njit
+def c_prop_4v_cpu(ops, c_locs, c, fault_line, fault_mask, fault_model, t0_idx, t1_idx):
+    t0 = c[c_locs[t0_idx]]
+    t1 = c[c_locs[t1_idx]]
+    for op, o0l, i0l, i1l, i2l, i3l in ops[:,:6]:
+        o0, i0, i1, i2, i3 = [c[c_locs[x]] for x in (o0l, i0l, i1l, i2l, i3l)]
+        if op == sim.BUF1 or op == sim.INV1:
+            # i0[0,1] -> o0[0,1]
+            # 0  0 0  -> 0  0 0
+            # -  0 1  -> X  1 0
+            # X  1 0  -> X  1 0
+            # 1  1 1  -> 1  1 1
+            o0[0] = i0[0] | i0[1]
+            o0[1] = i0[0] & i0[1]
+        elif op == sim.AND2 or op == sim.NAND2:
+            # i1[0,1] * i0[0,1] -> o0[0,1]
+            # 0  0 0    0  0 0  -> 0  0 0
+            # 0  0 0    -  0 1  -> 0  0 0
+            # 0  0 0    X  1 0  -> 0  0 0
+            # 0  0 0    1  1 1  -> 0  0 0
+            # -  0 1    0  0 0  -> 0  0 0
+            # -  0 1    -  0 1  -> X  1 0
+            # -  0 1    X  1 0  -> X  1 0
+            # -  0 1    1  1 1  -> X  1 0
+            # X  1 0    0  0 0  -> 0  0 0
+            # X  1 0    -  0 1  -> X  1 0
+            # X  1 0    X  1 0  -> X  1 0
+            # X  1 0    1  1 1  -> X  1 0
+            # 1  1 1    0  0 0  -> 0  0 0
+            # 1  1 1    -  0 1  -> X  1 0
+            # 1  1 1    X  1 0  -> X  1 0
+            # 1  1 1    1  1 1  -> 1  1 1
+            o0[0] = (i0[0]|i0[1]) & (i1[0]|i1[1])
+            o0[1] = i0[0] & i0[1] & i1[0] & i1[1]
+        elif op == sim.AND3 or op == sim.NAND3:
+            # i2[0,1] * i1[0,1] * i0[0,1] -> o0[0,1]
+            o0[0] = (i0[0]|i0[1]) & (i1[0]|i1[1]) & (i2[0]|i2[1])
+            o0[1] = i0[0] & i0[1] & i1[0] & i1[1] & i2[0] & i2[1]
+        elif op == sim.AND4 or op == sim.NAND4:
+            # i3[0,1] * i2[0,1] * i1[0,1] * i0[0,1] -> o0[0,1]
+            o0[0] = (i0[0]|i0[1]) & (i1[0]|i1[1]) & (i2[0]|i2[1]) & (i3[0]|i3[1])
+            o0[1] = i0[0] & i0[1] & i1[0] & i1[1] & i2[0] & i2[1] & i3[0] & i3[1]
+        elif op == sim.OR2 or op == sim.NOR2:
+            # i1[0,1] + i0[0,1] -> o0[0,1]
+            # 0  0 0    0  0 0  -> 0  0 0
+            # 0  0 0    -  0 1  -> X  1 0
+            # 0  0 0    X  1 0  -> X  1 0
+            # 0  0 0    1  1 1  -> 1  1 1
+            # -  0 1    0  0 0  -> X  1 0
+            # -  0 1    -  0 1  -> X  1 0
+            # -  0 1    X  1 0  -> X  1 0
+            # -  0 1    1  1 1  -> 1  1 1
+            # X  1 0    0  0 0  -> X  1 0
+            # X  1 0    -  0 1  -> X  1 0
+            # X  1 0    X  1 0  -> X  1 0
+            # X  1 0    1  1 1  -> 1  1 1
+            # 1  1 1    0  0 0  -> 1  1 1
+            # 1  1 1    -  0 1  -> 1  1 1
+            # 1  1 1    X  1 0  -> 1  1 1
+            # 1  1 1    1  1 1  -> 1  1 1
+            o0[0] = i0[0] | i0[1] | i1[0] | i1[1]
+            o0[1] = (i0[0]&i0[1]) | (i1[0]&i1[1])
+        elif op == sim.OR3 or op == sim.NOR3:
+            # i2[0,1] + i1[0,1] + i0[0,1] -> o0[0,1]
+            o0[0] = i0[0] | i0[1] | i1[0] | i1[1] | i2[0] | i2[1]
+            o0[1] = (i0[0]&i0[1]) | (i1[0]&i1[1]) | (i2[0]&i2[1])
+        elif op == sim.OR4 or op == sim.NOR4:
+            # i3[0,1] + i2[0,1] + i1[0,1] + i0[0,1] -> o0[0,1]
+            o0[0] = i0[0] | i0[1] | i1[0] | i1[1] | i2[0] | i2[1] | i3[0] | i3[1]
+            o0[1] = (i0[0]&i0[1]) | (i1[0]&i1[1]) | (i2[0]&i2[1]) | (i3[0]&i3[1])
+        elif op == sim.XOR2 or op == sim.XNOR2:
+            # i1[0,1] ^ i0[0,1] -> o0[0,1]
+            # 0  0 0    0  0 0  -> 0  0 0
+            # 0  0 0    -  0 1  -> X  1 0
+            # 0  0 0    X  1 0  -> X  1 0
+            # 0  0 0    1  1 1  -> 1  1 1
+            # -  0 1    0  0 0  -> X  1 0
+            # -  0 1    -  0 1  -> X  1 0
+            # -  0 1    X  1 0  -> X  1 0
+            # -  0 1    1  1 1  -> X  1 0
+            # X  1 0    0  0 0  -> X  1 0
+            # X  1 0    -  0 1  -> X  1 0
+            # X  1 0    X  1 0  -> X  1 0
+            # X  1 0    1  1 1  -> X  1 0
+            # 1  1 1    0  0 0  -> 1  1 1
+            # 1  1 1    -  0 1  -> X  1 0
+            # 1  1 1    X  1 0  -> X  1 0
+            # 1  1 1    1  1 1  -> 0  0 0
+            o0[0] = (i0[0]|i0[1]|i1[0]|i1[1]) & (~i0[0]|~i0[1]|~i1[0]|~i1[1])
+            o0[1] = (~i0[0]&~i0[1]&i1[0]&i1[1]) | (i0[0]&i0[1]&~i1[0]&~i1[1])
+        elif op == sim.XOR3 or op == sim.XNOR3:
+            # i2[0,1] ^ i1[0,1] ^ i0[0,1] -> o0[0,1]
+            t0[0] = (i0[0]|i0[1]|i1[0]|i1[1]) & (~i0[0]|~i0[1]|~i1[0]|~i1[1])
+            t0[1] = (~i0[0]&~i0[1]&i1[0]&i1[1]) | (i0[0]&i0[1]&~i1[0]&~i1[1])
+            o0[0] = (t0[0]|t0[1]|i2[0]|i2[1]) & (~t0[0]|~t0[1]|~i2[0]|~i2[1])
+            o0[1] = (~t0[0]&~t0[1]&i2[0]&i2[1]) | (t0[0]&t0[1]&~i2[0]&~i2[1])
+        elif op == sim.XOR4 or op == sim.XNOR4:
+            # i3[0,1] ^ i2[0,1] ^ i1[0,1] ^ i0[0,1] -> o0[0,1]
+            t0[0] = (i0[0]|i0[1]|i1[0]|i1[1]) & (~i0[0]|~i0[1]|~i1[0]|~i1[1])
+            t0[1] = (~i0[0]&~i0[1]&i1[0]&i1[1]) | (i0[0]&i0[1]&~i1[0]&~i1[1])
+            t1[0] = (t0[0]|t0[1]|i2[0]|i2[1]) & (~t0[0]|~t0[1]|~i2[0]|~i2[1])
+            t1[1] = (~t0[0]&~t0[1]&i2[0]&i2[1]) | (t0[0]&t0[1]&~i2[0]&~i2[1])
+            o0[0] = (t1[0]|t1[1]|i3[0]|i3[1]) & (~t1[0]|~t1[1]|~i3[0]|~i3[1])
+            o0[1] = (~t1[0]&~t1[1]&i3[0]&i3[1]) | (t1[0]&t1[1]&~i3[0]&~i3[1])
+        elif op == sim.AO21 or op == sim.AOI21:
+            # (i0[0,1] * i1[0,1]) + i2[0,1] -> o0[0,1]
+            t0[0] = (i0[0]|i0[1]) & (i1[0]|i1[1])
+            t0[1] = i0[0] & i0[1] & i1[0] & i1[1]
+            o0[0] = t0[0] | t0[1] | i2[0] | i2[1]
+            o0[1] = (t0[0]&t0[1]) | (i2[0]&i2[1])
+        elif op == sim.OA21 or op == sim.OAI21:
+            # (i0[0,1] + i1[0,1]) * i2[0,1] -> o0[0,1]
+            t0[0] = i0[0] | i0[1] | i1[0] | i1[1]
+            t0[1] = (i0[0]&i0[1]) | (i1[0]&i1[1])
+            o0[0] = (t0[0]|t0[1]) & (i2[0]|i2[1])
+            o0[1] = t0[0] & t0[1] & i2[0] & i2[1]
+        elif op == sim.AO22 or op == sim.AOI22:
+            # (i0[0,1] * i1[0,1]) + (i2[0,1] * i3[0,1]) -> o0[0,1]
+            t0[0] = (i0[0]|i0[1]) & (i1[0]|i1[1])
+            t0[1] = i0[0] & i0[1] & i1[0] & i1[1]
+            t1[0] = (i2[0]|i2[1]) & (i3[0]|i3[1])
+            t1[1] = i2[0] & i2[1] & i3[0] & i3[1]
+            o0[0] = t0[0] | t0[1] | t1[0] | t1[1]
+            o0[1] = (t0[0]&t0[1]) | (t1[0]&t1[1])
+        elif op == sim.OA22 or op == sim.OAI22:
+            # (i0[0,1] + i1[0,1]) * (i2[0,1] + i3[0,1]) -> o0[0,1]
+            t0[0] = i0[0] | i0[1] | i1[0] | i1[1]
+            t0[1] = (i0[0]&i0[1]) | (i1[0]&i1[1])
+            t1[0] = i2[0] | i2[1] | i3[0] | i3[1]
+            t1[1] = (i2[0]&i2[1]) | (i3[0]&i3[1])
+            o0[0] = (t0[0]|t0[1]) & (t1[0]|t1[1])
+            o0[1] = t0[0] & t0[1] & t1[0] & t1[1]
+        elif op == sim.AO211 or op == sim.AOI211:
+            # (i0[0,1] * i1[0,1]) + i2[0,1] + i3[0,1] -> o0[0,1]
+            t0[0] = (i0[0]|i0[1]) & (i1[0]|i1[1])
+            t0[1] = i0[0] & i0[1] & i1[0] & i1[1]
+            o0[0] = t0[0] | t0[1] | i2[0] | i2[1] | i3[0] | i3[1]
+            o0[1] = (t0[0]&t0[1]) | (i2[0]&i2[1]) | (i3[0]&i3[1])
+        elif op == sim.OA211 or op == sim.OAI211:
+            # (i0[0,1] + i1[0,1]) * i2[0,1] * i3[0,1] -> o0[0,1]
+            t0[0] = i0[0] | i0[1] | i1[0] | i1[1]
+            t0[1] = (i0[0]&i0[1]) | (i1[0]&i1[1])
+            o0[0] = (t0[0]|t0[1]) & (i2[0]|i2[1]) & (i3[0]|i3[1])
+            o0[1] = t0[0] & t0[1] & i2[0] & i2[1] & i3[0] & i3[1]
+        elif op == sim.MUX21:
+            # t1[0,1] = ~i2[0,1]
+            t1[0] = ~i2[0] | ~i2[1]
+            t1[1] = ~i2[0] & ~i2[1]
+            # t0 = i0 * t1
+            t0[0] = (i0[0]|i0[1]) & (t1[0]|t1[1])
+            t0[1] = i0[0] & i0[1] & t1[0] & t1[1]
+            # t1 = i1 * i2
+            t1[0] = (i1[0]|i1[1]) & (i2[0]|i2[1])
+            t1[1] = i1[0] & i1[1] & i2[0] & i2[1]
+            # o0 = t0 | t1
+            o0[0] = t0[0] | t0[1] | t1[0] | t1[1]
+            o0[1] = (t0[0]&t0[1]) | (t1[0]&t1[1])
+        else: print(f'unknown op {op}')
+        if (op == sim.INV1 or
+            op == sim.NAND2 or op == sim.NAND3 or op == sim.NAND4 or
+            op == sim.NOR2 or op == sim.NOR3 or op == sim.NOR4 or
+            op == sim.XNOR2 or op == sim.XNOR3 or op == sim.XNOR4 or
+            op == sim.AOI21 or op == sim.OAI21 or
+            op == sim.AOI22 or op == sim.OAI22 or
+            op == sim.AOI211 or op == sim.OAI211):
+            # o0[0,1] -> o0[0,1]
+            # 0  0 0  -> 1  1 1
+            # -  0 1  -> X  1 0
+            # X  1 0  -> X  1 0
+            # 1  1 1  -> 0  0 0
+            t0[0] = ~o0[0] | ~o0[1]
+            o0[1] = ~o0[0] & ~o0[1]
+            o0[0] = t0[0]
+        if fault_line >= 0 and o0l == fault_line:
+            if fault_model == 0:
+                o0[...] = o0 & ~fault_mask
+            elif fault_model == 1:
+                o0[...] = o0 | fault_mask
+            else:
+                t0[0] = ~(o0[0] ^ o0[1])  # mask for 0 and 1
+                o0[...] = o0 ^ (fault_mask & t0[0])  # leaves X and - untouched
+
+
+class LogicSim6V(sim.SimOps):
+    """A bit-parallel naïve combinational simulator for 6-valued logic.
+
+    Logic values: 0, 1, R, F, N, P
+
+    :param circuit: The circuit to simulate.
+    :param sims: The number of parallel logic simulations to perform.
+    :param c_reuse: If True, intermediate signal values may get overwritten when not needed anymore to save memory.
+    :param strip_forks: If True, forks are not included in the simulation model to save memory and simulation time.
+    """
+    def __init__(self, circuit: Circuit, sims: int = 8, c_reuse: bool = False, strip_forks: bool = False):
+        super().__init__(circuit, c_reuse=c_reuse, strip_forks=strip_forks)
+        self.sims = sims
+        nbytes = cdiv(sims, 8)
+
+        self.c = np.zeros((self.c_len, 3, nbytes), dtype=np.uint8)
+        self.s = np.full((2, self.s_len, self.sims), logic.UNASSIGNED, dtype=np.uint8)
+        """Logic values of the sequential elements (flip-flops) and ports.
+
+        It is a pair of arrays in mv storage format:
+
+        * ``s[0]`` Assigned values. Simulator will read (P)PI value from here.
+        * ``s[1]`` Result values. Simulator will write (P)PO values here.
+
+        Access this array to assign new values to the (P)PIs or read values from the (P)POs.
+        """
+        self.s_assign = self.s[0]
+        self.s_result = self.s[1]
+        self._full_mask = np.full(self.c.shape[-1], 255, dtype=np.uint8)
+
+    def __repr__(self):
+        return f'{{name: "{self.circuit.name}", sims: {self.sims}, c_bytes: {eng(self.c.nbytes)}}}'
+
+    def s_to_c(self):
+        """Assigns the values from ``s[0]`` to the inputs of the combinational portion.
+        """
+        self.c[self.pippi_c_locs] = logic.mv_to_bp(self.s[0, self.pippi_s_locs])
+
+    def c_prop(self, fault_line=-1, fault_model=2, fault_mask=None):
+        if fault_mask is None:
+            fault_mask = self._full_mask  # default: full mask
+        else:
+            if len(fault_mask) < self.c.shape[-1]:  # pad mask with 0's if necessary
+                fault_mask2 = np.full(self.c.shape[-1], 0, dtype=np.uint8)
+                fault_mask2[:len(fault_mask)] = fault_mask
+                fault_mask = fault_mask2
+        c_prop_6v_cpu(self.ops, self.c, self.c_locs, fault_line, fault_mask, fault_model, self.tmp_idx, self.tmp2_idx)
+
+    def c_to_s(self):
+        """Captures the results of the combinational portion into ``s[1]``.
+        """
+        self.s[1, self.poppo_s_locs] = logic.bp_to_mv(self.c[self.poppo_c_locs])[:,:self.sims]
+
+    def c_ppo_to_ppi(self):
+        """Copies the result data for all PPOs (flip-flops) to PPIs for a next simulation cycle.
+        """
+        self.c[self.ppi_c_locs] = self.c[self.ppo_c_locs].copy()  # copy prevents undefined behavior if (P)PIs are connected directly to (P)POs
+
+    def allocate(self):
+        """Allocates a new pattern array with appropriate dimensions ``(self.s_len, self.sims)`` for one-pass simulation.
+        """
+        return np.full((self.s_len, self.sims), logic.ZERO, dtype=np.uint8)
+
+    def simulate(self, patterns, responses = None, cycles:int = 1, fault_line=-1, fault_model=2, fault_mask=None):
+        assert cycles >= 1
+        for bo, bs in batchrange(patterns.shape[-1], self.sims):
+            self.s_assign[self.pippi_s_locs, :bs] = patterns[self.pippi_s_locs, bo:bo+bs]
+            self.s_to_c()
+            for cycle in range(cycles):
+                self.c_prop(fault_line=fault_line, fault_model=fault_model, fault_mask=fault_mask)
+                if cycle < (cycles-1): self.c_ppo_to_ppi()
+            self.c_to_s()
+            if responses is None:
+                patterns[self.poppo_s_locs, bo:bo+bs] = self.s_result[self.poppo_s_locs, :bs]
+            else:
+                responses[self.poppo_s_locs, bo:bo+bs] = self.s_result[self.poppo_s_locs, :bs]
+        return patterns
+
+
+@numba.njit
+def c_prop_6v_cpu(ops, c, c_locs, fault_line, fault_mask, fault_model, tmp_idx, tmp2_idx):
+    t0 = c[c_locs[tmp_idx]]
+    t1 = c[c_locs[tmp2_idx]]
+    inv_op = np.array([255, 255, 0], dtype=np.uint8)[np.newaxis, :, np.newaxis]
+    for op, o0l, i0l, i1l, i2l, i3l in ops[:,:6]:
+        o0, i0, i1, i2, i3 = [c[c_locs[x]] for x in (o0l, i0l, i1l, i2l, i3l)]
+        if op == sim.BUF1 or op == sim.INV1:
+            o0[...] = i0
+        elif op == sim.AND2 or op == sim.NAND2:
+            o0[0] = i0[0] & i1[0]
+            o0[1] = i0[1] & i1[1]
+            o0[2] = (i0[2]&(i1[0]|i1[1]|i1[2])|
+                     i1[2]&(i0[0]|i0[1]|i0[2]))
+        elif op == sim.AND3 or op == sim.NAND3:
+            o0[0] = i0[0] & i1[0] & i2[0]
+            o0[1] = i0[1] & i1[1] & i2[1]
+            o0[2] = (i0[2]&(i1[0]|i1[1]|i1[2])&(i2[0]|i2[1]|i2[2])|
+                     i1[2]&(i0[0]|i0[1]|i0[2])&(i2[0]|i2[1]|i2[2])|
+                     i2[2]&(i0[0]|i0[1]|i0[2])&(i1[0]|i1[1]|i1[2]))
+        elif op == sim.AND4 or op == sim.NAND4:
+            o0[0] = i0[0] & i1[0] & i2[0] & i3[0]
+            o0[1] = i0[1] & i1[1] & i2[1] & i3[1]
+            o0[2] = (i0[2]&(i1[0]|i1[1]|i1[2])&(i2[0]|i2[1]|i2[2])&(i3[0]|i3[1]|i3[2])|
+                     i1[2]&(i0[0]|i0[1]|i0[2])&(i2[0]|i2[1]|i2[2])&(i3[0]|i3[1]|i3[2])|
+                     i2[2]&(i0[0]|i0[1]|i0[2])&(i1[0]|i1[1]|i1[2])&(i3[0]|i3[1]|i3[2])|
+                     i3[2]&(i0[0]|i0[1]|i0[2])&(i1[0]|i1[1]|i1[2])&(i2[0]|i2[1]|i2[2]))
+        elif op == sim.OR2 or op == sim.NOR2:
+            o0[0] = i0[0] | i1[0]
+            o0[1] = i0[1] | i1[1]
+            o0[2] = (i0[2]&(~i1[0]|~i1[1]|i1[2])|
+                     i1[2]&(~i0[0]|~i0[1]|i0[2]))
+        elif op == sim.OR3 or op == sim.NOR3:
+            o0[0] = i0[0] | i1[0] | i2[0]
+            o0[1] = i0[1] | i1[1] | i2[1]
+            o0[2] = (i0[2]&(~i1[0]|~i1[1]|i1[2])&(~i2[0]|~i2[1]|i2[2])|
+                     i1[2]&(~i0[0]|~i0[1]|i0[2])&(~i2[0]|~i2[1]|i2[2])|
+                     i2[2]&(~i0[0]|~i0[1]|i0[2])&(~i1[0]|~i1[1]|i1[2]))
+        elif op == sim.OR4 or op == sim.NOR4:
+            o0[0] = i0[0] | i1[0] | i2[0] | i3[0]
+            o0[1] = i0[1] | i1[1] | i2[1] | i3[1]
+            o0[2] = (i0[2]&(~i1[0]|~i1[1]|i1[2])&(~i2[0]|~i2[1]|i2[2])&(~i3[0]|~i3[1]|i3[2])|
+                     i1[2]&(~i0[0]|~i0[1]|i0[2])&(~i2[0]|~i2[1]|i2[2])&(~i3[0]|~i3[1]|i3[2])|
+                     i2[2]&(~i0[0]|~i0[1]|i0[2])&(~i1[0]|~i1[1]|i1[2])&(~i3[0]|~i3[1]|i3[2])|
+                     i3[2]&(~i0[0]|~i0[1]|i0[2])&(~i1[0]|~i1[1]|i1[2])&(~i2[0]|~i2[1]|i2[2]))
+        elif op == sim.XOR2 or op == sim.XNOR2:
+            o0[0] = i0[0] ^ i1[0]
+            o0[1] = i0[1] ^ i1[1]
+            o0[2] = i0[2] | i1[2]
+        elif op == sim.XOR3 or op == sim.XNOR3:
+            o0[0] = i0[0] ^ i1[0] ^ i2[0]
+            o0[1] = i0[1] ^ i1[1] ^ i2[1]
+            o0[2] = i0[2] | i1[2] | i2[2]
+        elif op == sim.XOR4 or op == sim.XNOR4:
+            o0[0] = i0[0] ^ i1[0] ^ i2[0] ^ i3[0]
+            o0[1] = i0[1] ^ i1[1] ^ i2[1] ^ i3[1]
+            o0[2] = i0[2] | i1[2] | i2[2] | i3[2]
+        elif op == sim.AO21 or op == sim.AOI21:
+            # (i0[0,1] * i1[0,1]) + i2[0,1] -> o0[0,1]
+            t0[0] = i0[0] & i1[0]
+            t0[1] = i0[1] & i1[1]
+            t0[2] = (i0[2]&(i1[0]|i1[1]|i1[2])|
+                     i1[2]&(i0[0]|i0[1]|i0[2]))
+            o0[0] = t0[0] | i2[0]
+            o0[1] = t0[1] | i2[1]
+            o0[2] = (t0[2]&(~i2[0]|~i2[1]|i2[2])|
+                     i2[2]&(~t0[0]|~t0[1]|t0[2]))
+        elif op == sim.OA21 or op == sim.OAI21:
+            # (i0[0,1] + i1[0,1]) * i2[0,1] -> o0[0,1]
+            t0[0] = i0[0] | i1[0]
+            t0[1] = i0[1] | i1[1]
+            t0[2] = (i0[2]&(~i1[0]|~i1[1]|i1[2])|
+                     i1[2]&(~i0[0]|~i0[1]|i0[2]))
+            o0[0] = t0[0] & i2[0]
+            o0[1] = t0[1] & i2[1]
+            o0[2] = (t0[2]&(i2[0]|i2[1]|i2[2])|
+                     i2[2]&(t0[0]|t0[1]|t0[2]))
+        elif op == sim.AO22 or op == sim.AOI22:
+            # (i0[0,1] * i1[0,1]) + (i2[0,1] * i3[0,1]) -> o0[0,1]
+            t0[0] = i0[0] & i1[0]
+            t0[1] = i0[1] & i1[1]
+            t0[2] = (i0[2]&(i1[0]|i1[1]|i1[2])|
+                     i1[2]&(i0[0]|i0[1]|i0[2]))
+            t1[0] = i2[0] & i3[0]
+            t1[1] = i2[1] & i3[1]
+            t1[2] = (i2[2]&(i3[0]|i3[1]|i3[2])|
+                     i3[2]&(i2[0]|i2[1]|i2[2]))
+            o0[0] = t0[0] | t1[0]
+            o0[1] = t0[1] | t1[1]
+            o0[2] = (t0[2]&(~t1[0]|~t1[1]|t1[2])|
+                     t1[2]&(~t0[0]|~t0[1]|t0[2]))
+        elif op == sim.OA22 or op == sim.OAI22:
+            # (i0[0,1] + i1[0,1]) * (i2[0,1] + i3[0,1]) -> o0[0,1]
+            t0[0] = i0[0] | i1[0]
+            t0[1] = i0[1] | i1[1]
+            t0[2] = (i0[2]&(~i1[0]|~i1[1]|i1[2])|
+                     i1[2]&(~i0[0]|~i0[1]|i0[2]))
+            t1[0] = i2[0] | i3[0]
+            t1[1] = i2[1] | i3[1]
+            t1[2] = (i2[2]&(~i3[0]|~i3[1]|i3[2])|
+                     i3[2]&(~i2[0]|~i2[1]|i2[2]))
+            o0[0] = t0[0] & t1[0]
+            o0[1] = t0[1] & t1[1]
+            o0[2] = (t0[2]&(t1[0]|t1[1]|t1[2])|
+                     t1[2]&(t0[0]|t0[1]|t0[2]))
+        elif op == sim.AO211 or op == sim.AOI211:
+            # (i0[0,1] * i1[0,1]) + i2[0,1] + i3[0,1] -> o0[0,1]
+            t0[0] = i0[0] & i1[0]
+            t0[1] = i0[1] & i1[1]
+            t0[2] = (i0[2]&(i1[0]|i1[1]|i1[2])|
+                     i1[2]&(i0[0]|i0[1]|i0[2]))
+            o0[0] = t0[0] | i2[0] | i3[0]
+            o0[1] = t0[1] | i2[1] | i3[1]
+            o0[2] = (t0[2]&(~i2[0]|~i2[1]|i2[2])&(~i3[0]|~i3[1]|i3[2])|
+                     i2[2]&(~t0[0]|~t0[1]|t0[2])&(~i3[0]|~i3[1]|i3[2])|
+                     i3[2]&(~t0[0]|~t0[1]|t0[2])&(~i2[0]|~i2[1]|i2[2]))
+        elif op == sim.OA211 or op == sim.OAI211:
+            # (i0[0,1] + i1[0,1]) * i2[0,1] * i3[0,1] -> o0[0,1]
+            t0[0] = i0[0] | i1[0]
+            t0[1] = i0[1] | i1[1]
+            t0[2] = (i0[2]&(~i1[0]|~i1[1]|i1[2])|
+                     i1[2]&(~i0[0]|~i0[1]|i0[2]))
+            o0[0] = t0[0] & i2[0] & i3[0]
+            o0[1] = t0[1] & i2[1] & i3[1]
+            o0[2] = (t0[2]&(i2[0]|i2[1]|i2[2])&(i3[0]|i3[1]|i3[2])|
+                     i2[2]&(t0[0]|t0[1]|t0[2])&(i3[0]|i3[1]|i3[2])|
+                     i3[2]&(t0[0]|t0[1]|t0[2])&(i2[0]|i2[1]|i2[2]))
+        elif op == sim.MUX21:
+            # t1 = ~i2
+            t1[...] = i2 ^ inv_op
+            # t0 = i0 & t1
+            t0[0] = i0[0] & t1[0]
+            t0[1] = i0[1] & t1[1]
+            t0[2] = (i0[2]&(t1[0]|t1[1]|t1[2])|
+                     t1[2]&(i0[0]|i0[1]|i0[2]))
+            # t1 = i1 & i2
+            t1[0] = i1[0] & i2[0]
+            t1[1] = i1[1] & i2[1]
+            t1[2] = (i1[2]&(i2[0]|i2[1]|i2[2])|
+                     i2[2]&(i1[0]|i1[1]|i1[2]))
+            # o0 = t0 | t1
+            o0[0] = t0[0] | t1[0]
+            o0[1] = t0[1] | t1[1]
+            o0[2] = (t0[2]&(~t1[0]|~t1[1]|t1[2])|
+                     t1[2]&(~t0[0]|~t0[1]|t0[2]))
+        else: print(f'unknown op {op}')
+
+        if (op == sim.INV1 or
+            op == sim.NAND2 or op == sim.NAND3 or op == sim.NAND4 or
+            op == sim.NOR2 or op == sim.NOR3 or op == sim.NOR4 or
+            op == sim.XNOR2 or op == sim.XNOR3 or op == sim.XNOR4 or
+            op == sim.AOI21 or op == sim.OAI21 or
+            op == sim.AOI22 or op == sim.OAI22 or
+            op == sim.AOI211 or op == sim.OAI211):
+            o0[...] = o0 ^ inv_op

src/kyupy/sdf.py

@@ -59,22 +59,23 @@ class DelayFile:
             if name not in circuit.cells: name = name.replace('[', '_').replace(']', '_')
             return circuit.cells.get(name, None)
 
-        delays = np.zeros((len(circuit.lines), 2, 2, 3))  # dataset last during construction.
+        delays = np.zeros((len(circuit.lines), 2, 2, 3), dtype=np.float32)  # dataset last during construction.
 
-        for name, iopaths in self.cells.items():
+        with log.limit(50):
-            name = name.replace('\\', '')
+            for name, iopaths in self.cells.items():
-            if cell := circuit.cells.get(name, None):
+                name = name.replace('\\', '')
-                for i_pin_spec, o_pin_spec, *dels in iopaths:
+                if cell := circuit.cells.get(name, None):
-                    if i_pin_spec.startswith('(posedge '): i_pol_idxs = [0]
+                    for i_pin_spec, o_pin_spec, *dels in iopaths:
-                    elif i_pin_spec.startswith('(negedge '): i_pol_idxs = [1]
+                        if i_pin_spec.startswith('(posedge '): i_pol_idxs = [0]
-                    else: i_pol_idxs = [0, 1]
+                        elif i_pin_spec.startswith('(negedge '): i_pol_idxs = [1]
-                    i_pin_spec = re.sub(r'\((neg|pos)edge ([^)]+)\)', r'\2', i_pin_spec)
+                        else: i_pol_idxs = [0, 1]
-                    if line := cell.ins[tlib.pin_index(cell.kind, i_pin_spec)]:
+                        i_pin_spec = re.sub(r'\((neg|pos)edge ([^)]+)\)', r'\2', i_pin_spec)
-                        delays[line, i_pol_idxs] = [d if len(d) > 0 else [0, 0, 0] for d in dels]
+                        if line := cell.ins[tlib.pin_index(cell.kind, i_pin_spec)]:
-                    else:
+                            delays[line, i_pol_idxs] = [d if len(d) > 0 else [0, 0, 0] for d in dels]
-                        log.warn(f'No line to annotate in circuit: {i_pin_spec} for {cell}')
+                        else:
-            else:
+                            log.warn(f'No line to annotate in circuit: {i_pin_spec} for {cell}')
-                log.warn(f'Name from SDF not found in circuit: {name}')
+                else:
+                    log.warn(f'Name from SDF not found in circuit: {name}')
 
         return np.moveaxis(delays, -1, 0)
 
@@ -100,13 +101,19 @@ class DelayFile:
         * Axis 3: polarity of the transition, 0='rising/posedge', 1='falling/negedge'
         """
 
-        delays = np.zeros((len(circuit.lines), 2, 2, 3))  # dataset last during construction.
+        delays = np.zeros((len(circuit.lines), 2, 2, 3), dtype=np.float32)  # dataset last during construction.
 
+        def split_node_port(n_str):
+            if (split := n_str.rfind('.')) > 0 or (split := n_str.rfind('/')) > 0:
+                return n_str[:split], n_str[split+1:]
+            return (n_str, None)
+
+        nonfork_annotations = 0
         for n1, n2, *delvals in self._interconnects:
             delvals = [d if len(d) > 0 else [0, 0, 0] for d in delvals]
             if max(max(delvals)) == 0: continue
-            cn1, pn1 = n1.split('/') if '/' in n1 else (n1, None)
+            cn1, pn1 = split_node_port(n1)
-            cn2, pn2 = n2.split('/') if '/' in n2 else (n2, None)
+            cn2, pn2 = split_node_port(n2)
             cn1 = cn1.replace('\\','')
             cn2 = cn2.replace('\\','')
             c1, c2 = circuit.cells[cn1], circuit.cells[cn2]
@@ -119,19 +126,27 @@ class DelayFile:
                 log.warn(f'No line to annotate pin {pn2} of {c2}')
                 continue
             f1, f2 = c1.outs[p1].reader, c2.ins[p2].driver  # find the forks between cells.
-            assert f1.kind == '__fork__'
+            if f1 == c2 and f2 == c1:
-            assert f2.kind == '__fork__'
+                nonfork_annotations += 1
-            if f1 != f2:  # at least two forks, make sure f2 is a branchfork connected to f1
+                if nonfork_annotations < 10:
-                assert len(f2.outs) == 1
+                    log.warn(f'No fork between {c1.name}/{p1} and {c2.name}/{p2}, using {c2.name}/{p2}')
-                assert f1.outs[f2.ins[0].driver_pin] == f2.ins[0]
+                line = c2.ins[p2]
-                line = f2.ins[0]
-            elif len(f2.outs) == 1:  # f1==f2, only OK when there is no fanout.
-                line = f2.ins[0]
             else:
-                log.warn(f'No branchfork to annotate interconnect delay {c1.name}/{p1}->{c2.name}/{p2}')
+                assert f1.kind == '__fork__'
-                continue
+                assert f2.kind == '__fork__'
+                if len(f2.outs) == 1:
+                    assert f1 == f2 or f1.outs[f2.ins[0].driver_pin] == f2.ins[0]
+                    line = f2.ins[0]
+                else:
+                    nonfork_annotations += 1
+                    if nonfork_annotations < 10:
+                        log.warn(f'No branchfork between {c1.name}/{p1} and {c2.name}/{p2}, using {c2.name}/{p2}')
+                    line = c2.ins[p2]
             delays[line, :] = delvals
 
+        if nonfork_annotations > 0:
+            log.warn(f'{nonfork_annotations} interconnect annotations were moved to gate inputs due to missing forks.')
+
         return np.moveaxis(delays, -1, 0)
 
 
@@ -156,6 +171,10 @@ class SdfTransformer(Transformer):
         entries = [e for a in args if hasattr(a, 'children') for e in a.children]
         return name, entries
 
+    @staticmethod
+    def cond(args):  # ignore conditions
+        return args[1]
+
     @staticmethod
     def start(args):
         name = next((a for a in args if isinstance(a, str)), None)
@@ -180,9 +199,12 @@ GRAMMAR = r"""
         | "(INSTANCE" ID? ")"
         | "(TIMINGCHECK" _ignore* ")"
         | delay )* ")"
-    delay: "(DELAY" "(ABSOLUTE" (interconnect | iopath)* ")" ")"
+    delay: "(DELAY" "(ABSOLUTE" (interconnect | iopath | cond)* ")" ")"
     interconnect: "(INTERCONNECT" ID ID triple* ")"
     iopath: "(IOPATH" ID_OR_EDGE ID_OR_EDGE triple* ")"
+    cond: "(" "COND" cond_port_expr iopath ")"
+    ?cond_port_expr: ID | "(" cond_port_expr ")" | cond_port_expr BINARY_OP cond_port_expr
+    BINARY_OP: /&&/ | /==/
     NAME: /[^"]+/
     ID_OR_EDGE: ( /[^() ]+/ | "(" /[^)]+/ ")" )
     ID: ( /[^"() ]+/ | "\"" /[^"]+/ "\"" )
@@ -200,9 +222,9 @@ def parse(text):
     return Lark(GRAMMAR, parser="lalr", transformer=SdfTransformer()).parse(text)
 
 
-def load(file):
+def load(file) -> DelayFile:
     """Parses the contents of ``file`` and returns a :class:`DelayFile` object.
 
     Files with `.gz`-suffix are decompressed on-the-fly.
     """
-    return parse(readtext(file))
+    return parse(readtext(file)) # type: ignore

src/kyupy/sim.py

@@ -4,9 +4,16 @@ from bisect import bisect, insort_left
 
 import numpy as np
 
+from . import log
+from .circuit import Circuit
+from .techlib import KYUPY
+
 BUF1 = np.uint16(0b1010_1010_1010_1010)
 INV1 = ~BUF1
 
+__const0__ = BUF1
+__const1__ = INV1
+
 AND2 = np.uint16(0b1000_1000_1000_1000)
 AND3 = np.uint16(0b1000_0000_1000_0000)
 AND4 = np.uint16(0b1000_0000_0000_0000)
@@ -39,7 +46,10 @@ AOI211, OAI211 = ~AO211, ~OA211
 
 MUX21 = np.uint16(0b1100_1010_1100_1010)  # z = i1 if i2 else i0 (i2 is select)
 
-names = dict([(v, k) for k, v in globals().items() if isinstance(v, np.uint16)])
+names = dict([(v, k) for k, v in globals().items() if isinstance(v, np.uint16) and '__' not in k])
+
+prim2name = dict([(v, k) for k, v in globals().items() if isinstance(v, np.uint16) and '__' not in k])
+name2prim = dict([(k, v) for k, v in globals().items() if isinstance(v, np.uint16)])
 
 kind_prefixes = {
     'nand': (NAND4, NAND3, NAND2),
@@ -150,15 +160,15 @@ class SimOps:
     """A static scheduler that translates a Circuit into a topologically sorted list of basic logic operations (self.ops) and
     a memory mapping (self.c_locs, self.c_caps) for use in simulators.
 
-    :param circuit: The circuit to create a schedule for.
+    :param circuit: The circuit to create a schedule for. Must contain cells from KYUPY techlib only.
     :param strip_forks: If enabled, the scheduler will not include fork nodes to safe simulation time.
         Stripping forks will cause interconnect delay annotations of lines read by fork nodes to be ignored.
     :param c_reuse: If enabled, memory of intermediate signal waveforms will be re-used. This greatly reduces
         memory footprint, but intermediate signal waveforms become unaccessible after a propagation.
     """
-    def __init__(self, circuit, c_caps=1, c_caps_min=1, a_ctrl=None, c_reuse=False, strip_forks=False):
+    def __init__(self, circuit: Circuit, c_caps=1, c_caps_min=1, a_ctrl=None, c_reuse=False, strip_forks=False):
         self.circuit = circuit
-        self.s_len = len(circuit.s_nodes)
+        self.s_len = len(circuit.s_nodes(KYUPY))
 
         if isinstance(c_caps, int):
             c_caps = [c_caps] * (len(circuit.lines)+3)
@@ -175,84 +185,79 @@ class SimOps:
         self.ppo_offset = self.ppi_offset + self.s_len
         self.c_locs_len = self.ppo_offset + self.s_len
 
-        # translate circuit structure into self.ops
+        # ALAP-toposort the circuit into self.ops
-        ops = []
+        levels = []
-        interface_dict = dict((n, i) for i, n in enumerate(circuit.s_nodes))
-        for n in circuit.topological_order():
-            if n in interface_dict:
-                inp_idx = self.ppi_offset + interface_dict[n]
-                if len(n.outs) > 0 and n.outs[0] is not None:  # first output of a PI/PPI
-                    ops.append((BUF1, n.outs[0].index, inp_idx, self.zero_idx, self.zero_idx, self.zero_idx, *a_ctrl[n.outs[0]]))
-                if 'dff' in n.kind.lower():  # second output of DFF is inverted
-                    if len(n.outs) > 1 and n.outs[1] is not None:
-                        ops.append((INV1, n.outs[1].index, inp_idx, self.zero_idx, self.zero_idx, self.zero_idx, *a_ctrl[n.outs[1]]))
-                else:  # if not DFF, no output is inverted.
-                    for o_line in n.outs[1:]:
-                        if o_line is not None:
-                            ops.append((BUF1, o_line.index, inp_idx, self.zero_idx, self.zero_idx, self.zero_idx, *a_ctrl[o_line]))
-                continue
-            # regular node, not PI/PPI or PO/PPO
-            o0_idx = n.outs[0].index if len(n.outs) > 0 and n.outs[0] is not None else self.tmp_idx
-            i0_idx = n.ins[0].index if len(n.ins) > 0 and n.ins[0] is not None else self.zero_idx
-            i1_idx = n.ins[1].index if len(n.ins) > 1 and n.ins[1] is not None else self.zero_idx
-            i2_idx = n.ins[2].index if len(n.ins) > 2 and n.ins[2] is not None else self.zero_idx
-            i3_idx = n.ins[3].index if len(n.ins) > 3 and n.ins[3] is not None else self.zero_idx
-            kind = n.kind.lower()
-            if kind == '__fork__':
-                if not strip_forks:
-                    for o_line in n.outs:
-                        if o_line is not None:
-                            ops.append((BUF1, o_line.index, i0_idx, i1_idx, i2_idx, i3_idx, *a_ctrl[o_line]))
-                continue
-            sp = None
-            for prefix, prims in kind_prefixes.items():
-                if kind.startswith(prefix):
-                    sp = prims[0]
-                    if i3_idx == self.zero_idx:
-                        sp = prims[1]
-                        if i2_idx == self.zero_idx:
-                            sp = prims[2]
-                    break
-            if sp is None:
-                print('unknown cell type', kind)
-            else:
-                ops.append((sp, o0_idx, i0_idx, i1_idx, i2_idx, i3_idx, *a_ctrl[o0_idx]))
 
-        self.ops = np.asarray(ops, dtype='int32')
+        ppio2idx = dict((n, i) for i, n in enumerate(circuit.s_nodes(KYUPY)))
+        root_nodes = set([n for n in circuit.nodes if len(n.outs) == 0 or KYUPY.is_dff(n.kind)])  # start from POs (that are unused by circuit itself), PPOs, and any dangling nodes
+        readers = np.array([1 if l.reader in root_nodes else len(l.reader.outs) for l in circuit.lines], dtype=np.int32)  # for ref-counting forks
+
+        level_lines = [n.ins[0] for n in sorted(root_nodes, key=lambda n: n.index) if len(n.ins) > 0 ]
+        # FIXME: Should probably instanciate buffers for PPOs and attach DFF clocks
+
+        while len(level_lines) > 0:  # traverse the circuit level-wise back towards (P)PIs
+            level_ops = []
+            prev_level_lines = []
+
+            for l in level_lines:
+                n = l.driver
+                assert len(n.ins) <= 4, f'Too many input pins for node {n}. Map circuit to KYUPY techlib first using resolve_tlib_cells().'
+                in_idxs = [n.ins[x].index if len(n.ins) > x and n.ins[x] is not None else self.zero_idx for x in [0,1,2,3]]
+                if n in ppio2idx:
+                    in_idxs[1] = self.zero_idx
+                    if 'dff' in n.kind.lower() or in_idxs[0] == self.zero_idx:
+                        in_idxs[0] = self.ppi_offset + ppio2idx[n]
+                    else:  # this is a PO that is also used within the circuit. Treat as normal fork with ref counting. TODO: check strip_forks=True case for this.
+                        readers[n.ins[0]] -= 1
+                        if readers[n.ins[0]] == 0: prev_level_lines.append(n.ins[0])
+                    if l.driver_pin == 1 and 'dff' in n.kind.lower():  # second output of DFF is inverted
+                        level_ops.append((INV1, l.index, *in_idxs, *a_ctrl[l]))
+                    else:
+                        level_ops.append((BUF1, l.index, *in_idxs, *a_ctrl[l]))
+                elif n.kind == '__fork__':
+                    readers[n.ins[0]] -= 1
+                    if readers[n.ins[0]] == 0: prev_level_lines.append(n.ins[0])
+                    if not strip_forks: level_ops.append((BUF1, l.index, *in_idxs, *a_ctrl[l]))
+                else:
+                    prev_level_lines += n.ins
+                    sp = None
+                    kind = n.kind.lower()
+                    for prefix, prims in kind_prefixes.items():
+                        if kind.startswith(prefix):
+                            sp = prims[0]
+                            if in_idxs[3] == self.zero_idx:
+                                sp = prims[1]
+                                if in_idxs[2] == self.zero_idx:
+                                    sp = prims[2]
+                            break
+                    assert sp is not None, f'Unsupported node type {n}. Map circuit to KYUPY techlib first using resolce_tlib_cells().'
+                    level_ops.append((sp, l.index, *in_idxs, *a_ctrl[l]))
+
+            if len(level_ops) > 0: levels.append(level_ops)
+            level_lines = prev_level_lines
+
+        self.levels = [np.asarray(lv, dtype=np.int32) for lv in levels[::-1]]
+        level_sums = np.cumsum([0]+[len(lv) for lv in self.levels], dtype=np.int32)
+        self.level_starts, self.level_stops = level_sums[:-1], level_sums[1:]
+        # op format: [kind, out0, in0, in1, in2, in3, wsa_acc_pos, wsa_rise, wsa_fall]
+        self.ops = np.vstack(self.levels)
 
         # create a map from fanout lines to stem lines for fork stripping
-        stems = np.zeros(self.c_locs_len, dtype='int32') - 1  # default to -1: 'no fanout line'
+        stems = np.full(self.c_locs_len, -1, dtype=np.int32)  # default to -1: 'no fanout line'
         if strip_forks:
             for f in circuit.forks.values():
                 prev_line = f.ins[0]
                 while prev_line.driver.kind == '__fork__':
                     prev_line = prev_line.driver.ins[0]
-                stem_idx = prev_line.index
                 for ol in f.outs:
                     if ol is not None:
-                        stems[ol] = stem_idx
+                        stems[ol] = prev_line.index
 
-        # calculate level (distance from PI/PPI) and reference count for each line
+        ref_count = np.zeros(self.c_locs_len, dtype=np.int32)
-        levels = np.zeros(self.c_locs_len, dtype='int32')
+
-        ref_count = np.zeros(self.c_locs_len, dtype='int32')
+        for op in self.ops:
-        level_starts = [0]
+            for x in [2, 3, 4, 5]:
-        current_level = 1
+                ref_count[stems[op[x]] if stems[op[x]] >= 0 else op[x]] += 1
-        for i, op in enumerate(self.ops):
-            # if we fork-strip, always take the stems for determining fan-in level
-            i0_idx = stems[op[2]] if stems[op[2]] >= 0 else op[2]
-            i1_idx = stems[op[3]] if stems[op[3]] >= 0 else op[3]
-            i2_idx = stems[op[4]] if stems[op[4]] >= 0 else op[4]
-            i3_idx = stems[op[5]] if stems[op[5]] >= 0 else op[5]
-            if levels[i0_idx] >= current_level or levels[i1_idx] >= current_level or levels[i2_idx] >= current_level or levels[i3_idx] >= current_level:
-                current_level += 1
-                level_starts.append(i)
-            levels[op[1]] = current_level  # set level of the output line
-            ref_count[i0_idx] += 1
-            ref_count[i1_idx] += 1
-            ref_count[i2_idx] += 1
-            ref_count[i3_idx] += 1
-        self.level_starts = np.asarray(level_starts, dtype='int32')
-        self.level_stops = np.asarray(level_starts[1:] + [len(self.ops)], dtype='int32')
 
         # combinational signal allocation table. maps line and interface indices to self.c memory locations
         self.c_locs = np.full((self.c_locs_len,), -1, dtype=np.int32)
@@ -269,18 +274,18 @@ class SimOps:
         ref_count[self.tmp2_idx] += 1
 
         # allocate and keep memory for PI/PPI, keep memory for PO/PPO (allocated later)
-        for i, n in enumerate(circuit.s_nodes):
+        for i, n in enumerate(circuit.s_nodes(KYUPY)):
-            if len(n.outs) > 0:
+            if 'dff' in n.kind.lower() or len(n.ins) == 0:  # PPI or PI
                 self.c_locs[self.ppi_offset + i], self.c_caps[self.ppi_offset + i] = h.alloc(c_caps_min), c_caps_min
                 ref_count[self.ppi_offset + i] += 1
-            if len(n.ins) > 0:
+            if len(n.ins) > 0:  # PO
                 i0_idx = stems[n.ins[0]] if stems[n.ins[0]] >= 0 else n.ins[0]
                 ref_count[i0_idx] += 1
 
         # allocate memory for the rest of the circuit
-        for op_start, op_stop in zip(self.level_starts, self.level_stops):
+        for ops in self.levels:
             free_set = set()
-            for op in self.ops[op_start:op_stop]:
+            for op in ops:
                 # if we fork-strip, always take the stems
                 i0_idx = stems[op[2]] if stems[op[2]] >= 0 else op[2]
                 i1_idx = stems[op[3]] if stems[op[3]] >= 0 else op[3]
@@ -299,7 +304,8 @@ class SimOps:
                 self.c_locs[o_idx], self.c_caps[o_idx] = h.alloc(cap), cap
             if c_reuse:
                 for loc in free_set:
-                    h.free(loc)
+                    if loc >= 0:  # DFF clocks are not allocated. Ignore for now.
+                        h.free(loc)
 
         # copy memory location and capacity from stems to fanout lines
         for lidx, stem in enumerate(stems):
@@ -307,10 +313,24 @@ class SimOps:
                 self.c_locs[lidx], self.c_caps[lidx] = self.c_locs[stem], self.c_caps[stem]
 
         # copy memory location to PO/PPO area
-        for i, n in enumerate(circuit.s_nodes):
+        for i, n in enumerate(circuit.s_nodes(KYUPY)):
             if len(n.ins) > 0:
                 self.c_locs[self.ppo_offset + i], self.c_caps[self.ppo_offset + i] = self.c_locs[n.ins[0]], self.c_caps[n.ins[0]]
 
+        # sanity check
+        for oidx, io in zip(*np.nonzero(self.c_locs[self.ops[:, 1:6]] < 0)):
+            log.error(f'unallocated {["o0","i0","i1","i2","i3"][io]} in op {self.ops[oidx]}')
+        assert np.min(self.c_locs[self.ops[:, 1:6]]) >= 0, "Unallocated signal lines."
+
+        # line use information
+        self.line_use_start = np.full(self.c_locs_len, -1, dtype=np.int32)
+        self.line_use_stop = np.full(self.c_locs_len, len(self.levels), dtype=np.int32)
+        for i, lv in enumerate(self.levels):
+            for op in lv:
+                self.line_use_start[op[1]] = i
+                for x in [2, 3, 4, 5]:
+                    self.line_use_stop[op[x]] = i
+
         self.c_len = h.max_size
 
         d = defaultdict(int)
@@ -318,11 +338,16 @@ class SimOps:
         self.prim_counts = dict(d)
 
         self.pi_s_locs = np.flatnonzero(self.c_locs[self.ppi_offset+np.arange(len(self.circuit.io_nodes))] >= 0)
+        """Locations in pattern array (axis 0) of all primary inputs (excluding flip-flops)."""
         self.po_s_locs = np.flatnonzero(self.c_locs[self.ppo_offset+np.arange(len(self.circuit.io_nodes))] >= 0)
+        """Locations in pattern array (axis 0) of all primary outputs (excluding flip-flops)."""
         self.ppio_s_locs = np.arange(len(self.circuit.io_nodes), self.s_len)
+        """Locations in pattern array (axis 0) of all flip-flops. Pseudo-primary inputs (PPI) and pseudo-primary outputs (PPO) share the same location."""
 
         self.pippi_s_locs = np.concatenate([self.pi_s_locs, self.ppio_s_locs])
+        """Locations in pattern array (axis 0) of all primary and pseudo-primary inputs (input ports and flip-flops)."""
         self.poppo_s_locs = np.concatenate([self.po_s_locs, self.ppio_s_locs])
+        """Locations in pattern array (axis 0) of all primary and pseudo-primary outputs (output ports and flip-flops)."""
 
         self.pi_c_locs = self.c_locs[self.ppi_offset+self.pi_s_locs]
         self.po_c_locs = self.c_locs[self.ppo_offset+self.po_s_locs]
@@ -331,3 +356,12 @@ class SimOps:
 
         self.pippi_c_locs = np.concatenate([self.pi_c_locs, self.ppi_c_locs])
         self.poppo_c_locs = np.concatenate([self.po_c_locs, self.ppo_c_locs])
+
+    def cone_ops(self, origin_lines: set[int]):
+        active_lines = {l for l in origin_lines}
+        ops = []
+        for kind, out0, in0, in1, in2, in3, *params in self.ops:
+            if in0 in active_lines or in1 in active_lines or in2 in active_lines or in3 in active_lines or out0 in origin_lines:
+                ops.append((kind, out0, in0, in1, in2, in3, *params))
+                active_lines.add(out0)
+        return np.asarray(ops, dtype=np.int32)

src/kyupy/stil.py

@@ -41,7 +41,7 @@ class StilFile:
                 unload = {}
                 for so_port in self.so_ports:
                     if so_port in call.parameters:
-                        unload[so_port] = call.parameters[so_port].replace('\n', '').replace('N', '-')
+                        unload[so_port] = call.parameters[so_port]
                 if len(capture) > 0:
                     self.patterns.append(ScanPattern(sload, launch, capture, unload))
                     capture = {}
@@ -49,11 +49,9 @@ class StilFile:
                 sload = {}
                 for si_port in self.si_ports:
                     if si_port in call.parameters:
-                        sload[si_port] = call.parameters[si_port].replace('\n', '').replace('N', '-')
+                        sload[si_port] = call.parameters[si_port]
-            if call.name.endswith('_launch'):
+            if call.name.endswith('_launch'): launch = call.parameters
-                launch = dict((k, v.replace('\n', '').replace('N', '-')) for k, v in call.parameters.items())
+            if call.name.endswith('_capture'): capture = call.parameters
-            if call.name.endswith('_capture'):
-                capture = dict((k, v.replace('\n', '').replace('N', '-')) for k, v in call.parameters.items())
 
     def _maps(self, c):
         interface = list(c.io_nodes) + [n for n in c.nodes if 'DFF' in n.kind]
@@ -100,15 +98,15 @@ class StilFile:
         tests = np.full((len(interface), len(self.patterns)), logic.UNASSIGNED)
         for i, p in enumerate(self.patterns):
             for si_port in self.si_ports.keys():
-                pattern = logic.mvarray(p.load[si_port])
+                pattern = logic.mvarray(p.load[si_port][0])
                 inversions = np.choose((pattern == logic.UNASSIGNED) | (pattern == logic.UNKNOWN),
                                        [scan_inversions[si_port], logic.ZERO]).astype(np.uint8)
                 np.bitwise_xor(pattern, inversions, out=pattern)
                 tests[scan_maps[si_port], i] = pattern
-            tests[pi_map, i] = logic.mvarray(p.capture['_pi'])
+            tests[pi_map, i] = logic.mvarray(p.capture['_pi'][0])
         return tests
 
-    def tests_loc(self, circuit, init_filter=None, launch_filter=None):
+    def tests_loc(self, circuit, tlib, init_filter=None, launch_filter=None):
         """Assembles and returns a LoC scan test pattern set for given circuit.
 
         This function assumes a launch-on-capture (LoC) delay test.
@@ -134,14 +132,16 @@ class StilFile:
         for i, p in enumerate(self.patterns):
             # init.set_values(i, '0' * len(interface))
             for si_port in self.si_ports.keys():
-                pattern = logic.mvarray(p.load[si_port])
+                pattern = logic.mvarray(p.load[si_port][0])
                 inversions = np.choose((pattern == logic.UNASSIGNED) | (pattern == logic.UNKNOWN),
                                        [scan_inversions[si_port], logic.ZERO]).astype(np.uint8)
                 np.bitwise_xor(pattern, inversions, out=pattern)
                 init[scan_maps[si_port], i] = pattern
-            init[pi_map, i] = logic.mvarray(p.launch['_pi'] if '_pi' in p.launch else p.capture['_pi'])
+            init[pi_map, i] = logic.mvarray(p.launch['_pi'][0] if '_pi' in p.launch else p.capture['_pi'][0])
         if init_filter: init = init_filter(init)
-        sim8v = LogicSim(circuit, init.shape[-1], m=8)
+        circuit_resolved = circuit.copy()
+        circuit_resolved.resolve_tlib_cells(tlib)
+        sim8v = LogicSim(circuit_resolved, init.shape[-1], m=8)
         sim8v.s[0] = logic.mv_to_bp(init)
         sim8v.s_to_c()
         sim8v.c_prop()
@@ -149,12 +149,12 @@ class StilFile:
         launch = logic.bp_to_mv(sim8v.s[1])[..., :init.shape[-1]]
         for i, p in enumerate(self.patterns):
             # if there was no launch cycle or launch clock, then init = launch
-            if '_pi' not in p.launch or 'P' not in p.launch['_pi'] or 'P' not in p.capture['_pi']:
+            if '_pi' not in p.launch or 'P' not in p.launch['_pi'][0] or 'P' not in p.capture['_pi'][0]:
                 for si_port in self.si_ports.keys():
-                    pattern = logic.mv_xor(logic.mvarray(p.load[si_port]), scan_inversions[si_port])
+                    pattern = logic.mv_xor(logic.mvarray(p.load[si_port][0]), scan_inversions[si_port])
                     launch[scan_maps[si_port], i] = pattern
-            if '_pi' in p.capture and 'P' in p.capture['_pi']:
+            if '_pi' in p.capture and 'P' in p.capture['_pi'][0]:
-                launch[pi_map, i] = logic.mvarray(p.capture['_pi'])
+                launch[pi_map, i] = logic.mvarray(p.capture['_pi'][0])
             launch[po_map, i] = logic.UNASSIGNED
         if launch_filter: launch = launch_filter(launch)
 
@@ -171,9 +171,9 @@ class StilFile:
         interface, _, po_map, scan_maps, scan_inversions = self._maps(circuit)
         resp = np.full((len(interface), len(self.patterns)), logic.UNASSIGNED)
         for i, p in enumerate(self.patterns):
-            resp[po_map, i] = logic.mvarray(p.capture['_po'] if len(p.capture) > 0 else p.launch['_po'])
+            resp[po_map, i] = logic.mvarray(p.capture['_po'][0] if len(p.capture) > 0 else p.launch['_po'][0])
             for so_port in self.so_ports.keys():
-                pattern = logic.mv_xor(logic.mvarray(p.unload[so_port]), scan_inversions[so_port])
+                pattern = logic.mv_xor(logic.mvarray(p.unload[so_port][0]), scan_inversions[so_port])
                 resp[scan_maps[so_port], i] = pattern
         return resp
 
@@ -192,7 +192,7 @@ class StilTransformer(Transformer):
     def call(args): return Call(args[0], dict(args[1:]))
 
     @staticmethod
-    def call_parameter(args): return args[0], args[1].value
+    def call_parameter(args): return args[0], (args[1].value.replace('\n', '').replace('N', '-'), args[1].start_pos)
 
     @staticmethod
     def signal_group(args): return args[0], args[1:]
@@ -265,12 +265,12 @@ GRAMMAR = r"""
     """
 
 
-def parse(text):
+def parse(text) -> StilFile:
     """Parses the given ``text`` and returns a :class:`StilFile` object."""
-    return Lark(GRAMMAR, parser="lalr", transformer=StilTransformer()).parse(text)
+    return Lark(GRAMMAR, parser="lalr", transformer=StilTransformer()).parse(text) # type: ignore
 
 
-def load(file):
+def load(file) -> StilFile:
     """Parses the contents of ``file`` and returns a :class:`StilFile` object.
 
     Files with `.gz`-suffix are decompressed on-the-fly.

src/kyupy/techlib.py

@@ -10,51 +10,6 @@ from itertools import product
 
 from . import bench
 
-
-class TechLibOld:
-    @staticmethod
-    def pin_index(kind, pin):
-        if isinstance(pin, int):
-            return max(0, pin-1)
-        if kind[:3] in ('OAI', 'AOI'):
-            if pin[0] == 'A': return int(pin[1]) - 1
-            if pin == 'B': return int(kind[3])
-            if pin[0] == 'B': return int(pin[1]) - 1 + int(kind[3])
-        for prefix, pins, index in [('HADD', ('B0', 'SO'), 1),
-                                    ('HADD', ('A0', 'C1'), 0),
-                                    ('MUX21', ('S', 'S0'), 2),
-                                    ('MX2', ('S0',), 2),
-                                    ('TBUF', ('OE',), 1),
-                                    ('TINV', ('OE',), 1),
-                                    ('LATCH', ('D',), 0),
-                                    ('LATCH', ('QN',), 1),
-                                    ('DFF', ('D',), 0),
-                                    ('DFF', ('QN',), 1),
-                                    ('SDFF', ('D',), 0),
-                                    ('SDFF', ('QN',), 1),
-                                    ('SDFF', ('CLK',), 3),
-                                    ('SDFF', ('RSTB', 'RN'), 4),
-                                    ('SDFF', ('SETB',), 5),
-                                    ('ISOL', ('ISO',), 0),
-                                    ('ISOL', ('D',), 1)]:
-            if kind.startswith(prefix) and pin in pins: return index
-        for index, pins in enumerate([('A1', 'IN1', 'A', 'S', 'INP', 'I', 'Q', 'QN', 'Y', 'Z', 'ZN'),
-                                      ('A2', 'IN2', 'B', 'CK', 'CLK', 'CO', 'SE'),
-                                      ('A3', 'IN3', 'C', 'RN', 'RSTB', 'CI', 'SI'),
-                                      ('A4', 'IN4', 'D', 'SN', 'SETB'),
-                                      ('A5', 'IN5', 'E'),
-                                      ('A6', 'IN6', 'F')]):
-            if pin in pins: return index
-        raise ValueError(f'Unknown pin index for {kind}.{pin}')
-
-    @staticmethod
-    def pin_is_output(kind, pin):
-        if isinstance(pin, int):
-            return pin == 0
-        if 'MUX' in kind and pin == 'S': return False
-        return pin in ('Q', 'QN', 'Z', 'ZN', 'Y', 'CO', 'S', 'SO', 'C1')
-
-
 class TechLib:
     """Class for standard cell library definitions.
 
@@ -83,9 +38,11 @@ class TechLib:
                 else:
                     pin_dict[n.name] = (o_idx, True)
                     o_idx += 1
+            has_dff = 'DFF' in set(n.kind for n in c.cells.values())
+            has_latch = 'LATCH' in set(n.kind for n in c.cells.values())
             parts = [s[1:-1].split(',') if s[0] == '{' else [s] for s in re.split(r'({[^}]+})', c.name) if len(s) > 0]
             for name in [''.join(item) for item in product(*parts)]:
-                self.cells[name] = (c, pin_dict)
+                self.cells[name] = (c, pin_dict, has_dff, has_latch)
 
     def pin_index(self, kind, pin):
         """Returns a pin list position for a given node kind and pin name."""
@@ -93,12 +50,80 @@ class TechLib:
         assert pin in self.cells[kind][1], f'Unknown pin: {pin} for cell {kind}'
         return self.cells[kind][1][pin][0]
 
+    def pin_name(self, kind: str, pos: int, output=False) -> str:
+        """Returns the pin name for a given node kind, list position, and direction."""
+        assert kind in self.cells, f'Unknown cell: {kind}'
+        for name, (ppos, isout) in self.cells[kind][1].items():
+            if isout == output and ppos == pos:
+                return name
+        raise ValueError(f'Invalid {"output" if output else "input"} pin position {pos} for cell: {kind}')
+
     def pin_is_output(self, kind, pin):
         """Returns True, if given pin name of a node kind is an output."""
         assert kind in self.cells, f'Unknown cell: {kind}'
         assert pin in self.cells[kind][1], f'Unknown pin: {pin} for cell {kind}'
         return self.cells[kind][1][pin][1]
     
+    def is_dff(self, kind):
+        """Returns True, if given node kind is a d-flip-flop."""
+        if kind == '__fork__': return False
+        if kind == '__const0__': return False
+        if kind == '__const1__': return False
+        if kind == 'input': return False
+        if kind == 'output': return False
+        assert kind in self.cells, f'Unknown cell: {kind}'
+        return self.cells[kind][2]
+    
+    def is_latch(self, kind):
+        """Returns True, if given node kind is a latch."""
+        if kind == '__fork__': return False
+        if kind == '__const0__': return False
+        if kind == '__const1__': return False
+        if kind == 'input': return False
+        if kind == 'output': return False
+        assert kind in self.cells, f'Unknown cell: {kind}'
+        return self.cells[kind][3]
+
+
+KYUPY = TechLib(r"""
+BUF1    input(i0)          output(o) o=BUF1(i0)          ;
+INV1    input(i0)          output(o) o=INV1(i0)          ;
+AND2    input(i0,i1)       output(o) o=AND2(i0,i1)       ;
+AND3    input(i0,i1,i2)    output(o) o=AND3(i0,i1,i2)    ;
+AND4    input(i0,i1,i2,i3) output(o) o=AND4(i0,i1,i2,i3) ;
+NAND2   input(i0,i1)       output(o) o=NAND2(i0,i1)      ;
+NAND3   input(i0,i1,i2)    output(o) o=NAND3(i0,i1,i2)   ;
+NAND4   input(i0,i1,i2,i3) output(o) o=NAND4(i0,i1,i2,i3);
+OR2     input(i0,i1)       output(o) o=OR2(i0,i1)        ;
+OR3     input(i0,i1,i2)    output(o) o=OR3(i0,i1,i2)     ;
+OR4     input(i0,i1,i2,i3) output(o) o=OR4(i0,i1,i2,i3)  ;
+NOR2    input(i0,i1)       output(o) o=NOR2(i0,i1)       ;
+NOR3    input(i0,i1,i2)    output(o) o=NOR3(i0,i1,i2)    ;
+NOR4    input(i0,i1,i2,i3) output(o) o=NOR4(i0,i1,i2,i3) ;
+XOR2    input(i0,i1)       output(o) o=XOR2(i0,i1)       ;
+XOR3    input(i0,i1,i2)    output(o) o=XOR3(i0,i1,i2)    ;
+XOR4    input(i0,i1,i2,i3) output(o) o=XOR4(i0,i1,i2,i3) ;
+XNOR2   input(i0,i1)       output(o) o=XNOR2(i0,i1)      ;
+XNOR3   input(i0,i1,i2)    output(o) o=XNOR3(i0,i1,i2)   ;
+XNOR4   input(i0,i1,i2,i3) output(o) o=XNOR4(i0,i1,i2,i3)  ;
+AO21    input(i0,i1,i2)    output(o) o=AO21(i0,i1,i2)      ;
+AO22    input(i0,i1,i2,i3) output(o) o=AO22(i0,i1,i2,i3)   ;
+OA21    input(i0,i1,i2)    output(o) o=OA21(i0,i1,i2)      ;
+OA22    input(i0,i1,i2,i3) output(o) o=OA22(i0,i1,i2,i3)   ;
+AOI21   input(i0,i1,i2)    output(o) o=AOI21(i0,i1,i2)     ;
+AOI22   input(i0,i1,i2,i3) output(o) o=AOI22(i0,i1,i2,i3)  ;
+OAI21   input(i0,i1,i2)    output(o) o=OAI21(i0,i1,i2)     ;
+OAI22   input(i0,i1,i2,i3) output(o) o=OAI22(i0,i1,i2,i3)  ;
+AO211   input(i0,i1,i2,i3) output(o) o=AO211(i0,i1,i2,i3)  ;
+OA211   input(i0,i1,i2,i3) output(o) o=OA211(i0,i1,i2,i3)  ;
+AOI211  input(i0,i1,i2,i3) output(o) o=AOI211(i0,i1,i2,i3) ;
+OAI211  input(i0,i1,i2,i3) output(o) o=OAI211(i0,i1,i2,i3) ;
+MUX21   input(i0,i1,i2)    output(o) o=MUX21(i0,i1,i2)     ;
+DFF     input(D,CLK)       output(Q) Q=DFF(D,CLK)        ;
+LATCH   input(D,CLK)       output(Q) Q=LATCH(D,CLK)      ;
+""")
+"""A synthetic library of all KyuPy simulation primitives.
+"""
 
 GSC180 = TechLib(r"""
 BUFX{1,3}      input(A)    output(Y) Y=BUF1(A)    ;
@@ -125,7 +150,7 @@ OAI21X1 input(A0,A1,B0)          output(Y)    Y=OAI21(A0,A1,B0)    ;
 OAI22X1 input(A0,A1,B0,B1)       output(Y)    Y=OAI22(A0,A1,B0,B1) ;
 OAI33X1 input(A0,A1,A2,B0,B1,B2) output(Y)    AA=OR2(A0,A1) BB=OR2(B0,B1) Y=OAI22(AA,A2,BB,B2) ;
 ADDFX1  input(A,B,CI)            output(CO,S) AB=XOR2(A,B) CO=XOR2(AB,CI) S=AO22(AB,CI,A,B)    ;
-ADDHX1  input(A,B)               output(CO,S) CO=XOR2(A,B) S=AND2(A,B)                         ;
+ADDHX1  input(A,B)               output(CO,S) CO=AND2(A,B) S=XOR2(A,B)                         ;
 
 DFFX1    input(CK,D)             output(Q,QN) Q=DFF(D,CK) QN=INV1(Q) ;
 DFFSRX1  input(CK,D,RN,SN)       output(Q,QN) DR=AND2(D,RN) SET=INV1(SN) DRS=OR2(DR,SET) Q=DFF(DRS,CK) QN=INV1(Q) ;
@@ -138,21 +163,92 @@ TLATX1   input(C,D)       output(Q,QN) Q=LATCH(D,C) QN=INV1(Q) ;
 """
 
 
-_nangate_common = r"""
+NANGATE = TechLib(r"""
+FILLTIE ;
+FILL_X{1,2,4,8,16} ;
+ANTENNA input(I)   ;
+
+TIEH output(Z) Z=__const1__() ;
+TIEL output(ZN) ZN=__const0__() ;
+
+BUF_X{1,2,4,8,12,16}  input(I) output(Z)  Z=BUF1(I)  ;
+INV_X{1,2,4,8,12,16}  input(I) output(ZN) ZN=INV1(I) ;
+
+CLKBUF_X{1,2,4,8,12,16} input(I) output(Z)  Z=BUF1(I)  ;
+CLKGATETST_X1 input(CLK,E,TE) output(Q) Q=OA21(CLK,E,TE) ;
+
+AND2_X{1,2}  input(A1,A2)       output(Z)  Z=AND2(A1,A2)         ;
+AND3_X{1,2}  input(A1,A2,A3)    output(Z)  Z=AND3(A1,A2,A3)      ;
+AND4_X{1,2}  input(A1,A2,A3,A4) output(Z)  Z=AND4(A1,A2,A3,A4)   ;
+NAND2_X{1,2} input(A1,A2)       output(ZN) ZN=NAND2(A1,A2)       ;
+NAND3_X{1,2} input(A1,A2,A3)    output(ZN) ZN=NAND3(A1,A2,A3)    ;
+NAND4_X{1,2} input(A1,A2,A3,A4) output(ZN) ZN=NAND4(A1,A2,A3,A4) ;
+OR2_X{1,2}   input(A1,A2)       output(Z)  Z=OR2(A1,A2)          ;
+OR3_X{1,2}   input(A1,A2,A3)    output(Z)  Z=OR3(A1,A2,A3)       ;
+OR4_X{1,2}   input(A1,A2,A3,A4) output(Z)  Z=OR4(A1,A2,A3,A4)    ;
+NOR2_X{1,2}  input(A1,A2)       output(ZN) ZN=NOR2(A1,A2)        ;
+NOR3_X{1,2}  input(A1,A2,A3)    output(ZN) ZN=NOR3(A1,A2,A3)     ;
+NOR4_X{1,2}  input(A1,A2,A3,A4) output(ZN) ZN=NOR4(A1,A2,A3,A4)  ;
+XOR2_X1      input(A1,A2)       output(Z)  Z=XOR2(A1,A2)         ;
+XNOR2_X1     input(A1,A2)       output(ZN) ZN=XNOR2(A1,A2)       ;
+
+MUX2_X1 input(I0,I1,S) output(Z) Z=MUX21(I0,I1,S) ;
+
+HA_X1 input(A,B) output(CO,S) CO=AND2(A,B) S=XOR2(A,B) ;
+FA_X1 input(A,B,CI) output(CO,S) AB=XOR2(A,B) CO=XOR2(AB,CI) S=AO22(CI,A,B) ;
+
+AOI21_X{1,2} input(A1,A2,B)     output(ZN) ZN=AOI21(A1,A2,B)     ;
+OAI21_X{1,2} input(A1,A2,B)     output(ZN) ZN=OAI21(A1,A2,B)     ;
+AOI22_X{1,2} input(A1,A2,B1,B2) output(ZN) ZN=AOI22(A1,A2,B1,B2) ;
+OAI22_X{1,2} input(A1,A2,B1,B2) output(ZN) ZN=OAI22(A1,A2,B1,B2) ;
+
+DFFRNQ_X1  input(D,RN,CLK)    output(Q)  DR=AND2(D,RN) Q=DFF(DR,CLK) ;
+DFFSNQ_X1  input(D,SN,CLK)    output(Q)  S=INV1(SN) DS=OR2(D,S) Q=DFF(DS,CLK) ;
+
+SDFFRNQ_X1 input(D,RN,SE,SI,CLK) output(Q) DR=AND2(D,RN) DI=MUX21(DR,SI,SE) Q=DFF(DI,CLK) ;
+SDFFSNQ_X1 input(D,SE,SI,SN,CLK) output(Q) S=INV1(SN) DS=OR2(D,S) DI=MUX21(DS,SI,SE) Q=DFF(DI,CLK) ;
+
+TBUF_X{1,2,4,8,12,16} input(EN,I)   output(Z)  Z=BUF1(I) ;
+
+LHQ_X1 input(D,E) output(Q)  Q=LATCH(D,E) ;
+""")
+"""Nangate 15nm Open Cell Library (NanGate_15nm_OCL_v0.1_2014_06.A).
+"""
+
+
+NANGATE45 = TechLib(r"""
 FILLCELL_X{1,2,4,8,16,32} ;
+ANTENNA_X1 input(A) ;
 
 LOGIC0_X1 output(Z) Z=__const0__() ;
 LOGIC1_X1 output(Z) Z=__const1__() ;
 
 BUF_X{1,2,4,8,16,32}  input(A) output(Z)  Z=BUF1(A)  ;
-CLKBUF_X{1,2,3}       input(A) output(Z)  Z=BUF1(A)  ;
+INV_X{1,2,4,8,16,32}  input(A) output(ZN) ZN=INV1(A) ;
 
+CLKBUF_X{1,2,3}       input(A)       output(Z)   Z=BUF1(A)         ;
+CLKGATETST_X{1,2,4,8} input(CK,E,SE) output(GCK) GCK=OA21(CK,E,SE) ;
+CLKGATE_X{1,2,4,8}    input(CK,E)    output(GCK) GCK=AND2(CK,E)    ;
+
+AND2_X{1,2,4}  input(A1,A2)       output(ZN) ZN=AND2(A1,A2)        ;
+AND3_X{1,2,4}  input(A1,A2,A3)    output(ZN) ZN=AND3(A1,A2,A3)     ;
+AND4_X{1,2,4}  input(A1,A2,A3,A4) output(ZN) ZN=AND4(A1,A2,A3,A4)  ;
 NAND2_X{1,2,4} input(A1,A2)       output(ZN) ZN=NAND2(A1,A2)       ;
 NAND3_X{1,2,4} input(A1,A2,A3)    output(ZN) ZN=NAND3(A1,A2,A3)    ;
 NAND4_X{1,2,4} input(A1,A2,A3,A4) output(ZN) ZN=NAND4(A1,A2,A3,A4) ;
+OR2_X{1,2,4}   input(A1,A2)       output(ZN) ZN=OR2(A1,A2)         ;
+OR3_X{1,2,4}   input(A1,A2,A3)    output(ZN) ZN=OR3(A1,A2,A3)      ;
+OR4_X{1,2,4}   input(A1,A2,A3,A4) output(ZN) ZN=OR4(A1,A2,A3,A4)   ;
 NOR2_X{1,2,4}  input(A1,A2)       output(ZN) ZN=NOR2(A1,A2)        ;
 NOR3_X{1,2,4}  input(A1,A2,A3)    output(ZN) ZN=NOR3(A1,A2,A3)     ;
 NOR4_X{1,2,4}  input(A1,A2,A3,A4) output(ZN) ZN=NOR4(A1,A2,A3,A4)  ;
+XOR2_X{1,2}    input(A,B)         output(Z)  Z=XOR2(A,B)           ;
+XNOR2_X{1,2}   input(A,B)         output(ZN) ZN=XNOR2(A,B)         ;
+
+MUX2_X{1,2} input(A,B,S) output(Z) Z=MUX21(A,B,S) ;
+
+HA_X1 input(A,B) output(CO,S) CO=AND2(A,B) S=XOR2(A,B) ;
+FA_X1 input(A,B,CI) output(CO,S) AB=XOR2(A,B) CO=XOR2(AB,CI) S=AO22(CI,A,B) ;
 
 AOI21_X{1,2,4} input(A,B1,B2)     output(ZN) ZN=AOI21(B1,B2,A)     ;
 OAI21_X{1,2,4} input(A,B1,B2)     output(ZN) ZN=OAI21(B1,B2,A)     ;
@@ -162,8 +258,6 @@ OAI22_X{1,2,4} input(A1,A2,B1,B2) output(ZN) ZN=OAI22(A1,A2,B1,B2) ;
 OAI211_X{1,2,4} input(A,B,C1,C2) output(ZN) ZN=OAI211(C1,C2,A,B)   ;
 AOI211_X{1,2,4} input(A,B,C1,C2) output(ZN) ZN=AOI211(C1,C2,A,B)   ;
 
-MUX2_X{1,2} input(A,B,S) output(Z) Z=MUX21(A,B,S) ;
-
 AOI221_X{1,2,4} input(A,B1,B2,C1,C2) output(ZN) BC=AO22(B1,B2,C1,C2) ZN=NOR2(BC,A)  ;
 OAI221_X{1,2,4} input(A,B1,B2,C1,C2) output(ZN) BC=OA22(B1,B2,C1,C2) ZN=NAND2(BC,A) ;
 
@@ -172,14 +266,6 @@ OAI222_X{1,2,4} input(A1,A2,B1,B2,C1,C2) output(ZN) BC=OA22(B1,B2,C1,C2) ZN=OAI2
 
 OAI33_X1 input(A1,A2,A3,B1,B2,B3) output(ZN) AA=OR2(A1,A2) BB=OR2(B1,B2) ZN=OAI22(AA,A3,BB,B3) ;
 
-HA_X1 input(A,B) output(CO,S) CO=XOR2(A,B) S=AND2(A,B) ;
-
-FA_X1 input(A,B,CI) output(CO,S) AB=XOR2(A,B) CO=XOR2(AB,CI) S=AO22(CI,A,B) ;
-
-CLKGATE_X{1,2,4,8} input(CK,E) output(GCK) GCK=AND2(CK,E) ;
-
-CLKGATETST_X{1,2,4,8} input(CK,E,SE) output(GCK) GCK=OA21(CK,E,SE) ;
-
 DFF_X{1,2}   input(D,CK)       output(Q,QN)  Q=DFF(D,CK) QN=INV1(Q) ;
 DFFR_X{1,2}  input(D,RN,CK)    output(Q,QN)  DR=AND2(D,RN) Q=DFF(DR,CK) QN=INV1(Q) ;
 DFFS_X{1,2}  input(D,SN,CK)    output(Q,QN)  S=INV1(SN) DS=OR2(D,S) Q=DFF(DS,CK) QN=INV1(Q) ;
@@ -191,43 +277,16 @@ SDFFS_X{1,2}  input(D,SE,SI,SN,CK)    output(Q,QN)  S=INV1(SN) DS=OR2(D,S) DI=MU
 SDFFRS_X{1,2} input(D,RN,SE,SI,SN,CK) output(Q,QN)  S=INV1(SN) DS=OR2(D,S) DRS=AND2(DS,RN) DI=MUX21(DRS,SI,SE) Q=DFF(DI,CK) QN=INV1(Q) ;
 
 TBUF_X{1,2,4,8,16} input(A,EN)   output(Z)  Z=BUF1(A)    ;
+
 TINV_X1            input(I,EN)   output(ZN) ZN=INV1(I)   ;
 TLAT_X1            input(D,G,OE) output(Q)  Q=LATCH(D,G) ;
 
 DLH_X{1,2} input(D,G) output(Q)  Q=LATCH(D,G)            ;
 DLL_X{1,2} input(D,GN) output(Q) G=INV1(GN) Q=LATCH(D,G) ;
-"""
 
-
-NANGATE = TechLib(_nangate_common + r"""
-INV_X{1,2,4,8,16,32}  input(I) output(ZN) ZN=INV1(I) ;
-
-AND2_X{1,2,4}  input(A1,A2)       output(Z)  Z=AND2(A1,A2)        ;
-AND3_X{1,2,4}  input(A1,A2,A3)    output(Z)  Z=AND3(A1,A2,A3)     ;
-AND4_X{1,2,4}  input(A1,A2,A3,A4) output(Z)  Z=AND4(A1,A2,A3,A4)  ;
-OR2_X{1,2,4}   input(A1,A2)       output(Z)  Z=OR2(A1,A2)         ;
-OR3_X{1,2,4}   input(A1,A2,A3)    output(Z)  Z=OR3(A1,A2,A3)      ;
-OR4_X{1,2,4}   input(A1,A2,A3,A4) output(Z)  Z=OR4(A1,A2,A3,A4)   ;
-XOR2_X{1,2}    input(A1,A2)       output(Z)  Z=XOR2(A1,A2)        ;
-XNOR2_X{1,2}   input(A1,A2)       output(ZN) ZN=XNOR2(A1,A2)      ;
 """)
-"""An newer NANGATE-variant that uses 'Z' as output pin names for AND and OR gates.
+"""Nangate 45nm Open Cell Library (NangateOpenCellLibrary_PDKv1_3_v2010_12).
-"""
+This NANGATE-variant that uses 'ZN' as output pin names for AND and OR gates.
-
-
-NANGATE_ZN = TechLib(_nangate_common + r"""
-INV_X{1,2,4,8,16,32}  input(A) output(ZN) ZN=INV1(A) ;
-
-AND2_X{1,2,4}  input(A1,A2)       output(ZN) ZN=AND2(A1,A2)        ;
-AND3_X{1,2,4}  input(A1,A2,A3)    output(ZN) ZN=AND3(A1,A2,A3)     ;
-AND4_X{1,2,4}  input(A1,A2,A3,A4) output(ZN) ZN=AND4(A1,A2,A3,A4)  ;
-OR2_X{1,2,4}   input(A1,A2)       output(ZN) ZN=OR2(A1,A2)         ;
-OR3_X{1,2,4}   input(A1,A2,A3)    output(ZN) ZN=OR3(A1,A2,A3)      ;
-OR4_X{1,2,4}   input(A1,A2,A3,A4) output(ZN) ZN=OR4(A1,A2,A3,A4)   ;
-XOR2_X{1,2}    input(A,B)         output(Z)  Z=XOR2(A,B)           ;
-XNOR2_X{1,2}   input(A,B)         output(ZN) ZN=XNOR2(A,B)         ;
-""")
-"""An older NANGATE-variant that uses 'ZN' as output pin names for AND and OR gates.
 """
 
 
@@ -303,7 +362,7 @@ MUX41X{1,2}$ input(A1,A2,A3,A4,S0,S1) output(Y) A=MUX21(A1,A2,S0) B=MUX21(A3,A4,
 
 DEC24X{1,2}$ input(A0,A1) output(Y0,Y1,Y2,Y3) A0B=INV1(A0) A1B=INV1(A1) Y0=NOR2(A0,A1) Y1=AND(A0,A1B) Y2=AND(A0B,A1) Y3=AND(A0,A1) ;
 FADDX{1,2}$ input(A,B,CI) output(S,CO) AB=XOR2(A,B) CO=XOR2(AB,CI) S=AO22(AB,CI,A,B) ;
-HADDX{1,2}$ input(A0,B0) output(SO,C1) C1=XOR2(A0,B0) SO=AND2(A0,B0) ;
+HADDX{1,2}$ input(A0,B0) output(SO,C1) C1=AND2(A0,B0) SO=XOR2(A0,B0) ;
 
 {,AO}DFFARX{1,2}$ input(D,CLK,RSTB)      output(Q,QN) DR=AND2(D,RSTB) Q=DFF(DR,CLK) QN=INV1(Q) ;
 DFFASRX{1,2}$     input(D,CLK,RSTB,SETB) output(Q,QN) DR=AND2(D,RSTB) SET=INV1(SETB) DRS=OR2(DR,SET) Q=DFF(DRS,CLK) QN=INV1(Q) ;
@@ -328,7 +387,7 @@ It defines all cells except: negative-edge flip-flops, tri-state, latches, clock
 SAED90 = TechLib(r"""
 NBUFFX{2,4,8,16,32}$ input(INP) output(Z) Z=BUF1(INP) ;
 AOBUFX{1,2,4}$       input(INP) output(Z) Z=BUF1(INP) ;
-DELLN{1,2,3}X2$      input(INP) output(Z)Z=BUF1(INP) ;
+DELLN{1,2,3}X2$      input(INP) output(Z) Z=BUF1(INP) ;
 
 INVX{0,1,2,4,8,16,32}$ input(INP) output(ZN) ZN=INV1(INP) ;
 AOINVX{1,2,4}$         input(INP) output(ZN) ZN=INV1(INP) ;
@@ -393,7 +452,7 @@ MUX41X{1,2}$ input(IN1,IN2,IN3,IN4,S0,S1) output(Q) A=MUX21(IN1,IN2,S0) B=MUX21(
 
 DEC24X{1,2}$ input(IN1,IN2) output(Q0,Q1,Q2,Q3) IN1B=INV1(IN1) IN2B=INV1(IN2) Q0=NOR2(IN1,IN2) Q1=AND(IN1,IN2B) Q2=AND(IN1B,IN2) Q3=AND(IN1,IN2) ;
 FADDX{1,2}$ input(A,B,CI) output(S,CO) AB=XOR2(A,B) CO=XOR2(AB,CI) S=AO22(AB,CI,A,B) ;
-HADDX{1,2}$ input(A0,B0) output(SO,C1) C1=XOR2(A0,B0) SO=AND2(A0,B0) ;
+HADDX{1,2}$ input(A0,B0) output(SO,C1) C1=AND2(A0,B0) SO=XOR2(A0,B0) ;
 
 {,AO}DFFARX{1,2}$ input(D,CLK,RSTB)      output(Q,QN) DR=AND2(D,RSTB) Q=DFF(DR,CLK) QN=INV1(Q) ;
 DFFASRX{1,2}$     input(D,CLK,RSTB,SETB) output(Q,QN) DR=AND2(D,RSTB) SET=INV1(SETB) DRS=OR2(DR,SET) Q=DFF(DRS,CLK) QN=INV1(Q) ;
@@ -413,3 +472,196 @@ LATCHX{1,2}$ input(D,CLK) output(Q,QN) Q=LATCH(D,CLK) QN=INV1(Q) ;
 """The SAED 90nm educational technology library.
 It defines all cells except: negative-edge flip-flops, tri-state, latches, clock gating, level shifters
 """
+
+
+SKY130 = TechLib(r"""
+$decap_{3,4,6,8,12}             ;
+$fill_{1,2,4,8}                 ;
+$tap_{1,2}                      ;
+$tapvgnd_1                      ;
+$tapvgnd2_1                     ;
+$tapvpwrvgnd_1                  ;
+$lpflow_decapkapwr_{3,4,6,8,12} ;
+$lpflow_bleeder_1  input(SHORT) ;
+
+$conb_1  output(HI,LO) HI=__const1__() LO=__const0__() ;
+$macro_sparecell  output(LO) LO=__const0__() ;
+
+$diode_2      input(DIODE) ;
+$probe_p_8    input(A) output(X) X=BUF1(A) ;
+$probec_p_8   input(A) output(X) X=BUF1(A) ;
+
+$inv_{1,2,4,6,8,12,16}                   input(A) output(Y) Y=INV1(A) ;
+$buf_{1,2,4,6,8,12,16}                   input(A) output(X) X=BUF1(A) ;
+$bufbuf_{8,16}                           input(A) output(X) X=BUF1(A) ;
+$bufinv_{8,16}                           input(A) output(Y) Y=INV1(A) ;
+$clkbuf_{1,2,4,8,16}                     input(A) output(X) X=BUF1(A) ;
+$clkinv_{1,2,4,8,16}                     input(A) output(Y) Y=INV1(A) ;
+$clkinvlp_{2,4}                          input(A) output(Y) Y=INV1(A) ;
+$clkdlybuf4s15_{1,2}                     input(A) output(X) X=BUF1(A) ;
+$clkdlybuf4s18_{1,2}                     input(A) output(X) X=BUF1(A) ;
+$clkdlybuf4s25_{1,2}                     input(A) output(X) X=BUF1(A) ;
+$clkdlybuf4s50_{1,2}                     input(A) output(X) X=BUF1(A) ;
+$dlygate4sd1_1                           input(A) output(X) X=BUF1(A) ;
+$dlygate4sd2_1                           input(A) output(X) X=BUF1(A) ;
+$dlygate4sd3_1                           input(A) output(X) X=BUF1(A) ;
+$dlymetal6s2s_1                          input(A) output(X) X=BUF1(A) ;
+$dlymetal6s4s_1                          input(A) output(X) X=BUF1(A) ;
+$dlymetal6s6s_1                          input(A) output(X) X=BUF1(A) ;
+$lpflow_clkbufkapwr_{1,2,4,8,16}         input(A) output(X) X=BUF1(A) ;
+$lpflow_clkinvkapwr_{1,2,4,8,16}         input(A) output(Y) Y=INV1(A) ;
+$lpflow_lsbuf_lh_hl_isowell_tap_{1,2,4}  input(A) output(X) X=BUF1(A) ;
+$lpflow_lsbuf_lh_isowell_4               input(A) output(X) X=BUF1(A) ;
+$lpflow_lsbuf_lh_isowell_tap_{1,2,4}     input(A) output(X) X=BUF1(A) ;
+
+$and2_{0,1,2,4}    input(A,B)          output(X) X=AND2(A,B)                                    ;
+$and2b_{1,2,4}     input(A_N,B)        output(X) AN=INV1(A_N) X=AND2(AN,B)                      ;
+$and3_{1,2,4}      input(A,B,C)        output(X) X=AND3(A,B,C)                                  ;
+$and3b_{1,2,4}     input(A_N,B,C)      output(X) AN=INV1(A_N) X=AND3(AN,B,C)                    ;
+$and4_{1,2,4}      input(A,B,C,D)      output(X) X=AND4(A,B,C,D)                                ;
+$and4b_{1,2,4}     input(A_N,B,C,D)    output(X) AN=INV1(A_N) X=AND4(AN,B,C,D)                  ;
+$and4bb_{1,2,4}    input(A_N,B_N,C,D)  output(X) AN=INV1(A_N) BN=INV1(B_N) X=AND4(AN,BN,C,D)    ;
+
+$or2_{0,1,2,4}     input(A,B)          output(X) X=OR2(A,B)                                     ;
+$or2b_{1,2,4}      input(A,B_N)        output(X) BN=INV1(B_N) X=OR2(A,BN)                       ;
+$or3_{1,2,4}       input(A,B,C)        output(X) X=OR3(A,B,C)                                   ;
+$or3b_{1,2,4}      input(A,B,C_N)      output(X) CN=INV1(C_N) X=OR3(A,B,CN)                     ;
+$or4_{1,2,4}       input(A,B,C,D)      output(X) X=OR4(A,B,C,D)                                 ;
+$or4b_{1,2,4}      input(A,B,C,D_N)    output(X) DN=INV1(D_N) X=OR4(A,B,C,DN)                   ;
+$or4bb_{1,2,4}     input(A,B,C_N,D_N)  output(X) CN=INV1(C_N) DN=INV1(D_N) X=OR4(A,B,CN,DN)     ;
+
+$nand2_{1,2,4,8}   input(A,B)          output(Y) Y=NAND2(A,B)                                   ;
+$nand2b_{1,2,4}    input(A_N,B)        output(Y) AN=INV1(A_N) Y=NAND2(AN,B)                     ;
+$nand3_{1,2,4}     input(A,B,C)        output(Y) Y=NAND3(A,B,C)                                 ;
+$nand3b_{1,2,4}    input(A_N,B,C)      output(Y) AN=INV1(A_N) Y=NAND3(AN,B,C)                   ;
+$nand4_{1,2,4}     input(A,B,C,D)      output(Y) Y=NAND4(A,B,C,D)                               ;
+$nand4b_{1,2,4}    input(A_N,B,C,D)    output(Y) AN=INV1(A_N) Y=NAND4(AN,B,C,D)                 ;
+$nand4bb_{1,2,4}   input(A_N,B_N,C,D)  output(Y) AN=INV1(A_N) BN=INV1(B_N) Y=NAND4(AN,BN,C,D)   ;
+
+$nor2_{1,2,4,8}    input(A,B)          output(Y) Y=NOR2(A,B)                                    ;
+$nor2b_{1,2,4}     input(A,B_N)        output(Y) BN=INV1(B_N) Y=NOR2(A,BN)                      ;
+$nor3_{1,2,4}      input(A,B,C)        output(Y) Y=NOR3(A,B,C)                                  ;
+$nor3b_{1,2,4}     input(A,B,C_N)      output(Y) CN=INV1(C_N) Y=NOR3(A,B,CN)                    ;
+$nor4_{1,2,4}      input(A,B,C,D)      output(Y) Y=NOR4(A,B,C,D)                                ;
+$nor4b_{1,2,4}     input(A,B,C,D_N)    output(Y) DN=INV1(D_N) Y=NOR4(A,B,C,DN)                  ;
+$nor4bb_{1,2,4}    input(A,B,C_N,D_N)  output(Y) CN=INV1(C_N) DN=INV1(D_N) Y=NOR4(A,B,CN,DN)    ;
+
+$xor2_{1,2,4}      input(A,B)   output(X) X=XOR2(A,B)   ;
+$xor3_{1,2,4}      input(A,B,C) output(X) X=XOR3(A,B,C) ;
+$xnor2_{1,2,4}     input(A,B)   output(Y) Y=XNOR2(A,B)  ;
+$xnor3_{1,2,4}     input(A,B,C) output(X) X=XNOR3(A,B,C) ;
+
+$maj3_{1,2,4}  input(A,B,C) output(X) AB=AND2(A,B) BC=AND2(B,C) AC=AND2(A,C) X=OR3(AB,BC,AC) ;
+
+$mux2_{1,2,4,8}  input(A0,A1,S)           output(X) X=MUX21(A0,A1,S)                                    ;
+$mux2i_{1,2,4}   input(A0,A1,S)           output(Y) M=MUX21(A0,A1,S) Y=INV1(M)                          ;
+$mux4_{1,2,4}    input(A0,A1,A2,A3,S0,S1) output(X) M0=MUX21(A0,A1,S0) M1=MUX21(A2,A3,S0) X=MUX21(M0,M1,S1) ;
+
+$ha_{1,2,4}      input(A,B)     output(COUT,SUM) SUM=XOR2(A,B) COUT=AND2(A,B)                           ;
+$fa_{1,2,4}      input(A,B,CIN) output(COUT,SUM) AB=XOR2(A,B) SUM=XOR2(AB,CIN) COUT=AO22(A,B,AB,CIN)    ;
+$fah_1           input(A,B,CI)  output(COUT,SUM) AB=XOR2(A,B) SUM=XOR2(AB,CI)  COUT=AO22(A,B,AB,CI)     ;
+$fahcin_1        input(A,B,CIN) output(COUT,SUM) AB=XOR2(A,B) SUM=XOR2(AB,CIN) COUT=AO22(A,B,AB,CIN)    ;
+$fahcon_1        input(A,B,CI)  output(COUT_N,SUM) AB=XOR2(A,B) SUM=XOR2(AB,CI) T=AO22(A,B,AB,CI) COUT_N=INV1(T) ;
+
+$a21o_{1,2,4}     input(A1,A2,B1)              output(X) X=AO21(A1,A2,B1)                               ;
+$a21bo_{1,2,4}    input(A1,A2,B1_N)            output(X) BN=INV1(B1_N) X=AO21(A1,A2,BN)                 ;
+$a22o_{1,2,4}     input(A1,A2,B1,B2)           output(X) X=AO22(A1,A2,B1,B2)                            ;
+$a2bb2o_{1,2,4}   input(A1_N,A2_N,B1,B2)       output(X) AN1=INV1(A1_N) AN2=INV1(A2_N) X=AO22(AN1,AN2,B1,B2) ;
+$a211o_{1,2,4}    input(A1,A2,B1,C1)           output(X) X=AO211(A1,A2,B1,C1)                           ;
+$a221o_{1,2,4}    input(A1,A2,B1,B2,C1)        output(X) T=AO22(A1,A2,B1,B2) X=OR2(T,C1)                ;
+$a31o_{1,2,4}     input(A1,A2,A3,B1)           output(X) AA=AND3(A1,A2,A3) X=OR2(AA,B1)                 ;
+$a311o_{1,2,4}    input(A1,A2,A3,B1,C1)        output(X) T=AND2(A1,A2) X=AO211(T,A3,B1,C1)              ;
+$a32o_{1,2,4}     input(A1,A2,A3,B1,B2)        output(X) T=AND2(A1,A2) X=AO22(T,A3,B1,B2)               ;
+$a41o_{1,2,4}     input(A1,A2,A3,A4,B1)        output(X) T1=AND2(A1,A2) T2=AND2(T1,A3) X=AO21(T2,A4,B1) ;
+$a2111o_{1,2,4}   input(A1,A2,B1,C1,D1)        output(X) T=AO211(A1,A2,B1,C1) X=OR2(T,D1)               ;
+
+$a21oi_{1,2,4}    input(A1,A2,B1)              output(Y) Y=AOI21(A1,A2,B1)                               ;
+$a21boi_{0,1,2,4} input(A1,A2,B1_N)            output(Y) BN=INV1(B1_N) Y=AOI21(A1,A2,BN)                 ;
+$a22oi_{1,2,4}    input(A1,A2,B1,B2)           output(Y) Y=AOI22(A1,A2,B1,B2)                            ;
+$a2bb2oi_{1,2,4}  input(A1_N,A2_N,B1,B2)       output(Y) AN1=INV1(A1_N) AN2=INV1(A2_N) Y=AOI22(AN1,AN2,B1,B2) ;
+$a211oi_{1,2,4}   input(A1,A2,B1,C1)           output(Y) Y=AOI211(A1,A2,B1,C1)                           ;
+$a221oi_{1,2,4}   input(A1,A2,B1,B2,C1)        output(Y) T=AO22(A1,A2,B1,B2) Y=NOR2(T,C1)                ;
+$a222oi_1         input(A1,A2,B1,B2,C1,C2)     output(Y) AB=AO22(A1,A2,B1,B2) Y=AOI21(C1,C2,AB)          ;
+$a31oi_{1,2,4}    input(A1,A2,A3,B1)           output(Y) AA=AND3(A1,A2,A3) Y=NOR2(AA,B1)                 ;
+$a311oi_{1,2,4}   input(A1,A2,A3,B1,C1)        output(Y) T=AND2(A1,A2) Y=AOI211(T,A3,B1,C1)              ;
+$a32oi_{1,2,4}    input(A1,A2,A3,B1,B2)        output(Y) T=AND2(A1,A2) Y=AOI22(T,A3,B1,B2)               ;
+$a41oi_{1,2,4}    input(A1,A2,A3,A4,B1)        output(Y) T1=AND2(A1,A2) T2=AND2(T1,A3) Y=AOI21(T2,A4,B1) ;
+$a2111oi_{0,1,2,4} input(A1,A2,B1,C1,D1)       output(Y) T=AO211(A1,A2,B1,C1) Y=NOR2(T,D1)               ;
+
+$o21a_{1,2,4}     input(A1,A2,B1)              output(X) X=OA21(A1,A2,B1)                              ;
+$o21ba_{1,2,4}    input(A1,A2,B1_N)            output(X) BN=INV1(B1_N) X=OA21(A1,A2,BN)                ;
+$o22a_{1,2,4}     input(A1,A2,B1,B2)           output(X) X=OA22(A1,A2,B1,B2)                           ;
+$o2bb2a_{1,2,4}   input(A1_N,A2_N,B1,B2)       output(X) T=NAND2(A1_N,A2_N) X=OA21(B1,B2,T)            ;
+$o211a_{1,2,4}    input(A1,A2,B1,C1)           output(X) X=OA211(A1,A2,B1,C1)                          ;
+$o221a_{1,2,4}    input(A1,A2,B1,B2,C1)        output(X) T=OA22(A1,A2,B1,B2) X=AND2(T,C1)              ;
+$o31a_{1,2,4}     input(A1,A2,A3,B1)           output(X) T=OR3(A1,A2,A3) X=AND2(T,B1)                  ;
+$o311a_{1,2,4}    input(A1,A2,A3,B1,C1)        output(X) T=OR2(A1,A2) X=OA211(T,A3,B1,C1)              ;
+$o32a_{1,2,4}     input(A1,A2,A3,B1,B2)        output(X) T=OR3(A1,A2,A3) X=OA21(B1,B2,T)               ;
+$o41a_{1,2,4}     input(A1,A2,A3,A4,B1)        output(X) T=OR4(A1,A2,A3,A4) X=AND2(T,B1)               ;
+$o2111a_{1,2,4}   input(A1,A2,B1,C1,D1)        output(X) T=OA211(A1,A2,B1,C1) X=AND2(T,D1)             ;
+
+$o21ai_{0,1,2,4}  input(A1,A2,B1)              output(Y) Y=OAI21(A1,A2,B1)                             ;
+$o21bai_{1,2,4}   input(A1,A2,B1_N)            output(Y) BN=INV1(B1_N) Y=OAI21(A1,A2,BN)               ;
+$o22ai_{1,2,4}    input(A1,A2,B1,B2)           output(Y) Y=OAI22(A1,A2,B1,B2)                          ;
+$o2bb2ai_{1,2,4}  input(A1_N,A2_N,B1,B2)       output(Y) T=NAND2(A1_N,A2_N) Y=OAI21(B1,B2,T)           ;
+$o211ai_{1,2,4}   input(A1,A2,B1,C1)           output(Y) Y=OAI211(A1,A2,B1,C1)                         ;
+$o221ai_{1,2,4}   input(A1,A2,B1,B2,C1)        output(Y) T=OA22(A1,A2,B1,B2) Y=NAND2(T,C1)             ;
+$o31ai_{1,2,4}    input(A1,A2,A3,B1)           output(Y) T=OR3(A1,A2,A3) Y=NAND2(T,B1)                 ;
+$o311ai_{0,1,2,4} input(A1,A2,A3,B1,C1)        output(Y) T=OR2(A1,A2) Y=OAI211(T,A3,B1,C1)             ;
+$o32ai_{1,2,4}    input(A1,A2,A3,B1,B2)        output(Y) T=OR3(A1,A2,A3) Y=OAI21(B1,B2,T)              ;
+$o41ai_{1,2,4}    input(A1,A2,A3,A4,B1)        output(Y) T=OR4(A1,A2,A3,A4) Y=NAND2(T,B1)              ;
+$o2111ai_{1,2,4}  input(A1,A2,B1,C1,D1)        output(Y) T=OA211(A1,A2,B1,C1) Y=NAND2(T,D1)            ;
+
+$dfxtp_{1,2,4}    input(CLK,D)                 output(Q)     Q=DFF(D,CLK)                                                              ;
+$dfxbp_{1,2}      input(CLK,D)                 output(Q,Q_N) Q=DFF(D,CLK) Q_N=INV1(Q)                                                  ;
+$dfrtp_{1,2,4}    input(CLK,D,RESET_B)         output(Q)     DR=AND2(D,RESET_B) Q=DFF(DR,CLK)                                          ;
+$dfrbp_{1,2}      input(CLK,D,RESET_B)         output(Q,Q_N) DR=AND2(D,RESET_B) Q=DFF(DR,CLK) Q_N=INV1(Q)                              ;
+$dfstp_{1,2,4}    input(CLK,D,SET_B)           output(Q)     S=INV1(SET_B) DS=OR2(D,S) Q=DFF(DS,CLK)                                   ;
+$dfsbp_{1,2}      input(CLK,D,SET_B)           output(Q,Q_N) S=INV1(SET_B) DS=OR2(D,S) Q=DFF(DS,CLK) Q_N=INV1(Q)                       ;
+$dfbbp_1          input(CLK,D,SET_B,RESET_B)   output(Q,Q_N) DR=AND2(D,RESET_B) S=INV1(SET_B) DRS=OR2(DR,S) Q=DFF(DRS,CLK) Q_N=INV1(Q) ;
+$dfrtn_1          input(CLK_N,D,RESET_B)       output(Q)     CLKN=INV1(CLK_N) DR=AND2(D,RESET_B) Q=DFF(DR,CLKN)                        ;
+$dfbbn_{1,2}      input(CLK_N,D,SET_B,RESET_B) output(Q,Q_N) CLKN=INV1(CLK_N) DR=AND2(D,RESET_B) S=INV1(SET_B) DRS=OR2(DR,S) Q=DFF(DRS,CLKN) Q_N=INV1(Q) ;
+
+$sdfxtp_{1,2,4}   input(CLK,D,SCD,SCE)                 output(Q)     DI=MUX21(D,SCD,SCE) Q=DFF(DI,CLK)                                        ;
+$sdfxbp_{1,2}     input(CLK,D,SCD,SCE)                 output(Q,Q_N) DI=MUX21(D,SCD,SCE) Q=DFF(DI,CLK) Q_N=INV1(Q)                            ;
+$sdfrtp_{1,2,4}   input(CLK,D,SCD,SCE,RESET_B)         output(Q)     DR=AND2(D,RESET_B) DI=MUX21(DR,SCD,SCE) Q=DFF(DI,CLK)                    ;
+$sdfrbp_{1,2}     input(CLK,D,SCD,SCE,RESET_B)         output(Q,Q_N) DR=AND2(D,RESET_B) DI=MUX21(DR,SCD,SCE) Q=DFF(DI,CLK) Q_N=INV1(Q)        ;
+$sdfstp_{1,2,4}   input(CLK,D,SCD,SCE,SET_B)           output(Q)     S=INV1(SET_B) DS=OR2(D,S) DI=MUX21(DS,SCD,SCE) Q=DFF(DI,CLK)             ;
+$sdfsbp_{1,2}     input(CLK,D,SCD,SCE,SET_B)           output(Q,Q_N) S=INV1(SET_B) DS=OR2(D,S) DI=MUX21(DS,SCD,SCE) Q=DFF(DI,CLK) Q_N=INV1(Q) ;
+$sdfbbp_1         input(CLK,D,SCD,SCE,SET_B,RESET_B)   output(Q,Q_N) DR=AND2(D,RESET_B) S=INV1(SET_B) DRS=OR2(DR,S) DI=MUX21(DRS,SCD,SCE) Q=DFF(DI,CLK) Q_N=INV1(Q) ;
+$sdfrtn_1         input(CLK_N,D,SCD,SCE,RESET_B)       output(Q)     CLKN=INV1(CLK_N) DR=AND2(D,RESET_B) DI=MUX21(DR,SCD,SCE) Q=DFF(DI,CLKN) ;
+$sdfbbn_{1,2}     input(CLK_N,D,SCD,SCE,SET_B,RESET_B) output(Q,Q_N) CLKN=INV1(CLK_N) DR=AND2(D,RESET_B) S=INV1(SET_B) DRS=OR2(DR,S) DI=MUX21(DRS,SCD,SCE) Q=DFF(DI,CLKN) Q_N=INV1(Q) ;
+
+$edfxtp_1         input(CLK,D,DE)                    output(Q)     GCK=AND2(CLK,DE) Q=DFF(D,GCK)                                           ;
+$edfxbp_1         input(CLK,D,DE)                    output(Q,Q_N) GCK=AND2(CLK,DE) Q=DFF(D,GCK) Q_N=INV1(Q)                               ;
+$sedfxtp_{1,2,4}  input(CLK,D,DE,SCD,SCE)            output(Q)     DI=MUX21(D,SCD,SCE) GCK=AND2(CLK,DE) Q=DFF(DI,GCK)                      ;
+$sedfxbp_{1,2}    input(CLK,D,DE,SCD,SCE)            output(Q,Q_N) DI=MUX21(D,SCD,SCE) GCK=AND2(CLK,DE) Q=DFF(DI,GCK) Q_N=INV1(Q)          ;
+
+$dlxtp_1          input(D,GATE)                      output(Q)     Q=LATCH(D,GATE)                                                         ;
+$dlxtn_{1,2,4}    input(D,GATE_N)                    output(Q)     GN=INV1(GATE_N) Q=LATCH(D,GN)                                           ;
+$dlxbp_1          input(D,GATE)                      output(Q,Q_N) Q=LATCH(D,GATE) Q_N=INV1(Q)                                             ;
+$dlxbn_{1,2}      input(D,GATE_N)                    output(Q,Q_N) GN=INV1(GATE_N) Q=LATCH(D,GN) Q_N=INV1(Q)                               ;
+$dlrtp_{1,2,4}    input(D,GATE,RESET_B)              output(Q)     DR=AND2(D,RESET_B) Q=LATCH(DR,GATE)                                     ;
+$dlrtn_{1,2,4}    input(D,GATE_N,RESET_B)            output(Q)     GN=INV1(GATE_N) DR=AND2(D,RESET_B) Q=LATCH(DR,GN)                       ;
+$dlrbp_{1,2}      input(D,GATE,RESET_B)              output(Q,Q_N) DR=AND2(D,RESET_B) Q=LATCH(DR,GATE) Q_N=INV1(Q)                         ;
+$dlrbn_{1,2}      input(D,GATE_N,RESET_B)            output(Q,Q_N) GN=INV1(GATE_N) DR=AND2(D,RESET_B) Q=LATCH(DR,GN) Q_N=INV1(Q)           ;
+
+$dlclkp_{1,2,4}   input(CLK,GATE)                    output(GCLK) GCLK=AND2(CLK,GATE)                                                      ;
+$sdlclkp_{1,2,4}  input(CLK,GATE,SCE)                output(GCLK) G=OR2(GATE,SCE) GCLK=AND2(CLK,G)                                         ;
+
+$ebufn_{1,2,4,8}    input(A,TE_B) output(Z) TE=INV1(TE_B) Z=AND2(A,TE)               ;
+$einvn_{0,1,2,4,8}  input(A,TE_B) output(Z) TE=INV1(TE_B) ZI=INV1(A) Z=AND2(ZI,TE)   ;
+$einvp_{1,2,4,8}    input(A,TE)   output(Z) ZI=INV1(A) Z=AND2(ZI,TE)                 ;
+
+$lpflow_inputiso0n_1  input(A,SLEEP_B) output(X) X=AND2(A,SLEEP_B)                   ;
+$lpflow_inputiso0p_1  input(A,SLEEP)   output(X) SB=INV1(SLEEP) X=AND2(A,SB)         ;
+$lpflow_inputiso1n_1  input(A,SLEEP_B) output(X) X=AND2(A,SLEEP_B)                   ;
+$lpflow_inputiso1p_1  input(A,SLEEP)   output(X) SB=INV1(SLEEP) X=AND2(A,SB)         ;
+$lpflow_isobufsrc_{1,2,4,8,16}   input(A,SLEEP) output(X) AI=INV1(A) X=OR2(AI,SLEEP) ;
+$lpflow_isobufsrckapwr_16        input(A,SLEEP) output(X) AI=INV1(A) X=OR2(AI,SLEEP) ;
+$lpflow_inputisolatch_1  input(D,SLEEP_B) output(Q) Q=LATCH(D,SLEEP_B)               ;
+""".replace('$','sky130_fd_sc_hd__'))
+"""SkyWater 130nm High Density Digital Standard Cells (skywater-pdk-libs-sky130_fd_sc_hd).
+"""
+
+techlib_by_name = {k: v for k, v in globals().items() if isinstance(v, TechLib)}

src/kyupy/vcd.py

@@ -0,0 +1,218 @@
+"""A simple parser for Verilog Change Dump (VCD) files.
+
+This parser loads the changes into a ndarray of logic values.
+Axis 0 is the time step, axis 1 is the variable data.
+All variable values are flattened. Offsets are stored in each variable metadata.
+"""
+
+from collections import namedtuple
+from dataclasses import dataclass, field
+
+from lark import Lark, Transformer
+import numpy as np
+
+from . import readtext, logic
+
+
+@dataclass
+class Var:
+    type: str
+    width: int
+    idcode: str
+    reference: str
+    scope: 'Scope'
+    locs: list[int] = field(default_factory=list)
+
+
+class Scope:
+    def __init__(self, parent, name='', type='module') -> None:
+        self.name = name
+        self.type = type
+        self.parent_scope = parent
+        self.sub_scopes : list[Scope] = []
+        self.vars : list = []
+
+    @property
+    def path(self) -> str:
+        parts = []
+        s = self
+        while s is not None and s.name:
+            parts.append(s.name)
+            s = s.parent_scope
+        return '/'.join(reversed(parts))
+
+    def __repr__(self) -> str:
+        return self.path
+
+
+class VcdHeader:
+    def __init__(self) -> None:
+        self.comment: str = ''
+        self.date: str = ''
+        self.timescale: str = ''
+        self.version: str = ''
+        self.root_scope = Scope(None)
+        self.current_scope = self.root_scope
+
+    def scope(self, name, type):
+        new_scope = Scope(self.current_scope, name, type)
+        self.current_scope.sub_scopes.append(new_scope)
+        self.current_scope = new_scope
+
+    def upscope(self):
+        assert self.current_scope.parent_scope is not None, "root scope has no parent"
+        self.current_scope = self.current_scope.parent_scope
+
+    def add_var(self, type: str, width: int, idcode: str, reference: str):
+        self.current_scope.vars.append(Var(type, width, idcode, reference, self.current_scope))
+
+    def __repr__(self):
+        lines = [
+            'VcdHeader(',
+            f'  comment   = {self.comment!r}',
+            f'  date      = {self.date!r}',
+            f'  timescale = {self.timescale!r}',
+            f'  version   = {self.version!r}',
+            '  scopes:',
+        ]
+        def fmt_scope(scope, indent=4):
+            lines.append(f'{"  " * (indent // 2)}{scope.type}:{scope.name} ({len(scope.vars)} vars)')
+            for sub in scope.sub_scopes:
+                fmt_scope(sub, indent + 2)
+        for sub in self.root_scope.sub_scopes:
+            fmt_scope(sub)
+        lines.append(')')
+        return '\n'.join(lines)
+
+
+class VcdVarMap:
+    def __init__(self, header: VcdHeader, var_locs = lambda _: []) -> None:
+        self.var_list: list[Var] = []
+        def collect(scope):
+            for v in scope.vars:
+                v.locs = var_locs(v)
+                if v.locs is not None and len(v.locs) > 0:
+                    assert len(v.locs) == v.width, f'{v.reference}: len(locs)={len(v.locs)} != width={v.width}'
+                    self.var_list.append(v)
+            for sub in scope.sub_scopes:
+                collect(sub)
+        collect(header.root_scope)
+        self.total_width = max((max(v.locs) for v in self.var_list), default=-1) + 1
+        self.idcode2var = {var.idcode: var for var in self.var_list}
+
+
+class VcdData:
+    def __init__(self, var_map: VcdVarMap, steps, data):
+        self.var_map = var_map
+        self.steps = steps
+        self.data = data
+    def __repr__(self) -> str:
+        return f'{{variables: {len(self.var_map.var_list)}, bits: {self.data.shape[0]}, steps: {self.data.shape[1]}}}'
+
+
+class VcdHeaderTransformer(Transformer):
+    def __init__(self):
+        super().__init__()
+        self.header = VcdHeader()
+
+    def comment(self, args): self.header.comment = args[0].value.strip()
+    def date(self, args): self.header.date = args[0].value.strip()
+    def timescale(self, args): self.header.timescale = args[0].value.strip()
+    def version(self, args): self.header.version = args[0].value.strip()
+
+    def scope(self, args):
+        type, name = args
+        self.header.scope(name.value, type.value)
+
+    def upscope(self, args):
+        self.header.upscope()
+
+    def var(self, args):
+        type, width, idcode, reference = args
+        reference = reference.value.strip()
+        width = int(width)
+        self.header.add_var(type.value, width, idcode.value, reference)
+
+    def start(self, args):
+        return self.header
+        #return tuple(args) if len(args) > 1 else (args[0], None)
+
+
+GRAMMAR = r"""
+    start: (_declaration_command)*
+    _declaration_command: ( comment | date | scope | timescale | upscope | var | version ) "$end"
+    comment: "$comment" TEXT?
+    date: "$date" TEXT
+    scope: "$scope" /[^\s]+/ /[^\s]+/
+    timescale: "$timescale" TEXT
+    upscope: "$upscope"
+    var: "$var" /[^\s]+/ /[^\s]+/ /[^\s]+/ TEXT
+    version: "$version" TEXT
+    TEXT: /(?:(?!\$end)[\s\S])+/
+    %ignore ( /\r?\n/ )+
+    %ignore /[\t \f]+/
+    """
+
+
+def load(file, var_locs = lambda _: [], step_filter = lambda *_: True, s_len = None) -> VcdData:
+    """Parses the contents of ``file`` as Verilog Change Dump (VCD).
+
+    :param file: A file name or a file handle. Files with `.gz`-suffix are decompressed on-the-fly.
+    :param var_locs: A callback ``(var) -> list[int]`` mapping each variable to ndarray column indices. Empty list drops the variable.
+    :param step_filter: A callback ``(time, values, var_map) -> bool`` to select which timesteps to include.
+    :param s_len: length of the first axis of returned pattern data.
+    :return: A VcdData object with metadata and an ndarray with all values.
+    """
+    vcd = readtext(file)
+    header_size = vcd.find('$enddefinitions')
+    assert header_size > 0, "invalid VCD file: end of header not found"
+    vcd_header_str = vcd[:header_size]
+    vcd_header : VcdHeader = Lark(GRAMMAR, parser="lalr", lexer='contextual', transformer=VcdHeaderTransformer()).parse(vcd_header_str) # type: ignore
+    vcd_data = vcd[header_size:].splitlines()
+    var_map = VcdVarMap(vcd_header, var_locs)
+    width = s_len if s_len is not None else var_map.total_width
+
+    chunk_size = 10240
+    chunks = []
+    chunk = np.full((chunk_size, width), logic.UNASSIGNED, dtype=np.uint8)
+
+    _val_map = {'0': logic.ZERO, '1': logic.ONE,
+                'x': logic.UNKNOWN, 'X': logic.UNKNOWN,
+                'z': logic.UNASSIGNED, 'Z': logic.UNASSIGNED}
+    _pad_char = {'0': '0', '1': '0', 'x': 'x', 'X': 'x', 'z': 'z', 'Z': 'z'}
+
+    current_time = None
+    steps = []
+    step_idx = 0
+    for line in vcd_data:
+        if not line or line[0] == '$':
+            continue
+        if line[0] == '#':
+            if step_filter(current_time, chunk[step_idx], var_map):
+                step_idx += 1
+                steps.append(current_time)
+                if step_idx >= chunk_size:
+                    chunks.append(chunk)
+                    chunk = np.empty((chunk_size, width), dtype=np.uint8)
+                    chunk[0] = chunks[-1][-1]
+                    step_idx = 0
+                else:
+                    chunk[step_idx] = chunk[step_idx - 1]
+            current_time = int(line[1:])
+            continue
+        if line[0] in ('b', 'B'):
+            value_str, idcode = line[1:].split(None, 1)
+        else:
+            value_str, idcode = line[0], line[1:]
+        var = var_map.idcode2var.get(idcode)
+        if var is None:
+            continue
+        if len(value_str) < var.width:
+            value_str = _pad_char.get(value_str[0], '0') * (var.width - len(value_str)) + value_str
+        for i, c in enumerate(value_str):
+            if var.locs[i] >= 0:
+                chunk[step_idx, var.locs[i]] = _val_map.get(c, logic.UNKNOWN)
+
+    chunks.append(chunk[:step_idx])
+    data = np.concatenate(chunks, axis=0).T
+    return VcdData(var_map, steps, data)

src/kyupy/verilog.py

@@ -10,7 +10,7 @@ from lark import Lark, Transformer, Tree
 
 from . import log, readtext
 from .circuit import Circuit, Node, Line
-from .techlib import NANGATE
+from .techlib import KYUPY
 
 Instantiation = namedtuple('Instantiation', ['type', 'name', 'pins'])
 
@@ -35,7 +35,7 @@ class SignalDeclaration:
 
 
 class VerilogTransformer(Transformer):
-    def __init__(self, branchforks=False, tlib=NANGATE):
+    def __init__(self, branchforks=False, tlib=KYUPY):
         super().__init__()
         self.branchforks = branchforks
         self.tlib = tlib
@@ -52,15 +52,20 @@ class VerilogTransformer(Transformer):
     @staticmethod
     def instantiation(args):
         pinmap = {}
-        for idx, pin in enumerate(args[2:]):
+        if args[2] is not None:
-            p = pin.children[0]
+            for idx, pin in enumerate(args[2:]):
-            if isinstance(p, tuple):  # named pin
+                p = pin.children[0]
-                if p[1] is not None:
+                if isinstance(p, tuple):  # named pin
-                    pinmap[p[0]] = p[1]
+                    if p[1] is not None:
-            else:  # unnamed pin
+                        pinmap[p[0]] = p[1]
-                pinmap[idx] = p
+                else:  # unnamed pin
+                    pinmap[idx] = p
         return Instantiation(args[0], args[1], pinmap)
 
+    @staticmethod
+    def simple_dff(args):
+        return Instantiation('DFF', args[1], {'d': args[2], 'clk': args[0], 'q': args[1]} )
+
     def range(self, args):
         left = int(args[0].value)
         right = int(args[1].value) if len(args) > 1 else left
@@ -73,7 +78,7 @@ class VerilogTransformer(Transformer):
         elif "'" in args[0]:
             width, rest = args[0].split("'")
             width = int(width)
-            base, const = rest[0], rest[1:]
+            base, const = rest[0], rest[1:].replace('x','0')
             const = int(const, {'b': 2, 'd':10, 'h':16}[base.lower()])
             l = []
             for _ in range(width):
@@ -92,6 +97,13 @@ class VerilogTransformer(Transformer):
                 sigs.append(a)
         return sigs
 
+    def ternaryif(self, args):
+        sel = args[0]
+        ctrue = args[1]
+        cfalse = args[2]
+        log.warn(f"FIXME: not implemented: ternary if {args[0]} {args[1]}")
+        return args[1]
+
     def declaration(self, kind, args):
         rnge = None
         if isinstance(args[0], range):
@@ -103,6 +115,7 @@ class VerilogTransformer(Transformer):
     def output(self, args): return self.declaration("output", args)
     def inout(self, args): return self.declaration("input", args)  # just treat as input
     def wire(self, args): return self.declaration("wire", args)
+    def reg(self, args): pass #return self.declaration("wire", args)  # treat as wire
 
     def module(self, args):
         c = Circuit(args[0])
@@ -123,6 +136,9 @@ class VerilogTransformer(Transformer):
         assignments = []
         for stmt in args[2:]:  # pass 1: instantiate cells and driven signals
             if isinstance(stmt, Instantiation):
+                if stmt.type not in self.tlib.cells:
+                    log.warn(f'Ignoring cell of unknown kind "{stmt.type}"')
+                    continue
                 n = Node(c, stmt.name, kind=stmt.type)
                 for p, s in stmt.pins.items():
                     if self.tlib.pin_is_output(n.kind, p):
@@ -141,35 +157,50 @@ class VerilogTransformer(Transformer):
                         c.io_nodes[positions[name]] = n
                     if sd.kind == 'input':
                         Line(c, n, Node(c, name))
-        for target, source in assignments:  # pass 1.5: process signal assignments
+        deferred_assignments = set()
-            target_sigs = []
+        ignored = 0
-            if not isinstance(target, list): target = [target]
+        while len(assignments) > 0:
-            for s in target:
+            more_assignments = []
-                if s in sig_decls:
+            for target, source in assignments:  # pass 1.5: process signal assignments
-                    target_sigs += sig_decls[s].names
+                target_sigs = []
-                else:
+                if not isinstance(target, list): target = [target]
-                    target_sigs.append(s)
+                for s in target:
-            source_sigs = []
+                    if s in sig_decls:
-            if not isinstance(source, list): source = [source]
+                        target_sigs += sig_decls[s].names
-            for s in source:
+                    else:
-                if s in sig_decls:
+                        target_sigs.append(s)
-                    source_sigs += sig_decls[s].names
+                source_sigs = []
-                else:
+                if not isinstance(source, list): source = [source]
-                    source_sigs.append(s)
+                for s in source:
-            for t, s in zip(target_sigs, source_sigs):
+                    if s in sig_decls:
-                if t in c.forks:
+                        source_sigs += sig_decls[s].names
-                    assert s not in c.forks, 'assignment between two driven signals'
+                    else:
-                    Line(c, c.forks[t], Node(c, s))
+                        source_sigs.append(s)
-                elif s in c.forks:
+                for t, s in zip(target_sigs, source_sigs):
-                    assert t not in c.forks, 'assignment between two driven signals'
+                    if t in c.forks:
-                    Line(c, c.forks[s], Node(c, t))
+                        assert s not in c.forks, f'assignment between two driven signals: source={s} target={t}'
-                elif s.startswith("1'b"):
+                        Line(c, c.forks[t], Node(c, s))
-                    cnode = Node(c, f'__const{s[3]}_{const_count}__', f'__const{s[3]}__')
+                    elif s in c.forks:
-                    const_count += 1
+                        assert t not in c.forks, 'assignment between two driven signals'
-                    Line(c, cnode, Node(c, t))
+                        Line(c, c.forks[s], Node(c, t))
+                    elif s.startswith("1'b"):
+                        cnode = Node(c, f'__const{s[3]}_{const_count}__', f'__const{s[3]}__')
+                        const_count += 1
+                        Line(c, cnode, Node(c, t))
+                    else:
+                        if (t, s) in deferred_assignments:
+                            #log.info(f'ignoring: assign {t} = {s}')
+                            ignored += 1
+                        else:
+                            more_assignments.append((t, s))
+                            deferred_assignments.add((t, s)) 
+            assignments = more_assignments
+        if ignored > 0:
+            log.warn(f'ignored {ignored} assignments')
         for stmt in args[2:]:  # pass 2: connect signals to readers
             if isinstance(stmt, Instantiation):
                 for p, s in stmt.pins.items():
+                    if stmt.name not in c.cells: continue
                     n = c.cells[stmt.name]
                     if self.tlib.pin_is_output(n.kind, p): continue
                     if s.startswith("1'b"):
@@ -210,21 +241,24 @@ GRAMMAR = r"""
     start: (module)*
     module: "module" name parameters ";" (_statement)* "endmodule"
     parameters: "(" [ _namelist ] ")"
-    _statement: input | output | inout | tri | wire | assign | instantiation
+    _statement: input | output | inout | tri | wire | reg | assign | instantiation | simple_dff
     input: "input" range? _namelist ";"
     output: "output" range? _namelist ";"
     inout: "inout" range? _namelist ";"
     tri: "tri" range? _namelist ";"
     wire: "wire" range? _namelist ";"
+    reg: "reg" range? _namelist ";"
     assign: "assign" sigsel "=" sigsel ";"
     instantiation: name name "(" [ pin ( "," pin )* ] ")" ";"
+    simple_dff: "always" "@" "(" "posedge" name ")" sigsel "<=" sigsel ";"
     pin: namedpin | sigsel
     namedpin: "." name "(" sigsel? ")"
     range: "[" /[0-9]+/ (":" /[0-9]+/)? "]"
-    sigsel: name range? | concat
+    sigsel: name range? | concat | ternaryif
     concat: "{" sigsel ( "," sigsel )*  "}"
+    ternaryif: sigsel "?" sigsel ":" sigsel
     _namelist: name ( "," name )*
-    name: ( /[a-z_][a-z0-9_]*/i | /\\[^\t \r\n]+[\t \r\n]/i | /[0-9]+'[bdh][0-9a-f]+/i )
+    name: ( /[a-z_][a-z0-9_]*/i | /\\[^\t \r\n]+[\t \r\n]/i | /[0-9]+'[bdh][x0-9a-f]+/i )
     %import common.NEWLINE
     COMMENT: /\/\*(\*(?!\/)|[^*])*\*\// | /\(\*(\*(?!\))|[^*])*\*\)/ |  "//" /(.)*/ NEWLINE
     %ignore ( /\r?\n/ | COMMENT )+
@@ -232,7 +266,7 @@ GRAMMAR = r"""
     """
 
 
-def parse(text, tlib=NANGATE, branchforks=False):
+def parse(text, tlib=KYUPY, branchforks=False) -> Circuit:
     """Parses the given ``text`` as Verilog code.
 
     :param text: A string with Verilog code.
@@ -243,10 +277,10 @@ def parse(text, tlib=NANGATE, branchforks=False):
         (see :py:func:`~kyupy.sdf.DelayFile.interconnects()`).
     :return: A :py:class:`~kyupy.circuit.Circuit` object.
     """
-    return Lark(GRAMMAR, parser="lalr", transformer=VerilogTransformer(branchforks, tlib)).parse(text)
+    return Lark(GRAMMAR, parser="lalr", transformer=VerilogTransformer(branchforks, tlib)).parse(text) # type: ignore
 
 
-def load(file, tlib=NANGATE, branchforks=False):
+def load(file, tlib=KYUPY, branchforks=False):
     """Parses the contents of ``file`` as Verilog code.
 
     :param file: A file name or a file handle. Files with `.gz`-suffix are decompressed on-the-fly.

src/kyupy/wave_sim.py

@@ -13,10 +13,11 @@ Two simulators are available: :py:class:`WaveSim` runs on the CPU, and the deriv
 """
 
 import math
+from collections import defaultdict
 
 import numpy as np
 
-from . import numba, cuda, sim, cdiv
+from . import log, numba, cuda, sim, cdiv, eng
 
 
 TMAX = np.float32(2 ** 127)
@@ -59,8 +60,8 @@ class WaveSim(sim.SimOps):
         self.delays = np.zeros((len(delays), self.c_locs_len, 2, 2), dtype=delays.dtype)
         self.delays[:, :delays.shape[1]] = delays
 
-        self.c = np.zeros((self.c_len, sims), dtype=np.float32) + TMAX
+        self.c = np.full((self.c_len, self.sims), TMAX, dtype=np.float32)
-        self.s = np.zeros((11, self.s_len, sims), dtype=np.float32)
+        self.s = np.zeros((11, self.s_len, self.sims), dtype=np.float32)
         """Information about the logic values and transitions around the sequential elements (flip-flops) and ports.
 
         The first 3 values are read by :py:func:`s_to_c`.
@@ -98,12 +99,24 @@ class WaveSim(sim.SimOps):
         self.simctl_int[0] = range(sims)  # unique seed for each sim by default, zero this to pick same delays for all sims.
         self.simctl_int[1] = 2  # random picking by default.
 
-        self.nbytes = sum([a.nbytes for a in (self.c, self.s, self.c_locs, self.c_caps, self.ops, self.simctl_int)])
+        self.simctl_float = np.zeros((1, sims), dtype=np.float32) + 1
+        """Float array for per-simulation delay configuration.
+
+        * ``simctl_float[0]`` factor to be multiplied with each delay (default=1.0).
+        """
+
+        self.e = np.zeros((self.c_locs_len, sims, 2), dtype=np.uint8)  # aux data for each line and sim
+
+        self.error_counts = np.zeros(self.s_len, dtype=np.uint32)  # number of capture errors by PPO
+        self.lsts = np.zeros(self.s_len, dtype=np.float32)  # LST by PPO
+        self.overflows = np.zeros(self.s_len, dtype=np.uint32)  # Overflows by PPO
+
+        self.nbytes = sum([a.nbytes for a in (self.c, self.s, self.e, self.c_locs, self.c_caps, self.ops, self.simctl_int, self.simctl_float)])
 
     def __repr__(self):
         dev = 'GPU' if hasattr(self.c, 'copy_to_host') else 'CPU'
         return f'{{name: "{self.circuit.name}", device: "{dev}", sims: {self.sims}, ops: {len(self.ops)}, ' + \
-               f'levels: {len(self.level_starts)}, nbytes: {self.nbytes}}}'
+               f'levels: {len(self.level_starts)}, nbytes: {eng(self.nbytes)}}}'
 
     def s_to_c(self):
         """Transfers values of sequential elements and primary inputs to the combinational portion.
@@ -116,7 +129,7 @@ class WaveSim(sim.SimOps):
         self.c[self.pippi_c_locs+1] = np.choose(cond, [TMAX, TMAX, sins[1], TMAX])
         self.c[self.pippi_c_locs+2] = TMAX
 
-    def c_prop(self, sims=None, seed=1):
+    def c_prop(self, sims=None, seed=1, delta=0):
         """Propagates all waveforms from the (pseudo) primary inputs to the (pseudo) primary outputs.
 
         :param sims: Number of parallel simulations to execute. If None, all available simulations are performed.
@@ -124,7 +137,7 @@ class WaveSim(sim.SimOps):
         """
         sims = min(sims or self.sims, self.sims)
         for op_start, op_stop in zip(self.level_starts, self.level_stops):
-            level_eval_cpu(self.ops, op_start, op_stop, self.c, self.c_locs, self.c_caps, self.abuf, 0, sims, self.delays, self.simctl_int, seed)
+            level_eval_cpu(self.ops, op_start, op_stop, self.c, self.c_locs, self.c_caps, self.e, self.abuf, 0, sims, self.delays, self.simctl_int, self.simctl_float, seed, delta)
 
     def c_to_s(self, time=TMAX, sd=0.0, seed=1):
         """Simulates a capture operation at all sequential elements and primary outputs.
@@ -152,7 +165,7 @@ class WaveSim(sim.SimOps):
         self.s[2, self.ppio_s_locs] = self.s[8, self.ppio_s_locs]
 
 
-def _wave_eval(op, cbuf, c_locs, c_caps, sim, delays, simctl_int, seed=0):
+def _wave_eval(op, cbuf, c_locs, c_caps, ebuf, sim, delays, simctl_int, simctl_float, seed, delta):
     overflows = int(0)
 
     lut = op[0]
@@ -162,6 +175,18 @@ def _wave_eval(op, cbuf, c_locs, c_caps, sim, delays, simctl_int, seed=0):
     c_idx = op[4]
     d_idx = op[5]
 
+    input_epoch = (ebuf[a_idx, sim, 1]|
+                   ebuf[b_idx, sim, 1]|
+                   ebuf[c_idx, sim, 1]|
+                   ebuf[d_idx, sim, 1])
+
+    output_epoch = ebuf[z_idx, sim, 1]
+
+    if (delta):
+        if input_epoch == 0 and output_epoch == 0: return 0, 0
+
+    #out_changed = output_epoch
+
     if len(delays) > 1:
         if simctl_int[1] == 0:
             delays = delays[seed]
@@ -175,6 +200,8 @@ def _wave_eval(op, cbuf, c_locs, c_caps, sim, delays, simctl_int, seed=0):
     else:
         delays = delays[0]
         
+    
+
     a_mem = c_locs[a_idx]
     b_mem = c_locs[b_idx]
     c_mem = c_locs[c_idx]
@@ -192,10 +219,10 @@ def _wave_eval(op, cbuf, c_locs, c_caps, sim, delays, simctl_int, seed=0):
 
     z_val = z_cur
 
-    a = cbuf[a_mem + a_cur, sim] + delays[a_idx, 0, z_val]
+    a = cbuf[a_mem + a_cur, sim] + delays[a_idx, 0, z_val] * simctl_float[0]
-    b = cbuf[b_mem + b_cur, sim] + delays[b_idx, 0, z_val]
+    b = cbuf[b_mem + b_cur, sim] + delays[b_idx, 0, z_val] * simctl_float[0]
-    c = cbuf[c_mem + c_cur, sim] + delays[c_idx, 0, z_val]
+    c = cbuf[c_mem + c_cur, sim] + delays[c_idx, 0, z_val] * simctl_float[0]
-    d = cbuf[d_mem + d_cur, sim] + delays[d_idx, 0, z_val]
+    d = cbuf[d_mem + d_cur, sim] + delays[d_idx, 0, z_val] * simctl_float[0]
 
     previous_t = TMIN
 
@@ -206,27 +233,27 @@ def _wave_eval(op, cbuf, c_locs, c_caps, sim, delays, simctl_int, seed=0):
         if a == current_t:
             a_cur += 1
             inputs ^= 1
-            thresh = delays[a_idx, a_cur & 1, z_val]
+            thresh = delays[a_idx, a_cur & 1 ^ 1, z_val] * simctl_float[0]
-            a = cbuf[a_mem + a_cur, sim] + delays[a_idx, a_cur & 1, z_val]
+            a = cbuf[a_mem + a_cur, sim] + delays[a_idx, a_cur & 1, z_val] * simctl_float[0]
-            next_t = cbuf[a_mem + a_cur, sim] + delays[a_idx, (a_cur & 1) ^ 1, z_val ^ 1]
+            next_t = cbuf[a_mem + a_cur, sim] + delays[a_idx, (a_cur & 1) ^ 1, z_val ^ 1] * simctl_float[0]
         elif b == current_t:
             b_cur += 1
             inputs ^= 2
-            thresh = delays[b_idx, b_cur & 1, z_val]
+            thresh = delays[b_idx, b_cur & 1 ^ 1, z_val] * simctl_float[0]
-            b = cbuf[b_mem + b_cur, sim] + delays[b_idx, b_cur & 1, z_val]
+            b = cbuf[b_mem + b_cur, sim] + delays[b_idx, b_cur & 1, z_val] * simctl_float[0]
-            next_t = cbuf[b_mem + b_cur, sim] + delays[b_idx, (b_cur & 1) ^ 1, z_val ^ 1]
+            next_t = cbuf[b_mem + b_cur, sim] + delays[b_idx, (b_cur & 1) ^ 1, z_val ^ 1] * simctl_float[0]
         elif c == current_t:
             c_cur += 1
             inputs ^= 4
-            thresh = delays[c_idx, c_cur & 1, z_val]
+            thresh = delays[c_idx, c_cur & 1 ^ 1, z_val] * simctl_float[0]
-            c = cbuf[c_mem + c_cur, sim] + delays[c_idx, c_cur & 1, z_val]
+            c = cbuf[c_mem + c_cur, sim] + delays[c_idx, c_cur & 1, z_val] * simctl_float[0]
-            next_t = cbuf[c_mem + c_cur, sim] + delays[c_idx, (c_cur & 1) ^ 1, z_val ^ 1]
+            next_t = cbuf[c_mem + c_cur, sim] + delays[c_idx, (c_cur & 1) ^ 1, z_val ^ 1] * simctl_float[0]
         else:
             d_cur += 1
             inputs ^= 8
-            thresh = delays[d_idx, d_cur & 1, z_val]
+            thresh = delays[d_idx, d_cur & 1 ^ 1, z_val] * simctl_float[0]
-            d = cbuf[d_mem + d_cur, sim] + delays[d_idx, d_cur & 1, z_val]
+            d = cbuf[d_mem + d_cur, sim] + delays[d_idx, d_cur & 1, z_val] * simctl_float[0]
-            next_t = cbuf[d_mem + d_cur, sim] + delays[d_idx, (d_cur & 1) ^ 1, z_val ^ 1]
+            next_t = cbuf[d_mem + d_cur, sim] + delays[d_idx, (d_cur & 1) ^ 1, z_val ^ 1] * simctl_float[0]
 
         if (z_cur & 1) != ((lut >> inputs) & 1):
             # we generate an edge in z_mem, if ...
@@ -235,32 +262,45 @@ def _wave_eval(op, cbuf, c_locs, c_caps, sim, delays, simctl_int, seed=0):
                 or (current_t - previous_t) > thresh  # -OR- the generated hazard is wider than pulse threshold.
                 ):
                 if z_cur < (z_cap - 1):  # enough space in z_mem?
+                    #if delta and (cbuf[z_mem + z_cur, sim] != current_t):
+                    #    out_changed = 1
                     cbuf[z_mem + z_cur, sim] = current_t
                     previous_t = current_t
                     z_cur += 1
                 else:
                     overflows += 1
-                    previous_t = cbuf[z_mem + z_cur - 1, sim]
                     z_cur -= 1
+                    previous_t = cbuf[z_mem + z_cur, sim]
             else:
                 z_cur -= 1
                 previous_t = cbuf[z_mem + z_cur - 1, sim] if z_cur > 0 else TMIN
 
             # output value of cell changed. update all delayed inputs.
             z_val = z_val ^ 1
-            a = cbuf[a_mem + a_cur, sim] + delays[a_idx, a_cur & 1, z_val]
+            a = cbuf[a_mem + a_cur, sim] + delays[a_idx, a_cur & 1, z_val] * simctl_float[0]
-            b = cbuf[b_mem + b_cur, sim] + delays[b_idx, b_cur & 1, z_val]
+            b = cbuf[b_mem + b_cur, sim] + delays[b_idx, b_cur & 1, z_val] * simctl_float[0]
-            c = cbuf[c_mem + c_cur, sim] + delays[c_idx, c_cur & 1, z_val]
+            c = cbuf[c_mem + c_cur, sim] + delays[c_idx, c_cur & 1, z_val] * simctl_float[0]
-            d = cbuf[d_mem + d_cur, sim] + delays[d_idx, d_cur & 1, z_val]
+            d = cbuf[d_mem + d_cur, sim] + delays[d_idx, d_cur & 1, z_val] * simctl_float[0]
 
         current_t = min(a, b, c, d)
 
+    #if delta and (cbuf[z_mem + z_cur, sim] != TMAX):
+    #    out_changed = 1
+
     # generate or propagate overflow flag
     cbuf[z_mem + z_cur, sim] = TMAX_OVL if overflows > 0 else max(a, b, c, d)
 
     nrise = max(0, (z_cur+1) // 2 - (cbuf[z_mem, sim] == TMIN))
     nfall = z_cur // 2
 
+    e = int(((cbuf[z_mem, sim] == TMIN) << 1) & 2)  # initial value
+    e |= z_val  # final value
+    e |= (nrise + nfall)<<2  # number of transitions
+
+    ebuf[z_idx, sim, 0] = e
+
+    ebuf[z_idx, sim, 1] = input_epoch #& out_changed
+
     return nrise, nfall
 
 
@@ -268,11 +308,11 @@ wave_eval_cpu = numba.njit(_wave_eval)
 
 
 @numba.njit
-def level_eval_cpu(ops, op_start, op_stop, c, c_locs, c_caps, abuf, sim_start, sim_stop, delays, simctl_int, seed):
+def level_eval_cpu(ops, op_start, op_stop, c, c_locs, c_caps, ebuf, abuf, sim_start, sim_stop, delays, simctl_int, simctl_float, seed, delta):
     for op_idx in range(op_start, op_stop):
         op = ops[op_idx]
         for sim in range(sim_start, sim_stop):
-            nrise, nfall = wave_eval_cpu(op, c, c_locs, c_caps, sim, delays, simctl_int[:, sim], seed)
+            nrise, nfall = wave_eval_cpu(op, c, c_locs, c_caps, ebuf, sim, delays, simctl_int[:, sim], simctl_float[:, sim], seed, delta)
             a_loc = op[6]
             a_wr = op[7]
             a_wf = op[8]
@@ -344,47 +384,76 @@ class WaveSimCuda(WaveSim):
         self.c_caps = cuda.to_device(self.c_caps)
         self.delays = cuda.to_device(self.delays)
         self.simctl_int = cuda.to_device(self.simctl_int)
+        self.simctl_float = cuda.to_device(self.simctl_float)
         self.abuf = cuda.to_device(self.abuf)
+        self.e = cuda.to_device(self.e)
+        self.error_counts = cuda.to_device(self.error_counts)
+        self.lsts = cuda.to_device(self.lsts)
+        self.overflows = cuda.to_device(self.overflows)
+
+        self.aux = cuda.to_device(np.zeros(8*1024, dtype=np.int32))
 
         self._block_dim = (32, 16)
 
     def __getstate__(self):
         state = self.__dict__.copy()
-        state['c'] = np.array(self.c)
+        del state['c']
         state['s'] = np.array(self.s)
         state['ops'] = np.array(self.ops)
         state['c_locs'] = np.array(self.c_locs)
         state['c_caps'] = np.array(self.c_caps)
         state['delays'] = np.array(self.delays)
         state['simctl_int'] = np.array(self.simctl_int)
+        state['simctl_float'] = np.array(self.simctl_float)
         state['abuf'] = np.array(self.abuf)
+        state['e'] = np.array(self.e)
+        state['error_counts'] = np.array(self.error_counts)
+        state['lsts'] = np.array(self.lsts)
+        state['overflows'] = np.array(self.overflows)
+        state['aux'] = np.array(self.aux)
         return state
 
     def __setstate__(self, state):
         self.__dict__.update(state)
-        self.c = cuda.to_device(self.c)
+        self.c = cuda.to_device(np.full((self.c_len, self.sims), TMAX, dtype=np.float32))
         self.s = cuda.to_device(self.s)
         self.ops = cuda.to_device(self.ops)
         self.c_locs = cuda.to_device(self.c_locs)
         self.c_caps = cuda.to_device(self.c_caps)
         self.delays = cuda.to_device(self.delays)
         self.simctl_int = cuda.to_device(self.simctl_int)
+        self.simctl_float = cuda.to_device(self.simctl_float)
         self.abuf = cuda.to_device(self.abuf)
+        self.e = cuda.to_device(self.e)
+        self.error_counts = cuda.to_device(self.error_counts)
+        self.lsts = cuda.to_device(self.lsts)
+        self.overflows = cuda.to_device(self.overflows)
+        self.aux = cuda.to_device(self.aux)
 
     def s_to_c(self):
         grid_dim = self._grid_dim(self.sims, self.s_len)
         wave_assign_gpu[grid_dim, self._block_dim](self.c, self.s, self.c_locs, self.ppi_offset)
 
-    def _grid_dim(self, x, y): return cdiv(x, self._block_dim[0]), cdiv(y, self._block_dim[1])
+    def _grid_dim(self, x, y): return cdiv(int(x), self._block_dim[0]), cdiv(int(y), self._block_dim[1])
 
-    def c_prop(self, sims=None, seed=1):
+    def c_prop(self, sims=None, seed=1, op_from=0, op_to=None, delta=0):
         sims = min(sims or self.sims, self.sims)
         for op_start, op_stop in zip(self.level_starts, self.level_stops):
+            if op_from > op_start: continue
+            if op_to is not None and op_to <= op_start: break
             grid_dim = self._grid_dim(sims, op_stop - op_start)
-            wave_eval_gpu[grid_dim, self._block_dim](self.ops, op_start, op_stop, self.c, self.c_locs, self.c_caps, self.abuf, int(0),
+            wave_eval_gpu[grid_dim, self._block_dim](self.ops, op_start, op_stop, self.c, self.c_locs, self.c_caps, self.e, self.abuf, int(0),
-                sims, self.delays, self.simctl_int, seed)
+                sims, self.delays, self.simctl_int, self.simctl_float, seed, delta)
         cuda.synchronize()
 
+    def c_prop_level(self, level, sims=None, seed=1, delta=0):
+        sims = min(sims or self.sims, self.sims)
+        op_start = self.level_starts[level]
+        op_stop = self.level_stops[level]
+        grid_dim = self._grid_dim(sims, op_stop - op_start)
+        wave_eval_gpu[grid_dim, self._block_dim](self.ops, op_start, op_stop, self.c, self.c_locs, self.c_caps, self.e, self.abuf, int(0),
+            sims, self.delays, self.simctl_int, self.simctl_float, seed, delta)
+
     def c_to_s(self, time=TMAX, sd=0.0, seed=1):
         grid_dim = self._grid_dim(self.sims, self.s_len)
         wave_capture_gpu[grid_dim, self._block_dim](self.c, self.s, self.c_locs, self.c_caps, self.ppo_offset,
@@ -394,6 +463,77 @@ class WaveSimCuda(WaveSim):
         grid_dim = self._grid_dim(self.sims, self.s_len)
         ppo_to_ppi_gpu[grid_dim, self._block_dim](self.s, self.c_locs, time, self.ppi_offset, self.ppo_offset)
 
+    def acc_error_counts(self, sims=None):
+        sims = min(sims or self.sims, self.sims)
+        grid_dim = cdiv(self.s_len, 256)
+        acc_error_counts_gpu[grid_dim, 256](self.s, sims, self.error_counts)
+
+    def reset_error_counts(self):
+        self.error_counts[:] = 0
+
+    def get_error_counts(self):
+        return np.array(self.error_counts)
+
+    def acc_overflows(self, sims=None):
+        sims = min(sims or self.sims, self.sims)
+        grid_dim = cdiv(self.s_len, 256)
+        acc_overflows_gpu[grid_dim, 256](self.s, sims, self.overflows)
+
+    def reset_overflows(self):
+        self.overflows[:] = 0
+
+    def get_overflows(self):
+        return np.array(self.overflows)
+
+    def acc_lsts(self, sims=None):
+        sims = min(sims or self.sims, self.sims)
+        grid_dim = cdiv(self.s_len, 256)
+        acc_lsts_gpu[grid_dim, 256](self.s, sims, self.lsts)
+
+    def reset_lsts(self):
+        self.lsts[:] = 0.0
+
+    def get_lsts(self):
+        return np.array(self.lsts)
+
+
+@cuda.jit()
+def memcpy_gpu (src, dst, nitems):
+    tid = cuda.grid(1)
+    stride = cuda.gridDim.x * cuda.blockDim.x
+    for i in range(tid, nitems, stride):
+        dst.flat[i] = src.flat[i]
+
+
+@cuda.jit()
+def acc_error_counts_gpu(s, sims, error_counts):
+    x = cuda.grid(1)
+    if x >= s.shape[1]: return
+    cnt = 0
+    for i in range(sims):
+        cnt += (s[6,x,i] != s[8,x,i])
+    error_counts[x] += cnt
+
+
+@cuda.jit()
+def acc_overflows_gpu(s, sims, overflows):
+    x = cuda.grid(1)
+    if x >= s.shape[1]: return
+    cnt = 0
+    for i in range(sims):
+        cnt += s[10,x,i]
+    overflows[x] += cnt
+
+
+@cuda.jit()
+def acc_lsts_gpu(s, sims, lsts):
+    x = cuda.grid(1)
+    if x >= s.shape[1]: return
+    lst = 0
+    for i in range(sims):
+        lst = max(lst, s[5,x,i])
+    lsts[x] = max(lsts[x], lst)
+
 
 @cuda.jit()
 def wave_assign_gpu(c, s, c_locs, ppi_offset):
@@ -423,7 +563,7 @@ _wave_eval_gpu = cuda.jit(_wave_eval, device=True)
 
 
 @cuda.jit()
-def wave_eval_gpu(ops, op_start, op_stop, cbuf, c_locs, c_caps, abuf, sim_start, sim_stop, delays, simctl_int, seed):
+def wave_eval_gpu(ops, op_start, op_stop, cbuf, c_locs, c_caps, ebuf, abuf, sim_start, sim_stop, delays, simctl_int, simctl_float, seed, delta):
     x, y = cuda.grid(2)
     sim = sim_start + x
     op_idx = op_start + y
@@ -435,7 +575,7 @@ def wave_eval_gpu(ops, op_start, op_stop, cbuf, c_locs, c_caps, abuf, sim_start,
     a_wr = op[7]
     a_wf = op[8]
 
-    nrise, nfall = _wave_eval_gpu(op, cbuf, c_locs, c_caps, sim, delays, simctl_int[:, sim], seed)
+    nrise, nfall = _wave_eval_gpu(op, cbuf, c_locs, c_caps, ebuf, sim, delays, simctl_int[:, sim], simctl_float[:, sim], seed, delta)
 
     # accumulate WSA into abuf
     if a_loc >= 0:

tests/Makefile

@@ -0,0 +1,2 @@
+all:
+	abc -F map_to_minimal.abc

tests/all_kyupy_simprims.minimal.v

@@ -0,0 +1,91 @@
+// Benchmark "all_kyupy_primitives" written by ABC on Thu Nov  6 13:13:21 2025
+
+module all_kyupy_primitives ( 
+    i0, i1, i2, i3,
+    \o[0] , \o[1] , \o[2] , \o[3] , \o[4] , \o[5] , \o[6] , \o[7] , \o[8] ,
+    \o[9] , \o[10] , \o[11] , \o[12] , \o[13] , \o[14] , \o[15] , \o[16] ,
+    \o[17] , \o[18] , \o[19] , \o[20] , \o[21] , \o[22] , \o[23] , \o[24] ,
+    \o[25] , \o[26] , \o[27] , \o[28] , \o[29] , \o[30] , \o[31] , \o[32]   );
+  input  i0, i1, i2, i3;
+  output \o[0] , \o[1] , \o[2] , \o[3] , \o[4] , \o[5] , \o[6] , \o[7] ,
+    \o[8] , \o[9] , \o[10] , \o[11] , \o[12] , \o[13] , \o[14] , \o[15] ,
+    \o[16] , \o[17] , \o[18] , \o[19] , \o[20] , \o[21] , \o[22] , \o[23] ,
+    \o[24] , \o[25] , \o[26] , \o[27] , \o[28] , \o[29] , \o[30] , \o[31] ,
+    \o[32] ;
+  wire new_n45, new_n48, new_n51, new_n54, new_n55, new_n56, new_n57,
+    new_n60, new_n61, new_n62, new_n63, new_n64, new_n65, new_n66, new_n67,
+    new_n68, new_n69, new_n72, new_n73, new_n74, new_n75, new_n76, new_n77,
+    new_n78, new_n79, new_n80, new_n81, new_n86, new_n87, new_n91, new_n92,
+    new_n96, new_n99, new_n102, new_n103, new_n104;
+  INV1 g00(.i0(i1), .o(\o[1] ));
+  AND2 g01(.i0(i1), .i1(i0), .o(\o[2] ));
+  AND2 g02(.i0(\o[2] ), .i1(i2), .o(\o[3] ));
+  AND2 g03(.i0(\o[3] ), .i1(i3), .o(\o[4] ));
+  INV1 g04(.i0(\o[2] ), .o(\o[5] ));
+  INV1 g05(.i0(\o[3] ), .o(\o[6] ));
+  INV1 g06(.i0(\o[4] ), .o(\o[7] ));
+  INV1 g07(.i0(i0), .o(new_n45));
+  AND2 g08(.i0(\o[1] ), .i1(new_n45), .o(\o[11] ));
+  INV1 g09(.i0(\o[11] ), .o(\o[8] ));
+  INV1 g10(.i0(i2), .o(new_n48));
+  AND2 g11(.i0(\o[11] ), .i1(new_n48), .o(\o[12] ));
+  INV1 g12(.i0(\o[12] ), .o(\o[9] ));
+  INV1 g13(.i0(i3), .o(new_n51));
+  AND2 g14(.i0(\o[12] ), .i1(new_n51), .o(\o[13] ));
+  INV1 g15(.i0(\o[13] ), .o(\o[10] ));
+  AND2 g16(.i0(\o[1] ), .i1(i0), .o(new_n54));
+  INV1 g17(.i0(new_n54), .o(new_n55));
+  AND2 g18(.i0(i1), .i1(new_n45), .o(new_n56));
+  INV1 g19(.i0(new_n56), .o(new_n57));
+  AND2 g20(.i0(new_n57), .i1(new_n55), .o(\o[17] ));
+  INV1 g21(.i0(\o[17] ), .o(\o[14] ));
+  AND2 g22(.i0(i2), .i1(i1), .o(new_n60));
+  INV1 g23(.i0(new_n60), .o(new_n61));
+  AND2 g24(.i0(new_n48), .i1(\o[1] ), .o(new_n62));
+  INV1 g25(.i0(new_n62), .o(new_n63));
+  AND2 g26(.i0(new_n63), .i1(new_n61), .o(new_n64));
+  INV1 g27(.i0(new_n64), .o(new_n65));
+  AND2 g28(.i0(new_n65), .i1(i0), .o(new_n66));
+  INV1 g29(.i0(new_n66), .o(new_n67));
+  AND2 g30(.i0(new_n64), .i1(new_n45), .o(new_n68));
+  INV1 g31(.i0(new_n68), .o(new_n69));
+  AND2 g32(.i0(new_n69), .i1(new_n67), .o(\o[18] ));
+  INV1 g33(.i0(\o[18] ), .o(\o[15] ));
+  AND2 g34(.i0(i3), .i1(new_n48), .o(new_n72));
+  INV1 g35(.i0(new_n72), .o(new_n73));
+  AND2 g36(.i0(new_n51), .i1(i2), .o(new_n74));
+  INV1 g37(.i0(new_n74), .o(new_n75));
+  AND2 g38(.i0(new_n75), .i1(new_n73), .o(new_n76));
+  INV1 g39(.i0(new_n76), .o(new_n77));
+  AND2 g40(.i0(new_n77), .i1(\o[17] ), .o(new_n78));
+  INV1 g41(.i0(new_n78), .o(new_n79));
+  AND2 g42(.i0(new_n76), .i1(\o[14] ), .o(new_n80));
+  INV1 g43(.i0(new_n80), .o(new_n81));
+  AND2 g44(.i0(new_n81), .i1(new_n79), .o(\o[19] ));
+  INV1 g45(.i0(\o[19] ), .o(\o[16] ));
+  AND2 g46(.i0(\o[5] ), .i1(new_n48), .o(\o[24] ));
+  INV1 g47(.i0(\o[24] ), .o(\o[20] ));
+  AND2 g48(.i0(i3), .i1(i2), .o(new_n86));
+  INV1 g49(.i0(new_n86), .o(new_n87));
+  AND2 g50(.i0(new_n87), .i1(\o[5] ), .o(\o[25] ));
+  INV1 g51(.i0(\o[25] ), .o(\o[21] ));
+  AND2 g52(.i0(\o[8] ), .i1(i2), .o(\o[22] ));
+  AND2 g53(.i0(new_n51), .i1(new_n48), .o(new_n91));
+  INV1 g54(.i0(new_n91), .o(new_n92));
+  AND2 g55(.i0(new_n92), .i1(\o[8] ), .o(\o[23] ));
+  INV1 g56(.i0(\o[22] ), .o(\o[26] ));
+  INV1 g57(.i0(\o[23] ), .o(\o[27] ));
+  AND2 g58(.i0(\o[5] ), .i1(new_n51), .o(new_n96));
+  AND2 g59(.i0(new_n96), .i1(new_n48), .o(\o[30] ));
+  INV1 g60(.i0(\o[30] ), .o(\o[28] ));
+  AND2 g61(.i0(\o[8] ), .i1(i3), .o(new_n99));
+  AND2 g62(.i0(new_n99), .i1(i2), .o(\o[29] ));
+  INV1 g63(.i0(\o[29] ), .o(\o[31] ));
+  AND2 g64(.i0(new_n48), .i1(i0), .o(new_n102));
+  INV1 g65(.i0(new_n102), .o(new_n103));
+  AND2 g66(.i0(new_n103), .i1(new_n61), .o(new_n104));
+  INV1 g67(.i0(new_n104), .o(\o[32] ));
+  BUF1 g68(.i0(i0), .o(\o[0] ));
+endmodule
+
+

tests/all_kyupy_simprims.v

@@ -0,0 +1,53 @@
+module all_kyupy_primitives (i0, i1, i2, i3, o);
+    input i0;
+    input i1;
+    input i2;
+    input i3;
+    output [32:0] o;
+
+    BUF1 buf1_0 (.i0(i0), .o(o[0]));
+    INV1 inv1_0 (.i0(i1), .o(o[1]));
+
+    AND2 and2_0 (.i0(i0), .i1(i1), .o(o[2]));
+    AND3 and3_0 (.i0(i0), .i1(i1), .i2(i2), .o(o[3]));
+    AND4 and4_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[4]));
+
+    NAND2 nand2_0 (.i0(i0), .i1(i1), .o(o[5]));
+    NAND3 nand3_0 (.i0(i0), .i1(i1), .i2(i2), .o(o[6]));
+    NAND4 nand4_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[7]));
+
+    OR2 or2_0 (.i0(i0), .i1(i1), .o(o[8]));
+    OR3 or3_0 (.i0(i0), .i1(i1), .i2(i2), .o(o[9]));
+    OR4 or4_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[10]));
+
+    NOR2 nor2_0 (.i0(i0), .i1(i1), .o(o[11]));
+    NOR3 nor3_0 (.i0(i0), .i1(i1), .i2(i2), .o(o[12]));
+    NOR4 nor4_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[13]));
+
+    XOR2 xor2_0 (.i0(i0), .i1(i1), .o(o[14]));
+    XOR3 xor3_0 (.i0(i0), .i1(i1), .i2(i2), .o(o[15]));
+    XOR4 xor4_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[16]));
+
+    XNOR2 xnor2_0 (.i0(i0), .i1(i1), .o(o[17]));
+    XNOR3 xnor3_0 (.i0(i0), .i1(i1), .i2(i2), .o(o[18]));
+    XNOR4 xnor4_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[19]));
+
+    AO21 ao21_0 (.i0(i0), .i1(i1), .i2(i2), .o(o[20]));
+    AO22 ao22_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[21]));
+    OA21 oa21_0 (.i0(i0), .i1(i1), .i2(i2), .o(o[22]));
+    OA22 oa22_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[23]));
+
+    AOI21 aoi21_0 (.i0(i0), .i1(i1), .i2(i2), .o(o[24]));
+    AOI22 aoi22_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[25]));
+    OAI21 oai21_0 (.i0(i0), .i1(i1), .i2(i2), .o(o[26]));
+    OAI22 oai22_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[27]));
+
+    AO211 ao211_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[28]));
+    OA211 oa211_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[29]));
+
+    AOI211 aoi211_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[30]));
+    OAI211 oai211_0 (.i0(i0), .i1(i1), .i2(i2), .i3(i3), .o(o[31]));
+
+    MUX21 mux21_0 (.i0(i0), .i1(i1), .i2(i2), .o(o[32]));
+
+endmodule

tests/conftest.py

@@ -13,8 +13,44 @@ def b15_2ig_circuit(mydir):
     from kyupy.techlib import SAED32
     return verilog.load(mydir / 'b15_2ig.v.gz', branchforks=True, tlib=SAED32)
 
+@pytest.fixture(scope='session')
+def b15_2ig_circuit_resolved(b15_2ig_circuit):
+    from kyupy.techlib import SAED32
+    cr = b15_2ig_circuit.copy()
+    cr.resolve_tlib_cells(SAED32)
+    return cr
+
+@pytest.fixture(scope='session')
+def b15_4ig_circuit(mydir):
+    from kyupy import verilog
+    from kyupy.techlib import SAED32
+    return verilog.load(mydir / 'b15_4ig.v.gz', branchforks=True, tlib=SAED32)
+
+@pytest.fixture(scope='session')
+def b15_4ig_circuit_resolved(b15_4ig_circuit):
+    from kyupy.techlib import SAED32
+    cr = b15_4ig_circuit.copy()
+    cr.resolve_tlib_cells(SAED32)
+    return cr
+
 @pytest.fixture(scope='session')
 def b15_2ig_delays(mydir, b15_2ig_circuit):
     from kyupy import sdf
     from kyupy.techlib import SAED32
     return sdf.load(mydir / 'b15_2ig.sdf.gz').iopaths(b15_2ig_circuit, tlib=SAED32)[1:2]
+
+@pytest.fixture(scope='session')
+def b15_2ig_sa_nf_test_resp(mydir, b15_2ig_circuit_resolved):
+    from kyupy import stil
+    s = stil.load(mydir / 'b15_2ig.sa_nf.stil.gz')
+    tests = s.tests(b15_2ig_circuit_resolved)[:,1:]
+    resp = s.responses(b15_2ig_circuit_resolved)[:,1:]
+    return (tests, resp)
+
+@pytest.fixture(scope='session')
+def b15_4ig_sa_rf_test_resp(mydir, b15_4ig_circuit_resolved):
+    from kyupy import stil
+    s = stil.load(mydir / 'b15_4ig.sa_rf.stil.gz')
+    tests = s.tests(b15_4ig_circuit_resolved)[:,1:]
+    resp = s.responses(b15_4ig_circuit_resolved)[:,1:]
+    return (tests, resp)

tests/gates.sdf

@@ -7,22 +7,49 @@
 (TEMPERATURE 25.00:25.00:25.00)
 (TIMESCALE 1ns)
 (CELL
-  (CELLTYPE "NAND2X1")
+  (CELLTYPE "NAND2_X1")
   (INSTANCE nandgate)
   (DELAY
     (ABSOLUTE
-    (IOPATH IN1 QN (0.099:0.103:0.103) (0.122:0.127:0.127))
+    (IOPATH A1 ZN (0.099:0.103:0.103) (0.122:0.127:0.127))
-    (IOPATH IN2 QN (0.083:0.086:0.086) (0.100:0.104:0.104))
+    (IOPATH A2 ZN (0.083:0.086:0.086) (0.100:0.104:0.104))
     )
   )
 )
 (CELL
-  (CELLTYPE "AND2X1")
+  (CELLTYPE "AND2_X1")
   (INSTANCE andgate)
   (DELAY
     (ABSOLUTE
-    (IOPATH IN1 Q (0.367:0.378:0.378) (0.351:0.377:0.377))
+    (IOPATH A1 ZN (0.367:0.378:0.378) (0.351:0.377:0.377))
-    (IOPATH IN2 Q (0.366:0.375:0.375) (0.359:0.370:0.370))
+    (IOPATH A2 ZN (0.366:0.375:0.375) (0.359:0.370:0.370))
+    )
+  )
+)
+(CELL
+  (CELLTYPE "OAI21_X1")
+  (INSTANCE oai21gate)
+  (DELAY 
+    (ABSOLUTE
+    (IOPATH B1 ZN (0.000:0.025:0.025) (0.000:0.013:0.013))
+    (IOPATH B2 ZN (0.000:0.030:0.030) (0.000:0.016:0.016))
+    ( COND (B1 == 1'b0) && (B2 == 1'b1) (IOPATH A ZN (0.000:0.018:0.018)))
+    (COND (B1 == 1'b1) && (B2 == 1'b0) (IOPATH A ZN (0.000:0.018:0.018) (0.000:0.016:0.016)))
+    (COND (B1 == 1'b1) && (B2 == 1'b1) (IOPATH A ZN (0.000:0.019:0.019) (0.000:0.014:0.014)))
+    )
+  )
+)
+(CELL
+  (CELLTYPE "MUX2_X1")
+  (INSTANCE mux2gate)
+  (DELAY 
+    (ABSOLUTE
+    (COND (B == 1'b0) && (S == 1'b0) (IOPATH A Z (0.000:0.037:0.037) (0.000:0.058:0.058)))
+    (COND (B == 1'b1) && (S == 1'b0) (IOPATH A Z (0.000:0.037:0.037) (0.000:0.058:0.058)))
+    (COND (A == 1'b0) && (S == 1'b1) (IOPATH B Z (0.000:0.035:0.035) (0.000:0.056:0.056)))
+    (COND (A == 1'b1) && (S == 1'b1) (IOPATH B Z (0.000:0.035:0.035) (0.000:0.056:0.056)))
+    (COND (A == 1'b0) && (B == 1'b1) (IOPATH S Z (0.000:0.047:0.047) (0.000:0.073:0.073)))
+    (COND (A == 1'b1) && (B == 1'b0) (IOPATH S Z (0.000:0.072:0.072) (0.000:0.064:0.064)))
     )
   )
 )

tests/gates.v

@@ -1,11 +1,15 @@
-module gates (a, b, o0, o1 );
+module gates (a, b, c, o0, o1, o2, o3 );
 input a;
 input b;
+input c;
 output o0;
 output o1;
+output o2;
+output o3;
 
-AND2X1 andgate (.IN1 ( a ) , .IN2 ( b ) , .Q ( o0 ) ) ;
+AND2_X1 andgate (.A1 ( a ) , .A2 ( b ) , .ZN ( o0 ) ) ;
-NAND2X1 nandgate (.IN1 ( a ) , .IN2 ( b ) , .QN ( o1 ) ) ;
+NAND2_X1 nandgate (.A1 ( a ) , .A2 ( b ) , .ZN ( o1 ) ) ;
-
+OAI21_X1 oai21gate (.B1(a), .B2(b), .A(c), .ZN(o2) ) ;
+MUX2_X1 mux2gate (.A(a), .B(b), .S(c), .Z(o3)) ;
 
 endmodule

tests/kyupy_simprims.genlib

@@ -0,0 +1,34 @@
+# library of all KyuPy simulation primitives defined in kyupy.sim
+GATE BUF1     1  o=i0;                         PIN * NONINV 1 999 1 0 1 0
+GATE INV1     1  o=!i0;                        PIN * INV 1 999 1 0 1 0
+GATE AND2     1  o=i0*i1;                      PIN * NONINV 1 999 1 0 1 0
+GATE AND3     1  o=i0*i1*i2;                   PIN * NONINV 1 999 1 0 1 0
+GATE AND4     1  o=i0*i1*i2*i3;                PIN * NONINV 1 999 1 0 1 0
+GATE NAND2    1  o=!(i0*i1);                   PIN * INV 1 999 1 0 1 0
+GATE NAND3    1  o=!(i0*i1*i2);                PIN * INV 1 999 1 0 1 0
+GATE NAND4    1  o=!(i0*i1*i2*i3);             PIN * INV 1 999 1 0 1 0
+GATE OR2      1  o=i0+i1;                      PIN * NONINV 1 999 1 0 1 0
+GATE OR3      1  o=i0+i1+i2;                   PIN * NONINV 1 999 1 0 1 0
+GATE OR4      1  o=i0+i1+i2+i3;                PIN * NONINV 1 999 1 0 1 0
+GATE NOR2     1  o=!(i0+i1);                   PIN * INV 1 999 1 0 1 0
+GATE NOR3     1  o=!(i0+i1+i2);                PIN * INV 1 999 1 0 1 0
+GATE NOR4     1  o=!(i0+i1+i2+i3);             PIN * INV 1 999 1 0 1 0
+GATE XOR2     1  o=i0^i1;                      PIN * UNKNOWN 1 999 1 0 1 0
+GATE XOR3     1  o=i0^i1^i2;                   PIN * UNKNOWN 1 999 1 0 1 0
+GATE XOR4     1  o=i0^i1^i2^i3;                PIN * UNKNOWN 1 999 1 0 1 0
+GATE XNOR2    1  o=!(i0^i1);                   PIN * UNKNOWN 1 999 1 0 1 0
+GATE XNOR3    1  o=!(i0^i1^i2);                PIN * UNKNOWN 1 999 1 0 1 0
+GATE XNOR4    1  o=!(i0^i1^i2^i3);             PIN * UNKNOWN 1 999 1 0 1 0
+GATE AO21     1  o=(i0*i1)+i2;                 PIN * NONINV 1 999 1 0 1 0
+GATE AO22     1  o=(i0*i1)+(i2*i3);            PIN * NONINV 1 999 1 0 1 0
+GATE OA21     1  o=(i0+i1)*i2;                 PIN * NONINV 1 999 1 0 1 0
+GATE OA22     1  o=(i0+i1)*(i2+i3);            PIN * NONINV 1 999 1 0 1 0
+GATE AOI21    1  o=!( (i0*i1)+i2 );            PIN * INV 1 999 1 0 1 0
+GATE AOI22    1  o=!( (i0*i1)+(i2*i3) );       PIN * INV 1 999 1 0 1 0
+GATE OAI21    1  o=!( (i0+i1)*i2 );            PIN * INV 1 999 1 0 1 0
+GATE OAI22    1  o=!( (i0+i1)*(i2+i3) );       PIN * INV 1 999 1 0 1 0
+GATE AO211    1  o=(i0*i1)+i2+i3;              PIN * NONINV 1 999 1 0 1 0
+GATE OA211    1  o=(i0+i1)*i2*i3;              PIN * NONINV 1 999 1 0 1 0
+GATE AOI211   1  o=!( (i0*i1)+i2+i3 );         PIN * INV 1 999 1 0 1 0
+GATE OAI211   1  o=!( (i0+i1)*i2*i3 );         PIN * INV 1 999 1 0 1 0
+GATE MUX21    1  o=(i0*!i2)+(i1*i2);           PIN * UNKNOWN 1 999 1 0 1 0

tests/map_to_minimal.abc

@@ -0,0 +1,8 @@
+# abc -f map_to_minimal.abc
+
+read kyupy_simprims.genlib
+read -m all_kyupy_simprims.v
+fraig
+read minimal.genlib
+map
+write all_kyupy_simprims.minimal.v

tests/minimal.genlib

@@ -0,0 +1,5 @@
+# A minimal library for generating logically equivalent circuits with
+# minimum number of gate types.
+GATE BUF1     1  o=i0;                         PIN * NONINV 1 999 1 0 1 0
+GATE INV1     1  o=!i0;                        PIN * INV 1 999 1 0 1 0
+GATE AND2     1  o=i0*i1;                      PIN * NONINV 1 999 1 0 1 0

tests/test_atalanta.py

@@ -0,0 +1,29 @@
+import numpy as np
+
+from kyupy import atalanta, logic
+
+
+def test_parse():
+    text = '\n'.join([
+        '* Test pattern file',
+        '1: 0011',
+        '2: 1100',
+        '1010',
+    ])
+    tf = atalanta.parse(text)
+
+    # second-to-last axis = inputs, last axis = patterns
+    assert tf.patterns.shape == (4, 3)
+    assert tf.patterns.dtype == np.uint8
+    assert set(np.unique(tf.patterns)) <= {logic.ZERO, logic.ONE}
+
+    # patterns are stored along the last axis
+    assert list(tf.patterns[:, 0]) == [logic.ZERO, logic.ZERO, logic.ONE, logic.ONE]
+    assert list(tf.patterns[:, 1]) == [logic.ONE, logic.ONE, logic.ZERO, logic.ZERO]
+    assert list(tf.patterns[:, 2]) == [logic.ONE, logic.ZERO, logic.ONE, logic.ZERO]
+
+
+def test_empty():
+    tf = atalanta.parse('* only comments\n* nothing else\n')
+    assert tf.patterns.shape == (0, 0)
+    assert tf.patterns.dtype == np.uint8

tests/test_circuit.py

@@ -1,8 +1,29 @@
 import pickle
 
-from kyupy.circuit import Circuit, Node, Line
+from kyupy.circuit import GrowingList, Circuit, Node, Line
 from kyupy import verilog, bench
-from kyupy.techlib import SAED32
+from kyupy.techlib import KYUPY, SAED32
+
+def test_growing_list():
+    gl = GrowingList()
+    assert gl.free_idx == 0
+    gl[0] = 1
+    assert gl.free_idx == 1
+    gl[2] = 1
+    assert gl.free_idx == 1
+    gl[0] = None
+    assert gl.free_idx == 0
+    gl[0] = 1
+    assert gl.free_idx == 1
+    gl[1] = 1
+    assert gl.free_idx == 3
+    gl.append(1)
+    assert gl.free_idx == 4
+    gl[2] = None
+    assert gl.free_idx == 2
+    gl[2] = 1
+    gl[1] = None
+    assert gl.free_idx == 1
 
 def test_lines():
     c = Circuit()
@@ -110,7 +131,14 @@ def test_pickle(mydir):
     cs = pickle.dumps(c)
     assert cs is not None
     c2 = pickle.loads(cs)
-    assert c == c2
+    assert len(c.nodes) == len(c2.nodes)
+
+def test_equivalency():
+    c1 = bench.parse('input(a) output(z) z=not(a)')
+    c2 = bench.parse('input(a) output(z) z=not(a)')
+    assert c1 != c2
+    assert c1.lines[0] != c2.lines[0]
+    assert c1.nodes[0] != c2.nodes[0]
 
 
 def test_substitute():
@@ -127,7 +155,7 @@ def test_substitute():
 
 def test_resolve(mydir):
     c = verilog.load(mydir / 'b15_4ig.v.gz', tlib=SAED32)
-    s_names = [n.name for n in c.s_nodes]
+    s_names = [n.name for n in c.s_nodes(SAED32)]
     c.resolve_tlib_cells(SAED32)
-    s_names_prim = [n.name for n in c.s_nodes]
+    s_names_prim = [n.name for n in c.s_nodes(KYUPY)]
     assert s_names == s_names_prim, 'resolve_tlib_cells does not preserve names or order of s_nodes'

tests/test_logic_sim.py

@@ -1,14 +1,129 @@
 import numpy as np
 
-from kyupy.logic_sim import LogicSim
+from kyupy.logic_sim import LogicSim, LogicSim2V, LogicSim4V, LogicSim6V
-from kyupy import bench, logic, sim
+from kyupy import bench, logic, sim, verilog
-from kyupy.logic import mvarray, bparray, bp_to_mv, mv_to_bp
+from kyupy.logic import mv_str, mvarray, bparray, bp_to_mv, mv_to_bp
+from kyupy.techlib import SAED90, KYUPY
+
+def test_dangling():
+    c = verilog.parse('''
+        module test(i, o);
+            input i; output o; wire dangling;
+            INVX0 dff1 (.INP(i),.ZN(o));
+            assign dangling=o;
+        endmodule
+    ''', tlib=SAED90)
+    c.resolve_tlib_cells(SAED90)
+    sim = LogicSim2V(c)
+    p = sim.allocate()
+    p[...] = 0
+    sim.simulate(p)
+    assert p[1,0] == logic.ONE
+
+def test_used_output():
+    c = bench.parse('input(i0) output(o0) output(o1) o0=BUF1(i0) o1=BUF(o0)')
+    sim = LogicSim2V(c)
+    expect = mvarray('000', '111')
+    actual = sim.simulate(expect.copy())
+    np.testing.assert_array_equal(
+        expect[c.io_locs('o0')],
+        actual[c.io_locs('o0')])
+    np.testing.assert_array_equal(
+        expect[c.io_locs('o1')],
+        actual[c.io_locs('o1')])
 
 def assert_equal_shape_and_contents(actual, desired):
     desired = np.array(desired, dtype=np.uint8)
     assert actual.shape == desired.shape
     np.testing.assert_allclose(actual, desired)
 
+def test_LogicSim2V_minimal():
+    c = bench.parse('input(i0, i1) output(o0, o1, o2) o0=AND2(i0,i1) o1=INV1(i0) o2=BUF1(i0)')
+    sim = LogicSim2V(c)
+    expect = mvarray('00010',
+                     '01010',
+                     '10001',
+                     '11101')
+    actual = sim.simulate(expect.copy())
+    np.testing.assert_array_equal(
+        expect[c.io_locs('o')],
+        actual[c.io_locs('o')])
+
+def test_LogicSim4V_minimal():
+    c = bench.parse('input(i0, i1) output(o0, o1, o2) o0=AND2(i0,i1) o1=INV1(i0) o2=BUF1(i0)')
+    sim = LogicSim4V(c, 16)
+    expect = mvarray('00010', '01010', '0-010', '0X010',
+                     '10001', '11101', '1-X01', '1XX01',
+                     '-00XX', '-1XXX', '--XXX', '-XXXX',
+                     'X00XX', 'X1XXX', 'X-XXX', 'XXXXX',)
+    actual = sim.simulate(expect.copy())
+    np.testing.assert_array_equal(
+        expect[c.io_locs('o')],
+        actual[c.io_locs('o')])
+
+def test_LogicSim6V_minimal():
+    c = bench.parse('input(i0, i1) output(o0, o1, o2) o0=AND2(i0,i1) o1=INV1(i0) o2=BUF1(i0)')
+    sim = LogicSim6V(c, 36)
+    expect = mvarray('00010', '01010', '0R010', '0F010', '0P010', '0N010',
+                     '10001', '11101', '1RR01', '1FF01', '1PP01', '1NN01',
+                     'R00FR', 'R1RFR', 'RRRFR', 'RFPFR', 'RPPFR', 'RNRFR',
+                     'F00RF', 'F1FRF', 'FRPRF', 'FFFRF', 'FPPRF', 'FNFRF',
+                     'P00NP', 'P1PNP', 'PRPNP', 'PFPNP', 'PPPNP', 'PNPNP',
+                     'N00PN', 'N1NPN', 'NRRPN', 'NFFPN', 'NPPPN', 'NNNPN')
+    actual = sim.simulate(expect.copy())
+    np.testing.assert_array_equal(
+        expect[c.io_locs('o')],
+        actual[c.io_locs('o')])
+
+def test_LogicSim2V_simprims(mydir):
+    c1 = verilog.load(mydir / 'all_kyupy_simprims.v')
+    c2 = verilog.load(mydir / 'all_kyupy_simprims.minimal.v')
+    pi_count = sum([len(n.ins)==0 for n in c1.io_nodes])
+    sim1 = LogicSim2V(c1, sims=2**pi_count)
+    sim2 = LogicSim2V(c2, sims=2**pi_count)
+    tests1 = sim1.allocate()
+    tests1[sim1.pippi_s_locs] = np.fromfunction(lambda x, y: logic.ONE * ((2**x & y) != 0), (pi_count, sim1.sims), dtype=int)  # exhaustive test
+    tests2 = tests1.copy()
+    sim1.simulate(tests1)
+    sim2.simulate(tests2)
+    np.testing.assert_array_equal(
+        tests1[c1.io_locs('o')],
+        tests2[c2.io_locs('o')])
+
+def test_LogicSim4V_simprims(mydir):
+    c1 = verilog.load(mydir / 'all_kyupy_simprims.v')
+    c2 = verilog.load(mydir / 'all_kyupy_simprims.minimal.v')
+    pi_count = sum([len(n.ins)==0 for n in c1.io_nodes])
+    sim1 = LogicSim4V(c1, sims=4**pi_count)
+    sim2 = LogicSim4V(c2, sims=4**pi_count)
+    tests1 = sim1.allocate()
+    tests1[sim1.pippi_s_locs] = np.fromfunction(lambda x, y: (y//(4**x))%4, (pi_count, sim1.sims), dtype=int)  # exhaustive test
+    tests2 = tests1.copy()
+    sim1.simulate(tests1)
+    sim2.simulate(tests2)
+    np.testing.assert_array_equal(
+        tests1[c1.io_locs('o')],
+        tests2[c2.io_locs('o')])
+
+def test_LogicSim6V_simprims(mydir):
+    c1 = verilog.load(mydir / 'all_kyupy_simprims.v')
+    c2 = verilog.load(mydir / 'all_kyupy_simprims.minimal.v')
+    pi_count = sum([len(n.ins)==0 for n in c1.io_nodes])
+    sim1 = LogicSim6V(c1, sims=6**pi_count)
+    sim2 = LogicSim6V(c2, sims=6**pi_count)
+    tests1 = sim1.allocate()
+    tests1[sim1.pippi_s_locs] = logic.mvarray('01RFPN')[np.fromfunction(lambda x, y: (y//(6**x))%6, (pi_count, sim1.sims), dtype=int)]  # exhaustive test
+    tests2 = tests1.copy()
+    sim1.simulate(tests1)
+    sim2.simulate(tests2)
+    for loc1, loc2 in zip(c1.io_locs('o'), c2.io_locs('o')):
+        n = c1.s_nodes(KYUPY)[loc1]
+        if (tests1[loc1] != tests2[loc2]).any():
+            print(f'Mismatch at output {n}')
+    np.testing.assert_array_equal(
+        tests1[c1.io_locs('o')],
+        tests2[c2.io_locs('o')])
+
 def test_2v():
     c = bench.parse(f'''
         input(i3, i2, i1, i0)
@@ -47,14 +162,12 @@ def test_2v():
         o31=OAI211(i0,i1,i2,i3)
         o32=MUX21(i0,i1,i2)
     ''')
-    s = LogicSim(c, 16, m=2)
+    s = LogicSim2V(c, 16)
-    bpa = logic.bparray([f'{i:04b}'+('-'*(s.s_len-4)) for i in range(16)])
+    s.s_assign[...] = logic.mvarray([f'{i:04b}'+('-'*(s.s_len-4)) for i in range(16)])
-    s.s[0] = bpa
     s.s_to_c()
     s.c_prop()
     s.c_to_s()
-    mva = logic.bp_to_mv(s.s[1])
+    for res, exp in zip(logic.packbits(s.s_result[4:], dtype=np.uint32), [
-    for res, exp in zip(logic.packbits(mva[4:], dtype=np.uint32), [
             sim.BUF1, sim.INV1,
             sim.AND2, sim.AND3, sim.AND4,
             sim.NAND2, sim.NAND3, sim.NAND4,
@@ -75,26 +188,82 @@ def test_2v():
 
 def test_4v():
     c = bench.parse('input(x, y) output(a, o, n) a=and(x,y) o=or(x,y) n=not(x)')
-    s = LogicSim(c, 16, m=8)  # FIXME: m=4
+    s = LogicSim4V(c, 16)
     assert s.s_len == 5
-    bpa = bparray(
+    s.s_assign[...] = mvarray(
         '00---', '01---', '0----', '0X---',
         '10---', '11---', '1----', '1X---',
         '-0---', '-1---', '-----', '-X---',
         'X0---', 'X1---', 'X----', 'XX---')
-    s.s[0] = bpa
     s.s_to_c()
     s.c_prop()
     s.c_to_s()
-    mva = bp_to_mv(s.s[1])
+    assert_equal_shape_and_contents(s.s_result, mvarray(
-    assert_equal_shape_and_contents(mva, mvarray(
         '--001', '--011', '--0X1', '--0X1',
         '--010', '--110', '--X10', '--X10',
         '--0XX', '--X1X', '--XXX', '--XXX',
         '--0XX', '--X1X', '--XXX', '--XXX'))
 
+def test_4v_fault():
+    c = bench.parse('input(x, y) output(a) a=and(x,y)')
+    s = LogicSim4V(c, 16)
+    assert s.s_len == 3
+    s.s_assign[...] = mvarray(
+        '00-', '01-', '0--', '0X-',
+        '10-', '11-', '1--', '1X-',
+        '-0-', '-1-', '---', '-X-',
+        'X0-', 'X1-', 'X--', 'XX-')
+    s.s_to_c()
+    s.c_prop()
+    s.c_to_s()
+    assert_equal_shape_and_contents(s.s_result, mvarray(
+        '--0', '--0', '--0', '--0',
+        '--0', '--1', '--X', '--X',
+        '--0', '--X', '--X', '--X',
+        '--0', '--X', '--X', '--X'))
+    fault_line = s.circuit.cells['a'].ins[0]
+    s.s_to_c()
+    s.c_prop(fault_line=fault_line.index, fault_model=1)
+    s.c_to_s()
+    assert_equal_shape_and_contents(s.s_result, mvarray(
+        '--0', '--1', '--X', '--X',
+        '--0', '--1', '--X', '--X',
+        '--0', '--1', '--X', '--X',
+        '--0', '--1', '--X', '--X'))
+    s.s_to_c()
+    s.c_prop(fault_line=fault_line.index, fault_model=0)
+    s.c_to_s()
+    assert_equal_shape_and_contents(s.s_result, mvarray(
+        '--0', '--0', '--0', '--0',
+        '--0', '--0', '--0', '--0',
+        '--0', '--0', '--0', '--0',
+        '--0', '--0', '--0', '--0'))
 
-def test_8v():
+def test_6v():
+    c = bench.parse('input(x, y) output(a, o, n, xo, no) a=AND2(x,y) o=OR2(x,y) n=INV1(x) xo=XOR2(x,y) no=NOR2(x,y)')
+    s = LogicSim6V(c, 36)
+    assert s.s_len == 7
+    mva = mvarray(
+        '0000101', '0101110', '0R0R1RF', '0F0F1FR', '0P0P1PN', '0N0N1NP',
+        '1001010', '1111000', '1RR10F0', '1FF10R0', '1PP10N0', '1NN10P0',
+        'R00RFRF', 'R1R1FF0', 'RRRRFPF', 'RFPNFNP', 'RPPRFRF', 'RNRNFFP',
+        'F00FRFR', 'F1F1RR0', 'FRPNRNP', 'FFFFRPR', 'FPPFRFR', 'FNFNRRP',
+        'P00PNPN', 'P1P1NN0', 'PRPRNRF', 'PFPFNFR', 'PPPPNPN', 'PNPNNNP',
+        'N00NPNP', 'N1N1PP0', 'NRRNPFP', 'NFFNPRP', 'NPPNPNP', 'NNNNPPP')
+    tests = np.copy(mva)
+    tests[2:] = logic.ZERO
+    s.s[0] = tests
+    s.s_to_c()
+    s.c_prop()
+    s.c_to_s()
+    resp = s.s[1].copy()
+
+    exp_resp = np.copy(mva)
+    exp_resp[:2] = logic.UNASSIGNED
+    np.testing.assert_allclose(resp, exp_resp)
+
+
+def xtest_8v():
     c = bench.parse('input(x, y) output(a, o, n, xo) a=and(x,y) o=or(x,y) n=not(x) xo=xor(x,y)')
     s = LogicSim(c, 64, m=8)
     assert s.s_len == 6
@@ -121,8 +290,8 @@ def test_8v():
     np.testing.assert_allclose(resp, exp_resp)
 
 
-def test_loop():
+def xtest_loop():
-    c = bench.parse('q=dff(d) d=not(q)')
+    c = bench.parse('q=DFF(d) d=NOT(q)')
     s = LogicSim(c, 4, m=8)
     assert s.s_len == 1
     mva = mvarray([['0'], ['1'], ['R'], ['F']])
@@ -148,8 +317,8 @@ def test_loop():
     # assert resp[3] == 'F'
 
 
-def test_latch():
+def xtest_latch():
-    c = bench.parse('input(d, t) output(q) q=latch(d, t)')
+    c = bench.parse('input(d, t) output(q) q=LATCH(d, t)')
     s = LogicSim(c, 8, m=8)
     assert s.s_len == 4
     mva = mvarray('00-0', '00-1', '01-0', '01-1', '10-0', '10-1', '11-0', '11-1')
@@ -162,14 +331,66 @@ def test_latch():
     #     assert resp[i] == exp[i]
 
 
-def test_b01(mydir):
+def sim_and_compare(c, test_resp, m=8):
-    c = bench.load(mydir / 'b01.bench')
+    tests, resp = test_resp
+    lsim = LogicSim(c, m=m, sims=tests.shape[1])
+    lsim.s[0] = logic.mv_to_bp(tests)
+    lsim.s_to_c()
+    lsim.c_prop()
+    lsim.c_to_s()
+    resp_sim = logic.bp_to_mv(lsim.s[1])[:,:tests.shape[1]]
+    idxs, pats = np.nonzero(((resp == logic.ONE) & (resp_sim != logic.ONE)) | ((resp == logic.ZERO) & (resp_sim != logic.ZERO)))
+    for i, (idx, pat) in enumerate(zip(idxs, pats)):
+        if i >= 10:
+            print(f'...')
+            break
+        print(f'mismatch pattern:{pat} ppio:{idx} exp:{logic.mv_str(resp[idx,pat])} act:{logic.mv_str(resp_sim[idx,pat])}')
+    assert len(idxs) == 0
 
-    # 8-valued
+def sim_and_compare_cls(c, test_resp, SimCls):
-    s = LogicSim(c, 8, m=8)
+    tests, resp = test_resp
-    mva = np.zeros((s.s_len, 8), dtype=np.uint8)
+    lsim = SimCls(c, sims=tests.shape[1])
-    s.s[0] = mv_to_bp(mva)
+    lsim.s_assign[...] = tests
-    s.s_to_c()
+    lsim.s_to_c()
-    s.c_prop()
+    lsim.c_prop()
-    s.c_to_s()
+    lsim.c_to_s()
-    bp_to_mv(s.s[1])
+    resp_sim = lsim.s_result
+    idxs, pats = np.nonzero(((resp == logic.ONE) & (resp_sim != logic.ONE)) | ((resp == logic.ZERO) & (resp_sim != logic.ZERO)))
+    for i, (idx, pat) in enumerate(zip(idxs, pats)):
+        if i >= 10:
+            print(f'...')
+            break
+        print(f'mismatch pattern:{pat} ppio:{idx} exp:{logic.mv_str(resp[idx,pat])} act:{logic.mv_str(resp_sim[idx,pat])}')
+    assert len(idxs) == 0
+
+
+def test_b15_2ig_sa_2v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp):
+    sim_and_compare_cls(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp, LogicSim2V)
+
+
+def test_b15_2ig_sa_4v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp):
+    sim_and_compare_cls(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp, LogicSim4V)
+
+
+def test_b15_2ig_sa_6v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp):
+    sim_and_compare_cls(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp, LogicSim6V)
+
+
+def xtest_b15_2ig_sa_8v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp):
+    sim_and_compare(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp, m=8)
+
+
+def test_b15_4ig_sa_2v(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp):
+    sim_and_compare_cls(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp, LogicSim2V)
+
+
+def test_b15_4ig_sa_4v(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp):
+    sim_and_compare_cls(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp, LogicSim4V)
+
+
+def test_b15_4ig_sa_6v(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp):
+    sim_and_compare_cls(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp, LogicSim6V)
+
+
+def xtest_b15_4ig_sa_8v(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp):
+    sim_and_compare(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp, m=8)

tests/test_sdf.py

@@ -2,7 +2,7 @@ import numpy as np
 
 from kyupy import sdf, verilog, bench
 from kyupy.wave_sim import WaveSim, TMAX, TMIN
-from kyupy.techlib import SAED32, SAED90
+from kyupy.techlib import SAED32, NANGATE45
 
 def test_parse():
     test = '''
@@ -80,9 +80,9 @@ def test_b15(mydir):
 
 
 def test_gates(mydir):
-    c = verilog.load(mydir / 'gates.v', tlib=SAED90)
+    c = verilog.load(mydir / 'gates.v', tlib=NANGATE45)
     df = sdf.load(mydir / 'gates.sdf')
-    lt = df.iopaths(c, tlib=SAED90)[1]
+    lt = df.iopaths(c, tlib=NANGATE45)[1]
     nand_a = c.cells['nandgate'].ins[0]
     nand_b = c.cells['nandgate'].ins[1]
     and_a = c.cells['andgate'].ins[0]

tests/test_stil.py

@@ -14,7 +14,7 @@ def test_b15(mydir):
     assert len(resp) > 0
 
     s2 = stil.load(mydir / 'b15_2ig.tf_nf.stil.gz')
-    tests = s2.tests_loc(b15)
+    tests = s2.tests_loc(b15, SAED32)
     resp = s2.responses(b15)
     assert len(tests) > 0
     assert len(resp) > 0

tests/test_verilog.py

@@ -1,5 +1,5 @@
 from kyupy import verilog
-from kyupy.techlib import SAED90, SAED32
+from kyupy.techlib import SAED90, SAED32, NANGATE45
 
 def test_b01(mydir):
     with open(mydir / 'b01.v', 'r') as f:
@@ -9,7 +9,7 @@ def test_b01(mydir):
 
     assert len(c.nodes) == 139
     assert len(c.lines) == 203
-    stats = c.stats
+    stats = c.stats(SAED90)
     assert stats['input'] == 6
     assert stats['output'] == 3
     assert stats['__seq__'] == 5
@@ -19,19 +19,19 @@ def test_b15(mydir):
     c = verilog.load(mydir / 'b15_4ig.v.gz', tlib=SAED32)
     assert len(c.nodes) == 12067
     assert len(c.lines) == 20731
-    stats = c.stats
+    stats = c.stats(SAED32)
     assert stats['input'] == 40
     assert stats['output'] == 71
     assert stats['__seq__'] == 417
 
 
 def test_gates(mydir):
-    c = verilog.load(mydir / 'gates.v', tlib=SAED90)
+    c = verilog.load(mydir / 'gates.v', tlib=NANGATE45)
-    assert len(c.nodes) == 10
+    assert len(c.nodes) == 18
-    assert len(c.lines) == 10
+    assert len(c.lines) == 21
-    stats = c.stats
+    stats = c.stats(NANGATE45)
-    assert stats['input'] == 2
+    assert stats['input'] == 3
-    assert stats['output'] == 2
+    assert stats['output'] == 4
     assert stats['__seq__'] == 0
 
 
@@ -39,7 +39,7 @@ def test_halton2(mydir):
     c = verilog.load(mydir / 'rng_haltonBase2.synth_yosys.v', tlib=SAED90)
     assert len(c.nodes) == 146
     assert len(c.lines) == 210
-    stats = c.stats
+    stats = c.stats(SAED90)
     assert stats['input'] == 2
     assert stats['output'] == 12
     assert stats['__seq__'] == 12

tests/test_wave_sim.py

@@ -1,26 +1,60 @@
 import numpy as np
 
 from kyupy.wave_sim import WaveSim, WaveSimCuda, wave_eval_cpu, TMIN, TMAX
-from kyupy.logic_sim import LogicSim
+from kyupy.logic_sim import LogicSim6V
 from kyupy import logic, bench, sim
 from kyupy.logic import mvarray
 
+def test_xnor2_delays():
+    op = (sim.XNOR2, 2, 0, 1, 3, 3, -1, 0, 0)
+    #op = (0b0111, 4, 0, 1)
+    c = np.full((4*16, 1), TMAX, dtype=np.float32)  # 4 waveforms of capacity 16
+    c_locs = np.zeros((4,), dtype='int')
+    c_caps = np.zeros((4,), dtype='int')
+    ebuf = np.zeros((4, 1, 2), dtype=np.int32)
+
+    for i in range(4): c_locs[i], c_caps[i] = i*16, 16  # 1:1 mapping
+    delays = np.zeros((1, 4, 2, 2))
+    delays[0, 0, 0, 0] = 0.031  # A rise -> Z rise
+    delays[0, 0, 0, 1] = 0.027  # A rise -> Z fall
+    delays[0, 0, 1, 0] = 0.033  # A fall -> Z rise
+    delays[0, 0, 1, 1] = 0.037  # A fall -> Z fall
+    delays[0, 1, 0, 0] = 0.032  # B rise -> Z rise
+    delays[0, 1, 0, 1] = 0.030  # B rise -> Z fall
+    delays[0, 1, 1, 0] = 0.038  # B fall -> Z rise
+    delays[0, 1, 1, 1] = 0.036  # B fall -> Z fall
+
+    simctl_int = np.asarray([0], dtype=np.int32)
+    simctl_float = np.asarray([1], dtype=np.float32)
+
+    def wave_assert(inputs, output):
+        for i, a in zip(inputs, c.reshape(-1,16)): a[:len(i)] = i
+        wave_eval_cpu(op, c, c_locs, c_caps, ebuf, 0, delays, simctl_int, simctl_float, 0, 0)
+        for i, v in enumerate(output): np.testing.assert_allclose(c.reshape(-1,16)[2,i], v)
+
+    wave_assert([[TMIN,TMAX],[TMIN,TMAX]], [TMIN,TMAX])      # XNOR(1,1) => 1
+    wave_assert([[TMAX,TMAX],[TMIN,TMAX]], [TMAX])           # XNOR(0,1) => 0
+    # using Afall/Zfall for pulse length, bug: was using Arise/Zfall
+    #wave_assert([[0.07, 0.10, TMAX], [0.0, TMAX]], [TMIN, 0.03, 0.101, 0.137, TMAX])
+    wave_assert([[0.07, 0.10, TMAX], [0.0, TMAX]], [TMIN, 0.03, TMAX])
+    wave_assert([[0.06, 0.10, TMAX], [0.0, TMAX]], [TMIN, 0.03, 0.091, 0.137, TMAX])
+
 def test_nand_delays():
     op = (sim.NAND4, 4, 0, 1, 2, 3, -1, 0, 0)
-    #op = (0b0111, 4, 0, 1)
+    c = np.full((5*16, 1), TMAX, dtype=np.float32)  # 5 waveforms of capacity 16
-    c = np.full((5*16, 1), TMAX)  # 5 waveforms of capacity 16
     c_locs = np.zeros((5,), dtype='int')
     c_caps = np.zeros((5,), dtype='int')
+    ebuf = np.zeros((5, 1, 2), dtype=np.int32)
 
     for i in range(5): c_locs[i], c_caps[i] = i*16, 16  # 1:1 mapping
 
     # SDF specifies IOPATH delays with respect to output polarity
     # SDF pulse rejection value is determined by IOPATH causing last transition and polarity of last transition
     delays = np.zeros((1, 5, 2, 2))
-    delays[0, 0, 0, 0] = 0.1  # A -> Z rise delay
+    delays[0, 0, 0, 0] = 0.1  # A rise -> Z rise
-    delays[0, 0, 0, 1] = 0.2  # A -> Z fall delay
+    delays[0, 0, 0, 1] = 0.2  # A rise -> Z fall
-    delays[0, 0, 1, 0] = 0.1  # A -> Z negative pulse limit (terminate in rising Z)
+    delays[0, 0, 1, 0] = 0.1  # A fall -> Z rise
-    delays[0, 0, 1, 1] = 0.2  # A -> Z positive pulse limit
+    delays[0, 0, 1, 1] = 0.2  # A fall -> Z fall
     delays[0, 1, :, 0] = 0.3  # as above for B -> Z
     delays[0, 1, :, 1] = 0.4
     delays[0, 2, :, 0] = 0.5  # as above for C -> Z
@@ -29,10 +63,11 @@ def test_nand_delays():
     delays[0, 3, :, 1] = 0.8
 
     simctl_int = np.asarray([0], dtype=np.int32)
+    simctl_float = np.asarray([1], dtype=np.float32)
 
     def wave_assert(inputs, output):
         for i, a in zip(inputs, c.reshape(-1,16)): a[:len(i)] = i
-        wave_eval_cpu(op, c, c_locs, c_caps, 0, delays, simctl_int)
+        wave_eval_cpu(op, c, c_locs, c_caps, ebuf, 0, delays, simctl_int, simctl_float, 0, 0)
         for i, v in enumerate(output): np.testing.assert_allclose(c.reshape(-1,16)[4,i], v)
 
     wave_assert([[TMAX,TMAX],[TMAX,TMAX],[TMIN,TMAX],[TMIN,TMAX]], [TMIN,TMAX]) # NAND(0,0,1,1) => 1
@@ -140,12 +175,12 @@ def compare_to_logic_sim(wsim: WaveSim):
     resp |= ((resp ^ (resp >> 1)) & 1) << 2  # transitions
     resp[wsim.pi_s_locs] = logic.UNASSIGNED
 
-    lsim = LogicSim(wsim.circuit, tests.shape[-1])
+    lsim = LogicSim6V(wsim.circuit, tests.shape[-1])
-    lsim.s[0] = logic.mv_to_bp(tests)
+    lsim.s_assign[...] = tests
     lsim.s_to_c()
     lsim.c_prop()
     lsim.c_to_s()
-    exp = logic.bp_to_mv(lsim.s[1])
+    exp = lsim.s_result[:,:tests.shape[-1]]
 
     resp[resp == logic.PPULSE] = logic.ZERO
     resp[resp == logic.NPULSE] = logic.ONE
@@ -156,13 +191,13 @@ def compare_to_logic_sim(wsim: WaveSim):
     np.testing.assert_allclose(resp, exp)
 
 
-def test_b15(b15_2ig_circuit, b15_2ig_delays):
+def test_b15(b15_2ig_circuit_resolved, b15_2ig_delays):
-    compare_to_logic_sim(WaveSim(b15_2ig_circuit, b15_2ig_delays, 8))
+    compare_to_logic_sim(WaveSim(b15_2ig_circuit_resolved, b15_2ig_delays, 8))
 
 
-def test_b15_strip_forks(b15_2ig_circuit, b15_2ig_delays):
+def test_b15_strip_forks(b15_2ig_circuit_resolved, b15_2ig_delays):
-    compare_to_logic_sim(WaveSim(b15_2ig_circuit, b15_2ig_delays, 8, strip_forks=True))
+    compare_to_logic_sim(WaveSim(b15_2ig_circuit_resolved, b15_2ig_delays, 8, strip_forks=True))
 
 
-def test_b15_cuda(b15_2ig_circuit, b15_2ig_delays):
+def test_b15_cuda(b15_2ig_circuit_resolved, b15_2ig_delays):
-    compare_to_logic_sim(WaveSimCuda(b15_2ig_circuit, b15_2ig_delays, 8, strip_forks=True))
+    compare_to_logic_sim(WaveSimCuda(b15_2ig_circuit_resolved, b15_2ig_delays, 8, strip_forks=True))