sphinx config

* Techlib: fix NanGate variants
* Circuit: more visitor and manipulator utilities
* mv_transition: produce stable values when arguments are UNASSIGNED
* Verilog: better assign processing
* STIL parsing improvements
* SDF: conditional parsing (semantics ignored for now), better interconnect annotation, more robust name matching
* sim: use ALAP topo-sort
* logic_sim: improved fault injection, faster 6V sim
* wave_sim: use correct thresholds for pulse filtering, partial delta-sim support, per-sim delay factor
* log: fix limiter

.readthedocs.yaml

@@ -4,12 +4,17 @@ build:
   os: "ubuntu-20.04"
   tools:
     python: "3.8"
+  jobs:
+    post_create_environment:
+      - python -m pip install sphinx_rtd_theme
 
 sphinx:
   fail_on_warning: true
+  configuration: docs/conf.py
 
 python:
   # Install our python package before building the docs
   install:
     - method: pip
       path: .
+

pyproject.toml

@@ -0,0 +1,32 @@
+[project]
+name = "kyupy"
+version = "0.0.5"
+authors = [
+  { name="Stefan Holst", email="mail@s-holst.de" },
+]
+description = 'High-performance processing and analysis of non-hierarchical VLSI designs'
+readme = "README.rst"
+requires_python = ">=3.8"
+dependencies = [
+    "numpy>=1.17.0",
+    "lark-parser>=0.8.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"
+

setup.py

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

@@ -161,6 +161,8 @@ class Log:
     def __setstate__(self, state):
         self.logfile = sys.stdout
         self.indent = 0
+        self._limit = -1
+        self.filtered = 0
         self.start = time.perf_counter() - state['elapsed']
 
     def write(self, s, indent=0):
@@ -181,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):

src/kyupy/circuit.py

@@ -40,6 +40,11 @@ class GrowingList(list):
             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):
@@ -91,10 +96,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: list[Line] = GrowingList()
+        self.ins: GrowingList[Line] = GrowingList()
         """A list of input connections (:class:`Line` objects).
         """
-        self.outs: list[Line] = GrowingList()
+        self.outs: GrowingList[Line] = GrowingList()
         """A list of output connections (:class:`Line` objects).
         """
 
@@ -307,7 +312,7 @@ class Circuit:
     def _locs(self, prefix, nodes):
         d_top = dict()
         for i, n in enumerate(nodes):
-            if m := re.match(fr'({prefix}.*?)((?:[\d_\[\]])*$)', n.name):
+            if m := re.match(fr'({re.escape(prefix)}.*?)((?:[\d_\[\]])*$)', n.name):
                 path = [m[1]] + [int(v) for v in re.split(r'[_\[\]]+', m[2]) if len(v) > 0]
                 d = d_top
                 for j in path[:-1]:
@@ -610,6 +615,21 @@ 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

src/kyupy/logic.py

@@ -269,6 +269,18 @@ def mv_to_bp(mva):
     return np.packbits(unpackbits(mva)[...,:3], axis=-2, bitorder='little').swapaxes(-1,-2)
 
 
+def mv_init(mva):
+    """Returns the initial binary values for mva.
+    """
+    return (mva>>1) & ((mva>>2)|mva) & 1
+
+
+def mv_final(mva):
+    """Returns the final binary value of mva.
+    """
+    return mva & ((mva>>2)|(mva>>1)) & 1
+
+
 def bparray(*a):
     """Converts (lists of) Boolean values or strings into a bit-parallel array.
 

src/kyupy/logic_sim.py

@@ -10,9 +10,10 @@ import math
 
 import numpy as np
 
-from . import numba, logic, hr_bytes, sim, eng
+from . import numba, logic, hr_bytes, sim, eng, cdiv
 from .circuit import Circuit
 
+
 class LogicSim(sim.SimOps):
     """A bit-parallel naïve combinational simulator for 2-, 4-, or 8-valued logic.
 
@@ -28,7 +29,7 @@ class LogicSim(sim.SimOps):
         self.m = m
         self.mdim = math.ceil(math.log2(m))
         self.sims = sims
