line use and diff

fix readthedocs

.gitignore

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

.python-version

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

.readthedocs.yaml

@@ -3,18 +3,13 @@ version: 2
 build:
   os: "ubuntu-20.04"
   tools:
-    python: "3.12"
-  jobs:
-    post_create_environment:
-      - python -m pip install sphinx_rtd_theme myst-parser
+    python: "3.8"
 
 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-2025 Stefan Holst
+Copyright (c) 2020-2023 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

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

@@ -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.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 myst-parser
+# pip install sphinx sphinx-rtd-theme
 #
 # 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-2025, Stefan Holst'
+copyright = '2020-2023, Stefan Holst'
 author = 'Stefan Holst'
 
 # The full version, including alpha/beta/rc tags
-release = '0.0.6'
+release = '0.0.5'
 
 
 # -- General configuration ---------------------------------------------------
@@ -35,7 +35,6 @@ release = '0.0.6'
 extensions = [
     'sphinx.ext.autodoc',
     'sphinx_rtd_theme',
-    'myst_parser',
 ]
 
 # Add any paths that contain templates here, relative to this directory.

docs/index.rst

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

docs/simulators.rst

@@ -14,15 +14,6 @@ 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

@@ -1,40 +0,0 @@
-[project]
-name = "kyupy"
-version = "0.0.6"
-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",
-    "pytest>=9.0.1",
-    "sphinx>=8.1.3",
-    "sphinx-rtd-theme>=3.0.2",
-]
-

setup.py

@@ -0,0 +1,40 @@
+from setuptools import setup, find_packages
+
+with open('README.rst', 'r') as f:
+    long_description = f.read()
+
+setup(
+    name='kyupy',
+    version='0.0.5',
+    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

@@ -183,7 +183,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):
@@ -293,6 +293,8 @@ 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/bench.py

@@ -10,11 +10,7 @@ Besides loading these benchmarks, this module is also useful for easily construc
 from lark import Lark, Transformer
 
 from .circuit import Circuit, Node, Line
-from . import readtext, batchrange
-
-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)])
+from . import readtext
 
 
 class BenchTransformer(Transformer):
@@ -25,39 +21,13 @@ class BenchTransformer(Transformer):
 
     def start(self, _): return self.c
 
-    def parameters(self, args): return [self.c.get_or_add_fork(str(name)) for name in args if name is not None]
+    def parameters(self, args): return [self.c.get_or_add_fork(str(name)) for name in args]
 
     def interface(self, args): self.c.io_nodes.extend(args[0])
 
-    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, kind, drivers = args
-        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))
+        name, cell_type, drivers = args
+        cell = Node(self.c, str(name), str(cell_type))
         Line(self.c, cell, self.c.get_or_add_fork(str(name)))
         for d in drivers: Line(self.c, d, cell)
 
@@ -74,14 +44,14 @@ GRAMMAR = r"""
     """
 
 
-def parse(text, name=None) -> Circuit:
+def parse(text, name=None):
     """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) # type: ignore
+    return Lark(GRAMMAR, parser="lalr", transformer=BenchTransformer(name)).parse(text)
 
 
 def load(file, name=None):

src/kyupy/circuit.py

@@ -14,12 +14,12 @@ from __future__ import annotations
 
 from collections import deque, defaultdict
 import re
-from typing import Union, Any
+from typing import Union
 
 import numpy as np
 
 
-class GrowingList[T](list[T]):
+class GrowingList(list):
     def __setitem__(self, index, value):
         if value is None: self.has_nones = True
         if index == len(self): return super().append(value)
@@ -28,13 +28,9 @@ class GrowingList[T](list[T]):
             self.has_nones = True
         super().__setitem__(index, value)
 
-    # Override __getitem__ to return None when reading beyond the list
-    # 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
+    def __getitem__(self, index):
+        if isinstance(index, slice): return super().__getitem__(index)
+        return super().__getitem__(index) if index < len(self) else None
 
     @property
     def free_idx(self):
@@ -44,11 +40,6 @@ class GrowingList[T](list[T]):
             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):
     def __delitem__(self, index):
@@ -100,10 +91,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[Line] = GrowingList()
+        self.ins: list[Line] = GrowingList()
         """A list of input connections (:class:`Line` objects).
         """
-        self.outs: GrowingList[Line] = GrowingList()
+        self.outs: list[Line] = GrowingList()
         """A list of output connections (:class:`Line` objects).
         """
 
@@ -138,10 +129,10 @@ class Node:
 
         This is ok, because (name, kind) is unique within a circuit.
         """
-        return self.name == other.name and self.kind == other.kind and self.circuit == other.circuit
+        return self.name == other.name and self.kind == other.kind
 
     def __hash__(self):
-        return hash((self.name, self.kind, id(self.circuit)))
+        return hash((self.name, self.kind))
 
 
 class Line:
@@ -159,7 +150,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: Circuit, driver: Node | tuple[Node, None|int], reader: Node | tuple[Node, None|int]):
+    def __init__(self, circuit: Circuit, driver: Union[Node, tuple[Node, int]], reader: Union[Node, tuple[Node, int]]):
         self.circuit = circuit
         """The :class:`Circuit` object the line is part of.
         """
@@ -172,20 +163,20 @@ class Line:
         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_idx)
-        self.driver: Node = driver[0]
+        self.driver = driver[0]
         """The :class:`Node` object that drives this line.
         """
