Jacquard Documentation

Welcome to the documentation for Jacquard, a GPU-accelerated RTL logic simulator.

Use the sidebar to navigate between topics, or start with the Getting Started guide — it runs four bundled designs in a few seconds with no synthesis. To prepare and run your own RTL, see the Synthesis Flow; for UVM/cocotb/SVA questions, Testbench Interop.

Documents

Project Scope & Planning

Start here if you’re considering a feature contribution or want to understand Jacquard’s overall direction.

• Project Scope & Guarantees: Top-level contract — what Jacquard is for, what it isn’t, licensing and architecture constraints, stability tiers.
• Why Jacquard: Honest positioning vs. STA tools and event-driven simulators; what’s unique, what isn’t, and what output interface would let users extract the value.
• Timing Correctness: Scoped requirements for timing accuracy, validation, and the forthcoming timing IR.
• Timing Model Extensions: Pre-spike design notes for δ(T) dynamic delay, clock-tree skew, and wire delay at scale. Formalised in Decision 0007.
• Post-Phase-0 Roadmap: Sequencing of Phase 1+ work covering structured timing output (Decision 0008) and timing model fidelity (Decision 0007). (OpenTimer integration was originally Phase 1’s centrepiece; Decision 0003 was Superseded by the spike — OpenSTA out of process is now the sole STA path per Decision 0001.)
• Decisions: Design decisions and their rationale (numbered, per-decision). See the index for status and how the decisions relate.
• Implementation Plans: Phased implementation plans with entry and exit criteria. See the index for status and reading order.
• Spikes: Time-boxed experiments and their outcomes.

Core Documentation

• Simulation engine: Detailed explanation of Jacquard’s internal architecture

  • Pipeline stages (NetlistDB → AIG → StagedAIG → Partitions → FlattenedScript → GPU)
  • Data structures and representations
  • VCD input/output format requirements
  • Assertion and display support infrastructure
  • Performance characteristics
  • Known issues and limitations

• Timing Simulation: CPU-based timing simulation with Liberty/SDF delays

• Timing Violations: GPU-side setup/hold violation detection

Troubleshooting Guides

• Troubleshooting VCD: Debugging VCD input issues
  • VCD hierarchy requirements
  • Signal naming and matching
  • Solutions for flat VCD generation
  • Diagnostic checklist
  • Working examples

Quick Reference

VCD Input Requirements (Critical!)

Jacquard expects VCD signals at absolute top-level (no module hierarchy):

// ✓ Correct testbench
initial begin
    $dumpfile("output.vcd");
    $dumpvars(1, clk, reset, din, dout);  // Depth 1, explicit signals
end

// ✗ Incorrect testbench
initial begin
    $dumpfile("output.vcd");
    $dumpvars(0, testbench);  // Dumps entire hierarchy
end

Debug Commands

# Enable debug logging
RUST_LOG=debug cargo run -r --features metal --bin jacquard -- sim <args>

# Verify with CPU simulation
cargo run -r --features metal --bin jacquard -- sim <args> --check-with-cpu

# Check VCD structure
grep '\$scope\|\$var' input.vcd | head -20

Cosim (reactive peripherals)

jacquard cosim runs GPU-resident peripheral models (SPI flash, UART, Wishbone) alongside the design so inputs can react to outputs cycle-by-cycle. It runs on Metal, CUDA, and HIP (plus a CPU fallback).

# Drive the design from a JSON testbench config; write an output VCD
cargo run -r --features metal --bin jacquard -- cosim \
    design.v --config sim_config.json --output-vcd out.vcd

Key Statistics

When running Jacquard, look for these diagnostic outputs:

netlist has X pins, Y aig pins, Z and gates        # AIG complexity
current: N endpoints, try M parts                  # Partition count
Built script for B blocks, reg/io state size S     # Final script
WARN (GATESIM_VCDI_MISSING_PI) ...                 # VCD issues!

Investigation Methodology

This documentation was created through systematic investigation of Jacquard’s behavior:

1. Source Code Analysis: Examined src/aig.rs, src/flatten.rs, src/staging.rs
2. Debug Tracing: Used RUST_LOG=debug to capture internal state
3. Test Case Development: Created minimal reproducible examples
4. Comparative Testing: Compared Jacquard vs iverilog outputs
5. Third-Party Validation: Tested with real-world examples (sva-playground)

Known Issues

Tracked live on GitHub — see the open issues and the priority:high label. The long-standing ones:

• VCD hierarchy mismatch — Jacquard expects a flat top-level VCD; most testbenches emit hierarchical ones. Workaround: --input-vcd-scope (see Troubleshooting VCD). Tracking: #142.
• Complex FSM simulation — some FSM designs (e.g. safe.v) don’t simulate correctly; under investigation. Tracking: #143.
• Format-string preservation — Yosys may drop gem_format attributes, so $display messages show placeholders. This is an upstream Yosys limitation; the workaround is to extract format strings from the pre-synthesis JSON.

Contributing

When adding documentation:

1. Be specific: Include actual commands, file paths, code snippets
2. Show examples: Both working and non-working cases
3. Link related docs: Cross-reference other documentation files
4. Date updates: Update version and date at bottom of documents
5. Test instructions: Verify all commands actually work

Future Documentation Needs

Dedicated guides not yet written (coverage today is scattered across decisions and reference docs):