-        nbytes = (sims - 1) // 8 + 1
+        nbytes = cdiv(sims, 8)
 
         self.c = np.zeros((self.c_len, self.mdim, nbytes), dtype=np.uint8)
         self.s = np.zeros((2, self.s_len, 3, nbytes), dtype=np.uint8)
@@ -51,7 +52,7 @@ class LogicSim(sim.SimOps):
         """
         self.c[self.pippi_c_locs] = self.s[0, self.pippi_s_locs, :self.mdim]
 
-    def c_prop(self, inject_cb=None):
+    def c_prop(self, sims=None, inject_cb=None, fault_line=-1, fault_mask=None, fault_model=2):
         """Propagate the input values through the combinational circuit towards the outputs.
 
         Performs all logic operations in topological order.
@@ -67,10 +68,17 @@ class LogicSim(sim.SimOps):
         t1 = self.c_locs[self.tmp2_idx]
         if self.m == 2:
             if inject_cb is None:
-                _prop_cpu(self.ops, self.c_locs, self.c)
+                if fault_mask is None:
+                    fault_mask = np.full(self.c.shape[-1], 255, dtype=np.uint8)
                 else:
-                for op, o0, i0, i1, i2, i3 in self.ops[:,:6]:
+                    if len(fault_mask) < self.c.shape[-1]:
-                    o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0, i0, i1, i2, i3)]
+                        fault_mask2 = np.full(self.c.shape[-1], 0, dtype=np.uint8)
+                        fault_mask2[:len(fault_mask)] = fault_mask
+                        fault_mask = fault_mask2
+                _prop_cpu(self.ops, self.c_locs, self.c, int(fault_line), fault_mask, int(fault_model))
+            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)]
                     if op == sim.BUF1: self.c[o0]=self.c[i0]
                     elif op == sim.INV1: self.c[o0] = ~self.c[i0]
                     elif op == sim.AND2: self.c[o0] = self.c[i0] & self.c[i1]
@@ -105,10 +113,10 @@ class LogicSim(sim.SimOps):
                     elif op == sim.OAI211:self.c[o0] = ~((self.c[i0] | self.c[i1]) & self.c[i2] & self.c[i3])
                     elif op == sim.MUX21: self.c[o0] = (self.c[i0] & ~self.c[i2]) | (self.c[i1] & self.c[i2])
                     else: print(f'unknown op {op}')
-                    inject_cb(o0, self.c[o0])
+                    inject_cb(o0l, self.c[o0])
         elif self.m == 4:
-            for op, o0, i0, i1, i2, i3 in self.ops[:,:6]:
+            for op, o0l, i0l, i1l, i2l, i3l in self.ops[:,:6]:
-                o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0, i0, i1, i2, i3)]
+                o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0l, i0l, i1l, i2l, i3l)]
                 if op == sim.BUF1: self.c[o0]=self.c[i0]
                 elif op == sim.INV1: logic.bp4v_not(self.c[o0], self.c[i0])
                 elif op == sim.AND2: logic.bp4v_and(self.c[o0], self.c[i0], self.c[i1])
@@ -181,10 +189,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(o0, self.c[o0])
+                if inject_cb is not None: inject_cb(o0l, self.c[o0])
         else:
-            for op, o0, i0, i1, i2, i3 in self.ops[:,:6]:
+            for op, o0l, i0l, i1l, i2l, i3l in self.ops[:,:6]:
-                o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0, i0, i1, i2, i3)]
+                o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0l, i0l, i1l, i2l, i3l)]
                 if op == sim.BUF1: self.c[o0]=self.c[i0]
                 elif op == sim.INV1: logic.bp8v_not(self.c[o0], self.c[i0])
                 elif op == sim.AND2: logic.bp8v_and(self.c[o0], self.c[i0], self.c[i1])
@@ -257,7 +265,7 @@ class LogicSim(sim.SimOps):
                     logic.bp8v_and(self.c[t1], self.c[i1], self.c[i2])
                     logic.bp8v_or(self.c[o0], self.c[t0], self.c[t1])
                 else: print(f'unknown op {op}')
