Compare commits

...

5 Commits

  1. 1720
      examples/Introduction.ipynb
  2. 2
      pyproject.toml
  3. 2
      src/kyupy/__init__.py
  4. 2
      src/kyupy/bench.py
  5. 9
      src/kyupy/circuit.py
  6. 27
      src/kyupy/logic_sim.py
  7. 40
      tests/test_logic_sim.py
  8. 6
      tests/test_wave_sim.py

1720
examples/Introduction.ipynb

File diff suppressed because it is too large Load Diff

2
pyproject.toml

@ -9,7 +9,7 @@ readme = "README.rst"
requires_python = ">=3.8" requires_python = ">=3.8"
dependencies = [ dependencies = [
"numpy>=1.17.0", "numpy>=1.17.0",
"lark-parser>=0.8.0", "lark>=1.3.0",
] ]
classifiers = [ classifiers = [
"Development Status :: 3 - Alpha", "Development Status :: 3 - Alpha",

2
src/kyupy/__init__.py

@ -293,8 +293,6 @@ if importlib.util.find_spec('numba') is not None:
try: try:
list(numba.cuda.gpus) list(numba.cuda.gpus)
from numba import cuda from numba import cuda
from numba.core import config
config.CUDA_LOW_OCCUPANCY_WARNINGS = False
except CudaSupportError: except CudaSupportError:
log.warn('Cuda unavailable. Falling back to pure Python.') log.warn('Cuda unavailable. Falling back to pure Python.')
cuda = MockCuda() cuda = MockCuda()

2
src/kyupy/bench.py

@ -21,7 +21,7 @@ class BenchTransformer(Transformer):
def start(self, _): return self.c def start(self, _): return self.c
def parameters(self, args): return [self.c.get_or_add_fork(str(name)) for name in args] def parameters(self, args): return [self.c.get_or_add_fork(str(name)) for name in args if name is not None]
def interface(self, args): self.c.io_nodes.extend(args[0]) def interface(self, args): self.c.io_nodes.extend(args[0])

9
src/kyupy/circuit.py

@ -669,13 +669,16 @@ class Circuit:
ins = '|'.join([f'<i{i}>{i}' for i in range(len(n.ins))]) ins = '|'.join([f'<i{i}>{i}' for i in range(len(n.ins))])
outs = '|'.join([f'<o{i}>{i}' for i in range(len(n.outs))]) outs = '|'.join([f'<o{i}>{i}' for i in range(len(n.outs))])
io = f' [{s_dict[n]}]' if n in s_dict else '' io = f' [{s_dict[n]}]' if n in s_dict else ''
s.node(name=str(n.index), label = f'{{{{{ins}}}|{n.index}{io}\n{n.kind}\n{n.name}|{{{outs}}}}}', shape='record') s.node(name=str(n.index), label = fr'{{{{{ins}}}|{n.index}{io}\n{n.kind}\n{n.name}|{{{outs}}}}}', shape='record')
for l in self.lines: for l in self.lines:
driver, reader = f'{l.driver.index}:o{l.driver_pin}', f'{l.reader.index}:i{l.reader_pin}' driver, reader = f'{l.driver.index}:o{l.driver_pin}', f'{l.reader.index}:i{l.reader_pin}'
if node_level[l.driver] >= node_level[l.reader]: if node_level[l.driver] == node_level[l.reader]:
s.node(f'_{l.index}_')
dot.edge(driver, f'_{l.index}_', style='dotted', label=str(l.index))
dot.edge(f'_{l.index}_', reader, style='dotted', label=str(l.index))
elif node_level[l.driver] > node_level[l.reader]:
dot.edge(driver, reader, style='dotted', label=str(l.index)) dot.edge(driver, reader, style='dotted', label=str(l.index))
pass
else: else:
dot.edge(driver, reader, label=str(l.index)) dot.edge(driver, reader, label=str(l.index))

27
src/kyupy/logic_sim.py

@ -43,6 +43,7 @@ class LogicSim(sim.SimOps):
Access this array to assign new values to the (P)PIs or read values from the (P)POs. Access this array to assign new values to the (P)PIs or read values from the (P)POs.
""" """
self.s[:,:,1,:] = 255 # unassigned self.s[:,:,1,:] = 255 # unassigned
self._full_mask = np.full(self.c.shape[-1], 255, dtype=np.uint8)
def __repr__(self): def __repr__(self):
return f'{{name: "{self.circuit.name}", sims: {self.sims}, m: {self.m}, c_bytes: {eng(self.c.nbytes)}}}' return f'{{name: "{self.circuit.name}", sims: {self.sims}, m: {self.m}, c_bytes: {eng(self.c.nbytes)}}}'
@ -61,20 +62,22 @@ class LogicSim(sim.SimOps):
:param inject_cb: A callback function for manipulating intermediate signal values. :param inject_cb: A callback function for manipulating intermediate signal values.
This function is called with a line and its new logic values (in bit-parallel format) after This function is called with a line and its new logic values (in bit-parallel format) after
evaluation of a node. The callback may manipulate the given values in-place, the simulation evaluation of a node. The callback may manipulate the given values in-place, the simulation
resumes with the manipulated values after the callback returns. resumes with the manipulated values after the callback returns. Specifying this callback
may reduce performance as it disables jit compilation.
:type inject_cb: ``f(Line, ndarray)`` :type inject_cb: ``f(Line, ndarray)``
""" """
t0 = self.c_locs[self.tmp_idx] fault_line = int(fault_line)
t1 = self.c_locs[self.tmp2_idx]
if self.m == 2:
if inject_cb is None:
if fault_mask is None: if fault_mask is None:
fault_mask = np.full(self.c.shape[-1], 255, dtype=np.uint8) fault_mask = self._full_mask # default: full mask
else: else:
if len(fault_mask) < self.c.shape[-1]: if len(fault_mask) < self.c.shape[-1]: # pad mask with 0's if necessary
fault_mask2 = np.full(self.c.shape[-1], 0, dtype=np.uint8) fault_mask2 = np.full(self.c.shape[-1], 0, dtype=np.uint8)
fault_mask2[:len(fault_mask)] = fault_mask fault_mask2[:len(fault_mask)] = fault_mask
fault_mask = fault_mask2 fault_mask = fault_mask2
t0 = self.c_locs[self.tmp_idx]
t1 = self.c_locs[self.tmp2_idx]
if self.m == 2:
if inject_cb is None:
_prop_cpu(self.ops, self.c_locs, self.c, int(fault_line), fault_mask, int(fault_model)) _prop_cpu(self.ops, self.c_locs, self.c, int(fault_line), fault_mask, int(fault_model))
else: else:
for op, o0l, i0l, i1l, i2l, i3l in self.ops[:,:6]: for op, o0l, i0l, i1l, i2l, i3l in self.ops[:,:6]:
@ -190,6 +193,15 @@ class LogicSim(sim.SimOps):
logic.bp4v_or(self.c[o0], self.c[t0], self.c[t1]) logic.bp4v_or(self.c[o0], self.c[t0], self.c[t1])
else: print(f'unknown op {op}') else: print(f'unknown op {op}')
if inject_cb is not None: inject_cb(o0l, self.c[o0]) if inject_cb is not None: inject_cb(o0l, self.c[o0])
if fault_line >= 0 and o0l == fault_line:
if fault_model == 0:
self.c[o0] = self.c[o0] & ~fault_mask[np.newaxis]
elif fault_model == 1:
self.c[o0] = self.c[o0] | fault_mask[np.newaxis]
else:
self.c[t0, 0] = ~(self.c[o0, 0] & self.c[o0, 1] & fault_mask)
self.c[o0, 1] = ~self.c[o0, 0] & ~self.c[o0, 1] & fault_mask
self.c[o0, 0] = self.c[t0, 0]
else: else:
for op, o0l, i0l, i1l, i2l, i3l 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 (o0l, i0l, i1l, i2l, i3l)] o0, i0, i1, i2, i3 = [self.c_locs[x] for x in (o0l, i0l, i1l, i2l, i3l)]
@ -343,7 +355,6 @@ def _prop_cpu(ops, c_locs, c, fault_line, fault_mask, fault_model):
elif op == sim.MUX21: c[o0] = (c[i0] & ~c[i2]) | (c[i1] & c[i2]) elif op == sim.MUX21: c[o0] = (c[i0] & ~c[i2]) | (c[i1] & c[i2])
else: print(f'unknown op {op}') else: print(f'unknown op {op}')
if fault_line >= 0 and o0l == fault_line: if fault_line >= 0 and o0l == fault_line:
#n = len(fault_mask)
if fault_model == 0: if fault_model == 0:
c[o0] = c[o0] & ~fault_mask c[o0] = c[o0] & ~fault_mask
elif fault_model == 1: elif fault_model == 1:

40
tests/test_logic_sim.py

@ -75,7 +75,7 @@ def test_2v():
def test_4v(): def test_4v():
c = bench.parse('input(x, y) output(a, o, n) a=and(x,y) o=or(x,y) n=not(x)') c = bench.parse('input(x, y) output(a, o, n) a=and(x,y) o=or(x,y) n=not(x)')
s = LogicSim(c, 16, m=8) # FIXME: m=4 s = LogicSim(c, 16, m=4)
assert s.s_len == 5 assert s.s_len == 5
bpa = bparray( bpa = bparray(
'00---', '01---', '0----', '0X---', '00---', '01---', '0----', '0X---',
@ -93,6 +93,44 @@ def test_4v():
'--0XX', '--X1X', '--XXX', '--XXX', '--0XX', '--X1X', '--XXX', '--XXX',
'--0XX', '--X1X', '--XXX', '--XXX')) '--0XX', '--X1X', '--XXX', '--XXX'))
def test_4v_fault():
c = bench.parse('input(x, y) output(a) a=and(x,y)')
s = LogicSim(c, 16, m=4)
assert s.s_len == 3
bpa = bparray(
'00-', '01-', '0--', '0X-',
'10-', '11-', '1--', '1X-',
'-0-', '-1-', '---', '-X-',
'X0-', 'X1-', 'X--', 'XX-')
s.s[0] = bpa
s.s_to_c()
s.c_prop()
s.c_to_s()
mva = bp_to_mv(s.s[1])
assert_equal_shape_and_contents(mva, mvarray(
'--0', '--0', '--0', '--0',
'--0', '--1', '--X', '--X',
'--0', '--X', '--X', '--X',
'--0', '--X', '--X', '--X'))
fault_line = s.circuit.cells['a'].ins[0]
s.s_to_c()
s.c_prop(fault_line=fault_line, fault_model=1)
s.c_to_s()
mva = bp_to_mv(s.s[1])
assert_equal_shape_and_contents(mva, mvarray(
'--0', '--1', '--X', '--X',
'--0', '--1', '--X', '--X',
'--0', '--1', '--X', '--X',
'--0', '--1', '--X', '--X'))
s.s_to_c()
s.c_prop(fault_line=fault_line, fault_model=0)
s.c_to_s()
mva = bp_to_mv(s.s[1])
assert_equal_shape_and_contents(mva, mvarray(
'--0', '--0', '--0', '--0',
'--0', '--0', '--0', '--0',
'--0', '--0', '--0', '--0',
'--0', '--0', '--0', '--0'))
def test_6v(): 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)') 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)')