• Performance tuning guide (choosing NUM_BLOCKS, --level-split)
• SRAM modeling & synthesis (synthesis flow + preload + observability in one place)
• Multi-clock domain user guide (config examples; cf. #87 for test coverage)
• GPU kernel optimization internals (profiling, backend-specific tuning)

Now covered: custom cell libraries → Adding a New PDK + Decision 0010/0011; VCD scope behaviour → Troubleshooting VCD.

Related Resources

• Main README: ../README.md - Project overview and quick start
• CLAUDE.md: ../CLAUDE.md - Development guidelines and architecture overview
• Test Suite: ../tests/ - Examples and regression tests
• Third-Party Tests: ../tests/regression/third_party/ - Real-world examples with attribution

Last Updated: 2026-06-26 Maintained By: gpu-eda community

Jacquard — Project Scope and Guarantees

Status: Draft — under review.

This document states what Jacquard is for, what it is deliberately not for, and the constraints under which it is built. It is the top-level contract that scoped requirements docs (timing-correctness.md, future equivalents) inherit from.

If you are a contributor deciding whether a feature or change fits Jacquard, start here.

Purpose

Jacquard is a GPU-accelerated gate-level simulator for synthesized digital circuits. It exists to make pre-silicon functional verification of synthesized designs substantially faster than CPU-based simulators, by mapping the design onto a GPU’s massively parallel execution model.

It is a descendant of NVIDIA Research’s GEM project, maintained by Rob Taylor and community contributors.

Jacquard ships with AIGPDK, a simple and-inverter standard cell library aligned to the GPU schema’s internal representation. It also accepts Liberty-described cells from open-source PDKs (SKY130) and commercial PDKs (under a private test track; see Decision 0004).

In scope

• Gate-level simulation of synthesized Verilog netlists against AIGPDK, SKY130, and similar Liberty-described cell libraries.
• Stimulus from static VCD input, from co-simulated GPU-resident peripheral models (SPI flash, UART, and similar), or from combinations of both for SoC-scale functional verification.
• Synchronous designs clocked by one or more clocks of known frequency.
• GPU-backed simulation across CUDA (NVIDIA), HIP (AMD), and Metal (Apple Silicon).
• Back-annotated timing simulation from SDF, including setup/hold violation detection.
• Correctness validation against permissive third-party simulators and STA tools.

Non-goals

Jacquard does not aim to cover the following areas. Feature work requiring cross-cutting changes to the GPU schema or core architecture is declined on scope grounds unless this document is first amended. Contributions that interface Jacquard with external tools covering these areas are welcome where noted.

• Mixed-signal or analog simulation. Digital gate-level only. Interfacing to external analog / mixed-signal simulators via cosim hooks is in scope for future contribution.
• RTL-level simulation. Input is a synthesized netlist, not behavioural RTL. Synthesis (via Yosys or commercial tools) happens upstream.
• Sign-off STA. Jacquard performs functional simulation with setup/hold guardrails. Sign-off timing analysis is delegated to dedicated STA tools; Jacquard validates its results against them but does not aim to replace them.
• Incremental or interactive editing of running simulations. Each Jacquard run starts from a complete netlist and stimulus. Contributions that embed Jacquard’s engine into interactive or incremental workflows (REPLs, debuggers, IDE integrations) as an external driver are welcome and considered outside this project’s direct scope rather than non-goals.

Constraints

Licensing

All code linked into the Jacquard binary must be under a permissive license (MIT / Apache-2 / BSD-3 or equivalent). GPL tools may be invoked as subprocesses. This keeps Jacquard commercially usable by downstream integrators.

Platforms

Jacquard commits to three GPU backends: CUDA, HIP, and Metal. A change that lands on one backend without a plan for the others is a regression in product surface. Feature-parity timing is negotiable; eventual feature-parity is not.

Design assumptions baked into the architecture

These are structural properties of the current GPU schema:

• Sequential logic is currently edge-triggered and synchronous: a raw latch in the logic, and asynchronous sequential (self-timed) logic, are not modelled today — the GPU schema’s scheduling assumes synchronous clocking. (Async set/reset on flip-flops is supported; so are clock gating via CKLNQD and latch-based memory mapped to RAM — see the latch note in architecture/simulation-engine.md.) Extending it to support async sequential approaches is open territory; contributions with a viable approach are welcome.
• Circuits fit the boomerang block shape: 8191-signal input/output and 4095 intermediate-pin limits per partition, 64 SRAM output groups. Very wide designs may require manual --level-split tuning.
• Numerics are 4-state at partition granularity (X-capable or not), not per-bit.

Validation

Results are validated against at least one independent third-party tool per format. No single parse path (Jacquard’s or otherwise) is its own reference. See timing-correctness.md for the detailed validation contract on timing.

Stability

Treat the following as honest characterisations, not marketing claims:

• Stable. Core GPU simulation of AIGPDK designs. NVDLA / Rocket / Gemmini regression path.
• Stable with caveats. SKY130 flow; known to work on the MCU SoC reference design, with a history of PDK-specific issues resolved over time.
• Evolving. Timing simulation: sim-path arrival tracking and setup/hold violation detection now work across Metal, CUDA, and HIP (Decision 0008), with per-DFF clock-arrival skew folded in (Decision 0007 Pillar B); cosim-path timing output (--timing-report) is not yet wired. Multi-clock scheduling. X-propagation semantics. SDF parser is hand-rolled and has received multiple reactive fixes.
• Experimental. GPU-resident peripheral models. HIP-on-NVIDIA path exists primarily to unblock CI.
• Planned. Private commercial-PDK test track (see Decision 0004).

Contributors can expect stable-tier behaviour to remain stable across releases. Evolving and experimental tiers may change shape between releases; reasonable migration notes will be provided.

Decision principles

When scope conflicts arise:

1. Product surface before performance. Jacquard’s speed advantage is its value proposition; optimizations that compromise correctness, validation, or portability are declined.
2. Permissive-license contributors win. Any change that pushes Jacquard toward GPL-contamination is rejected; subprocess-based integrations with GPL tools are acceptable.
3. Non-goals hold until amended. Extending into latches, interactive stimulus, or RTL simulation is not a bug fix; it is a scope change and requires this document to be updated first.
4. Validation is not optional. New parsers, new formats, new simulation modes require an independent third-party reference at least for representative cases.
5. Honest stability labels. Moving a feature from “evolving” to “stable” requires evidence (regression coverage, oracle-backed validation, absence of known silent-failure modes), not just time in the codebase.

References

• README.md — project overview, quick start.
• CLAUDE.md — repository conventions and architecture overview for contributors working with AI assistance.
• docs/architecture/simulation-engine.md — internal pipeline and data structures.
• docs/timing-correctness.md — scoped contract for timing accuracy, validation, and IR requirements.
• docs/architecture/decisions/ — architectural decision records.
• docs/plans/ — phased implementation plans.

Last updated: 2026-06-26 (v0.2.x; CUDA/HIP sim-timing + cosim backends shipped; reflects Decision 0007/0008/0013/0017/0018 amendments).

Installation

Jacquard is three tools; install only what your task needs:

Availability. The prebuilt-binary, Homebrew, and PyPI channels go live with the first tagged release (v0.1.0). Until then, build the simulator from source and run netlist-graph from the repo with uv run. The design behind this layout is Decision 0018.

The simulator (jacquard + opensta-to-ir)

Homebrew — macOS / Apple Silicon (Metal)

brew install gpu-eda/tap/jacquard      # installs jacquard + opensta-to-ir

The cleanest path on a Mac. Requires an Apple Silicon machine with a Metal GPU. The Homebrew formula is built with --features synth, so behavioral RTL input works out of the box — see Accepted RTL surface.

To try a release candidate before it ships, install from the prerelease tap instead (it tracks the latest -rc tag):

brew install gpu-eda/tap-prerelease/jacquard

cargo binstall — prebuilt binary, no toolchain build

brew install llvm     # runtime dependency — see note below
cargo binstall --git https://github.com/gpu-eda/Jacquard jacquard-sim \
  --disable-strategies compile,quick-install

Installs the jacquard binary (+ timing_analysis) on macOS/Metal. The crate is the jacquard-sim package (the binary it installs is still jacquard); the --git form is required because it’s not on crates.io (its dependencies are a vendored fork carrying in-flight patches), so binstall reads the [package.metadata.binstall] pkg-url straight from the repo.

Why jacquard-sim, not jacquard. The package is named jacquard-sim because the crate name jacquard is taken on crates.io by an unrelated project (an AT-Protocol client library) — so cargo install jacquard would build that, not this. Naming our package jacquard-sim makes resolution unambiguous. --disable-strategies compile,quick-install is kept as belt-and-suspenders: it turns a missing prebuilt binary into a clean hard error rather than any source-build fallback. (validate-install.yml uses the same guard.) Linux is not binstall-able: there are two GPU backends (CUDA, HIP) for one target triple, so it can’t be auto-selected — use the release tarball for your backend, a container, or build from source.

Release binaries are built with --features synth and support behavioral RTL input directly.

Runtime dependency — Homebrew LLVM. The prebuilt macOS binary links Homebrew LLVM’s libc++ and libomp (the build uses LLVM clang for OpenMP, via the mt-kahypar partitioner), so it needs brew install llvm to run. The Homebrew install handles this automatically (depends_on "llvm"); binstall and the raw tarball do not, so install LLVM first.

Prebuilt release tarball

Download jacquard-<version>-<target>.tar.gz from the releases page, extract, and put jacquard, timing_analysis, and opensta-to-ir on your PATH. The GPU kernel is embedded, but the binary still needs Homebrew LLVM at runtime (brew install llvm) — see the note above.

From source (any backend)

The portable path, and the only one for Linux CUDA / HIP today. Needs the Rust toolchain and the GPU SDK for your backend.

git clone https://github.com/gpu-eda/Jacquard.git
cd Jacquard
git submodule update --init --recursive

cargo build -r --features metal --bin jacquard         # macOS / Apple Silicon
cargo build -r --features cuda  --bin jacquard         # NVIDIA (CUDA toolkit)
cargo build -r --features hip   --bin jacquard         # AMD (ROCm)

The binary lands at target/release/jacquard. See the README’s Dependencies table for optional tooling (flatc, mdbook, OpenSTA).

Behavioral RTL on-ramp (jacquard sim design.v …) requires the synth feature (embedded YoWASP Yosys engine). Add it to your build:

cargo build -r --features metal,synth --bin jacquard   # macOS + RTL synthesis
cargo build -r --features cuda,synth  --bin jacquard   # NVIDIA + RTL synthesis
cargo build -r --features hip,synth   --bin jacquard   # AMD + RTL synthesis

A binary built without synth still simulates pre-synthesized gate-level netlists; it gives an actionable error if handed behavioral RTL.

Providing yosys.wasm for RTL synthesis — the synth engine needs the YoWASP Yosys wasm module. Discovery order:

1. --yosys-wasm <path> flag on sim/cosim (overrides the rest).

2. JACQUARD_YOSYS_WASM=/path/to/yosys.wasm environment variable.

3. Installed yowasp-yosys Python package, found automatically:

   pip install yowasp-yosys
   

4. Fetch-from-release is a planned follow-up (not yet implemented).

Prebuilt and Homebrew binaries include the synth feature; the wasm is still found via the flag, env var, or installed Python package until auto-fetch ships.

The signal-analysis companion (netlist-graph)

Pure Python — install from PyPI, no GPU or Rust needed:

uvx netlist-graph search design.gv psel     # one-off, no install
pip install netlist-graph                    # or install it

From a Jacquard checkout you can also run it without installing: uv run netlist-graph … (it’s a workspace member). See signal tracing for what it’s used for.

The timing path (opensta-to-ir + a PDK)

For post-PnR timing simulation you also need PDK Liberty files, fetched with volare/ciel (pinned in the root pyproject.toml). opensta-to-ir converts SDF to the Jacquard timing IR (.jtir) that jacquard sim --timing-ir / cosim --timing-ir consume. Pure functional (pre-PnR) runs need none of this — see signal tracing § pre-PnR functional runs.

Verify

jacquard --version
# A quick self-contained cosim (from a Jacquard checkout):
jacquard cosim tests/apb_trace/apb_trace_synth.gv \
    --config tests/apb_trace/sim_config.json \
    --top-module apb_trace --max-clock-edges 200 \
    --bus-trace-csv /tmp/apb.csv

Then head to Getting Started to run bundled designs, Accepted RTL surface to simulate your own behavioral RTL, or Synthesis Flow to prepare a high-QoR gate-level netlist.

Getting Started

This guide gets you from a fresh clone to four working simulations, in increasing order of realism. Every command here runs against designs that ship in the repository — no synthesis, no large downloads, no extra EDA tools. Each one is also a CI job, so it is known-green.

Already have your own RTL? Pass it directly to jacquard sim design.v … — the simulator auto-detects behavioral Verilog / SystemVerilog and synthesizes it transparently (see Accepted RTL surface). For peak GPU performance, or to use a commercial synthesizer (DC) for best mapping quality, see Synthesis Flow. For UVM/cocotb/SVA questions see Testbench Interop.

Prerequisites

A built jacquard binary. Either install it (see Installation) or build from a clone:

git clone https://github.com/gpu-eda/Jacquard.git
cd Jacquard
git submodule update --init --recursive
cargo build -r --features metal --bin jacquard   # macOS / Apple Silicon

The examples below use cargo run -r --features metal --bin jacquard -- … so they work straight from a clone. If you installed the binary (e.g. brew install gpu-eda/homebrew-tap/jacquard), replace that prefix with just jacquard. On Linux, swap --features metal for --features cuda (NVIDIA) or --features hip (AMD), and replace the trailing 1 (NUM_BLOCKS) with 2× your GPU’s SM/CU count.

Tier 1 — Your first simulation (a flip-flop)

The smallest possible run: a single D flip-flop. jacquard sim reads a gate-level netlist plus a static input waveform (VCD) and writes an output VCD.

cargo run -r --features metal --bin jacquard -- sim \
    tests/timing_test/dff_test_synth.gv \
    tests/timing_test/dff_test.vcd \
    /tmp/dff_out.vcd \
    1

It finishes in well under a second (6 cycles). Confirm the captured q matches the golden waveform:

diff <(grep -E '^[01]!' tests/timing_test/dff_test.vcd) \
     <(grep -E '^[01]!' /tmp/dff_out.vcd) && echo "MATCH"

That MATCH is your “Jacquard runs on this machine” checkpoint. The last argument, 1, is NUM_BLOCKS — always 1 for Metal; on CUDA/HIP set it to 2× your GPU’s SM/CU count.

Tier 2 — A reactive testbench (cosim)

jacquard sim drives the design from a fixed input waveform. Real testbenches need to react to the design’s outputs. That is what jacquard cosim is for: it runs peripheral models (UART, SPI flash, JTAG, bus monitors) as GPU kernels alongside the design, so inputs can depend on outputs cycle-by-cycle.

The dual_uart design is two independent UART transmitters that send "Hi" and "OK". Cosim decodes the serial bits back into bytes — no input VCD required; clock and reset come from the config.

cargo run -r --features metal --bin jacquard -- cosim \
    tests/dual_uart/dual_uart_synth.gv \
    --config tests/dual_uart/sim_config.json \
    --top-module dual_uart_top \
    --max-clock-edges 10000

Verify both channels decoded correctly:

python3 tests/dual_uart/check_pass.py target/test-out/dual_uart_events.json
# console: b'Hi' OK
# debug:   b'OK' OK
# PASS: both UART channels decoded correctly

The peripherals, pin mapping, baud rates, and clock are all declared in tests/dual_uart/sim_config.json. See Cosim execution model and Bus Transaction Tracing for the full peripheral set.

Tier 3 — A real open-source PDK

Tiers 1 and 2 use Jacquard’s abstract AIG cell library. Real designs come out of synthesis mapped to a real PDK (process design kit). Jacquard decomposes real standard cells — SKY130 and GF180MCU today — into its AIG representation directly, so a synthesized netlist simulates as-is.

3a. Pre-P&R (post-synthesis) — runs from the repo

A pre-place-and-route netlist is the output of logic synthesis, before layout and before any SDF timing annotation. This is the fastest, most common gate-level verification step. logic_cone is a small SKY130 netlist (real sky130_fd_sc_hd__* cells) that ships in the repo:

cargo run -r --features metal --bin jacquard -- sim \
    tests/timing_test/sky130_timing/logic_cone.v \
    tests/timing_test/sky130_timing/logic_cone.vcd \
    /tmp/logic_cone_out.vcd \
    1 --input-vcd-scope logic_cone

No --sdf flag: this is pure functional simulation of a real-PDK netlist. The circuit is 4 DFFs → nand2/nor2/and2/inv cone → DFF, computing Q = a & !b & !c & !d (registered). Its sibling inv_chain.v (16 SKY130 inverters = identity) is a cleaner “output tracks input, delayed two cycles” check if you want one.

Add --sdf logic_cone.sdf --sdf-corner typ to layer Liberty cell delays on top — that moves you from functional to timing-aware simulation, covered in Timing Simulation.

3b. A full chip on GF180MCU

To see Jacquard handle a real, large design on a real PDK end-to-end, the wafer.space chess chip_top (~227,000 cells, GlobalFoundries open-source GF180MCU PDK) is wired up as a smoke-test recipe. The post-P&R netlist is ~200 MB so it is not committed — you regenerate it from the upstream LibreLane flow. Full step-by-step recipe, expected timings, and what it exercises: tests/gf180mcu_chess_chip_top/README.md.

This is the bridge to your own designs: synthesize to SKY130 or GF180MCU, then run jacquard sim exactly as in 3a.

Where to next

Input: netlist language & RTL

What Jacquard reads

Jacquard is a gate-level emulator — it maps a synthesized and-inverter graph onto a virtual manycore GPU processor (see Decision 0014). So the direct input to jacquard sim / cosim is a gate-level Verilog netlist: structural Verilog whose leaf cells are aigpdk, SKY130, or GF180MCU standard cells.

Behavioral RTL is the intended design input — via a synthesis step. You bring your RTL, synthesize it to a gate-level netlist (memory mapping + logic synthesis to aigpdk.lib cells), and Jacquard emulates the result at 5–40×. Synthesis is currently a separate, documented step — see Synthesis Flow — because synthesis quality directly sets Jacquard’s performance, so it’s a deliberate knob (open-source Yosys works; a commercial tool like DC gives better QoR).

Roadmap: an integrated jacquard build design.v on-ramp (Yosys via YoWASP, no manual synthesis) is planned — see Decision 0021 and #162. Until then, synthesize first, then sim.

Behavioral constructs (always, if/case, reg, parameters, generate, function/task, #delay) are not read by Jacquard directly — the synthesizer elaborates them away. If you feed raw behavioral RTL to sim, it will fail to parse; run it through synthesis first.

Supported netlist syntax

The netlist parser (sverilogparse, consumed by netlistdb) accepts the structural subset a synthesizer emits:

Structure

• Multiple module … endmodule per file; // and /* */ comments; (* … *) attributes (parsed and ignored).
• Both port-list styles: non-ANSI (module m(a, b); with body input/output decls) and ANSI (module m(input [7:0] a, output b);). A wire/reg net-type keyword in an ANSI header is accepted and ignored.

Declarations

• input / output / inout / wire, scalar or bus ([hi:lo]), with comma-separated names (wire a, b, c;).

Cell instantiations

• CELL_TYPE inst (.pin(expr), …); — named-port connections only. An empty .pin() (unconnected) is accepted and dropped.

Assigns

• assign lhs = rhs;, both sides full wire expressions.

Wire expressions (in assigns, port connections, and .name(expr) hookups)

Identifiers — standard ([A-Za-z_][A-Za-z0-9_$]*) and escaped (\name-until-whitespace ).

Not supported in the netlist

Beyond the behavioral constructs above (which belong in pre-synthesis RTL), the parser does not accept:

• Positional port connections — CELL inst(a, b, c); use .pin(expr).
• Parameters — parameter / localparam, or #(…) overrides on instances.
• Preprocessor — `define / `ifdef (no preprocessor).
• Net types beyond input/output/inout/wire — no supply0/1, tri, wand/wor, trireg.
• Multi-dimensional / memory arrays — one [hi:lo] range per declaration.
• Bare unsized decimal literals — write 1'b1, not 1.
• Replication — {4{a}}.
• ~ inside a concatenation — {~a, b} (use an explicit inverter cell).
• Operators — &, |, ?:, arithmetic; and base-10 literals with x/z (8'dx).

Design-level constraints (netlistdb)

• Single top module — auto-detected as the module no other instantiates; pass --top-module <name> if ambiguous. Cyclic instantiation is rejected.
• Hierarchy is flattened — multi-level module hierarchy is supported and flattened at load.
• Leaf-cell pin directions come from the cell library (aigpdk / SKY130 / GF180MCU); an unknown cell’s pins default to Unknown with a warning.
• inout on the design boundary is parsed but modelled as Unknown (not fully supported); prefer split input/output where possible.
• Edge-triggered flops only — a raw LATCH cell is rejected (async set/reset is fine). Clock gating uses CKLNQD; latch/register-file memory maps to RAM through the memory-synthesis step (see Synthesis Flow).

SystemVerilog & assertions (SVA)

• Immediate assertions are already lowered through synthesis today, as GEM_ASSERT cells (see the assertion handling in src/aigpdk.rs).
• Broader SVA (SystemVerilog Assertions) is planned — see Testbench Interop § Roadmap. The concrete near-term slice is the X-barrier $isunknown assertion work: a spec-file-driven !$isunknown check against the X-mask (#106) and lowering an $isunknown SVA subset from RTL to that spec (#107), building on selective X-propagation (Decision 0016).

See also

• Synthesis Flow — RTL → gate-level netlist (the step above).
• Testbench Interop — UVM / cocotb / SVA, and record-and-replay.
• Decision 0014 — why the input is a synthesized AIG.
• Decision 0021 — the planned integrated RTL on-ramp.

Accepted Behavioral RTL Surface

jacquard sim and jacquard cosim accept behavioral Verilog / SystemVerilog directly — no separate synthesis step, no external toolchain. This page describes what the on-ramp synthesizes, what it drops, and where the coverage boundary lies.

One-line summary: the accepted surface is whatever YoWASP Yosys (with the yosys-slang read_slang frontend) can synthesize to an aigpdk netlist. Jacquard delegates the elaboration and synthesis; it does not implement a Verilog front-end itself.

Invoking the on-ramp

# Pass behavioral RTL directly — synthesis is transparent and cached:
jacquard sim design.v in.vcd out.vcd NUM_BLOCKS

# Cosim works the same way:
jacquard cosim design.v --config sim_config.json

sim/cosim classify the input file automatically:

Override flags:

• --rtl — force the synthesis path (useful if detection misclassifies a file).
• --netlist — force direct gate-level loading, skip detection.
• --emit-synth <path> — also write the intermediate synthesized netlist for inspection or as a fixture.

The simulator always logs its decision:

design.v: behavioral RTL → synthesized [YoWASP Yosys, functional QoR] → <cache>

Providing yosys.wasm

The synthesis engine requires yosys.wasm (the YoWASP Yosys WebAssembly module). It is located in this order:

1. --yosys-wasm /path/to/yosys.wasm — CLI flag on sim/cosim (highest priority; overrides the env var and discovery).

2. JACQUARD_YOSYS_WASM=/path/to/yosys.wasm — environment variable.

3. Installed yowasp-yosys Python package — discovered automatically via python3 -c "import yowasp_yosys …". Install with:

   pip install yowasp-yosys          # or: uvx yowasp-yosys
   

4. Fetch from a pinned release (implemented; the default when the options above aren’t set): a sha256-verified yosys.wasm is downloaded from the gpu-eda/yowasp-yosys fork and cached under $XDG_CACHE_HOME/jacquard. The RTL on-ramp reference describes the full resolution order and why the fork wasm is the default.

Version caveat. Pin to yowasp-yosys==0.64.0.0.post1131 (the version in the project’s uv.lock, verified to carry read_slang). Newer wheels ship a wasm built with the WebAssembly exception-handling proposal, which the current embedded wasmtime engine cannot load yet (pip install yowasp-yosys gets the latest and may fail with “exception refs not supported”). Loading newer modules is a tracked Phase-4 follow-up; until then, pip install 'yowasp-yosys==0.64.0.0.post1131'.

Released jacquard binaries are built with --features synth and include the Yosys WASM runtime. Source builds must add --features synth explicitly: cargo build -r --features metal,synth --bin jacquard.

A binary built without --features synth gives an actionable error when handed behavioral RTL, pointing at the synth-enabled build path.

The compiled Yosys module is cached under $XDG_CACHE_HOME/jacquard by content hash — only the first run of a given wasm pays the ~20 s cranelift compile. The synthesized netlist is also cached keyed by (design source + synth script + wasm), so repeat sim runs skip synthesis entirely.

SystemVerilog frontend

The embedded wasm (pinned yowasp-yosys 0.64.0.0.post1131) bundles yosys-slang — a near-complete SystemVerilog-2017 elaborator. Jacquard probes for read_slang at startup and uses it when present; it falls back gracefully to Yosys’s built-in read_verilog -sv for older wasm modules.

read_slang coverage is substantially broader than read_verilog -sv: packages, interfaces, structs, enums, and most SV-2017 constructs are handled by the slang elaborator before Yosys sees any RTL. If you observe a “not supported” message, check whether it names slang or read_verilog — the fallback has a narrower surface.

Supported language subset

The following are synthesizable and produce correct simulation output:

Verilog-2005 (synthesizable subset):

• All combinational logic: assign, always @(*), always @(posedge/negedge)
• Synchronous flip-flops (positive or negative edge)
• Synchronous reset DFFs (if (rst) Q <= 0; else Q <= D;)
• Asynchronous reset / set DFFs (async reset is supported; latches are not)
• Case and if/else chains
• Parameterized modules and generate blocks
• Inferred memories (mapped to RAMGEM by memlib_yosys.txt — see below)

SystemVerilog-2017 (via yosys-slang read_slang):

• Packages and package imports
• Interfaces (including modports)
• always_ff, always_comb, always_latch
• logic, enum, struct, union (synthesizable forms)
• Parameters with complex types; advanced generate (if, for, case)
• Casting and type conversions (synthesizable)

This is not the narrow subset accepted by Yosys’s built-in read_verilog -sv; it is the slang elaborator’s synthesizable coverage, which tracks the SystemVerilog-2017 standard closely.

Project-specific mappings

The synthesis script applies three project-specific transformations on top of standard Yosys synthesis:

The techmap rules are in aigpdk/gem_formal.v. Pass --strip-assertions (if available on your build) to use chformal -remove and drop GEM_ASSERT cells instead, for a pure logic netlist.

Known limits

Concurrent SVA → synthesizable checker synthesis is partial. Turning concurrent SystemVerilog Assertions (property / sequence bound to assert property with a clock) into gate-level checkers is a Yosys formal-flow capability — independent of slang’s ability to parse SVA. This is tracked in issues #106 and #107. Immediate assertions (assert (cond) inside procedural blocks) synthesize correctly via the GEM_ASSERT mapping above.

Testbench-only constructs are dropped. Synthesis discards constructs that have no synthesizable meaning: #delay timing controls, most initial blocks (other than register initialization), $finish, $stop, and similar testbench procedural code. These are silently dropped by Yosys during elaboration, not simulated. If a $display is inside an initial or inside a #delay-driven block, it may be dropped too.

Latches are not supported. Jacquard’s GPU emulator core is synchronous only — see docs/synthesis-flow.md for the latch constraint. A latch inferred from RTL causes a synthesis or AIG-load error.

QoR note

YoWASP Yosys produces functional-grade QoR — correct results but not optimized for GPU speed. Synthesis quality is the primary tuner of simulation throughput (the AIG the GPU emulates is only as good as the mapping). For peak performance, synthesize your design with a commercial synthesizer (DC) or a native Yosys install, and point jacquard sim at the resulting gate-level netlist directly. See docs/synthesis-flow.md for the performance path.

Authoritative surface: empirical coverage (Phase 4)

This prose describes the intended surface. Because the accepted RTL surface is whatever the embedded Yosys frontend synthesizes, the authoritative measure is an empirical pass/fail coverage table driven through sv-tests (or a curated subset) — a Phase 4 follow-up tracked in docs/plans/rtl-onramp-sim-integration.md. Until that table exists, hand-claimed feature lists for a delegated frontend are inherently approximate. Report gaps as bugs; they will be tracked in the coverage table when it ships.

See also

• Getting Started — run bundled designs to verify your build first.
• Synthesis Flow — the performance path (DC / native Yosys, peak QoR).
• Decision 0021 — the design decision behind the on-ramp.
• docs/plans/rtl-onramp-sim-integration.md — implementation plan and Phase 4 roadmap.

Synthesis Flow — Running Your Own Design

New to Jacquard? Start with Getting Started, which runs bundled designs in a few seconds with no synthesis. This page is the deeper path: preparing your own RTL for peak GPU performance using a commercial synthesizer (DC) or a native Yosys install.

Behavioral RTL shortcut: jacquard sim design.v in.vcd out.vcd N accepts behavioral Verilog / SystemVerilog directly — synthesis is automatic and transparent. See Accepted RTL surface. This page covers the performance path where synthesis quality is the primary tuner of simulation throughput.

Caveats: jacquard sim drives the design from a static input waveform (e.g. VCD); for reactive testbenches use jacquard cosim. Storage must be edge-triggered flip-flops — latches (level-sensitive storage) are not supported. Asynchronous set/reset on flip-flops is fine (async reset is not the restriction); what’s excluded is latch-based / level-sensitive sequential logic.

Dataset: Some (namely, netlists after AIG transformation in Steps 1-2 below, and reference VCDs) input data is available here .

Step 0. Download the AIG Process Kit

Go to aigpdk directory where you can download aigpdk.lib, aigpdk_nomem.lib, aigpdk.v, and memlib_yosys.txt. You will need them later in the flow.

Before continuing, make sure your design’s logic storage is edge-triggered flip-flops — a raw latch left in the gate-level logic is rejected (async set/reset on flip-flops is fine; async reset is not the restriction). Two structured latch uses are still supported: clock gates map to the CKLNQD integrated clock-gating cell — replace any RTL clock gates manually with CKLNQD instantiations from aigpdk.v — and latch-based register files / memory are captured by the memory-synthesis step below (memory_libmap → RAM), not left as raw latches. Also, you are advised to be familiar with where memory blocks (e.g., caches) are implemented in your design so you can check that the memory blocks are mapped correctly later.

Step 1. Memory Synthesis with Yosys

This step makes use of the open-source Yosys synthesizer to recognize and map the memory blocks automatically.

Download and compile the latest version of Yosys. Then run yosys shell with the following synthesis script.

# replace this with paths to your RTL code, and add `-I`, `-D`, `-sv` etc when necessary
read_verilog xx.v yy.v top.v

# replace TOP_MODULE with your top module name
hierarchy -check -top TOP_MODULE

# simplify design before mapping
proc;;
opt_expr; opt_dff; opt_clean
memory -nomap

# map the rams
# point -lib path to your downloaded memlib_yosys.txt
memory_libmap -lib path/to/memlib_yosys.txt -logic-cost-rom 100 -logic-cost-ram 100

The memory_libmap command will output a list of RAMs it found and mapped.

• If you see $__RAMGEM_SYNC_ (naming inherited from GEM), it means the mapping is successful.
• If you see $__RAMGEM_ASYNC_, it means this RAM is found to have asynchronous READ port. You need to confirm if it is the case.
  • If it is a synchronous one but accidentally recognized as asynchronous, you might need to patch the RTL code to fix it. There might be multiple reasons it cannot be recognized as synchronous. For example, when the read and write clocks are different.
  • If it is indeed asynchronous, check its size. If its size is very small and affordable to be synthesized using registers and mux trees (which is very expensive for large RAM banks), you can remove the $__RAMGEM_ASYNC_ block in memlib_yosys.txt, re-run Yosys to force the use of registers.
• If you see using FF mapping for memory, it means the memory is recognized, but due to it being nonstandard (e.g., special global reset or nontrivial initialization), Jacquard will fall back to registers and mux trees. If the size of the memory is small, this is usually not an issue. Otherwise, you are advised to try other implementations.

After a successful mapping, use the following command to write out the mapped RTL as a single Verilog file.

write_verilog memory_mapped.v

Check the correctness of this step by simulating memory_mapped.v with your reference CPU simulator.

Step 2. Logic Synthesis

This step maps all combinational and sequential logic into a special set of standard cells we defined in aigpdk.lib. The quality of synthesis is directly tied to Jacquard’s final performance, so we suggest you use a commercial synthesis tool like DC. You can also use Yosys to complete this if you do not have access to a commercial synthesis tool.

Check the correctness of this step by simulating gatelevel.gv with your reference CPU simulator.

Use Synopsys DC

First, you need to compile aigpdk.lib to aigpdk.db using Library Compiler.

With that, you synthesize the memory_mapped.v obtained before under aigpdk.db.

Some key commands you may use on top of your existing DC flow:

# change path/to/aigpdk.db to a correct path. same for other commands.
set_app_var link_path path/to/aigpdk.db
set_app_var target_library path/to/aigpdk.db
read_file -format db $target_library

# elaborate TOP_MODULE
# current_design TOP_MODULE

# timing settings like create_clock ... are recommended. Jacquard benefits from timing-driven synthesis.

compile_ultra -no_seq_output_inversion -no_autoungroup
optimize_netlist -area

write -format verilog -hierarchy -out gatelevel.gv

Use Yosys: Example script

# if you exited Yosys in step 2, you can read back in your memory_mapped.v yourself.
# read_verilog memory_mapped.v
# hierarchy -check -top TOP_MODULE

# synthesis
synth -flatten
delete t:$print

# change path/to/aigpdk_nomem.lib to a correct path. same for other commands.
dfflibmap -liberty path/to/aigpdk_nomem.lib
opt_clean -purge
abc -liberty path/to/aigpdk_nomem.lib
opt_clean -purge
techmap
abc -liberty path/to/aigpdk_nomem.lib
opt_clean -purge

# write out
write_verilog gatelevel.gv

Step 3. Download and Compile Jacquard

Download and install the Rust toolchain. This is as simple as a one-liner in your terminal. We recommend https://rustup.rs.

Clone Jacquard along with its dependencies.

git clone https://github.com/gpu-eda/Jacquard.git
cd Jacquard
git submodule update --init --recursive

Jacquard supports two GPU backends: CUDA (NVIDIA GPUs on Linux) and Metal (Apple Silicon Macs).

All functionality is accessed through the jacquard CLI, which provides sim, cosim, dump-paths, and other subcommands:

# Analyze timing / validate a netlist (no GPU features needed)
cargo run -r --bin jacquard -- dump-paths --help

# Simulation (Metal - macOS)
cargo run -r --features metal --bin jacquard -- sim --help

# Simulation (CUDA - Linux, requires CUDA toolkit)
cargo run -r --features cuda --bin jacquard -- sim --help

Simulate the Design

Jacquard automatically partitions the design at startup using mt-kahypar-sc hypergraph partitioning.

If partitioning fails due to deep circuits (which often shows as trying to partition a circuit with only 0 or 1 endpoints), try adding a --level-split option to force a stage split. For example --level-split 30 or --level-split 20,40.

Metal (macOS)

Use NUM_BLOCKS=1 for Metal.

cargo run -r --features metal --bin jacquard -- sim path/to/gatelevel.gv path/to/input.vcd path/to/output.vcd 1

CUDA (Linux)

Replace NUM_BLOCKS with twice the number of physical streaming multiprocessors (SMs) of your GPU.

cargo run -r --features cuda --bin jacquard -- sim path/to/gatelevel.gv path/to/input.vcd path/to/output.vcd NUM_BLOCKS

VCD Scope Handling

Jacquard automatically detects the correct VCD scope containing your design’s ports. In most cases, you don’t need to specify --input-vcd-scope. If auto-detection fails or you need to override it, use:

# Metal
cargo run -r --features metal --bin jacquard -- sim path/to/gatelevel.gv path/to/input.vcd path/to/output.vcd 1 --input-vcd-scope "testbench/dut"

# CUDA
cargo run -r --features cuda --bin jacquard -- sim path/to/gatelevel.gv path/to/input.vcd path/to/output.vcd NUM_BLOCKS --input-vcd-scope "testbench/dut"

Use slash separators (/) for hierarchical paths, not dots. See troubleshooting-vcd.md for details.

The simulated output ports value will be stored in output.vcd.

Caveat: The actual GPU simulation runtime will also be outputted. You might see a long time before GPU enters due to reading and parsing input.vcd. You are recommended to develop your own pipeline to feed the input waveform into Jacquard’s GPU kernels.

Timing-Aware Simulation

Jacquard supports two ways to feed timing data into the simulator:

1. --timing-ir <path.jtir> — pre-converted Jacquard timing IR. This is the canonical path and requires no external tools at run time. Generate the IR ahead of time with the standalone opensta-to-ir tool (see crates/opensta-to-ir/).
2. --sdf <path.sdf> --liberty <path.lib> — raw SDF, converted to IR on the fly. This subprocesses OpenSTA, which must be installed on the user’s machine.

OpenSTA dependency

When using --sdf, Jacquard locates OpenSTA in this order:

1. JACQUARD_OPENSTA_BIN environment variable.
2. <repo-root>/scripts/build-opensta.sh --print-binary (the canonical install path during development; the script builds the version vendored at vendor/opensta/).
3. sta on PATH.

Jacquard requires OpenSTA 3.1.0 or newer, matching the commit pinned at vendor/opensta/. The pinned version is the only one with end-to-end test coverage; newer OpenSTA versions are accepted with a warning, older versions are a hard error.

The simplest way to get a known-good OpenSTA is to build the vendored copy from the Jacquard repo:

git submodule update --init --recursive
./scripts/build-opensta.sh

Then either set JACQUARD_OPENSTA_BIN to the path printed by ./scripts/build-opensta.sh --print-binary, or just let Jacquard find it automatically — the build script’s output is searched by default.

Error messages

For licensing context (Jacquard is permissively-licensed, OpenSTA is GPL-3, and Jacquard’s runtime subprocess invocation is permitted but bundling is not), see architecture/decisions/0006-sdf-preprocessing-model.md.

Testbench Interop (UVM, cocotb, SVA)

A common first question: “Can I point my existing UVM / cocotb / SVA testbench at Jacquard?”

Short answer: not today. Jacquard is a gate-level engine, not a drop-in replacement for a SystemVerilog or Python simulator’s testbench runtime. This page explains what works now, what is on the roadmap, and the fallback that works for any flow today.

What Jacquard drives today

So reactive stimulus is supported — but through Jacquard’s own peripheral model architecture and cosim execution model, not through an external testbench framework.

The fallback that works for any flow: record-and-replay

You do not need framework integration to use Jacquard with a UVM, cocotb, or plain-Verilog testbench. The universal path:

1. Run your existing testbench against any simulator that can dump a VCD (Verilator, Icarus, a commercial simulator, …).
2. Dump the design’s top-level input pins to a VCD.
3. Replay that VCD through Jacquard with jacquard sim.

This is exactly what jacquard sim already is — a recorded-waveform replay — so it works regardless of how the stimulus was generated. The trade-off is that the replay is open-loop: it reproduces the recorded inputs, so it can’t react to divergent design behaviour. For closed-loop reactive stimulus, model the peripheral as a cosim kernel instead.

Roadmap

These are directions under consideration, not commitments or dated milestones:

• SVA (SystemVerilog Assertions) — planned. Jacquard already lowers a class of immediate assertions through synthesis (GEM_ASSERT cells; see the assertion handling in aigpdk.rs). Broader SVA support is the next step here.
• Running UVM test suites — design settled, unbuilt. The shape is Decision 0022: split the testbench the way emulators have since SCE-MI. Sequences, randomisation and checking stay on the host; the driver’s timed half is rewritten as a synthesizable transactor and compiled into the AIG beside the design (possible since the RTL on-ramp shipped), so it wiggles pins at GPU speed. Existing UVM drivers don’t port as-is, and that isn’t a gap we can close: they drive every cycle, and cosim runs 1024 edges per dispatch with the CPU absent, which is where the speed comes from.
• cocotb — needs more work. A naive bridge would marshal Python ↔ GPU every cycle, which would dominate runtime and erase the GPU speedup. Same cliff as above, same fix: exchange transactions, not cycles. Making cocotb performant against Jacquard needs more design thinking, not just a shim.

If any of these is blocking for you, the record-and-replay path above is the recommended interim approach.

See also

• Getting Started — sim and cosim walkthroughs
• Bus Transaction Tracing — protocol-aware observation of on-chip buses
• Cosim execution model — how peripheral kernels run

Why Jacquard — positioning and output interface

Status: Honest assessment of where Jacquard fits in an EDA flow alongside dedicated STA tools (OpenTimer/OpenSTA) and event-driven simulators (Verilator, iverilog, CVC). Includes a survey of what timing information Jacquard exposes today and what would let users actually consume it.

This is not a marketing document. The goal is for a contributor or user to read it and decide accurately whether Jacquard helps them — and, if it does, how to extract the answer they need.

TL;DR

Jacquard’s unique value is vector-driven timing analysis at GPU scale: answering “did this stimulus violate setup/hold at any DFF, on which cycle, on which signal?” for designs large enough that SDF-annotated event-driven sim is too slow to finish in useful time.

Everything else Jacquard offers is offered, often better, by the standard flow:

• For functional sim: Verilator is faster on small designs.
• For timing: OpenSTA gives more accurate answers than Jacquard, vector-independent.
• For glitch / metastability: event-driven sim with SDF (CVC, iverilog) sees behaviours Jacquard’s lockstep kernel structurally cannot.

Jacquard becomes the right tool when (design size × vector length) exceeds what event-driven SDF-annotated sim can handle, and you specifically want vector-driven timing answers.

STA is not optional even with Jacquard. Jacquard does not replace OpenSTA; it complements it. The right framing is “STA proves no bad vectors exist; Jacquard proves your real workload runs cleanly within those bounds.” OpenSTA is also a hard runtime dependency for any timing-aware Jacquard flow — the timing IR is produced by opensta-to-ir, which subprocesses OpenSTA. See Decision 0001.

What’s actually unique

The intersection where Jacquard wins is narrow but real:

1. Activity-driven setup/hold sweep at scale. Run a long workload (boot trace, architectural validation, NoC congestion stimulus) on a large design at GPU speed; get a per-cycle violation report. STA can’t tell you “this real workload trips violation X at cycle 12,847”; CVC can but won’t finish in time on big designs.

2. Arrival-time distributions for power/activity analysis. Per-signal arrival histograms across millions of cycles → useful for worst-case-power analysis informed by actual switching activity. STA gives you nothing here; CVC could but slowly.

3. Failure forensics. When a functional test fails, answering “was this a timing issue?” without rerunning under a different simulator. Jacquard’s timing-VCD output ties violations to cycle/signal/path — useful when you already have it from the same run.

4. Fast iteration during timing closure. Change a constraint, resynthesise, re-run a long test — Jacquard’s loop time is short enough to make this practical on big designs in a way iverilog+SDF isn’t.

What dedicated STA (OpenSTA) gives you that Jacquard doesn’t

This list is long and you should know it:

• Worst-case path enumeration. STA tells you the top-N critical paths over all possible inputs. Jacquard sees only what your stimulus exercises. If your testbench misses a critical path, Jacquard’s “no violations” report is silent on it; OpenSTA would flag it.
• True min-delay analysis. OpenSTA does proper min-delay path search. Jacquard’s hold check is per-DFF against actual stimulus only.
• Per-pair CRPR. OpenSTA applies common-path-pessimism removal as a launch/capture credit on each path. Jacquard consumes per-DFF clock arrival from opensta-to-ir and folds it into setup/hold (see timing-model-extensions.md, Part B Stages 1+2 — landed), but treats the launch reference as 0 — i.e. the per-pair CRPR credit is intentionally not modelled at this stage. Stage 3 in the same doc is the lever if Stage 1+2 pessimism turns out to matter on a real design.
• SDC-aware constraint handling. False paths, multi-cycle paths, generated clocks, async groups — OpenSTA reads SDC and respects it. Jacquard doesn’t read SDC at the timing layer.
• Coverage by construction. STA covers every path by definition. Dynamic sim covers only what’s exercised.
• Vector-independent confidence. “This design meets timing” is something STA can claim; Jacquard can only claim “this design met timing on these vectors.”

What event-driven SDF sim (CVC/iverilog) gives you that Jacquard doesn’t

The honest comparison isn’t “Jacquard vs. Verilator + OpenTimer.” It’s “Jacquard vs. iverilog/CVC-with-SDF + OpenTimer.” On the timing-sim side specifically:

• Glitch propagation. CVC/iverilog with inertial or transport delay see intra-cycle pulses. Jacquard’s lockstep cycle-accurate kernel does not.
• Per-pin wire delay fidelity. CVC consumes SDF interconnect records per-receiver, per-edge, with rise/fall distinction. Jacquard collapses to per-cell-max (see timing-model-extensions.md, Part C).
• Per-DFF setup/hold without per-word collapse pessimism. Jacquard collapses all DFFs in a 32-bit state word to min(setup), min(hold); CVC checks each flop individually.
• Async event handling. Real $setup/$hold checks across asynchronous control. Jacquard explicitly assumes synchronous designs.

So today, accuracy-per-vector goes to CVC; throughput goes to Jacquard.

When to choose what

The trajectory

Jacquard’s timing fidelity gap with CVC is closeable. The work in timing-model-extensions.md — δ(T), clock-tree skew, per-receiver wire delay — closes much of it while preserving GPU throughput. The further along that path the project goes, the more “Jacquard” looks like “GPU-accelerated SDF-annotated event-driven sim, with the inherent limits the cycle-accurate kernel imposes (no glitches, lockstep cycles)” — i.e. CVC’s report quality at Verilator’s speed, on designs where neither alone suffices.

Output interface — what Jacquard exposes today

Jacquard’s unique value depends on getting the timing information out of a run in a form users can act on. Phase 1 of the post-Phase-0 roadmap (Decision 0008) closed the gap between “data Jacquard has” and “answers users want” for setup/hold violations.

Symbolic stderr violation messages

The kernel writes setup/hold violation events to a per-block event buffer (csrc/kernel_v1.metal:554-576). The host drains the buffer each cycle (src/event_buffer.rs), resolves the state-word index to a hierarchical DFF site name via WordSymbolMap, and emits:

[cycle 12847] SETUP VIOLATION at top/cpu/regs[7][bit 22] [word=412]: arrival=2150ps setup=80ps slack=-30ps
[cycle 12847] HOLD VIOLATION at top/cpu/state[bit 3] [word=412]: arrival=12ps hold=20ps slack=-8ps

The bare [word=N] suffix is preserved for grep/tooling compatibility; up to four DFFs per word are named, with +N more truncation beyond that.

Structured timing report (--timing-report <path.json>)

Schema-versioned JSON document written at end of run. Contents:

• Per-cycle violation list (cycle, kind, word, site, arrival, constraint, slack).
• Per-word aggregate: violation counts and worst slack (sorted by total violations).
• Top-N worst-slack ranking per kind (setup, hold).
• Run metadata: design, vector source, timing source, clock period, cycles run, Jacquard version.
• Aggregate stats: setup/hold totals, dropped events.

Machine-readable, CI-friendly. Sample at tests/timing_ir/sample_reports/two_violations.json; full schema in src/timing_report.rs (SCHEMA_VERSION = "1.0.0"). Stability contract per Decision 0008: additive-only extensions, breaking changes bump the major.

Text summary (--timing-summary)

One-screen human summary on stdout. Same data as the JSON report, different channel; either or both flags can be set:

=== Jacquard Timing Summary ===
Design:        my_cpu.gv
Vectors:       boot.vcd (1000 cycles)
Clock period:  1000 ps
Timing source: my_cpu.jtir

Violations:
  Setup: 5
  Hold:  2
  Total: 7

Worst slack:
  Setup: -150ps  at top/cpu/regs[7][bit 22] [word=5]  (cycle 87)
  Hold:   -40ps  at top/cpu/state[bit 3] [word=12]  (cycle 91)

Top 2 by violation count (of 2 total words with violations):
  top/cpu/regs[7][bit 22] [word=5] (5 violations): worst setup=-150ps hold=- arrival=950ps
  top/cpu/state[bit 3] [word=12] (2 violations): worst setup=- hold=-40ps arrival=10ps

Format is for human inspection — explicitly not a stable parseable contract. Tools should use --timing-report JSON.

Timed VCD (--timed)

Annotates the output VCD with per-signal arrival times. Largest, most detailed output; suitable for waveform-level inspection.

• What you get: per-signal arrival ps at each writeout cycle.
• Caveat: the VCD doesn’t carry slack relative to the clock edge — you compute it yourself.
• Cost: doubles VCD size. Not appropriate for long workloads on large designs.

SimStats aggregate counts (in-process)

SimStats { setup_violations, hold_violations, ... } is available to in-process consumers (src/event_buffer.rs). Only counts; full detail flows through the structured report path.

Still on the wishlist

Items captured in Decision 0008’s “Optional / later outputs” plus a few caveats on what shipped. Demand-driven; not scheduled.

Closest-to-violation tracking when no violation occurred

The shipped worst_slack ranking is populated only from observed violation events. Surfacing “where am I close to the edge” on a run that passed timing requires GPU-side near-miss instrumentation (emit slack events whenever |slack| falls below a configurable threshold). Useful for proactive signoff regression. Separate workstream — needs a kernel change.

Arrival histogram (--arrival-histogram <pattern>)

Per-signal arrival histogram dump for matched signal patterns, as JSON or CSV. Foundation for activity-based power analysis and “is my actual timing margin healthy” reporting.

STA cross-reference (--sta-cross-reference <opensta-paths.txt>)

Read OpenSTA’s worst-N critical-path report and produce coverage output: of those paths, which were exercised by the stimulus, at what observed arrival. Closes the loop between vector-driven and static analysis.

Path back-trace from worst-arrival DFF

Given a flagged DFF, walk the max-of-fanin chain backward to the source AIG pin / primary input, emitting per-edge contributions. Most expensive item on the wishlist; only useful once symbolic names are in place (which they now are).

CUDA / HIP / cosim runtime violation routing

The current Metal sim path routes runtime violations through process_events (which is what feeds the resolver, structured report, and text summary). The CUDA, HIP, and cosim paths don’t yet share that plumbing — they detect violations on the GPU but don’t drain through process_events. Independent plumbing follow-up; doesn’t affect the Metal user experience.

Per-signal activity / transition counts

Listed in Decision 0008 as part of the JSON report’s wishlist. Not in v1.0.0 of the schema; will be added (additively) when the GPU kernel emits transition events.

“Corner” and “margin percentage” in the text summary

Decision 0008’s summary template includes both. Corner is missing because the metadata struct doesn’t carry it through from the IR yet; margin percentage is trivially derivable from slack_ps / clock_period_ps and was omitted to keep the v1 summary terse.

Related artefacts

• project-scope.md — what Jacquard is for and not for; the formal contract this doc operates inside.
• timing-correctness.md — forward-looking validation requirements.
• timing-violations.md — current GPU-side violation detection mechanics.
• timing-validation.md — how Jacquard’s timing output is validated against CVC/iverilog.
• timing-model-extensions.md — proposed accuracy improvements (δ(T), clock-tree skew, wire delay).

Timing Correctness Requirements

Status: Draft — under review.

This document is the contract for Jacquard’s timing correctness story. It defines what must be true; decisions under docs/architecture/decisions/ define how; plans under docs/plans/ execute them.

Scope

Inherits constraints and non-goals from project-scope.md — in particular the permissive-license requirement for linked code and the synchronous-only design assumption.

Addresses three known weaknesses from the April 2026 architecture review:

1. Timing-model pessimism from the packed-32 ALU (no single fix; needs refinement path for sign-off use).
2. Weak correctness observability (CPU reference and GPU kernel share Rust source, so representation bugs pass both sides).
3. Hand-rolled SDF parser fragility on real post-P&R output from production tools.

Out of scope (tracked elsewhere, addressed later): hard schema limits (no latches, static VCD), three-backend maintenance cost, GPU-side IO model scaling, silent-failure engineering patterns beyond parsers, partitioner non-determinism.

Principles

P1 — OpenSTA is the oracle

When Jacquard’s results disagree with OpenSTA, Jacquard is wrong until demonstrated otherwise. In the shipped release, OpenSTA is never invoked from the jacquard runtime binary, and never linked (it is GPL; Jacquard stays permissive). Subprocess invocation from CI, test harnesses, and the standalone opensta-to-ir preprocessing tool is acceptable. During development before first release, a runtime subprocess invocation may exist as a contributor-ergonomics convenience (see Decision 0006) — this is explicitly interim and is removed before release. Divergence is never accepted silently: it is fixed, explicitly justified in-doc, or filed as a bug.

P2 — No single parse path is its own reference

Every format we consume (SDF, Liberty, SPEF, Verilog) is validated against at least one third-party tool’s parse of the same file. This is the primary defense against representation bugs that would otherwise affect Jacquard’s primary and reference paths simultaneously.

P3 — Fail loud on silent failures

Every parser and pipeline stage emits a success count (cells parsed, arcs matched, annotations applied). The test harness asserts thresholds. Zero-match is never silently acceptable. A pipeline that quietly succeeds with the wrong data is a bug.

P4 — Multi-corner from day one

Any timing-value representation natively supports multiple PVT corners. Single-corner shortcuts become tech debt the moment commercial flows arrive; the shape is enforced upfront.

Functional requirements

R1 — Timing IR

A canonical intermediate representation (IR) for SDF-equivalent timing annotations exists. It is:

• Schema-versioned with compatible evolution rules.
• Lossless for the subset of information Jacquard consumes downstream.
• Preserves vendor-specific annotations as typed passthrough (never silently dropped).
• Multi-corner by default.
• Tags each arc with provenance: source tool, source file, and origin category (asserted / computed / defaulted).

Details: docs/architecture/decisions/0002-timing-ir.md.

R2 — In-process reference STA — deferred

The original requirement was an in-process STA engine that computes per-endpoint arrival/slack from Liberty + SPEF, linked directly (subject to the permissive-license constraint in project-scope.md), cross-checking Jacquard’s SDF-derived timing at load time and on demand during sim.

Preferred implementation was OpenTimer; the SKY130 spike (docs/architecture/decisions/spikes/opentimer-sky130.md) found OpenTimer’s input pipeline unfit for OpenROAD-flow outputs, and Decision 0003 was Superseded (commit d002bde). The cross-check role is now performed out of process by OpenSTA via opensta-to-ir (Decision 0001 — sole STA path); the in-process variant is parked until a fit-for-purpose permissive option appears (libreda-sta or in-house walker, both behind a future decision).

R3 — Oracle-backed CI

The OpenSTA test corpus is vendored (or submoduled) into the repository and used as a regression fixture. Every IR converter runs against the corpus; converter-to-converter diffs must be explained or fixed before merge.

Jacquard’s own regression designs are run through OpenSTA (subprocess) and compared against Jacquard’s output. Runs nightly or pre-release, not per-PR, due to runtime cost.

R4 — Critical-path refinement reporting

For any sim run that reports timing violations, or for any user-requested critical-path report, top-K paths are reported with:

• Full per-stage path trace with per-edge delay.
• Pessimism delta: the gap between the packed-thread max arrival and the actual arrival along this specific path (exposes the magnitude of the packed-32 ALU’s worst-case accounting).
• Provenance of each delay (SDF-asserted / Liberty-computed / defaulted).

R5 — Private PDK testing

A private test track for commercial PDKs exists, gated on per-PDK environment variables (e.g. <VENDOR>_PDK_PATH). Tests skip cleanly when PDK files are unavailable; CI runs with PDK access execute them. No PDK-derived artifacts are committed. Details: docs/architecture/decisions/0004-private-pdk-testing.md.

Non-functional requirements

N1 — Startup parse speed

IR consumption at sim startup is effectively O(1) of IR size (binary, memory-mapped, zero-copy). The expensive SDF→IR conversion is a one-time preprocessing step, not an every-sim cost.

N2 — Reproducibility

Given identical inputs and tool versions, every converter produces byte-identical IR. Tool version bumps may change output but must be explicit (recorded in IR metadata).

N3 — Modularity

Each converter (SDF→IR, Liberty→IR, OpenSTA→IR, reference-STA→IR) is testable in isolation, without the full Jacquard build. Converters are separate crates or binaries. Consumers of IR need not know which converter produced it.

Acceptance criteria — phase 0

Phase 0 (see docs/plans/archive/phase-0-ir-and-oracle.md) is complete when:

1. Timing IR schema defined, with both a binary encoding and a JSON sidecar for human/CI diffs.
2. OpenSTA → IR converter implemented as a subprocess-driven tool.
3. Jacquard’s existing SDF parser emits IR alongside its native representation.
4. CI harness runs both converters on the OpenSTA test corpus and reports structured diffs.
5. At least one representative Jacquard test design produces matching IR between both paths, within a declared tolerance.
6. Parser-success assertions (per P3) in place for the SDF and Liberty paths.

Phases 1 and beyond are planned at the start of each phase, not all up front.

Open questions

Items not settled by this document; they resolve in decisions, spike outcomes, or phase plans:

• Exact IR schema format (FlatBuffers / Cap’n Proto / other). Tracked in Decision 0002.
• Whether OpenTimer handles SKY130 Liberty robustly. Resolved: spike failed Q2 on SKY130; Decision 0003 Superseded.
• Whether SPEF gets its own IR or is embedded in the timing IR. Deferred; likely separate.
• Whether the IR is Jacquard-local or shared across a broader tooling ecosystem. External decision; answer affects investment level and schema stability requirements.

Supersession

This document supersedes the “±5% arrival tolerance” convention from docs/timing-validation.md once phase 0 ships. Until then, the existing convention remains in effect for designs not yet covered by oracle-backed CI.

References

• docs/architecture/simulation-engine.md — current pipeline.
• docs/timing-simulation.md — current timing-sim usage.
• docs/timing-validation.md — current validation methodology (to be superseded).
• docs/architecture/decisions/ — decisions executed through this document.
• docs/plans/ — phased implementation.
• docs/architecture/decisions/spikes/ — time-boxed experiments.

Last updated: 2026-04-23

Timing Simulation in GEM

See also: timing-correctness.md — forward-looking validation contract and timing IR requirements (in progress). The document below describes current behaviour.

This document explains GEM’s boomerang evaluation architecture and how timing simulation with per-gate delays can be implemented efficiently on GPU.

Current status (what ships today)

Phase 1 closed 2026-05-02; see plans/post-phase-0-roadmap.md.

What “HIP” means above. HIP is two quite different things, and for most of this project’s life the table only ever meant the first:

• HIP-over-CUDA — hip-runtime-nvidia, built with the CUDA toolkit and executed on an NVIDIA GPU (HIP Tests (NVIDIA backend), tesla4-runner). Source-compatible with CUDA and, in practice, CUDA semantics throughout.
• HIP on real ROCm — an actual AMD GPU (HIP Tests (ROCm backend)). This had never been built, let alone run, until 2026-07-15, and sim did not run there at all until the non-cooperative fallback landed: the device has no cooperative launch, so the grid-wide barrier sim relies on is unavailable and the host drives one launch per (cycle, stage) instead. See architecture/decisions/spikes/amd-laptop-backend.md.

Both are now covered by CI and both pass, so the ✅s above are accurate for each. The distinction still matters when reading a green tick: it is what hid #203 — an out-of-bounds read that nvcc silently truncated back into range, so CUDA and HIP-over-CUDA were accidentally correct while ROCm faulted. Arrival/violation arithmetic itself is host-side Rust (#195) and therefore backend-independent.

See timing-violations.md for the full violation-output interface and why-jacquard.md for positioning.

Background: The Simulation Challenge

GEM simulates And-Inverter Graphs (AIGs) where every node is either:

• A primary input (value comes from VCD stimulus)
• An AND gate with two inputs (possibly inverted)

Traditional simulation evaluates gates in topological order, which is inherently serial. GPUs excel at massive parallelism - thousands of threads doing the same operation on different data. GEM bridges this gap with the boomerang architecture.

Boomerang Evaluation

Core Concept

The boomerang structure is a hierarchical reduction tree that maps an AIG onto GPU threads. It’s called “boomerang” because data flows down the tree during reduction, then results are written back out at various levels - like a boomerang going out and returning.

Hierarchy Structure

GEM uses BOOMERANG_NUM_STAGES = 13, meaning the tree has 2^13 = 8192 leaf positions:

Level 0 (inputs):   8192 positions
Level 1:            4096 positions  (8192 / 2)
Level 2:            2048 positions
Level 3:            1024 positions
Level 4:             512 positions
Level 5:             256 positions
Level 6:             128 positions
Level 7:              64 positions
Level 8:              32 positions
Level 9:              16 positions
Level 10:              8 positions
Level 11:              4 positions
Level 12:              2 positions
Level 13 (output):     1 position

Each level halves the number of positions by computing AND gates that combine pairs.

Thread Organization

A GPU block has 256 threads (threadIdx.x = 0..255). Each thread holds a 32-bit word where each bit represents an independent Boolean signal:

Thread 0:   [bit0, bit1, bit2, ... bit31]  = 32 Boolean signals
Thread 1:   [bit0, bit1, bit2, ... bit31]  = 32 Boolean signals
...
Thread 255: [bit0, bit1, bit2, ... bit31]  = 32 Boolean signals
            ─────────────────────────────
            Total: 256 × 32 = 8192 signals per level

Thread position refers to threadIdx.x - which of the 256 threads we’re addressing. Each thread position processes 32 signals in parallel using SIMD operations.

Memory Layout

__shared__ u32 shared_metadata[256];   // Partition configuration
__shared__ u32 shared_writeouts[256];  // Output staging area
__shared__ u32 shared_state[256];      // Working state (8192 bits)

The shared_state array holds the current level’s values during reduction.

The Reduction Process

Phase 1: Level 0 → Level 1 (hier[0])

Only threads 128-255 are active. Each computes 32 AND gates in parallel:

if(threadIdx.x >= 128) {
    u32 hier_input_a = shared_state[threadIdx.x - 128];  // From threads 0-127
    u32 hier_input_b = hier_input;                        // This thread's data

    // 32 AND gates computed simultaneously (one per bit)
    u32 ret = (hier_input_a ^ hier_flag_xora) &
              ((hier_input_b ^ hier_flag_xorb) | hier_flag_orb);

    shared_state[threadIdx.x] = ret;
}

The xora, xorb, and orb flags encode:

• xora/xorb: Input inversions (for AND-inverter graph)
• orb: Passthrough mode (when output equals input A, skip the AND)

Visual representation:

Before:  [T0][T1]...[T127] [T128][T129]...[T255]
              │                  │
              └───────┬──────────┘
                      │
                   AND gates (128 threads × 32 bits = 4096 gates)
                      │
                      ▼
After:   [----unused----] [T128][T129]...[T255]
                          (128 × 32 = 4096 results)

Phase 2: Levels 1-3 (Shared Memory)

for(int hi = 1; hi <= 3; ++hi) {
    int hier_width = 1 << (7 - hi);  // 64, 32, 16
    if(threadIdx.x >= hier_width && threadIdx.x < hier_width * 2) {
        u32 hier_input_a = shared_state[threadIdx.x + hier_width];
        u32 hier_input_b = shared_state[threadIdx.x + hier_width * 2];
        u32 ret = (hier_input_a ^ xora) & ((hier_input_b ^ xorb) | orb);
        shared_state[threadIdx.x] = ret;
    }
    __syncthreads();  // Barrier between levels
}

Each level activates fewer threads:

• Level 1: threads 64-127 (64 threads → 2048 gates)
• Level 2: threads 32-63 (32 threads → 1024 gates)
• Level 3: threads 16-31 (16 threads → 512 gates)

Phase 3: Levels 4-7 (Warp Shuffle)

Within a single warp (32 threads), data exchange uses fast shuffle instructions instead of shared memory:

if(threadIdx.x < 32) {
    for(int hi = 4; hi <= 7; ++hi) {
        int hier_width = 1 << (7 - hi);  // 8, 4, 2, 1
        u32 hier_input_a = __shfl_down_sync(0xffffffff, tmp_cur_hi, hier_width);
        u32 hier_input_b = __shfl_down_sync(0xffffffff, tmp_cur_hi, hier_width * 2);
        if(threadIdx.x >= hier_width && threadIdx.x < hier_width * 2) {
            tmp_cur_hi = (hier_input_a ^ xora) & ((hier_input_b ^ xorb) | orb);
        }
    }
}

No synchronization needed - warp shuffle is implicitly synchronized.

Phase 4: Levels 8-12 (Bit Operations)

The final levels operate on bits within a single u32, computed by thread 0 only:

if(threadIdx.x == 0) {
    // Level 8: 32 → 16 (operates on upper/lower halves)
    u32 r8 = ((v1 << 16) ^ xora) & ((v1 ^ xorb) | orb) & 0xffff0000;

    // Level 9: 16 → 8
    u32 r9 = ((r8 >> 8) ^ xora) & (((r8 >> 16) ^ xorb) | orb) & 0xff00;

    // Level 10: 8 → 4
    u32 r10 = ((r9 >> 4) ^ xora) & (((r9 >> 8) ^ xorb) | orb) & 0xf0;

    // Level 11: 4 → 2
    u32 r11 = ((r10 >> 2) ^ xora) & (((r10 >> 4) ^ xorb) | orb) & 0b1100;

    // Level 12: 2 → 1
    u32 r12 = ((r11 >> 1) ^ xora) & (((r11 >> 2) ^ xorb) | orb) & 0b10;

    tmp_cur_hi = r8 | r9 | r10 | r11 | r12;
}

Write-Outs

Results are captured at various levels (not just the final output) and written to global memory:

if((writeout_hook_i >> 8) == bs_i) {
    shared_writeouts[threadIdx.x] = shared_state[writeout_hook_i & 255];
}

This is the “return” part of the boomerang - results flow back from intermediate levels.

Timing Simulation Approaches

Approach Comparison

Recommended: Levelized with Delay Accumulation

This approach piggybacks on the existing boomerang structure with minimal changes.

Data Structure Addition

// Add to shared memory (256 bytes additional)
__shared__ u8 shared_arrival[256];  // One arrival time per thread position

Each thread position stores a single 8-bit arrival time representing the maximum arrival across all 32 bits in that position.

Modified AND Gate Evaluation

// Current (value only):
u32 ret = (hier_input_a ^ xora) & ((hier_input_b ^ xorb) | orb);
shared_state[threadIdx.x] = ret;

// With timing (add ~4 instructions):
u32 ret = (hier_input_a ^ xora) & ((hier_input_b ^ xorb) | orb);
shared_state[threadIdx.x] = ret;

u8 arr_a = shared_arrival[threadIdx.x - offset_a];
u8 arr_b = shared_arrival[threadIdx.x - offset_b];
u8 arr_ret = min(max(arr_a, arr_b) + GATE_DELAY, 255);  // Saturating add
shared_arrival[threadIdx.x] = arr_ret;

Complexity Analysis

• Same number of kernel launches as zero-delay simulation
• O(levels × cycles) - identical to current
• ~256 bytes additional shared memory per partition
• Estimated 10-20% performance overhead

The Approximation Trade-off

What We Track

One arrival time per thread position (256 values) instead of per signal (8192 values).

Implications

If thread position 50 contains signals A, B, C with different true arrivals:

Signal A: 15ps (shortest path)
Signal B: 23ps (longest path)
Signal C: 8ps  (medium path)

We store only: arrival[50] = 23ps (the maximum).

Why This Works

1. Conservative: We might report false violations, but never miss real ones
2. Correlated signals: Signals at the same thread position are often topologically nearby with similar timing
3. Endpoint focus: We ultimately only care about arrivals at DFF D inputs

When Full Accuracy is Needed

For bit-accurate timing, you would need:

// 8KB additional shared memory (may exceed limits)
__shared__ u8 shared_arrival[256][32];  // Per-bit arrivals

This is feasible but significantly increases memory pressure and computation.

Implementation Phases

Phase 1: CPU Timing Analysis (Completed)

• Liberty parser for delay extraction
• Static timing analysis on AIG
• CPU reference simulation with delays
• Timing violation detection

Phase 2: Hybrid GPU+CPU (Completed)

• GPU performs zero-delay value simulation
• CPU performs timing analysis on results
• Validates infrastructure without kernel changes

Phase 3: GPU Arrival Tracking (Completed)

• Added shared_arrival[256] (u16) to Metal and CUDA kernels
• Arrivals tracked during boomerang reduction at all hierarchy levels
• Per-gate delays injected via script padding slots from SDF data
• DFF timing constraint checking at cycle boundaries (setup/hold)
• Timing-aware VCD output (--timed flag)
• Validated against CVC reference simulator (88ps / 7.1% conservative overestimate)

Phase 4: Full Integration (Partial)

• Timing violation events via event buffer (completed)
• Per-cycle timing reports (completed)
• Integration with output VCD (completed via --timed)
• Timing-aware bit packing for reduced approximation error (future)

Conservative Timing Model: Sources of Overestimation

Jacquard’s GPU timing is intentionally conservative — it may over-estimate arrival times but will never under-estimate them. This is important for setup violation detection: false positives are safe, false negatives would miss real bugs.

There are three independent sources of conservatism, each adding to the overestimate:

Source 1: max(rise, fall) per cell

The GPU kernel tracks a single u16 arrival per thread position. It cannot distinguish between rising and falling signal transitions because each thread processes 32 packed Boolean signals simultaneously — there’s no per-bit transition direction available.

How it works: For each cell, inject_timing_to_script() computes:

#![allow(unused)]
fn main() {
delay = max(gate_delays[pin].rise_ps, gate_delays[pin].fall_ps)
}

Impact: For the SKY130 inv_chain test (16 inverters), rise delays average ~10ps larger than fall delays. In a real inverter chain, transitions alternate (rise→fall→rise), so half the cells use the smaller fall delay. Jacquard uses the larger rise delay for all.

Measured: 80ps overestimate on 1235ps (6.5%) for 16 inverters with ~10ps rise/fall asymmetry per cell.

Source 2: max wire delay across all input pins

For multi-input cells (AND gates, MUXes), INTERCONNECT delays to different input pins may differ significantly. Jacquard takes the maximum across all input pins:

#![allow(unused)]
fn main() {
// wire_delays_per_cell: dest_cellid → max(all input wire delays)
entry.rise_ps = entry.rise_ps.max(ic.delay.rise_ps);
entry.fall_ps = entry.fall_ps.max(ic.delay.fall_ps);
}

Impact: If an AND gate has input A arriving via a 10ps wire and input B via a 200ps wire, Jacquard assigns 200ps to the cell regardless of which input is on the critical path. An event-driven simulator would correctly propagate the 10ps arrival on input A independently.

When this matters: Designs with highly asymmetric routing (e.g., one input is local, another crosses the chip). Well-routed designs typically have balanced wire delays to multi-input cells.

Source 3: max arrival across 32 packed signals per thread

Each thread position holds 32 independent Boolean signals. Jacquard tracks one arrival per thread position (the maximum across all 32 signals):

Thread 50: [signal_A: 5ps, signal_B: 23ps, signal_C: 8ps, ...]
Tracked:   arrival[50] = 23ps (max of all 32)

Impact: If signals with very different timing are packed into the same thread, the fastest signals inherit the slowest signal’s arrival time.

Mitigation: The bit-packing algorithm can sort signals by estimated timing before assignment (see “Timing-Aware Bit Packing” section). This keeps similar-timing signals together, reducing the max approximation error.

Combined Effect

These sources are multiplicative in the worst case. For the inv_chain test:

For larger designs with more routing asymmetry and denser bit packing, the combined overestimate could be larger. The bit-packing sort (Source 3) is the most actionable mitigation.

CVC Reference Validation

The inv_chain design (2 DFFs + 16 SKY130 inverters) was validated against CVC (open-src-cvc), an event-driven Verilog simulator with native SDF back-annotation:

CVC:  clk_to_q=350ps  chain=885ps  total=1235ps  (transition-accurate)
Jacquard: clk_to_q=350ps  chain=973ps  total=1323ps  (conservative max)
Difference: 88ps (7.1% overestimate)

Both simulators agree on CLK→Q delay (350ps) because the DFF has a single output transition direction per clock edge. The chain delay differs because CVC tracks actual rise/fall polarity through each inverter.

To run the CVC comparison locally:

bash tests/timing_test/cvc/run_cvc.sh

Requires Docker (builds CVC from source on first run).

Delay Data Encoding

Script Format

The existing boomerang section has padding that can store delay data:

Current format per thread per stage:
  [xora: u32]
  [xorb: u32]
  [orb:  u32]
  [padding: u32]  ← Can store delay here

PackedDelay Structure

#![allow(unused)]
fn main() {
#[repr(C)]
pub struct PackedDelay {
    pub rise_ps: u16,  // Rising edge delay in picoseconds
    pub fall_ps: u16,  // Falling edge delay in picoseconds
}
}

For simplified timing, a single uniform delay constant can be used instead of per-gate delays.

Timing Violation Detection

At Each Cycle Boundary

The GPU kernel checks timing constraints per state word (32 signals) after the boomerang evaluation completes. Arrivals and constraints use u16 picosecond values (range 0–65535 ps). Arithmetic is performed in u32 to avoid overflow when summing arrival + setup:

// After boomerang completes, before next cycle
// arrival: u16 max accumulated delay for this 32-signal group
// constraint_word: packed [setup_ps:16][hold_ps:16]
u16 setup_ps = constraint_word >> 16;
u16 hold_ps  = constraint_word & 0xFFFF;

// Setup check: skip when arrival == 0 (no data propagated, e.g. first cycle
// or DFF with constant inputs)
if (arrival > 0 && (u32)arrival + (u32)setup_ps > clock_period_ps) {
    int slack = (int)clock_period_ps - (int)arrival - (int)setup_ps;
    write_event(event_buffer, EVENT_TYPE_SETUP_VIOLATION,
                cycle, io_offset + threadIdx.x,
                (u32)slack, (u32)arrival, (u32)setup_ps);
}

// Hold check: no arrival > 0 guard (hold violations matter even at cycle 0)
if ((u32)arrival < (u32)hold_ps) {
    int slack = (int)arrival - (int)hold_ps;
    write_event(event_buffer, EVENT_TYPE_HOLD_VIOLATION,
                cycle, io_offset + threadIdx.x,
                (u32)slack, (u32)arrival, (u32)hold_ps);
}

Event Buffer Integration

#![allow(unused)]
fn main() {
pub enum EventType {
    Stop = 0,
    Finish = 1,
    Display = 2,
    AssertFail = 3,
    SetupViolation = 4,   // Timing events
    HoldViolation = 5,
}
}

For full details on interpreting violation reports and tracing violations to source signals, see docs/timing-violations.md.

Timing-Aware Bit Packing

The Problem

Each thread position holds 32 signals packed into a u32. When tracking timing with one arrival value per thread position, we approximate all 32 signals as having the same arrival time (the maximum).

This approximation is accurate when signals in the same thread have similar timing. But the default placement algorithm uses first-fit for bit assignment:

#![allow(unused)]
fn main() {
// Default: first available slot
for i in 0..hier[selected_level].len() {
    if hier[selected_level][i] == usize::MAX {
        slot_at_level = i;  // First-fit, not timing-aware
        break;
    }
}
}

This can result in signals with very different timing sharing a thread:

Thread 50 (accidental grouping):
  bit 0: level 5,  ~5ps arrival
  bit 1: level 12, ~12ps arrival  ← 7ps difference!
  bit 2: level 6,  ~6ps arrival

Thread 50 (timing-aware grouping):
  bit 0: level 5, ~5ps arrival
  bit 1: level 5, ~5ps arrival    ← similar timing
  bit 2: level 6, ~6ps arrival

Current Timing Correlation

The placement algorithm already computes logic levels:

#![allow(unused)]
fn main() {
// Level = max(level of inputs) + 1
level[node] = max(level[input_a], level[input_b]) + 1;
}

Logic level correlates with timing (more levels = more gate delays), but signals at the same level can still have different actual delays due to:

• Different gate types (AND2_00_0 vs AND2_11_1)
• Different wire loads
• Path reconvergence

Solution: Sort by Timing Before Packing

Before assigning bit positions, sort signals by their estimated arrival time:

#![allow(unused)]
fn main() {
// Collect nodes at this level
let mut nodes_to_place: Vec<_> = candidates
    .filter(|n| level[n] == selected_level)
    .collect();

// Sort by arrival time (level as proxy, or actual timing if available)
nodes_to_place.sort_by_key(|n| arrival_estimate[n]);

// Place in sorted order - similar timing ends up in same thread
for (slot, node) in nodes_to_place.iter().enumerate() {
    place_bit(..., slot, *node);
}
}

Alternative Approaches

Timing Bands

Group signals into arrival time bands:

Band 0: 0-10ps   → Threads 0-63
Band 1: 10-20ps  → Threads 64-127
Band 2: 20-30ps  → Threads 128-191
Band 3: 30+ps    → Threads 192-255

Measuring Packing Quality

Diagnostic to measure timing variance per thread:

#![allow(unused)]
fn main() {
fn analyze_timing_packing(hier: &Hierarchy, arrivals: &[u64]) {
    for thread in 0..256 {
        let times: Vec<_> = get_bits_in_thread(hier, thread)
            .map(|b| arrivals[b])
            .collect();

        let range = times.iter().max() - times.iter().min();
        let variance = compute_variance(&times);

        if range > threshold {
            warn!("Thread {} has {}ps timing spread", thread, range);
        }
    }
}
}

Impact on Approximation Accuracy

With timing-aware packing:

• Reduced false positives: Fewer spurious timing violations from max approximation
• Tighter bounds: Per-thread arrival closer to actual signal arrivals
• Better critical path identification: Max arrival more accurately reflects true critical path

Performance Expectations

The overhead is modest because:

1. Timing operations are simple (max, add)
2. Memory access pattern is identical
3. No additional synchronization needed
4. Same parallelism structure

References

• src/pe.rs - Partition executor and boomerang stage construction
• csrc/kernel_v1_impl.cuh - GPU kernel implementation
• src/flatten.rs - Script generation with timing data
• src/event_buffer.rs - GPU→CPU event communication
• src/liberty_parser.rs - Timing library parsing

Timing Violation Detection

See also: timing-correctness.md — forward-looking validation contract and timing IR requirements (in progress). The document below describes current behaviour.

Guide to enabling, reading, and debugging setup/hold timing violations in GEM.

Overview

Setup and hold violations occur when data arrives too late (setup) or too early (hold) relative to the clock edge at a flip-flop. GEM checks for these violations during GPU simulation by tracking arrival times — the accumulated gate delay from primary inputs or DFF outputs through combinational logic to the next DFF data input.

Approximation model: GEM tracks one arrival time per 32-signal group (one GPU thread position). The arrival is the maximum across all 32 signals in the group. This is conservative: it may over-report violations but will never miss a real one. See Reducing False Positives for details.

Enabling Timing Checks

Prerequisites

1. SDF file with back-annotated delays from your place-and-route tool
2. Gate-level netlist synthesized to aigpdk.lib cells

Step-by-step

1. Generate SDF from your P&R tool (or use scripts/generate_sdf.py for test designs):

   # Example: OpenROAD flow output
   ls my_build/6_final.sdf
   

2. Run the simulator with --sdf and a clock period:

   Metal (macOS):

   cargo run -r --features metal --bin jacquard -- sim \
       design.gv input.vcd output.vcd 1 \
       --sdf design.sdf \
       --sdf-corner typ
   

   CUDA (NVIDIA):

   cargo run -r --features cuda --bin jacquard -- sim \
       design.gv input.vcd output.vcd 8 \
       --sdf design.sdf \
       --sdf-corner typ \
       --enable-timing \
       --timing-clock-period 1200
   

   cosim (co-simulation):

   cargo run -r --features metal --bin jacquard -- cosim \
       design.gv \
       --config testbench.json \
       --sdf design.sdf \
       --sdf-corner typ
   

CLI Flags Reference

Example: inv_chain_pnr Test Case

# Run with SDF timing
cargo run -r --features metal --bin jacquard -- sim \
    tests/timing_test/inv_chain_pnr/6_final.v \
    tests/timing_test/inv_chain_pnr/input.vcd \
    tests/timing_test/inv_chain_pnr/output.vcd 1 \
    --sdf tests/timing_test/inv_chain_pnr/6_final.sdf

Reading Violation Reports

Setup Violation Format

[cycle 42] SETUP VIOLATION at top/cpu/regs[7][bit 22] [word=5]: arrival=900ps setup=200ps slack=-100ps

(WS-P1.1.a, 2026-05-02: state-word indices are now resolved to symbolic hierarchical signal names. The bare [word=N] suffix is preserved for grep compatibility. Words packing more than 4 DFFs truncate with a +N more suffix.)

Hold Violation Format

[cycle 11] HOLD VIOLATION at top/cpu/state[bit 3] [word=3]: arrival=10ps hold=50ps slack=-40ps

Summary Statistics

At the end of simulation, GEM prints totals:

Simulation complete: 1000 cycles, 5 setup violations, 0 hold violations

Text Summary (--timing-summary)

A one-screen human summary printed to stdout at end of run. Reuses the same data the JSON report builds (so --timing-report and --timing-summary cost the same; only the output channel differs). Sample output:

=== Jacquard Timing Summary ===
Design:        my_cpu.gv
Vectors:       boot.vcd (1000 cycles)
Clock period:  1000 ps
Timing source: my_cpu.jtir

Violations:
  Setup: 5
  Hold:  2
  Total: 7

Worst slack:
  Setup: -150ps  at top/cpu/regs[7][bit 22] [word=5]  (cycle 87)
  Hold:   -40ps  at top/cpu/state[bit 3] [word=12]  (cycle 91)
  Total negative: setup (TNS)=-620ps  hold (THS)=-80ps

Top 2 by violation count (of 2 total words with violations):
  top/cpu/regs[7][bit 22] [word=5] (5 violations): worst setup=-150ps hold=- arrival=950ps
  top/cpu/state[bit 3] [word=12] (2 violations): worst setup=- hold=-40ps arrival=10ps

The Worst slack block reports the four standard signoff metrics:

TNS/THS reflect every observed violation regardless of the per-cycle list cap, and are never optimistic: if the GPU event buffer overflows, dropped events only make the true totals more negative.

The format is for human inspection — explicitly not a stable parseable contract. Tools that need to script against the data should use --timing-report JSON.

Structured JSON Report (--timing-report <path.json>)

For CI integration and downstream tooling, pass --timing-report <path> to get an end-of-run JSON document. The schema is versioned (Decision 0008’s stability contract: additive-only extensions, breaking changes bump the major). Sample at tests/timing_ir/sample_reports/two_violations.json; authoritative type definitions in src/timing_report.rs.

Top-level shape:

{
  "schema_version": "1.0.0",
  "metadata": { "design": "...", "cycles_run": 1000, "clock_period_ps": 1000, "...": "..." },
  "stats": { "setup_violations": 5, "hold_violations": 0, "events_dropped": 0 },
  "violations": [
    { "cycle": 42, "kind": "setup", "word_id": 5, "site": "top/cpu/regs[7][bit 22] [word=5]",
      "arrival_ps": 900, "constraint_ps": 200, "slack_ps": -100 }
  ],
  "per_word": [
    { "word_id": 5, "site": "...", "setup_violations": 5, "hold_violations": 0,
      "worst_setup_slack_ps": -100, "worst_hold_slack_ps": null, "worst_arrival_ps": 900 }
  ],
  "worst_slack": {
    "setup": [ /* top-N most-negative slacks across the run */ ],
    "hold":  [ /* same shape */ ]
  }
}

per_word is sorted by total violation count desc, then by word_id. worst_slack.setup / .hold are top-10 by closest-to-violation slack (most negative first). Caveats:

• The “even when no violation occurred” half of WS-P1.1.d (per-DFF closest-to-violation tracking when the design never tripped a violation) needs GPU-side near-miss instrumentation and is not in v1.0.0; for now, worst_slack is populated only from actual violation events.
• --timing-report is wired on all three sim backends (Metal, CUDA, HIP). The cosim path does not yet route runtime violations through process_events — bringing it in is independent plumbing.
• The violations array is capped at 100,000 records by default (~8 MB JSON). Override or disable the cap with --timing-report-max-violations <N> (0 = unbounded). Setup/hold totals, events_dropped, and worst_slack rankings always reflect every observed event; only the per-cycle list is bounded. stats.violations_truncated reports how many records were dropped because the cap was reached.

Tracing Violations to Source Signals

When you see a violation on a specific word, follow this workflow to identify the offending signals and their logic cone.

1. Get the Word Index

From the log: word 5 means state word index 5.

2. Map Word to DFF Signals

Each word covers 32 bits of state. The DFFs in that word have data_state_pos / 32 == word_index. To find which DFFs:

• Look at the dff_constraints entries in the FlattenedScriptV1:

  dff_constraints entries where data_state_pos / 32 == 5
  → cell_id values → netlist cell names
  

• In gpu_sim, violations are logged with word IDs that map directly to the output_map positions. Each word covers bit positions word * 32 through word * 32 + 31.

3. Trace Backwards with netlist_graph

Use the netlist_graph tool to trace the combinational logic cone feeding the DFF. After uv sync --group dev, the netlist-graph console script is on the workspace’s uv run path — no cd required:

# Find the DFF data input driver chain
uv run netlist-graph drivers design.v "dff_name.D" -d 10

# Search for DFFs matching a pattern
uv run netlist-graph search design.v "dff_out*"

Discovered signal names can be passed directly into jacquard sim --trace-signals <file> / jacquard cosim --trace-signals <file> (one name per line) to surface them in the output VCD alongside top-level IO.

4. Detailed Timing Analysis with CVC

For per-signal accuracy (no 32-signal approximation), use CVC (open-src-cvc) with SDF back-annotation:

# Run CVC with SDF timing
cvc64 +typdelays tb.v design.v
./cvcsim

CVC provides event-driven simulation with full SDF support (IOPATH + INTERCONNECT delays), allowing you to pinpoint exactly which path is critical.

The Approximation Caveat

GEM tracks one arrival time per 32-signal group (one GPU thread position). The tracked value is the maximum arrival across all 32 signals in that thread. This means:

• Conservative: If any signal in the group has a long path, the arrival for the entire group reflects that worst case. Violations may be reported for signals that individually meet timing.
• Never misses real violations: A real violation always results in a reported violation (the max is >= any individual signal’s arrival).

Reducing False Positives

If a violation is reported but you suspect it’s a false positive from the approximation:

1. Use CVC for per-signal accuracy (see Detailed Timing Analysis with CVC above).
2. Timing-aware bit packing groups signals with similar arrival times into the same thread, reducing the approximation error. See docs/timing-simulation.md § “Timing-Aware Bit Packing” for details.

Common Scenarios

Setup violations on many words, same cycle: The clock period is likely too tight for the design. The combinational logic depth exceeds what can settle in one clock period. Try increasing the clock period.

Setup violation on a single word: A critical path through one specific logic cone. Use netlist_graph drivers to trace the path and identify the bottleneck.

Hold violation: Rare with SKY130 process (negative hold times clamp to 0 in the SDF). If seen, the design likely has minimum-delay paths that are too short. Check for direct connections between DFF outputs and nearby DFF inputs with minimal combinational logic.

Violations only on first cycle: The arrival > 0 guard in the GPU kernel skips setup checks when arrival is zero (meaning no data has propagated through combinational logic yet). If you see violations on cycle 0, they are hold violations — setup violations on cycle 0 are suppressed by design.

Timing Validation Methodology

Note: The ±5% tolerance convention described here will be superseded by timing-correctness.md once phase 0 ships (OpenSTA oracle + timing IR + diff harness). Content below remains accurate for current behaviour.

This document describes how GEM validates timing simulation accuracy against reference simulators (CVC, Icarus Verilog). It covers test cases, comparison metrics, known simulator differences, and acceptance criteria.

Overview

GEM’s timing simulation (--timed and --enable-timing flags) must be validated against independent reference simulators to ensure correctness. We use two reference tools:

• CVC (Cadence Verilog Compiler): For post-layout SDF-annotated designs (commercial tool, high confidence)
• Icarus Verilog: For structural Verilog and Liberty-based timing (open source, good for pre-layout)

What We Validate

1. Functional Correctness (Primary Validation)

Definition: Both GEM and reference simulators produce identical output values at each clock cycle.

Signals compared: Primary design outputs (e.g., gpio_out[43:0] for MCU SoC), NOT internal timing signals.

Tolerance: Exact match required. If output values differ, timing simulation has a correctness bug.

Tool: compare_outputs.py script compares VCD outputs cycle-by-cycle.

# Example: MCU SoC functional comparison
uv run tests/mcu_soc/cvc/compare_outputs.py \
    jacquard_output.vcd cvc_output.vcd \
    --skip-cycles 5  # Skip first 5 cycles (reset/initialization)

2. Timing Accuracy (Secondary Validation)

Definition: Gate-level delays and arrival times computed by GEM match reference simulator within acceptable margins.

Known differences (simulator-specific semantics):

Why the difference? GEM’s boomerang architecture computes cumulative delays through pipeline stages, resulting in conservative estimates that over-predict delays by including full combinational paths. CVC only reports final DFF gate delays.

Tolerance:

• Functional output: Exact match required
• Arrival times: ±5% acceptable (due to architectural differences)
• Setup/hold margins: ±10ps acceptable (platform-dependent rounding)

Test Cases

inv_chain_pnr (Simple Reference Case)

Location: tests/timing_test/inv_chain_pnr/

Description: Single inverted AND gate followed by a chain of inverters. Provides ground truth for:

• Basic gate delay accuracy
• Multi-stage combinational path accumulation
• Clock-to-Q propagation

Size: 9 cells, ~40ps full path delay

Validation:

# Generate CVC reference
cd tests/timing_test/inv_chain_pnr
cvc64 +typdelays tb_cvc.v inv_chain.v 2>&1 | tee cvc.log
./cvcsim 2>&1 | grep "RESULT:" # Extract: clk_to_q, chain_delay, total_delay

# Generate GEM output
cargo run -r --features metal --bin jacquard -- sim \
    6_final.v stimulus.vcd output.vcd 1 \
    --sdf 6_final.sdf \
    --timed

Expected results:

• Jacquard Q arrival ≈ CVC total_delay (within ±5%)
• Both simulators show monotonic delay increase with each inverter stage

MCU SoC post-layout (Complex Case)

Location: tests/mcu_soc/data/6_final.v + 6_final.sdf

Description: Full MCU SoC design with SKY130 cells, post-P&R netlist with SDF timing.

Size: 2.7k cells, ~19MB netlist, ~18MB SDF

Validation:

# In CI: .github/workflows/ci.yml mcu-soc-metal job

# 1. Strip SDF timing checks (remove malformed TIMINGCHECK directives)
uv run tests/mcu_soc/cvc/strip_sdf_checks.py \
    tests/mcu_soc/data/6_final.sdf \
    tests/mcu_soc/data/6_final_stripped.sdf

# 2. Generate Jacquard timing VCD
cargo run -r --features metal --bin jacquard -- sim \
    tests/mcu_soc/data/6_final.v \
    tests/mcu_soc/stimulus.vcd \
    tests/mcu_soc/jacquard_timed_mcu.vcd 1 \
    --sdf tests/mcu_soc/data/6_final_stripped.sdf \
    --sdf-corner typ \
    --timed \
    --max-clock-edges 10000

# 3. Generate CVC reference
cvc64 +typdelays tests/mcu_soc/cvc/tb_cvc.v \
    tests/mcu_soc/data/6_final.v \
    tests/mcu_soc/cvc/sky130_cells.v \
    2>&1 | tee cvc_compile.log
./cvcsim > cvc_output.vcd 2>&1

# 4. Compare functional outputs
uv run tests/mcu_soc/cvc/compare_outputs.py \
    tests/mcu_soc/jacquard_timed_mcu.vcd \
    tests/mcu_soc/cvc/cvc_output.vcd \
    --skip-cycles 5

Expected results:

• Functional output: Exact match (or documented difference with explanation)
• Arrival times: Jacquard values ≥ CVC due to conservative path accumulation
• CI: Comparison completes without errors

Pre-layout Library Timing

Location: tests/timing_test/sky130_timing/

Description: Synthesized designs with Library-only timing (no SDF from P&R). Pre-layout tests validate timing accuracy early, before place-and-route.

Test circuits:

1. inv_chain: DFF → 16 inverters → DFF

   • Expected combo delay: 16 × 28ps = 448ps (from Liberty inv_1 tpd)
   • Full path: CLK→Q (310ps) + chain (448ps) = 758ps

2. logic_cone: 4 input DFFs → nand2/nor2/and2 tree → DFF

   • Critical path: 4 gates × 35ps avg = 140ps
   • Full path: CLK→Q (310ps) + logic (140ps) = 450ps

Validation:

# Generate Liberty-only SDF
cd tests/timing_test/sky130_timing
python3 gen_liberty_sdf.py inv_chain.v
python3 gen_liberty_sdf.py logic_cone.v

# Run CVC reference
cvc64 +typdelays tb_inv_chain.v inv_chain.v 2>&1 | tee cvc.log
./cvcsim

# Run Jacquard GPU simulation
cargo run -r --features metal --bin jacquard -- sim \
    inv_chain.v stimulus.vcd jacquard_output.vcd 1 \
    --sdf inv_chain.sdf \
    --timed

# Compare outputs
uv run ../../mcu_soc/cvc/compare_outputs.py jacquard_output.vcd cvc_output.vcd

Expected results:

• Functional output: Exact match (logic correctness)
• Arrival times: Jacquard ≥ CVC (conservative accumulation expected)
• No SDF parsing errors or malformed delays

Purpose: Pre-layout tests catch timing issues early without P&R turnaround (30+ min). Library-only delays provide floor (lower bound); post-layout adds routing parasitics for final validation.

Multi-corner workflow (local Mac dev)

Location: crates/opensta-to-ir/tests/opensta_integration.rs — sky130_multi_corner_emits_per_corner_values.

The test drives opensta-to-ir --liberty NAME=PATH (WS2.4) three times with the genuine sky130 typ/slow/fast Liberty files and asserts the emitted timing IR carries three corners with different setup values (slow > typ > fast). This is the load-bearing check that --timing-corner actually means something — the companion aigpdk_dff_emits_per_corner_timing_values only validates IR shape.

One-time setup on a fresh Mac:

uv sync --group dev                                     # installs volare
uv run volare enable c6d73a35f524070e85faff4a6a9eef49553ebc2b
# Optional override if the PDK lives elsewhere:
export SKY130_LIBERTY_DIR=/path/to/sky130_fd_sc_hd/lib

The pinned volare hash lives in pyproject.toml::[tool.jacquard.pdks.sky130]. The test skips cleanly with an install hint when the PDK isn’t found, so it’s safe to run on machines without sky130 installed.

Running:

cd crates/opensta-to-ir
cargo test --test opensta_integration sky130_multi_corner -- --nocapture

GF180MCU PDK (in progress — phased per docs/plans/gf180mcu-enablement.md):

uv sync --group dev
uv run volare enable --pdk gf180mcu 559a117b163cef2f920f33f30f6f690aa0b47e4c

Liberty lands under ~/.volare/volare/gf180mcu/versions/<hash>/gf180mcuC/libs.ref/gf180mcu_fd_sc_mcu{7,9}t5v0/liberty/. Both 7-track and 9-track standard-cell libraries are pulled together by volare. Cell models are vendored as submodules at vendor/gf180mcu_fd_sc_mcu7t5v0/ and vendor/gf180mcu_fd_sc_mcu9t5v0/.

CI Integration

MCU SoC Timing Comparison Workflow

The main CI pipeline (.github/workflows/ci.yml) includes:

1. mcu-soc-metal job: Generates Jacquard timing VCD

   • Includes SDF stripping step (strip_sdf_checks.py)
   • Produces jacquard_timed_mcu.vcd with arrival time annotations

2. mcu-soc-cvc job: Generates CVC reference output

   • Uses stripped SDF (6_final_nocheck.sdf)
   • Produces cvc_output.vcd for comparison

3. mcu-soc-comparison job: Validates both produce same functional output

   • Runs compare_outputs.py
   • Reports pass/fail in CI summary
   • Skips gracefully if either simulator fails

SDF Stripping

SDF files from post-P&R tools may contain:

• Malformed TIMINGCHECK directives (setup/hold specs with syntax errors)
• INTERCONNECT entries with escaped port names that parsers reject

Solution: tests/mcu_soc/cvc/strip_sdf_checks.py removes:

• TIMINGCHECK blocks (parser errors, not needed for gate-level functional sim)
• INTERCONNECT lines (wire delays, optional)
• Empty DELAY blocks (after INTERCONNECT removal)
• CELL blocks with escaped $ in instance names

Result: Stripped SDF retains ~402k lines of useful IOPATH (gate delay) data, removing ~131k problematic lines.

Comparison Tolerances

Pre-Layout (Liberty-Only) Tests

Pre-layout tests compare designs without P&R parasitics, using only library cell timing:

Why loose tolerances? Pre-layout designs have no P&R uncertainty. Functional correctness is critical (exact match); timing metrics can vary due to simulator implementation differences.

Post-Layout (SDF-Annotated) Tests

Post-layout designs include routing delays and P&R context:

Known Issues & Limitations

1. Boomerang Path Accumulation

Issue: GEM sums gate delays across all pipeline stages leading to a node, producing cumulative arrival times. CVC only reports final DFF gate delay.

Example: inv_chain test:

• CVC reports: CLK→Q = 350ps (final gate only)
• GEM reports: Q arrival = 1323ps (CLK→dff_in.Q + inverter chain + interconnect)

Why: Boomerang architecture evaluates in hierarchical stages. To know Q arrival, you sum delays from all stages.

Mitigation: Compare Jacquard Q arrival against CVC RESULT: total_delay (if available), or accept that arrival times differ but functional output matches.

2. SDF Parser Robustness

Issue: Post-P&R SDF files from various tools contain edge-case syntax that doesn’t parse correctly.

Examples:

• Empty delay specs: (IOPATH A B () (0:0:0)) — treated as 0ps
• COND-qualified pins: (SETUP (COND x==1 D) (posedge CLK) (180:200:220)) — not yet supported
• Malformed TIMINGCHECK: cause parser errors in byte-range 18M+

Current approach: Strip SDF timing checks before use, preserving IOPATH data.

Future: Improve SDF parser to handle more edge cases without stripping.

3. Timing Model Accuracy

Conservative model: GEM accumulates delays pessimistically. This is intentional—designs validated under the GEM model are guaranteed to meet timing in actual P&R.

Trade-off: Over-predicting delays means GEM may flag timing violations that won’t occur in silicon. This is preferable to under-predicting (which would give false confidence).

Debugging Timing Failures

When Functional Output Doesn’t Match

1. Check testbench stimulus: Does both simulators receive same input?

   • Generate stimulus VCD: --stimulus-vcd stimulus_out.vcd
   • Compare against reference testbench

2. Check SDF parsing: Did parser successfully load all timing data?

   • Enable debug logging: --sdf-debug
   • Look for “unmatched SDF instances” warnings

3. Check initialization: Are both simulators starting from same state?

   • Compare first 5-10 cycles after reset
   • Verify reset logic is synchronized

4. Compare against Jacquard non-timed version:

   # Run without timing VCD
   cargo run -r --features metal --bin jacquard -- sim \
       design.v stimulus.vcd functional_output.vcd 1
   
   # Does functional output match CVC?
   compare_outputs.py functional_output.vcd cvc_output.vcd
   

When Arrival Times Seem Wrong

1. Verify SDF was loaded: Check logs for “Failed to load SDF” warnings
2. Check corner selection: Confirm --sdf-corner typ|min|max matches reality
3. Compare against CVC RESULT lines: If available, extract CVC timing measurements
4. Check for pipeline stage effects: Arrival time = sum of delays across stages

Acceptance Criteria

References

• Timing Simulation in GEM — Architecture details
• Timing Violation Detection — Setup/hold checks
• SDF Parser — Parsing implementation
• CVC Integration — Reference simulator setup

Timing-model extensions — design notes

Status: Idea / pre-spike. Not scheduled. Captured here so the architecture sketch survives the next session-clear.

Scope: Three related extensions to Jacquard’s timing model, all aimed at making setup/hold reporting more honest without abandoning the cycle-accurate boomerang kernel.

1. Dynamic delay — per-gate δ(T) inspired by the Involution Delay Model (Maier 2021, arXiv:2107.06814). Captures pulse-width-dependent delay degradation that fixed δ∞ misses on near-threshold paths.
2. Clock-tree skew — per-DFF clock arrival accounting. Today every DFF on a clock is treated as if it captures simultaneously; SDF clock-buffer arcs and clock-net interconnect are silently dropped during AIG construction.
3. Wire delay at scale — per-receiver interconnect delay applied to the right edge in the AIG, and explicit modelling of inter-partition wires. Today wire delay is collapsed to a max-per-destination-cell scalar — fine for sky130 short routes, increasingly wrong as we move to faster clocks, finer processes, and large many-core/NoC designs.

All three share the same insight: the data the model needs is already in the TimingIR. The work is at the consumer layer (flatten.rs, aig.rs, the kernel arrival math), not the IR or the partitioner.

Background — what the timing pipeline does today

.sdf  ─┬─► opensta-to-ir ──► TimingIR (.jtir, FlatBuffers)
.jtir ─┘                          │
                                  ▼
              flatten.rs::load_timing_from_ir   (per-cell arc → AIG-pin delay)
                                  │
                                  ▼
                       gate_delays: Vec<PackedDelay>     (rise/fall ps per AIG pin)
                       dff_constraints: Vec<DFFConstraint>  (setup/hold ps per DFF)
                                  │
                                  ▼
              flatten.rs::inject_timing_to_script   (bake max ps into u16 script slot)
                                  │
                                  ▼
                       kernel_v1.metal at runtime:
                       per-AND:    new_arr = max(arr_a, arr_b) + gate_delay
                       per-DFF:    check arrival vs setup/hold per word

Reference points:

• IR schema: crates/timing-ir/schemas/timing_ir.fbs
• IR consumer: src/flatten.rs:1768 (load_timing_from_ir), src/flatten.rs:1686 (inject_timing_to_script)
• Setup/hold buffer: src/flatten.rs:1732 (build_timing_constraint_buffer)
• GPU arrival math: csrc/kernel_v1.metal:220-255 (AND gates), csrc/kernel_v1.metal:547-580 (setup/hold)

Per-AIG-pin arrival is a single ushort accumulated by max through the boomerang reduction. There is no event scheduling — arrival is a scalar that rides alongside the Boolean evaluation in lockstep with cycle ticks.

Part A — Dynamic delay (IDM-style δ(T))

What IDM is, briefly

A per-gate dynamic delay model that makes δ a function of T (time since the gate’s last output transition). The distinguishing property: input pulses with Δᵢ → 0 have diminishing effect on the output. The model handles pulse-width degradation faithfully and is the only model proven to solve the short-pulse-filtration problem. The paper notes ~80–590% CPU overhead vs. inertial delay on a CPU event-driven simulator.

The architectural wall

True IDM needs event scheduling and intra-cycle pulse observability — neither is available in Jacquard’s lockstep cycle-accurate kernel. We cannot model glitch suppression or metastability oscillation traces without either sub-cycle ticks or a different kernel architecture.

What we can do is enrich the per-gate delay used in arrival propagation so setup/hold reporting reflects realistic pulse degradation on marginal paths.

Five hook points

For dynamic-IDM the kernel needs two new persistent buffers:

• last_transition_ps[aig_pin] — when the gate’s output last switched (absolute ps).
• last_value[aig_pin] — to detect transitions across cycles.

Memory cost ~4 bytes per AIG pin per partition. For NVDLA-scale designs (~hundreds of thousands of pins) this is MB-scale — fine.

eval_idm on GPU

The paper uses exp/log per gate. On GPU replace with a 16-entry LUT indexed by quantised T. Cheap, branch-free, smooth enough.

Characterisation

The δ(T) parameters have to come from per-cell SPICE characterisation. For sky130 we’d characterise each sky130_fd_sc_hd__*_* cell once, check the result into the repo, and ship it as a sidecar table consumed by the IR builder. This is the expensive one-off — the paper flags characterisation cost as the unsolved part of making IDM “truly competitive.”

Staged plan

Stage 1 is a 1–2 day spike with no kernel risk. Stage 2 is the honest implementation. Stage 3 is a different simulator.

What we get / don’t get from dynamic δ(T)

Achievable

• Per-corner δ(T) propagating through arrival → setup/hold reports that distinguish “just meets timing under δ∞” from “fails under realistic pulse degradation”.
• Stays inside cycle-accurate boomerang. ~1.5–2× memory growth on arrival data, ~10–20% kernel slowdown (estimate).

Not achievable

• Glitch suppression (Δᵢ → 0 → no transition).
• Metastable oscillation traces.
• Combinational-loop behaviours (loops are forbidden in the AIG anyway).

Why sky130 is the right vehicle

sky130_pdk.rs decomposes vendor functional Verilog into AIG nodes while preserving cell identity through aigpin_cell_origins. We can attach δ(T) at the original sky130 cell granularity even after AIG flattening — that structural property is what makes any of this tractable. Cells from a hand-coded library without origin tracking would be much harder.

Part B — Clock-tree skew

Status: Stages 1 + 2 implemented (2026-05-01). Per-DFF clock arrival is carried through the IR (ClockArrival table) and folded into per-DFF setup/hold via DFFConstraint::effective_setup_hold before the per-word collapse. Producer landed in c403cc8; consumer fold-in in 6767c3e. The narrative below describes the original motivation; the Staged plan at the end of this part records what shipped and what remains (Stage 3, conditional).

Where the information is — and where we drop it

Clocks in Jacquard are walked back from each DFF through buffers/inverters/clock-gates, terminating at an InputClockFlag(pinid, is_negedge) (src/aig.rs:441, :477, :495-560). Recognised cells: INV/BUF/CKLNQD and the sky130 equivalents inv*, clkinv*, buf*, clkbuf*, clkdlybuf*, lpflow_*.

Two consequences:

1. Clock-tree cells produce no AIG pin. They collapse into a polarity flag on the DFF. Since aigpin_cell_origins only lists cells that produced AIG pins, the timing-IR arcs on those cells (IOPATH records on clkbuf_8, etc.) match no AIG pin in load_timing_from_ir and are silently discarded.
2. Clock-net interconnect is dropped the same way. interconnect_delays records keyed by net endpoints have no destination cell to attach to, so they fall on the floor.

Net effect: every DFF on a given clock domain is treated as having identical clock arrival, i.e. perfect skew. The current setup/hold check is honest about combinational-path delay but blind to clock-tree topology.

For a sky130 MCU SoC at ~25 ns clock period this is fine functionally; for any timing claim near the period boundary it’s misleading. Intra-domain clock-tree skew on sky130 is typically O(50–200 ps) — small relative to a 25 ns period, but exactly the order of magnitude that determines whether a path “barely meets” or “barely fails” setup.

Do we have the information?

Yes, in three places, in increasing fidelity:

1. TimingIR arcs on clock cells (.jtir already contains them; we just don’t consume them).

2. The AIG clock walk in aig.rs:495–560 already iterates the clock-side cells of each DFF in order. It just doesn’t accumulate their delays. Adding a dff_clock_origins: Vec<Vec<cellid>> parallel structure costs O(num_dffs × clock_depth) memory — negligible.

3. OpenSTA can compute per-DFF clock arrival end-to-end. (OpenTimer was the original primary STA candidate per Decision 0003 but the spike Superseded it; Decision 0001 makes OpenSTA the sole STA path, called out of process via opensta-to-ir.) Per-pair common-path-pessimism removal (CRPR) is fundamentally a launch/capture credit, not a per-DFF property — so what shipped is per-DFF capture-side arrival, treating launch as a 0-reference. This is the form in the IR today:

   table ClockArrival {
       cell_instance: string;     // DFF instance path
       clk_pin: string;           // local pin name
       arrival: [TimingValue];    // per-corner clock arrival ps
       provenance: Provenance;
   }
   

   Populated by opensta-to-ir’s Tcl driver via [all_registers -clock_pins] + [::sta::vertex_worst_arrival_path]. Consumer code never touches the netlist — it just looks up each DFF’s clock arrival.

Consumer change (shipped)

DFFConstraint carries the field now:

#![allow(unused)]
fn main() {
pub struct DFFConstraint {
    pub setup_ps: u16,
    pub hold_ps: u16,
    pub clock_arrival_ps: i16,   // signed — capture-side arrival, launch ref = 0
    pub data_state_pos: u32,
    pub cell_id: u32,
}
}

The setup/hold formula for per-pair skew is:

• Setup margin = (clock_period + clock_arr_capture - clock_arr_launch) - data_arrival - setup
• Hold margin = data_arrival - (clock_arr_capture - clock_arr_launch + hold)

Per-launch/per-capture pairing is awkward in the current per-word-collapsed constraint buffer, so the implementation folds the capture-side clock arrival into the per-DFF effective setup/hold before packing, via DFFConstraint::effective_setup_hold:

• effective_setup = setup - clock_arrival_capture (clamped to [0, u16::MAX])
• effective_hold = hold + clock_arrival_capture (clamped to [0, u16::MAX])

The GPU kernel runs unchanged — the same packed (setup<<16)|hold word it already consumes now carries skew-aware values. Launch arrival is treated as zero (ref) — pessimistic for paths whose launch DFF also has a long clock path, but a clean first cut. Stage 3 below addresses that pessimism if measurement justifies it.

Partitioning question

“could we partition a design effectively to do this somewhat accurately without sacrificing too much?”

Today partitioning (src/repcut.rs) is hypergraph-cut on logic connectivity. DFFs co-located by logic affinity may have very different clock arrivals.

The pessimism cost: build_timing_constraint_buffer collapses all DFFs in a 32-bit state word to min(setup) and min(hold). If a word holds DFFs with clock arrival 50 ps and 200 ps, the per-word effective setup is the worst of both — i.e. we report timing as if every DFF in that word saw the worst skew in the word. That’s a 150 ps pessimism for the lucky DFF.

Three options, ranked:

1. Do nothing. For typical sky130 SoCs at ≥10 ns clock periods, intra-word skew (≤200 ps worst-case) vs. period (10 000+ ps) is ≤2%. Worth-it threshold for the optimisation: when designs run close enough to the period that 2% pessimism flips genuine passes into reported violations. Likely never for sky130. Plausibly relevant for designs running at ≥1 GHz on a more aggressive PDK.

2. Skew-bucket the DFF constraint packing, not the partitioning. Group DFFs into clock-arrival buckets after partitioning, and emit one constraint word per bucket-within-partition rather than collapsing everything in the word. Increases constraint-buffer size by O(num_buckets) but doesn’t disturb the partitioner. Probably the right answer if we ever need to.

3. Skew-aware partitioning. Add a soft objective to repcut.rs that prefers grouping DFFs by clock arrival. Degrades cut quality (more inter-partition logic edges → more state shuffling). Almost certainly worse than option 2 for the same accuracy gain.

So: yes we have the info, no we probably don’t need to repartition, and the constraint-collapsing pessimism is the real lever — either accept it (option 1) or break it bucket-wise (option 2).

Staged plan for clock tree

Part C — Wire delay at scale

Why this gets more important as designs grow

In sky130 at 25 ns clock periods, wire delay is a small perturbation on gate delay and the lumped model is fine. The picture changes in two regimes:

• Faster clocks. Wire delay is a fixed physical quantity (RC-dominated); period shrinks; wire fraction of the budget grows.
• Finer processes (e.g. 22nm and below). Gate delays scale down with feature size; wire RC scales unfavourably (resistance per square goes up, capacitance per length stays roughly flat). The classic “reverse scaling” inflection: gates get faster, long wires don’t. Typical 22nm: inverter delay 5–15 ps, local short wires 5–20 ps, global routes 50–500 ps, multi-mm wires 1+ ns without repeaters.
• Large many-core/NoC SoCs. Inter-tile mesh links can span multiple millimetres; chip-level signals have wire delays comparable to or larger than entire combinational stages.

For a many-small-core NoC at 22nm, wire delay on inter-core links is typically the dominant timing factor. Any model that can’t represent it accurately will misreport the critical paths.

What Jacquard does today

The IR side is already in shape. crates/timing-ir/schemas/timing_ir.fbs carries InterconnectDelay { net, from_pin, to_pin, delay[corner] } per receiver, and opensta-to-ir populates it from SDF.

The lossy step is the consumer in src/flatten.rs:1850-1872:

#![allow(unused)]
fn main() {
let mut wire_delays_per_cell: HashMap<usize, (u64, u64)> = HashMap::new();
// ... for each InterconnectDelay record:
let entry = wire_delays_per_cell.entry(dest_cellid).or_insert((0, 0));
entry.0 = entry.0.max(d);   // rise
entry.1 = entry.1.max(d);   // fall (same value!)
}

Three layers of pessimism stacked here:

1. Keyed by destination cell, not destination pin. A cell with two inputs from very different routes loses per-pin fidelity.
2. Max across inputs of the same cell. Worst-case incoming wire is applied to every output of the cell.
3. No rise/fall distinction on wire delay. SDF carries both; we collapse to one number.

Then in arrival propagation (csrc/kernel_v1.metal:220-255):

new_arr = max(arr_a, arr_b) + gate_delay

where gate_delay = intrinsic + max_wire_into_cell. The mathematically correct propagation is:

new_arr = max(arr_a + wire_a, arr_b + wire_b) + intrinsic

These are equivalent only when the input with the worst arrival also has the worst wire. When they don’t coincide — common on a NoC node where one input comes from a long mesh hop and another from local logic — the current model over-reports by max_wire − actual_wire_on_critical_input.

For sky130 small designs this gap is in the noise. For 22nm with 10× variation between local and global wire delays, it’s the difference between “this path meets timing” and “STA reports a violation that doesn’t exist.”

Inter-partition wires — the architectural wrinkle

A NoC tile naturally maps to one (or a few) partition(s). The inter-tile links — the long, wire-dominated, timing-critical ones — are precisely the partition-crossing signals. Today wire delay sits on the destination cell’s gate_delays slot, evaluated inside the destination partition’s boomerang reduction. The wire is a property of the crossing, not the destination cell, and should ideally be modelled at the partition I/O boundary, where src/sim/cosim_metal.rs already shuffles state between partitions.

This is the inverse alignment of the clock-tree case. There partitioning didn’t help with skew accounting. Here partitioning is load-bearing: tile-aligned partitions naturally expose the small set of edges that deserve careful wire-delay modelling, and let intra-partition logic stay on the fast lumped path.

Three fidelity tiers

Tier 1 is mostly a flatten.rs rewrite. Tier 2 needs cosim_metal.rs extension and a new field in the inter-partition transfer format.

Information availability

Yes, it’s there:

• InterconnectDelay records exist per receiver. SDF carries them. opensta-to-ir emits them.
• Per-input-pin granularity is in the IR (to_pin includes the local pin name). The consumer just discards it via to_pin.rfind('/') to derive dest_inst.
• Rise/fall distinction is in the schema (delay: [TimingValue] per corner; rise/fall could be on top via the same pattern as TimingArc). For SDF-back-annotated flows the rise/fall split usually comes from the SDF; we’d need to confirm opensta-to-ir preserves both edges.

What’s missing today:

• Tier-1 plumbing: AIG-pin-level wire delay per fanout. Current gate_delays: Vec<PackedDelay> is keyed by AIG pin (the output side); to do per-input-edge correctly we want delay attached to the edge, not the node. Either add a parallel wire_delays: HashMap<(src_aigpin, dst_aigpin), PackedDelay> or refactor toward an edge-attributed AIG.
• Tier-2 plumbing: a “partition-crossing arc” concept in cosim_metal.rs. Currently inter-partition state shuffle moves bits with no associated arrival bump. Adding a per-edge ps adjustment is straightforward in principle; finding the right place in the shuffle pipeline matters.

IR scale

The IR-size concern bites here. InterconnectDelay is roughly 100–200 bytes per record; a 22nm SoC with 10⁶–10⁷ nets is a .jtir file in the hundreds-of-MB to multi-GB range.

Mitigations:

• Streaming load: today TimingIrFile::from_path reads the whole buffer. Could mmap and lazy-decode, since FlatBuffers is offset-based.
• Sharding: split IR per partition or per top-level module. Adds a build-time step but bounds memory per process.
• Drop intra-cell wires from IR generation: SDF often has microscopic interconnect records that lump into the destination’s own pin-cap. Filter these out at the opensta-to-ir builder. Loss is genuinely negligible.

Worth measuring before committing to mitigations — sky130 NVDLA-scale today is fine; the question is what 22nm + N-tile mesh looks like.

Partitioning question — the other direction

For NoC designs partitioning becomes a positive lever (unlike the clock-tree case where it was neutral). Two specific levers:

1. Tile-aligned partitions. If repcut.rs finds tile-aligned cuts naturally (likely, given typical tile-to-tile connectivity sparsity), inter-partition arcs are a small, well-defined set of NoC links. Worth verifying with a representative design — a partitioning report keyed by signal name pattern (*_link_*, noc_*, configurable) would expose whether the partitioner’s logic-affinity score is already aligned with tile boundaries or whether we need to bias it.
2. NoC-link partitioning hint. Add a soft bias to repcut that prefers cutting nets matching a configured regex. Same partitioning machinery, configurable input. Cost: degrades cut quality if the hint conflicts with logic affinity. Likely worth it for explicitly tile-decomposed designs where the user knows the tile boundaries; not worth it for flat designs.

The point of any of this is to make Tier-2 cheap: if the inter-partition arc set is small, per-edge wire delay on those crossings costs almost nothing.

Crosstalk and OCV

These are upstream concerns. SDF from a crosstalk-aware STA flow already carries pessimistic delays; OCV (on-chip variation) is similarly baked into the chosen corner. Jacquard consumes whatever the IR was generated against. Worth a one-line note in the user-facing docs that the timing report’s accuracy is bounded by the SDF/STA flow it was built from — Jacquard does not invent crosstalk pessimism.

Staged plan for wire delay

For a sky130 use case Stage 1+2 likely covers everything you’d notice. For 22nm NoC, Stages 1–3 are the meaningful set; Stage 5 is the optimisation that makes Stage 3 cheap.

What we get / don’t get

Achievable

• Setup/hold accuracy on long routes that today gets clobbered by max-collapse pessimism.
• Honest reporting on NoC inter-tile links — the paths that actually matter for many-core SoC timing closure.
• All of the above without changing Jacquard’s cycle-accurate kernel architecture.

Not achievable from this work alone

• Crosstalk-driven delay uncertainty (handled upstream in STA).
• Variation-aware (statistical) timing — would need OCV-corner sweeping or SSTA, neither of which is on the roadmap.
• Process variation modelling beyond the corners the SDF/IR was generated against.

Open questions

1. δ(T) characterisation cost. One-off SPICE per cell-type per corner. Cheaper if we lean on existing ECSM/CCSM data already in vendor Liberty rather than re-running SPICE. Worth investigating before committing to Stage 2.
2. Whose clock arrival is authoritative? Resolved by Pillar B Stage 1+2: OpenSTA-computed per-DFF arrival via opensta-to-ir, treating launch as 0-reference. Per-pair CRPR credit is intentionally not modelled at this stage (see Stage 3 in the staged plan above).
3. Interaction. Does δ(T) on clock-tree buffers matter? Probably not enough to model — clock buffers are sized for fast edges and operate far from their pulse-degradation regime. But the framework should be able to express “ignore δ(T) on clock domain” cleanly.
4. Validation oracle. CVC and Icarus already serve as functional oracles; for skew-aware and wire-aware reporting OpenSTA’s slack report (via opensta-to-ir / direct subprocess) is the ground truth for unit tests. (Decision 0003 originally nominated OpenTimer for this role; superseded by the spike outcome — OpenSTA carries the role end-to-end now.)
5. IR size at 22nm scale. Open question whether .jtir for a representative many-core NoC fits in available memory under the current eager-load model. Needs measurement before committing to streaming mitigations.
6. Edge-attributed AIG. Per-receiver wire delay wants delay attached to AIG edges, not nodes. Today the AIG is node-attributed (gate_delays: Vec<PackedDelay> indexed by aigpin). A clean Tier-1 implementation may push toward edge attribution, with downstream effects on the boomerang reduction script layout. Worth a small spike before the main implementation.
7. Partition-crossing format. Adding per-edge wire delay to cosim_metal.rs inter-partition transfers needs a precise place in the existing pipeline. Currently the shuffle moves Boolean state words without arrival; the natural place is alongside the writeout-arrival path that already exists for setup/hold checking, but the alignment isn’t 1:1 because partition crossings happen at logic boundaries, not capture-DFF boundaries.

Related artefacts

• docs/timing-correctness.md — forward-looking validation contract; this doc extends rather than replaces.
• docs/timing-simulation.md — boomerang architecture; the kernel-side context.
• docs/timing-validation.md — current ±5% acceptance criteria; would tighten under δ(T).
• docs/architecture/decisions/0002-timing-ir.md — IR design rationale; schema additions here follow the “lossless extension” principle.
• docs/architecture/decisions/0001-opensta-as-oracle.md — STA path; OpenSTA out of process is committed (post-supersedure of Decision 0003).
• docs/architecture/decisions/archive/0003-opentimer-primary-sta.md — Superseded. Original in-process STA proposal; spike Q2 fail moved Jacquard to OpenSTA-only. See docs/architecture/decisions/spikes/opentimer-sky130.md.

In-Design Signal Tracing (--trace-signals)

Overview

By default a Jacquard output VCD contains only top-level IO. --trace-signals <FILE> surfaces user-selected internal nets in that VCD alongside the top-level ports — so you can watch a DFF’s Q, a controller state bit, or an SRAM port wire without re-synthesizing or exposing it as a port.

It is available on both jacquard sim and jacquard cosim, and is observe-only: traced nets are read out each tick, never driven.

Each name in the file is resolved against the netlist, registered as a primary output before partitioning (so it gets a state-buffer slot), and emitted on the same path as the top-level IO. It works uniformly for sequential (DFF Q) and combinational nets — anything that has a name in the netlist database.

This is the raw-wire counterpart to bus transaction tracing: --trace-signals gives you per-cycle waveforms of individual nets; bus tracing gives you decoded transaction records. Use this when you want a waveform; use bus tracing when you want READ 0x40 => 0x1.

File format

One hierarchical signal name per line:

# JTAG debug-module state (comments and blank lines are ignored)
chip_core.dm.haltreq_q[0]
chip_core.dm.haltreq_q[1]

# Yosys-internal nets — same syntax works
chip_core.sram_u._00147_

# A whole bus, one bit per line
data0_obs[0]
data0_obs[1]

• Blank lines and lines whose first non-whitespace character is # are skipped.
• Hierarchy uses . as the separator; a trailing [N] selects a bus bit.
• A leading backslash (Verilog escaped-identifier syntax) is stripped.

A real example ships in tests/jtag_minimal/trace_signals.txt.

Name resolution

Post-synthesis net names are ambiguous — Yosys may flatten a hierarchy into one escaped identifier (\soc.sram.read_port__data), expand a bus into per-bit scalars (soc.bus__addr[3]), or preserve real structural hierarchy. Rather than guess, the resolver tries multiple candidate interpretations of each name and takes the first that matches the netlist database, so the same syntax works across all three conventions.

• Unresolved names warn, they don’t abort. A bad name logs a warning and is skipped; the rest of the list still registers. A trailing summary line reports how many signals registered vs. were dropped, so a mistyped list surfaces clearly at startup:

  --trace-signals: registered 34 signal(s), dropped 2 (file: trace.txt)
  

• Names that resolve to a constant (tied 0/1) are skipped — there’s nothing to observe at runtime.

Where the output lands

Traced nets appear in whichever VCD the run already emits:

They show up as ordinary VCD wires next to the top-level IO, named by the string you put in the trace file.

cosim: traced nets land in --output-vcd only. The --stimulus-vcd carries primary inputs and does not include them, so if you trace a net and look in the stimulus VCD you’ll see nothing. --output-vcd does not require timing data — see Pre-PnR functional runs.

Pre-PnR functional runs

--output-vcd is the functional output path too — it does not require --timing-ir/SDF. Run a synthesized (pre-PnR) netlist through cosim with --output-vcd out.vcd and you get chip outputs and traced nets per cycle, with transitions at clock edges (no arrival-time offsets). This is the right mode for functional / 4-state X-pessimism debugging, where there is no timing data to supply yet. Adding --timing-ir later only adds arrival-time offsets to the same VCD.

Top-level inout (bidir) pads

A top-level inout pad is split into two observables in the output VCD: <pad>__out (the value the core drives) and <pad>__oe (the pad’s output-enable). The raw <pad> net reads the pad’s input side, so on an output-only or undriven cycle it can look flat — watch <pad>__out / <pad>__oe to see what the design is driving. Example: bidir_PAD[12]__out, bidir_PAD[12]__oe. These appear automatically; you don’t need to list them in the trace file.

Finding signal names

Use the netlist-graph tool (see the project README) to discover the exact post-synthesis names:

# Search for nets matching a pattern
uv run netlist-graph search <netlist.v> "haltreq"

# Trace what drives / loads a signal (to find nearby observable nets)
uv run netlist-graph drivers <netlist.v> "soc.cpu.state" -d 5
uv run netlist-graph loads   <netlist.v> "soc.cpu.ack"   -d 5

# Emit a ready-to-use trace file
uv run netlist-graph watchlist <netlist.v> out.json signal1 signal2 ...

SRAM observability workflow

The recommended way to observe SRAM port activity is wire-level tracing rather than the env-var-gated JACQUARD_SRAM_DUMP. netlist-graph can discover the port wires and emit a trace file directly:

# 1. Discover SRAM port wire names from the netlist
uv run netlist-graph sram-ports design.v --cell-type SRAM -o sram_trace.txt

# 2. Surface them in the VCD with full per-tick accuracy
jacquard cosim design.v --config sim.json \
    --trace-signals sram_trace.txt --output-vcd out.vcd

# 3. Post-process the VCD to reconstruct bus values

Example

tests/jtag_minimal/ uses --trace-signals to surface the debug module’s observable outputs (dmactive_obs, haltreq_obs, data0_obs[0..31]) so the test’s pass criterion can check that the magic value 0xCAFEBABE lands in data0_obs:

jacquard cosim tests/jtag_minimal/data/top.pnl.v \
    --config tests/jtag_minimal/sim_config.json \
    --trace-signals tests/jtag_minimal/trace_signals.txt \
    --jtag-replay tests/jtag_minimal/data/bitbang.rec \
    --output-vcd out.vcd

Troubleshooting

Implementation notes

Registration happens at AIG construction, before partitioning, which is why the list must be supplied via the CLI flag (not a runtime env var). The mechanism lives in src/sim/trace_signals.rs; emission piggybacks on emit_extra_observables in src/sim/vcd_io.rs. The same multi-candidate resolver backs bus-trace pin binding (see bus tracing and Decision 0013).

Bus Transaction Tracing (AHB / APB)

Overview

jacquard cosim can decode on-chip bus transactions and emit them in a compact, transaction-level form — one row per transfer, rather than raw per-cycle waveforms. You declare the bus interfaces to watch in sim_config.json; cosim observes their pins on the GPU each tick and runs the protocol decode on the CPU, writing decoded transactions to a CSV file.

This is observe-only: the tracer watches signals the design already drives, it never drives anything. It adds no measurable simulation overhead when no buses are configured.

The design rationale lives in Decision 0013; the roadmap is in plans/bus-transaction-tracing.md.

Bus tracing is the structured, protocol-aware counterpart to --trace-signals, which surfaces raw internal nets in the output VCD. Use --trace-signals when you want waveforms of individual wires; use bus tracing when you want decoded READ 0x40 => 0x1 records.

Configuring a bus

Add a bus_traces array to sim_config.json. Each entry names one bus interface:

{
    "netlist_path": "build/soc.gv",
    "clock_gpio": 0,
    "reset_gpio": 1,
    "num_cycles": 100000,
    "clock_period_ps": 40000,

    "bus_traces": [
        {
            "name": "dmi",
            "protocol": "apb3",
            "prefix": "soc.dm.",
            "addr_bits": 9,
            "data_bits": 32
        }
    ]
}

Pin names

By default each protocol pin is resolved as {prefix}{pin}. For APB3:

So a bus with "prefix": "soc.dm." looks for soc.dm.psel, soc.dm.paddr[0], …, soc.dm.prdata[31].

If your design’s pins don’t follow that convention, remap individual logical pins with signals:

{
    "name": "periph",
    "protocol": "apb3",
    "prefix": "soc.apb.",
    "signals": {
        "psel": "soc.apb_decode.sel_periph",
        "prdata": "soc.apb_mux.readback"
    }
}

Running

cargo run -r --features metal --bin jacquard -- cosim \
    build/soc.gv \
    --config sim_config.json \
    --bus-trace-csv bus.csv

At startup each bus logs whether it resolved:

bus-trace `dmi` (APB3): psel/penable resolved, addr 9/9 bits, pready=true pslverr=true

and at the end:

bus-trace: decoded 12 transaction(s) across 1 bus(es)
bus-trace: wrote 12 transaction(s) to bus.csv

CSV output

tick,bus,protocol,dir,addr,data,resp,burst
24,dmi,apb3,WR,0x10,0xCAFEBABE,OK,
30,dmi,apb3,RD,0x10,0xCAFEBABE,OK,

Pin resolution

For the GPU to read a bus pin each tick, that net must (1) exist in the post-synthesis netlist under a resolvable name and (2) survive into the simulation’s output state. Two consequences:

• Names must survive synthesis. The resolver uses the same multi-candidate matcher as --trace-signals, so Yosys-flattened (\soc.dm.psel), scalar-expanded (soc.dm.paddr[3]), and structurally-hierarchical names all work. But synthesis is free to rename or delete combinational nets. The robust pattern is to make the bus signals registers (their DFF Q outputs keep their names), or to annotate the RTL nets with (* keep *).

• Constant-folded bits read as 0 — correctly. If a design only ever drives, say, addresses 0x00 and 0x04, synthesis folds every address bit except paddr[2] to a constant. The startup log then shows e.g. addr 1/8 bits. This is expected: the tracer reconstructs the full value correctly because the dropped bits are genuinely 0.

pready / pslverr are allowed to be absent. A common case is an always-ready slave that ties pready high — it folds to a constant, fails to resolve, and the tracer correctly treats the bus as always-ready.

Worked example

tests/apb_trace/ is a self-contained, synthesizable APB3 system used as the CI regression. Its master issues a fixed program — two writes then two reads — to a register-file slave, and check.py asserts the decoded CSV. See tests/apb_trace/README.md.

yosys -s tests/apb_trace/synth.tcl          # (from tests/apb_trace/)
cargo run -r --features metal --bin jacquard -- cosim \
    tests/apb_trace/apb_trace_synth.gv \
    --config tests/apb_trace/sim_config.json \
    --top-module apb_trace \
    --max-clock-edges 200 \
    --bus-trace-csv apb.csv
python3 tests/apb_trace/check.py apb.csv

Troubleshooting

Limitations

• APB3 only for now; AHB-Lite / AHB5 and annotated-VCD output are the next phases (see the plan).
• Up to 4 buses per run, addresses/data up to 32 bits.
• Cosim is now backend-portable (Metal, CUDA, HIP, plus a CPU fallback); bus tracing is wired across the GPU backends (cuda.rs/hip.rs/metal.rs).
• The legacy hardcoded Wishbone trace (a separate, SoC-specific path) is unaffected; folding it onto this general mechanism is a planned follow-up.

Cosim Perf Report (--cosim-perf-json)

Overview

jacquard cosim reports a per-edge timing breakdown of the co-simulation loop, including ground-truth GPU-execution time from device timestamps — free of the instrumentation overhead that a full GPU trace (Metal System Trace / nsys) imposes on a workload of thousands of tiny dispatches per batch.

The breakdown always prints in the human-readable profiling summary at the end of a run. Passing --cosim-perf-json <PATH> additionally writes a flat JSON record for CI to log and track.

The measurement

The cosim loop batches edges into GPU command buffers (BATCH_SIZE edges each). Per batch, the CPU: builds the command buffer (encode), commits it, then waits for the GPU to finish (spin), then drains ring buffers (VCD/UART). The existing summary times these CPU phases with Instant deltas. This report adds the GPU-side truth:

• Metal: MTLCommandBuffer.GPUStartTime/GPUEndTime — host-clock timestamps the driver records for free; GPUEndTime − GPUStartTime is the batch’s GPU-execution wall. Read after waitUntilCompleted (the driver posts the timestamps only at full completion).
• CUDA / HIP: cudaEvent / hipEvent elapsed time — follow-up (the trait seam exists; the impls land in a later PR validated on the GPU CI runner).
• CPU backend: none — gpu_* fields are 0 / gpu_exec_recorded: false.

The seam is one CosimBackend trait method, last_batch_gpu_seconds() -> Option<f64>, called by the orchestration after each batch’s wait(). The generic loop owns the CPU timing and report assembly; each backend supplies only its GPU clock.

Human summary

=== Profiling Breakdown ===
  Batch encode + commit             7.6μs/tick   10.3%
  GPU wait (spin)                  66.3μs/tick   89.3%
  ...
  TOTAL (instrumented)             74.2μs/tick  100.0%
  GPU exec (device ts)             65.8μs/tick   88.7%  (GPU-busy share of wall)

GPU exec is a separate measure from the CPU categories — it overlaps the GPU wait (spin) line (the CPU wall spent waiting is the GPU execution). Its % is the GPU-busy share of the instrumented wall.

JSON schema (report-only)

{
  "backend": "metal",
  "edges": 500000,
  "wall_seconds": 40.1,
  "total_us_per_edge": 74.24,
  "gpu_exec_us_per_edge": 65.85,
  "gpu_util_pct": 88.7,
  "gpu_exec_recorded": true,
  "cpu_encode_us_per_edge": 7.56,
  "gpu_wait_spin_us_per_edge": 66.31,
  "drain_us_per_edge": 0.001,
  "output_vcd_us_per_edge": 0.36,
  "batches": 550,
  "mean_batch": 909.1
}

This is a report format, not a stable contract — it may gain fields.

CI

The mcu-soc-metal job runs the mcu_soc cosim with --cosim-perf-json and posts the JSON to the GitHub step summary (report-only — no perf gate). This builds a timing history before we commit to a regression threshold: GPU/thermal noise on shared runners needs a measured noise floor first, and a premature gate would flake. Adding a gate (fail if total_us_per_edge regresses beyond a tuned threshold vs a committed baseline) is the natural follow-up once the history exists, alongside the CUDA/HIP last_batch_gpu_seconds impls so every tested architecture reports.

Why not just use a GPU trace?

Metal System Trace / nsys instrument every Metal/CUDA call. For a batch of ~2048 tiny dispatches that overhead is large and systematically distorts the picture — it inflated an early measurement to make the GPU look ~88% idle when device timestamps show it ~88% busy. Command-buffer timestamps cost nothing and are the trustworthy signal for steady-state throughput; reserve full traces for one-off per-dispatch investigation (see docs/gpu-capture.md).

GPU Frame Capture (.gputrace)

Overview

jacquard cosim (Metal backend) can emit an Xcode .gputrace document capturing the exact Metal command stream of the GPU simulation — every compute dispatch, buffer binding, and blit for a bounded window of edge batches. This is the detailed counterpart to a system-level xctrace --template "Metal System Trace" recording: the system trace shows when GPU work runs relative to the CPU, while a .gputrace shows what each dispatch does (threadgroup sizes, buffer arguments, per-dispatch timing, dependency DAG, and — after an Xcode replay — shader occupancy and memory-bandwidth counters).

Capture is env-gated and zero-overhead when off: with no JACQUARD_GPU_CAPTURE set, none of the capture code runs.

Quick start

# Capture one steady-state batch (skip the reset/warm-up batch 0).
# METAL_CAPTURE_ENABLED=1 is REQUIRED — Metal refuses programmatic
# .gputrace capture without it.
METAL_CAPTURE_ENABLED=1 \
JACQUARD_GPU_CAPTURE=/tmp/cosim.gputrace \
JACQUARD_GPU_CAPTURE_SKIP=1 \
  cargo run -r --features metal --bin jacquard -- cosim \
    tests/qspi_psram/qspi_psram_dut_synth.gv \
    --config tests/qspi_psram/sim_config.json \
    --top-module qspi_psram_dut \
    --max-clock-edges 200 \
    --output-vcd target/test-out/qspi_psram.vcd

open /tmp/cosim.gputrace   # opens in Xcode's GPU debugger

Environment variables

How it works

The Metal backend commits one command buffer per edge batch (MetalSimulator::encode_and_commit_gpu_batch). Each batch encodes, for every edge in it, the full per-edge dispatch chain:

state_prep → apply_flash_din → simulate_v1_stage × N → flash_model_step → io_step → VCD blit

Capture brackets an MTLCaptureManager scope (destination GpuTraceDocument, scoped to the simulator’s command queue) around the window [SKIP, SKIP + BATCHES):

• Start — immediately before the first in-window batch’s command buffer is created, so the trace opens on a clean batch boundary.
• Stop — after the last in-window batch is committed; that command buffer is waited to completion first so its GPU work is recorded before stopCapture() finalizes the document.

A stale bundle at the target path is removed first (Metal refuses to overwrite). Any failure (missing METAL_CAPTURE_ENABLED, unsupported destination) is logged and the simulation continues normally, uncaptured.

Choosing a batch to capture

Batch 0 is the reset/warm-up batch and is usually small and unrepresentative. JACQUARD_GPU_CAPTURE_SKIP=1 (or higher) captures a steady-state batch. The batch size is BATCH_SIZE edges (see src/sim/cosim/mod.rs), so a single captured batch already contains the complete per-edge dispatch stream repeated across the batch — enough to profile the kernel without a giant multi-batch trace.

Shader performance counters

The .gputrace records the command stream and per-dispatch timing out of the box. Occupancy, bandwidth, and cache-hit counters require an Xcode replay with shader profiling enabled: open the trace in Xcode, press Replay, and Xcode writes the counter streamData back into the bundle. Only then do the profiler_gpu_counters / profiler_gpu_scheduling analyses have data to read.

Relationship to system-level tracing

For CPU↔GPU timeline correlation (where the sim spends wall-clock, GPU duty cycle, inter-batch CPU gaps) use a Metal System Trace instead:

xctrace record --template "Metal System Trace" --output cosim.trace \
  --launch -- jacquard cosim ...

That .trace answers “is the sim CPU- or GPU-bound?”; the .gputrace here answers “is each GPU dispatch efficient?”. They are complementary.

Selective X-Propagation for Jacquard

Summary

This document proposes adding selective unknown-value (X) propagation to Jacquard’s gate-level simulator. Rather than uniformly upgrading the entire simulator to four-state logic (2x storage, ~2-3x ALU cost), we use static analysis at compile time to identify which signals can carry X values, and only apply X-aware simulation to the affected partitions. The rest of the design continues to run with the existing fast two-state kernel.

Motivation

Jacquard currently simulates in pure two-state Boolean logic (0/1). All DFFs and SRAMs start at 0. This is fast but has two significant drawbacks:

1. Undetected initialisation bugs: If a design reads a register before it has been written through a proper reset sequence, the simulator silently returns 0 instead of flagging the value as unknown. Real hardware would produce an arbitrary bit pattern.

2. SRAM initialisation masking: SRAMs read as all-zeros before being written. A design that depends on SRAM initialisation may appear to work in simulation but fail on silicon.

3. No RTL/gate-level mismatch detection: RTL simulators (Icarus, VCS, Questa) propagate X values that expose these bugs. When comparing Jacquard’s gate-level results against an RTL reference, false mismatches arise because Jacquard resolves unknowns to zero.

Naively upgrading the entire simulator to four-state would halve throughput (double storage per signal, double ALU per gate). The key insight is that in a well-designed SoC after reset, typically <5% of signals are genuinely X-capable (uninitialised memories, clock-domain crossings, registers without reset). We should only pay the overhead where it matters.

Background: X in And-Inverter Graphs

Jacquard’s core IR is an And-Inverter Graph (AIG). Every gate is an AND with optional input inversions. The boomerang reduction tree computes:

ret = (a ^ xora) & ((b ^ xorb) | orb)

Where xora/xorb encode inversions and orb encodes pass-through (when orb = 0xFFFFFFFF, input b is forced to 1, making the gate a buffer for a).

The AND gate has a favourable property for X propagation:

A known-zero on either input forces the output to known-zero regardless of X on the other input. This means X does not spread as aggressively through AND logic as it would through XOR or MUX logic. AIG-based designs have a natural tendency to “absorb” X values at AND gates with known-zero inputs.

NOT (inversion) simply preserves the X mask: NOT(X) = X, NOT(0) = 1, NOT(1) = 0.

Design

Phase 1: Static X-Source Analysis (Compile Time)

At AIG construction time, identify all X sources – signals whose initial value is unknown:

1. DFF Q outputs: Every DFF output is an X source at cycle 0. In real hardware, flip-flop power-on state is indeterminate.

2. SRAM read ports: All 32 data-output pins of each RAMBlock are X sources. Memory contents are undefined until written.

3. Undriven primary inputs: If a primary input is not driven by the testbench VCD, it should be marked X. (Currently Jacquard warns about missing PI signals; this would upgrade that warning to X propagation.)

These are identified directly from AIG.drivers (variant DFF(_) and SRAM(_)) and from AIG.dffs / AIG.srams.

Phase 2: Forward Cone Computation (Compile Time)

Compute the forward cone of influence from all X sources through the AIG. Since AIG pins are guaranteed to be in topological order, this is a single linear-time forward pass:

x_capable = BitVec::zeros(num_aigpins + 1)

// Mark X sources
for each DFF:
    x_capable.set(dff.q)
for each SRAM:
    for each read data pin:
        x_capable.set(pin)

// Forward propagate (pins are in topological order)
for aigpin in 1..=num_aigpins:
    if let AndGate(a_iv, b_iv) = drivers[aigpin]:
        a = a_iv >> 1   // strip inversion bit
        b = b_iv >> 1
        if x_capable[a] || x_capable[b]:
            x_capable.set(aigpin)

This is O(V + E) on the AIG – negligible compared to partitioning.

Sequential propagation: X-capability also propagates through DFF feedback loops. If a DFF’s D input is X-capable, its Q output remains X-capable (even after the first clock edge, because it may have captured an X value). The analysis iterates until a fixpoint:

loop:
    changed = false
    for each DFF:
        d_pin = dff.d_iv >> 1
        if x_capable[d_pin] && !x_capable[dff.q]:
            x_capable.set(dff.q)
            changed = true
    if !changed:
        break
    // Re-run forward cone from newly-marked DFF Q outputs

In practice this converges in 1-2 iterations because most feedback loops go through DFFs that are already marked.

Phase 3: Partition Classification

After partitioning (mt-kahypar), classify each partition:

• X-capable: Contains at least one X-capable aigpin, OR reads input state from an X-capable partition’s output. Run with the X-aware kernel variant.

• X-free: All signals are provably not-X. Run with the existing fast two-state kernel.

The classification must account for inter-partition communication:

loop:
    changed = false
    for each partition P:
        if P is already X-capable:
            continue
        for each global-read in P's script:
            source_word = identify source partition and state word
            if source partition is X-capable:
                mark P as X-capable
                changed = true
                break
    if !changed:
        break

Phase 4: X-Mask Representation

For X-capable partitions, each signal carries a sideband X mask alongside its value:

• v (value bit): The Boolean value. When x = 1, this is a “best guess” (we use 0 by convention, matching Jacquard’s current behaviour for backwards-compatible output).
• x (X-mask bit): 1 = unknown, 0 = known.

This doubles the per-signal storage within X-capable partitions only.

State Buffer Layout

State buffer (current):
  [word 0] [word 1] ... [word N-1]   ← N u32 words, 32 signals each

State buffer (with X sideband):
  [word 0] [word 1] ... [word N-1]   ← value words (unchanged)
  [word N] [word N+1] ... [word 2N-1] ← X-mask words (new, same layout)

X-free partitions only read/write the value section. X-capable partitions read/write both sections. The sideband occupies the same state-buffer address space with a fixed offset of state_size words.

Phase 5: X-Aware Boomerang Kernel

The existing boomerang gate computation:

ret_v = (a ^ xora) & ((b ^ xorb) | orb)

The X-aware version computes in parallel:

// Effective inputs after inversion and OR-bypass
a_eff   = a_v ^ xora
b_eff   = (b_v ^ xorb) | orb
b_eff_x = b_x & ~orb        // OR-bypass forces bits to known-1

// Value: same as before (X bits treated as 0 in value lane)
ret_v = a_eff & b_eff

// X mask: result is X when both inputs are not-known-zero AND at
// least one input is X
//
//   known_zero_a = ~a_x & ~a_eff  (not X, and effective value is 0)
//   known_zero_b = ~b_eff_x & ~b_eff
//   ret_x = (a_x | b_eff_x) & ~known_zero_a & ~known_zero_b
//
// Expanded:
ret_x = (a_x | b_eff_x) & (a_eff | a_x) & (b_eff | b_eff_x)

Verification of the X-mask formula:

Wait – row for a=X, b=*, orb=1 (pass-through): a_x=1, a_eff=0, b_eff=1, b_eff_x=0. ret_x = (1|0) & (0|1) & (1|0) = 1 & 1 & 1 = 1. Correct.

The X-mask computation adds 5 extra bitwise operations per boomerang level (2 OR, 2 AND, 1 AND-NOT – plus loading the X mask values). This is roughly 1.5-1.6x the CPU ALU work per gate, but only within X-capable partitions.

GPU Resource Impact

Per X-capable partition:

Metal threadgroup memory limit is 32 KB; current usage is ~4 KB. The extra 1 KB is well within budget.

Phase 6: Boundary Protocol

When signals cross from an X-capable partition to another partition:

• X-capable -> X-capable: Both value and X-mask words are communicated through the state buffer. The reading partition loads both.

• X-capable -> X-free: At the boundary, we assert that no X values cross. If an X-mask bit is set for a signal read by an X-free partition, this is a simulation error (the design has an X reaching a region we statically proved should be X-free). Report it and optionally halt.

• X-free -> X-capable: The reading partition loads the value word normally and treats the X-mask as all-zero for those inputs. No overhead.

Phase 7: Dynamic X Narrowing (Optional Enhancement)

After the design’s reset sequence completes, most DFFs will hold known values and the X-mask will be all-zeros across most of the state. The simulator can detect this:

every K cycles (e.g., K = 1000):
    for each X-capable partition:
        if all X-mask words in this partition's state are zero:
            switch partition to fast two-state kernel

This gives the best of both worlds: full X-propagation during initialisation (when it matters most), then automatic fallback to maximum throughput once the design is in steady state.

For designs with a clear reset phase, this means the performance overhead of X-propagation is confined to the first few thousand cycles.

Expected Performance Impact

Compile Time

The static analysis adds:

• X-source identification: O(|DFFs| + |SRAMs|) – negligible
• Forward cone computation: O(|AIG pins|) – one linear pass
• Partition classification: O(|partitions| x |global reads|) – negligible
• Fixpoint iteration: 1-2 rounds of the above

Total: well under 1% of the partitioning time.

Simulation Time

For a typical SoC with reset:

The overall impact on a full simulation run is estimated at 5-15% throughput reduction compared to current two-state, in exchange for catching initialisation bugs that would otherwise escape to silicon.

Implementation Status

All stages are implemented. Enable with --xprop on jacquard sim.

Stage 1: Static X-Source Analysis (src/aig.rs)

• compute_x_sources() -> Vec<bool> – identifies DFF Q outputs and SRAM read data ports
• compute_x_capable_pins() -> (Vec<bool>, XPropStats) – forward cone + fixpoint iteration

Stage 2: Partition Classification (src/flatten.rs)

• FlattenedScriptV1 fields: xprop_enabled, partition_x_capable, xprop_state_offset
• effective_state_size() returns 2 * reg_io_state_size when xprop enabled
• Metadata words 8 (is_x_capable) and 9 (xmask_state_offset) encode per-partition xprop info
• X-propagation is configured at simulation startup, no separate mapping step needed

Stage 3: X-Aware CPU Reference Kernel (src/sim/cpu_reference.rs)

• simulate_block_v1_xprop() – dual-lane value + X-mask computation
• sanity_check_cpu_xprop() – CPU vs GPU comparison for both lanes
• Non-X-capable partitions delegate to simulate_block_v1 (zero overhead)

Stage 4: CLI and VCD Integration (src/bin/jacquard.rs, src/sim/vcd_io.rs)

• jacquard sim --xprop runs static analysis, logs X-capable pin/partition stats, and enables X-aware simulation with doubled state buffer
• write_output_vcd_xprop() emits Value::X for X-masked primary outputs
• expand_states_for_xprop() / split_xprop_states() for state buffer management

Stage 5: GPU Kernels (csrc/kernel_v1.metal, csrc/kernel_v1_impl.cuh)

• Dual-lane X-mask tracking through all kernel phases (global read, boomerang, SRAM, DFF writeout)
• is_x_capable branch is uniform per threadgroup (zero warp/SIMD divergence)
• Shared memory: shared_state_x[256] + shared_writeouts_x[256] (+2KB, within 32KB limit)
• SRAM X-mask shadow via sram_xmask buffer (buffer slot 7 on Metal)

Stage 6: Diagnostics

• Compile-time: X-source count, X-capable pin %, X-capable partition %
• Runtime: first-cycle-X-free detection, final-cycle X warning
• CPU sanity check uses xprop variant when enabled

Stage 7: Benchmarks (benches/xprop.rs)

• Criterion micro-benchmarks: two_state vs xprop_xfree vs xprop_xcapable
• X-free partitions: zero overhead confirmed
• X-capable partitions: ~1.5-1.6x CPU overhead (well within 2x budget)

Dynamic Narrowing (Future Enhancement)

• Periodic X-mask scan on CPU between GPU batches
• Partition kernel hot-swapping from X-aware to fast mode
• Statistics reporting: “X cleared after N cycles”

Prior Art

• Mixed 2-4 State Simulation with VCS (Chaudhry et al., 1997): Proved that mixed two-state / four-state simulation is viable within a single Verilog simulator, validating the core approach.

• A Two-State Methodology for RTL Logic Simulation (1999): Eliminated X entirely using random two-state initialisation, arguing it catches more bugs than X-optimistic RTL simulation. Their technique for handling Z-state boundaries is relevant to our boundary protocol.

• Essent (Beamer, 2020-2021): Demonstrated that partitioning a compiled simulator by signal activity yields significant speedups. Their activity-proportional execution is analogous to our dynamic X narrowing – skip work for partitions where nothing interesting is happening.

• Synopsys VCS T-Prop: Taint propagation in VCS uses a parallel sideband bit per signal to track data flow for security verification. Our X-mask is architecturally identical – a sideband bit that propagates alongside the value using different rules.

• Chris Drake, “Improving Verilog Four State Logic” (2024): Argues that Verilog conflates “uninitialised” and “don’t care” uses of X, and proposes splitting them. Our approach naturally supports this: the X-mask tracks only uninitialised/unknown values, not synthesis don’t-cares (which are already resolved by the synthesis tool before Jacquard sees the netlist).

What This Does NOT Address

• Z (high-impedance): Jacquard does not support tri-state buses and this proposal does not add support. Z would require a third state bit and bus resolution logic.

• X-optimism in RTL control flow: Since Jacquard operates on gate-level netlists (post-synthesis), there are no if/case statements to be X-optimistic about. Gate-level X propagation through AND/OR truth tables is naturally correct (though potentially X-pessimistic – see below).

• X-pessimism reduction: Gate-level simulation is inherently X-pessimistic in some cases (e.g., a XOR a should be 0 even if a is X, but standard gate-level propagation gives X). Since AIGs decompose XOR into AND/NOT, this pessimism is present. Addressing this would require symbolic analysis or reconvergence detection, which is out of scope.

• Strength modelling: No weak/strong drive strength tracking. All known values are “strong.”

Design Decisions

1. SRAM X granularity: Conservative whole-SRAM approach – all reads return X until any write occurs. This is pessimistic but correct and simple. Per-address tracking (requiring 8192 x 32 bits of shadow state per SRAM block) is a potential follow-up for reduced pessimism.

2. Reset-aware analysis: Skipped for v1 – all DFF Q outputs start as X regardless of whether they have async reset. Identifying reset nets is fragile (varies by synthesis tool and coding style), and the fixpoint iteration naturally resolves DFFs that become non-X after reset propagates.

3. VCD X output: Yes – write_output_vcd_xprop() emits Value::X when the X-mask bit is set for a primary output. This is the primary user-visible benefit. The vcd-ng crate already supports Value::X. Downstream tools that only handle two-state VCD would need updating, but Verilog-standard four-state VCD is widely supported.

4. Partition-level granularity: X-awareness is applied at partition granularity (entire partition runs X-aware kernel or not). This provides better steady-state performance than per-signal tracking when most partitions are X-free (~95% typical), at the cost of slightly more pessimism at partition boundaries.

5. Runtime CLI flag: X-propagation is controlled by --xprop on jacquard sim, rather than a compile-time feature flag. No new Cargo dependencies are needed.

6. State buffer layout: When xprop is enabled, the state buffer doubles. Value words occupy [0 .. reg_io_state_size) and X-mask words occupy [reg_io_state_size .. 2*reg_io_state_size) per cycle. Per-partition metadata word 8 stores the is_x_capable flag, and word 9 stores the xmask_state_offset (equal to reg_io_state_size for X-capable partitions). X-free partitions ignore both fields and run unchanged.

Debugging X values in cosim (xsources + xroots)

When you run jacquard cosim --xprop and a signal reads x (unknown), the question is always why — which uninitialised state or undriven input is feeding that X forward? This guide covers the two-command workflow that answers it as a static query, without the trace→guess→re-run loop.

• jacquard xsources — enumerate where X originates in the design.
• netlist-graph xroots — given a signal, find which of those X-sources reach it (its backward “X-root” frontier).

See Decision 0016 for the X-propagation semantics and issue #98 for the design rationale.

Background: where X comes from

Under --xprop, an X value originates at one of three X-sources and then propagates forward through combinational logic:

A signal reads X iff at least one X-source lies in its backward logic cone. That is a pure netlist property — no simulation needed to compute it.

Step 1 — enumerate X-sources: jacquard xsources

jacquard xsources <netlist> --config <sim_config.json> -o xsources.json

This builds the static AIG (no GPU) and writes a JSON manifest of every X-source. --config is required only to classify undriven-input sources (the driven-set complement); without it, only DFF and SRAM sources are emitted.

{
  "schema_version": "1.0",
  "netlist": "design.v",
  "undriven_inputs_classified": true,
  "x_sources": [
    { "net": "q_unreset",  "kind": "unreset-dff",    "cell": "_29_" },
    { "net": "_10_[1]",    "kind": "unreset-dff",    "cell": "_30_" },
    { "net": "q_reset",    "kind": "reset-dff",      "cell": "_28_" },
    { "net": "unconnected","kind": "undriven-input" }
  ]
}

Notes:

• net is jacquard’s canonical net name (after assign-merging), which may differ from the raw name in the Verilog. cell is the driving instance and is stable across tools — xroots resolves DFF/SRAM sources by cell, so the naming difference does not matter.
• The undriven-input classification reflects the config’s declared drivers: clock(s), reset, constant_inputs, constant_ports, and each peripheral’s pins (flash/uart/gpio/jtag). An input outside that set is reported as undriven.
• A reset-connected DFF is emitted as reset-dff. It is X at power-up but resolves once reset asserts, so xroots traverses through it rather than stopping (see below).

Step 2 — find a signal’s X-roots: netlist-graph xroots

netlist-graph xroots <netlist> <signal> --xsources xsources.json

xroots walks backward from <signal> through the driver cone — and through DFF data (D) pins, the reverse of the forward X-prop fixpoint — skipping clock/set/reset pins. It stops at each persistent X-source (unreset-dff / sram-read / undriven-input), reporting the nearest frontier:

$ netlist-graph xroots design.v top.cpu.stall --xsources xsources.json
X-source frontier of top.cpu.stall (2 root(s), nearest first):
  [unreset-dff] top.cpu.state[3]   (depth 5)
  [undriven-input] cfg_mode        (depth 8)

Reset DFFs (reset-dff) are traversed through (their data cone is followed), because reset defines them — so the frontier surfaces the real persistent roots behind them, not the reset flop itself.

If no X-source reaches the signal, xroots says so — under --xprop that signal should be X-free.

Without a manifest

xroots also runs without --xsources, classifying X-sources from the netlist alone (DFF Q by reset-pin presence, SRAM reads, and genuinely undriven internal nets). This cannot classify undriven-input sources (it does not know the cosim driven set), so pass jacquard xsources --config output when you need those.

Step 3 — confirm with a targeted --xprop run

--emit-trace writes the frontier as a --trace-signals file, so confirming the X actually originates where xroots says is one command:

netlist-graph xroots design.v top.cpu.stall \
    --xsources xsources.json --emit-trace xroots.txt

jacquard cosim design.v --config sim.json \
    --xprop --trace-signals xroots.txt --output-vcd out.vcd

The traced frontier nets appear in out.vcd; watch them carry x and resolve (or not) over the reset window to confirm the diagnosis.

Worked example

tests/xprop_cosim/ is the X-propagation demo. q_unreset is a self-holding unreset register (stays X); q_reset resolves after reset:

jacquard xsources tests/xprop_cosim/xprop_demo_synth.gv \
    --config tests/xprop_cosim/sim_config.json -o /tmp/xsrc.json

# Why is the counter's next-state X?
netlist-graph xroots tests/xprop_cosim/xprop_demo_synth.gv "_11_[3]" \
    --xsources /tmp/xsrc.json
#   [unreset-dff] unreset_count[3]  (depth 2)
#   [unreset-dff] unreset_count[2]  (depth 3)
#   ...

Limitations

• The frontier is a static over-approximation: it reports X-sources whose cone reaches the signal, not whether the X is live on a given cycle (an unreset DFF may resolve once its own data cone settles from known inputs). It tells you what to initialise/drive to make the signal defined.
• Name round-tripping is exact for flattened post-synthesis netlists (the target of these tools). DFF/SRAM sources resolve by cell instance, which is robust; undriven-input sources resolve by name.
• Dominator analysis (X-sources that every path passes through — guaranteed roots) is a planned follow-up; v1 reports the reachable frontier.

Interactive JTAG Debug (--jtag-server)

Overview

jacquard cosim --jtag-server <PORT> opens a live remote_bitbang JTAG socket alongside a running GPU co-simulation. An external debugger — OpenOCD, then gdb on top of it — attaches and inspects the design through its RISC-V Debug Module (DM): halt/resume/step the hart, read/write GPRs, CSRs, PC and memory, and load firmware — exactly as the same firmware would be debugged on real silicon.

It is the interactive sibling of --jtag-replay. Where --jtag-replay <FILE> plays back a recorded remote_bitbang byte stream (deterministic, one-directional), --jtag-server <PORT> opens a live socket and lets the connected client drive the same configured TCK/TMS/TDI/(TRST) pins, stepping the design in lock-step with each debug transaction and answering TDO reads from the design’s live output.

The two are mutually exclusive — pick a recorded stream or a live client, not both.

Configuration

The JTAG TAP pins are declared once in sim_config.json, shared by both the replay and server paths. TCK/TMS/TDI/TRST are design inputs; TDO is a design output, so the server needs tdo_gpio to answer TDO reads:

{
  "jtag": {
    "tck_gpio": 2,
    "tms_gpio": 3,
    "tdi_gpio": 4,
    "trst_gpio": 5,
    "trst_active_low": true,
    "tdo_gpio": 6
  }
}

• trst_gpio should be set for a RISC-V Debug TAP: the DTM resets on negedge trst_n, and the server injects a one-shot power-on TRST pulse at startup so a debugger that never asserts TRST (stock OpenOCD with reset_config none) still resets the DTM — mirroring a real chip’s power-on TAP reset. (Tunable/disable via the JACQUARD_JTAG_TRST_PULSE env var; see How it works.) TAPs that genuinely reset via a five-cycle TMS=1 sequence can omit it.
• tdo_gpio is optional for --jtag-replay (replay never answers R) but required for useful --jtag-server use: without it, every TDO read returns 0 and the debugger cannot read the design back. cosim warns at startup if it is missing.

The TCK clock domain must also appear in clocks[] so the multi-clock scheduler allocates it a tick slot; the model overrides TCK at each byte transition. See tests/jtag_minimal/sim_config.json for a complete, working example.

Launching the server

jacquard cosim design.pnl.v \
    --config sim_config.json \
    --top-module top \
    --jtag-server 9999 \
    --jtag-hold-cycles 8 \
    --output-vcd debug.vcd \
    --max-clock-edges 100000000

cosim builds the design, binds 127.0.0.1:9999, and blocks waiting for one client before the run starts:

JTAG server `jtag_0`: listening on 127.0.0.1:9999, hold_edges=8 (...);
  waiting for a remote_bitbang client (e.g. OpenOCD)…

Pass --jtag-server 0 to bind an OS-assigned free port instead of a fixed one — read the actual port from the listening on 127.0.0.1:<port> line. Use this when several jacquard instances debug concurrently (or in CI) so they never collide; if a fixed port is already taken, the bind fails with an actionable error.

Once a client connects, the design steps forward one remote_bitbang transaction at a time, paced by the client. By default a single connection is served; when the client disconnects (or sends Q), the session ends and the simulation free-runs to its edge budget.

Add --jtag-reconnect to keep the server alive across disconnects: when the debugger detaches, the server preserves the design state and waits to accept() the next client — so you can restart OpenOCD without re-running the slow cosim setup (each fresh attach gets a clean DTM reset). The run then never stops on its own; Ctrl-C/kill to end it.

Edge budget. --max-clock-edges does not stop the run while a debug client is attached — the session is paced by the client and would otherwise die mid-inspect. The cap is honoured again once the client disconnects (so the post-session free-run is still bounded). You can leave --max-clock-edges at its default for interactive use.

OpenOCD configuration

The boilerplate is the most error-prone part of attaching a debugger, so generate it with the built-in helper instead of hand-writing it:

jacquard jtag-openocd-config --port 9999 --expected-id 0xdeadbeef \
    -o openocd.cfg

--irlen (default 5), --host, --chipname and --dmi-timeout are all configurable; jacquard jtag-openocd-config --help lists them. The emitted config wires up the remote_bitbang driver, a RISC-V Debug TAP, and the target, and ends with init. For reference, the equivalent hand-written config (mirroring tests/jtag_minimal/scripts/openocd.cfg):

adapter driver remote_bitbang
remote_bitbang host localhost
remote_bitbang port 9999
transport select jtag
adapter speed 1

set _CHIPNAME jtag
jtag newtap $_CHIPNAME cpu -irlen 5
set _TARGETNAME $_CHIPNAME.cpu
target create $_TARGETNAME riscv -chain-position $_TARGETNAME

# The JTAG protocol is bit-serial and the GPU server is per-edge while
# stepping — keep the DMI timeout generous.
riscv set_command_timeout_sec 60

Run it against the launched server:

openocd -f openocd.cfg

OpenOCD examines the TAP, brings up the DM (dmactive=1), and exposes a gdb remote on :3333 by default.

Attaching gdb

riscv32-unknown-elf-gdb firmware.elf
(gdb) target remote :3333
(gdb) info registers          # read GPRs / CSRs / PC via the DM
(gdb) x/16xw 0x20000000       # read memory
(gdb) load                    # write firmware through Program Buffer / sysbus
(gdb) break main
(gdb) continue

The design runs on the GPU; gdb sees it through the DM exactly as it would see silicon over a real JTAG probe.

How it works

The cosim main loop is synchronous (step_edge → run_edges → wait → repeat). A live debug session inverts time control: the client is the clock source, so step_edge blocking on a socket read is the correct synchronisation — no async executor, no background thread.

• While a client is connected the JTAG model reports is_active() == true, which forces the existing single-edge (batch=1) dispatch — the same fine-grained stepping --jtag-replay already uses. Each cosim edge processes one remote_bitbang step.
• On each R (read-TDO) command the model samples the design’s live TDO from the output state (output_state[tdo_pos]) and writes the ASCII '0'/'1' back over the socket — the only response the protocol requires.
• TCK/TMS/TDI/TRST drive through the usual overrides → BitOps → state_prep path, unchanged from replay.
• Power-on TRST pulse. A RISC-V DTM resets on negedge trst_n. On real silicon the power-on reset supplies that edge; in a recorded capture the harness pulses TRST before the debugger runs. But stock OpenOCD with reset_config none never asserts TRST — so without help the DTM would never reset and examine would fail with “dtmcontrol is 0” / “scan chain interrogation: all zeroes”. The live server therefore injects one brief TRST pulse at startup (deasserted → asserted → released, timed in cosim edges), reproducing the recorded stream’s leading u…r. Replay is unaffected (it drives TRST from the stream), and an explicit client TRST assertion still wins after the pulse. The window is [hold_edges, 5·hold_edges) by default; override or disable it with JACQUARD_JTAG_TRST_PULSE="<from>,<to>" / ="off" (edges) if a design needs different timing.

The batched fast path is untouched when no client is attached: a design with a jtag config but no --jtag-server/--jtag-replay flag simply leaves the JTAG inputs floating and runs at full throughput.

This works on any cosim backend (CPU / Metal / CUDA / HIP) — it is the CPU-side model plus batch=1 of the GPU backend, with no per-backend kernel work. See Decision 0017 (Amendment 2026-06-21) for the execution-model rationale and Decision 0013 (Amendment 2026-06-21) for the tdo_gpio config surface.

Caveats and limitations

• One client at a time. A single connection is served at once; --jtag-reconnect accepts a new one after the previous disconnects. Simultaneous clients and multi-tap chains are future work.
• Performance. An attached session loses edge batching for its duration — inherent and acceptable, since interactive debug is slow relative to free-running simulation. Throughput is unaffected when no client is attached.
• X-propagation. Under --xprop, the debug read path is two-state in v1 (TDO resolves through the design as 0/1). X-aware debug reads are a planned refinement (J5 in the plan).
• Reset interplay. The design’s own reset (reset_gpio) and the JTAG TAP reset (driven by the client) proceed concurrently, mirroring real hardware where the debugger drives JTAG while the chip resets.

Validation

Three layers, increasing in fidelity:

• A model-level loopback unit test (live_source_loops_back_tdo_over_socket in src/sim/models/jtag.rs) drives the FSM through a real TcpStream and reads back a sampled TDO bit — GPU-free and deterministic, runs in cargo test.

• The jtag-minimal-cosim-server CI job streams the same bitbang.rec the --jtag-replay gate uses over the live socket (via tests/jtag_minimal/scripts/bitbang_client.py) and asserts the design reaches the same data0_obs == 0xCAFEBABE. This pins live-vs-replay drive equivalence with zero external tooling — but reaches 0xCAFEBABE via a DMI write, so it does not exercise a correct TDO read.

• The jtag-minimal-openocd CI job drives the live server with real OpenOCD (tests/jtag_minimal/scripts/openocd_debug.sh): examine reads IDCODE/DTMCS back over TDO, then it writes and reads back DATA0, asserting IDCODE == 0xdeadbeef and the DMI read-back == 0xcafebabe. This is the only gate that exercises the live TDO read path (and the power-on TRST-pulse fix). Run it locally with any port:

  jacquard cosim … --jtag-server 9824 &
  bash tests/jtag_minimal/scripts/openocd_debug.sh 9824
  

  Fake hart. jtag_minimal is a bare DTM+DM with no real CPU, so full RISC-V target examination stops at Failed to read MISA and OpenOCD exits non-zero — expected. DMI access (write/read DATA0) is the contract the gate checks. A design with a real hart examines fully.

References

• Issue #124.
• Implementation staging: docs/plans/jtag-debug-server.md.
• Decision 0017 — Cosim execution model (interactive, externally-paced peripheral models; output_state wiring).
• Decision 0013 — Cosim peripheral model architecture (tdo_gpio config surface).
• Replay path & fixture: src/sim/models/jtag.rs, tests/jtag_minimal/.

Adding a New PDK for Post-Layout Simulation

This guide documents the process of enabling a new process design kit (PDK) for gate-level simulation in Jacquard. It is based on the SKY130 enablement and captures every integration point.

Overview

Jacquard natively supports AIGPDK (its own synthesis library of AND gates, DFFs, and SRAMs). Supporting a foundry PDK like SKY130 requires teaching the simulator how to interpret the PDK’s standard cells: their pin directions, their boolean function, and which ones are sequential.

There are three pathways for enabling new cells; pick based on what you’re adding:

• Cell-model IR descriptor (recommended — the modern path, decision 0019). A standard-cell library — pin directions, boolean functions, sequential roles, and timing — is captured in one generated JSON descriptor produced from the library’s Liberty by the liberty-to-cellir converter. Jacquard consumes the descriptor at runtime; no per-PDK Rust is required. This is how a brand-new PDK is added (vendor the library + generate a descriptor) and how a proprietary library you cannot vendor is simulated (--cell-descriptor). See “Pathway A” below. This supersedes the legacy Rust workflow (Steps 1–8) for all combinational logic today, with sequential consumption completing across the built-in PDKs in the Decision 0019 C3 series.
• Runtime cell library (--cell-library + .cells.toml manifest). For third-party IP, hard macros, foundry memories, and any other cells that don’t need new AIG decomposition rules — i.e. cells that act as opaque outputs (RAM macros), filler/cap blocks, or IO pads. See Decision 0010 and docs/plans/archive/declarative-cell-metadata.md for the recipe. No Jacquard PR required — users ship a manifest alongside their netlist. See “Adding third-party IP via runtime manifest” at the end.
• Legacy per-PDK Rust (Steps 1–8 below). The original pathway: pin tables, classifiers, decomposition functions, and AIG builder hooks hand-written in Rust. Being retired by the cell-model IR cutover (Decision 0019) — retained here as reference for the machinery the descriptor replaces. Do not add a new PDK this way; use Pathway A.

If you’re adding just a memory macro or other behaviourally-opaque IP, skip ahead to “Adding third-party IP via runtime manifest” at the end of this document — it’s a 6-line TOML entry, not a Rust PR.

Pathway A: Cell-model IR descriptor (recommended)

A cell-model IR descriptor is a portable, versioned, generated JSON file carrying everything per-cell-type about a library — L1 pin directions, L2 combinational logic (as a pre-decomposed AIG), L3 sequential roles/classification, and L4 timing characterization — from one source: the library’s Liberty. Adding a PDK is no longer a Jacquard code change; it is a data generation step. (Decision 0019.)

Generate a descriptor from Liberty

The liberty-to-cellir converter reads a Liberty .lib (Liberty-first; a functional .v is consulted only as a fallback / cross-check) and emits the descriptor:

cargo run --release --manifest-path crates/liberty-to-cellir/Cargo.toml -- \
    path/to/mylib_typ_1p20V_25C.lib \
    --functional-v path/to/mylib_stdcell.v \
    -o mylib.cellir.json

The converter derives the corner (PVT) from the Liberty’s operating_conditions/default_operating_conditions, compiles each cell’s function into an AIG, extracts ff/latch sequential roles, and (where a .v is present and per-cell) cross-checks logic + timing-arc topology, surfacing any disagreement. It prints a summary (cells=… combinational=… l3_sequential=… l4_timing=… corners=… cross_check_mismatches=…). A monolithic single-.lib-per-corner commercial library and a per-cell split library are both handled.

Selection: how Jacquard picks a descriptor for a netlist

Each descriptor declares the cell-name prefix(es) it covers (D8). Jacquard auto-matches a netlist’s cell types against the bundled descriptors by prefix; --cell-descriptor <file.json> is the explicit override (and the path for proprietary libraries). Precedence: --cell-descriptor (explicit file) > --bundled-descriptor <name> (explicit bundled) > auto-match by prefix > the default-fallback descriptor (used when no vendor prefix matches — this is how AIGPDK, whose cells have no common prefix, is selected).

Worked example — a built-in PDK (IHP SG13G2, zero per-PDK Rust)

IHP’s open SG13G2 PDK was added as a built-in with no Rust:

1. Vendor the library as a sparse/shallow git submodule (only the stdcell Liberty + .v are checked out — ~10 MB, not the multi-GB PDK):

   git submodule add https://github.com/IHP-GmbH/IHP-Open-PDK.git vendor/IHP-Open-PDK
   git -C vendor/IHP-Open-PDK sparse-checkout set \
       ihp-sg13g2/libs.ref/sg13g2_stdcell/lib \
       ihp-sg13g2/libs.ref/sg13g2_stdcell/verilog
   

2. Generate + embed at build time. build.rs runs the converter over the pinned vendored .lib and embeds the descriptor into the binary (descriptors are not checked in — CI regenerates them deterministically, D7). Add the descriptor to bundled_descriptors::ALL with its declared prefix (sg13g2_).

3. That’s it. A SG13G2 netlist auto-selects the descriptor by prefix; combinational cells splice from it with no IHP-specific code. Adding the PDK touched no ihp_*.rs — only a submodule, a build.rs generation line, and a registry entry.

Worked example — a proprietary library you cannot vendor (GF130)

A foundry/NDA library that can never live in vendor/ is simulated entirely at runtime — the Liberty never leaves your machine:

# 1. Generate a descriptor from your own Liberty (offline, on your machine):
cargo run --release --manifest-path crates/liberty-to-cellir/Cargo.toml -- \
    /path/to/pdk-gf130/.../GF013bcd_sc6_1p5_a0_TT_1P50V_25C_max.lib \
    --functional-v /path/to/pdk-gf130/.../GF013bcd_sc6_1p5_a0.v \
    -o gf130.cellir.json

# 2. Point jacquard at it — no Jacquard build, no vendoring:
jacquard sim my_chip.gv stim.vcd out.vcd 1 --cell-descriptor gf130.cellir.json

The generator is the only tool that touches raw foundry files; the released jacquard binary is library-agnostic. (A --corner <name> flag selects among a multi-corner descriptor’s corners, defaulting to the descriptor’s default_corner.)

Current status & limits

• Combinational logic is fully descriptor-driven for every built-in PDK (GF180, SKY130, AIGPDK, IHP) and for proprietary libraries via --cell-descriptor.
• Sequential consumption is descriptor-driven for GF180 today and is being wired for the other PDKs in the Decision 0019 C3 series; the descriptor already carries the L3 sequential data for all of them.
• A library whose functional .v is a single flat-module file (common for commercial PDKs) skips the optional .v logic cross-check at generation (the descriptor is still emitted; validate via simulation) — improving the flat-.v indexer is tracked follow-up.
• RAM/SRAM macros are not covered by the descriptor — they use the runtime manifest (Decision 0011), below.

Prerequisites

You need:

• The PDK’s Verilog cell library (behavioral or functional models)
• A post-synthesis or post-P&R netlist using those cells
• The cell naming convention (prefix, drive strength suffix format)

For SKY130, the PDK data lives in vendor/sky130_fd_sc_hd/ as a git submodule.

Legacy pathway: per-PDK Rust (Steps 1–8)

This section documents the original hand-written-Rust workflow, which the cell-model IR (Pathway A, above) is retiring under Decision 0019. It is kept as a reference for the machinery the descriptor replaces and for the still-live paths during the cutover. To add a new PDK, use Pathway A — do not write per-PDK Rust. The RAM/macro manifest section that follows Step 8 remains fully current.

Step 1: Library Detection

Reference: src/sky130.rs – is_sky130_cell(), detect_library(), detect_library_from_file()

Jacquard scans the netlist to determine which cell library is in use. Each PDK needs a name-matching function:

#![allow(unused)]
fn main() {
// src/sky130.rs:535
pub fn is_sky130_cell(name: &str) -> bool {
    name.starts_with("sky130_fd_sc_")
        || name.starts_with("CF_SRAM_")
}
}

The CellLibrary enum tracks known libraries. detect_library() iterates cell names and returns the detected library (or Mixed if cells from multiple libraries are found – this is an error).

For a new PDK: Add a variant to CellLibrary, write an is_<pdk>_cell() function, and update detect_library().

Step 2: Cell Type Extraction

Reference: src/sky130.rs – extract_cell_type()

PDK cell names follow a convention: <prefix>__<type>_<drive>. The simulator needs to strip the prefix and drive strength to get the base cell type:

sky130_fd_sc_hd__nand2_4  -->  nand2
sky130_fd_sc_hd__dfxtp_1  -->  dfxtp

This function must handle all library variants (hd, hs, ms, ls, lp, hdll, hvl for SKY130) and any custom macros (CF_SRAM_*).

For a new PDK: Write an equivalent extract_cell_type() for the PDK’s naming scheme.

Step 3: Pin Direction Provider

Reference: src/sky130.rs – SKY130LeafPins implementing LeafPinProvider

The netlist parser (from eda-infra-rs/netlistdb) needs to know pin directions and widths for every cell type. This is implemented as a trait:

#![allow(unused)]
fn main() {
impl LeafPinProvider for SKY130LeafPins {
    fn direction_of(&self, macro_name, pin_name, pin_idx) -> Direction;
    fn width_of(&self, macro_name, pin_name) -> Option<SVerilogRange>;
}
}

For SKY130, direction_of() is a large match statement covering ~80 cell types with all their pin names. This is tedious but straightforward – for each cell, list which pins are inputs and which are outputs.

Sources for pin directions:

• The PDK’s Liberty (.lib) files list pin directions
• The PDK’s behavioral Verilog models declare input/output ports
• LEF files also contain pin direction information

For a new PDK: Implement the trait for all cells that appear in your target netlists. You can start with just the cells used in your design and add others as needed.

Step 4: Cell Classification

Reference: src/sky130_pdk.rs – is_sequential_cell(), is_tie_cell(), is_multi_output_cell()

Three classification functions control how cells are processed during AIG construction:

Sequential cells (DFFs and latches)

These are handled specially in the AIG builder – their outputs become state elements rather than combinational logic.

Critical: Use an explicit whitelist, not prefix matching. PDK naming collisions will silently break simulation if you guess wrong (e.g., SKY130’s dlygate4sd3 starts with “dl” but is a combinational delay buffer, not a latch).

Derivation method: Grep the PDK’s behavioral Verilog models for DFF/latch primitives:

for cell in $(ls vendor/<pdk>/cells/); do
    vfile="vendor/<pdk>/cells/$cell/<pdk>__${cell}.behavioral.v"
    if [ -f "$vfile" ] && grep -qE 'udp_dff|udp_dlatch' "$vfile"; then
        echo "$cell"
    fi
done

For PDKs that don’t use Verilog UDPs, look for always @(posedge blocks or check the Liberty file’s ff and latch groups.

Tie cells

Cells that produce constant 0 or 1 (e.g., SKY130’s conb with HI/LO pins).

Multi-output cells

Cells with more than one output (e.g., half-adder ha with SUM and COUT, full-adder fa). These need special handling because the AIG builder processes one output pin at a time.

Step 5: Behavioral Model Loading

Reference: src/sky130_pdk.rs – load_pdk_models(), parse_functional_model(), parse_udp()

Jacquard decomposes PDK cells to AIG primitives (AND gates and inversions) by parsing their functional Verilog models. The expected file structure:

vendor/<pdk>/
  cells/
    <cell_type>/
      <pdk>__<cell_type>.functional.v    # Gate-level behavioral model
  models/
    <udp_name>/
      <pdk>__<udp_name>.v               # Verilog UDP definitions

Functional models

These are gate-level Verilog using primitives like and, or, nand, nor, not, xor, xnor, buf. The parser (parse_functional_model()) extracts these into a topologically-ordered list of BehavioralGate structures.

Example (sky130_fd_sc_hd__o21ai.functional.v):

module sky130_fd_sc_hd__o21ai (Y, A1, A2, B1);
    output Y;
    input  A1, A2, B1;
    wire or0_out;
    or  or0  (or0_out, A2, A1);
    nand nand0 (Y, B1, or0_out);
endmodule

UDP models

Some cells (typically muxes) use Verilog User-Defined Primitives with truth tables. The parser (parse_udp()) converts these to a row-based representation, which is then evaluated as sum-of-products during AIG decomposition.

What’s loaded

Only models for cell types actually present in the design are loaded. Sequential cells are skipped (their behavior is hardcoded in the AIG builder). Tie cells are also skipped (constant generation is trivial).

For a new PDK: If the PDK uses the same Verilog gate primitive syntax, the existing parsers should work. If it uses behavioral Verilog (assign statements, always blocks), the parser would need extension.

Step 6: AIG Decomposition

Reference: src/sky130_pdk.rs – decompose_with_pdk(), decompose_from_behavioral()

The decomposition converts each combinational cell to a set of 2-input AND gates with optional inversions:

1. Map the cell’s input pin names to AIG pin indices via CellInputs
2. Walk the behavioral model’s gate list in topological order
3. For each gate, build the equivalent AIG sub-graph:
   • and/nand -> AND gate (with optional output inversion)
   • or/nor -> De Morgan’s: OR(a,b) = NOT(AND(NOT a, NOT b))
   • xor/xnor -> Four AND gates: XOR(a,b) = NOT(AND(NOT(AND(a, NOT b)), NOT(AND(NOT a, b))))
   • buf/not -> Pass-through with optional inversion
   • UDP -> Sum-of-products from truth table
4. Record the output with cell origin (for SDF timing annotation)

CellInputs struct

CellInputs has named fields for all possible input pins across all SKY130 cells (A, B, C, D, A_N, B_N, S, S0, S1, CIN, SET_B, RESET_B, etc.). The set_pin() method maps netlist pin names to AIG pin indices.

For a new PDK: If the PDK introduces pin names not in the current struct, add new fields.

Step 7: AIG Builder Integration

Reference: src/aig.rs – get_sky130_dependencies(), sky130_preprocess(), sky130_postprocess()

The AIG builder processes cells in three phases during topological traversal:

Dependencies (what must be built before this cell)

• Tie cells: No dependencies
• Sequential cells: Only SET_B and RESET_B pins (the data input D is handled by the DFF mechanism, not combinational decomposition)
• Combinational cells: All input pins

Preprocessing (before dependencies are resolved)

• Sequential cells: Create a DFF output AIG pin. This establishes the state element before the combinational cone driving it is built.

Postprocessing (after all dependencies are resolved)

• Tie cells: Wire HI to constant-1, LO to constant-0
• Sequential cells: Apply reset/set logic: Q = AND(OR(Q_state, NOT SET_B), RESET_B) (active-low semantics)
• Combinational cells: Call decompose_with_pdk() and wire the resulting AND gates into the AIG

For a new PDK: The three-phase structure is reusable. You need PDK-specific implementations of each phase that handle the new cell types’ pin names and reset/set conventions.

Step 8: CLI Integration

Reference: src/bin/jacquard.rs

The load_design function detects the library and creates the netlist with the appropriate pin provider:

#![allow(unused)]
fn main() {
let lib = detect_library_from_file(&args.netlist_verilog)?;
let netlistdb = match lib {
    CellLibrary::SKY130 => NetlistDB::from_sverilog_file(&paths, &SKY130LeafPins),
    CellLibrary::AIGPDK => NetlistDB::from_sverilog_file(&paths, &AIGPDKLeafPins()),
    CellLibrary::Mixed => panic!("Mixed libraries not supported"),
};
}

For a new PDK: Add a match arm for the new library.

Testing Strategy

Unit tests

1. Cell type extraction: Verify prefix/suffix stripping

2. Pin directions: Spot-check common cells

3. Behavioral model parsing: Parse each cell type, verify gate count

4. Decomposition correctness: For each combinational cell, exhaustively test all input combinations against the PDK’s truth table:

   #![allow(unused)]
   fn main() {
   #[test]
   fn test_all_cells_vs_pdk() {
       let pdk = load_test_pdk();
       for (cell_type, model) in &pdk.models {
           // For each input combination:
           //   1. Evaluate behavioral model directly
           //   2. Decompose to AIG and evaluate AIG
           //   3. Assert outputs match
       }
   }
   }

   This test exists in src/sky130_pdk.rs as test_all_cells_vs_pdk and covers every combinational cell against every input combination.

Integration tests

1. Small test circuit: Synthesize a simple design (DFF + some gates) to the new PDK and verify simulation output matches a reference (e.g., iverilog)
2. Flash boot test: If targeting an SoC, verify the CPU boots and reads from flash (this exercises sequential logic, combinational cones, and IO)

File Checklist

For a complete PDK integration, you need:

Common Pitfalls

• Cell name collisions: Do not use prefix matching for cell classification. dlygate4sd3 starts with “dl” but is not a latch. Always derive the exhaustive list from behavioral models.

• Active-low vs active-high resets: SKY130 uses active-low RESET_B and SET_B. Other PDKs may use active-high. Get this wrong and every DFF will be stuck.

• Multi-output cells: The AIG builder processes one output pin at a time. If a cell has both Q and Q_N outputs (e.g., dfbbp), the second output must be derived from the first (Q_N = NOT Q), not decomposed independently.

• Liberty file size: SKY130’s liberty files are 12MB+. If your PDK has similarly large files, ensure the parser doesn’t OOM or timeout.

• Power/ground pins: Post-layout netlists often include VPWR/VGND pins. Use the unpowered netlist variant (.nl.v not .pnl.v in OpenLane2) or handle power pins as constants in the pin provider.

• Hold-time repair buffers: P&R tools insert delay buffers (like dlygate4sd3) that must be treated as combinational. If your PDK’s delay cells have names that collide with sequential cell prefixes, the whitelist approach prevents misclassification.

Adding third-party IP via runtime manifest

If you’re adding a memory macro, IO pad, hard block, or filler library — anything that doesn’t need new AIG decomposition rules — the runtime cell-library pathway (Decision 0010) is the right route. No Jacquard PR required. Ship a Verilog blackbox file plus a TOML manifest alongside your design.

Step 1: Provide the cell’s Verilog interface

The blackbox just declares the cell’s module + port directions. The foundry typically ships this (<library>__blackbox.v). Example for the OCD GF180MCU SRAM:

module gf180mcu_ocd_ip_sram__sram1024x8m8wm1 (CLK, CEN, GWEN, WEN, A, D, Q);
  input CLK;
  input CEN;
  input GWEN;
  input [7:0] WEN;
  input [9:0] A;
  input [7:0] D;
  output [7:0] Q;
endmodule

Step 2: Write the TOML manifest

Co-locate <library>.cells.toml next to <library>.v (it autoloads when present) or pass it via --cell-manifest:

schema_version = "1.0"

[cells.gf180mcu_ocd_ip_sram__sram1024x8m8wm1]
kind = "ram"

Recognised kind values in v1.0: std, dff, latch, clock_gate, ram, filler, endcap, tap, io_pad_input, io_pad_output, io_pad_bidir, delay, multi_output, tie_high, tie_low.

Step 3: Invoke jacquard with the manifest

jacquard sim my_chip.v stim.vcd out.vcd 1 \
    --cell-library deps/gf180mcu_ocd_ip_sram/cells/gf180mcu_ocd_ip_sram__sram1024x8m8wm1/gf180mcu_ocd_ip_sram__sram1024x8m8wm1__blackbox.v

The --cell-library flag is repeatable for multi-IP designs.

What kind = "ram" delivers — opaque vs explicit-port modes

There are two modes depending on whether the manifest includes a [cells.NAME.ram] port-mapping sub-table:

Opaque mode (no ram sub-table, schema v1.0+): the cell’s output pins become X-source slots in the AIG. The SRAM’s internal memory behaviour is not modelled. Sufficient for designs whose CPU executes from boot ROM / register file and never reads SRAM contents at the timescales Jacquard simulates.

Explicit-port mode (with ram sub-table, schema v1.1+, decision 0011): outputs are wired to a real AIG-backed RAMBlock, writes populate per-entry storage, reads return what was written. Real memory modelling end-to-end. Use this when the CPU reads its own SRAM (the common case for any design beyond heartbeat verification).

Schema (full example, mirroring the upstream OCD GF180MCU SRAM):

schema_version = "1.1"

[cells.gf180mcu_ocd_ip_sram__sram1024x8m8wm1]
kind = "ram"

[cells.gf180mcu_ocd_ip_sram__sram1024x8m8wm1.ram]
depth = 1024
width = 8
clock        = { pin = "CLK", edge = "pos" }
chip_enable  = { pin = "CEN", polarity = "low" }
write_enable = { pin = "GWEN", polarity = "low" }
write_mask   = { pin = "WEN", polarity = "low", granularity = "bit" }
address      = "A"
data_in      = "D"
data_out     = "Q"

Field semantics, defaults, and the multi-port-SRAM/async/wider-than-32-bit out-of-scope items are documented in Decision 0011. Polarity defaults to low; clock edge defaults to pos; mask granularity defaults to bit. All three control pins (chip_enable / write_enable / write_mask) are optional — omit them for sync SRAMs without those signals.

Preloading SRAM contents at sim start

Once a SRAM is in explicit-port mode, its contents can be preloaded from an ELF file via sim_config.json:

{
  "sram_init": {
    "elf_path": "build/firmware.elf"
  }
}

The ELF’s PT_LOAD segments are packed into the SRAM’s backing storage before tick 0; the lowest loadable virtual address is taken as SRAM address 0. Single-SRAM designs only — multi-SRAM instance-targeting is a future schema extension (issue #80).

Other kinds

• filler, endcap, tap — physical-only, contribute no logic.
• io_pad_input / io_pad_output / io_pad_bidir — pad-level behaviour (parallel to the built-in gf180mcu_ws_io__* family).
• dff, latch, clock_gate, delay, multi_output — recognised but the v1.0 schema doesn’t yet expose enough port semantics to drive AIG construction for these. Coming in the port-mapping schema (future decision). For now, declaring these kinds documents intent without changing behaviour.

Troubleshooting VCD Input Issues

This guide helps debug VCD input problems where GEM simulations produce incorrect results or warn about missing signals.

VCD Scope Auto-Detection (Recommended)

NEW: GEM now automatically detects the correct VCD scope containing your design’s ports. In most cases, you don’t need to specify --input-vcd-scope manually.

How Auto-Detection Works

When you run jacquard sim without specifying --input-vcd-scope, GEM:

1. Extracts the list of required input ports from your synthesized design
2. Searches the VCD file for scopes containing all required ports
3. Tries common DUT scope names first: dut, uut, DUT, UUT, or your module name
4. Falls back to any scope that contains all required ports
5. Logs which scope was selected for transparency

Example Output

INFO No VCD scope specified - attempting auto-detection
DEBUG Searching for VCD scope containing 4 input ports
DEBUG Required ports: {"din_valid", "clk", "reset", "din"}
INFO Auto-detected VCD scope: safe_tb/uut (matched common pattern 'uut')

Manual Override

If auto-detection fails or selects the wrong scope, use --input-vcd-scope to specify manually:

# Slash-separated path to the DUT scope
jacquard sim design.gv input.vcd output.vcd 8 \
    --input-vcd-scope "testbench/dut"

# For nested hierarchies
jacquard sim design.gv input.vcd output.vcd 8 \
    --input-vcd-scope "top_tb/subsystem/my_module"

Note: Use slash separators (/), not dots (.).

Symptom: Missing Primary Input Warnings

WARN (GATESIM_VCDI_MISSING_PI) Primary input port (HierName(), "reset", None) not present in the VCD input
WARN (GATESIM_VCDI_MISSING_PI) Primary input port (HierName(), "din", Some(3)) not present in the VCD input

Root Cause

GEM expects VCD signals at absolute top-level with no module hierarchy prefix. The signal names must exactly match the synthesized module’s port names.

How to Check

1. Inspect your VCD file:

grep '\$var' your_input.vcd | head -20

2. Look for module scopes:

grep '\$scope module' your_input.vcd

3. Check synthesized module ports:

head -20 your_design_synth.gv

What GEM Expects

Correct - Signals at top level:

$timescale 1ns/1ns
$var reg 1 ! clk $end
$var reg 1 " reset $end
$var reg 4 # din [3:0] $end
$var reg 1 $ din_valid $end
$var wire 1 % unlocked $end
$enddefinitions $end
$dumpvars
0"
0$
0%
1!
#10
1"
#20
b1100 #
1$

Incorrect - Signals scoped under module:

$scope module testbench $end
  $scope module dut $end
    $var wire 1 ! clk $end
    $var wire 1 " reset $end
    $var wire 4 # din [3:0] $end
    ...
  $upscope $end
$upscope $end

Solution 1: Flat VCD Generation

Create a testbench that dumps signals at absolute top level:

module testbench;

reg clk = 0;
reg reset;
reg [3:0] din;
reg din_valid = 0;
wire unlocked;

// DUT instantiation
your_module dut (
    .clk(clk),
    .reset(reset),
    .din(din),
    .din_valid(din_valid),
    .unlocked(unlocked)
);

always #10 clk = !clk;

initial begin
    // CRITICAL: Dump signals at top level (depth 1)
    // NOT inside module hierarchy!
    $dumpfile("output.vcd");
    $dumpvars(1, clk, reset, din, din_valid, unlocked);

    // Test sequence
    reset = 1;
    #60;
    reset = 0;

    // ... your test stimulus ...

    #200;
    $finish;
end

endmodule

Key Point: $dumpvars(1, signal1, signal2, ...) dumps individual signals at the current scope level, not inside child modules.

Compile and Run

# For synthesis-compatible testbench
iverilog -DSYNTHESIS -o sim your_design.v testbench.v
./sim

# Check VCD structure
grep '\$scope' output.vcd  # Should be minimal or none
grep '\$var' output.vcd | head -10

Solution 2: Post-Process VCD (Advanced)

If you can’t change the testbench, post-process the VCD to flatten hierarchy:

#!/usr/bin/env python3
"""Flatten VCD hierarchy to top level"""

import sys

def flatten_vcd(input_vcd, output_vcd):
    with open(input_vcd) as inf, open(output_vcd, 'w') as outf:
        in_scope = False
        scope_depth = 0

        for line in inf:
            # Track scope depth
            if line.strip().startswith('$scope'):
                scope_depth += 1
                if scope_depth == 1:
                    continue  # Keep root scope
                in_scope = True
                continue
            elif line.strip().startswith('$upscope'):
                scope_depth -= 1
                if in_scope and scope_depth == 0:
                    in_scope = False
                continue

            # Skip signals inside nested scopes, keep only top-level
            if in_scope and line.strip().startswith('$var'):
                continue  # Skip nested module signals

            outf.write(line)

if __name__ == '__main__':
    flatten_vcd(sys.argv[1], sys.argv[2])

Usage:

python3 flatten_vcd.py hierarchical.vcd flat.vcd

Solution 3: VCD Scope Option (Experimental)

GEM provides --input-vcd-scope to specify which module hierarchy to read:

cargo run -r --features metal --bin jacquard -- sim \
    design.gv input.vcd output.vcd 48 \
    --input-vcd-scope module_name

Known Issue: Currently, signal matching still fails even with correct scope specified. This is under investigation.

Diagnostic Checklist

1. Verify Signal Names Match

Synthesized Module:

grep "^module\|input\|output" design_synth.gv

Output:

module safe(clk, reset, din, din_valid, unlocked);
  input clk;
  input reset;
  input [3:0] din;
  input din_valid;
  output unlocked;

VCD Signals:

grep '\$var.*\(clk\|reset\|din\|unlocked\)' input.vcd

Output should match synthesized port names exactly.

2. Check Signal Bit Widths

Multi-bit signals must have correct indices:

Synthesized: input [3:0] din;

VCD:

$var reg 4 # din [3:0] $end

GEM expects separate indices: din[3], din[2], din[1], din[0]

3. Verify Timestamp Format

GEM expects integer timestamps (not real numbers):

Correct:

#0
#10
#20

Incorrect:

#0.0
#10.5
#20.25

4. Check Timescale

Ensure VCD timescale matches simulation expectations:

$timescale 1ns $end

or

$timescale 1ps $end

Clock periods in testbench should use same time unit.

Validation Steps

After fixing VCD issues, validate GEM is reading inputs correctly:

1. Run with CPU Verification

cargo run -r --features metal --bin jacquard -- sim \
    design.gv input.vcd output.vcd 48 \
    --check-with-cpu

This compares GPU results against CPU gate-level simulation. Should print:

[INFO] sanity test passed!

2. Compare Output VCD with Reference

Run same design with iverilog:

iverilog -o reference_sim design.v testbench.v
./reference_sim  # Generates reference.vcd

Compare outputs:

# Check if unlocked signal toggles the same in both
grep '^[01]!' gem_output.vcd
grep '^[01]!' reference.vcd

3. Check Cycle Count

cargo run -r --features metal --bin jacquard -- sim \
    design.gv input.vcd output.vcd 48 \
    2>&1 | grep "total number of cycles"

Should match your testbench’s simulation time / clock period.

Common Pitfalls

1. Testbench Inside `ifndef SYNTHESIS

If testbench is only compiled when SYNTHESIS is not defined:

`ifndef SYNTHESIS
module testbench;
  // ...
endmodule
`endif

You must compile without -DSYNTHESIS for VCD generation:

iverilog -o sim design.v testbench.v  # No -DSYNTHESIS!

But the DUT must be compiled with -DSYNTHESIS if it has non-synthesizable constructs:

# Separate compilation
iverilog -DSYNTHESIS -c design.v
iverilog -o sim design.v testbench.v

2. X/Z Values in VCD

GEM may not handle unknown (X) or high-impedance (Z) values correctly:

$dumpvars
x"  # reset = X
bxxxx #  # din = XXXX

Solution: Initialize all inputs in testbench:

initial begin
    reset = 0;  // Don't leave uninitialized
    din = 4'h0;
    din_valid = 0;
end

3. Missing Clock Signal

If VCD doesn’t include clock:

WARN (GATESIM_VCDI_MISSING_PI) Primary input port (HierName(), "clk", None) not present

Ensure:

• Clock is generated in testbench
• Clock is included in $dumpvars
• Clock signal name matches synthesized netlist exactly

Example: Working Flat VCD Testbench

// testbench_flat.v - Generates GEM-compatible VCD
module testbench_flat;

// Declare all signals at top level
reg clk = 0;
reg reset = 1;
reg [3:0] din = 4'h0;
reg din_valid = 0;
wire unlocked;

// DUT instantiation
safe dut (
    .clk(clk),
    .reset(reset),
    .din(din),
    .din_valid(din_valid),
    .unlocked(unlocked)
);

// Clock generation
always #10 clk = !clk;  // 20ns period = 50MHz

// Test sequence
initial begin
    // CRITICAL: Dump at top level (depth 1)
    $dumpfile("safe_flat.vcd");
    $dumpvars(1, clk, reset, din, din_valid, unlocked);

    // Reset phase
    reset = 1;
    #60;  // 3 clock cycles
    reset = 0;
    #11;  // Small offset from clock edge

    // Apply test stimulus
    din = 4'hc;
    din_valid = 1;
    #20;

    din = 4'h0;
    #20;

    din = 4'hd;
    #20;

    din = 4'he;
    #20;

    din_valid = 0;
    #40;

    $finish;
end

endmodule

Compile and test:

# Compile (DUT must be SYNTHESIS-compatible)
iverilog -DSYNTHESIS -o sim safe.v testbench_flat.v

# Run simulation
./sim

# Verify VCD structure
echo "=== VCD Scopes ==="
grep '\$scope' safe_flat.vcd

echo -e "\n=== VCD Signals ==="
grep '\$var' safe_flat.vcd

# Should show signals at top level, no nested $scope modules

Still Having Issues?

1. Enable debug logging:

   RUST_LOG=debug,vcd_ng=trace cargo run -r --features metal --bin jacquard -- sim <args> 2>&1 | tee debug.log
   

2. Check with minimal test:

   • Create simplest possible design (single DFF)
   • Generate flat VCD
   • Verify GEM can read it correctly

3. Report issue with:

   • Synthesized .gv file
   • Input VCD file
   • GEM command line
   • Error messages or unexpected output

Document Version: 1.0 Last Updated: 2025-01-08 Related: architecture/simulation-engine.md

Authoring architecture diagrams

Diagrams in the docs are written as D2 text and rendered to inline SVG at build time. They are theme-adaptive: colours and fonts come from the active mdBook theme, so a diagram reads correctly in light, navy, coal, ayu, and rust, and recolours on a live theme switch. This page is the how-to; the why (D2 over mermaid, the layout choices) is in the redesign spike.

Adding a diagram

Write a fenced d2 block in any docs/*.md file:

```d2
vars: { d2-config: { pad: 16 } }
direction: down

a: First box
b: Second box
a -> b: does a thing
```

The mdbook-d2 preprocessor (configured in book.toml) shells out to the d2 binary and replaces the block with an inline SVG. Nothing else is needed — no image files, no per-diagram wiring.

Authoring rules

These keep diagrams legible and consistent. They come out of the spike; follow them unless you have a specific reason not to.

• Lay regions out vertically; reserve horizontal flow for genuine pipelines. On-screen text size is the native font scaled by content_column_width / diagram_native_width — a wide diagram shrinks its text into an unreadable hairline when fit to the column. A tall, narrow diagram keeps text near body size and scrolls vertically, which is natural on the web. Set direction: down at the top level and only use direction: right inside a container that is a real left-to-right pipeline.
• One abstraction level per diagram. A single overview, then per-stage detail if needed — never one diagram that tries to show everything.
• Show what, not why. The rationale belongs in the decision record the diagram’s page links to, not in the picture.
• Keep pad small. vars: { d2-config: { pad: 16 } } trims D2’s default 100px canvas border (mdbook-d2 exposes no pad config, so set it in the source).
• Do not set colours or fonts in the D2 source. The stylesheet owns them (next section). The one exception is a sentinel fill used as a colour hook.

The global layout engine is elk (set in book.toml), which handles containers and edge routing better than the default dagre.

How theming works

D2 bakes one fixed palette into the SVG. theme/d2-diagrams.css repaints it from the active mdBook theme’s CSS variables (--fg, --bg) — author CSS overrides the SVG’s baked fill=/stroke= attributes, so no SVG post-processing is needed. The same stylesheet strips the <pre> code-block chrome that mdbook-d2 wraps each SVG in, and scales the diagram to the content column.

Three details in that stylesheet are load-bearing, and each fixes a real bug — know them before you touch it:

• Shape colours need !important. D2 sets shape colour through an embedded <style> using .d2-<hash> .stroke-B1-style selectors, specificity (0,2,0), with a per-diagram hash we can’t match directly. Element selectors score lower and lose, so shape fill/stroke overrides use !important. (Text scores high enough via .text to win without it.)
• Match rect[class], never bare rect. D2’s connection arrows use an SVG <mask> whose unclassed <rect>s mean “visible here / hole there”. A bare rect rule repaints those mask rects too, turning the mask opaque and clipping every connection line to nothing — invisible arrows whose stroke still computes correctly. Rendered shapes always carry a D2 colour class; mask rects never do, so rect[class] skips them.
• Region accents use a sentinel fill. D2 gives no stable per-shape hook (user classes aren’t emitted; the auto-generated classes are base64 of the shape path). To colour a specific region, set a sentinel hex in the D2 source (style.fill: "#4f8bff") and remap it in CSS (rect[fill="#4f8bff" i] { fill: rgba(...) !important; ... }). The remap is a translucent tint, which is what makes it read over both light and dark backgrounds. The cosim diagram uses #4f8bff / #3cbe78 / #aa78f0 for its three regions.

Verifying a diagram’s styling

A D2 PNG render (d2 file.d2 out.png) shows the baked palette, not the themed result, so it cannot confirm the styling. Render the built page in headless Chrome instead:

mdbook build
chrome=/Applications/Google\ Chrome.app/Contents/MacOS/Google\ Chrome
"$chrome" --headless=new --disable-gpu --force-device-scale-factor=2 \
  --window-size=1400,2600 --screenshot=/tmp/page.png \
  "file://$PWD/book/architecture/<page>.html"

To force a theme, seed localStorage before mdBook’s boot script runs by injecting <script>localStorage.setItem('mdbook-theme','navy')</script> right after <head> in a copy of the built HTML. To read the computed style of an element (e.g. to confirm a connection line’s stroke actually resolved), inject a getComputedStyle probe that writes the result into document.title and read it back with --dump-dom. This is how the invisible-arrows bug was found — the PNG render hid it completely.

Build dependencies

The docs build needs the d2 binary and the mdbook-d2 preprocessor. Locally:

brew install d2            # or: curl -fsSL https://d2lang.com/install.sh | sh
cargo install mdbook-d2

CI installs both in the Documentation job (see .github/workflows/ci.yml): d2 from its release tarball into ~/.cargo/bin, and mdbook-d2 from crates.io.

Development

For working on Jacquard itself. If you only want to run it, start at Installation and Getting Started.

Build

You need the Rust toolchain (2021 edition) and the GPU SDK for the backend you’re building. Nothing else.

git submodule update --init --recursive

The submodule init is load-bearing, not boilerplate — and its failure is silent. build.rs bakes the standard-cell pin tables into the binary by reading the vendored PDK cell libraries (vendor/sky130_fd_sc_hd, vendor/gf180mcu_fd_sc_mcu7t5v0, vendor/gf180mcu_fd_sc_mcu9t5v0). If one of those submodules is absent, build.rs skips it and the build succeeds — but the resulting binary carries an incomplete pin table and panics at runtime on any netlist using an omitted cell (Unknown SKY130 cell (no pin-table entry for type '…')), which looks like a Jacquard bug rather than a missing checkout. So the --recursive above is required, and any script that inits submodules selectively — a CI matrix that only needs vendor/eda-infra-rs, or a historical-build / bisect harness — must also init the PDK cell submodules, or it will silently produce a crippled binary.

The GPU backend is a feature, and you pick exactly one:

cargo build -r --features metal --bin jacquard   # macOS, Apple Silicon
cargo build -r --features cuda  --bin jacquard   # NVIDIA, CUDA toolkit
cargo build -r --features hip   --bin jacquard   # AMD, ROCm

Two features are worth knowing about:

• synth adds the embedded YoWASP Yosys engine, which is what lets sim/cosim take behavioral RTL instead of a gate-level netlist (see Accepted RTL Surface). It costs real compile time — wasmtime and cranelift — so it is opt-in for local builds, and on for release binaries. Combine it with a backend: --features metal,synth.
• No GPU feature at all still builds dump-paths, so timing analysis and netlist validation work on a machine with no GPU SDK installed.

Test

cargo test                       # library + integration tests, no GPU needed
cargo bench --bench event_buffer # criterion micro-benchmarks, no GPU needed
cargo bench --bench xprop

For the GPU path, two flags matter more than the test suite:

jacquard sim ... --check-with-cpu        # run a CPU baseline and compare
jacquard sim ... --max-clock-edges 1000  # bound a long run while bisecting
                                         # (edges, not cycles: 2 per cycle)

--check-with-cpu is the one to reach for when a kernel change produces plausible-looking waveforms: it is the difference between “it ran” and “it is right”.

Benchmark designs (NVDLA, Rocket, Gemmini) live in benchmarks/dataset/, a submodule. NVDLA is the smallest and the usual first thing to try.

Where things live

The pipeline reads left to right:

NetlistDB → AIG → StagedAIG → Partitions → FlattenedScript → GPU kernel

Simulation engine walks each stage. For why it is shaped this way, Decision 0014 explains the AIG choice and Decision 0015 the boomerang execution model.

The constraint that shapes everything

A GPU block is a fixed budget, and src/pe.rs has to fit each partition inside it:

• at most 8191 unique inputs and 8191 unique outputs per partition (for SRAMs and DFFs, outputs count every enable and bus pin, and holes mean the effective capacity can be as low as 4095);
• at most 4095 intermediate pins alive at any stage;
• at most 64 SRAM output groups — that is 8192 / (32 × 4).

When a design doesn’t fit you get single endpoint cannot map. The answer is usually --level-split, which forces more stages so each one is smaller:

jacquard sim ... --level-split 30
jacquard sim ... --level-split 20,40

This is the first thing most people hit on a large design, and it is a mapping limit rather than a bug.

Optional tooling

Each of these is needed by one workflow and nothing else. Editing Rust, C++ or kernel sources needs none of them: the timing-IR bindings are checked in, and OpenSTA is only for the timing-correctness corpus.

Python tooling (PDK fetchers, harness utilities) belongs in the workspace’s uv dev group rather than an ad-hoc pip install: add it under [dependency-groups] in the root pyproject.toml, uv sync --group dev, and run it with uv run.

Debugging

Jacquard’s own tools, each with a page:

• Signal Tracing — surface internal nets in the output VCD. netlist-graph finds the names to trace: uv run netlist-graph drivers <netlist> <signal>.
• Debugging X Values — find why a signal went x, statically, instead of trace-guess-rerun.
• Bus Transaction Tracing — decode on-chip bus transfers rather than reading raw wires.
• Timing Violations — GPU-side setup/hold checks.
• Cosim Perf Report and GPU Frame Capture — where the time actually goes in a cosim run.

Docs

The docs are an mdBook, sourced from docs/:

mdbook serve   # http://localhost:3000

SUMMARY.md is the table of contents, and a page missing from it isn’t rendered at all — which is why docs/scripts/check_doc_links.py validates links against the rendered page set rather than against files on disk. Run it before pushing docs; CI does too.

The published site keeps main’s docs at the root and a frozen copy of each release under /vX.Y.Z/, with a picker to move between them. Release notes link to the pinned copy, so a link in an old release keeps meaning what it meant. See Release Process.

Conventions

• Decisions record choices worth understanding later. A record is immutable and past-tense: when a decision changes, a new record supersedes the old rather than rewriting it. A record’s status is a claim about the code, so check it against the code. Current state lives in the architecture reference, not the record.
• Plans hold work in progress; deferred work gets a plan and an issue rather than being lost.
• Handoffs are working memory, not history. One per active thread, folded into the decisions or plans and deleted when resolved.
• Prose style is the house style for the docs’ writing.
• Release Process covers cutting a release, the RC-first flow, and what’s automated.

The through-line: a sentence that says how the tool behaves is a verifiable claim. That applies to docs, --help text, and a decision’s status alike; check it against the code before writing it.

Prose style

House style for Jacquard’s prose docs: the docs/ book, decision records, and plans. It covers how sentences read. How docs are structured — the reference / decision / spike split, the three kinds of claim, stable-ID cross-links — lives in the decisions README and the redesign spike; this page is only about the writing.

Voice

Be direct and have opinions. State the point first, then support it. Use specific names and numbers, not vague claims (BATCH_SIZE = 1024, not “a fixed batch size”). Trust the reader to see what matters without labels like “significant” or “important”. Write to be understood by a non-native speaker and by an LLM reading a single retrieved section, so keep each section self-contained and restate its subject rather than leaning on “as described above”.

Claims track the code

A present-tense sentence in a reference doc is a verifiable claim about the code — check it before writing it, and keep it current when the code moves. Rationale is past-tense and lives in decision records, which don’t get rewritten when the code changes. This is the whole point of the layer split; see the decisions README.

Sentences

• Use contractions: “it’s”, “don’t”, “won’t”.
• Vary sentence and paragraph length. Don’t write uniform blocks.
• Avoid mid-sentence interjections. Don’t bracket a qualifier inside a sentence with paired dashes or parentheses. Promote it to its own sentence or move it to the end. Write “Only X-capable partitions run the X-aware kernel. In a typical SoC after reset, under 5% are X-capable.” — not “partitions — typically under 5% after reset — run the kernel”.
• A single dash introducing an appositive or definition is fine, and is the house pattern for naming a thing: “a hierarchical binary tree — the boomerang”. Keep it to the end of a clause, not the middle of one.

Lists

Bold-term definition lists are house style for a genuine enumeration where the reader should weigh each item:

• NetlistDB parses structural Verilog into a flattened database.
• AIG rewrites combinational logic to one uniform node type.

Use them when the terms are the subject. Don’t use them as a crutch to avoid writing a paragraph, and don’t pad a two-item list into bullets.

Banned words and tics

These are the most flagged AI-writing markers. Never use them:

delve, dive into, navigate (figurative), underscore, bolster, foster, harness (figurative), leverage, unpack (figurative), shed light on, pave the way, pivotal, groundbreaking, cutting-edge, transformative, game-changing, seamless, intricate, multifaceted, holistic, testament, landscape (figurative), realm.

And these structures, which mimic insight without providing it:

• “It’s not just X — it’s Y” / “Not only X, but Y”
• “This isn’t about X. It’s about Y.”
• “It’s important / worth noting that…”, “At its core…”, “When it comes to…”
• Opening with a sweeping statement about the field, or closing with an inspirational wrap-up. Start and end on substance.

Before committing a doc

1. Read it aloud. Does a sentence sound like a press release or a marketing page? Rewrite it.
2. Are you saying the same thing twice in different words? Say it once.
3. Is every present-tense claim true of the code as it stands today?

Release Process

Lightweight by design. Jacquard is a single-binary Rust project with vendor/ submodules; releases are git tags + a CHANGELOG entry. No crates.io publication, no pre-built binaries (until/unless that demand surfaces).

When to release

Cut a release when:

• A user-visible feature or fix lands that you want consumers to be able to pin against.
• Schema or CLI changes happened (--timing-report JSON, CLI flags) and consumers need a stable reference point.
• A meaningful chunk of work in docs/plans/ is closed (e.g. a Phase exits all criteria).

There is no fixed cadence.

Versioning

SemVer, starting once the first numbered release ships. Pre-1.0 versions (0.x.0) carry the standard SemVer caveat: minor bumps may include breaking changes; the public contracts (--timing-report schema, IR layout) are documented in their own decisions and follow stricter rules.

Stable contracts (additive-only, breaking changes require a major bump and a deprecation window):

• --timing-report JSON schema — src/timing_report.rs::SCHEMA_VERSION, governed by Decision 0008.
• Timing IR FlatBuffers schema — crates/timing-ir/schemas/timing_ir.fbs, governed by Decision 0002.

CLI flags, log message formats, and --timing-summary text output are not stable parseable contracts; consumers that need to script against them should use --timing-report JSON.

Steps

For maintainers cutting a release:

1. Verify CI is green on main for all three GPU backends (Metal, CUDA, HIP) plus the unit-test, opensta-to-ir, and lint jobs. If any GPU runner is offline, hold the release until it’s restored — see .github/workflows/ci.yml. Do not ship a binary the CI hasn’t built.

2. Roll the [Unreleased] section in CHANGELOG.md into a numbered version block. Format follows Keep a Changelog. Update the link references at the bottom of the file. Leave a fresh empty [Unreleased] section at the top.

3. Bump the Rust crate version to match: python3 scripts/bump_version.py <X.Y.Z>. The three first-party Rust crates (jacquard, opensta-to-ir, timing-ir) ship together in one tarball and carry a single shared version, so this script sets all three at once — never edit their [package].version by hand. Then cargo build to update Cargo.lock. (netlist-graph versions independently — see its own netlist-graph-v* tag flow.) The release workflow re-runs bump_version.py --check <tag> as a verify-guard and aborts before publishing if the tag and the crates disagree.

4. Commit: chore: release v<X.Y.Z> with the standard Co-developed-by trailer.

5. Tag: git tag -a v<X.Y.Z> -m "v<X.Y.Z>" then git push --tags.

6. Create a GitHub release from the tag. Body = the CHANGELOG section for that version. No artefacts attached unless someone has asked for them.

Homebrew tap (automated)

The Homebrew formula is auto-bumped by release.yml — no manual step (this closes the drift that had left the tap stale at 0.2.3). On every release tag the bump-tap job rewrites packaging/homebrew/jacquard.rb’s url/version/sha256 from the just-published tarball + .sha256 and pushes Formula/jacquard.rb to the tap:

• final release → gpu-eda/homebrew-tap (brew install gpu-eda/tap/jacquard);
• prerelease (RC) → gpu-eda/homebrew-tap-prerelease (brew install gpu-eda/tap-prerelease/jacquard), so RCs are brew install-able for staging without touching the stable channel.

packaging/homebrew/jacquard.rb is the template — edit it only to change the formula’s structure (deps, install, test); its version pin is a placeholder the job overwrites. Requires the one-time org setup: secrets.HOMEBREW_TAP_TOKEN (a token with contents: write on both tap repos) and the gpu-eda/homebrew-tap-prerelease repo.

Release notes & versioned docs (automated)

Release notes come from the CHANGELOG, and doc links are pinned to the release — both automatic in release.yml:

• Notes body = the CHANGELOG section for the tag’s version. A prerelease (X.Y.Z-rc.N) has no dated section, so the extractor falls back to [Unreleased] — RCs ship the same curated draft you’ll ship at promotion. So: write the notes in [Unreleased].
• Lead with a user-facing overview. Before the technical ### Added / ### Changed sections, open with a short “What this means for you” block — a few benefit-framed bullets (what a user can now do, not just what changed). The technical changelog then gives the detail. This becomes the release intro and is the first thing a reader sees.
• Doc links are version-pinned. The extractor rewrites `docs/foo.md` references into [docs/foo.md](https://gpu-eda.github.io/Jacquard/<tag>/foo.html) — the mdBook page frozen for this release. So keep CHANGELOG doc references as repo-relative `docs/foo.md` (clean in-repo); the workflow does the URL rewrite.
• Versioned docs deploy. The docs-version job publishes the book to gh-pages /(tag)/ on every release tag (keep_files: true), while the main push keeps deploying “latest” to the site root. Version subdirs accumulate side by side; the pinned links above resolve to them.
• The version picker (theme/version-picker.js) reads versions.json at the site root and lets a reader move between main’s docs and the frozen releases. main is the default and stays at the root, so anyone arriving without a version in the URL reads main HEAD; a pinned build is flagged in the control, since someone who followed a link out of a release note has no other cue that the page is frozen. versions.json is regenerated after each deploy by docs/scripts/refresh_doc_versions.sh, derived from the directories actually published rather than accumulated — so pruning an old version directory also removes it from the picker, and the control can never offer a 404. Release candidates are deliberately excluded: their docs stay published (RC notes link into them) but they would crowd out the releases people want.

Staging validation (release candidates)

Optional but recommended before a user-facing release: prove the documented install commands work against a staging artifact before promoting. There is no “test crates registry”, so a GitHub prerelease is the staging tier for the binary channels.

Sequence matters — do NOT roll main to the final version before the RC validates. During the RC window main is the candidate: it stays at X.Y.Z-rc.N with the changelog still under [Unreleased]. Only the Promote step (below) rolls [Unreleased] → [X.Y.Z] and bumps to the final X.Y.Z. Rolling the release onto main first leaves main advertising a version that has no release — the version pin and dated changelog claim X.Y.Z is shipped while only the RC tag exists, which also breaks cargo binstall --git (it reads main’s version and fetches a non-existent vX.Y.Z tarball). Each new RC is a main commit bumped to the next -rc.N; main == the latest RC tag throughout.

1. Cut an RC. On main at X.Y.Z-rc.N (changelog under [Unreleased]), python3 scripts/bump_version.py <X.Y.Z>-rc.<N>, commit, tag v<X.Y.Z>-rc.<N>, push the tag. release.yml detects the SemVer pre-release suffix and publishes a GitHub prerelease (never shown as “Latest”) with the Metal tarball attached, and bump-tap pushes the formula to gpu-eda/homebrew-tap-prerelease.

2. Validate. Dispatch the Validate install (staging) workflow (validate-install.yml) with that tag. It runs the real install commands against the prerelease asset on macOS:

   • cargo binstall (asset-fetch via the [package.metadata.binstall] override, compile fallback disabled so a missing asset fails hard);
   • brew install of an RC formula (the source-of-truth formula repointed at the prerelease tarball + its .sha256, installed from a throwaway local tap).

3. Promote. A green run means the channels work. Bump to the final <X.Y.Z> (drop the -rc.<N>), commit, tag v<X.Y.Z>, push — the same flow as a normal release below.

The netlist-graph (PyPI) channel validates separately via its own workflow_dispatch → TestPyPI path (see publish-netlist-graph.yml); it versions independently of the Rust crates.

What does NOT need to change at release time

• Submodule pins (unless deliberately bumping a vendored dep).
• The vendor/opensta/ submodule pin is the version named in crates/opensta-to-ir::MIN_TESTED_OPENSTA_VERSION. If you bump the submodule, also bump the constant and the version-probe test — see WS-RH.1 in docs/plans/post-phase-0-roadmap.md.
• LICENSE (unless re-licensing).

Pre-release checklist (one-time, before the first numbered release)

These items are tracked in docs/plans/post-phase-0-roadmap.md § Release hardening; this section is the visible punch-list:

• Phase 1 (Decision 0008 required outputs) closed.
• WS-RH.1 (OpenSTA detection + version check) shipped.
• Metal CI on macos-runner-1 green (re-enabled in commit 12e98df, 2026-05-12).
• CUDA CI on nvidia-runner-1 green on main. Currently disabled in .github/workflows/ci.yml (if: ${{ false }}, ~line 268). Re-enable when hardware lands.
• HIP CI on the AMD runner green on main. Currently disabled in .github/workflows/ci.yml (if: ${{ false }}, ~line 357). Re-enable when the AMD runner is online.
• Prebuilt CUDA/HIP binaries (Decision 0018 Phase 4), when produced, must build with JACQUARD_CUDA_ARCH=all-major so the kernel ships portable SASS for every major arch (sm_50…sm_120 on CUDA ≥ 12.8, Blackwell included) plus PTX for the newest — see the README § CUDA target architecture. Local dev uses JACQUARD_CUDA_ARCH=native instead. The nvidia1.local Blackwell box (sm_120, CUDA 12.8) is a candidate self-hosted CUDA release runner.
• Vendored-dep license posture confirmed (gzz2000/eda-infra-rs#2 — sverilogparse AGPL declaration acknowledged as a typo; workspace Apache-2.0 governs).
• Cargo.toml::license = "Apache-2.0" set.
• NOTICE file enumerating vendored deps + their licenses.
• Bump vendor/eda-infra-rs submodule once upstream pushes the sverilogparse Cargo.toml correction; remove the corresponding footnote in NOTICE. Maintainer acknowledged the typo on 2026-05-02 but hasn’t pushed the fix as of 2026-05-13. Verify with git -C vendor/eda-infra-rs fetch && git log origin/master --oneline.
• CUDA / HIP runtime violation routing through process_events — done (commit 24723b5, issue #104). sim_cuda / sim_hip now dispatch the timed-batched path and drain violation events, so --timing-report / --timing-summary / --timed are no longer Metal-only.
• Bounded violations array (--timing-report-max-violations, default 100k).
• End-to-end --timing-report test on Metal CI. The inv_chain_pnr sim step uses --timing-ir (pre-generated .jtir checked in) + --timing-report + --timing-summary; a follow-up step validates the JSON shape (top-level keys, semver schema version, metadata, stats, arrays).
• GF180MCU enablement (Phases 0–6) shipped. See docs/plans/gf180mcu-enablement.md. Phase 7 (wafer.space test-run-1 design integration) deferred pending design availability; not release-blocking.

License posture

Project license is Apache-2.0 (LICENSE). Vendored-dep posture is enumerated in NOTICE. Summary:

• vendor/eda-infra-rs/ — Apache-2.0 (workspace). The sverilogparse sub-crate’s stale AGPL-3.0-only declaration in Cargo.toml is a typo per upstream maintainer (gzz2000/eda-infra-rs#2); governed by the workspace LICENSE. Submodule pin will be bumped when upstream pushes the correction.
• vendor/sky130_fd_sc_hd/ — Apache-2.0.
• vendor/opensta/ — GPL-3 (subprocess only per Decision 0001 + Decision 0006 § Amendment; never linked, never bundled).

Cross-references

• CHANGELOG.md — release log.
• docs/architecture/decisions/0008-structured-timing-output.md — --timing-report stability contract.
• docs/architecture/decisions/0002-timing-ir.md — IR schema versioning.
• docs/architecture/decisions/0006-sdf-preprocessing-model.md — OpenSTA bundling rules.
• docs/project-scope.md — license posture contract.

Handoff discipline

Handoffs in this project are ephemeral working memory, not historical record. They exist to bridge a single session boundary — when you stop working and someone else (Claude or human) picks up — and they are deleted once the work they describe is resolved.

This document defines what a handoff is, what it isn’t, when to write one, and exactly what to do when one is resolved.

Why this discipline exists

Decision rationale, technical context, and project state all have natural homes:

• Decisions (docs/architecture/decisions/) capture architectural choices and their why.
• Design docs (docs/timing-model-extensions.md, etc.) capture how things work.
• Plan docs (docs/plans/archive/phase-0-ir-and-oracle.md, post-phase-0-roadmap.md) capture what’s left and the next workstream slices.

When that content lives in a handoff instead, two things go wrong:

1. It’s not where contributors look. A new contributor reading the README → SUMMARY → decision chain shouldn’t have to dig through a stack of resolved handoff docs to find load-bearing decisions or the current state of a workstream.
2. It rots out of sync with reality. Handoffs are point-in-time snapshots. A “STATUS: RESOLVED” banner doesn’t help when the thing referenced has moved or changed; the canonical doc is what should hold the current truth.

The discipline closes this gap by forcing migration before deletion. Every load-bearing piece of a handoff lands in its proper home (decision / design doc / plan doc) before the handoff file is removed.

What a handoff IS

A handoff lives in its own dedicated directory, separate from the persistent plan docs whose content it eventually feeds: a single markdown file at docs/handoffs/<topic>-handoff.md containing exactly what the next session needs to pick up where you left off:

• Goal & next-up — what this session was trying to do, and what the very next concrete action is.
• Done this session — commits landed, with one-line summaries.
• Open follow-ups — the work that wasn’t done, with enough scope detail to start cold.
• Critical context — gotchas, surprising findings, environment specifics that aren’t obvious from the code or docs yet.
• Verification — the command(s) the next session runs to confirm the work is in the state you say it is.

One handoff per active thread of work — not one globally. A thread is a distinct arc someone could pick up cold (the RTL on-ramp, the cell-model-IR migration, a triage sweep). Concurrent threads get concurrent handoffs, each named by topic (<topic>-handoff.md). There’s still no chain: a thread holds at most one live handoff, and a resolved one is folded-and-deleted (see below), not archived behind its successor.

Treat the handoff count as a WIP signal. If the number of live handoffs climbs past roughly 2× the number of people actively working, that’s a smell, not a badge — more parallel threads than the team can hold context on. Read it as a prompt to resolve, consolidate, or drop threads, not to open more. (Earlier revisions of this doc mandated “exactly one handoff at a time”; that didn’t match engineering reality, where several independent arcs are legitimately in flight at once. The per-thread rule plus the WIP heuristic replaces it.)

What a handoff IS NOT

• Not a decision log. Decisions go in decision records. If you find yourself writing “we chose X over Y because Z” in a handoff, that paragraph belongs in a decision record (or an existing record’s “Consequences” / “Walk-back” section).
• Not a design doc. “How clock arrival flows from OpenSTA Tcl through the IR into the GPU constraint buffer” is a design topic; it lives in docs/timing-model-extensions.md Part B, not in a handoff’s “Critical context” section.
• Not a status dashboard for the project. Workstream status lives in plan docs — phase-0-ir-and-oracle.md for current-phase WS state, post-phase-0-roadmap.md for forward-looking sequencing. A handoff cites those, doesn’t reproduce them.
• Not a historical record. git log is the historical record. Handoffs that survive past their resolution turn into noise that misleads new contributors.

When to write one

Write a handoff at the end of any session that:

1. Leaves work in a partial state that someone else might pick up cold.
2. Captures non-obvious context the next session needs (e.g. “the OpenSTA Tcl find_timing proc rejects -full_update; use ::sta::find_timing_cmd 1 directly”).
3. Documents the next concrete step with enough scope to start without re-discovering it.

If the session ended at a clean stopping point (everything merged, all decisions documented in decision records or plans, nothing surprising), don’t write a handoff. The plan doc already says what’s next.

Resolution: fold, then delete

The two-location split is deliberate: handoffs live at docs/handoffs/<topic>-handoff.md while in flight; their content migrates into the persistent docs (docs/architecture/decisions/, docs/plans/, design docs under docs/) at resolution. The handoff file then gets removed; nothing about the work is lost because everything load-bearing has a permanent home elsewhere.

When a handoff’s work is done — whether in the next session or several sessions later — every load-bearing piece of it must be migrated to its proper home before the handoff file is deleted:

After migration, the handoff file is removed in the same commit as the migration:

git rm docs/handoffs/<topic>-handoff.md
git add <files-receiving-the-migrated-content>
git commit -m "$(cat <<'EOF'
docs: resolve <topic> handoff — fold into <where-it-went>

<one-paragraph summary of what was migrated and where>

Co-developed-by: Claude Code v<version> (<model-id>)
EOF
)"

The commit message records what migrated where — that’s the audit trail. git log -- docs/handoffs/ then shows the project’s handoff history (one add, one delete per session) without needing the files themselves to live forever.

Template

When you do need to write one, use this skeleton. Replace placeholders inline; delete sections that don’t apply (better to omit a section than fill it with “N/A”).

# Handoff — <Topic> (one-line summary of what this session left open)

**Created:** YYYY-MM-DD
**Working tree:** clean | <state if not clean>
**Branch:** main | <branch>

## Goal & next-up

**Goal of this session:** <what you were trying to do, in 1–3 sentences>

**Next session should pick up:** <the very next concrete action, by name. Reference the plan doc section if applicable.>

**Verification command:**
```sh
<commands the next session runs to confirm this handoff's claimed state>
# Expect: <what success looks like>

Done this session

Open follow-ups (priority-ordered)

1. ()

<Concrete scope. Enough detail to start cold. Link to existing plan/decision/design-doc sections rather than reproducing them.>

2. …

Critical context

<Things the next session needs to know that aren’t yet in the code/docs. Be honest about what’s truly load-bearing — anything obvious from git log or a quick grep doesn’t belong here.>

References

• <predecessor-handoff if any> — predecessor (if relevant)
• <plan doc> — current workstream state
• <decision> — relevant decision

Resume in a new session with: ``` /resume_handoff docs/handoffs/-handoff.md ```

## Tooling

The `create_handoff` and `resume_handoff` skills (from various Claude Code orchestration toolkits) generate and consume handoffs. They're optional — the discipline above is the load-bearing artifact. A handoff written by hand following this template is just as valid.

If you use one of those skills, expect it to default to YAML format under `thoughts/shared/handoffs/` with database indexing. **That doesn't apply to this project.** Override it: produce markdown at `docs/handoffs/<topic>-handoff.md` and skip the database step. The skill activation is informational; the project's convention takes precedence.

Architecture

The current design of Jacquard, by area. Each area is a reference doc: what the code does today, present tense, no rationale and no dead history. The why behind each choice lives in the decision records, linked from the section it explains. The not-yet-built work lives in each doc’s Implementation-status section and in plans.

If you’re an agent or a new reader, start here: this page maps every area to its current-state doc and the decisions behind it.

Decision 0019 (cell-model IR) is cross-cutting: it supplies both a cell’s logic (PDK enablement) and its timing characterization (timing correctness), so both rows list it. See the redesign spike for the design.

Timing correctness

Reference — the current design as it stands in the code. The why lives in the decision records linked from each section; the not-yet-built work is under Implementation status. Present tense means “true today”; if a claim here is stale, the code moved and this doc didn’t.

Jacquard’s timing story has two halves. Getting a number is the simulation engine’s job: per-gate delays ride the boomerang’s arrival propagation and come out as setup/hold reports. Trusting that number is this doc’s job — where the delays come from, how they’re validated, and what they mean. Static timing analysis (STA) is a solved problem; Jacquard doesn’t re-solve it. It treats an existing STA tool as ground truth and builds a narrow, auditable pipe from that tool’s output into the GPU kernel. The accuracy tolerances and validation requirements this pipe must meet are the Timing Correctness Requirements contract.

OpenSTA as the oracle and the only STA path

OpenSTA is both jobs at once: the tool Jacquard’s own numbers are checked against, and the only static timing analyzer Jacquard talks to at all. An earlier plan paired OpenSTA with an in-process reference engine (OpenTimer), so a divergence between Jacquard and OpenSTA could be triangulated against a third source; a spike found OpenTimer’s input pipeline unfit for OpenROAD-flow output, and that plan is superseded. Where Jacquard’s results disagree with OpenSTA’s past a declared tolerance, Jacquard is wrong until proven otherwise — the divergence gets fixed, justified in writing, or filed as a bug. Decision 0001.

OpenSTA is GPL-3.0; Jacquard is permissively licensed, so nothing links it. Every call to OpenSTA crosses a subprocess boundary, driven by opensta-to-ir, a standalone preprocessing tool that also owns the Liberty/Verilog/SDF/SPEF/SDC parsing on OpenSTA’s side. opensta-to-ir may be invoked ahead of time to produce a .jtir file, or subprocessed live by jacquard sim --sdf <path> — both are permitted in the shipped runtime, as long as OpenSTA itself is discovered on PATH rather than bundled into a Jacquard release artefact. Decision 0006.

OpenSTA is vendored as a git submodule at vendor/opensta/, pinned to a specific revision for CI reproducibility, but never built as part of Jacquard — subprocess calls use whatever OpenSTA binary the environment provides. The submodule’s real payload is its test corpus. A small, curated primary regression corpus (tests/timing_ir/corpus/) runs on every CI job; a much larger stress corpus, referenced by manifest into the submodule rather than duplicated, runs nightly to catch parser crashes and malformed IR on dialects Jacquard doesn’t otherwise see. As shipped, the primary corpus holds one entry (aigpdk_dff_chain); SKY130/NVDLA/MCU-SoC entries are pending on a CI SKY130-Liberty install strategy, and the stress corpus’s manifest is still empty. Decision 0005.

The timing IR

Everything OpenSTA computes for a design reaches Jacquard through one channel: the timing IR, a schema-versioned FlatBuffers format with a JSON sidecar for CI diffs and human inspection. opensta-to-ir is the producer; flatten.rs::load_timing_from_ir is the consumer, turning per-cell arcs into the PackedDelay (rise/fall picoseconds) and DFFConstraint (setup/hold) structures the simulation engine bakes into its script. The IR is multi-corner natively — values are min/typ/max across a declared PVT corner set, selectable at sim time via --timing-corner — and every arc carries provenance (asserted from the input file, computed by OpenSTA, or defaulted). Unrecognised vendor annotations pass through as typed VendorExtension blobs rather than being silently dropped. The IR’s scope is deliberately narrow: per-design timing annotation only, never a netlist representation, a timing graph, or cell characterization. Decision 0002.

Per-cell-type characterization — setup/hold, clock-to-Q, the numbers that don’t vary by instance — is a different axis and, today, comes from a separate runtime path: liberty_parser::TimingLibrary, parsed from a .lib file at sim startup. Decision 0019 proposes folding that parse, plus a cell’s logic and sequential classification, into one generated cell-model IR per library, with the simulation corner picked explicitly via --corner rather than read from the netlist or SDF. That descriptor doesn’t exist yet; it’s Proposed, not built — see Implementation status. Decision 0019.

From SDF to a running gate

The hand-rolled src/sdf_parser.rs is gone. Every path that once consumed SDF directly now consumes timing IR instead, which is what makes OpenSTA a required runtime dependency for any timing-aware flow, not just a CI concern: jacquard sim --sdf <path> and the standalone opensta-to-ir tool both subprocess OpenSTA to produce the IR that flatten.rs then loads.

OpenSTA’s own read_verilog accepts only structural Verilog — cell instantiations and bare-net assigns, no RTL operators or bit-selects. Some flows wrap a clean post-P&R netlist in a thin RTL shim (an active-low OEB patch, for instance), which OpenSTA’s reader rejects outright. opensta-to-ir filters each --verilog input at invocation time, extracting just the module <top> … endmodule block before handing it to OpenSTA, so wrapper modules are simply never seen. What opensta-to-ir cannot check is whether the Verilog handed to it is the right stage for a given SDF — a pre-P&R synthesis netlist looks structurally identical to the post-P&R body but is missing the hundreds of thousands of P&R-inserted cells the SDF actually references, and OpenSTA quietly drops SDF entries whose endpoints aren’t in the loaded design. Picking the right netlist stays the caller’s job. Decision 0009.

Once the IR is loaded, flatten.rs overlays it on the AIG: each cell’s delay lands on the AIG pin it drives, and inject_timing_to_script bakes that value into the flattened script’s per-gate padding slot as picoseconds the kernel reads directly during boomerang reduction. DFFConstraint values become the per-word setup/hold check the kernel runs at each cycle boundary. See ../timing-simulation.md for the full mechanics of that propagation, including the three independent sources of conservative overestimate in today’s packed-32 model (max-of-rise/fall per cell, max wire delay across a cell’s input pins, max arrival across the 32 signals sharing a thread) — the model never under-reports a violation, which is what setup checking needs, at the cost of some false positives.

Structured output

Getting a timing number out of a run used to mean grepping stderr for a state-word index. Symbolic violation messages now name the hierarchical signal directly (top/cpu/regs[7][bit 22]); --timing-report <path.json> emits a schema-versioned, end-of-run document with per-DFF worst arrival and slack, a per-cycle violation list, and top-N worst-slack rankings even where no violation occurred; --timing-summary is the fast human-readable counterpart for scripts and quick inspection. All three ship for sim on Metal, CUDA, and HIP. Decision 0008.

Constraints

• OpenSTA is the only STA path. There is no in-process reference engine; a run that needs timing data needs an OpenSTA-produced (or pre-converted) IR. Decision 0001.
• The Verilog fed to opensta-to-ir must be the post-P&R structural body, matching what the SDF was generated against. opensta-to-ir strips RTL wrapper modules automatically; it cannot detect a wrong-design-stage substitution, since a synthesis netlist and a P&R netlist can share a module name while differing by orders of magnitude in cell count. Decision 0009.
• The timing model is conservative, not exact. Rise/fall collapse, max wire delay per cell, and max arrival per packed thread all push arrival times up, never down. See ../timing-simulation.md.
• Timed (arrival-annotated) cosim runs on Metal only, and the structured --timing-report JSON is sim-only; the cosim runtime hasn’t wired either through yet.
• Synchronous, edge-triggered logic only, inherited from the simulation engine — there’s no timing model for a level-sensitive latch or asynchronous sequential feedback to annotate in the first place.
• No linking of GPL code, ever, and no bundling of OpenSTA in a Jacquard release artefact. Subprocess invocation of a user-installed OpenSTA is the only sanctioned integration path. Decision 0006.

Implementation status

Built and in use: OpenSTA as the sole out-of-process STA oracle; the FlatBuffers + JSON timing IR with multi-corner support and provenance tagging; the opensta-to-ir subprocess path (both pre-converted --timing-ir and live sim --sdf); the structural-Verilog wrapper filter; per-gate delay and per-DFF setup/hold injection into the flattened script; GPU-side setup/hold violation detection on Metal, CUDA, and HIP for sim; and the structured output surface (symbolic violations, --timing-report, --timing-summary, worst-slack ranking) for those same three backends. Pillar B (per-DFF clock-tree skew via a ClockArrival IR table, folded into the setup/hold check) is also landed, ahead of its own roadmap’s schedule.

Decided but not yet built:

• Native Rust SDF→IR converter (Phase 3), which would remove the OpenSTA subprocess from the shipped runtime entirely. No longer release-gating; deferred indefinitely until bandwidth or commercial demand reopens it. Decision 0006.
• cosim’s structured timing report and timed simulation on CPU/CUDA/HIP. Today only Metal cosim carries arrival annotations, and none of the --timing-report/--timing-summary surface reaches cosim. Decision 0008.
• The cell-model IR (Decision 0019) — Proposed, not started. Per-cell timing characterization stays on the runtime liberty_parser::TimingLibrary path until this lands; --corner-based cell-library selection doesn’t exist yet. Decision 0019.
• The timing-model fidelity roadmap (Decision 0007) — Proposed. Of its three pillars, only Pillar B (clock-tree skew, above) has landed. Still open: Pillar C Tier 1 (per-receiver wire delay, gated on Decision 0007 acceptance), Pillar A (dynamic per-gate delay, δ(T), gated on per-cell SPICE characterisation and explicitly scheduled after Pillars B and C), and Pillar C Tiers 2–3 (inter-partition wire delay, NoC-aware partitioning hints). Scheduling detail in ../plans/post-phase-0-roadmap.md.
• Decision 0008’s optional outputs — arrival histograms, an OpenSTA-critical-path cross-reference, and DFF path-back-trace. Scheduled per user demand once the required items above see use.

Decisions behind this

• Decision 0001 — OpenSTA as the timing-correctness oracle and the sole STA path; the in-process alternative (OpenTimer) was superseded by a spike.
• Decision 0002 — the timing IR: FlatBuffers
  • JSON, multi-corner, provenance-tagged, per-design annotation only.
• Decision 0005 — OpenSTA vendored as a pinned submodule; primary vs. stress test corpus split.
• Decision 0006 — the SDF preprocessing model: deleting the hand-rolled parser, opensta-to-ir as the subprocess bridge, and the (now relaxed) rules on runtime subprocess invocation.
• Decision 0008 — structured timing output as a first-class deliverable: symbolic violations, --timing-report, --timing-summary, worst-slack ranking.
• Decision 0009 — OpenSTA’s structural-only Verilog reader, and where the wrapper-stripping and design-stage responsibilities sit.
• Decision 0019 — Proposed: a generated cell-model IR unifying per-cell-type logic, classification, and L4 timing characterization, replacing the runtime Liberty parse.
• Decision 0007 — Proposed: the three-pillar roadmap (dynamic delay, clock-tree skew, wire delay at scale) for closing the accuracy gap with an event-driven reference simulator.

Simulation engine

Reference — the current design as it stands in the code. The why lives in the decision records linked from each section; the not-yet-built work is under Implementation status. Present tense means “true today”; if a claim here is stale, the code moved and this doc didn’t.

The simulation engine compiles a gate-level netlist into a packed GPU script and evaluates it as a branch-free Boolean processor. A design is converted once, at the start of a run, into an and-inverter graph; the graph is partitioned to fit a fixed GPU block geometry, and each partition becomes a stream of u32 instructions the kernel executes with no per-gate dispatch. This is the path behind jacquard sim (static-input replay) and the substrate the cosim runtime drives reactively.

The pipeline

Compilation is a fixed sequence of lowering stages, each narrowing the design toward the GPU’s fixed shape:

• NetlistDB parses structural (gate-level) Verilog (sverilogparse) into a flattened database of cells, pins, and nets with CSR connectivity. Behavioural RTL is synthesised to this form first; see the RTL on-ramp.
• AIG rewrites all combinational logic to one uniform AND-with-invert node type, so the kernel needs no opcode decode (Decision 0014).
• StagedAIG splits combinational cones too deep for one boomerang tree into sequential major stages.
• Partitions distribute the AIG’s endpoint groups across GPU blocks under fixed resource limits, via hypergraph partitioning.
• FlattenedScriptV1 packs each partition into the u32 instruction stream the kernel reads directly (Decision 0015).

Partitioning and script generation run automatically at simulation start; the user only sets NUM_BLOCKS (2× the GPU’s SM/CU count; 1 for Metal).

The AIG

All combinational logic is one and-inverter graph. Every node in a combinational cone is a DriverType:

#![allow(unused)]
fn main() {
pub enum DriverType {
    AndGate(usize, usize),     // two operands, each carrying an inversion bit
    InputPort(usize),          // primary input
    InputClockFlag(usize, u8), // clock edge (posedge / negedge) detector
    DFF(usize),                // D flip-flop output
    SRAM(usize),               // memory block output
    Tie0,                      // constant zero
}
}

Only AndGate has combinational fan-in. Each operand is encoded aigpin << 1 | invert, so one node type covers the whole {AND, NAND, OR, NOR} family and inverters/buffers fold into the inversion bit rather than becoming nodes. This uniformity is the load-bearing property: because every combinational node is the same (a XOR xa) AND (b XOR xb) operation, the boomerang tree evaluates them all with a single instruction pattern. Rationale, and why AIG over BDDs / LUTs / MIG / direct netlist execution: Decision 0014.

Construction (src/aig.rs, AIG::from_netlistdb) is technology-independent and descriptor-driven. Every cell type is looked up in a cell-model-IR descriptor that carries its pre-decomposed AIG, and that AIG is spliced straight in (Decision 0019). The built-in PDKs (AIGPDK, SKY130, GF180MCU, IHP SG13G2) ship generated descriptors, and cells outside them supply the same kind of declarative metadata (Decision 0010); both take the one descriptor path. How a PDK is onboarded: PDK enablement. A legacy per-PDK behavioural-model decomposition (decompose_with_pdk) survives only as a fallback for cell types no descriptor covers. A structural cache deduplicates identical sub-expressions, and AIG pins come out in topological order, which the downstream level computation and scheduling rely on.

The AIG’s outputs are grouped into endpoint groups — the units of work a partition must realise: a primary output, a DFF (data + clock-enable), a RAM block (address/data/enables), a $stop/$finish control node, a $display node, or an inter-stage boundary pin from staging. Each group bundles the signals that must be evaluated together; the partitioner uses the group’s input set for connectivity and the partition executor uses it for resource accounting.

Staging deep circuits

When a design’s combinational depth exceeds one 8192-wide boomerang tree, src/staging.rs splits the AIG into major stages at user-specified level thresholds (--level-split 30, or --level-split 20,40 for two cuts). Each stage carries its own primary inputs (the pins produced by earlier stages, or the design’s real inputs for the first), the live boundary pins it must forward, and the endpoint groups whose depth falls within it. Major stages run sequentially — the kernel loops over them, writing boundary values to the state buffer and re-reading them in the next stage. Staging trades extra sequential dispatches for fitting the fixed boomerang; without it, designs with cones deeper than ~50 levels fail partitioning outright. Mapping details: Decision 0015.

Partitioning

Endpoint groups are distributed across GPU blocks by RepCut (src/repcut.rs), which builds a weighted hypergraph and cuts it with mt-kahypar. A hypergraph, not a graph, because one shared AIG node reached by many endpoint groups is a single hyperedge across all of them — cutting it costs one global-memory read shared by every block that needs the signal, which pairwise edges cannot express. The cut minimises Sum-of-External-Degrees, which is exactly that cross-block read count.

mt-kahypar deliberately over-partitions (≈2× the block count); process_partitions() in src/pe.rs then greedily merges partitions back, scoring merge candidates by AIG-node-bitset overlap, trying merges speculatively in parallel with cancel-on-success, and rejecting any merge that would add boomerang stages beyond a degradation bound. The result is 2–4× fewer partitions than the raw solution, each validated to fit the boomerang’s resource limits (Constraints).

The boomerang execution model

One GPU block (CUDA/HIP) or threadgroup (Metal) evaluates one partition by reducing its endpoint groups’ fan-in cones through a hierarchical binary tree — the boomerang. The tree has BOOMERANG_NUM_STAGES = 13 levels (2^13 = 8192 leaf positions); each thread owns 32 bits, so a block is 8192 / 32 = 256 threads (NUM_THREADS_V1). The 13 levels map to three GPU mechanisms:

Every position computes (a XOR xora) AND (b XOR xorb) OR orb — the AIG’s AND-with-invert, with orb all-ones turning a position into a pass-through. One instruction pattern, zero branch divergence, maximal SIMT utilisation across all backends. When a partition needs more than one 8192-wide tree, a shuffle permutation (16-bit index pairs in the script) re-routes signals from shared memory back into thread registers between stages. Full derivation: Decision 0015.

The flattened script

src/flatten.rs packs each partition into FlattenedScriptV1, a u32 stream read sequentially by the kernel, with four sections: a metadata block (256 u32 of per-partition control fields at fixed indices, plus the write-out hook table that maps each thread to the boomerang stage+position where it captures its output); a global-read permutation that gathers each thread’s input bits from the state buffer; the boomerang sections (per stage, the shuffle permutation plus the xora/xorb/orb AND-gate flags, with a padding slot reused to carry per-gate gate-delay picoseconds for timing runs); and the global write-out that commits results, SRAM ports, and output duplicates back to the state buffer. The metadata index layout is the load-bearing contract between flatten.rs and every kernel.

The global-read permutation carries one subtle cross-backend detail. A read’s word index has bit 31 set when the value is an inter-stage intermediate rather than previous-cycle state. Every decoder must strip that flag from the index before indexing: flatten.rs, the CPU reference, the shared CUDA/HIP kernel_v1_impl.cuh, and kernel_v1.metal. Cancelling the flag by offsetting the base pointer instead forms an out-of-bounds address. That bug placed staged reads 2^32 words past the buffer on ROCm, the one backend that computes the full address rather than truncating it to 32 bits (#203).

Backends and the reference model

The same FlattenedScriptV1 runs on four backends behind one seam. The GPU kernels each evaluate one partition per block: csrc/kernel_v1.cu and csrc/kernel_v1.hip.cpp share kernel_v1_impl.cuh, and csrc/kernel_v1.metal is the Metal port. sim takes a grid-wide barrier between stages, which needs cooperative launch; the cosim runtime sidesteps that so its CUDA/HIP backends run without it. A CPU reference kernel (--check-with-cpu) evaluates the same script bit-for-bit and is the cross-backend equivalence oracle. Every backend is expected to produce identical results on the same design, and that equivalence is checked in CI.

Selective X-propagation

By default the engine is two-state: uninitialised DFF and SRAM outputs resolve to 0, which hides init bugs and causes false mismatches against four-state RTL simulators. --xprop turns on selective four-state simulation. Static analysis identifies X-source signals (uninitialised DFFs, SRAM reads), and forward-cone analysis marks each partition X-capable or X-free. Only X-capable partitions run the X-aware kernel variant and pay its ~2× storage and ALU cost, held in X-mask words appended to the state buffer; the rest keep the fast two-state path. In a typical SoC after reset, under ~5% of partitions are X-capable. Output VCD then carries x values, and --check-with-cpu has an X-aware reference. Design choices and the seven-phase implementation are in Decision 0016 and docs/selective-x-propagation.md. The reactive extension belongs to the cosim runtime: it adds undriven input pads as a third X-source and maintains the X-mask per edge.

Assertions and display

assert() and $display/$write survive synthesis as GEM_ASSERT / GEM_DISPLAY cells (via the gem_formal.v techmap over Yosys $check/$print cells). The AIG records their bit positions in the script; after each GPU step the CPU reads those positions and acts on them. An assertion fires a configurable action (log, pause, or terminate, bounded by max_failures). A display reconstructs its message from format strings held in JSON metadata and argument bits read from the state buffer.

Constraints

• Synchronous, edge-triggered logic only. Sequential state is D flip-flops capturing on clock edges. What is not modelled is a raw level-sensitive latch left in the logic, and asynchronous sequential (self-timed) feedback. Three common latch-derived structures are supported: asynchronous set/reset on flip-flops (it lowers to an AIG overlay), clock gating via the CKLNQD integrated clock-gating cell, and latch-based register files mapped to $__RAMGEM_SYNC_ SRAM cells by memory synthesis. See the latch note in Decision 0014.
• 8191 unique inputs and 8191 unique outputs per partition — the 8192 boomerang leaf/write-out slots minus Tie0. Wide buses or highly-connected cones force finer partitioning, which raises inter-block reads.
• 4095 intermediate pins alive per stage, and 64 SRAM output groups per partition (each SRAM consumes 4 of the 256 write-out slots). SRAM-heavy designs may need finer partitioning than gate count alone implies.
• Fixed boomerang geometry. The 8192-leaf, 256-thread, 13-stage tree is one hardcoded design, uniform across CUDA, HIP, and Metal rather than tuned per GPU architecture. It assumes 32-bit lanes and a 32-wide SIMD reduction at the shuffle tier (levels 4–7). There is no occupancy-tuning knob; retargeting the geometry to a different architecture would mean redesigning the hierarchy depth and the script bit-packing.

When a single endpoint cannot be mapped, --level-split forces stage splits.

Implementation status

Built and in use: the full NetlistDB → AIG → StagedAIG → Partitions → FlattenedScriptV1 → kernel pipeline; all four backends (CUDA, HIP, Metal, CPU reference) with cross-backend equivalence in CI; RepCut hypergraph partitioning with greedy merge-back; selective X-propagation (all seven phases, including the CPU reference and Criterion overhead benchmarks); and assertion/display support.

Decided but not yet built:

• Dynamic X narrowing — periodically re-scanning the X-mask and hot-swapping a partition from the X-aware kernel back to the fast path once its cone is X-free. Distinct from the shipped static X analysis. Decision 0016.
• Per-address SRAM X — reads currently return X until the first write to the memory; finer per-address tracking is deferred. Decision 0016.

Decisions behind this

• Decision 0014 — AIG as the simulation IR: the uniform AND-invert node, the inversion-bit encoding, the decomposition path, and the EndpointGroup abstraction.
• Decision 0015 — the boomerang execution model: the 13-level reduction tree, the GPU resource limits, RepCut hypergraph partitioning, and the FlattenedScriptV1 format.
• Decision 0016 — selective X-propagation: partition-level four-state simulation behind --xprop.

Cosim runtime

Reference — the current design as it stands in the code. The why lives in the decision records linked from each section; the not-yet-built work is under Implementation status. Present tense means “true today”; if a claim here is stale, the code moved and this doc didn’t.

Cosim runs a GPU-simulated design alongside reactive peripheral models that drive and observe the design’s pins every clock edge. Unlike sim, which replays a fixed input VCD, cosim lets an input depend on a design output cycle by cycle: a SPI flash serves firmware in response to the addresses the design issues, a UART decodes the bytes it transmits, a JTAG client single-steps it. Rationale for the execution model: Decision 0017.

The per-edge pipeline

The unit of time is a scheduler edge, not a clock cycle (see Time: edges, not cycles). Each edge runs this sequence on the GPU:

state_prep            apply clock, GPIO, JTAG pin drives via BitOps
  → inject            bidirectional peripherals write design inputs (flash MISO)
    → simulate ×N     combinational logic evaluation, N pipeline stages
  → sample + advance  bidirectional peripherals read outputs, step their FSM
  → observe           output-only peripherals decode into ring buffers (UART, bus trace)

CPU-side peripheral models run between batches of edges, not within them, so they see the design’s output state only at batch boundaries.

Batch dispatch

Consecutive scheduler edges are grouped into batches of up to BATCH_SIZE = 1024 and dispatched as one GPU command buffer. Between batches the CPU steps its peripheral models, drains the ring buffers, and patches the next batch’s pin drives. A batch runs with no CPU interaction inside it, which is why cosim is GPU-bound rather than round-trip-bound.

The batch collapses to a single edge whenever any peripheral model reports is_active(). A JTAG replay mid-transmission, or a connected interactive debug client, holds the design at one edge per dispatch for as long as it needs per-edge coupling. This is the force_single_edge path; it is correct but slow, and it is the exception, not the norm. The why, and the measured batch-utilisation that justifies the fixed size, are in Decision 0017.

Backends

One orchestration driver runs above a CosimBackend trait; each backend owns the design state and runs N consecutive edges per call.

• CpuBackend — the per-edge reference model. It is the cross-backend equivalence oracle and the reason cosim regressions run on free Linux CI. Throughput is not its job.
• MetalBackend — runs a batch in one command buffer with GPU peripherals inside it. Unified memory makes the pin drives zero-copy: the CPU’s write to an edge’s ops is the upload.
• CudaBackend / HipBackend — run the same simulate kernel over managed memory, sidestepping the cooperative-launch grid sync that only sim needs. They ship with their GPU peripherals so reactive designs batch from the start; ops upload only for edges the CPU actually changed.

Every backend produces a byte-identical output VCD on the same reactive design. That equivalence is the correctness contract for the whole runtime, checked in CI against CpuBackend goldens.

Peripherals

A peripheral is one shape on either substrate: observe some design-output bits, advance an FSM over persistent state, then drive some design-input bits and/or emit decoded records. Where it runs follows a rule: a model that drives inputs each edge (must react to output) can run on the CPU; a model that only observes outputs, or exchanges data bidirectionally, runs on the GPU for zero-copy access to the state buffer. Rationale and the full contract: Decision 0013.

GPU peripherals come in two patterns:

• Observe-only — one post-simulate kernel per edge, reading output state into a ring buffer. UART TX decode and bus trace.
• Bidirectional — a pre-simulate kernel injecting into input state plus a post-simulate kernel sampling and advancing the FSM. SPI flash and QSPI PSRAM.

Three tiers back each peripheral. Tier 1 is the CPU PeripheralModel (always present, the semantic ground truth). Tier 2 is a hand-written GPU kernel for the core set — flash, UART, bus trace — one shared *_impl.cuh for CUDA and HIP plus one .metal. Tier 3, single-source peripheral compilation for user-defined peripherals, is not yet built.

Config is plural where instances can repeat, following one back-compat pattern (a legacy singular key folds into instance 0 of a new Vec):

GPU→CPU ring buffers

GPU peripherals write into fixed-size ring buffers in device memory; the CPU drains each after a batch completes, from a local read_head up to the GPU-written write_head. No synchronisation beyond command-buffer completion is needed. This is the only channel out of the batch today, and it is the GPU→CPU half of the flow-controlled I/O work. When VCD output is on, a separate per-edge snapshot ring lets a batch stay larger than one edge without a CPU read after every edge.

Time: edges, not cycles

The MultiClockScheduler computes a deterministic interleaving of edges across clock domains. The scheduler tick is the GCD of all half-periods and phase offsets; the schedule repeats at the LCM of all full periods. A scheduler edge is one tick; a clock cycle is two half-periods of one domain. UART baud dividers, reset duration, and the --max-clock-edges flag all count edges, not cycles — conflating the two is a real bug class, not a pedantic distinction. The schedule length is capped at one million ticks, which rejects non-commensurable clock ratios rather than building an unbounded schedule.

Constraints

• Synchronous logic only, inherited from the simulation engine.
• A CPU-side model cannot observe intra-batch state; anything needing per-edge visibility must set is_active() and pay the single-edge dispatch.
• The scheduler rejects clock sets that don’t repeat within one million ticks.
• Timed (arrival-annotated) cosim runs on Metal only; the structured --timing-report JSON is sim-only.

Implementation status

Built and in use: batch dispatch, the multi-clock scheduler, all four backends with cross-backend equivalence, the CPU and GPU peripheral tiers (1 and 2), plural configs, and the interactive JTAG server. These are the sections above without a “not yet.”

Decided but not yet built:

• Flow-controlled external I/O across the batch boundary — bounded pipes, an adaptive commit window replacing the fixed BATCH_SIZE, per-tap projection, a socket reactor, and a CPU→GPU input pipe. Proposed in Decision 0022; nothing built.
• Tier 3 single-source peripherals — one peripheral definition compiled to CPU and every GPU backend. Decision 0017.
• Bus trace beyond APB3 — AHB-Lite/AHB5 decode and annotated-VCD output; migrating the hardcoded Wishbone trace onto the config-driven monitor. Decision 0013.
• Timed cosim on CPU/CUDA/HIP, and the cosim structured timing report. Decision 0017.
• Min-heap scheduler to remove the one-million-tick cap. Multi-clock and stimulus architecture.

Scheduling for the larger items lives in cosim backend portability and the multi-clock plan.

Decisions behind this

• Decision 0013 — peripheral model architecture: the CPU/GPU split, the two GPU patterns, ring buffers, plural configs.
• Decision 0017 — execution model: batch dispatch, the multi-clock scheduler, edges versus cycles, the backend seam.
• Decision 0012 — CDC jitter injection, which uses the scheduler’s edge timestamps as its injection point.
• Decision 0022 — flow-controlled I/O (proposed; the future of the batch boundary).

RTL on-ramp

Reference — the current design as it stands in the code. The why lives in the decision records linked from each section; the not-yet-built work is under Implementation status. Present tense means “true today”; if a claim here is stale, the code moved and this doc didn’t.

jacquard sim and jacquard cosim accept behavioral RTL directly — no separate build step, no external toolchain. Pass a .v/.sv file where you’d otherwise pass a gate-level netlist, and the same command classifies the input, synthesizes it if needed, caches the result, and simulates it on the GPU. This is a front end onto the simulation engine: the emulator core is unchanged, and only ever consumes a synthesized structural netlist. Rationale for embedding a synthesis engine at all: Decision 0021.

Input dispatch

sim/cosim decide what they were handed by attempting a structural parse (sverilogparse) and matching the result against known cell families:

The command always prints what it decided, for example design.v: behavioral RTL → synthesized [YoWASP Yosys, functional QoR] → <cache>, so synthesis is never silent. --rtl and --netlist override auto-detection when it guesses wrong; --emit-synth <path> dumps the intermediate gate-level netlist for inspection or fixture authoring. Full detection table and CLI surface: the accepted RTL surface.

The embedded synthesis engine

Synthesis runs as Yosys compiled to WebAssembly, executed in-process from Rust (src/synth.rs), not a vendored native build and not a Python subprocess. YoWASP ships Yosys as a single self-contained yosys.wasm, with abc compiled in-tree into that module and called in-process — WASI has no exec, so there is no separate abc process to shell out to. src/synth.rs embeds wasmtime, loads the wasm module, preopens the design, the aigpdk support files, and a temp work directory under WASI, and runs a generated synthesis script through it. No Python interpreter and no external toolchain: jacquard sim design.v … goes from RTL to waveform out of the single jacquard binary. This is behind the opt-in synth feature (cargo build --features metal,synth) because wasmtime and cranelift are heavy to compile; a binary built without it gives an actionable error on behavioral input instead of failing to build at all.

The SystemVerilog front end is yosys-slang (read_slang), a near-complete SystemVerilog-2017 elaborator bundled in the pinned wasm. src/synth.rs probes for read_slang once per wasm module (help read_slang, cached for the process) and uses it when present; on an older wasm without it, the script falls back to Yosys’s built-in read_verilog -sv, whose accepted subset is narrower. Which frontend ran is logged, and it factors into the cache key below.

Two independent caches make repeat runs cheap. The compiled wasm module is cached under $XDG_CACHE_HOME/jacquard, keyed by content hash, so only the first run of a given yosys.wasm pays the cranelift compile — later runs deserialize the cached module. The synthesized netlist is cached separately, keyed by a hash of the design source, the generated script, the wasm bytes, and the embedded aigpdk support files (aigpdk_nomem.lib, memlib_yosys.txt, gem_formal.v), so editing any of those correctly invalidates a hit — a repeat sim/cosim run of unchanged RTL skips synthesis entirely.

The yosys.wasm module itself is fetched on first use rather than bundled into the binary. locate_yosys_wasm resolves it in order: an explicit --yosys-wasm path, the JACQUARD_YOSYS_WASM environment variable, or a pinned release fetched (sha256-verified) from the gpu-eda/yowasp-yosys fork and cached under $XDG_CACHE_HOME/jacquard/wasm/<tag>. That fork, not a stock YoWASP wheel, is the default source: it carries the patched abc &origins path the on-ramp’s provenance mapping (below) depends on, and a stock wheel would silently drop source locations instead of failing loudly.

Two synthesis tracks

• On-ramp — YoWASP Yosys, the default whenever sim/cosim are handed behavioral RTL. Easy: one command, nothing to install beyond the synth feature. Functional-grade QoR.
• Performance — bring your own synthesizer (Synopsys DC, or a native Yosys install) to produce a gate-level .gv, then point jacquard sim at it directly (the “gate-level, built-in PDK” row above). Synthesis quality is the primary tuner of simulation throughput, since the AIG the GPU emulates is only as good as the mapping. Full flow: Synthesis flow.

The two tracks share the same aigpdk libraries and memory-mapping rules (memlib_yosys.txt); the on-ramp runs the same synthesis shape the performance-path docs describe manually, just from an embedded engine instead of a standalone Yosys or DC invocation.

Accepted language surface

The accepted behavioral subset is whatever the embedded frontend can synthesize to an aigpdk netlist — Jacquard delegates elaboration to Yosys rather than implementing a Verilog front end itself. Verilog-2005 (assign, always @(*)/@(posedge/negedge), synchronous and async-reset flip-flops, case/if chains, parameterized modules and generate, inferred memories) is supported through Yosys’s built-in reader. Through read_slang, the SystemVerilog-2017 surface is substantially broader: packages, interfaces (including modports), always_ff/always_comb/always_latch, logic/enum/struct/union, and advanced generate. Full subset table and known gaps: the accepted RTL surface.

Three project-specific techmaps run on top of standard synthesis, defined in aigpdk/gem_formal.v:

• GEM_ASSERT — $assert/$assume/immediate assert property, visible in simulation as assertion failures.
• GEM_DISPLAY — $display (via Yosys $print); the cell is emitted by the on-ramp, though the display-info JSON companion the CPU side needs to decode format strings isn’t written yet on this path.
• RAMGEM ($__RAMGEM_SYNC_) — inferred synchronous memories, mapped via aigpdk/memlib_yosys.txt, the same memory-synthesis step the performance path runs by hand.

Testbench-only constructs (#delay, most initial blocks, $finish, $stop) have no synthesizable meaning and are dropped during elaboration. Latches aren’t supported: the emulator core is synchronous only, and a latch inferred from RTL is a synthesis or AIG-load error.

RTL-source provenance

The on-ramp maps to aigpdk cells via abc_new, abc9’s AIGER/XAIGER path, rather than the classic abc -liberty flow. XAIGER round-trips the design through an in-memory AIGER “y” channel that carries Yosys’s (* src *) attributes through std-cell mapping, where the classic BLIF-based flow drops them — BLIF has no object identity to hang provenance off. sverilogparse captures those \src locations into netlistdb.cell_src, and AIG::aigpin_src_locations threads them into the AIG, so xsources and --trace-signals can report an RTL file:line instead of a flattened gate name. A signal resolving to binary_to_gray.sv:15 instead of $auto$…$1234 is the visible result. Full write-up: RTL-source provenance.

Constraints

• The emulator core does not change. The input to the AIG/boomerang pipeline is always a synthesized structural netlist; behavioral elaboration stays Yosys’s job. Jacquard does not reimplement an RTL front end, so the accepted surface is bounded by what the embedded Yosys accepts, not by anything Jacquard-specific.
• A QoR ceiling on the easy path. YoWASP Yosys is functional-grade — correct, not tuned for GPU throughput. Peak performance still wants DC or a native Yosys install on the performance path.
• Concurrent-SVA synthesis is partial. Turning property/sequence-based assert property into gate-level checkers is a separate Yosys formal-flow capability, independent of read_slang’s ability to parse SVA, and it’s still incomplete (issues #106, #107). Immediate assertions synthesize correctly through the GEM_ASSERT mapping regardless.
• Released binaries must be built --features synth. The feature is opt-in at compile time because wasmtime/cranelift are heavy to build, but because behavioral input is a capability of sim/cosim themselves rather than a separate command, a released binary without it fails behavioral input with a build-without-synth message rather than silently rejecting RTL.

Implementation status

Built and in use: the on-ramp itself (behavioral input to sim/cosim, three-way auto-detection, the embedded wasmtime-hosted Yosys engine, read_slang with fallback, both caching layers, fetch-on-first-use of the pinned provenance wasm), and RTL-source provenance — abc_new/XAIGER mapping, \src capture through sverilogparse/netlistdb, and its surfacing in xsources and --trace-signals. These are the sections above without a “not yet.”

Decided but not yet built:

• An empirical sv-tests coverage table. The accepted-language-surface prose above describes the intended surface. The authoritative measure is pass/fail coverage driven through sv-tests, which hasn’t shipped. Decision 0021.
• Concurrent-SVA → synthesizable-checker synthesis (#106, #107) — a Yosys formal-flow gap, not something src/synth.rs can work around.
• Source locations in --timing-report. Provenance reaches --trace-signals and xsources/xroots today; threading \src into the structured timing report is the one deferred piece of the provenance plan. RTL-source provenance.

Decisions behind this

• Decision 0021 — behavioral RTL support: folding synthesis into sim/cosim, the embedded wasmtime/Yosys engine, the two synthesis tracks, and the RTL-source-provenance roadmap.
• Decision 0014 — the emulator model: why synthesis is a front end at all, and why Jacquard doesn’t elaborate behavioral RTL itself.
• Decision 0019 — the cell-model IR: how a built-in PDK’s descriptor supplies logic and timing together, and how --corner selects among them.

PDK enablement

Reference — the current design as it stands in the code. The why lives in the decision records linked from each section; the not-yet-built work is under Implementation status. Present tense means “true today”; if a claim here is stale, the code moved and this doc didn’t.

A PDK gives Jacquard two things about every cell type it contains: the combinational or sequential logic the simulation engine splices into the design’s AIG, and the timing characterization a --timing-report run reads. Both come from one source, a per-library cell-model IR descriptor (liberty-to-cellir converter, Decision 0019), and both the built-in PDKs and a user’s own foundry library go through the same descriptor mechanism. Cells a descriptor doesn’t cover, such as third-party IP, memory macros, and filler and pad cells, get the same kind of declarative treatment through a second, smaller path: a hand-authored TOML manifest (Decision 0010). One mechanism for generated, per-library data; one for hand-declared, per-cell data. Neither requires a Jacquard code change to onboard a new cell.