-                if inject_cb is not None: inject_cb(o0, self.c[o0])
+                if inject_cb is not None: inject_cb(o0l, self.c[o0])
 
     def c_to_s(self):
         """Copies (captures) the results of the combinational portion to ``s[1]``.
@@ -297,9 +305,9 @@ class LogicSim(sim.SimOps):
 
 
 @numba.njit
-def _prop_cpu(ops, c_locs, c):
+def _prop_cpu(ops, c_locs, c, fault_line, fault_mask, fault_model):
-    for op, o0, i0, i1, i2, i3 in ops[:,:6]:
+    for op, o0l, i0l, i1l, i2l, i3l in ops[:,:6]:
-        o0, i0, i1, i2, i3 = [c_locs[x] for x in (o0, i0, i1, i2, i3)]
+        o0, i0, i1, i2, i3 = [c_locs[x] for x in (o0l, i0l, i1l, i2l, i3l)]
         if 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]
@@ -334,3 +342,129 @@ def _prop_cpu(ops, c_locs, c):
         elif op == sim.OAI211: c[o0] = ~((c[i0] | c[i1]) & c[i2] & c[i3])
         elif op == sim.MUX21: c[o0] = (c[i0] & ~c[i2]) | (c[i1] & c[i2])
         else: print(f'unknown op {op}')
+        if fault_line >= 0 and o0l == fault_line:
+            #n = len(fault_mask)
+            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 LogicSim6V(sim.SimOps):
+    """A bit-parallel naïve combinational simulator for 6-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.
+    """
+    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.
+        """
+
+    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):
+        c_prop_cpu(self.ops, self.c, self.c_locs, 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]
+
+
+@numba.njit
+def c_prop_cpu(ops, c, c_locs, 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.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):
+            o0[...] = o0 ^ inv_op

src/kyupy/sdf.py

@@ -103,11 +103,12 @@ 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.split('/') if '/' in n1 else (n1, None)
+            cn1, pn1 = (n1, None) if (slash := n1.rfind('/')) < 0 else (n1[:slash], n1[slash+1:])
-            cn2, pn2 = n2.split('/') if '/' in n2 else (n2, None)
+            cn2, pn2 = (n2, None) if (slash := n2.rfind('/')) < 0 else (n2[:slash], n2[slash+1:])
             cn1 = cn1.replace('\\','')
             cn2 = cn2.replace('\\','')
             c1, c2 = circuit.cells[cn1], circuit.cells[cn2]
@@ -120,19 +121,27 @@ class DelayFile:
                 log.warn(f'No line to annotate pin {pn2} of {c2}')
                 continue
             f1, f2 = c1.outs[p1].reader, c2.ins[p2].driver  # find the forks between cells.
+            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 f1 != f2:  # at least two forks, make sure f2 is a branchfork connected to f1
+                if len(f2.outs) == 1:
-                assert len(f2.outs) == 1
+                    assert f1 == f2 or f1.outs[f2.ins[0].driver_pin] == f2.ins[0]
-                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:
-                log.warn(f'No branchfork to annotate interconnect delay {c1.name}/{p1}->{c2.name}/{p2}')
+                    nonfork_annotations += 1
-                continue
+                    if nonfork_annotations < 10:
+                        log.warn(f'No branchfork between {c1.name}/{p1} and {c2.name}/{p2}, using {c2.name}/{p2}')
+                    line = c2.ins[p2]
             delays[line, :] = delvals
 
+        if nonfork_annotations > 0:
+            log.warn(f'{nonfork_annotations} interconnect annotations were moved to gate inputs due to missing forks.')
+
         return np.moveaxis(delays, -1, 0)
 
 

src/kyupy/sim.py

@@ -187,7 +187,6 @@ class SimOps:
         levels = []
 
         ppio2idx = dict((n, i) for i, n in enumerate(circuit.s_nodes))
-        pis = set([n for n in circuit.s_nodes if len(n.ins) == 0])
         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
 
@@ -311,6 +310,15 @@ class SimOps:
             if len(n.ins) > 0:
                 self.c_locs[self.ppo_offset + i], self.c_caps[self.ppo_offset + i] = self.c_locs[n.ins[0]], self.c_caps[n.ins[0]]
 
