fix cuda ppo_to_ppi

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

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