-        self.driver_pin = driver[1] if driver[1] is not None else self.driver.outs.free_idx
+        self.driver_pin = driver[1]
         """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_idx)
-        self.reader: Node = reader[0]
+        self.reader = reader[0]
         """The :class:`Node` object that reads this line.
         """
-        self.reader_pin = reader[1] if reader[1] is not None else self.reader.ins.free_idx
+        self.reader_pin = reader[1]
         """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:
@@ -207,6 +198,8 @@ 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):
@@ -311,10 +304,10 @@ class Circuit:
         """
         return self._locs(prefix, self.s_nodes)
 
-    def _locs(self, prefix, nodes:list[Node]) -> Node|list[Any]:  # can return list[list[...]]
+    def _locs(self, prefix, nodes):
         d_top = dict()
         for i, n in enumerate(nodes):
-            if m := re.match(fr'({re.escape(prefix)}.*?)((?:[\d_\[\]])*$)', n.name):
+            if m := re.match(fr'({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]:
@@ -326,7 +319,7 @@ class Circuit:
         def sorted_values(d): 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]
-        return l #None if isinstance(l, list) and len(l) == 0 else l
+        return None if isinstance(l, list) and len(l) == 0 else l
 
     @property
     def stats(self):
@@ -541,6 +534,9 @@ 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)}}}'
 
@@ -614,21 +610,6 @@ class Circuit:
             if marks[n]:
                 yield n
 
-    def fanout(self, origin_nodes):
-        """Generator function to iterate over the fan-out cone of a given list of origin nodes.
-
-        Nodes are yielded in topological order.
-        """
-        marks = [False] * len(self.nodes)
-        for n in origin_nodes:
-            marks[n] = True
-        for n in self.topological_order():
-            if not marks[n]:
-                for line in n.ins.without_nones():
-                    marks[n] |= marks[line.driver]
-            if marks[n]:
-                yield n
-
     def fanout_free_regions(self):
         for stem in self.reversed_topological_order():
             if len(stem.outs) == 1 and 'dff' not in stem.kind.lower(): continue
@@ -650,9 +631,9 @@ class Circuit:
                 region.append(n)
             yield stem, region
 
-    def dot(self, format='svg', graph_attr={}, line_labels={}):
+    def dot(self, format='svg'):
         from graphviz import Digraph
-        dot = Digraph(format=format, graph_attr={'rankdir': 'LR', 'splines': 'true', 'size': '10', 'ranksep': '0.1'} | graph_attr)
+        dot = Digraph(format=format, graph_attr={'rankdir': 'LR', 'splines': 'true'})
 
         s_dict = dict((n, i) for i, n in enumerate(self.s_nodes))
         node_level = np.zeros(len(self.nodes), dtype=np.uint32)
@@ -665,24 +646,17 @@ 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))]) + '}|' if len(n.ins) > 1 else ''
-                    outs = '|{' + '|'.join([f'<o{i}>{i}' for i in range(len(n.outs))]) + '}' if len(n.outs) > 1 else ''
+                    ins = '|'.join([f'<i{i}>{i}' for i in range(len(n.ins))])
+                    outs = '|'.join([f'<o{i}>{i}' for i in range(len(n.outs))])
                     io = f' [{s_dict[n]}]' if n in s_dict else ''
-                    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)
+                    s.node(name=str(n.index), label = f'{{{{{ins}}}|{n.index}{io}\n{n.kind}\n{n.name}|{{{outs}}}}}', shape='record')
 
         for l in self.lines:
-            driver = f'{l.driver.index}:o{l.driver_pin}' if len(l.driver.outs)>1 else f'{l.driver.index}'
-            reader = f'{l.reader.index}:i{l.reader_pin}' if len(l.reader.ins)>1 else f'{l.reader.index}'
-            label = str(line_labels.get(l, l.index))
-            if node_level[l.driver] == node_level[l.reader]:
-                s.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)
+            driver, reader = f'{l.driver.index}:o{l.driver_pin}', f'{l.reader.index}:i{l.reader_pin}'
+            if node_level[l.driver] >= node_level[l.reader]:
+                dot.edge(driver, reader, style='dotted', label=str(l.index))
+                pass
             else:
-                dot.edge(driver, reader, label=label)
+                dot.edge(driver, reader, label=str(l.index))
 
         return dot

src/kyupy/logic.py

@@ -47,7 +47,6 @@ 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
 
@@ -115,7 +114,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.str_))
+    sa = np.choose(mva, np.array([*'0X-1PRFN'], dtype=np.unicode_))
     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)])
@@ -270,18 +269,6 @@ 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.
 
@@ -434,7 +421,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:DTypeLike=np.uint8):
+def packbits(a, dtype=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

@@ -10,9 +10,8 @@ import math
 
 import numpy as np
 
-from . import numba, logic, hr_bytes, sim, eng, cdiv, batchrange
-from .circuit import Circuit, Line
-
+from . import numba, logic, hr_bytes, sim, eng
+from .circuit import Circuit
 
 class LogicSim(sim.SimOps):
     """A bit-parallel naïve combinational simulator for 2-, 4-, or 8-valued logic.
@@ -29,7 +28,7 @@ class LogicSim(sim.SimOps):
         self.m = m
         self.mdim = math.ceil(math.log2(m))
         self.sims = sims