+        # line use information
+        self.line_use_start = np.full(self.c_locs_len, -1, dtype=np.int32)
+        self.line_use_stop = np.full(self.c_locs_len, len(self.levels), dtype=np.int32)
+        for i, lv in enumerate(self.levels):
+            for op in lv:
+                self.line_use_start[op[1]] = i
+                for x in [2, 3, 4, 5]:
+                    self.line_use_stop[op[x]] = i
+
         self.c_len = h.max_size
 
         d = defaultdict(int)

src/kyupy/techlib.py

@@ -11,50 +11,6 @@ from itertools import product
 from . import bench
 
 
-class TechLibOld:
-    @staticmethod
-    def pin_index(kind, pin):
-        if isinstance(pin, int):
-            return max(0, pin-1)
-        if kind[:3] in ('OAI', 'AOI'):
-            if pin[0] == 'A': return int(pin[1]) - 1
-            if pin == 'B': return int(kind[3])
-            if pin[0] == 'B': return int(pin[1]) - 1 + int(kind[3])
-        for prefix, pins, index in [('HADD', ('B0', 'SO'), 1),
-                                    ('HADD', ('A0', 'C1'), 0),
-                                    ('MUX21', ('S', 'S0'), 2),
-                                    ('MX2', ('S0',), 2),
-                                    ('TBUF', ('OE',), 1),
-                                    ('TINV', ('OE',), 1),
-                                    ('LATCH', ('D',), 0),
-                                    ('LATCH', ('QN',), 1),
-                                    ('DFF', ('D',), 0),
-                                    ('DFF', ('QN',), 1),
-                                    ('SDFF', ('D',), 0),
-                                    ('SDFF', ('QN',), 1),
-                                    ('SDFF', ('CLK',), 3),
-                                    ('SDFF', ('RSTB', 'RN'), 4),
-                                    ('SDFF', ('SETB',), 5),
-                                    ('ISOL', ('ISO',), 0),
-                                    ('ISOL', ('D',), 1)]:
-            if kind.startswith(prefix) and pin in pins: return index
-        for index, pins in enumerate([('A1', 'IN1', 'A', 'S', 'INP', 'I', 'Q', 'QN', 'Y', 'Z', 'ZN'),
-                                      ('A2', 'IN2', 'B', 'CK', 'CLK', 'CO', 'SE'),
-                                      ('A3', 'IN3', 'C', 'RN', 'RSTB', 'CI', 'SI'),
-                                      ('A4', 'IN4', 'D', 'SN', 'SETB'),
-                                      ('A5', 'IN5', 'E'),
-                                      ('A6', 'IN6', 'F')]):
-            if pin in pins: return index
-        raise ValueError(f'Unknown pin index for {kind}.{pin}')
-
-    @staticmethod
-    def pin_is_output(kind, pin):
-        if isinstance(pin, int):
-            return pin == 0
-        if 'MUX' in kind and pin == 'S': return False
-        return pin in ('Q', 'QN', 'Z', 'ZN', 'Y', 'CO', 'S', 'SO', 'C1')
-
-
 class TechLib:
     """Class for standard cell library definitions.
 
@@ -93,6 +49,14 @@ 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}'

src/kyupy/verilog.py

@@ -73,7 +73,7 @@ class VerilogTransformer(Transformer):
         elif "'" in args[0]:
             width, rest = args[0].split("'")
             width = int(width)
-            base, const = rest[0], rest[1:]
+            base, const = rest[0], rest[1:].replace('x','0')
             const = int(const, {'b': 2, 'd':10, 'h':16}[base.lower()])
             l = []
             for _ in range(width):
@@ -92,6 +92,14 @@ class VerilogTransformer(Transformer):
                 sigs.append(a)
         return sigs
 
+    def ternaryif(self, args):
+        sel = args[0]
+        ctrue = args[1]
+        cfalse = args[2]
+        print(f"got ternary if {args[0]} {args[1]}")
+        
+        return args[1]
+
     def declaration(self, kind, args):
         rnge = None
         if isinstance(args[0], range):
@@ -123,6 +131,9 @@ class VerilogTransformer(Transformer):
         assignments = []
         for stmt in args[2:]:  # pass 1: instantiate cells and driven signals
             if isinstance(stmt, Instantiation):
+                if stmt.type not in self.tlib.cells:
+                    log.warn(f'Ignoring cell of unknown kind "{stmt.type}"')
+                    continue
                 n = Node(c, stmt.name, kind=stmt.type)
                 for p, s in stmt.pins.items():
                     if self.tlib.pin_is_output(n.kind, p):
@@ -141,6 +152,10 @@ 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]
@@ -158,7 +173,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, 'assignment between two driven signals'
+                        assert s not in c.forks, f'assignment between two driven signals: source={s} target={t}'
                         Line(c, c.forks[t], Node(c, s))
                     elif s in c.forks:
                         assert t not in c.forks, 'assignment between two driven signals'
@@ -167,9 +182,20 @@ 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"):
@@ -221,10 +247,11 @@ GRAMMAR = r"""
     pin: namedpin | sigsel
     namedpin: "." name "(" sigsel? ")"
     range: "[" /[0-9]+/ (":" /[0-9]+/)? "]"
