From 4e2022291eee642fc6d8cada0ddadff755c7f9f2 Mon Sep 17 00:00:00 2001 From: Stefan Holst Date: Wed, 15 Mar 2023 15:15:22 +0900 Subject: [PATCH] fix cuda ppo_to_ppi --- src/kyupy/logic_sim.py | 11 ++++++----- src/kyupy/wave_sim.py | 4 ++-- 2 files changed, 8 insertions(+), 7 deletions(-) diff --git a/src/kyupy/logic_sim.py b/src/kyupy/logic_sim.py index 422cc91..8eace75 100644 --- a/src/kyupy/logic_sim.py +++ b/src/kyupy/logic_sim.py @@ -140,7 +140,7 @@ class LogicSim(sim.SimOps): self.s[0, self.ppio_s_locs, 0] = self.s[1, self.ppio_s_locs, 0] # final value is newly captured final value self.s[0, self.ppio_s_locs, 2] = self.s[0, self.ppio_s_locs, 0] ^ self.s[0, self.ppio_s_locs, 1] # TODO: not correct for X, - - def cycle(self, inject_cb=None): + def cycle(self, cycles=1, inject_cb=None): """Assigns the given state, propagates it and captures the new state. :param state: A bit-parallel array in a compatible shape holding the current circuit state. @@ -150,10 +150,11 @@ class LogicSim(sim.SimOps): :param inject_cb: A callback function for manipulating intermediate signal values. See :py:func:`propagate`. :returns: The given state object. """ - self.s_to_c() - self.c_prop(inject_cb) - self.c_to_s() - self.s_ppo_to_ppi() + for _ in range(cycles): + self.s_to_c() + self.c_prop(inject_cb) + self.c_to_s() + self.s_ppo_to_ppi() @numba.njit diff --git a/src/kyupy/wave_sim.py b/src/kyupy/wave_sim.py index 61054fc..0796968 100644 --- a/src/kyupy/wave_sim.py +++ b/src/kyupy/wave_sim.py @@ -605,8 +605,8 @@ def wave_capture_gpu(c, s, c_locs, c_caps, ppo_offset, time, s_sqrt2, seed): @cuda.jit() def ppo_to_ppi_gpu(s, c_locs, time, ppi_offset, ppo_offset): x, y = cuda.grid(2) - if y >= s.shape[0]: return - if x >= s.shape[1]: return + if y >= s.shape[1]: return + if x >= s.shape[2]: return if c_locs[ppi_offset + y] < 0: return if c_locs[ppo_offset + y] < 0: return