-        nbytes = cdiv(sims, 8)
+        nbytes = (sims - 1) // 8 + 1
 
         self.c = np.zeros((self.c_len, self.mdim, nbytes), dtype=np.uint8)
         self.s = np.zeros((2, self.s_len, 3, nbytes), dtype=np.uint8)
@@ -43,7 +42,6 @@ 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: {eng(self.c.nbytes)}}}'
@@ -53,7 +51,7 @@ class LogicSim(sim.SimOps):
         """
         self.c[self.pippi_c_locs] = self.s[0, self.pippi_s_locs, :self.mdim]
 
-    def c_prop(self, sims=None, inject_cb=None, fault_line:Line|int=-1, fault_mask=None, fault_model=2):
+    def c_prop(self, inject_cb=None):
         """Propagate the input values through the combinational circuit towards the outputs.
 
         Performs all logic operations in topological order.
@@ -62,26 +60,17 @@ 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. Specifying this callback
-            may reduce performance as it disables jit compilation.
+            resumes with the manipulated values after the callback returns.
         :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:
-                c_prop_2v_cpu(self.ops, self.c_locs, self.c, int(fault_line), fault_mask, int(fault_model))
+                _prop_cpu(self.ops, self.c_locs, self.c)
             else:
-                for op, o0l, i0l, i1l, i2l, i3l in self.ops[:,:6]:
-                    o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0l, i0l, i1l, i2l, i3l)]
+                for op, o0, i0, i1, i2, i3 in self.ops[:,:6]:
+                    o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0, i0, i1, i2, i3)]
                     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]
@@ -116,10 +105,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(o0l, self.c[o0])
+                    inject_cb(o0, self.c[o0])
         elif self.m == 4:
-            for op, o0l, i0l, i1l, i2l, i3l in self.ops[:,:6]:
-                o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0l, i0l, i1l, i2l, i3l)]
+            for op, o0, i0, i1, i2, i3 in self.ops[:,:6]:
+                o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0, i0, i1, i2, i3)]
                 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])
@@ -192,19 +181,10 @@ 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]
+                if inject_cb is not None: inject_cb(o0, self.c[o0])
         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, o0l, i0l, i1l, i2l, i3l in self.ops[:,:6]:
-                o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0l, i0l, i1l, i2l, i3l)]
+            for op, o0, i0, i1, i2, i3 in self.ops[:,:6]:
+                o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0, i0, i1, i2, i3)]
                 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])
@@ -277,7 +257,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(o0l, self.c[o0])
+                if inject_cb is not None: inject_cb(o0, self.c[o0])
 
     def c_to_s(self):
         """Copies (captures) the results of the combinational portion to ``s[1]``.
@@ -316,98 +296,10 @@ 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.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, 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 c_prop_2v_cpu(ops, c_locs, c, fault_line, fault_mask, fault_model):
-    for op, o0l, i0l, i1l, i2l, i3l in ops[:,:6]:
-        o0, i0, i1, i2, i3 = [c_locs[x] for x in (o0l, i0l, i1l, i2l, i3l)]
+def _prop_cpu(ops, c_locs, c):
+    for op, o0, i0, i1, i2, i3 in ops[:,:6]:
+        o0, i0, i1, i2, i3 = [c_locs[x] for x in (o0, i0, i1, i2, i3)]
         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]
@@ -442,518 +334,3 @@ def c_prop_2v_cpu(ops, c_locs, c, fault_line, fault_mask, fault_model):
         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
-            else:
-                c[o0] = c[o0] ^ fault_mask
-
-
-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.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])[:,: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.zeros((2, self.s_len, self.sims), 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

@@ -103,12 +103,11 @@ class DelayFile:
 
         delays = np.zeros((len(circuit.lines), 2, 2, 3))  # dataset last during construction.
 
-        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, None) if (slash := n1.rfind('/')) < 0 else (n1[:slash], n1[slash+1:])
-            cn2, pn2 = (n2, None) if (slash := n2.rfind('/')) < 0 else (n2[:slash], n2[slash+1:])
+            cn1, pn1 = n1.split('/') if '/' in n1 else (n1, None)
+            cn2, pn2 = n2.split('/') if '/' in n2 else (n2, None)
             cn1 = cn1.replace('\\','')
             cn2 = cn2.replace('\\','')
             c1, c2 = circuit.cells[cn1], circuit.cells[cn2]
@@ -121,27 +120,19 @@ 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.
-            if f1 == c2 and f2 == c1:
-                nonfork_annotations += 1
-                if nonfork_annotations < 10:
-                    log.warn(f'No fork between {c1.name}/{p1} and {c2.name}/{p2}, using {c2.name}/{p2}')
-                line = c2.ins[p2]
-            else:
             assert f1.kind == '__fork__'
             assert f2.kind == '__fork__'
-                if len(f2.outs) == 1:
-                    assert f1 == f2 or f1.outs[f2.ins[0].driver_pin] == f2.ins[0]
+            if f1 != f2:  # at least two forks, make sure f2 is a branchfork connected to f1
+                assert len(f2.outs) == 1
+                assert f1.outs[f2.ins[0].driver_pin] == f2.ins[0]
+                line = f2.ins[0]
+            elif len(f2.outs) == 1:  # f1==f2, only OK when there is no fanout.
                 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]
+                log.warn(f'No branchfork to annotate interconnect delay {c1.name}/{p1}->{c2.name}/{p2}')
+                continue
             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)
 
 

src/kyupy/sim.py

@@ -4,7 +4,6 @@ from bisect import bisect, insort_left
 
 import numpy as np
 
-from . import log
 from .circuit import Circuit
 
 BUF1 = np.uint16(0b1010_1010_1010_1010)
@@ -188,10 +187,10 @@ class SimOps:
         levels = []
 
         ppio2idx = dict((n, i) for i, n in enumerate(circuit.s_nodes))
-        root_nodes = set([n for n in circuit.s_nodes if len(n.ins) > 0] + [n for n in circuit.nodes if len(n.outs) == 0])  # start from POs, 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
+        ppos = set([n for n in circuit.s_nodes if len(n.ins) > 0])
+        readers = np.array([1 if l.reader in ppos 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 root_nodes]
+        level_lines = [n.ins[0] for n in ppos]  # start from PPOs
         # 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
@@ -200,8 +199,6 @@ class SimOps:
 
             for l in level_lines:
                 n = l.driver
-                if len(n.ins) > 4:
-                    log.warn(f'too many input pins: {n}')
                 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[0] = self.ppi_offset + ppio2idx[n]
@@ -226,7 +223,7 @@ class SimOps:
                                     sp = prims[2]
                             break
                     if sp is None:
-                        log.warn(f'ignored cell of unknown type: {n}')
+                        print('unknown cell type', kind)
                     else:
                         level_ops.append((sp, l.index, *in_idxs, *a_ctrl[l]))
 
@@ -236,7 +233,6 @@ class SimOps:
         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
@@ -272,10 +268,10 @@ class SimOps:
 
         # allocate and keep memory for PI/PPI, keep memory for PO/PPO (allocated later)
         for i, n in enumerate(circuit.s_nodes):
-            if 'dff' in n.kind.lower() or len(n.ins) == 0:  # PPI or PI
+            if len(n.outs) > 0:
                 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:  # PO
+            if len(n.ins) > 0:
                 i0_idx = stems[n.ins[0]] if stems[n.ins[0]] >= 0 else n.ins[0]
                 ref_count[i0_idx] += 1
 

src/kyupy/techlib.py

@@ -11,6 +11,50 @@ 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.
 
@@ -49,14 +93,6 @@ 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, pos, output=False):
-        """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
-        return None
-
     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}'