-    sigsel: name range? | concat
+    sigsel: name range? | concat | ternaryif
     concat: "{" sigsel ( "," sigsel )*  "}"
+    ternaryif: sigsel "?" sigsel ":" sigsel
     _namelist: name ( "," name )*
-    name: ( /[a-z_][a-z0-9_]*/i | /\\[^\t \r\n]+[\t \r\n]/i | /[0-9]+'[bdh][0-9a-f]+/i )
+    name: ( /[a-z_][a-z0-9_]*/i | /\\[^\t \r\n]+[\t \r\n]/i | /[0-9]+'[bdh][x0-9a-f]+/i )
     %import common.NEWLINE
     COMMENT: /\/\*(\*(?!\/)|[^*])*\*\// | /\(\*(\*(?!\))|[^*])*\*\)/ |  "//" /(.)*/ NEWLINE
     %ignore ( /\r?\n/ | COMMENT )+

src/kyupy/wave_sim.py

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

tests/conftest.py

@@ -20,8 +20,37 @@ 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/test_logic_sim.py

@@ -1,6 +1,6 @@
 import numpy as np
 
-from kyupy.logic_sim import LogicSim
+from kyupy.logic_sim import LogicSim, LogicSim6V
 from kyupy import bench, logic, sim
 from kyupy.logic import mvarray, bparray, bp_to_mv, mv_to_bp
 
@@ -94,6 +94,30 @@ def test_4v():
         '--0XX', '--X1X', '--XXX', '--XXX'))
 
 
+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)')
     s = LogicSim(c, 64, m=8)
@@ -175,14 +199,8 @@ def test_b01(mydir):
     bp_to_mv(s.s[1])
 
 
-def sim_and_compare(v_file, stil_file, m=8):
+def sim_and_compare(c, test_resp, m=8):
-    from kyupy import verilog, stil
+    tests, resp = test_resp
-    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()
@@ -197,26 +215,46 @@ def sim_and_compare(v_file, stil_file, 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):
-def test_b15_2ig_sa_2v(mydir):
+    tests, resp = test_resp
-    sim_and_compare(mydir / 'b15_2ig.v.gz', mydir / 'b15_2ig.sa_nf.stil.gz', m=2)
+    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_4v(mydir):
+def test_b15_2ig_sa_2v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp):
-    sim_and_compare(mydir / 'b15_2ig.v.gz', mydir / 'b15_2ig.sa_nf.stil.gz', m=4)
+    sim_and_compare(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp, m=2)
 
 
-def test_b15_2ig_sa_8v(mydir):
+def test_b15_2ig_sa_4v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp):
-    sim_and_compare(mydir / 'b15_2ig.v.gz', mydir / 'b15_2ig.sa_nf.stil.gz', m=8)
+    sim_and_compare(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp, m=4)
 
 
-def test_b15_4ig_sa_2v(mydir):
+def test_b15_2ig_sa_6v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp):
-    sim_and_compare(mydir / 'b15_4ig.v.gz', mydir / 'b15_4ig.sa_rf.stil.gz', m=2)
+    sim_and_compare_6v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp)
 
 
-def test_b15_4ig_sa_4v(mydir):
+def test_b15_2ig_sa_8v(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp):
-    sim_and_compare(mydir / 'b15_4ig.v.gz', mydir / 'b15_4ig.sa_rf.stil.gz', m=4)
+    sim_and_compare(b15_2ig_circuit_resolved, b15_2ig_sa_nf_test_resp, m=8)
 
 
-def test_b15_4ig_sa_8v(mydir):
+def test_b15_4ig_sa_2v(b15_4ig_circuit_resolved, b15_4ig_sa_rf_test_resp):
-    sim_and_compare(mydir / 'b15_4ig.v.gz', mydir / 'b15_4ig.sa_rf.stil.gz', m=8)
+    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)