6
tests/test_wave_sim.py

@ -25,10 +25,11 @@ def test_xnor2_delays():
delays[0, 1, 1, 1] = 0.036 # B fall -> Z fall delays[0, 1, 1, 1] = 0.036 # B fall -> Z fall
simctl_int = np.asarray([0], dtype=np.int32) simctl_int = np.asarray([0], dtype=np.int32)
simctl_float = np.asarray([1], dtype=np.float32)
def wave_assert(inputs, output): def wave_assert(inputs, output):
for i, a in zip(inputs, c.reshape(-1,16)): a[:len(i)] = i for i, a in zip(inputs, c.reshape(-1,16)): a[:len(i)] = i
wave_eval_cpu(op, c, c_locs, c_caps, ebuf, 0, delays, simctl_int, 0, 0) wave_eval_cpu(op, c, c_locs, c_caps, ebuf, 0, delays, simctl_int, simctl_float, 0, 0)
for i, v in enumerate(output): np.testing.assert_allclose(c.reshape(-1,16)[2,i], v) for i, v in enumerate(output): np.testing.assert_allclose(c.reshape(-1,16)[2,i], v)
wave_assert([[TMIN,TMAX],[TMIN,TMAX]], [TMIN,TMAX]) # XNOR(1,1) => 1 wave_assert([[TMIN,TMAX],[TMIN,TMAX]], [TMIN,TMAX]) # XNOR(1,1) => 1
@ -63,10 +64,11 @@ def test_nand_delays():
delays[0, 3, :, 1] = 0.8 delays[0, 3, :, 1] = 0.8
simctl_int = np.asarray([0], dtype=np.int32) simctl_int = np.asarray([0], dtype=np.int32)
simctl_float = np.asarray([1], dtype=np.float32)
def wave_assert(inputs, output): def wave_assert(inputs, output):
for i, a in zip(inputs, c.reshape(-1,16)): a[:len(i)] = i for i, a in zip(inputs, c.reshape(-1,16)): a[:len(i)] = i
wave_eval_cpu(op, c, c_locs, c_caps, ebuf, 0, delays, simctl_int, 0, 0) wave_eval_cpu(op, c, c_locs, c_caps, ebuf, 0, delays, simctl_int, simctl_float, 0, 0)
for i, v in enumerate(output): np.testing.assert_allclose(c.reshape(-1,16)[4,i], v) 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 wave_assert([[TMAX,TMAX],[TMAX,TMAX],[TMIN,TMAX],[TMIN,TMAX]], [TMIN,TMAX]) # NAND(0,0,1,1) => 1

Loading…
Cancel
Save