@@ -64,45 +100,6 @@ class TechLib:
         return self.cells[kind][1][pin][1]
 
 
-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)        ;
-""")
-"""A synthetic library of all KyuPy simulation primitives.
-"""
-
 GSC180 = TechLib(r"""
 BUFX{1,3}      input(A)    output(Y) Y=BUF1(A)    ;
 CLKBUFX{1,2,3} input(A)    output(Y) Y=BUF1(A)    ;

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 KYUPY
+from .techlib import NANGATE
 
 Instantiation = namedtuple('Instantiation', ['type', 'name', 'pins'])
 
@@ -35,7 +35,7 @@ class SignalDeclaration:
 
 
 class VerilogTransformer(Transformer):
-    def __init__(self, branchforks=False, tlib=KYUPY):
+    def __init__(self, branchforks=False, tlib=NANGATE):
         super().__init__()
         self.branchforks = branchforks
         self.tlib = tlib
@@ -61,10 +61,6 @@ class VerilogTransformer(Transformer):
                 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
@@ -77,7 +73,7 @@ class VerilogTransformer(Transformer):
         elif "'" in args[0]:
             width, rest = args[0].split("'")
             width = int(width)
-            base, const = rest[0], rest[1:].replace('x','0')
+            base, const = rest[0], rest[1:]
             const = int(const, {'b': 2, 'd':10, 'h':16}[base.lower()])
             l = []
             for _ in range(width):
@@ -96,13 +92,6 @@ 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):
@@ -114,7 +103,6 @@ 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])
@@ -135,9 +123,6 @@ 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):
@@ -156,10 +141,6 @@ class VerilogTransformer(Transformer):
                         c.io_nodes[positions[name]] = n
                     if sd.kind == 'input':
                         Line(c, n, Node(c, name))
-        deferred_assignments = set()
-        ignored = 0
-        while len(assignments) > 0:
-            more_assignments = []
         for target, source in assignments:  # pass 1.5: process signal assignments
             target_sigs = []
             if not isinstance(target, list): target = [target]
@@ -177,7 +158,7 @@ class VerilogTransformer(Transformer):
                     source_sigs.append(s)
             for t, s in zip(target_sigs, source_sigs):
                 if t in c.forks:
-                        assert s not in c.forks, f'assignment between two driven signals: source={s} target={t}'
+                    assert s not in c.forks, 'assignment between two driven signals'
                     Line(c, c.forks[t], Node(c, s))
                 elif s in c.forks:
                     assert t not in c.forks, 'assignment between two driven signals'
@@ -186,20 +167,9 @@ class VerilogTransformer(Transformer):
                     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"):
@@ -240,24 +210,21 @@ GRAMMAR = r"""
     start: (module)*
     module: "module" name parameters ";" (_statement)* "endmodule"
     parameters: "(" [ _namelist ] ")"
-    _statement: input | output | inout | tri | wire | reg | assign | instantiation | simple_dff
+    _statement: input | output | inout | tri | wire | assign | instantiation
     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 | ternaryif
+    sigsel: name range? | concat
     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][x0-9a-f]+/i )
+    name: ( /[a-z_][a-z0-9_]*/i | /\\[^\t \r\n]+[\t \r\n]/i | /[0-9]+'[bdh][0-9a-f]+/i )
     %import common.NEWLINE
     COMMENT: /\/\*(\*(?!\/)|[^*])*\*\// | /\(\*(\*(?!\))|[^*])*\*\)/ |  "//" /(.)*/ NEWLINE
     %ignore ( /\r?\n/ | COMMENT )+
@@ -265,7 +232,7 @@ GRAMMAR = r"""
     """
 
 
-def parse(text, tlib=KYUPY, branchforks=False) -> Circuit:
+def parse(text, tlib=NANGATE, branchforks=False):
     """Parses the given ``text`` as Verilog code.
 
     :param text: A string with Verilog code.