tests/test_wave_sim.py

@@ -8,9 +8,10 @@ from kyupy.logic import mvarray
 def test_xnor2_delays():
     op = (sim.XNOR2, 2, 0, 1, 3, 3, -1, 0, 0)
     #op = (0b0111, 4, 0, 1)
-    c = np.full((4*16, 1), TMAX)  # 4 waveforms of capacity 16
+    c = np.full((4*16, 1), TMAX, dtype=np.float32)  # 4 waveforms of capacity 16
     c_locs = np.zeros((4,), dtype='int')
     c_caps = np.zeros((4,), dtype='int')
+    ebuf = np.zeros((4, 1, 2), dtype=np.int32)
 
     for i in range(4): c_locs[i], c_caps[i] = i*16, 16  # 1:1 mapping
     delays = np.zeros((1, 4, 2, 2))
@@ -27,7 +28,7 @@ def test_xnor2_delays():
 
     def wave_assert(inputs, output):
         for i, a in zip(inputs, c.reshape(-1,16)): a[:len(i)] = i
-        wave_eval_cpu(op, c, c_locs, c_caps, 0, delays, simctl_int)
+        wave_eval_cpu(op, c, c_locs, c_caps, ebuf, 0, delays, simctl_int, 0, 0)
         for i, v in enumerate(output): np.testing.assert_allclose(c.reshape(-1,16)[2,i], v)
 
     wave_assert([[TMIN,TMAX],[TMIN,TMAX]], [TMIN,TMAX])      # XNOR(1,1) => 1
@@ -40,19 +41,20 @@ 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)  # 5 waveforms of capacity 16
+    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((4, 1, 2), dtype=np.int32)
 
     for i in range(5): c_locs[i], c_caps[i] = i*16, 16  # 1:1 mapping
 
     # SDF specifies IOPATH delays with respect to output polarity
     # SDF pulse rejection value is determined by IOPATH causing last transition and polarity of last transition
     delays = np.zeros((1, 5, 2, 2))
-    delays[0, 0, 0, 0] = 0.1  # A -> Z rise delay
+    delays[0, 0, 0, 0] = 0.1  # A rise -> Z rise
-    delays[0, 0, 0, 1] = 0.2  # A -> Z fall delay
+    delays[0, 0, 0, 1] = 0.2  # A rise -> Z fall
-    delays[0, 0, 1, 0] = 0.1  # A -> Z negative pulse limit (terminate in rising Z)
+    delays[0, 0, 1, 0] = 0.1  # A fall -> Z rise
-    delays[0, 0, 1, 1] = 0.2  # A -> Z positive pulse limit
+    delays[0, 0, 1, 1] = 0.2  # A fall -> Z fall
     delays[0, 1, :, 0] = 0.3  # as above for B -> Z
     delays[0, 1, :, 1] = 0.4
     delays[0, 2, :, 0] = 0.5  # as above for C -> Z
@@ -64,7 +66,7 @@ def test_nand_delays():
 
     def wave_assert(inputs, output):
         for i, a in zip(inputs, c.reshape(-1,16)): a[:len(i)] = i
-        wave_eval_cpu(op, c, c_locs, c_caps, 0, delays, simctl_int)
+        wave_eval_cpu(op, c, c_locs, c_caps, ebuf, 0, delays, simctl_int, 0, 0)
         for i, v in enumerate(output): np.testing.assert_allclose(c.reshape(-1,16)[4,i], v)
 
     wave_assert([[TMAX,TMAX],[TMAX,TMAX],[TMIN,TMAX],[TMIN,TMAX]], [TMIN,TMAX]) # NAND(0,0,1,1) => 1