fix(cosim): CS-gate QSPI MISO injection for shared-bus arbitration

.gitignore

@@ -28,6 +28,8 @@ __pycache__/
 !tests/xprop_cosim/expected/*.vcd
 # Stage C flash cosim golden (mcu_soc; CpuBackend == Metal, #105).
 !tests/mcu_soc/expected/*.vcd
+# Shared-bus QSPI arbitration golden (CpuBackend == Metal/CUDA/HIP).
+!tests/qspi_shared_bus/expected/*.vcd
 *.tar.gz
 debug_trace.csv
 *_watchlist.json

csrc/kernel_v1.metal

@@ -976,6 +976,11 @@ kernel void gpu_apply_flash_din(
 
     for (uint f = 0; f < params.n_flashes && f < MAX_QSPI_MEMS; f++) {
         constant FlashDinParams& fp = params.flashes[f];
+        // Gate on selection: a deselected instance (prev_csn high) presents
+        // high-Z and must not drive, else on a shared SIO bus (instances sharing
+        // d_in_pos) it clobbers the selected driver. Mirrors
+        // CpuBackend::apply_flash_din_gated.
+        if (flash_state[f].prev_csn != 0u) continue;
         uchar d_i = flash_state[f].d_i;
         u32 xmask_off = fp.xmask_state_offset;
 

csrc/kernel_v1_impl.cuh

@@ -1325,6 +1325,10 @@ __global__ void gpu_apply_flash_din(
 
   for (u32 f = 0; f < params->n_flashes && f < MAX_QSPI_MEMS; f++) {
     const FlashDinParams &fp = params->flashes[f];
+    // Gate on selection: a deselected instance (prev_csn high) presents high-Z
+    // and must not drive, else on a shared SIO bus (instances sharing d_in_pos)
+    // it clobbers the selected driver. Mirrors CpuBackend::apply_flash_din_gated.
+    if (flash_state[f].prev_csn != 0u) continue;
     u8 d_i = flash_state[f].d_i;
     u32 xmask_off = fp.xmask_state_offset;
 

docs/adr/0013-plural-peripheral-configs.md

@@ -15,6 +15,21 @@ refactors; the conventions it establishes are already followed.
 > receives `&backend.state()[state_size..]` (`cosim/mod.rs`). The original
 > text below is retained for the record.
 
+> **Amendment (2026-07-09):** Plural QSPI memories may now share one SCK/SIO
+> bus, selected by distinct CS lines (config: same `clk_gpio`/`d0_gpio`,
+> distinct `csn_gpio`). This requires CS-gated MISO injection: a deselected
+> instance (`prev_csn` high) presents high-Z and must not drive. Previously
+> `gpu_apply_flash_din` (and the CpuBackend mirror) wrote every instance's
+> `d_i` to its `d_in_pos` unconditionally, so on a shared bus the
+> last-iterated (typically deselected) memory clobbered the selected driver.
+> The gate lives in `gpu_apply_flash_din` (Metal + the shared CUDA/HIP
+> `kernel_v1_impl.cuh`) and `CpuBackend::apply_flash_din_gated`; regression:
+> `tests/qspi_shared_bus/` (golden + content, `all`/`qspi` scopes) plus the
+> `flash_din_tests` unit tests. Independent-pin plurality (the original Stage
+> B/C design, `tests/multi_mem_cosim/`) is unaffected — a deselected flash on
+> its own pins was always ignored by the design; the gate only changes the
+> now-shared case.
+
 ## Context
 
 Jacquard's cosim mode runs reactive peripheral models alongside the

scripts/ci/cosim_cpu_check.sh

@@ -150,6 +150,24 @@ if [ "$SCOPE" = all ] || [ "$SCOPE" = qspi ]; then
     else
         fail "qspi_psram_content"
     fi
+
+    # qspi_shared_bus: two flashes on ONE shared SCK/SIO bus, distinct CS. The
+    # master reads addr 0 from each in turn; a deselected flash must present
+    # high-Z on the shared MISO (CS-gated din injection) or it clobbers the
+    # other's read. Regression for the shared-bus arbitration fix. Golden
+    # captured on CpuBackend, byte-identical to the Metal/CUDA/HIP kernels.
+    # See tests/qspi_shared_bus/README.md.
+    run qspi_shared_bus "$BIN" cosim tests/qspi_shared_bus/shared_bus_dut_synth.gv \
+        --config tests/qspi_shared_bus/sim_config.json --top-module shared_bus_dut \
+        --max-clock-edges 2000 --output-vcd "$OUT/qspi_shared_bus.vcd" || true
+    check qspi_shared_bus "$OUT/qspi_shared_bus.vcd" tests/qspi_shared_bus/expected/qspi_shared_bus.vcd
+    # Content assertion (each flash returned its own byte), independent of the diff.
+    total=$((total + 1))
+    if python3 tests/qspi_shared_bus/check.py "$OUT/qspi_shared_bus.vcd" >/dev/null 2>&1; then
+        pass "qspi_shared_bus_content"
+    else
+        fail "qspi_shared_bus_content"
+    fi
 fi
 
 # --- flash fixture (mcu_soc) — `flash` scope only ------------------------

src/sim/cosim/mod.rs

@@ -4534,6 +4534,51 @@ struct CpuFlashInstance {
     prev_d_out: UnsafeCell<u8>,
 }
 
+/// Inject one QSPI instance's MISO nibble into its `d_in` bit positions —
+/// but only when the instance is **selected** (`csn` low). A deselected
+/// instance (csn high) presents high-Z and must not drive: otherwise, on a
+/// shared SIO bus where instances share `d_in_pos`, a deselected memory
+/// clobbers the selected driver's data (the raw per-instance loop is
+/// last-writer-wins). `xmask_off == 0` means xprop is disabled.
+///
+/// This is the CS gate mirrored by `gpu_apply_flash_din` (Metal / CUDA / HIP),
+/// which skips a flash whose `FlashState.prev_csn` is high. Kept as a free fn
+/// so the shared-bus arbitration is unit-testable without a GPU or a netlist.
+///
+/// `csn` is the **delayed** (`prev_csn`) chip-select, not the live output-slot
+/// value: `d_i` was produced by the setup-delay dual-step using that same
+/// `prev_csn`, so gating injection on it keeps the driven value and the
+/// selection decision in phase. Gating on live `csn` would skew them by one
+/// edge — the very skew the delay model exists to avoid.
+fn apply_flash_din_gated(
+    state: &mut [u32],
+    xmask_off: usize,
+    d_i: u8,
+    csn: u32,
+    d_in_pos: &[u32; 4],
+) {
+    if csn != 0 {
+        return; // deselected → high-Z, don't drive the (possibly shared) line
+    }
+    for i in 0..4usize {
+        let pos = d_in_pos[i];
+        if pos == 0xFFFFFFFF {
+            continue;
+        }
+        let word_idx = (pos >> 5) as usize;
+        let bit_mask = 1u32 << (pos & 31);
+        if (d_i >> i) & 1 != 0 {
+            state[word_idx] |= bit_mask;
+        } else {
+            state[word_idx] &= !bit_mask;
+        }
+        // Driven ⇒ known: clear the X-mask so a shared MISO bit isn't X-seeded.
+        if xmask_off != 0 {
+            state[xmask_off + word_idx] &= !bit_mask;
+        }
+    }
+}
+
 struct CpuBackend {
     /// Full design state `[input_state | output_state]`, `2 × state_size` words.
     state: UnsafeCell<Vec<u32>>,
@@ -5007,26 +5052,14 @@ impl CosimBackend for CpuBackend {
             // ── CPU apply_flash_din (mirror gpu_apply_flash_din, shader:904) ─
             // Inject each memory's current MISO nibble into ITS input-state
             // d_in bits BEFORE simulate, clearing their X-mask (driven ⇒ known,
-            // #95 ph3). Every QSPI instance drives its own lanes independently.
+            // #95 ph3). Gated on selection (prev_csn low) so a deselected memory
+            // presents high-Z and can't clobber a shared SIO bus — see
+            // `apply_flash_din_gated`.
             for inst in &self.flashes {
                 // SAFETY: sequential dispatch — no concurrent borrow.
                 let d_i = unsafe { *inst.d_i.get() };
-                for i in 0..4usize {
-                    let pos = inst.d_in_pos[i];
-                    if pos == 0xFFFFFFFF {
-                        continue;
-                    }
-                    let word_idx = (pos >> 5) as usize;
-                    let bit_mask = 1u32 << (pos & 31);
-                    if (d_i >> i) & 1 != 0 {
-                        state[word_idx] |= bit_mask;
-                    } else {
-                        state[word_idx] &= !bit_mask;
-                    }
-                    if xmask_off != 0 {
-                        state[xmask_off + word_idx] &= !bit_mask;
-                    }
-                }
+                let csn = unsafe { *inst.prev_csn.get() };
+                apply_flash_din_gated(state, xmask_off, d_i, csn, &inst.d_in_pos);
             }
 
             // ── Simulate every partition for this edge ─────────────────────
@@ -5297,3 +5330,68 @@ pub use cuda::run_cosim_cuda;
 mod hip;
 #[cfg(feature = "hip")]
 pub use hip::run_cosim_hip;
+
+#[cfg(test)]
+mod flash_din_tests {
+    use super::apply_flash_din_gated;
+    use crate::sim::models::read_bit;
+
+    #[test]
+    fn selected_instance_drives_its_lane() {
+        let mut state = vec![0u32; 8];
+        // Instance selected (csn=0), d_i = 0b0010 → lane 1 high. d0 at pos 33.
+        apply_flash_din_gated(&mut state, 0, 0b0010, 0, &[33, 34, 35, 36]);
+        assert_eq!(read_bit(&state, 34), 1); // lane 1 driven high
+        assert_eq!(read_bit(&state, 33), 0); // lane 0 low
+    }
+
+    #[test]
+    fn deselected_instance_does_not_drive() {
+        let mut state = vec![0u32; 8];
+        // Pre-set the shared lane high, then a DESELECTED instance (csn=1)
+        // with d_i=0 must NOT pull it low.
+        state[1] |= 1 << (34 & 31); // pos 34 = word 1, bit 2
+        apply_flash_din_gated(&mut state, 0, 0b0000, 1, &[33, 34, 35, 36]);
+        assert_eq!(read_bit(&state, 34), 1); // untouched — high-Z
+    }
+
+    #[test]
+    fn shared_bus_selected_wins_regardless_of_order() {
+        // Two instances share the SAME d_in positions (shared SIO bus).
+        // A is SELECTED with data (lane1=1); B is DESELECTED with stale idle
+        // (all lanes high). The historical bug was last-writer-wins, so assert
+        // both application orders leave the shared line reflecting A, not B.
+        let shared = [33u32, 34, 35, 36];
+
+        // Order 1: selected A first, then deselected B.
+        let mut s1 = vec![0u32; 8];
+        apply_flash_din_gated(&mut s1, 0, 0b0010, 0, &shared); // A selected
+        apply_flash_din_gated(&mut s1, 0, 0b1111, 1, &shared); // B deselected
+        assert_eq!(read_bit(&s1, 34), 1);
+        assert_eq!(read_bit(&s1, 33), 0);
+        assert_eq!(read_bit(&s1, 36), 0); // B's idle-high did NOT leak onto lane 3
+
+        // Order 2: deselected B first, then selected A. Same result.
+        let mut s2 = vec![0u32; 8];
+        apply_flash_din_gated(&mut s2, 0, 0b1111, 1, &shared); // B deselected
+        apply_flash_din_gated(&mut s2, 0, 0b0010, 0, &shared); // A selected
+        assert_eq!(read_bit(&s2, 34), 1);
+        assert_eq!(read_bit(&s2, 33), 0);
+        assert_eq!(read_bit(&s2, 36), 0);
+    }
+
+    #[test]
+    fn xmask_cleared_only_when_driven() {
+        // xprop enabled: selected drive clears the x-mask bit; deselected leaves
+        // it untouched. xmask_off = 4 words in.
+        let mut state = vec![0u32; 8];
+        state[4 + 1] |= 1 << (34 & 31); // x-mask bit set for pos 34
+        apply_flash_din_gated(&mut state, 4, 0b0010, 0, &[33, 34, 35, 36]);
+        assert_eq!((state[4 + 1] >> (34 & 31)) & 1, 0); // cleared (driven known)
+
+        let mut state2 = vec![0u32; 8];
+        state2[4 + 1] |= 1 << (34 & 31);
+        apply_flash_din_gated(&mut state2, 4, 0b0010, 1, &[33, 34, 35, 36]);
+        assert_eq!((state2[4 + 1] >> (34 & 31)) & 1, 1); // deselected → untouched
+    }
+}

src/testbench.rs

@@ -333,8 +333,9 @@ impl TestbenchConfig {
     ///
     /// Merges the legacy singular `flash` (prepended, so it is instance 0) with
     /// the plural `qspi_memory` list, mirroring `effective_uarts()`. This is the
-    /// single surface every backend consumes: the CPU backend steps every entry
-    /// independently; the GPU backends (Stage A) require `len() <= 1`.
+    /// single surface every backend consumes: every backend steps each entry
+    /// independently, up to `MAX_QSPI_MEMS` (4) instances. Instances may share
+    /// one SCK/SIO bus with distinct CS (see ADR 0013's 2026-07-09 amendment).
     pub fn effective_qspi_memory(&self) -> Vec<QspiMemoryConfig> {
         let mut out = self.qspi_memory.clone();
         if let Some(ref f) = self.flash {

tests/qspi_shared_bus/README.md

@@ -0,0 +1,40 @@
+# qspi_shared_bus — shared SCK/SIO bus, distinct CS
+
+Regression for the CS-gated MISO arbitration on a **shared** QSPI bus (ADR 0013).
+
+Two flash instances share one SCK + SIO bus and are selected by distinct chip
+selects. In `sim_config.json` both memories use the same `clk_gpio` (2) and
+`d0_gpio` (5) but distinct `csn_gpio` (3 vs 4). The DUT (`shared_bus_dut.v`) is
+a single SPI master that reads address 0 from flash A, then from flash B, over
+the shared bus.
+
+The point: a **deselected** flash must present high-Z on the shared MISO. If it
+drives (the pre-fix behaviour), it clobbers the selected flash's read data —
+`gpu_apply_flash_din` / `CpuBackend::apply_flash_din_gated` iterate every
+instance, so the last-iterated (often deselected) memory wins. The fix gates the
+MISO injection on selection (`prev_csn` low). Flash A holds `0xA1`, flash B
+holds `0xB2`; distinct reads prove no clobbering.
+
+## Pass criterion
+
+Byte-exact golden VCD diff (`expected/qspi_shared_bus.vcd`, captured on
+CpuBackend, byte-identical to the Metal/CUDA/HIP kernels) plus a content
+assertion (`check.py`): `fa == 0xA1`, `fb == 0xB2`, `all_match == 1`.
+
+Runs in the `all` and `qspi` scopes of `scripts/ci/cosim_cpu_check.sh`, so it
+gates CpuBackend (Linux) and the CUDA/HIP/Metal GPU suites alike.
+
+## Regenerating
+
+```bash
+# Netlist (system Yosys):
+cd tests/qspi_shared_bus && yosys -q -s synth.tcl
+
+# Golden VCD (any backend — output is byte-identical across backends):
+target/release/jacquard cosim tests/qspi_shared_bus/shared_bus_dut_synth.gv \
+    --config tests/qspi_shared_bus/sim_config.json --top-module shared_bus_dut \
+    --max-clock-edges 2000 --output-vcd tests/qspi_shared_bus/expected/qspi_shared_bus.vcd
+```
+
+Firmware (`fw/flashA.bin` = `0xA1…`, `fw/flashB.bin` = `0xB2…`) is 16 bytes each;
+only byte 0 is read.

tests/qspi_shared_bus/check.py

@@ -0,0 +1,70 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: Apache-2.0
+"""Pass criterion for the shared-bus QSPI cosim test.
+
+Two flashes share one SCK + SIO bus, selected by distinct CS lines. The master
+reads address 0 from each in turn. Correct CS-gated MISO arbitration returns
+each flash's own byte; the pre-fix bug lets the deselected flash drive the
+shared MISO, corrupting the second read.
+
+  fa == 0xA1   (flash A, selected during the first read)
+  fb == 0xB2   (flash B, selected during the second read)
+  all_match == 1, done == 1
+
+Usage: check.py <output.vcd>
+"""
+import re
+import sys
+
+EXPECT = {"fa": 0xA1, "fb": 0xB2}
+
+
+def parse_final(path):
+    id2name = {}
+    for line in open(path):
+        m = re.match(r"\$var \w+ \d+ (\S+) (.+?) \$end", line)
+        if m:
+            id2name[m.group(1)] = m.group(2).strip()
+    last = {}
+    for line in open(path):
+        line = line.rstrip("\n")
+        if not line:
+            continue
+        if line[0] in "01xz":
+            last[id2name.get(line[1:], line[1:])] = line[0]
+        elif line[0] == "b":
+            m = re.match(r"b([01xz]+) (\S+)", line)
+            if m:
+                last[id2name.get(m.group(2), m.group(2))] = m.group(1)
+    return last
+
+
+def busval(last, prefix, w=8):
+    bits = "".join(last.get(f"{prefix}[{b}]", "x") for b in range(w - 1, -1, -1))
+    return int(bits, 2) if set(bits) <= set("01") else None
+
+
+def main():
+    last = parse_final(sys.argv[1])
+    ok = True
+    for name, exp in EXPECT.items():
+        got = busval(last, name)
+        status = "OK" if got == exp else "MISMATCH"
+        if got != exp:
+            ok = False
+        gs = f"0x{got:02X}" if got is not None else "x/undriven"
+        print(f"  {name}: got {gs}, expect 0x{exp:02X}  [{status}]")
+    for scalar in ("done", "all_match"):
+        got = last.get(scalar, "?")
+        if got != "1":
+            ok = False
+        print(f"  {scalar}: {got}  [{'OK' if got == '1' else 'FAIL'}]")
+    if not ok:
+        print("FAIL: shared-bus MISO arbitration wrong "
+              "(deselected flash clobbered the shared line)")
+        sys.exit(1)
+    print("PASS: shared SCK/SIO + distinct CS — each flash returned its own byte")
+
+
+if __name__ == "__main__":
+    main()