@@ -276,10 +243,10 @@ def parse(text, tlib=KYUPY, branchforks=False) -> Circuit:
         (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) # type: ignore
+    return Lark(GRAMMAR, parser="lalr", transformer=VerilogTransformer(branchforks, tlib)).parse(text)
 
 
-def load(file, tlib=KYUPY, branchforks=False):
+def load(file, tlib=NANGATE, 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

@@ -99,19 +99,26 @@ 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.simctl_float = np.zeros((1, sims), dtype=np.float32) + 1
-        """Float array for per-simulation delay configuration.
+        # flat array for line use information
+        line_use = defaultdict(list)
+        for lidx in range(len(self.circuit.lines)):
+            if self.line_use_start[lidx] < 0: continue
+            if self.line_use_stop[lidx] < 0:
+                log.warn(f'line {lidx} never read?')
+            for i in range(self.line_use_start[lidx], self.line_use_stop[lidx]):
+                line_use[i].append(lidx)
 
-        * ``simctl_float[0]`` factor to be multiplied with each delay (default=1.0).
-        """
+        self.line_use_counts = np.array([len(line_use[i]) for i in range(len(self.levels))], dtype=np.int32)
+        self.line_use_offsets = np.zeros_like(self.line_use_counts)
+        self.line_use_offsets[1:] = self.line_use_counts.cumsum()[:-1]
+        self.line_use = np.hstack([line_use[i] for i in range(len(self.levels))])
 
-        self.e = np.zeros((self.c_locs_len, sims, 2), dtype=np.uint8)  # aux data for each line and sim
+        self.h = np.zeros((self.c_locs_len, sims), dtype=np.float32)  # hashes of generated waveforms
+        self.h_base = np.zeros_like(self.h)  # base hashes to compare to
 
         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)])
+        self.nbytes = sum([a.nbytes for a in (self.c, self.s, self.h, self.c_locs, self.c_caps, self.ops, self.simctl_int)])
 
     def __repr__(self):
         dev = 'GPU' if hasattr(self.c, 'copy_to_host') else 'CPU'
@@ -129,7 +136,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, delta=0):
+    def c_prop(self, sims=None, seed=1):
         """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.
@@ -137,7 +144,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.e, self.abuf, 0, sims, self.delays, self.simctl_int, self.simctl_float, seed, delta)
+            level_eval_cpu(self.ops, op_start, op_stop, self.c, self.c_locs, self.c_caps, self.h, self.abuf, 0, sims, self.delays, self.simctl_int, seed)
 
     def c_to_s(self, time=TMAX, sd=0.0, seed=1):
         """Simulates a capture operation at all sequential elements and primary outputs.
@@ -165,7 +172,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, ebuf, sim, delays, simctl_int, simctl_float, seed, delta):
+def _wave_eval(op, cbuf, c_locs, c_caps, hbuf, sim, delays, simctl_int, seed):
     overflows = int(0)
 
     lut = op[0]
@@ -175,18 +182,6 @@ def _wave_eval(op, cbuf, c_locs, c_caps, ebuf, sim, delays, simctl_int, simctl_f
     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]
@@ -200,8 +195,6 @@ def _wave_eval(op, cbuf, c_locs, c_caps, ebuf, sim, delays, simctl_int, simctl_f
     else:
         delays = delays[0]
 
-    
-
     a_mem = c_locs[a_idx]
     b_mem = c_locs[b_idx]
     c_mem = c_locs[c_idx]
@@ -209,6 +202,8 @@ def _wave_eval(op, cbuf, c_locs, c_caps, ebuf, sim, delays, simctl_int, simctl_f
     z_mem = c_locs[z_idx]
     z_cap = c_caps[z_idx]
 
+    h = np.float32(0)
+
     a_cur = int(0)
     b_cur = int(0)
     c_cur = int(0)
@@ -219,10 +214,10 @@ def _wave_eval(op, cbuf, c_locs, c_caps, ebuf, sim, delays, simctl_int, simctl_f
 
     z_val = z_cur
 
-    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] * simctl_float[0]
-    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] * simctl_float[0]
+    a = cbuf[a_mem + a_cur, sim] + delays[a_idx, 0, z_val]
+    b = cbuf[b_mem + b_cur, sim] + delays[b_idx, 0, z_val]
+    c = cbuf[c_mem + c_cur, sim] + delays[c_idx, 0, z_val]
+    d = cbuf[d_mem + d_cur, sim] + delays[d_idx, 0, z_val]
 
     previous_t = TMIN
 
@@ -233,37 +228,36 @@ def _wave_eval(op, cbuf, c_locs, c_caps, ebuf, sim, delays, simctl_int, simctl_f
         if a == current_t:
             a_cur += 1
             inputs ^= 1
-            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] * simctl_float[0]
-            next_t = cbuf[a_mem + a_cur, sim] + delays[a_idx, (a_cur & 1) ^ 1, z_val ^ 1] * simctl_float[0]
+            thresh = delays[a_idx, a_cur & 1 ^ 1, z_val]
+            a = cbuf[a_mem + a_cur, sim] + delays[a_idx, a_cur & 1, z_val]
+            next_t = cbuf[a_mem + a_cur, sim] + delays[a_idx, (a_cur & 1) ^ 1, z_val ^ 1]
         elif b == current_t:
             b_cur += 1
             inputs ^= 2
-            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] * simctl_float[0]
-            next_t = cbuf[b_mem + b_cur, sim] + delays[b_idx, (b_cur & 1) ^ 1, z_val ^ 1] * simctl_float[0]
+            thresh = delays[b_idx, b_cur & 1 ^ 1, z_val]
+            b = cbuf[b_mem + b_cur, sim] + delays[b_idx, b_cur & 1, z_val]
+            next_t = cbuf[b_mem + b_cur, sim] + delays[b_idx, (b_cur & 1) ^ 1, z_val ^ 1]
         elif c == current_t:
             c_cur += 1
             inputs ^= 4
