Browse Source
- MVArray for multi-valued logic - BPArray for bit-parallel storage layout - Started documenting with Sphinx - Migrated simulators to new BPArraymain v0.0.2
32 changed files with 1847 additions and 1476 deletions
@ -1,56 +0,0 @@
@@ -1,56 +0,0 @@
|
||||
KyuPy - Processing VLSI Circuits With Ease |
||||
========================================== |
||||
|
||||
KyuPy is a python package for high-performance processing and analysis of |
||||
non-hierarchical VLSI designs. Its purpose is to provide a rapid prototyping |
||||
platform to aid and accelerate research in the fields of VLSI test, diagnosis |
||||
and reliability. KyuPy is freely available under the MIT license. |
||||
|
||||
|
||||
Main Features |
||||
------------- |
||||
|
||||
* Partial [lark](https://github.com/lark-parser/lark) parsers for common files used with synthesized designs: |
||||
bench, gate-level verilog, standard delay format (SDF), standard test interface language (STIL) |
||||
* 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 requires Python 3.6 or newer. |
||||
Install the latest release by running: |
||||
```commandline |
||||
pip3 install --user kyupy |
||||
``` |
||||
For best performance, ensure you have [numba](https://pypi.org/project/numba) installed: |
||||
```commandline |
||||
pip3 install --user numba |
||||
``` |
||||
GPU/CUDA support may [require some additional setup](https://numba.pydata.org/numba-doc/latest/cuda/index.html). |
||||
If CUDA or numba is not available, KyuPy will automatically fall back to slow, pure python execution. |
||||
|
||||
The Jupyter Notebook [UsageExamples.ipynb](https://github.com/s-holst/kyupy/blob/main/UsageExamples.ipynb) on GitHub |
||||
contains some useful examples to get familiar with the API. |
||||
|
||||
|
||||
Development |
||||
----------- |
||||
|
||||
To contribute to KyuPy or simply explore the source code, clone the KyuPy [repository](https://github.com/s-holst/kyupy) on GitHub. |
||||
Within your local checkout, run: |
||||
```commandline |
||||
pip3 install --user -e . |
||||
``` |
||||
to make the kyupy package available in your python environment. |
||||
The source code comes with tests that can be run with: |
||||
``` |
||||
pytest |
||||
``` |
||||
|
||||
KyuPy depends on the following packages: |
||||
* [lark-parser](https://pypi.org/project/lark-parser) |
||||
* [numpy](https://pypi.org/project/numpy) |
||||
* [numba](https://pypi.org/project/numba) (optional, required only for GPU/CUDA support) |
||||
@ -0,0 +1,28 @@
@@ -0,0 +1,28 @@
|
||||
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). |
||||
* 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.6 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.pydata.org/numba-doc/latest/cuda/index.html>`_. |
||||
If numba is not available, KyuPy will automatically fall back to slow, pure Python execution. |
||||
|
||||
The Jupyter Notebook `Demo.ipynb <https://github.com/s-holst/kyupy/blob/main/Demo.ipynb>`_ contains some useful examples to get familiar with the API. |
||||
|
||||
To work with the latest pre-release source code, clone the `KyuPy GitHub repository <https://github.com/s-holst/kyupy>`_. |
||||
Run ``pip3 install --user -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``. |
||||
@ -0,0 +1,20 @@
@@ -0,0 +1,20 @@
|
||||
# Minimal makefile for Sphinx documentation
|
||||
#
|
||||
|
||||
# You can set these variables from the command line, and also
|
||||
# from the environment for the first two.
|
||||
SPHINXOPTS ?= |
||||
SPHINXBUILD ?= sphinx-build |
||||
SOURCEDIR = . |
||||
BUILDDIR = _build |
||||
|
||||
# Put it first so that "make" without argument is like "make help".
|
||||
help: |
||||
@$(SPHINXBUILD) -M help "$(SOURCEDIR)" "$(BUILDDIR)" $(SPHINXOPTS) $(O) |
||||
|
||||
.PHONY: help Makefile |
||||
|
||||
# Catch-all target: route all unknown targets to Sphinx using the new
|
||||
# "make mode" option. $(O) is meant as a shortcut for $(SPHINXOPTS).
|
||||
%: Makefile |
||||
@$(SPHINXBUILD) -M $@ "$(SOURCEDIR)" "$(BUILDDIR)" $(SPHINXOPTS) $(O) |
||||
@ -0,0 +1,64 @@
@@ -0,0 +1,64 @@
|
||||
# Configuration file for the Sphinx documentation builder. |
||||
# |
||||
# This file only contains a selection of the most common options. For a full |
||||
# list see the documentation: |
||||
# https://www.sphinx-doc.org/en/master/usage/configuration.html |
||||
|
||||
# -- Path setup -------------------------------------------------------------- |
||||
|
||||
# If extensions (or modules to document with autodoc) are in another directory, |
||||
# add these directories to sys.path here. If the directory is relative to the |
||||
# documentation root, use os.path.abspath to make it absolute, like shown here. |
||||
# |
||||
import os |
||||
import sys |
||||
import sphinx_rtd_theme |
||||
sys.path.insert(0, os.path.abspath('../src')) |
||||
|
||||
|
||||
|
||||
# -- Project information ----------------------------------------------------- |
||||
|
||||
project = 'KyuPy' |
||||
copyright = '2020, Stefan Holst' |
||||
author = 'Stefan Holst' |
||||
|
||||
# The full version, including alpha/beta/rc tags |
||||
release = '0.0.2' |
||||
|
||||
|
||||
# -- General configuration --------------------------------------------------- |
||||
|
||||
# Add any Sphinx extension module names here, as strings. They can be |
||||
# extensions coming with Sphinx (named 'sphinx.ext.*') or your custom |
||||
# ones. |
||||
extensions = [ |
||||
'sphinx.ext.autodoc', |
||||
'sphinx_rtd_theme', |
||||
] |
||||
|
||||
# Add any paths that contain templates here, relative to this directory. |
||||
templates_path = ['_templates'] |
||||
|
||||
# List of patterns, relative to source directory, that match files and |
||||
# directories to ignore when looking for source files. |
||||
# This pattern also affects html_static_path and html_extra_path. |
||||
exclude_patterns = ['_build', 'Thumbs.db', '.DS_Store'] |
||||
|
||||
|
||||
# -- Options for HTML output ------------------------------------------------- |
||||
|
||||
# The theme to use for HTML and HTML Help pages. See the documentation for |
||||
# a list of builtin themes. |
||||
# |
||||
# html_theme = 'alabaster' |
||||
html_theme = 'sphinx_rtd_theme' |
||||
|
||||
# Add any paths that contain custom static files (such as style sheets) here, |
||||
# relative to this directory. They are copied after the builtin static files, |
||||
# so a file named "default.css" will overwrite the builtin "default.css". |
||||
html_static_path = ['_static'] |
||||
|
||||
autodoc_default_options = { |
||||
'member-order': 'bysource', |
||||
} |
||||
@ -0,0 +1,29 @@
@@ -0,0 +1,29 @@
|
||||
Data Structures |
||||
=============== |
||||
|
||||
KyuPy provides two types of core data structures, one for gate-level circuits, and a few others for representing and storing logic data and signal values. |
||||
The data structures are designed to work together nicely with numpy arrays. |
||||
For example, all the nodes and connections in the circuit graph have consecutive integer indices that can be used to access ndarrays with associated data. |
||||
Circuit graphs also define an ordering of inputs, outputs and other nodes to easily process test vector data and alike. |
||||
|
||||
Circuit Graph - :mod:`kyupy.circuit` |
||||
------------------------------------ |
||||
|
||||
.. automodule:: kyupy.circuit |
||||
|
||||
.. autoclass:: kyupy.circuit.Node |
||||
:members: |
||||
|
||||
.. autoclass:: kyupy.circuit.Line |
||||
:members: |
||||
|
||||
.. autoclass:: kyupy.circuit.Circuit |
||||
:members: |
||||
|
||||
Multi-Valued Logic - :mod:`kyupy.logic` |
||||
--------------------------------------- |
||||
|
||||
.. automodule:: kyupy.logic |
||||
:members: |
||||
|
||||
|
||||
@ -0,0 +1,12 @@
@@ -0,0 +1,12 @@
|
||||
.. include:: ../README.rst |
||||
|
||||
API Reference |
||||
------------- |
||||
|
||||
.. toctree:: |
||||
:maxdepth: 2 |
||||
|
||||
datastructures |
||||
parsers |
||||
simulators |
||||
|
||||
@ -0,0 +1,42 @@
@@ -0,0 +1,42 @@
|
||||
Parsers |
||||
======= |
||||
|
||||
KyuPy contains simple (and often incomplete) parsers for common file formats. |
||||
These parsers are tailored to the most common use-cases to keep the grammars and the code-base as simple as possible. |
||||
|
||||
Each of the modules export a function ``parse()`` for parsing a string directly and a function |
||||
``load()`` for loading a file. Files with a '.gz' extension are uncompressed on-the-fly. |
||||
|
||||
|
||||
Verilog - :mod:`kyupy.verilog` |
||||
------------------------------ |
||||
|
||||
.. automodule:: kyupy.verilog |
||||
:members: parse, load |
||||
|
||||
|
||||
Bench Format - :mod:`kyupy.bench` |
||||
--------------------------------- |
||||
|
||||
.. automodule:: kyupy.bench |
||||
:members: parse, load |
||||
|
||||
|
||||
Standard Test Interface Language - :mod:`kyupy.stil` |
||||
---------------------------------------------------- |
||||
|
||||
.. automodule:: kyupy.stil |
||||
:members: parse, load |
||||
|
||||
.. autoclass:: kyupy.stil.StilFile |
||||
:members: |
||||
|
||||
|
||||
Standard Delay Format - :mod:`kyupy.sdf` |
||||
---------------------------------------- |
||||
|
||||
.. automodule:: kyupy.sdf |
||||
:members: parse, load |
||||
|
||||
.. autoclass:: kyupy.sdf.DelayFile |
||||
:members: |
||||
@ -0,0 +1,20 @@
@@ -0,0 +1,20 @@
|
||||
Simulators |
||||
========== |
||||
|
||||
Logic Simulation - :mod:`kyupy.logic_sim` |
||||
----------------------------------------- |
||||
|
||||
.. autoclass:: kyupy.logic_sim.LogicSim |
||||
:members: |
||||
|
||||
|
||||
Timing Simulation - :mod:`kyupy.wave_sim` |
||||
----------------------------------------- |
||||
|
||||
.. automodule:: kyupy.wave_sim |
||||
|
||||
.. autoclass:: kyupy.wave_sim.WaveSim |
||||
:members: |
||||
|
||||
.. autoclass:: kyupy.wave_sim.WaveSimCuda |
||||
:members: |
||||
@ -1,23 +0,0 @@
@@ -1,23 +0,0 @@
|
||||
import numpy as np |
||||
import importlib.util |
||||
if importlib.util.find_spec('numba') is not None: |
||||
import numba |
||||
else: |
||||
from . import numba |
||||
print('Numba unavailable. Falling back to pure python') |
||||
|
||||
|
||||
_pop_count_lut = np.asarray([bin(x).count('1') for x in range(256)]) |
||||
|
||||
|
||||
def popcount(a): |
||||
return np.sum(_pop_count_lut[a]) |
||||
|
||||
|
||||
_bit_in_lut = np.array([2 ** x for x in range(7, -1, -1)], dtype='uint8') |
||||
|
||||
|
||||
@numba.njit |
||||
def bit_in(a, pos): |
||||
return a[pos >> 3] & _bit_in_lut[pos & 7] |
||||
|
||||
@ -0,0 +1,402 @@
@@ -0,0 +1,402 @@
|
||||
"""This module contains definitions and data structures for 2-, 4-, and 8-valued logic operations. |
||||
|
||||
8 logic values are defined as integer constants. |
||||
|
||||
* For 2-valued logic: ``ZERO`` and ``ONE`` |
||||
* 4-valued logic adds: ``UNASSIGNED`` and ``UNKNOWN`` |
||||
* 8-valued logic adds: ``RISE``, ``FALL``, ``PPULSE``, and ``NPULSE``. |
||||
|
||||
The bits in these constants have the following meaning: |
||||
|
||||
* bit 0: Final/settled binary value of a signal |
||||
* bit 1: Initial binary value of a signal |
||||
* bit 2: Activity or transitions are present on a signal |
||||
|
||||
Special meaning is given to values where bits 0 and 1 differ, but bit 2 (activity) is 0. |
||||
These values are interpreted as ``UNKNOWN`` or ``UNASSIGNED`` in 4-valued and 8-valued logic. |
||||
|
||||
In general, 2-valued logic only considers bit 0, 4-valued logic considers bits 0 and 1, and 8-valued logic |
||||
considers all 3 bits. |
||||
The only exception is constant ``ONE=0b11`` which has two bits set for all logics including 2-valued logic. |
||||
""" |
||||
|
||||
import math |
||||
from collections.abc import Iterable |
||||
|
||||
import numpy as np |
||||
|
||||
from . import numba |
||||
|
||||
|
||||
ZERO = 0b000 |
||||
"""Integer constant ``0b000`` for logic-0. ``'0'``, ``0``, ``False``, ``'L'``, and ``'l'`` are interpreted as ``ZERO``. |
||||
""" |
||||
UNKNOWN = 0b001 |
||||
"""Integer constant ``0b001`` for unknown or conflict. ``'X'``, or any other value is interpreted as ``UNKNOWN``. |
||||
""" |
||||
UNASSIGNED = 0b010 |
||||
"""Integer constant ``0b010`` for unassigned or high-impedance. ``'-'``, ``None``, ``'Z'``, and ``'z'`` are |
||||
interpreted as ``UNASSIGNED``. |
||||
""" |
||||
ONE = 0b011 |
||||
"""Integer constant ``0b011`` for logic-1. ``'1'``, ``1``, ``True``, ``'H'``, and ``'h'`` are interpreted as ``ONE``. |
||||
""" |
||||
PPULSE = 0b100 |
||||
"""Integer constant ``0b100`` for positive pulse, meaning initial and final values are 0, but there is some activity |
||||
on a signal. ``'P'``, ``'p'``, and ``'^'`` are interpreted as ``PPULSE``. |
||||
""" |
||||
RISE = 0b101 |
||||
"""Integer constant ``0b110`` for a rising transition. ``'R'``, ``'r'``, and ``'/'`` are interpreted as ``RISE``. |
||||
""" |
||||
FALL = 0b110 |
||||
"""Integer constant ``0b101`` for a falling transition. ``'F'``, ``'f'``, and ``'\\'`` are interpreted as ``FALL``. |
||||
""" |
||||
NPULSE = 0b111 |
||||
"""Integer constant ``0b111`` for negative pulse, meaning initial and final values are 1, but there is some activity |
||||
on a signal. ``'N'``, ``'n'``, and ``'v'`` are interpreted as ``NPULSE``. |
||||
""" |
||||
|
||||
|
||||
def interpret(value): |
||||
if isinstance(value, Iterable) and not (isinstance(value, str) and len(value) == 1): |
||||
return list(map(interpret, value)) |
||||
if value in [0, '0', False, 'L', 'l']: |
||||
return ZERO |
||||
if value in [1, '1', True, 'H', 'h']: |
||||
return ONE |
||||
if value in [None, '-', 'Z', 'z']: |
||||
return UNASSIGNED |
||||
if value in ['R', 'r', '/']: |
||||
return RISE |
||||
if value in ['F', 'f', '\\']: |
||||
return FALL |
||||
if value in ['P', 'p', '^']: |
||||
return PPULSE |
||||
if value in ['N', 'n', 'v']: |
||||
return NPULSE |
||||
return UNKNOWN |
||||
|
||||
|
||||
_bit_in_lut = np.array([2 ** x for x in range(7, -1, -1)], dtype='uint8') |
||||
|
||||
|
||||
@numba.njit |
||||
def bit_in(a, pos): |
||||
return a[pos >> 3] & _bit_in_lut[pos & 7] |
||||
|
||||
|
||||
def mv_cast(*args, m=8): |
||||
return [a if isinstance(a, MVArray) else MVArray(a, m=m) for a in args] |
||||
|
||||
|
||||
def mv_getm(*args): |
||||
return max([a.m for a in args if isinstance(a, MVArray)] + [0]) or 8 |
||||
|
||||
|
||||
def _mv_not(m, out, inp): |
||||
np.bitwise_xor(inp, 0b11, out=out) # this also exchanges UNASSIGNED <-> UNKNOWN |
||||
if m > 2: |
||||
np.putmask(out, (inp == UNKNOWN), UNKNOWN) # restore UNKNOWN |
||||
|
||||
|
||||
def mv_not(x1, out=None): |
||||
m = mv_getm(x1) |
||||
x1 = mv_cast(x1, m=m)[0] |
||||
out = out or MVArray(x1.data.shape, m=m) |
||||
_mv_not(m, out.data, x1.data) |
||||
return out |
||||
|
||||
|
||||
def _mv_or(m, out, *ins): |
||||
if m > 2: |
||||
any_unknown = (ins[0] == UNKNOWN) | (ins[0] == UNASSIGNED) |
||||
for inp in ins[1:]: any_unknown |= (inp == UNKNOWN) | (inp == UNASSIGNED) |
||||
any_one = (ins[0] == ONE) |
||||
for inp in ins[1:]: any_one |= (inp == ONE) |
||||
|
||||
out[...] = ZERO |
||||
np.putmask(out, any_one, ONE) |
||||
for inp in ins: |
||||
np.bitwise_or(out, inp, out=out, where=~any_one) |
||||
np.putmask(out, (any_unknown & ~any_one), UNKNOWN) |
||||
else: |
||||
out[...] = ZERO |
||||
for inp in ins: np.bitwise_or(out, inp, out=out) |
||||
|
||||
|
||||
def mv_or(x1, x2, out=None): |
||||
m = mv_getm(x1, x2) |
||||
x1, x2 = mv_cast(x1, x2, m=m) |
||||
out = out or MVArray(np.broadcast(x1.data, x2.data).shape, m=m) |
||||
_mv_or(m, out.data, x1.data, x2.data) |
||||
return out |
||||
|
||||
|
||||
def _mv_and(m, out, *ins): |
||||
if m > 2: |
||||
any_unknown = (ins[0] == UNKNOWN) | (ins[0] == UNASSIGNED) |
||||
for inp in ins[1:]: any_unknown |= (inp == UNKNOWN) | (inp == UNASSIGNED) |
||||
any_zero = (ins[0] == ZERO) |
||||
for inp in ins[1:]: any_zero |= (inp == ZERO) |
||||
|
||||
out[...] = ONE |
||||
np.putmask(out, any_zero, ZERO) |
||||
for inp in ins: |
||||
np.bitwise_and(out, inp | 0b100, out=out, where=~any_zero) |
||||
if m > 4: np.bitwise_or(out, inp & 0b100, out=out, where=~any_zero) |
||||
np.putmask(out, (any_unknown & ~any_zero), UNKNOWN) |
||||
else: |
||||
out[...] = ONE |
||||
for inp in ins: np.bitwise_and(out, inp, out=out) |
||||
|
||||
|
||||
def mv_and(x1, x2, out=None): |
||||
m = mv_getm(x1, x2) |
||||
x1, x2 = mv_cast(x1, x2, m=m) |
||||
out = out or MVArray(np.broadcast(x1.data, x2.data).shape, m=m) |
||||
_mv_and(m, out.data, x1.data, x2.data) |
||||
return out |
||||
|
||||
|
||||
def _mv_xor(m, out, *ins): |
||||
if m > 2: |
||||
any_unknown = (ins[0] == UNKNOWN) | (ins[0] == UNASSIGNED) |
||||
for inp in ins[1:]: any_unknown |= (inp == UNKNOWN) | (inp == UNASSIGNED) |
||||
|
||||
out[...] = ZERO |
||||
for inp in ins: |
||||
np.bitwise_xor(out, inp & 0b011, out=out) |
||||
if m > 4: np.bitwise_or(out, inp & 0b100, out=out) |
||||
np.putmask(out, any_unknown, UNKNOWN) |
||||
else: |
||||
out[...] = ZERO |
||||
for inp in ins: np.bitwise_xor(out, inp, out=out) |
||||
|
||||
|
||||
def mv_xor(x1, x2, out=None): |
||||
m = mv_getm(x1, x2) |
||||
x1, x2 = mv_cast(x1, x2, m=m) |
||||
out = out or MVArray(np.broadcast(x1.data, x2.data).shape, m=m) |
||||
_mv_xor(m, out.data, x1.data, x2.data) |
||||
return out |
||||
|
||||
|
||||
def mv_transition(init, final, out=None): |
||||
m = mv_getm(init, final) |
||||
init, final = mv_cast(init, final, m=m) |
||||
init = init.data |
||||
final = final.data |
||||
out = out or MVArray(np.broadcast(init, final).shape, m=8) |
||||
out.data[...] = (init & 0b010) | (final & 0b001) |
||||
out.data[...] |= ((out.data << 1) ^ (out.data << 2)) & 0b100 |
||||
unknown = (init == UNKNOWN) | (init == UNASSIGNED) | (final == UNKNOWN) | (final == UNASSIGNED) |
||||
unassigned = (init == UNASSIGNED) & (final == UNASSIGNED) |
||||
np.putmask(out.data, unknown, UNKNOWN) |
||||
np.putmask(out.data, unassigned, UNASSIGNED) |
||||
return out |
||||
|
||||
|
||||
class MVArray: |
||||
"""An n-dimensional array of m-valued logic values. |
||||
|
||||
This class wraps a numpy.ndarray of type uint8 and adds support for encoding and |
||||
interpreting 2-valued, 4-valued, and 8-valued logic values. |
||||
Each logic value is stored as an uint8, value manipulations are cheaper than in BPArray. |
||||
|
||||
An MVArray always has 2 axes: |
||||
|
||||
* Axis 0 is PI/PO/FF position, the length of this axis is called "width". |
||||
* Axis 1 is vector/pattern, the length of this axis is called "length". |
||||
|
||||
""" |
||||
|
||||
def __init__(self, a, m=None): |
||||
self.m = m or 8 |
||||
assert self.m in [2, 4, 8] |
||||
|
||||
# Try our best to interpret given a. |
||||
if isinstance(a, MVArray): |
||||
self.data = a.data.copy() |
||||
self.m = m or a.m |
||||
elif hasattr(a, 'data'): # assume it is a BPArray. Can't use isinstance() because BPArray isn't declared yet. |
||||
self.data = np.zeros((a.width, a.length), dtype=np.uint8) |
||||
self.m = m or a.m |
||||
for i in range(a.data.shape[-2]): |
||||
self.data[...] <<= 1 |
||||
self.data[...] |= np.unpackbits(a.data[..., -i-1, :], axis=1)[:, :a.length] |
||||
if a.data.shape[-2] == 1: |
||||
self.data *= 3 |
||||
elif isinstance(a, int): |
||||
self.data = np.full((a, 1), UNASSIGNED, dtype=np.uint8) |
||||
elif isinstance(a, tuple): |
||||
self.data = np.full(a, UNASSIGNED, dtype=np.uint8) |
||||
else: |
||||
if isinstance(a, str): a = [a] |
||||
self.data = np.asarray(interpret(a), dtype=np.uint8) |
||||
self.data = self.data[:, np.newaxis] if self.data.ndim == 1 else np.moveaxis(self.data, -2, -1) |
||||
|
||||
# Cast data to m-valued logic. |
||||
if self.m == 2: |
||||
self.data[...] = ((self.data & 0b001) & ((self.data >> 1) & 0b001) | (self.data == RISE)) * ONE |
||||
elif self.m == 4: |
||||
self.data[...] = (self.data & 0b011) & ((self.data != FALL) * ONE) | ((self.data == RISE) * ONE) |
||||
elif self.m == 8: |
||||
self.data[...] = self.data & 0b111 |
||||
|
||||
self.length = self.data.shape[-1] |
||||
self.width = self.data.shape[-2] |
||||
|
||||
def __repr__(self): |
||||
return f'<MVArray length={self.length} width={self.width} m={self.m} nbytes={self.data.nbytes}>' |
||||
|
||||
def __str__(self): |
||||
return str([self[idx] for idx in range(self.length)]) |
||||
|
||||
def __getitem__(self, vector_idx): |
||||
chars = ["0", "X", "-", "1", "P", "R", "F", "N"] |
||||
return ''.join(chars[v] for v in self.data[:, vector_idx]) |
||||
|
||||
def __len__(self): |
||||
return self.length |
||||
|
||||
|
||||
def bp_buf(out, inp): |
||||
md = out.shape[-2] |
||||
assert md == inp.shape[-2] |
||||
if md > 1: |
||||
unknown = inp[..., 0, :] ^ inp[..., 1, :] |
||||
if md > 2: unknown &= ~inp[..., 2, :] |
||||
out[..., 0, :] = inp[..., 0, :] | unknown |
||||
out[..., 1, :] = inp[..., 1, :] & ~unknown |
||||
if md > 2: out[..., 2, :] = inp[..., 2, :] & ~unknown |
||||
else: |
||||
out[..., 0, :] = inp[..., 0, :] |
||||
|
||||
|
||||
def bp_not(out, inp): |
||||
md = out.shape[-2] |
||||
assert md == inp.shape[-2] |
||||
if md > 1: |
||||
unknown = inp[..., 0, :] ^ inp[..., 1, :] |
||||
if md > 2: unknown &= ~inp[..., 2, :] |
||||
out[..., 0, :] = ~inp[..., 0, :] | unknown |
||||
out[..., 1, :] = ~inp[..., 1, :] & ~unknown |
||||
if md > 2: out[..., 2, :] = inp[..., 2, :] & ~unknown |
||||
else: |
||||
out[..., 0, :] = ~inp[..., 0, :] |
||||
|
||||
|
||||
def bp_or(out, *ins): |
||||
md = out.shape[-2] |
||||
for inp in ins: assert md == inp.shape[-2] |
||||
out[...] = 0 |
||||
if md == 1: |
||||
for inp in ins: out[..., 0, :] |= inp[..., 0, :] |
||||
elif md == 2: |
||||
any_unknown = ins[0][..., 0, :] ^ ins[0][..., 1, :] |
||||
for inp in ins[1:]: any_unknown |= inp[..., 0, :] ^ inp[..., 1, :] |
||||
any_one = ins[0][..., 0, :] & ins[0][..., 1, :] |
||||
for inp in ins[1:]: any_one |= inp[..., 0, :] & inp[..., 1, :] |
||||
for inp in ins: |
||||
out[..., 0, :] |= inp[..., 0, :] | any_unknown |
||||
out[..., 1, :] |= inp[..., 1, :] & (~any_unknown | any_one) |
||||
else: |
||||
any_unknown = (ins[0][..., 0, :] ^ ins[0][..., 1, :]) & ~ins[0][..., 2, :] |
||||
for inp in ins[1:]: any_unknown |= (inp[..., 0, :] ^ inp[..., 1, :]) & ~inp[..., 2, :] |
||||
any_one = ins[0][..., 0, :] & ins[0][..., 1, :] & ~ins[0][..., 2, :] |
||||
for inp in ins[1:]: any_one |= inp[..., 0, :] & inp[..., 1, :] & ~inp[..., 2, :] |
||||
for inp in ins: |
||||
out[..., 0, :] |= inp[..., 0, :] | any_unknown |
||||
out[..., 1, :] |= inp[..., 1, :] & (~any_unknown | any_one) |
||||
out[..., 2, :] |= inp[..., 2, :] & (~any_unknown | any_one) & ~any_one |
||||
|
||||
|
||||
def bp_and(out, *ins): |
||||
md = out.shape[-2] |
||||
for inp in ins: assert md == inp.shape[-2] |
||||
out[...] = 0xff |
||||
if md == 1: |
||||
for inp in ins: out[..., 0, :] &= inp[..., 0, :] |
||||
elif md == 2: |
||||
any_unknown = ins[0][..., 0, :] ^ ins[0][..., 1, :] |
||||
for inp in ins[1:]: any_unknown |= inp[..., 0, :] ^ inp[..., 1, :] |
||||
any_zero = ~ins[0][..., 0, :] & ~ins[0][..., 1, :] |
||||
for inp in ins[1:]: any_zero |= ~inp[..., 0, :] & ~inp[..., 1, :] |
||||
for inp in ins: |
||||
out[..., 0, :] &= inp[..., 0, :] | (any_unknown & ~any_zero) |
||||
out[..., 1, :] &= inp[..., 1, :] & ~any_unknown |
||||
else: |
||||
any_unknown = (ins[0][..., 0, :] ^ ins[0][..., 1, :]) & ~ins[0][..., 2, :] |
||||
for inp in ins[1:]: any_unknown |= (inp[..., 0, :] ^ inp[..., 1, :]) & ~inp[..., 2, :] |
||||
any_zero = ~ins[0][..., 0, :] & ~ins[0][..., 1, :] & ~ins[0][..., 2, :] |
||||
for inp in ins[1:]: any_zero |= ~inp[..., 0, :] & ~inp[..., 1, :] & ~inp[..., 2, :] |
||||
out[..., 2, :] = 0 |
||||
for inp in ins: |
||||
out[..., 0, :] &= inp[..., 0, :] | (any_unknown & ~any_zero) |
||||
out[..., 1, :] &= inp[..., 1, :] & ~any_unknown |
||||
out[..., 2, :] |= inp[..., 2, :] & (~any_unknown | any_zero) & ~any_zero |
||||
|
||||
|
||||
def bp_xor(out, *ins): |
||||
md = out.shape[-2] |
||||
for inp in ins: assert md == inp.shape[-2] |
||||
out[...] = 0 |
||||
if md == 1: |
||||
for inp in ins: out[..., 0, :] ^= inp[..., 0, :] |
||||
elif md == 2: |
||||
any_unknown = ins[0][..., 0, :] ^ ins[0][..., 1, :] |
||||
for inp in ins[1:]: any_unknown |= inp[..., 0, :] ^ inp[..., 1, :] |
||||
for inp in ins: out[...] ^= inp |
||||
out[..., 0, :] |= any_unknown |
||||
out[..., 1, :] &= ~any_unknown |
||||
else: |
||||
any_unknown = (ins[0][..., 0, :] ^ ins[0][..., 1, :]) & ~ins[0][..., 2, :] |
||||
for inp in ins[1:]: any_unknown |= (inp[..., 0, :] ^ inp[..., 1, :]) & ~inp[..., 2, :] |
||||
for inp in ins: |
||||
out[..., 0, :] ^= inp[..., 0, :] |
||||
out[..., 1, :] ^= inp[..., 1, :] |
||||
out[..., 2, :] |= inp[..., 2, :] |
||||
out[..., 0, :] |= any_unknown |
||||
out[..., 1, :] &= ~any_unknown |
||||
out[..., 2, :] &= ~any_unknown |
||||
|
||||
|
||||
class BPArray: |
||||
"""An n-dimensional array of m-valued logic values that uses bit-parallel storage. |
||||
|
||||
The primary use of this format is in aiding efficient bit-parallel logic simulation. |
||||
The secondary benefit over MVArray is its memory efficiency. |
||||
Accessing individual values is more expensive than with :py:class:`MVArray`. |
||||
It is advised to first construct a MVArray, pack it into a :py:class:`BPArray` for simulation and unpack the results |
||||
back into a :py:class:`MVArray` for value access. |
||||
|
||||
The values along the last axis (vectors/patterns) are packed into uint8 words. |
||||
The second-last axis has length ceil(log2(m)) for storing all bits. |
||||
All other axes stay the same as in MVArray. |
||||
""" |
||||
|
||||
def __init__(self, a, m=None): |
||||
if not isinstance(a, MVArray) and not isinstance(a, BPArray): |
||||
a = MVArray(a, m) |
||||
self.m = a.m |
||||
if isinstance(a, MVArray): |
||||
if m is not None and m != a.m: |
||||
a = MVArray(a, m) # cast data |
||||
self.m = a.m |
||||
assert self.m in [2, 4, 8] |
||||
nwords = math.ceil(math.log2(self.m)) |
||||
nbytes = (a.data.shape[-1] - 1) // 8 + 1 |
||||
self.data = np.zeros(a.data.shape[:-1] + (nwords, nbytes), dtype=np.uint8) |
||||
for i in range(self.data.shape[-2]): |
||||
self.data[..., i, :] = np.packbits((a.data >> i) & 1, axis=-1) |
||||
else: # we have a BPArray |
||||
self.data = a.data.copy() # TODO: support conversion to different m |
||||
self.m = a.m |
||||
self.length = a.length |
||||
self.width = a.width |
||||
|
||||
def __repr__(self): |
||||
return f'<BPArray length={self.length} width={self.width} m={self.m} bytes={self.data.nbytes}>' |
||||
|
||||
def __len__(self): |
||||
return self.length |
||||
@ -1,299 +0,0 @@
@@ -1,299 +0,0 @@
|
||||
import numpy as np |
||||
from .bittools import popcount, bit_in |
||||
|
||||
|
||||
class PackedVectors: |
||||
def __init__(self, nvectors=8, width=1, vdim=1, from_cache=None): |
||||
if from_cache is not None: |
||||
self.bits = np.array(from_cache) |
||||
self.width, self.vdim, nbytes = self.bits.shape |
||||
else: |
||||
self.bits = np.zeros((width, vdim, (nvectors - 1) // 8 + 1), dtype='uint8') |
||||
self.vdim = vdim |
||||
self.width = width |
||||
self.nvectors = nvectors |
||||
m1 = np.array([2 ** x for x in range(7, -1, -1)], dtype='uint8') |
||||
m0 = ~m1 |
||||
self.mask = np.rollaxis(np.vstack((m0, m1)), 1) |
||||
|
||||
@classmethod |
||||
def from_pair(cls, init, final): |
||||
assert init.nvectors == final.nvectors |
||||
assert len(init.bits) == len(final.bits) |
||||
init_v = init.bits[:, 0] |
||||
if init.vdim == 3: |
||||
init_c = (init.bits[:, 0] ^ init.bits[:, 1]) | init.bits[:, 2] |
||||
elif init.vdim == 2: |
||||
init_c = init.bits[:, 1] |
||||
else: |
||||
init_c = ~np.zeros_like(init.bits[:, 0]) |
||||
final_v = final.bits[:, 0] |
||||
if final.vdim == 3: |
||||
final_c = (final.bits[:, 0] ^ final.bits[:, 1]) | final.bits[:, 2] |
||||
final_v = ~final.bits[:, 1] |
||||
elif final.vdim == 2: |
||||
final_c = final.bits[:, 1] |
||||
else: |
||||
final_c = ~np.zeros_like(final.bits[:, 0]) |
||||
c = init_c & final_c |
||||
a0 = init_v & c |
||||
a1 = ~final_v & c |
||||
a2 = (init_v ^ final_v) & c |
||||
p = PackedVectors(init.nvectors, len(init.bits), 3) |
||||
p.bits[:, 0] = a0 |
||||
p.bits[:, 1] = a1 |
||||
p.bits[:, 2] = a2 |
||||
return p |
||||
|
||||
def transition_vectors(self): |
||||
a = PackedVectors(self.nvectors-1, self.width, 3) |
||||
for pos in range(self.width): |
||||
for vidx in range(self.nvectors-1): |
||||
tr = self.get_value(vidx, pos) + self.get_value(vidx+1, pos) |
||||
if tr == '00': |
||||
a.set_value(vidx, pos, '0') |
||||
elif tr == '11': |
||||
a.set_value(vidx, pos, '1') |
||||
elif tr == '01': |
||||
a.set_value(vidx, pos, 'R') |
||||
elif tr == '10': |
||||
a.set_value(vidx, pos, 'F') |
||||
elif tr == '--': |
||||
a.set_value(vidx, pos, '-') |
||||
else: |
||||
a.set_value(vidx, pos, 'X') |
||||
return a |
||||
|
||||
def __add__(self, other): |
||||
a = PackedVectors(self.nvectors + other.nvectors, self.width, max(self.vdim, other.vdim)) |
||||
# a.bits[:self.bits.shape[0], 0] = self.bits[:, 0] |
||||
# if self.vdim == 2: |
||||
# a.bits[:self.bits.shape[0], 1] = self.care_bits |
||||
# elif self.vdim == 3: |
||||
# a.bits[:self.bits.shape[0], 1] = ~self.value_bits |
||||
# a.bits[:self.bits.shape[0], 2] = self.toggle_bits |
||||
for i in range(self.nvectors): |
||||
a[i] = self[i] |
||||
for i in range(len(other)): |
||||
a[self.nvectors+i] = other[i] |
||||
return a |
||||
|
||||
def __len__(self): |
||||
return self.nvectors |
||||
|
||||
def randomize(self, one_probability=0.5): |
||||
for data in self.bits: |
||||
data[0] = np.packbits((np.random.rand(self.nvectors) < one_probability).astype(int)) |
||||
if self.vdim == 2: |
||||
data[1] = 255 |
||||
elif self.vdim == 3: |
||||
data[1] = ~np.packbits((np.random.rand(self.nvectors) < one_probability).astype(int)) |
||||
data[2] = data[0] ^ ~data[1] |
||||
|
||||
def copy(self, selection_mask=None): |
||||
if selection_mask is not None: |
||||
cpy = PackedVectors(popcount(selection_mask), len(self.bits), self.vdim) |
||||
cur = 0 |
||||
for vidx in range(self.nvectors): |
||||
if bit_in(selection_mask, vidx): |
||||
cpy[cur] = self[vidx] |
||||
cur += 1 |
||||
else: |
||||
cpy = PackedVectors(self.nvectors, len(self.bits), self.vdim) |
||||
np.copyto(cpy.bits, self.bits) |
||||
return cpy |
||||
|
||||
@property |
||||
def care_bits(self): |
||||
if self.vdim == 1: |
||||
return self.bits[:, 0] | 255 |
||||
elif self.vdim == 2: |
||||
return self.bits[:, 1] |
||||
elif self.vdim == 3: |
||||
return (self.bits[:, 0] ^ self.bits[:, 1]) | self.bits[:, 2] |
||||
|
||||
@property |
||||
def initial_bits(self): |
||||
return self.bits[:, 0] |
||||
|
||||
@property |
||||
def value_bits(self): |
||||
if self.vdim == 3: |
||||
return ~self.bits[:, 1] |
||||
else: |
||||
return self.bits[:, 0] |
||||
|
||||
@property |
||||
def toggle_bits(self): |
||||
if self.vdim == 3: |
||||
return self.bits[:, 2] |
||||
else: |
||||
return self.bits[:, 0] & 0 |
||||
|
||||
def get_value(self, vector, position): |
||||
if vector >= self.nvectors: |
||||
raise IndexError(f'vector out of range: {vector} >= {self.nvectors}') |
||||
a = self.bits[position, :, vector // 8] |
||||
m = self.mask[vector % 8] |
||||
if self.vdim == 1: |
||||
return '1' if a[0] & m[1] else '0' |
||||
elif self.vdim == 2: |
||||
if a[0] & m[1]: |
||||
return '1' if a[1] & m[1] else 'X' |
||||
else: |
||||
return '0' if a[1] & m[1] else '-' |
||||
elif self.vdim == 3: |
||||
if a[2] & m[1]: |
||||
if a[0] & m[1]: |
||||
return 'F' if a[1] & m[1] else 'N' |
||||
else: |
||||
return 'P' if a[1] & m[1] else 'R' |
||||
else: |
||||
if a[0] & m[1]: |
||||
return 'X' if a[1] & m[1] else '1' |
||||
else: |
||||
return '0' if a[1] & m[1] else '-' |
||||
|
||||
def get_values_for_position(self, position): |
||||
return ''.join(self.get_value(x, position) for x in range(self.nvectors)) |
||||
|
||||
def set_value(self, vector, position, v): |
||||
if vector >= self.nvectors: |
||||
raise IndexError(f'vector out of range: {vector} >= {self.nvectors}') |
||||
a = self.bits[position, :, vector // 8] |
||||
m = self.mask[vector % 8] |
||||
if self.vdim == 1: |
||||
self._set_value_vd1(a, m, v) |
||||
elif self.vdim == 2: |
||||
self._set_value_vd2(a, m, v) |
||||
elif self.vdim == 3: |
||||
self._set_value_vd3(a, m, v) |
||||
|
||||
def set_values(self, vector, v, mapping=None, inversions=None): |
||||
if vector >= self.nvectors: |
||||
raise IndexError(f'vector out of range: {vector} >= {self.nvectors}') |
||||
if not mapping: |
||||
mapping = [y for y in range(len(v))] |
||||
if inversions is None: |
||||
inversions = [False] * len(v) |
||||
for i, c in enumerate(v): |
||||
if inversions[i]: |
||||
if c == '1': |
||||
c = '0' |
||||
elif c == '0': |
||||
c = '1' |
||||
elif c == 'H': |
||||
c = 'L' |
||||
elif c == 'L': |
||||
c = 'H' |
||||
elif c == 'R': |
||||
c = 'F' |
||||
elif c == 'F': |
||||
c = 'R' |
||||
self.set_value(vector, mapping[i], c) |
||||
|
||||
def set_values_for_position(self, position, values): |
||||
for i, v in enumerate(values): |
||||
self.set_value(i, position, v) |
||||
|
||||
def __setitem__(self, vector, value): |
||||
for i, c in enumerate(value): |
||||
self.set_value(vector, i, c) |
||||
|
||||
def __getitem__(self, vector): |
||||
if isinstance(vector, slice): |
||||
first = self.get_values_for_position(0)[vector] |
||||
ret = PackedVectors(len(first), self.width, self.vdim) |
||||
ret.set_values_for_position(0, first) |
||||
for pos in range(1, self.width): |
||||
ret.set_values_for_position(pos, self.get_values_for_position(pos)[vector]) |
||||
return ret |
||||
return ''.join(self.get_value(vector, pos) for pos in range(len(self.bits))) |
||||
|
||||
@staticmethod |
||||
def _set_value_vd1(a, m, v): |
||||
if v in [True, 1, '1', 'H', 'h']: |
||||
a[0] |= m[1] |
||||
else: |
||||
a[0] &= m[0] |
||||
|
||||
@staticmethod |
||||
def _set_value_vd2(a, m, v): |
||||
if v in [True, 1, '1', 'H', 'h']: |
||||
a[0] |= m[1] |
||||
a[1] |= m[1] |
||||
elif v in [False, 0, '0', 'L', 'l']: |
||||
a[0] &= m[0] |
||||
a[1] |= m[1] |
||||
elif v in ['X', 'x']: |
||||
a[0] |= m[1] |
||||
a[1] &= m[0] |
||||
else: |
||||
a[0] &= m[0] |
||||
a[1] &= m[0] |
||||
|
||||
# i fb act |
||||
# a 0 1 2 |
||||
# - 0 0 0 None, '-' |
||||
# 0 0 1 0 False, 0, '0', 'l', 'L' |
||||
# 1 1 0 0 True, 1, '1', 'h', 'H' |
||||
# X 1 1 0 'x', 'X' |
||||
# / 0 0 1 '/', 'r', 'R' |
||||
# ^ 0 1 1 '^', 'p', 'P' |
||||
# v 1 0 1 'v', 'n', 'N' |
||||
# \ 1 1 1 '\', 'f', 'F' |
||||
@staticmethod |
||||
def _set_value_vd3(a, m, v): |
||||
if v in [False, 0, '0', 'L', 'l']: |
||||
a[0] &= m[0] |
||||
a[1] |= m[1] |
||||
a[2] &= m[0] |
||||
elif v in [True, 1, '1', 'H', 'h']: |
||||
a[0] |= m[1] |
||||
a[1] &= m[0] |
||||
a[2] &= m[0] |
||||
elif v in ['X', 'x']: |
||||
a[0] |= m[1] |
||||
a[1] |= m[1] |
||||
a[2] &= m[0] |
||||
elif v in ['/', 'r', 'R']: |
||||
a[0] &= m[0] |
||||
a[1] &= m[0] |
||||
a[2] |= m[1] |
||||
elif v in ['^', 'p', 'P']: |
||||
a[0] &= m[0] |
||||
a[1] |= m[1] |
||||
a[2] |= m[1] |
||||
elif v in ['v', 'n', 'N']: |
||||
a[0] |= m[1] |
||||
a[1] &= m[0] |
||||
a[2] |= m[1] |
||||
elif v in ['\\', 'f', 'F']: |
||||
a[0] |= m[1] |
||||
a[1] |= m[1] |
||||
a[2] |= m[1] |
||||
else: |
||||
a[0] &= m[0] |
||||
a[1] &= m[0] |
||||
a[2] &= m[0] |
||||
|
||||
def __repr__(self): |
||||
return f'<PackedVectors nvectors={self.nvectors}, width={self.width}, vdim={self.vdim}>' |
||||
|
||||
def __str__(self): |
||||
lst = [] |
||||
for p in range(self.nvectors): |
||||
lst.append(''.join(self.get_value(p, w) for w in range(len(self.bits)))) |
||||
if len(lst) == 0: return '' |
||||
if len(lst[0]) > 64: |
||||
lst = [s[:32] + '...' + s[-32:] for s in lst] |
||||
if len(lst) <= 16: |
||||
return '\n'.join(lst) |
||||
else: |
||||
return '\n'.join(lst[:8]) + '\n...\n' + '\n'.join(lst[-8:]) |
||||
|
||||
def diff(self, other, out=None): |
||||
if out is None: |
||||
out = np.zeros((self.width, self.bits.shape[-1]), dtype='uint8') |
||||
out[...] = (self.value_bits ^ other.value_bits) & self.care_bits & other.care_bits |
||||
return out |
||||
@ -1,317 +0,0 @@
@@ -1,317 +0,0 @@
|
||||
import numpy as np |
||||
import math |
||||
from .wave_sim import WaveSim |
||||
from . import cuda |
||||
|
||||
TMAX = np.float32(2 ** 127) # almost np.PINF for 32-bit floating point values |
||||
TMAX_OVL = np.float32(1.1 * 2 ** 127) # almost np.PINF with overflow mark |
||||
TMIN = np.float32(-2 ** 127) # almost np.NINF for 32-bit floating point values |
||||
|
||||
|
||||
class WaveSimCuda(WaveSim): |
||||
def __init__(self, circuit, timing, sims=8, wavecaps=16, strip_forks=False, keep_waveforms=True): |
||||
super().__init__(circuit, timing, sims, wavecaps, strip_forks, keep_waveforms) |
||||
|
||||
self.tdata = np.zeros((len(self.interface), 3, (sims - 1) // 8 + 1), dtype='uint8') |
||||
|
||||
self.d_state = cuda.to_device(self.state) |
||||
self.d_sat = cuda.to_device(self.sat) |
||||
self.d_ops = cuda.to_device(self.ops) |
||||
self.d_timing = cuda.to_device(self.timing) |
||||
self.d_tdata = cuda.to_device(self.tdata) |
||||
self.d_cdata = cuda.to_device(self.cdata) |
||||
|
||||
self._block_dim = (32, 16) |
||||
|
||||
def get_line_delay(self, line, polarity): |
||||
return self.d_timing[line, 0, polarity] |
||||
|
||||
def set_line_delay(self, line, polarity, delay): |
||||
self.d_timing[line, 0, polarity] = delay |
||||
|
||||
def assign(self, vectors, time=0.0, offset=0): |
||||
assert (offset % 8) == 0 |
||||
byte_offset = offset // 8 |
||||
assert byte_offset < vectors.bits.shape[-1] |
||||
pdim = min(vectors.bits.shape[-1] - byte_offset, self.tdata.shape[-1]) |
||||
|
||||
self.tdata[..., 0:pdim] = vectors.bits[..., byte_offset:pdim + byte_offset] |
||||
if vectors.vdim == 1: |
||||
self.tdata[:, 1, 0:pdim] = ~self.tdata[:, 1, 0:pdim] |
||||
self.tdata[:, 2, 0:pdim] = 0 |
||||
cuda.to_device(self.tdata, to=self.d_tdata) |
||||
|
||||
grid_dim = self._grid_dim(self.sims, len(self.interface)) |
||||
assign_kernel[grid_dim, self._block_dim](self.d_state, self.d_sat, self.ppi_offset, |
||||
len(self.interface), self.d_tdata, time) |
||||
|
||||
def _grid_dim(self, x, y): |
||||
gx = math.ceil(x / self._block_dim[0]) |
||||
gy = math.ceil(y / self._block_dim[1]) |
||||
return gx, gy |
||||
|
||||
def propagate(self, sims=None, sd=0.0, seed=1): |
||||
if sims is None: |
||||
sims = self.sims |
||||
else: |
||||
sims = min(sims, self.sims) |
||||
for op_start, op_stop in zip(self.level_starts, self.level_stops): |
||||
grid_dim = self._grid_dim(sims, op_stop - op_start) |
||||
wave_kernel[grid_dim, self._block_dim](self.d_ops, op_start, op_stop, self.d_state, self.sat, int(0), |
||||
sims, self.d_timing, sd, seed) |
||||
cuda.synchronize() |
||||
self.lst_eat_valid = False |
||||
|
||||
def wave(self, line, vector): |
||||
if line < 0: |
||||
return None |
||||
mem, wcap, _ = self.sat[line] |
||||
if mem < 0: |
||||
return None |
||||
return self.d_state[mem:mem + wcap, vector] |
||||
|
||||
def capture(self, time=TMAX, sd=0, seed=1, cdata=None, offset=0): |
||||
grid_dim = self._grid_dim(self.sims, len(self.interface)) |
||||
capture_kernel[grid_dim, self._block_dim](self.d_state, self.d_sat, self.ppo_offset, |
||||
self.d_cdata, time, sd * math.sqrt(2), seed) |
||||
self.cdata[...] = self.d_cdata |
||||
if cdata is not None: |
||||
assert offset < cdata.shape[1] |
||||
cap_dim = min(cdata.shape[1] - offset, self.sims) |
||||
cdata[:, offset:cap_dim + offset] = self.cdata[:, 0:cap_dim] |
||||
self.lst_eat_valid = True |
||||
return self.cdata |
||||
|
||||
def reassign(self, time=0.0): |
||||
grid_dim = self._grid_dim(self.sims, len(self.interface)) |
||||
reassign_kernel[grid_dim, self._block_dim](self.d_state, self.d_sat, self.ppi_offset, self.ppo_offset, |
||||
self.d_cdata, time) |
||||
cuda.synchronize() |
||||
|
||||
def wavecaps(self): |
||||
gx = math.ceil(len(self.circuit.lines) / 512) |
||||
wavecaps_kernel[gx, 512](self.d_state, self.d_sat, self.sims) |
||||
self.sat[...] = self.d_sat |
||||
return self.sat[..., 2] |
||||
|
||||
|
||||
@cuda.jit() |
||||
def wavecaps_kernel(state, sat, sims): |
||||
idx = cuda.grid(1) |
||||
if idx >= len(sat): return |
||||
|
||||
lidx, lcap, _ = sat[idx] |
||||
if lidx < 0: return |
||||
|
||||
wcap = 0 |
||||
for sidx in range(sims): |
||||
for tidx in range(lcap): |
||||
t = state[lidx + tidx, sidx] |
||||
if tidx > wcap: |
||||
wcap = tidx |
||||
if t >= TMAX: break |
||||
|
||||
sat[idx, 2] = wcap + 1 |
||||
|
||||
|
||||
@cuda.jit() |
||||
def reassign_kernel(state, sat, ppi_offset, ppo_offset, cdata, ppi_time): |
||||
vector, y = cuda.grid(2) |
||||
if vector >= state.shape[-1]: return |
||||
if ppo_offset + y >= len(sat): return |
||||
|
||||
ppo, ppo_cap, _ = sat[ppo_offset + y] |
||||
ppi, ppi_cap, _ = sat[ppi_offset + y] |
||||
if ppo < 0: return |
||||
if ppi < 0: return |
||||
|
||||
ppo_val = int(cdata[y, vector, 1]) |
||||
ppi_val = int(0) |
||||
for tidx in range(ppi_cap): |
||||
t = state[ppi + tidx, vector] |
||||
if t >= TMAX: break |
||||
ppi_val ^= 1 |
||||
|
||||
# make new waveform at PPI |
||||
toggle = 0 |
||||
if ppi_val: |
||||
state[ppi + toggle, vector] = TMIN |
||||
toggle += 1 |
||||
if ppi_val != ppo_val: |
||||
state[ppi + toggle, vector] = ppi_time |
||||
toggle += 1 |
||||
state[ppi + toggle, vector] = TMAX |
||||
|
||||
|
||||
@cuda.jit() |
||||
def capture_kernel(state, sat, ppo_offset, cdata, time, s_sqrt2, seed): |
||||
x, y = cuda.grid(2) |
||||
if ppo_offset + y >= len(sat): return |
||||
line, tdim, _ = sat[ppo_offset + y] |
||||
if line < 0: return |
||||
if x >= state.shape[-1]: return |
||||
vector = x |
||||
m = 0.5 |
||||
acc = 0.0 |
||||
eat = TMAX |
||||
lst = TMIN |
||||
tog = 0 |
||||
ovl = 0 |
||||
val = int(0) |
||||
final = int(0) |
||||
for tidx in range(tdim): |
||||
t = state[line + tidx, vector] |
||||
if t >= TMAX: |
||||
if t == TMAX_OVL: |
||||
ovl = 1 |
||||
break |
||||
m = -m |
||||
final ^= 1 |
||||
if t < time: |
||||
val ^= 1 |
||||
if t <= TMIN: continue |
||||
if s_sqrt2 > 0: |
||||
acc += m * (1 + math.erf((t - time) / s_sqrt2)) |
||||
eat = min(eat, t) |
||||
lst = max(lst, t) |
||||
tog += 1 |
||||
if s_sqrt2 > 0: |
||||
if m < 0: |
||||
acc += 1 |
||||
if acc >= 0.99: |
||||
val = 1 |
||||
elif acc > 0.01: |
||||
seed = (seed << 4) + (vector << 20) + (y << 1) |
||||
seed = int(0xDEECE66D) * seed + 0xB |
||||
seed = int(0xDEECE66D) * seed + 0xB |
||||
rnd = float((seed >> 8) & 0xffffff) / float(1 << 24) |
||||
val = rnd < acc |
||||
else: |
||||
val = 0 |
||||
else: |
||||
acc = val |
||||
|
||||
cdata[y, vector, 0] = acc |
||||
cdata[y, vector, 1] = val |
||||
cdata[y, vector, 2] = final |
||||
cdata[y, vector, 3] = (val != final) |
||||
cdata[y, vector, 4] = eat |
||||
cdata[y, vector, 5] = lst |
||||
cdata[y, vector, 6] = ovl |
||||
|
||||
|
||||
@cuda.jit() |
||||
def assign_kernel(state, sat, ppi_offset, intf_len, tdata, time): |
||||
x, y = cuda.grid(2) |
||||
if y >= intf_len: return |
||||
line = sat[ppi_offset + y, 0] |
||||
if line < 0: return |
||||
sdim = state.shape[-1] |
||||
if x >= sdim: return |
||||
vector = x |
||||
a0 = tdata[y, 0, vector // 8] |
||||
a1 = tdata[y, 1, vector // 8] |
||||
a2 = tdata[y, 2, vector // 8] |
||||
m = np.uint8(1 << (7 - (vector % 8))) |
||||
toggle = 0 |
||||
if a0 & m: |
||||
state[line + toggle, x] = TMIN |
||||
toggle += 1 |
||||
if (a2 & m) and ((a0 & m) == (a1 & m)): |
||||
state[line + toggle, x] = time |
||||
toggle += 1 |
||||
state[line + toggle, x] = TMAX |
||||
|
||||
|
||||
@cuda.jit(device=True) |
||||
def rand_gauss(seed, sd): |
||||
clamp = 0.5 |
||||
if sd <= 0.0: |
||||
return 1.0 |
||||
while True: |
||||
x = -6.0 |
||||
for i in range(12): |
||||
seed = int(0xDEECE66D) * seed + 0xB |
||||
x += float((seed >> 8) & 0xffffff) / float(1 << 24) |
||||
x *= sd |
||||
if abs(x) <= clamp: |
||||
break |
||||
return x + 1.0 |
||||
|
||||
|
||||
@cuda.jit() |
||||
def wave_kernel(ops, op_start, op_stop, state, sat, st_start, st_stop, line_times, sd, seed): |
||||
x, y = cuda.grid(2) |
||||
st_idx = st_start + x |
||||
op_idx = op_start + y |
||||
if st_idx >= st_stop: return |
||||
if op_idx >= op_stop: return |
||||
lut = ops[op_idx, 0] |
||||
z_idx = ops[op_idx, 1] |
||||
a_idx = ops[op_idx, 2] |
||||
b_idx = ops[op_idx, 3] |
||||
overflows = int(0) |
||||
|
||||
_seed = (seed << 4) + (z_idx << 20) + (st_idx << 1) |
||||
|
||||
a_mem = sat[a_idx, 0] |
||||
b_mem = sat[b_idx, 0] |
||||
z_mem, z_cap, _ = sat[z_idx] |
||||
|
||||
a_cur = int(0) |
||||
b_cur = int(0) |
||||
z_cur = lut & 1 |
||||
if z_cur == 1: |
||||
state[z_mem, st_idx] = TMIN |
||||
|
||||
a = state[a_mem, st_idx] + line_times[a_idx, 0, z_cur] * rand_gauss(_seed ^ a_mem ^ z_cur, sd) |
||||
b = state[b_mem, st_idx] + line_times[b_idx, 0, z_cur] * rand_gauss(_seed ^ b_mem ^ z_cur, sd) |
||||
|
||||
previous_t = TMIN |
||||
|
||||
current_t = min(a, b) |
||||
inputs = int(0) |
||||
|
||||
while current_t < TMAX: |
||||
z_val = z_cur & 1 |
||||
if b < a: |
||||
b_cur += 1 |
||||
b = state[b_mem + b_cur, st_idx] |
||||
b += line_times[b_idx, 0, z_val ^ 1] * rand_gauss(_seed ^ b_mem ^ z_val ^ 1, sd) |
||||
thresh = line_times[b_idx, 1, z_val] * rand_gauss(_seed ^ b_mem ^ z_val, sd) |
||||
inputs ^= 2 |
||||
next_t = b |
||||
else: |
||||
a_cur += 1 |
||||
a = state[a_mem + a_cur, st_idx] |
||||
a += line_times[a_idx, 0, z_val ^ 1] * rand_gauss(_seed ^ a_mem ^ z_val ^ 1, sd) |
||||
thresh = line_times[a_idx, 1, z_val] * rand_gauss(_seed ^ a_mem ^ z_val, sd) |
||||
inputs ^= 1 |
||||
next_t = a |
||||
|
||||
if (z_cur & 1) != ((lut >> inputs) & 1): |
||||
# we generate a toggle in z_mem, if: |
||||
# ( it is the first toggle in z_mem OR |
||||
# following toggle is earlier OR |
||||
# pulse is wide enough ) AND enough space in z_mem. |
||||
if z_cur == 0 or next_t < current_t or (current_t - previous_t) > thresh: |
||||
if z_cur < (z_cap - 1): |
||||
state[z_mem + z_cur, st_idx] = current_t |
||||
previous_t = current_t |
||||
z_cur += 1 |
||||
else: |
||||
overflows += 1 |
||||
previous_t = state[z_mem + z_cur - 1, st_idx] |
||||
z_cur -= 1 |
||||
else: |
||||
z_cur -= 1 |
||||
if z_cur > 0: |
||||
previous_t = state[z_mem + z_cur - 1, st_idx] |
||||
else: |
||||
previous_t = TMIN |
||||
current_t = min(a, b) |
||||
|
||||
if overflows > 0: |
||||
state[z_mem + z_cur, st_idx] = TMAX_OVL |
||||
else: |
||||
state[z_mem + z_cur, st_idx] = a if a > b else b # propagate overflow flags by storing biggest TMAX from input |
||||
@ -0,0 +1,214 @@
@@ -0,0 +1,214 @@
|
||||
import kyupy.logic as lg |
||||
|
||||
|
||||
def test_mvarray(): |
||||
|
||||
# instantiation with shape |
||||
|
||||
ary = lg.MVArray(4) |
||||
assert ary.length == 1 |
||||
assert len(ary) == 1 |
||||
assert ary.width == 4 |
||||
|
||||
ary = lg.MVArray((3, 2)) |
||||
assert ary.length == 2 |
||||
assert len(ary) == 2 |
||||
assert ary.width == 3 |
||||
|
||||
# instantiation with single vector |
||||
|
||||
ary = lg.MVArray([1, 0, 1]) |
||||
assert ary.length == 1 |
||||
assert ary.width == 3 |
||||
assert str(ary) == "['101']" |
||||
assert ary[0] == '101' |
||||
|
||||
ary = lg.MVArray("10X-") |
||||
assert ary.length == 1 |
||||
assert ary.width == 4 |
||||
assert str(ary) == "['10X-']" |
||||
assert ary[0] == '10X-' |
||||
|
||||
ary = lg.MVArray("1") |
||||
assert ary.length == 1 |
||||
assert ary.width == 1 |
||||
|
||||
ary = lg.MVArray(["1"]) |
||||
assert ary.length == 1 |
||||
assert ary.width == 1 |
||||
|
||||
# instantiation with multiple vectors |
||||
|
||||
ary = lg.MVArray([[0, 0], [0, 1], [1, 0], [1, 1]]) |
||||
assert ary.length == 4 |
||||
assert ary.width == 2 |
||||
|
||||
ary = lg.MVArray(["000", "001", "110", "---"]) |
||||
assert ary.length == 4 |
||||
assert ary.width == 3 |
||||
assert str(ary) == "['000', '001', '110', '---']" |
||||
assert ary[2] == '110' |
||||
|
||||
# casting to 2-valued logic |
||||
|
||||
ary = lg.MVArray([0, 1, 2, None], m=2) |
||||
assert ary.data[0] == lg.ZERO |
||||
assert ary.data[1] == lg.ONE |
||||
assert ary.data[2] == lg.ZERO |
||||
assert ary.data[3] == lg.ZERO |
||||
|
||||
ary = lg.MVArray("0-X1PRFN", m=2) |
||||
assert ary.data[0] == lg.ZERO |
||||
assert ary.data[1] == lg.ZERO |
||||
assert ary.data[2] == lg.ZERO |
||||
assert ary.data[3] == lg.ONE |
||||
assert ary.data[4] == lg.ZERO |
||||
assert ary.data[5] == lg.ONE |
||||
assert ary.data[6] == lg.ZERO |
||||
assert ary.data[7] == lg.ONE |
||||
|
||||
# casting to 4-valued logic |
||||
|
||||
ary = lg.MVArray([0, 1, 2, None, 'F'], m=4) |
||||
assert ary.data[0] == lg.ZERO |
||||
assert ary.data[1] == lg.ONE |
||||
assert ary.data[2] == lg.UNKNOWN |
||||
assert ary.data[3] == lg.UNASSIGNED |
||||
assert ary.data[4] == lg.ZERO |
||||
|
||||
ary = lg.MVArray("0-X1PRFN", m=4) |
||||
assert ary.data[0] == lg.ZERO |
||||
assert ary.data[1] == lg.UNASSIGNED |
||||
assert ary.data[2] == lg.UNKNOWN |
||||
assert ary.data[3] == lg.ONE |
||||
assert ary.data[4] == lg.ZERO |
||||
assert ary.data[5] == lg.ONE |
||||
assert ary.data[6] == lg.ZERO |
||||
assert ary.data[7] == lg.ONE |
||||
|
||||
# casting to 8-valued logic |
||||
|
||||
ary = lg.MVArray([0, 1, 2, None, 'F'], m=8) |
||||
assert ary.data[0] == lg.ZERO |
||||
assert ary.data[1] == lg.ONE |
||||
assert ary.data[2] == lg.UNKNOWN |
||||
assert ary.data[3] == lg.UNASSIGNED |
||||
assert ary.data[4] == lg.FALL |
||||
|
||||
ary = lg.MVArray("0-X1PRFN", m=8) |
||||
assert ary.data[0] == lg.ZERO |
||||
assert ary.data[1] == lg.UNASSIGNED |
||||
assert ary.data[2] == lg.UNKNOWN |
||||
assert ary.data[3] == lg.ONE |
||||
assert ary.data[4] == lg.PPULSE |
||||
assert ary.data[5] == lg.RISE |
||||
assert ary.data[6] == lg.FALL |
||||
assert ary.data[7] == lg.NPULSE |
||||
|
||||
# copy constructor and casting |
||||
|
||||
ary8 = lg.MVArray(ary, m=8) |
||||
assert ary8.length == 1 |
||||
assert ary8.width == 8 |
||||
assert ary8.data[7] == lg.NPULSE |
||||
|
||||
ary4 = lg.MVArray(ary, m=4) |
||||
assert ary4.data[1] == lg.UNASSIGNED |
||||
assert ary4.data[7] == lg.ONE |
||||
|
||||
ary2 = lg.MVArray(ary, m=2) |
||||
assert ary2.data[1] == lg.ZERO |
||||
assert ary2.data[7] == lg.ONE |
||||
|
||||
|
||||
def test_mv_operations(): |
||||
x1_2v = lg.MVArray("0011", m=2) |
||||
x2_2v = lg.MVArray("0101", m=2) |
||||
x1_4v = lg.MVArray("0000XXXX----1111", m=4) |
||||
x2_4v = lg.MVArray("0X-10X-10X-10X-1", m=4) |
||||
x1_8v = lg.MVArray("00000000XXXXXXXX--------11111111PPPPPPPPRRRRRRRRFFFFFFFFNNNNNNNN", m=8) |
||||
x2_8v = lg.MVArray("0X-1PRFN0X-1PRFN0X-1PRFN0X-1PRFN0X-1PRFN0X-1PRFN0X-1PRFN0X-1PRFN", m=8) |
||||
|
||||
assert lg.mv_not(x1_2v)[0] == '1100' |
||||
assert lg.mv_not(x1_4v)[0] == '1111XXXXXXXX0000' |
||||
assert lg.mv_not(x1_8v)[0] == '11111111XXXXXXXXXXXXXXXX00000000NNNNNNNNFFFFFFFFRRRRRRRRPPPPPPPP' |
||||
|
||||
assert lg.mv_or(x1_2v, x2_2v)[0] == '0111' |
||||
assert lg.mv_or(x1_4v, x2_4v)[0] == '0XX1XXX1XXX11111' |
||||
assert lg.mv_or(x1_8v, x2_8v)[0] == '0XX1PRFNXXX1XXXXXXX1XXXX11111111PXX1PRFNRXX1RRNNFXX1FNFNNXX1NNNN' |
||||
|
||||
assert lg.mv_and(x1_2v, x2_2v)[0] == '0001' |
||||
assert lg.mv_and(x1_4v, x2_4v)[0] == '00000XXX0XXX0XX1' |
||||
assert lg.mv_and(x1_8v, x2_8v)[0] == '000000000XXXXXXX0XXXXXXX0XX1PRFN0XXPPPPP0XXRPRPR0XXFPPFF0XXNPRFN' |
||||
|
||||
assert lg.mv_xor(x1_2v, x2_2v)[0] == '0110' |
||||
assert lg.mv_xor(x1_4v, x2_4v)[0] == '0XX1XXXXXXXX1XX0' |
||||
assert lg.mv_xor(x1_8v, x2_8v)[0] == '0XX1PRFNXXXXXXXXXXXXXXXX1XX0NFRPPXXNPRFNRXXFRPNFFXXRFNPRNXXPNFRP' |
||||
|
||||
|
||||
def test_bparray(): |
||||
|
||||
ary = lg.BPArray(4) |
||||
assert ary.length == 1 |
||||
assert len(ary) == 1 |
||||
assert ary.width == 4 |
||||
|
||||
ary = lg.BPArray((3, 2)) |
||||
assert ary.length == 2 |
||||
assert len(ary) == 2 |
||||
assert ary.width == 3 |
||||
|
||||
assert lg.MVArray(lg.BPArray("01", m=2))[0] == '01' |
||||
assert lg.MVArray(lg.BPArray("0X-1", m=4))[0] == '0X-1' |
||||
assert lg.MVArray(lg.BPArray("0X-1PRFN", m=8))[0] == '0X-1PRFN' |
||||
|
||||
x1_2v = lg.BPArray("0011", m=2) |
||||
x2_2v = lg.BPArray("0101", m=2) |
||||
x1_4v = lg.BPArray("0000XXXX----1111", m=4) |
||||
x2_4v = lg.BPArray("0X-10X-10X-10X-1", m=4) |
||||
x1_8v = lg.BPArray("00000000XXXXXXXX--------11111111PPPPPPPPRRRRRRRRFFFFFFFFNNNNNNNN", m=8) |
||||
x2_8v = lg.BPArray("0X-1PRFN0X-1PRFN0X-1PRFN0X-1PRFN0X-1PRFN0X-1PRFN0X-1PRFN0X-1PRFN", m=8) |
||||
|
||||
out_2v = lg.BPArray((4, 1), m=2) |
||||
out_4v = lg.BPArray((16, 1), m=4) |
||||
out_8v = lg.BPArray((64, 1), m=8) |
||||
|
||||
lg.bp_buf(out_2v.data, x1_2v.data) |
||||
lg.bp_buf(out_4v.data, x1_4v.data) |
||||
lg.bp_buf(out_8v.data, x1_8v.data) |
||||
|
||||
assert lg.MVArray(out_2v)[0] == '0011' |
||||
assert lg.MVArray(out_4v)[0] == '0000XXXXXXXX1111' |
||||
assert lg.MVArray(out_8v)[0] == '00000000XXXXXXXXXXXXXXXX11111111PPPPPPPPRRRRRRRRFFFFFFFFNNNNNNNN' |
||||
|
||||
lg.bp_not(out_2v.data, x1_2v.data) |
||||
lg.bp_not(out_4v.data, x1_4v.data) |
||||
lg.bp_not(out_8v.data, x1_8v.data) |
||||
|
||||
assert lg.MVArray(out_2v)[0] == '1100' |
||||
assert lg.MVArray(out_4v)[0] == '1111XXXXXXXX0000' |
||||
assert lg.MVArray(out_8v)[0] == '11111111XXXXXXXXXXXXXXXX00000000NNNNNNNNFFFFFFFFRRRRRRRRPPPPPPPP' |
||||
|
||||
lg.bp_or(out_2v.data, x1_2v.data, x2_2v.data) |
||||
lg.bp_or(out_4v.data, x1_4v.data, x2_4v.data) |
||||
lg.bp_or(out_8v.data, x1_8v.data, x2_8v.data) |
||||
|
||||
assert lg.MVArray(out_2v)[0] == '0111' |
||||
assert lg.MVArray(out_4v)[0] == '0XX1XXX1XXX11111' |
||||
assert lg.MVArray(out_8v)[0] == '0XX1PRFNXXX1XXXXXXX1XXXX11111111PXX1PRFNRXX1RRNNFXX1FNFNNXX1NNNN' |
||||
|
||||
lg.bp_and(out_2v.data, x1_2v.data, x2_2v.data) |
||||
lg.bp_and(out_4v.data, x1_4v.data, x2_4v.data) |
||||
lg.bp_and(out_8v.data, x1_8v.data, x2_8v.data) |
||||
|
||||
assert lg.MVArray(out_2v)[0] == '0001' |
||||
assert lg.MVArray(out_4v)[0] == '00000XXX0XXX0XX1' |
||||
assert lg.MVArray(out_8v)[0] == '000000000XXXXXXX0XXXXXXX0XX1PRFN0XXPPPPP0XXRPRPR0XXFPPFF0XXNPRFN' |
||||
|
||||
lg.bp_xor(out_2v.data, x1_2v.data, x2_2v.data) |
||||
lg.bp_xor(out_4v.data, x1_4v.data, x2_4v.data) |
||||
lg.bp_xor(out_8v.data, x1_8v.data, x2_8v.data) |
||||
|
||||
assert lg.MVArray(out_2v)[0] == '0110' |
||||
assert lg.MVArray(out_4v)[0] == '0XX1XXXXXXXX1XX0' |
||||
assert lg.MVArray(out_8v)[0] == '0XX1PRFNXXXXXXXXXXXXXXXX1XX0NFRPPXXNPRFNRXXFRPNFFXXRFNPRNXXPNFRP' |
||||
@ -1,161 +1,96 @@
@@ -1,161 +1,96 @@
|
||||
from kyupy.logic_sim import LogicSim |
||||
from kyupy import bench |
||||
from kyupy.packed_vectors import PackedVectors |
||||
from kyupy.logic import MVArray, BPArray |
||||
|
||||
|
||||
def test_vd1(): |
||||
def test_2v(): |
||||
c = bench.parse('input(x, y) output(a, o, n) a=and(x,y) o=or(x,y) n=not(x)') |
||||
s = LogicSim(c, 4) |
||||
s = LogicSim(c, 4, m=2) |
||||
assert len(s.interface) == 5 |
||||
p = PackedVectors(4, len(s.interface)) |
||||
p[0] = '00000' |
||||
p[1] = '01000' |
||||
p[2] = '10000' |
||||
p[3] = '11000' |
||||
s.assign(p) |
||||
mva = MVArray(['00000', '01000', '10000', '11000'], m=2) |
||||
bpa = BPArray(mva) |
||||
s.assign(bpa) |
||||
s.propagate() |
||||
s.capture(p) |
||||
assert p[0] == '00001' |
||||
assert p[1] == '01011' |
||||
assert p[2] == '10010' |
||||
assert p[3] == '11110' |
||||
s.capture(bpa) |
||||
mva = MVArray(bpa) |
||||
assert mva[0] == '00001' |
||||
assert mva[1] == '01011' |
||||
assert mva[2] == '10010' |
||||
assert mva[3] == '11110' |
||||
|
||||
|
||||
def test_vd2(): |
||||
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, 2) |
||||
s = LogicSim(c, 16, m=4) |
||||
assert len(s.interface) == 5 |
||||
p = PackedVectors(16, len(s.interface), 2) |
||||
p[0] = '00000' |
||||
p[1] = '01000' |
||||
p[2] = '0-000' |
||||
p[3] = '0X000' |
||||
p[4] = '10000' |
||||
p[5] = '11000' |
||||
p[6] = '1-000' |
||||
p[7] = '1X000' |
||||
p[8] = '-0000' |
||||
p[9] = '-1000' |
||||
p[10] = '--000' |
||||
p[11] = '-X000' |
||||
p[12] = 'X0000' |
||||
p[13] = 'X1000' |
||||
p[14] = 'X-000' |
||||
p[15] = 'XX000' |
||||
s.assign(p) |
||||
mva = MVArray(['00000', '01000', '0-000', '0X000', |
||||
'10000', '11000', '1-000', '1X000', |
||||
'-0000', '-1000', '--000', '-X000', |
||||
'X0000', 'X1000', 'X-000', 'XX000'], m=4) |
||||
bpa = BPArray(mva) |
||||
s.assign(bpa) |
||||
s.propagate() |
||||
s.capture(p) |
||||
assert p[0] == '00001' |
||||
assert p[1] == '01011' |
||||
assert p[2] == '0-0X1' |
||||
assert p[3] == '0X0X1' |
||||
assert p[4] == '10010' |
||||
assert p[5] == '11110' |
||||
assert p[6] == '1-X10' |
||||
assert p[7] == '1XX10' |
||||
assert p[8] == '-00XX' |
||||
assert p[9] == '-1X1X' |
||||
assert p[10] == '--XXX' |
||||
assert p[11] == '-XXXX' |
||||
assert p[12] == 'X00XX' |
||||
assert p[13] == 'X1X1X' |
||||
assert p[14] == 'X-XXX' |
||||
assert p[15] == 'XXXXX' |
||||
s.capture(bpa) |
||||
mva = MVArray(bpa) |
||||
assert mva[0] == '00001' |
||||
assert mva[1] == '01011' |
||||
assert mva[2] == '0-0X1' |
||||
assert mva[3] == '0X0X1' |
||||
assert mva[4] == '10010' |
||||
assert mva[5] == '11110' |
||||
assert mva[6] == '1-X10' |
||||
assert mva[7] == '1XX10' |
||||
assert mva[8] == '-00XX' |
||||
assert mva[9] == '-1X1X' |
||||
assert mva[10] == '--XXX' |
||||
assert mva[11] == '-XXXX' |
||||
assert mva[12] == 'X00XX' |
||||
assert mva[13] == 'X1X1X' |
||||
assert mva[14] == 'X-XXX' |
||||
assert mva[15] == 'XXXXX' |
||||
|
||||
|
||||
def test_vd3(): |
||||
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, 3) |
||||
s = LogicSim(c, 64, m=8) |
||||
assert len(s.interface) == 6 |
||||
p = PackedVectors(64, len(s.interface), 3) |
||||
p[0] = '000010' |
||||
p[1] = '010111' |
||||
p[2] = '0-0X1X' |
||||
p[3] = '0X0X1X' |
||||
p[4] = '0R0R1R' |
||||
p[5] = '0F0F1F' |
||||
p[6] = '0P0P1P' |
||||
p[7] = '0N0N1N' |
||||
p[8] = '100101' |
||||
p[9] = '111100' |
||||
p[10] = '1-X10X' |
||||
p[11] = '1XX10X' |
||||
p[12] = '1RR10F' |
||||
p[13] = '1FF10R' |
||||
p[14] = '1PP10N' |
||||
p[15] = '1NN10P' |
||||
p[16] = '-00XXX' |
||||
p[17] = '-1X1XX' |
||||
p[18] = '--XXXX' |
||||
p[19] = '-XXXXX' |
||||
p[20] = '-RXXXX' |
||||
p[21] = '-FXXXX' |
||||
p[22] = '-PXXXX' |
||||
p[23] = '-NXXXX' |
||||
p[24] = 'X00XXX' |
||||
p[25] = 'X1X1XX' |
||||
p[26] = 'X-XXXX' |
||||
p[27] = 'XXXXXX' |
||||
p[28] = 'XRXXXX' |
||||
p[29] = 'XFXXXX' |
||||
p[30] = 'XPXXXX' |
||||
p[31] = 'XNXXXX' |
||||
p[32] = 'R00RFR' |
||||
p[33] = 'R1R1FF' |
||||
p[34] = 'R-XXFX' |
||||
p[35] = 'RXXXFX' |
||||
p[36] = 'RRRRFP' |
||||
p[37] = 'RFPNFN' |
||||
p[38] = 'RPPRFR' |
||||
p[39] = 'RNRNFF' |
||||
p[40] = 'F00FRF' |
||||
p[41] = 'F1F1RR' |
||||
p[42] = 'F-XXRX' |
||||
p[43] = 'FXXXRX' |
||||
p[44] = 'FRPNRN' |
||||
p[45] = 'FFFFRP' |
||||
p[46] = 'FPPFRF' |
||||
p[47] = 'FNFNRR' |
||||
p[48] = 'P00PNP' |
||||
p[49] = 'P1P1NN' |
||||
p[50] = 'P-XXNX' |
||||
p[51] = 'PXXXNX' |
||||
p[52] = 'PRPRNR' |
||||
p[53] = 'PFPFNF' |
||||
p[54] = 'PPPPNP' |
||||
p[55] = 'PNPNNN' |
||||
p[56] = 'N00NPN' |
||||
p[57] = 'N1N1PP' |
||||
p[58] = 'N-XXPX' |
||||
p[59] = 'NXXXPX' |
||||
p[60] = 'NRRNPF' |
||||
p[61] = 'NFFNPR' |
||||
p[62] = 'NPPNPN' |
||||
p[63] = 'NNNNPP' |
||||
expect = p.copy() |
||||
s.assign(p) |
||||
mva = MVArray(['000010', '010111', '0-0X1X', '0X0X1X', '0R0R1R', '0F0F1F', '0P0P1P', '0N0N1N', |
||||
'100101', '111100', '1-X10X', '1XX10X', '1RR10F', '1FF10R', '1PP10N', '1NN10P', |
||||
'-00XXX', '-1X1XX', '--XXXX', '-XXXXX', '-RXXXX', '-FXXXX', '-PXXXX', '-NXXXX', |
||||
'X00XXX', 'X1X1XX', 'X-XXXX', 'XXXXXX', 'XRXXXX', 'XFXXXX', 'XPXXXX', 'XNXXXX', |
||||
'R00RFR', 'R1R1FF', 'R-XXFX', 'RXXXFX', 'RRRRFP', 'RFPNFN', 'RPPRFR', 'RNRNFF', |
||||
'F00FRF', 'F1F1RR', 'F-XXRX', 'FXXXRX', 'FRPNRN', 'FFFFRP', 'FPPFRF', 'FNFNRR', |
||||
'P00PNP', 'P1P1NN', 'P-XXNX', 'PXXXNX', 'PRPRNR', 'PFPFNF', 'PPPPNP', 'PNPNNN', |
||||
'N00NPN', 'N1N1PP', 'N-XXPX', 'NXXXPX', 'NRRNPF', 'NFFNPR', 'NPPNPN', 'NNNNPP'], m=8) |
||||
bpa = BPArray(mva) |
||||
s.assign(bpa) |
||||
s.propagate() |
||||
s.capture(p) |
||||
resp_bp = BPArray(bpa) |
||||
s.capture(resp_bp) |
||||
resp = MVArray(resp_bp) |
||||
|
||||
for i in range(64): |
||||
assert p[i] == expect[i] |
||||
assert resp[i] == mva[i] |
||||
|
||||
|
||||
def test_b01(mydir): |
||||
c = bench.parse(mydir / 'b01.bench') |
||||
c = bench.load(mydir / 'b01.bench') |
||||
|
||||
# 2-valued |
||||
s = LogicSim(c, 8) |
||||
s = LogicSim(c, 8, m=2) |
||||
assert len(s.interface) == 9 |
||||
t = PackedVectors(8, len(s.interface)) |
||||
t.randomize() |
||||
s.assign(t) |
||||
mva = MVArray((len(s.interface), 8), m=2) |
||||
# mva.randomize() |
||||
bpa = BPArray(mva) |
||||
s.assign(bpa) |
||||
s.propagate() |
||||
s.capture(t) |
||||
s.capture(bpa) |
||||
|
||||
# 8-valued |
||||
s = LogicSim(c, 8, 3) |
||||
t = PackedVectors(8, len(s.interface), 3) |
||||
t.randomize() |
||||
s.assign(t) |
||||
s = LogicSim(c, 8, m=8) |
||||
mva = MVArray((len(s.interface), 8), m=8) |
||||
# mva.randomize() |
||||
bpa = BPArray(mva) |
||||
s.assign(bpa) |
||||
s.propagate() |
||||
s.capture(t) |
||||
s.capture(bpa) |
||||
|
||||
@ -1,88 +0,0 @@
@@ -1,88 +0,0 @@
|
||||
from kyupy.packed_vectors import PackedVectors |
||||
|
||||
|
||||
def test_basic(): |
||||
ba = PackedVectors(8, 1, 1) |
||||
assert '0\n0\n0\n0\n0\n0\n0\n0' == str(ba) |
||||
ba.set_value(0, 0, 1) |
||||
ba.set_value(1, 0, 'H') |
||||
ba.set_value(2, 0, 'h') |
||||
ba.set_value(3, 0, True) |
||||
ba.set_value(4, 0, 0) |
||||
ba.set_value(5, 0, 'L') |
||||
ba.set_value(6, 0, 'l') |
||||
ba.set_value(7, 0, False) |
||||
assert '1\n1\n1\n1\n0\n0\n0\n0' == str(ba) |
||||
ba.set_value(1, 0, '0') |
||||
ba.set_value(5, 0, '1') |
||||
assert '1\n0\n1\n1\n0\n1\n0\n0' == str(ba) |
||||
ba = PackedVectors(8, 1, 2) |
||||
assert '-\n-\n-\n-\n-\n-\n-\n-' == str(ba) |
||||
ba.set_value(0, 0, 1) |
||||
ba.set_value(7, 0, 0) |
||||
ba.set_value(4, 0, 'X') |
||||
assert '1\n-\n-\n-\nX\n-\n-\n0' == str(ba) |
||||
ba.set_value(4, 0, '-') |
||||
assert '1\n-\n-\n-\n-\n-\n-\n0' == str(ba) |
||||
ba = PackedVectors(8, 2, 2) |
||||
assert '--\n--\n--\n--\n--\n--\n--\n--' == str(ba) |
||||
ba.set_value(0, 0, '1') |
||||
ba.set_value(7, 1, '0') |
||||
ba.set_values(1, 'XX') |
||||
assert '1-\nXX\n--\n--\n--\n--\n--\n-0' == str(ba) |
||||
|
||||
|
||||
def test_8v(): |
||||
ba = PackedVectors(1, 8, 3) |
||||
assert '--------' == str(ba) |
||||
ba.set_values(0, r'-x01^v\/') |
||||
assert r'-X01PNFR' == str(ba) |
||||
ba.set_values(0, '-XLHPNFR') |
||||
assert r'-X01PNFR' == str(ba) |
||||
ba.set_values(0, '-xlhpnfr') |
||||
assert r'-X01PNFR' == str(ba) |
||||
p1 = PackedVectors(1, 8, 1) |
||||
p2 = PackedVectors(1, 8, 1) |
||||
p1.set_values(0, '01010101') |
||||
p2.set_values(0, '00110011') |
||||
p = PackedVectors.from_pair(p1, p2) |
||||
assert r'0FR10FR1' == str(p) |
||||
p1 = PackedVectors(1, 8, 2) |
||||
p2 = PackedVectors(1, 8, 2) |
||||
p1.set_values(0, '0101-X-X') |
||||
p2.set_values(0, '00110011') |
||||
p = PackedVectors.from_pair(p1, p2) |
||||
assert r'0FR1----' == str(p) |
||||
p1.set_values(0, '0101-X-X') |
||||
p2.set_values(0, '-X-X--XX') |
||||
p = PackedVectors.from_pair(p1, p2) |
||||
assert r'--------' == str(p) |
||||
|
||||
|
||||
def test_slicing(): |
||||
lv = PackedVectors(3, 2, 1) |
||||
assert '00\n00\n00' == str(lv) |
||||
lv.set_value(1, 0, '1') |
||||
lv.set_value(1, 1, '1') |
||||
assert '00' == lv[0] |
||||
assert '11' == lv[1] |
||||
assert 3 == len(lv) |
||||
lv2 = lv[1:3] |
||||
assert 2 == len(lv2) |
||||
assert '11' == lv2[0] |
||||
assert '00' == lv2[1] |
||||
|
||||
|
||||
def test_copy(): |
||||
lv1 = PackedVectors(8, 1, 1) |
||||
lv1.set_values_for_position(0, '01010101') |
||||
lv2 = PackedVectors(8, 1, 1) |
||||
lv2.set_values_for_position(0, '00100101') |
||||
diff = lv1.diff(lv2) |
||||
lv3 = lv1.copy(selection_mask=diff) |
||||
assert str(lv3) == '1\n0\n1' |
||||
lv4 = lv1.copy(selection_mask=~diff) |
||||
assert str(lv4) == '0\n0\n1\n0\n1' |
||||
lv5 = lv3 + lv4 |
||||
assert str(lv5) == '1\n0\n1\n0\n0\n1\n0\n1' |
||||
|
||||
Loading…
Reference in new issue