-            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] * simctl_float[0]
-            next_t = cbuf[c_mem + c_cur, sim] + delays[c_idx, (c_cur & 1) ^ 1, z_val ^ 1] * simctl_float[0]
+            thresh = delays[c_idx, c_cur & 1 ^ 1, z_val]
+            c = cbuf[c_mem + c_cur, sim] + delays[c_idx, c_cur & 1, z_val]
+            next_t = cbuf[c_mem + c_cur, sim] + delays[c_idx, (c_cur & 1) ^ 1, z_val ^ 1]
         else:
             d_cur += 1
             inputs ^= 8
-            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] * simctl_float[0]
-            next_t = cbuf[d_mem + d_cur, sim] + delays[d_idx, (d_cur & 1) ^ 1, z_val ^ 1] * simctl_float[0]
+            thresh = delays[d_idx, d_cur & 1 ^ 1, z_val]
+            d = cbuf[d_mem + d_cur, sim] + delays[d_idx, d_cur & 1, z_val]
+            next_t = cbuf[d_mem + d_cur, sim] + delays[d_idx, (d_cur & 1) ^ 1, z_val ^ 1]
 
         if (z_cur & 1) != ((lut >> inputs) & 1):
+            h += h*3 + max(current_t, -10)  # hash based on generated transitions before filtering
             # we generate an edge in z_mem, if ...
             if (z_cur == 0                            # it is the first edge in z_mem ...
                 or next_t < current_t                 # -OR- the next edge on SAME input is EARLIER (need current edge to filter BOTH in next iteration) ...
                 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
@@ -277,30 +271,21 @@ def _wave_eval(op, cbuf, c_locs, c_caps, ebuf, sim, delays, simctl_int, simctl_f
 
             # 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] * simctl_float[0]
-            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] * simctl_float[0]
-            d = cbuf[d_mem + d_cur, sim] + delays[d_idx, d_cur & 1, z_val] * simctl_float[0]
+            a = cbuf[a_mem + a_cur, sim] + delays[a_idx, a_cur & 1, z_val]
+            b = cbuf[b_mem + b_cur, sim] + delays[b_idx, b_cur & 1, z_val]
+            c = cbuf[c_mem + c_cur, sim] + delays[c_idx, c_cur & 1, z_val]
+            d = cbuf[d_mem + d_cur, sim] + delays[d_idx, d_cur & 1, z_val]
 
         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)
 
+    hbuf[z_idx, sim] = h
+
     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
 
 
@@ -308,11 +293,11 @@ wave_eval_cpu = numba.njit(_wave_eval)
 
 
 @numba.njit
-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):
+def level_eval_cpu(ops, op_start, op_stop, c, c_locs, c_caps, hbuf, abuf, sim_start, sim_stop, delays, simctl_int, seed):
     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, ebuf, sim, delays, simctl_int[:, sim], simctl_float[:, sim], seed, delta)
+            nrise, nfall = wave_eval_cpu(op, c, c_locs, c_caps, hbuf, sim, delays, simctl_int[:, sim], seed)
             a_loc = op[6]
             a_wr = op[7]
             a_wf = op[8]
@@ -384,14 +369,13 @@ 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.h = cuda.to_device(self.h)
+        self.h_base = cuda.to_device(self.h_base)
+        self.line_use = cuda.to_device(self.line_use)
         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.retval_int = cuda.to_device(np.array([0], dtype=np.int32))
 
         self._block_dim = (32, 16)
 
@@ -404,13 +388,12 @@ class WaveSimCuda(WaveSim):
         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['h'] = np.array(self.h)
+        state['h_base'] = np.array(self.h_base)
+        state['line_use'] = np.array(self.line_use)
         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)
+        state['retval_int'] = np.array(self.retval_int)
         return state
 
     def __setstate__(self, state):
@@ -422,13 +405,12 @@ 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.h = cuda.to_device(self.h)
+        self.h_base = cuda.to_device(self.h_base)
+        self.line_use = cuda.to_device(self.line_use)
         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)
+        self.retval_int = cuda.to_device(self.retval_int)
 
     def s_to_c(self):
         grid_dim = self._grid_dim(self.sims, self.s_len)
@@ -436,23 +418,23 @@ class WaveSimCuda(WaveSim):
 
     def _grid_dim(self, x, y): return cdiv(x, self._block_dim[0]), cdiv(y, self._block_dim[1])
 
-    def c_prop(self, sims=None, seed=1, op_from=0, op_to=None, delta=0):
+    def c_prop(self, sims=None, seed=1, op_from=0, op_to=None):
         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.e, self.abuf, int(0),
-                sims, self.delays, self.simctl_int, self.simctl_float, seed, delta)
+            wave_eval_gpu[grid_dim, self._block_dim](self.ops, op_start, op_stop, self.c, self.c_locs, self.c_caps, self.h, self.abuf, int(0),
+                sims, self.delays, self.simctl_int, seed)
         cuda.synchronize()
 
-    def c_prop_level(self, level, sims=None, seed=1, delta=0):
+    def c_prop_level(self, level, sims=None, seed=1):
         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)
+        wave_eval_gpu[grid_dim, self._block_dim](self.ops, op_start, op_stop, self.c, self.c_locs, self.c_caps, self.h, self.abuf, int(0),
+            sims, self.delays, self.simctl_int, seed)
 
     def c_to_s(self, time=TMAX, sd=0.0, seed=1):
         grid_dim = self._grid_dim(self.sims, self.s_len)
@@ -463,39 +445,24 @@ 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):
+    def set_base_hashes(self):
+        nitems = self.h_base.shape[0] * self.h_base.shape[1]
+        grid_dim = cdiv(nitems, 256)
+        memcpy_gpu[grid_dim, 256](self.h, self.h_base, nitems)
+
+    def compare_hashes_level(self, lv):
+        self.retval_int[0] = 0
+        grid_dim = self._grid_dim(self.sims, self.line_use_counts[lv])
+        diff_hash_gpu[grid_dim, self._block_dim](self.h, self.h_base, self.line_use, self.line_use_offsets[lv],
+                                                 self.line_use_counts[lv], self.retval_int)
+        return self.retval_int[0]
+
+    def calc_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):
+        calc_error_counts_gpu[grid_dim, 256](self.s, sims, self.error_counts)
         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):
@@ -506,33 +473,23 @@ def memcpy_gpu (src, dst, nitems):
 
 
 @cuda.jit()
-def acc_error_counts_gpu(s, sims, error_counts):
+def diff_hash_gpu(hbuf1, hbuf2, h_locs, h_locs_offset, h_locs_cnt, differs):
+    x, y = cuda.grid(2)
+    if x >= hbuf1.shape[1]: return
+    if y >= h_locs_cnt: return
+    h_loc = h_locs[h_locs_offset+y]
+    if hbuf1[h_loc, x] != hbuf2[h_loc, x]:
+        differs[0] = 1
+
+
+@cuda.jit()
+def calc_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)
+    error_counts[x] = cnt
 
 
 @cuda.jit()
@@ -563,7 +520,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, ebuf, abuf, sim_start, sim_stop, delays, simctl_int, simctl_float, seed, delta):
+def wave_eval_gpu(ops, op_start, op_stop, cbuf, c_locs, c_caps, hbuf, abuf, sim_start, sim_stop, delays, simctl_int, seed):
     x, y = cuda.grid(2)
     sim = sim_start + x
     op_idx = op_start + y
@@ -575,7 +532,7 @@ def wave_eval_gpu(ops, op_start, op_stop, cbuf, c_locs, c_caps, ebuf, abuf, sim_
     a_wr = op[7]
     a_wf = op[8]
 
-    nrise, nfall = _wave_eval_gpu(op, cbuf, c_locs, c_caps, ebuf, sim, delays, simctl_int[:, sim], simctl_float[:, sim], seed, delta)
+    nrise, nfall = _wave_eval_gpu(op, cbuf, c_locs, c_caps, hbuf, sim, delays, simctl_int[:, sim], seed)
 
     # accumulate WSA into abuf
     if a_loc >= 0:

tests/Makefile

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

tests/all_kyupy_simprims.minimal.v

@@ -1,91 +0,0 @@
-// 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

@@ -1,53 +0,0 @@
-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

@@ -20,37 +20,8 @@ def b15_2ig_circuit_resolved(b15_2ig_circuit):
     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/kyupy_simprims.genlib

@@ -1,34 +0,0 @@
-# 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

@@ -1,8 +0,0 @@
-# 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

@@ -1,5 +0,0 @@
-# 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_circuit.py

@@ -131,14 +131,7 @@ def test_pickle(mydir):
     cs = pickle.dumps(c)
     assert cs is not None
     c2 = pickle.loads(cs)
-    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]
+    assert c == c2
 
 
 def test_substitute():

tests/test_logic_sim.py

@@ -1,117 +1,14 @@
 import numpy as np
 
-from kyupy.logic_sim import LogicSim, LogicSim2V, LogicSim4V, LogicSim6V
-from kyupy import bench, logic, sim, verilog
+from kyupy.logic_sim import LogicSim
+from kyupy import bench, logic, sim
 from kyupy.logic import mvarray, bparray, bp_to_mv, mv_to_bp
-from kyupy.techlib import SAED90
-
-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 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[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)
@@ -178,7 +75,7 @@ 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=4)
+    s = LogicSim(c, 16, m=8)  # FIXME: m=4
     assert s.s_len == 5
     bpa = bparray(
         '00---', '01---', '0----', '0X---',
@@ -196,68 +93,6 @@ def test_4v():
         '--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 = LogicSim(c, 16, m=4)
-    assert s.s_len == 3
-    bpa = bparray(
-        '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(mva, 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, fault_model=1)
-    s.c_to_s()
-    mva = bp_to_mv(s.s[1])
-    assert_equal_shape_and_contents(mva, 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, fault_model=0)
-    s.c_to_s()
-    mva = bp_to_mv(s.s[1])
-    assert_equal_shape_and_contents(mva, mvarray(
-        '--0', '--0', '--0', '--0',
-        '--0', '--0', '--0', '--0',
-        '--0', '--0', '--0', '--0',
-        '--0', '--0', '--0', '--0'))
-
-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.ZERO
-    np.testing.assert_allclose(resp, exp_resp)
-
 
 def test_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)')
@@ -340,8 +175,14 @@ def test_b01(mydir):
     bp_to_mv(s.s[1])
 
 
-def sim_and_compare(c, test_resp, m=8):
-    tests, resp = test_resp
+def sim_and_compare(v_file, stil_file, m=8):
+    from kyupy import verilog, stil
+    from kyupy.techlib import SAED32
+    c = verilog.load(v_file, branchforks=True, tlib=SAED32)
+    c.resolve_tlib_cells(SAED32)
+    s = stil.load(stil_file)
+    tests = s.tests(c)[:,1:]
+    resp = s.responses(c)[:,1:]
     lsim = LogicSim(c, m=m, sims=tests.shape[1])
     lsim.s[0] = logic.mv_to_bp(tests)
     lsim.s_to_c()
@@ -356,46 +197,26 @@ def sim_and_compare(c, test_resp, m=8):
         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 sim_and_compare_6v(c, test_resp):
-    tests, resp = test_resp
-    lsim = LogicSim6V(c, sims=tests.shape[1])
-    lsim.s[0] = tests
-    lsim.s_to_c()
-    lsim.c_prop()
-    lsim.c_to_s()
-    resp_sim = lsim.s[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
+
+def test_b15_2ig_sa_2v(mydir):
+    sim_and_compare(mydir / 'b15_2ig.v.gz', mydir / 'b15_2ig.sa_nf.stil.gz', m=2)
 
 
-def test_b15_2ig_sa_2v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp):
-    sim_and_compare(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp, m=2)
+def test_b15_2ig_sa_4v(mydir):
+    sim_and_compare(mydir / 'b15_2ig.v.gz', mydir / 'b15_2ig.sa_nf.stil.gz', m=4)
 
 
-def test_b15_2ig_sa_4v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp):
-    sim_and_compare(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp, m=4)
+def test_b15_2ig_sa_8v(mydir):
+    sim_and_compare(mydir / 'b15_2ig.v.gz', mydir / 'b15_2ig.sa_nf.stil.gz', m=8)
 
 
-def test_b15_2ig_sa_6v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp):
-    sim_and_compare_6v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp)
+def test_b15_4ig_sa_2v(mydir):
+    sim_and_compare(mydir / 'b15_4ig.v.gz', mydir / 'b15_4ig.sa_rf.stil.gz', m=2)
 
 
-def test_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_4v(mydir):
+    sim_and_compare(mydir / 'b15_4ig.v.gz', mydir / 'b15_4ig.sa_rf.stil.gz', m=4)
 
 
-def test_b15_4ig_sa_2v(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp):
-    sim_and_compare(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp, m=2)
-
-
-def test_b15_4ig_sa_4v(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp):
-    sim_and_compare(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp, m=4)
-
-
-def test_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)
+def test_b15_4ig_sa_8v(mydir):
+    sim_and_compare(mydir / 'b15_4ig.v.gz', mydir / 'b15_4ig.sa_rf.stil.gz', m=8)

tests/test_wave_sim.py

@@ -11,7 +11,7 @@ def test_xnor2_delays():
     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)
+    h = np.zeros((4, 1), dtype=np.float32)
 
     for i in range(4): c_locs[i], c_caps[i] = i*16, 16  # 1:1 mapping
     delays = np.zeros((1, 4, 2, 2))
@@ -25,11 +25,10 @@ def test_xnor2_delays():
     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)
+        wave_eval_cpu(op, c, c_locs, c_caps, h, 0, delays, simctl_int, 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
@@ -41,20 +40,21 @@ def test_xnor2_delays():
 
 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_locs = np.zeros((5,), dtype='int')
     c_caps = np.zeros((5,), dtype='int')
-    ebuf = np.zeros((5, 1, 2), dtype=np.int32)
+    h = np.zeros((5, 1), dtype=np.float32)
 
     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 rise -> Z rise
-    delays[0, 0, 0, 1] = 0.2  # A rise -> Z fall
-    delays[0, 0, 1, 0] = 0.1  # A fall -> Z rise
-    delays[0, 0, 1, 1] = 0.2  # A fall -> Z fall
+    delays[0, 0, 0, 0] = 0.1  # A -> Z rise delay
+    delays[0, 0, 0, 1] = 0.2  # A -> Z fall delay
+    delays[0, 0, 1, 0] = 0.1  # A -> Z negative pulse limit (terminate in rising Z)
+    delays[0, 0, 1, 1] = 0.2  # A -> Z positive pulse limit
     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
@@ -63,11 +63,10 @@ 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, ebuf, 0, delays, simctl_int, simctl_float, 0, 0)
+        wave_eval_cpu(op, c, c_locs, c_caps, h, 0, delays, simctl_int, 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